From bd49c93703e264bf855b81e4e805c0028de67251 Mon Sep 17 00:00:00 2001 From: Yangwook Kang Date: Fri, 22 May 2026 10:52:42 -0700 Subject: [PATCH] =?UTF-8?q?adr:=20add=20ADR-0050-0053=20=E2=80=94=20close?= =?UTF-8?q?=20/report's=20second-pass=20G4=20candidates?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Documents four cross-cutting surfaces one layer deeper than the prior G4 batch: - 0050 par-ccl-algorithm-module-contract: how to author a new CCL algorithm in src/kernbench/ccl/algorithms/. Pairs with ADR-0045's bench-module contract. Pins the four required public symbols (kernel, kernel_args, TOPO_NAME_TO_KIND constants, kernel alias), the 9 + tl standardized kernel signature, the kernel_args tuple format, sip_topo_kind dispatch, and the ccl.yaml entry workflow. - 0051 lat-routing-helper-api: every public method of AddressResolver (resolve, find_m_cpu, find_pcie_ep, find_io_cpu, find_all_pcie_eps) and PathRouter (find_path, find_path_with_distance, find_mcpu_dma_path, find_memory_path, find_node_path + 2 shims). Pins the four adjacency graphs (_adj_all / _adj / _adj_mcpu_dma / _adj_local) and the edge-kind exclusion sets they use, plus the single-owner naming convention. - 0052 dev-oplog-memory-store-schemas: OpRecord's 7 fields, the per-op_name params matrix (dma_read, dma_write, gemm_*, math, math reduction, composite_gemm, ipcq_copy, unknown), snapshot timing rules (math = all inputs, dma_write = HBM-only — ADR-0027 race avoidance), TileToken stage_type capture, and MemoryStore's (space, addr) two-level dict with reference-store semantics. - 0053 dev-topology-builder-algorithms: the 6-stage compile pipeline, cube_mesh.yaml's source_hash cache and its 5 input fields, the cube NoC auto-layout algorithm (row/col placement, HBM exclusion zone, PE/M_CPU/SRAM attachment via nearest-router, UCIe N/S/E/W distribution), the node naming convention (single-owner with router.py), the edge-kind catalog, the 4 view projections, and a table of spec-field changes vs mesh regeneration. Bilingual pair verifier passes for all four EN/KO pairs. Co-Authored-By: Claude Opus 4.7 (1M context) --- ...-0050-par-ccl-algorithm-module-contract.md | 308 +++++++++++++++ .../adr-ko/ADR-0051-lat-routing-helper-api.md | 267 +++++++++++++ ...ADR-0052-dev-oplog-memory-store-schemas.md | 352 +++++++++++++++++ ...DR-0053-dev-topology-builder-algorithms.md | 307 +++++++++++++++ ...-0050-par-ccl-algorithm-module-contract.md | 322 +++++++++++++++ docs/adr/ADR-0051-lat-routing-helper-api.md | 288 ++++++++++++++ ...ADR-0052-dev-oplog-memory-store-schemas.md | 371 ++++++++++++++++++ ...DR-0053-dev-topology-builder-algorithms.md | 351 +++++++++++++++++ 8 files changed, 2566 insertions(+) create mode 100644 docs/adr-ko/ADR-0050-par-ccl-algorithm-module-contract.md create mode 100644 docs/adr-ko/ADR-0051-lat-routing-helper-api.md create mode 100644 docs/adr-ko/ADR-0052-dev-oplog-memory-store-schemas.md create mode 100644 docs/adr-ko/ADR-0053-dev-topology-builder-algorithms.md create mode 100644 docs/adr/ADR-0050-par-ccl-algorithm-module-contract.md create mode 100644 docs/adr/ADR-0051-lat-routing-helper-api.md create mode 100644 docs/adr/ADR-0052-dev-oplog-memory-store-schemas.md create mode 100644 docs/adr/ADR-0053-dev-topology-builder-algorithms.md diff --git a/docs/adr-ko/ADR-0050-par-ccl-algorithm-module-contract.md b/docs/adr-ko/ADR-0050-par-ccl-algorithm-module-contract.md new file mode 100644 index 0000000..e6308df --- /dev/null +++ b/docs/adr-ko/ADR-0050-par-ccl-algorithm-module-contract.md @@ -0,0 +1,308 @@ +# ADR-0050: CCL Algorithm Module Contract — `ccl/algorithms/*.py` + +## Status + +Accepted (2026-05-22). + +`src/kernbench/ccl/algorithms/` 디렉터리 안의 모듈이 AHBM CCL backend +(ADR-0047) 에서 collective algorithm 으로 사용되려면 갖춰야 할 인터페이스, +kernel 시그너처, 그리고 새 알고리즘 추가 절차를 명시한다. ADR-0047 D3 가 +"algorithm 모듈은 `kernel`, `kernel_args`, optional `TOPO_NAME_TO_KIND` 를 +expose 해야 한다" 라고만 한 줄로 언급하나, **algorithm 모듈 작성자가 따라야 +할 contract** 는 ADR-level 에서 정리된 적이 없다. ADR-0045 가 bench 모듈 +contract 를 다루는 것과 짝을 이룬다. + +## First action (제일 처음에 하는 일) + +알고리즘 모듈이 import 되는 시점은 두 가지다: + +1. **AHBM backend 진입**: 사용자 코드가 `dist.init_process_group(backend="ahbm")` + 를 호출하면, `AhbmCCLBackend.__init__` 안에서 `self._algo_module = + importlib.import_module(self._merged["module"])` 가 실행된다. 이때 모듈 + 레벨에서 가장 먼저 일어나는 일: + - `SIP_TOPO_RING/TORUS/MESH` 같은 정수 상수가 모듈 namespace 에 노출. + - `TOPO_NAME_TO_KIND` 사전이 모듈 namespace 에 노출 — backend 가 + `topo_map = getattr(self._algo_module, "TOPO_NAME_TO_KIND", None)` 로 + 조회. + - `kernel_args` 함수 정의 — 호출 시 호출자가 사용. + - `allreduce_intercube_multidevice` 같은 알고리즘 함수 정의. + - 모듈 마지막 줄에서 `kernel = allreduce_intercube_multidevice` 로 + alias 가 노출. + +2. **ccl.yaml install 단계**: `kernbench.ccl.install.install_ipcq` 가 호출 + 되어 IPCQ neighbor table 을 푸시할 때 같은 알고리즘 모듈이 import 됨. + +즉, **algorithm 모듈의 첫 일은 "topology-kind 상수, `TOPO_NAME_TO_KIND` +사전, `kernel_args` 함수, 그리고 `kernel` alias 를 모듈 namespace 에 노출 +하는 것"** 이다. 모든 노출은 import-time 부수효과로 충분하며 별도 초기화 +함수 호출이 필요하지 않다. + +## Context + +`AhbmCCLBackend` (ADR-0047) 는 process group 초기화 시점에 `ccl.yaml` 의 +`defaults.algorithm` (또는 사용자가 지정한 알고리즘 이름) 으로부터 모듈 +경로를 얻어 dynamic import 한다. backend 는 그 모듈로부터 다음 4 가지를 +기대한다: + +- `kernel`: collective 의 진입 함수. +- `kernel_args(world_size, n_elem, cube_w=, cube_h=) -> tuple`: kernel 에 + 넘길 위치 인자 묶음. +- `TOPO_NAME_TO_KIND` (optional): `topology.yaml` 의 `sips.topology` + 문자열 (예: `"ring_1d"`, `"torus_2d"`, `"mesh_2d_no_wrap"`) 을 정수 + 상수로 매핑하는 dict. +- (간접) IPCQ neighbor table 설치: `configure_sfr_intercube_multisip` 가 + 알고리즘 모듈의 `TOPO_NAME_TO_KIND` 와 `cube_w/h` 를 보고 SFR 을 결정. + +현재 코퍼스의 유일한 algorithm 모듈은 `lrab_hierarchical_allreduce.py` +(248 줄) 이다. 이름은 "**l**eft-**r**ight **a**lternating **b**roadcast +**hierarchical allreduce**". 향후 `ring_allreduce`, `tree_allreduce`, +`broadcast` 같은 모듈이 추가될 때마다 이 contract 를 따라야 일관된 +디스패치가 가능하다. + +이 contract 가 ADR-level 에 없으면: + +- 새 algorithm 작성자가 ADR-0047 D3 의 한 줄 만으로 시그너처를 추론해야. +- kernel 함수 인자 순서 (특히 `t_ptr, n_elem, cube_w, cube_h, n_sips, + sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h, tl`) 의 의미가 코드 + grep 없이는 명확하지 않다. +- `kernel_args` 가 어떤 인자를 받고 어떤 tuple 을 돌려줘야 하는지 관례 + 로만 굳어진다. + +## Decision + +### D1. algorithm 모듈은 4 가지 public symbol 을 노출한다 + +```python +# src/kernbench/ccl/algorithms/.py + +from __future__ import annotations + +# (필수) topology-kind 상수 — 알고리즘 내부에서 사용 +SIP_TOPO_RING = 0 +SIP_TOPO_TORUS = 1 +SIP_TOPO_MESH = 2 + +# (선택) topology 이름 → kind 매핑. backend 가 ccl.yaml/topology 의 +# 문자열 SIP topology 를 정수로 변환하는 데 사용. +TOPO_NAME_TO_KIND = { + "ring_1d": SIP_TOPO_RING, + "torus_2d": SIP_TOPO_TORUS, + "mesh_2d_no_wrap": SIP_TOPO_MESH, +} + +# (필수) kernel 인자 빌더 +def kernel_args(world_size: int, n_elem: int, *, cube_w: int = 4, cube_h: int = 4) -> tuple: + return (n_elem, cube_w, cube_h, world_size) + +# (필수) kernel 함수 (`tl=...` 키워드를 통해 TLContext 가 주입됨) +def my_allreduce_kernel(t_ptr, n_elem, cube_w, cube_h, n_sips, + sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h, *, tl): + ... + +# (필수) kernel alias — backend 가 `module.kernel` 로 접근 +kernel = my_allreduce_kernel +``` + +- `kernel` alias 는 backend 가 직접 호출하는 entry point 다. 함수 이름이 + 무엇이든 (`allreduce_intercube_multidevice` 처럼) `module.kernel = fn` + 으로 노출해야 한다. +- `kernel_args` 가 없으면 backend 가 알고리즘 인자를 만들 방법이 없다. + signature 는 D2 참고. +- `TOPO_NAME_TO_KIND` 가 없으면 backend 는 `sip_topo_kind = 0` 으로 + fallback 한다. 단일 topology 만 지원하는 알고리즘이라면 생략 가능. + +### D2. `kernel_args` 시그너처 — `(world_size, n_elem, *, cube_w, cube_h)` + +```python +def kernel_args(world_size: int, n_elem: int, *, + cube_w: int = 4, cube_h: int = 4) -> tuple: + return (n_elem, cube_w, cube_h, world_size) +``` + +- **위치 인자**: `world_size` (= rank 수), `n_elem` (= 단일 shard 의 + element 수, f16 기준). +- **키워드 인자**: `cube_w`, `cube_h` (= cube mesh 크기). default 는 + 4×4 — `topology.yaml` 의 `sip.cube_mesh` 기본값과 정합. +- **반환**: kernel 의 위치 인자 순서대로 묶은 tuple. + +backend 의 `all_reduce` 가 호출 시: + +```python +kernel_args_tuple = self._algo_module.kernel_args( + self._world_size, n_elem, cube_w=eff_cube_w, cube_h=eff_cube_h, +) +extra_args = (sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h) +pending = self.ctx.launch( + self._merged["algorithm"], kernel_fn, tensor, + *kernel_args_tuple, *extra_args, _defer_wait=True, +) +``` + +즉 kernel 의 최종 위치 인자는: `(tensor_ptr, *kernel_args_tuple, +sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)` 이며, 거기에 `tl=...` 가 +키워드로 자동 주입된다. `kernel_args` 가 돌려주는 tuple 의 길이/순서는 +**kernel signature 와 1:1 일치** 해야 한다. + +### D3. `kernel` 함수 시그너처 — 정형화된 9 + tl 인자 + +권장 시그너처: + +```python +def my_kernel( + t_ptr: int, # VA base of the row-wise-sharded tensor on this SIP + n_elem: int, # element count per cube tile (or per shard) + cube_w: int, # cube mesh width (kernel_args 에서 옴) + cube_h: int, # cube mesh height (kernel_args 에서 옴) + n_sips: int, # world_size 와 동일 (rank = SIP, ADR-0024) + sip_rank: int, # 이 SIP 의 rank + sip_topo_kind: int, # TOPO_NAME_TO_KIND lookup 결과 + sip_topo_w: int, # SIP mesh width (ring_1d 면 0) + sip_topo_h: int, # SIP mesh height (ring_1d 면 0) + *, tl, # TLContext (auto-injected) +) -> None: +``` + +`kernel_args` 가 다른 위치 인자 순서를 채택하더라도, kernel 의 **마지막 +4 개 위치 인자는 항상 `(sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)`** +이며 backend 가 `extra_args` 로 append 한다 (ADR-0047 D5). 이 4 개 인자는 +사용자 정의 algorithm 도 받아야 하지만, 알고리즘이 single-SIP 이라면 +그냥 무시하면 된다. + +`tl` 은 위치 인자가 아닌 키워드로 주입된다 — `RuntimeContext.launch` 가 +kernel 호출 직전에 `tl=tl_ctx` 를 추가한다. 따라서 kernel signature 의 +`tl` 은 keyword-only (`*, tl`) 또는 마지막 키워드 매개변수 형태여야 +한다. + +### D4. kernel body 의 자유도와 제약 + +kernel body 안에서 사용 가능한 표면: ADR-0046 D3 의 모든 `tl.*` primitive. + +특히 자주 쓰이는 패턴: + +- `cube_id = tl.program_id(axis=1)` — 이 PE 가 속한 cube 인덱스. +- `pe_addr = t_ptr + cube_id * nbytes` — cube-별 tile 의 VA 계산. +- `acc = tl.load(pe_addr, shape=(n_elem,), dtype="f16")` — local 데이터 + 로드. +- `tl.send(dir=...)` / `tl.recv(dir=..., shape=, dtype=)` — IPCQ + collective. +- `acc = acc + recv` — TensorHandle 산술 연산자 (ADR-0046 D4). +- `tl.store(pe_addr, acc)` — 결과 저장. + +kernel body 는 일반 Python 함수이며, branching/looping 자유. 단: + +- SimPy `yield` 또는 `async` 금지 (ADR-0046 D1). +- TensorHandle 의 `.data` 직접 접근 금지 — phase 1 timing 모델은 + 데이터 의존을 모른다 (ADR-0020 의 2-pass 분리). +- kernel 실행은 deterministic 해야 한다 — 같은 입력으로 두 번 실행하면 + 같은 op 시퀀스 발사. random / external IO 금지. + +### D5. SIP topology semantics — `sip_topo_kind` 의 의미 + +backend 가 `topology.yaml` 의 `system.sips.topology` 문자열을 algorithm +모듈의 `TOPO_NAME_TO_KIND` 로 lookup 하여 `sip_topo_kind` 정수로 변환. +algorithm 은 이 정수를 보고 분기: + +```python +if sip_topo_kind == SIP_TOPO_RING: + acc = _inter_sip_ring(...) +elif sip_topo_kind == SIP_TOPO_TORUS: + acc = _inter_sip_torus_2d(...) +elif sip_topo_kind == SIP_TOPO_MESH: + acc = _inter_sip_mesh_2d(...) +``` + +각 topology branch 는 IPCQ direction 이름 (예: `"global_E"`, `"W"`, `"S"`, +`"N"`) 을 통해 peer 와 통신. direction 의 의미는 ADR-0023/0025 가 정의 +하며, `configure_sfr_intercube_multisip` 가 IPCQ neighbor table 을 그에 +맞춰 설치한다. + +algorithm 모듈은 자기가 지원하지 않는 topology kind 가 들어오면 silent +no-op 으로 두기보다 명시적으로 `raise ValueError(f"unsupported topology +kind {sip_topo_kind}")` 하는 것을 권장 — 실수로 backend 에 잘못 dispatch +된 경우 빠르게 fail. + +### D6. ccl.yaml 의 algorithm entry 구조 + +algorithm 모듈은 `ccl.yaml` 의 entry 와 짝을 이룬다 (ADR-0023 D10 + +ADR-0047 D3): + +```yaml +defaults: + algorithm: lrab_hierarchical_allreduce + n_elem: 8 + +algorithms: + lrab_hierarchical_allreduce: + module: kernbench.ccl.algorithms.lrab_hierarchical_allreduce + # optional: world_size override + # optional: per-algorithm parameters consumed by configure_sfr_intercube_multisip +``` + +- `module`: full Python module path. backend 의 `importlib.import_module` + 가 이 문자열을 그대로 사용. +- `world_size` (optional): 명시되면 topology fallback 을 override + (ADR-0047 D2). +- algorithm-specific parameters 는 `configure_sfr_intercube_multisip` 가 + 소비. + +새 algorithm 추가 시: + +1. `src/kernbench/ccl/algorithms/.py` 작성 (D1 컨벤션). +2. `ccl.yaml` 의 `algorithms` 섹션에 entry 추가. +3. (필요 시) `kernbench.ccl.sfr_config` 에 SFR 설치 분기 추가. +4. test 추가 (예: `tests/sccl/test_.py`, ADR-0043 의 eval harness + 확장). + +### D7. legacy "rank = flat PE index" 모드 + +ADR-0047 D2 가 명시한 `ccl.yaml` 의 `world_size` override 경로는 legacy +"rank = flat PE index" 테스트가 사용한다. algorithm 모듈은 이 모드 에서도 +`n_sips=world_size` 만큼의 rank 가 들어옴을 가정하면 된다 — backend 가 +rank↔(SIP, cube, PE) 매핑을 사전에 분리해 두므로 algorithm 본체에서는 +modal 분기가 필요 없다. + +단, single-cube workload 에서는 `cube_w=cube_h=1` 이 들어와 mesh-기반 +phase 들이 skip 되도록 작성해야 한다 (`lrab_hierarchical_allreduce.py` +의 `single_cube = (cube_w == 1 and cube_h == 1)` 패턴 참고). + +## Alternatives Considered + +### A1. algorithm 모듈을 class 로 구조화 (`class Allreduce: kernel(...)` 등) + +기각. Python 모듈 namespace 자체가 algorithm 의 identity 로 사용 중이며 +(ADR-0047 D3 의 `importlib.import_module`), class 한 겹은 추가 indirection +만 늘리고 dispatch 측 코드를 두텁게 만든다. 모듈-레벨 free function ++ `kernel` alias 패턴이 충분히 명확. + +### A2. `kernel_args` 를 명시적 dataclass 로 typing + +기각 (현재). algorithm 마다 인자 갯수가 다른 것이 정상이며, dataclass 한 +종류를 강제하면 다양한 algorithm 간 호환이 어려워진다. tuple 반환은 simple +하고 backend 측 `*kernel_args_tuple` unpacking 과 깨끗이 맞물린다. +algorithm 별 자체 타입 강도가 필요해지면 그 algorithm 모듈 안에서 NamedTuple +사용은 자유. + +### A3. SFR 설치를 algorithm 모듈 안으로 + +기각. SFR 설치 (`configure_sfr_intercube_multisip`) 는 topology + algorithm +모두를 보고 IPCQ neighbor table 을 설치하는 cross-module 결정이라, algorithm +모듈 내부보다 `kernbench.ccl.sfr_config` 같은 전용 위치가 자연스럽다. D6 의 +"필요 시 sfr_config 분기 추가" 워크플로우가 책임 분리 측면에서 더 명확. + +### A4. algorithm name 을 모듈 namespace 에 자동 등록 (decorator) + +기각. ADR-0045 (bench) 와 달리 algorithm 은 ccl.yaml entry 와 직접 묶여 +있어 추가 등록 레지스트리가 중복이다. `module` 문자열 매핑 하나면 충분. + +## Consequences + +- ADR-0047 D3 의 한 줄 contract 가 D1–D7 의 작성자-친화적 가이드로 확장 + 되어, 새 algorithm 추가 시 시그너처를 grep 으로 추론할 필요 없음. +- D3 의 9 + tl 인자 시그너처가 표준화되어, backend 의 `extra_args` append + (ADR-0047 D5) 와 자연스럽게 맞물림. 향후 single-SIP-only algorithm 도 + 4 개의 sip_* 인자를 받아야 함이 명시. +- D5 의 fail-loud 권장으로, ccl.yaml 의 topology 가 algorithm 미지원 + topology 로 잘못 설정되면 backend 가 silent wrong-result 가 아닌 + ValueError 로 fail. +- D6 의 단계별 추가 절차가 명시되어, 새 algorithm 추가가 sfr_config / + test / ccl.yaml 어디까지 손대야 하는지 분명. diff --git a/docs/adr-ko/ADR-0051-lat-routing-helper-api.md b/docs/adr-ko/ADR-0051-lat-routing-helper-api.md new file mode 100644 index 0000000..2b01721 --- /dev/null +++ b/docs/adr-ko/ADR-0051-lat-routing-helper-api.md @@ -0,0 +1,267 @@ +# ADR-0051: Routing Helper API — `AddressResolver` + `PathRouter` + +## Status + +Accepted (2026-05-22). + +`policy/routing/router.py` 가 노출하는 두 helper 클래스 +(`AddressResolver`, `PathRouter`) 의 모든 public API, 인자, 반환 값, +그리고 네 가지 다른 adjacency graph 의 사용처를 명시한다. ADR-0002 가 +routing distance 와 ordering, bypass 규칙을 정의하나, **helper API 표면 +자체** 는 ADR-level 에 정리된 적이 없다. + +## First action (제일 처음에 하는 일) + +### `AddressResolver(graph)` + +생성 즉시 다음 두 가지를 캐시한다: + +1. `self._node_ids = set(graph.nodes)` — 모든 node id 의 set (lookup 용). +2. `self._hbm_slice_bytes = hbm_total_gb * (1 << 30) // slices_per_cube` — + `graph.spec.cube.memory_map` 으로부터 산출 (기본 `48 GB / 8 slices = 6 + GB`). 이 값이 `resolve()` 가 HBM PA 의 `hbm_offset` 에서 `pe_id` 를 + 복원하는 데 쓰인다. + +즉, **AddressResolver 의 첫 일은 "전체 node id 집합과 HBM slice 크기를 +미리 계산해 두는 것"** 이다. graph 자체는 보유하지 않는다. + +### `PathRouter(graph)` + +생성 즉시 **네 개의 별도 adjacency graph 를 동시 구축**한다: + +1. `self._adj_all`: 모든 edge 포함 (component-to-component routing 용). +2. `self._adj`: `kind != "command"` 인 edge 만 (PE DMA / 일반 data path). +3. `self._adj_mcpu_dma`: `_MCPU_DMA_EXCLUDE = {"pe_internal", + "pe_to_router"}` 를 제외 (M_CPU DMA 가 PE pipeline 노드로 잘못 라우팅 + 되지 않게). +4. `self._adj_local`: `_UCIE_KINDS` 8 종을 제외 (cube-local routing 용 — + UCIe 가 zero-distance bus 처럼 보여 Dijkstra 가 mesh 보다 선호하는 + 것을 막음). + +각 그래프는 `defaultdict(list)` of `(neighbor, weight)` 형태이며, +`edge.routing_weight_mm or edge.distance_mm` 이 weight 로 쓰인다. + +즉, **PathRouter 의 첫 일은 "topology edge 들을 4개의 다른 정책으로 동시 +분류하여 4 개의 인접 리스트로 구축하는 것"**. 매 `find_*()` 호출 시 적절 +한 그래프를 골라 Dijkstra 를 돌린다. + +## Context + +`policy/routing/router.py` 는 다음 두 책임을 함께 수행한다: + +- **이름 매핑**: 토폴로지 명명 규칙 (`sip{S}.cube{C}.`, + `sip{S}.io{I}.pcie_ep` 등) 의 단일 소유자. 컴포넌트 / probe / IPCQ + install / runtime API 가 이름 문자열을 직접 만들지 않고 helper 를 호출. +- **경로 결정**: edge 의 `kind` 에 따른 정책 분리. 같은 src→dst 라도 + routing 의도 (PE DMA vs M_CPU DMA vs general component routing) 에 따라 + 다른 adjacency 를 사용해야 결과가 달라진다. + +이 helper API 가 코드 전반에서 광범위하게 소비되는데도 (probe.py / +distributed.py / install.py / 각종 component / tests), ADR-level 에서 +**정확한 시그너처 / 반환 의미 / 어떤 adjacency 를 쓰는지** 가 한 곳에 +정리되어 있지 않다. 본 ADR 이 그 빈자리를 채운다. + +## Decision + +### D1. `AddressResolver` 의 5 개 public API + +#### D1.1. `resolve(addr: PhysAddr) -> str` + +`PhysAddr` 인스턴스를 토폴로지의 destination node id 로 변환. + +``` +addr.kind == "hbm" → f"sip{s}.cube{d}.hbm_ctrl.pe{pe_id}" + where pe_id = addr.hbm_offset // self._hbm_slice_bytes (ADR-0017 D4/D9) + +addr.kind == "pe_resource": + addr.unit_type == PE → f"sip{s}.cube{d}.pe{addr.pe_id}.pe_tcm" + addr.unit_type == SRAM → f"sip{s}.cube{d}.sram" + addr.unit_type == MCPU → f"sip{s}.cube{d}.m_cpu" + 그 외 → RoutingError("unsupported unit_type") + +다른 kind → RoutingError("unsupported address kind") +``` + +산출된 node id 가 `self._node_ids` 에 없으면 `RoutingError(f"node {node_id} +not found in topology")`. 즉, address 의 syntax 가 valid 해도 topology 에 +실제로 매핑되는 노드가 없으면 fail-loud. + +#### D1.2. `find_m_cpu(sip, cube) -> str` + +`f"sip{sip}.cube{cube}.m_cpu"`. 없으면 `RoutingError`. + +#### D1.3. `find_pcie_ep(sip, io_id="io0") -> str` + +`f"sip{sip}.{io_id}.pcie_ep"`. 없으면 `RoutingError`. + +#### D1.4. `find_io_cpu(sip, io_id="io0") -> str` + +`f"sip{sip}.{io_id}.io_cpu"`. 없으면 `RoutingError`. + +#### D1.5. `find_all_pcie_eps() -> list[str]` + +전 SIP 의 PCIE_EP node id 를 정렬된 리스트로 반환. `endswith(".pcie_ep")` +필터링. cross-SIP IPCQ 가 모든 PCIE_EP 를 enumerate 할 때 사용. + +명명 규칙 (`sip{S}.cube{C}.`, `sip{S}.{io_id}.`) 의 단일 +소유자가 이 클래스다 (ADR-0015 D4). 토폴로지 빌더가 같은 명명 규칙으로 +노드를 만들고, 컴포넌트는 이름 문자열을 절대 직접 구성하지 않는다 — +모두 helper 를 거친다. + +### D2. `PathRouter` 의 4 개 adjacency graph + +생성자가 한 번에 구축. edge `kind` 가 정책을 결정: + +| graph | 제외 edge kinds | 용도 | +|-------------------|-----------------------------------------------|--------------------------------------------| +| `_adj_all` | (none) | M_CPU↔NOC command 포함, IO_CPU/M_CPU routing | +| `_adj` | `"command"` | PE DMA / 일반 data path | +| `_adj_mcpu_dma` | `"pe_internal"`, `"pe_to_router"` | M_CPU DMA (PE pipeline 우회) | +| `_adj_local` | `_UCIE_KINDS` (`ucie_internal`, `ucie_conn_to_router`, `router_to_ucie_conn`, `ucie_conn_to_noc`, `noc_to_ucie_conn`, `ucie_mesh`, `io_to_cube`, `cube_to_io`) | same-cube routing (UCIe bus 우회) | + +각 그래프는 `dict[node_id, list[(neighbor, weight)]]` 이며, weight 는 +`edge.routing_weight_mm or edge.distance_mm`. command edge 의 routing +영향력을 명시적으로 가르고, UCIe 의 "0-distance bus" 가 mesh 보다 선호 +되는 것을 막기 위한 `_adj_local` 분리가 ADR-0017 D7 의 cross-PE-slice +mesh-distance 요구와 정합. + +### D3. `PathRouter` 의 6 개 public API (+ 2 backward-compat) + +#### D3.1. `find_path(src_pe: str, dst_node: str) -> list[str]` + +**PE DMA routing**. `src_pe` 는 PE prefix (예: `"sip0.cube0.pe0"`) 이며, +함수가 `.pe_dma` 를 자동으로 prepend 하여 실제 시작 노드를 +`"sip0.cube0.pe0.pe_dma"` 로 설정. + +cube-local 여부 (`_same_cube`) 에 따라 adjacency 선택: + +- **same-cube** (src 와 dst 가 `sip{S}.cube{C}.` prefix 공유): + `_adj_local` 사용. UCIe 우회를 막아 cross-PE-slice 가 mesh 거리를 정확 + 히 지불 (ADR-0017 D7). +- **cross-cube**: `_adj` 사용. UCIe 가 자연스럽게 cross-cube path 의 + 최적 선택지로 포함됨. + +#### D3.2. `find_path_with_distance(src_pe, dst_node) -> tuple[list[str], float]` + +D3.1 과 동일한 adjacency 정책을 사용하나, 결과로 `(path, total_distance)` +를 함께 반환. probe / 분석 도구에서 distance 메트릭이 필요할 때 사용. + +#### D3.3. `find_mcpu_dma_path(m_cpu_id: str, dst_hbm_id: str) -> list[str]` + +**M_CPU DMA path**. cube 가 같으면 `_adj_local` (mesh 안에서 마무리), 다르 +면 `_adj_all` (UCIe 경유). `_MCPU_DMA_EXCLUDE` 가 PE pipeline 노드를 자동 +배제하므로, M_CPU 가 PE 의 내부 stage 를 거쳐 routing 되는 잘못된 경로가 +나오지 않는다. + +#### D3.4. `find_memory_path(src: str, dst: str) -> list[str]` + +`pcie_ep → io_noc → cube → router mesh → hbm_ctrl` 같은 직접 메모리 +경로. `_adj_mcpu_dma` 를 사용하여 `pe_internal` 및 `pe_to_router` edge +를 제외 — host-issued read/write 가 PE pipeline 으로 새지 않게 보장. +probe (ADR-0049 D1 의 H2D/D2H case) 에서 직접 호출. + +#### D3.5. `find_node_path(src: str, dst: str) -> list[str]` + +임의의 두 node 사이의 path. **command edge 포함** (`_adj_all` 사용). M_CPU +↔ NOC 같은 command-kind link 를 거쳐야 하는 IoCpuComponent / +MCpuComponent 등이 호출. + +#### D3.6. backward-compat shims + +- `_dijkstra(start, goal) -> list[str]` — `_run_dijkstra(self._adj, …)` + 의 thin wrapper. +- `_dijkstra_with_dist(start, goal) -> tuple[list[str], float]` — distance + 포함 버전. + +언더스코어 prefix 에서 보듯이 내부 API 인 척이지만 기존 테스트가 직접 +호출. 새 코드는 D3.1–D3.5 를 사용하고, 이 두 shim 은 deprecation 후보. + +### D4. Dijkstra 알고리즘 — single-source shortest path + +`_run_dijkstra_with_dist(adj, start, goal)`: + +- `heapq` priority queue. +- `best: dict[node, distance]` — 노드별 최단 거리 캐시. +- `prev: dict[node, predecessor]` — path reconstruction. +- weight 는 `routing_weight_mm or distance_mm`. UCIe 처럼 routing_weight 가 + 명시되어 distance 와 다른 edge 가 있으므로 weight 분리가 의도된 것. + +`start == goal` 은 빠른 path `([start], 0.0)` 반환. 도달 불가는 +`RoutingError(f"no path from {start} to {goal}")`. + +이 알고리즘은 **deterministic** 하다 — 같은 graph + start/goal 이면 같은 +경로. 이는 SPEC R1 의 "Routing MUST be deterministic" 요구와 정합. tie- +break 는 `heapq` 의 push 순서를 따른다 (Python list 순서가 deterministic). + +### D5. helper API 의 단일 소유자 원칙 + +다음 정보는 오직 router.py 안에서만 결정된다: + +- 명명 규칙: `sip{S}.cube{C}.`, `sip{S}.{io_id}.`, + `sip{S}.cube{C}.hbm_ctrl.pe{pe_id}`. +- adjacency 정책: 어떤 edge kind 가 어떤 그래프에 포함되는가. +- HBM slice 크기로부터 PE id 복원 방법. +- Dijkstra의 weight 결정 (`routing_weight_mm or distance_mm`). + +이 단일 소유자 원칙이 깨지면 (예: 컴포넌트가 자체적으로 `f"sip{s}..."` 를 +구성하기 시작하면) 명명 규칙 변경 시 영향 범위가 폭발한다. ADR-0015 D4 의 +정신과 정렬. + +### D6. helper API consumer 의 목록 + +본 helper 가 노출하는 메소드를 호출하는 곳을 명시 (현재 코퍼스 기준): + +- `probes/probe.py` (ADR-0049): `find_pcie_ep`, `find_io_cpu`, + `find_m_cpu`, `find_node_path`, `find_mcpu_dma_path`, + `find_memory_path`, `find_path`, `resolve`. +- `runtime_api/distributed.py` (ADR-0047): 간접 (engine 내부 routing). +- `ccl/install.py` (ADR-0023): `find_all_pcie_eps`, `resolve`. +- `sim_engine/event_log.py`: probe 와 유사하게 `find_pcie_ep`, + `find_memory_path`. +- `components/builtin/m_cpu.py`, `components/builtin/io_cpu.py`: + `find_node_path`, `find_mcpu_dma_path`. +- 각종 tests (test_routing.py, test_cross_sip_routing.py 등): D3.1–D3.5 + 대부분. + +새 consumer 가 추가될 때 본 ADR 의 D1/D3 가 그 의도에 맞는 메소드가 +이미 있는지 / 새 메소드를 추가해야 하는지 1차 판단의 기준이 된다. + +## Alternatives Considered + +### A1. 단일 adjacency graph + edge-kind filter 동적 적용 + +기각. 매 `find_*()` 마다 graph filtering 을 다시 하면 Dijkstra 의 cache +locality 와 성능이 떨어진다. 4 개 그래프 동시 구축 (D2) 은 메모리 비용 +이 작고 (edge ≤ 수만 건 규모), 호출 시점에 정책 선택이 O(1) 로 결정. + +### A2. adjacency 분리를 edge 의 `kind` 가 아닌 별도 metadata 로 + +기각. edge `kind` 는 이미 topology builder 가 부여하며 (ADR-0015 D4 + +ADR-0017), 별도 metadata 를 도입하면 두 시스템이 동기화되어야 하는 +중복이 생긴다. + +### A3. Dijkstra 대신 BFS + uniform weight + +기각. routing_weight_mm 이 edge 별로 다른 (mesh link / UCIe / IO-internal) +현실에서 BFS 는 hop 수 최소화일 뿐 latency / distance 최단을 보장하지 +않는다. SPEC R1 + R2 의 결정적·정확한 routing 요구에 어긋남. + +### A4. helper API 를 클래스 메서드가 아닌 모듈 함수로 + +기각. 두 클래스 (`AddressResolver`, `PathRouter`) 가 각각 cache 상태 +(`_node_ids`, `_hbm_slice_bytes`, 4 adjacency graphs) 를 보유해야 하며, +같은 graph 인스턴스에 여러 routing 질의가 발생한다. 모듈 함수는 매 호출 +시 state 를 다시 만들거나 global 로 두어야 해서 안전성/성능 저하. + +## Consequences + +- 컴포넌트 / probe / IPCQ install / runtime API 가 모두 router.py 의 + helper 만 호출하면 명명 규칙 변경 (예: `.io0.` → `.iochiplet0.`) 이 + 단 한 파일 수정으로 끝남 (D5). +- D2 의 4 그래프 분리가 ADR 에 굳어져, 새 edge kind 가 추가될 때 (예: + Inter-die UCIe link 의 새 kind) 어느 그래프에 포함시킬지 결정의 명확 + 한 기준 제공. +- D3.1 의 cube-local vs cross-cube 분기 (ADR-0017 D7) 가 명시되어, 향후 + routing 동작을 변경하려는 사람이 어느 adjacency 를 건드려야 할지 안다. +- D6 의 consumer 목록이 명시되어, helper API 변경 시 PR review 범위가 + 분명. backward-compat shim (D3.6) 의 deprecation 후보가 식별됨. diff --git a/docs/adr-ko/ADR-0052-dev-oplog-memory-store-schemas.md b/docs/adr-ko/ADR-0052-dev-oplog-memory-store-schemas.md new file mode 100644 index 0000000..23aced6 --- /dev/null +++ b/docs/adr-ko/ADR-0052-dev-oplog-memory-store-schemas.md @@ -0,0 +1,352 @@ +# ADR-0052: OpLog + MemoryStore Schemas — sim_engine internals + +## Status + +Accepted (2026-05-22). + +`sim_engine/op_log.py` 의 `OpRecord` 스키마와 `OpLogger` 의 record_start / +record_end / record_copy 동작, 그리고 `sim_engine/memory_store.py` 의 +`MemoryStore` 가 사용하는 (space, addr) 주소공간 namespace 와 read/write +의미를 명시한다. ADR-0020 (2-pass data execution) 가 두 인프라의 존재를 +선언하나, **레코드의 정확한 필드와 의미** 는 ADR-level 에서 정리되지 +않았고 ADR-0046 D3.2 (`tl.store` visibility), ADR-0023 D9 (IPCQ copy +record) 등 여러 ADR 이 이들의 동작에 의존하고 있다. + +## First action (제일 처음에 하는 일) + +### `OpLogger(memory_store=None)` + +생성 즉시 다음 3 가지 필드 초기화: + +1. `self._records: list[OpRecord] = []` — 누적된 op record. +2. `self._pending: dict[int, dict] = {}` — `id(msg)` 키로 partial record + (record_start 시점에 만들어졌고 record_end 가 아직 안 온 것). +3. `self._memory_store = memory_store` — 옵션 MemoryStore reference. + math op 의 input 스냅샷 + dma_write 의 HBM source 스냅샷 캡처에 사용. + +생성 시점에는 records / pending 모두 비어 있으며, `record_*` 호출이 +순차적으로 데이터를 누적한다. + +### `MemoryStore()` + +생성 즉시 `self._storage: dict[str, dict[int, np.ndarray]] = {}` 단 하나 +의 필드 초기화. 두 단계 dict (`space → addr → ndarray`) 이며 lazy 하게 +필요한 space 가 생길 때마다 inner dict 가 채워진다. + +즉, **두 인프라의 첫 일은 "비어 있는 누적 buffer + space-별 sparse dict +를 만들어 두는 것"** 이다. 첫 record / write 가 실제로 도착하면 그때 +필드가 채워지기 시작한다. + +## Context + +ADR-0020 (2-pass data execution) 의 D2/D5/D7 가 다음을 선언: + +- Phase 1 (timing) 동안 `ComponentBase._on_process_start/end` hook 이 + `OpLogger.record_start/end` 를 호출하여 모든 data op 의 시간 + 메타 + 데이터를 기록. +- Phase 2 (data) 가 op_log 를 t_start 순으로 재생하여 실 데이터 결과를 + 계산. +- 데이터 페이로드 자체는 `MemoryStore` 에 (space, addr) 키로 보관. + +ADR-0023 D9 (IPCQ atomic write), ADR-0027 (Megatron TP scratch +overwrite 회피), ADR-0046 D3.2 (`tl.store` visibility) 등 후속 ADR 들이 +op_log 와 MemoryStore 의 동작에 의존하지만, **정확한 record 필드 / space +이름 / 스냅샷 시점** 은 코드 grep 으로만 확인 가능하다. 본 ADR 이 이를 +정리한다. + +## Decision + +### D1. `OpRecord` 스키마 — 7 개 필드 + +```python +@dataclass +class OpRecord: + t_start: float + t_end: float + component_id: str + op_kind: str # "memory" | "gemm" | "math" | "unknown" + op_name: str # e.g. "dma_read", "gemm_f16", "exp", + # "TileToken/DMA_READ", "composite_gemm", + # "ipcq_copy" + params: dict[str, Any] + dependency_ids: list[int] = field(default_factory=list) +``` + +- **`t_start` / `t_end`**: SimPy 시간 (float ns). `t_start` 는 component + 가 op 를 시작한 시점, `t_end` 는 완료 시점. duration = `t_end - t_start`. +- **`component_id`**: op 가 발생한 node id (예: + `"sip0.cube0.pe0.pe_dma"`). +- **`op_kind`**: 4 가지 중 하나. Phase 2 DataExecutor 가 이 값으로 분기. +- **`op_name`**: 디버깅 / 분석용 사람-친화 이름. TileToken 일 경우 + `"TileToken/{stage_type}"` (예: `"TileToken/DMA_READ"`) 로 stage 를 + 구분. +- **`params`**: op-종속 메타데이터 dict (D3 참고). +- **`dependency_ids`**: 현재 사용되지 않음 (default `[]`). 향후 cross-op + dependency 추적이 필요해질 때를 위한 자리. + +### D2. `OpLogger.records` — t_start 정렬 보장 + +```python +@property +def records(self) -> list[OpRecord]: + self._records.sort(key=lambda r: r.t_start) + return self._records +``` + +매 접근 시 `t_start` 로 stable sort. 즉 같은 t_start 인 record 들은 insertion +순서를 유지. ADR-0020 D5 의 "t_start stable ordering" 요구와 정합. + +Phase 2 DataExecutor 는 항상 `records` property 를 통해 접근하므로, +record_end 호출이 t_start 와 다른 순서로 도착해도 (예: 짧은 op 가 긴 +op 보다 늦게 시작했으나 먼저 끝남) 재정렬되어 일관된 시퀀스를 받는다. + +### D3. op_name 별 `params` 스키마 (`_extract_op_info` 매핑) + +#### D3.1. `op_kind="memory", op_name="dma_read"` (DmaReadCmd) + +```python +{"src_addr": int, "nbytes": int, "handle_id": str} +``` + +#### D3.2. `op_kind="memory", op_name="dma_write"` (DmaWriteCmd) + +```python +{ + "src_space": str, # handle.space ("tcm"|"hbm"|"sram"), default "tcm" + "src_addr": int, # handle.addr + "shape": tuple, "dtype": str, + "dst_space": "hbm", # DmaWrite 는 항상 HBM 으로 + "dst_addr": int, + "nbytes": int, + "handle_id": str, + # record_end 시점에 src_space == "hbm" 이면 snapshot 추가 (D4) + "snapshot": np.ndarray | None, +} +``` + +#### D3.3. `op_kind="gemm", op_name=f"gemm_{dtype_a}"` (GemmCmd) + +```python +{ + "src_a_addr": int, "src_b_addr": int, "dst_addr": int, + "shape_a": tuple, "shape_b": tuple, "shape_out": tuple, + "dtype_in": str, "dtype_out": str, + "m": int, "k": int, "n": int, + # ADR-0027: per-operand + output spaces 보존 + "src_a_space": str, "src_b_space": str, "dst_space": str, +} +``` + +#### D3.4. `op_kind="math", op_name=msg.op` (MathCmd; op = "exp", "sum", "add", "where" 등) + +```python +{ + "input_addrs": list[int], # 입력 핸들들의 addr + "input_shapes": list[tuple], + "input_spaces": list[str], + "input_dtypes": list[str], + "dst_addr": int, "dst_space": str, + "shape_out": tuple, "dtype": str, + "axis": int | None, # reduction 인 경우만 의미 있음 + # record_end 시점에 모든 input 의 스냅샷이 채워짐 (D4) + "input_snapshots": list[np.ndarray | None], +} +``` + +#### D3.5. `op_kind="gemm" or "math", op_name=f"composite_{op}"` (CompositeCmd) + +```python +{ + "op": str, # "gemm" | "math" + "out_addr": int, "out_nbytes": int, + # op == "gemm" 인 경우 GemmCmd 와 같은 필드 추가: + "src_a_addr": int, "src_b_addr": int, + "shape_a": tuple, "shape_b": tuple, + "dtype_in": str, "dtype_out": str, + "src_a_space": str, "src_b_space": str, + "dst_space": "hbm", "dst_addr": int, # = out_addr +} +``` + +`op == "gemm"` 이면 `op_kind = "gemm"`, 아니면 `"math"`. Phase 2 측에서 +GemmCmd 와 동일 path 로 재생되도록 alias. + +#### D3.6. `op_kind="memory", op_name="ipcq_copy"` (record_copy 전용 경로) + +```python +{ + "src_space": str, "src_addr": int, + "dst_space": str, "dst_addr": int, + "shape": tuple, "dtype": str, "nbytes": int, + "snapshot": np.ndarray | None, # 호출자가 전달, 없으면 record_copy 가 fresh read +} +``` + +`PE_DMA._handle_ipcq_inbound` (ADR-0023 D9) 가 이 record 를 발사하여 IPCQ +slot 의 inbound copy 를 Phase 2 가 재생 가능하게 한다. 이 record 는 +`record_start` / `record_end` 를 거치지 않고 직접 `record_copy()` 로 push. + +#### D3.7. `op_kind="unknown", op_name=type(msg).__name__` + +`_extract_op_info` 가 인식 못 한 message 의 fallback. params = `{}`. +DataExecutor 가 이 op_kind 를 만나면 skip — Phase 2 replay 에 영향 없음. + +### D4. snapshot 캡처 시점 + +`OpLogger._memory_store` 가 set 되어 있을 때 record_end 가 다음을 수행: + +- **math op**: 모든 input addr/shape/space/dtype 으로 + `self._memory_store.read(...)` 를 호출하여 `params["input_snapshots"]` 에 + ndarray copy 첨부. read 실패 시 None. +- **dma_write op**: `src_space == "hbm"` 인 경우에만 source HBM 의 + 스냅샷을 `params["snapshot"]` 에 첨부. TCM source 는 **명시적으로 + 스킵** — TCM (PE scratch) 은 Phase 2 math/gemm 재생이 다시 채우므로, + Phase-1-time snapshot 을 잡으면 이전 kernel 의 stale 데이터를 잡을 위험 + (ADR-0027 postmortem: TP gemm → all_reduce race). +- **ipcq_copy**: `record_copy` 호출자가 `snapshot=token.data` 같이 in-flight + 스냅샷을 전달. 없으면 record_copy 가 fresh read 로 대체 시도. + +스냅샷은 `.copy()` 가 호출되어 (`ndarray.copy()` 가 fresh allocation) 이후 +storage mutation 으로부터 안전. ADR-0027 의 "cross-PE Phase 2 ordering" +race 회피의 근간. + +`memory_store` 가 None 인 경우 (Phase 1 timing-only 모드) 스냅샷 단계는 +전부 skip. record 의 timing 정보만 보존되며 데이터 replay 는 불가능. + +### D5. TileToken 처리 — record_start 가 stage 정보를 캡처 + +ADR-0014 D6 의 self-routing tile token (pipeline 모드) 은 stage_idx 가 +record_end 시점에 이미 advance 되어 있을 수 있다 (TileToken 이 다음 +component 로 이동하면서 next stage 의 params 를 캐시). 따라서: + +`record_start` 가 다음을 `pending[id(msg)]["snap"]` 에 미리 저장: + +```python +snap["stage_type"] = stage.stage_type.name # "DMA_READ", "GEMM", 등 +snap["stage_params"] = dict(stage.params) # 시점의 params 복사본 +``` + +`record_end` 에서 이 snap 을 꺼내 params 에 merge: + +- `params["stage_type"]` 가 final params 에 추가. +- `stage_params` 의 key 들이 (이미 있으면 보존) merge. +- `op_name == "TileToken"` 이면 `op_name = f"TileToken/{stage_type}"` 로 + rewrite (예: `"TileToken/DMA_READ"`) — 같은 component 에서 발생한 서로 + 다른 stage 의 record 를 disambiguate. + +이 메커니즘 덕분에 DMA_READ vs DMA_WRITE, FETCH vs STORE 가 같은 component +(예: pe_dma) 에서 발생하더라도 reporting 측에서 구분 가능. + +### D6. `MemoryStore` — (space, addr) 두 단계 dict + +```python +class MemoryStore: + def __init__(self) -> None: + self._storage: dict[str, dict[int, np.ndarray]] = {} + + def write(self, space, addr, data): self._storage[space][addr] = data + def read(self, space, addr, shape=None, dtype=None) -> np.ndarray: ... + def has(self, space, addr) -> bool: ... + def snapshot(self) -> MemoryStore: ... +``` + +#### D6.1. space namespace + +문자열 키. 표준 값: + +- `"hbm"`: HBM 데이터 (deploy_tensor + Phase 2 dma_write 결과). +- `"tcm"`: PE-로컬 TCM (Phase 2 math/gemm 결과). +- `"sram"`: cube-level SRAM (ADR-0023 D9.7 IPCQ slot tier). + +다른 space (예: `"reg"`) 도 자유롭게 허용 — `_storage` 가 lazy dict 라 +새 space 가 write 호출과 함께 자동 생성. + +#### D6.2. address keying + +`addr` 는 정수. **physical address (PA) 또는 virtual address (VA)** 일 수 +있다 — MemoryStore 자체는 address space 의 의미를 모르고 그저 키로 쓴다. +Phase 1 의 `MemoryWriteMsg` 는 PA + VA 둘 다 write (`_create_tensor` 에서 +PA 로 zero-init, VA base 로도 zero-init), Phase 2 는 op_log 가 captured +한 address 로 read/write. + +`addr` 의 의미는 호출자가 결정한다 — `MemoryStore` 는 lookup 만 제공. + +#### D6.3. read/write 의미 — reference store (no copy) + +`write(space, addr, data)`: `data` ndarray 의 reference 를 저장. **copy +하지 않음**. 호출자가 같은 ndarray 를 이후 mutate 하면 stored value 도 +변경된다. + +`read(space, addr, shape=None, dtype=None)`: 저장된 ndarray 의 reference +반환. `shape` 또는 `dtype` 이 제공되면: + +- `dtype != stored.dtype`: `arr.view(np_dtype)` 로 reinterpret cast (no + copy). +- `shape != stored.shape`: `nbytes` 가 일치하면 `arr.reshape(shape)` (view). +- `nbytes` 불일치: `ValueError`. + +데이터를 안전하게 분리하려면 호출자가 `arr.copy()` 호출. ADR-0027 의 +race 회피가 op_log snapshot 단계에서 명시적 copy 를 강제하는 이유. + +#### D6.4. `has(space, addr) -> bool` + +해당 키의 존재 여부만 확인. 데이터 인스턴스화는 안 함. + +#### D6.5. `snapshot() -> MemoryStore` + +shallow copy. inner dict 의 새 인스턴스를 만들되 ndarray reference 는 +공유. Phase 2 초기화 시점에 Phase 1 의 store 를 fork 하여 Phase 2 의 +mutation 이 Phase 1 의 다른 사용처에 영향을 주지 않게 분리하는 데 사용. + +### D7. op_log 가 SimPy 단일-스레드를 가정한다 + +`OpLogger` 의 `_records`, `_pending` 은 lock 없이 사용. SimPy 가 single- +threaded 라 `record_start` → `record_end` 사이에 다른 thread 가 끼어들 +수 없다는 가정. + +향후 multi-process kernbench (ADR-0047 D6) 가 도입되면 OpLogger 도 process +별로 분리되어야 함이 명시. 단일 OpLogger 인스턴스가 multiple process 의 +record 를 받지 못한다. + +## Alternatives Considered + +### A1. op_log 를 SQLite / parquet 같은 외부 store 로 + +기각 (현재). in-memory list 가 Phase 1 → Phase 2 의 핸드오프 latency 를 +최소화한다. 외부화는 long-running batch run 에서 의미가 있겠으나, 현재 +single-run 워크로드 에서는 overhead 만 추가. + +### A2. snapshot 을 record_start 시점에 캡처 + +기각. record_start 시점은 input 이 아직 채워지지 않은 상황 (예: math +op 의 input 이 직전 op 의 output 일 때) 이 흔하다. record_end 가 정확한 +시점. + +### A3. MemoryStore 를 component-별 store 로 분리 + +기각. (space, addr) 키가 이미 충분히 disambiguation 을 제공하며, component +별 분리는 cross-PE IPCQ copy (ADR-0023 D9) 가 source/destination 양쪽 +store 를 접근해야 하는 케이스를 복잡하게 만든다. + +### A4. op_log 에 cross-op dependency edge 명시 + +부분 채택. `dependency_ids` 필드가 OpRecord 에 자리 잡고 있지만 현재 +사용되지 않음 (D1). Phase 2 DataExecutor 가 t_start 정렬 + secondary sort +(memory ops before math at same t_start) 로 ordering 을 결정하며, 명시적 +dependency graph 가 필요해지면 이 필드가 채워질 자리. 현재는 ordering rule +이 충분하므로 미사용. + +## Consequences + +- ADR-0020 의 op_log / MemoryStore 선언이 D1–D6 의 구체 schema 로 확장 + 되어, Phase 2 DataExecutor 작성/수정 시 정확한 필드 의미를 grep 없이 + ADR 에서 확인 가능. +- D3 의 op_name 별 params 스키마가 명시되어, 새 op (예: 새 reduction + type) 추가 시 `_extract_op_info` 분기 어디에 끼울지 명확. +- D4 의 snapshot 시점 차이 (math = input snapshot, dma_write = HBM-only + snapshot) 가 ADR 에 굳어져, ADR-0027 의 cross-PE race 회피 결정이 향후 + refactor 에서 silently 깨지지 않음. +- D6.3 의 reference-store 의미가 명시되어, 호출자가 mutation safety 책임 + 을 인지. ADR-0027 의 explicit `.copy()` 패턴이 정당화됨. +- D7 의 single-thread 가정이 명시되어, multi-process kernbench (ADR-0047 + D6 supersession 후보) 도입 시 OpLogger 분리가 필요함이 분명. diff --git a/docs/adr-ko/ADR-0053-dev-topology-builder-algorithms.md b/docs/adr-ko/ADR-0053-dev-topology-builder-algorithms.md new file mode 100644 index 0000000..19d4389 --- /dev/null +++ b/docs/adr-ko/ADR-0053-dev-topology-builder-algorithms.md @@ -0,0 +1,307 @@ +# ADR-0053: Topology Builder + Visualizer Algorithms + +## Status + +Accepted (2026-05-22). + +`topology/builder.py`, `topology/mesh_gen.py`, `topology/visualizer.py` 가 +함께 수행하는 토폴로지 컴파일·시각화 파이프라인의 핵심 알고리즘 선택 +(placement-driven router attachment, mesh auto-layout, source_hash 캐시, +view projection, SVG rendering) 을 명시한다. ADR-0006 가 topology +compilation 의 high-level intent (compiled topology, distance extraction, +automatic diagram generation) 를 정의하나, **builder 가 실제로 어떤 +알고리즘을 사용하는지** 는 코드 grep 으로만 확인 가능했다. + +## First action (제일 처음에 하는 일) + +`resolve_topology(path_str)` 가 호출되면 다음 4 단계가 순서대로 일어난다: + +1. **경로 검증** (`builder.py::resolve_topology`): + `Path(path_str).expanduser().resolve()`, 존재 확인, file 여부 확인. + 실패 시 `FileNotFoundError` 또는 `ValueError`. +2. **YAML 파싱** (`_read_spec`): `yaml.safe_load`. parse error 면 line/ + column 정보 포함한 `ValueError`. dict 가 아니면 reject. +3. **mesh 자동 생성** (`mesh_gen.ensure_mesh_file`): topology yaml 과 + 같은 디렉터리에 `cube_mesh.yaml` 을 만들거나 (캐시 invalid 시) 재사용 + (캐시 hit 시). 이 단계가 cube NoC 의 라우터 grid 와 부착 정보를 결정. +4. **graph 컴파일** (`_compile_graph`): system → IO chiplets → cubes → + inter-cube edges → IO↔cube edges → system↔IO edges 순으로 nodes/edges + 를 누적, 그 다음 4 개의 view projection (system, sip, cube, pe) 을 + 생성하여 `TopologyGraph` 로 묶음. + +즉, **topology compile 의 첫 일은 "topology.yaml 을 dict 로 읽고, 동일 +디렉터리에 cube_mesh.yaml 을 생성/검증한 뒤, system→sip→cube→pe 순으로 +flat graph + 4-view projection 을 만드는 것"** 이다. + +## Context + +`topology/` 패키지의 책임: + +- **builder.py** (1207 줄): topology.yaml 을 받아 `TopologyGraph` (nodes + + edges + 4 view projections) 를 컴파일. +- **mesh_gen.py** (305 줄): cube NoC 의 라우터 grid 와 PE/UCIe/M_CPU/SRAM + 부착 위치를 자동 결정하여 `cube_mesh.yaml` 로 캐시. +- **visualizer.py** (887 줄): `TopologyGraph` 로부터 SVG 다이어그램 4종 + (system / sip / cube / pe) 을 생성. + +ADR-0006 가 "topology compilation 의 결과는 distance metadata 와 diagram +generation 의 single source" 라는 high-level 결정을 정의하나, 구체 알고리즘 +(예: placement-driven nearest-router attachment, HBM 제외 zone 산출, +source_hash 의 어떤 필드가 invalidation 을 트리거하는가) 은 ADR 에 없다. + +특히 다음 결정들이 ADR-level 에 부재: + +- 왜 mesh_gen 이 별도 파일 (`cube_mesh.yaml`) 로 캐시되는가? +- source_hash 가 어떤 필드를 포함하며, 어떤 변경이 재생성을 강제하는가? +- placement coordinate 가 cube 좌표가 아닌 mm 단위인 이유? +- HBM zone 제외와 UCIe N/S/E/W 분배가 mesh 안에서 어떻게 결정되는가? +- view projection 4 개 (system/sip/cube/pe) 의 추상화 레벨 차이? + +이 ADR 이 이 결정들을 한 곳에 정리한다. + +## Decision + +### D1. compile 파이프라인 — 6 단계 + +`_compile_graph(spec)`: + +1. **시스템 노드 생성** (`_instantiate_system`): `fabric.switch0`, host CPU + 등 system-level 노드 추가. +2. **per-SIP loop** (`for sip_id in range(system.sips.count)`): + - **IO chiplets** (`_instantiate_io_chiplets`): pcie_ep / io_cpu / + io_noc / io_ucie PHY / conn 노드 + 내부 양방향 edge 생성. + - **cube instantiation** (`_instantiate_cube`): cube_mesh.yaml 의 router + grid 를 토대로 cube-별 라우터, PE sub-components (pe_cpu, pe_dma, + pe_fetch_store, pe_gemm, pe_math, pe_mmu, pe_tcm, pe_scheduler, + pe_ipcq), m_cpu, sram, hbm_ctrl 인스턴스화 + 내부 edge 깔기. + - **inter-cube edges** (`_add_inter_cube_edges`): UCIe N/S/E/W mesh + edge. + - **IO ↔ cube edges** (`_add_io_to_cube_edges`): io_noc 와 cube 의 + edge UCIe phy 사이 연결. +3. **switch ↔ IO edges** (`_add_system_to_io_edges`): `fabric.switch0` + 와 각 SIP 의 `pcie_ep` 사이 양방향 edge (ADR-0038 D3 + ADR-0010 의 + cross-SIP IPCQ 경로). +4. **view projections** 4 종 build: + - `_build_system_view(spec)` — Tray 레벨, SIP 들과 system switch. + - `_build_sip_view(spec)` — SIP 안의 cube mesh + IO chiplet. + - `_build_cube_view(spec)` — 단일 cube 안의 router grid + PE/M_CPU/SRAM/ + HBM_CTRL 부착. + - `_build_pe_view(spec)` — 단일 PE 안의 9 sub-components + 내부 edge. +5. **TopologyGraph 리턴**: `TopologyGraph(spec, nodes, edges, system_view, + sip_view, cube_view, pe_view)`. + +이 6 단계는 **순서가 의미를 가진다**: cubes 가 만들어진 후에야 inter-cube +edges 가 valid 한 src/dst 를 갖고, IO chiplet 이 먼저 만들어져야 IO ↔ cube +edge 가 그를 참조할 수 있다. 새 노드 종류를 끼울 때는 의존 관계를 보고 +적절한 위치에 삽입해야 한다. + +### D2. `cube_mesh.yaml` — 별도 파일 + source_hash 캐시 + +`mesh_gen.ensure_mesh_file(cube_spec, mesh_path)`: + +1. `source_hash = _compute_source_hash(cube_spec)` 산출. 입력 필드: + - `geometry` (cube_mm.w/h 등). + - `pe_layout` (corners, pe_per_corner). + - `ucie.n_connections`. + - `memory_map.hbm_mapping_mode`. + - `placement` (m_cpu/sram pos_mm). +2. `mesh_path` (= `topology.yaml` 와 같은 디렉터리의 `cube_mesh.yaml`) 이 + 존재하고 `existing.source_hash == source_hash` 면 재사용 (캐시 hit). +3. 아니면 `_generate_mesh(cube_spec, source_hash)` 로 새 mesh 생성 후 + yaml 로 저장. + +별도 파일로 캐시하는 이유: + +- mesh 생성은 PE/UCIe/router 부착 계산이 들어가 매번 다시 하기 무거움. +- 같은 cube spec 으로 여러 번 실행 시 동일 mesh 가 보장되어야 함. +- 사람이 직접 mesh 를 inspect / debug 할 수 있는 artifact 가 됨. + +`source_hash` 가 list 한 5 개 필드가 mesh 형상을 결정하는 핵심이며, 그 +외 (예: bandwidth, overhead_ns) 변경은 mesh 재생성을 트리거하지 않는다. + +### D3. cube NoC mesh auto-layout 알고리즘 + +`_generate_mesh(cube_spec)`: + +#### D3.1. 행/열 결정 + +- `pe_positions = _corner_pe_positions(cube_w, cube_h)`: 4 corner (NW/NE/ + SW/SE) 마다 PE center 좌표 (mm). hardcoded `(1.5, 1.5)` / `(cube_w-1.5, + cube_h-1.5)` 패턴 + `pe_per_corner=2` 면 각 corner 에 2 PE 위치. +- `col_xs = _compute_col_positions(...)`: PE 들의 x 좌표 union + `max_spacing + = 3.0 mm` 보다 큰 gap 에 relay 컬럼 삽입. +- `row_ys, rows_per_half = _compute_row_positions(cube_h, n_connections, + pe_positions)`: + - `n_conn = max(n_connections, 2)` (hot path minimum). + - `rows_per_half = ceil(n_conn / 2)`. + - top 절반 + HBM 두 row + bottom 절반. HBM 은 `(cube_h/2 - 1.5, cube_h/2 + + 1.5)` 에 위치. PE rows 와 HBM rows 사이 `hbm_gap = 1.5 mm`. + +#### D3.2. HBM 제외 zone + +`hbm_row_start = rows_per_half`, `hbm_row_end = rows_per_half + 1`. +`hbm_col_start = n_cols // 2 - 1`, `hbm_col_end = n_cols // 2`. + +이 (row, col) 사각형 안의 router 슬롯은 `None` 으로 마킹 (라우터 없음). +실제 HBM 컨트롤러는 별도 `hbm_ctrl.pe{X}` 노드로 ADR-0017 D9 의 per-PE +파티션 패턴을 따라 부착. + +#### D3.3. PE 부착 + +각 corner 의 PE 들은 다음 row 에 매핑: + +- Top half: NW → row 0, NE → row 1 (top_corners 안의 index). +- Bottom half: SW → row `hbm_row_end + 1`, SE → row `hbm_row_end + 2`. + +각 PE 의 x 좌표가 가장 가까운 col 의 router 에 부착 (`min(range(n_cols), +key=lambda c: abs(col_xs[c] - pe_x))`). 부착 항목은 `pe{pe_idx}.dma`, +`pe{pe_idx}.cpu`, `pe{pe_idx}.hbm` 세 가지 (router 별 attach list 에 push). + +#### D3.4. M_CPU / SRAM 부착 — nearest router by Euclidean distance + +`placement.m_cpu.pos_mm` (default `[1.5, 5.5]`) 와 `placement.sram.pos_mm` +(default `[1.5, 8.5]`) 의 좌표에서 가장 가까운 router 를 Euclidean +distance 로 찾아 attach list 에 `"m_cpu"` / `"sram"` 추가. + +#### D3.5. UCIe N/S/E/W 분배 + +`ucie_pe_rows = top_pe_rows + bot_pe_rows` (총 `2 * rows_per_half` 개). + +- UCIe-E: 매 PE row 마다 rightmost col 의 router 에 `ucie_e.c{i}`. +- UCIe-W: leftmost col 의 router 에 `ucie_w.c{i}` (E 의 mirror). +- UCIe-N/S: PE column 들 중 절반을 좌측, 절반을 우측으로 나눠 top row / + bottom row 의 해당 col 에 부착. + +각 UCIe connection 은 `c{i}` index 가 붙어 ucie_n_connections 만큼의 PHY +가 분산된다 (ADR-0017 D5+). + +### D4. node 명명 규칙 — 단일 소유자 + +builder.py 는 다음 명명 규칙으로 노드를 만든다 (ADR-0051 D5 의 단일 +소유자 원칙): + +- `fabric.switch0` — system-level switch. +- `sip{S}.{io_id}.{pcie_ep|io_cpu|io_noc|io_ucie.{dir}|conn.{id}}` — IO + chiplet. +- `sip{S}.cube{C}.{m_cpu|sram|hbm_ctrl.pe{X}|noc.r{R}c{C}|...}` — cube 내부. +- `sip{S}.cube{C}.pe{P}.{pe_cpu|pe_dma|pe_fetch_store|pe_gemm|pe_math|pe_mmu|pe_tcm|pe_scheduler|pe_ipcq}` — PE sub-components. + +이 명명 규칙을 변경하려면 builder.py 와 router.py (ADR-0051) 의 helper +양쪽이 함께 갱신되어야 한다. 컴포넌트는 명명 규칙을 직접 알지 못하고 +helper 만 호출한다. + +### D5. edge `kind` 분류 + +각 edge 가 부여받는 `kind` 가 라우팅 정책 (ADR-0051 D2) 의 입력. 주요 +kind 값: + +- `"pe_internal"` — PE 내부 sub-component 간. +- `"pe_to_router"` — PE_DMA ↔ cube NoC router. +- `"router_mesh"` — cube NoC router 간. +- `"router_to_hbm"`, `"router_to_mcpu"`, `"router_to_sram"`, + `"sram_to_router"` 등 — cube-attached component 간. +- `"ucie_internal"`, `"ucie_conn_to_router"`, `"router_to_ucie_conn"`, + `"ucie_conn_to_noc"`, `"noc_to_ucie_conn"`, `"ucie_mesh"` — UCIe 관련. +- `"io_internal"` — IO chiplet 내부. +- `"io_to_cube"`, `"cube_to_io"` — IO ↔ cube 경계. +- `"pcie"` — switch ↔ pcie_ep. +- `"command"` — control-plane only edges (M_CPU ↔ NOC 등; PE DMA path 에서 + 제외). + +새 edge kind 를 추가하면 router.py 의 4 adjacency graph (ADR-0051 D2) 의 +어느 카테고리에 속할지 결정해야 한다 — 그렇지 않으면 default 로 `_adj_all` +에만 포함되어 의도와 다른 routing 발생 가능. + +### D6. view projection — 4 추상화 레벨 + +`TopologyGraph` 는 flat (nodes + edges) 외에 4 개의 view projection 을 +보유: + +- **system_view** (`_build_system_view`): Tray 레벨. SIP 박스들 + `fabric. + switch0`. PCIE 링크 표시. 외부 발표용 high-level overview. +- **sip_view** (`_build_sip_view`): 한 SIP 안. cube mesh + IO chiplet + (pcie_ep + io_cpu + io_noc). UCIe N/S/E/W 가 cube 간 연결로 보임. +- **cube_view** (`_build_cube_view`): 한 cube 안. router grid + PE/M_CPU/ + SRAM/HBM_CTRL 부착 + UCIe PHY edge 부분. cube 내부 라우팅 / placement + 진단용. +- **pe_view** (`_build_pe_view`): 한 PE 안. 9 sub-components + 내부 edge + (pe_internal kind). 자세한 PE 내부 dataflow 검토용. + +view 는 spec 에서 `visualization.emit_views: [system, sip, cube]` 같이 +선택적으로 출력 (ADR-0006). pe view 는 기본 출력에서 빠져 있으나 코드는 +유지 (자세한 디버그용). + +### D7. visualizer.py — SVG 다이어그램 출력 + +`emit_diagrams(graph, out_dir)` 가 모든 view 를 SVG 로 렌더. 핵심 함수: + +- `_render_view_svg(view)` — 일반적인 view 렌더 (router grid 가 없는 + 경우). +- `_render_cube_view_svg(view, spec)` — cube view 전용 (HBM block 그리기, + router grid layout, PE/M_CPU/SRAM/HBM positioning). +- `_draw_node`, `_draw_edge` — 노드 / edge 의 시각적 표현. +- `_pick_scale`, `_compute_node_sizes` — 자동 스케일링. + +visualizer 는 **derived artifact** (ADR-0006) 로 분류되며, 코드 변경 시 +production check 대상이 아니다. CLAUDE.md 의 "Derived Artifacts" 항목과 +정합. + +### D8. spec 변경의 영향 범위 + +| spec 필드 | 영향 | mesh 재생성 | +|---------------------------------------|-------------------|-------------| +| `system.sips.count` | SIP 갯수, node 수 | No | +| `sip.cube_mesh.w/h` | cube mesh 형상 | No | +| `cube.geometry.cube_mm.w/h` | cube 크기 (mm) | **Yes** | +| `cube.pe_layout.corners/pe_per_corner`| PE 부착 위치 | **Yes** | +| `cube.ucie.n_connections` | UCIe PHY 분배 | **Yes** | +| `cube.memory_map.hbm_mapping_mode` | HBM 분배 모드 | **Yes** | +| `cube.placement` | M_CPU/SRAM 위치 | **Yes** | +| `cube.memory_map.*` (위 제외) | HBM 용량 / BW | No | +| `*.links.*.bw_gbs` | edge bandwidth | No | +| `*.attrs.overhead_ns` | 컴포넌트 latency | No | + +위 표가 D2 의 `_compute_source_hash` 입력과 일치. mesh 재생성이 필요한 +변경은 `cube_mesh.yaml` 의 source_hash 가 자동 invalidate. + +## Alternatives Considered + +### A1. mesh 를 별도 캐시 파일 없이 매 compile 시 재생성 + +기각. 같은 spec 으로 여러 번 호출되는 케이스 (CLI run, probe, test) 마다 +mesh 생성 비용을 다시 지불. 또한 사람이 mesh 를 inspect 할 수 있는 artifact +가 사라짐. + +### A2. mesh 생성을 builder.py 에 합치기 + +기각 (현재). 305 줄 짜리 자체 알고리즘이며, mesh layout 의 결정 (placement- +driven router attachment, HBM exclusion zone) 이 builder 의 일반적인 +node/edge 생성 책임과 다르다. 분리 유지가 단일 책임 원칙에 더 부합. + +### A3. placement coordinate 를 cube 좌표 (col/row) 로 표현 + +기각. mm 단위 좌표가 시각화 측 (visualizer) 과 mesh layout 측 (nearest- +router 산출) 양쪽에서 일관되게 쓰인다. cube 좌표는 router grid 가 결정 +되기 전까지는 정의되지 않으므로 placement 입력에 부적절. + +### A4. view projection 을 lazy 하게 생성 + +기각 (현재). 4 개 view 의 생성 비용이 작고 (보통 < 100 ms), eager 생성이 +`TopologyGraph` 를 통한 single source of truth 를 보장. + +### A5. visualizer 출력 형식을 SVG 외 (PNG/PDF) 도 + +기각. SVG 가 vector + 텍스트 검색 가능 + 브라우저 직접 렌더가 가능한 가장 +유연한 형식. PNG 변환이 필요하면 별도 도구 (rsvg-convert 등) 로 후처리. + +## Consequences + +- ADR-0006 의 high-level intent 가 D1–D7 로 구체화되어, topology 변경 + 영향을 D8 표로 빠르게 가늠 가능. +- D3 의 mesh auto-layout 알고리즘이 ADR-level 에서 굳어져, 추후 새 PE + 부착 패턴 (예: HBM 의 6-zone 분할) 도입 시 어느 단계가 영향받는지 명확. +- D5 의 edge kind 목록과 D7 의 view 구조가 명시되어, 새 component 종류 + 추가 시 (builder + router + visualizer) 어디까지 손대야 하는지 PR + reviewer 가 한눈에 파악 가능. +- D2 의 source_hash invalidation 규칙이 명시되어, cube_mesh.yaml 이 stale + 하게 남는 경우 (예: bw 값만 바꿨을 때) 가 정상 동작임이 분명. diff --git a/docs/adr/ADR-0050-par-ccl-algorithm-module-contract.md b/docs/adr/ADR-0050-par-ccl-algorithm-module-contract.md new file mode 100644 index 0000000..b86a02e --- /dev/null +++ b/docs/adr/ADR-0050-par-ccl-algorithm-module-contract.md @@ -0,0 +1,322 @@ +# ADR-0050: CCL Algorithm Module Contract — `ccl/algorithms/*.py` + +## Status + +Accepted (2026-05-22). + +Pins down the interface, kernel signature, and addition workflow that a +module under `src/kernbench/ccl/algorithms/` must satisfy in order to be +used as a collective algorithm by the AHBM CCL backend (ADR-0047). +ADR-0047 D3 states only that "the algorithm module must expose `kernel`, +`kernel_args`, optionally `TOPO_NAME_TO_KIND`"; **the contract an +algorithm-module author needs to follow** has had no ADR-level coverage. +This ADR pairs with ADR-0045's bench-module contract. + +## First action + +An algorithm module is imported at two moments: + +1. **AHBM backend entry**: when user code calls + `dist.init_process_group(backend="ahbm")`, + `AhbmCCLBackend.__init__` runs + `self._algo_module = importlib.import_module(self._merged["module"])`. + At module level, the following occur first: + - Topology-kind integer constants like `SIP_TOPO_RING/TORUS/MESH` + are bound in the module namespace. + - The `TOPO_NAME_TO_KIND` dict is bound; the backend reads it via + `getattr(self._algo_module, "TOPO_NAME_TO_KIND", None)`. + - `kernel_args` function is defined for the caller. + - The actual algorithm function (e.g., + `allreduce_intercube_multidevice`) is defined. + - At the bottom of the module, `kernel = allreduce_intercube_multidevice` + publishes the alias. + +2. **ccl.yaml install stage**: + `kernbench.ccl.install.install_ipcq` imports the same algorithm + module while pushing the IPCQ neighbor table. + +In short, **the algorithm module's first act is "publish topology-kind +constants, the `TOPO_NAME_TO_KIND` dict, the `kernel_args` function, and +the `kernel` alias into the module namespace"** — all as import-time +side effects, no separate initialization call. + +## Context + +`AhbmCCLBackend` (ADR-0047), at process-group creation, dynamically +imports a module path obtained from `ccl.yaml`'s `defaults.algorithm` (or +a user-specified algorithm). The backend expects four things from the +module: + +- `kernel`: the collective's entry function. +- `kernel_args(world_size, n_elem, cube_w=, cube_h=) -> tuple`: a tuple + packing the kernel's positional arguments. +- `TOPO_NAME_TO_KIND` (optional): a dict mapping `topology.yaml`'s + `sips.topology` string (e.g., `"ring_1d"`, `"torus_2d"`, + `"mesh_2d_no_wrap"`) to the integer kind constants. +- (Indirectly) IPCQ neighbor-table install: + `configure_sfr_intercube_multisip` reads + the module's `TOPO_NAME_TO_KIND` plus cube dimensions to decide the + SFR. + +The current corpus has one algorithm module: +`lrab_hierarchical_allreduce.py` (248 lines). The name expands to +"**l**eft-**r**ight **a**lternating **b**roadcast hierarchical allreduce". +When future modules like `ring_allreduce`, `tree_allreduce`, or +`broadcast` are added, they must follow this contract for the backend's +dispatch path to keep working. + +Without an ADR-level contract: + +- A new algorithm author has to infer the signature from ADR-0047 D3's + one-liner. +- The kernel-function argument order (especially `t_ptr, n_elem, + cube_w, cube_h, n_sips, sip_rank, sip_topo_kind, sip_topo_w, + sip_topo_h, tl`) is unclear without grep. +- It is conventional, but not documented, what `kernel_args` takes as + inputs and what tuple it must return. + +## Decision + +### D1. The algorithm module exposes four public symbols + +```python +# src/kernbench/ccl/algorithms/.py + +from __future__ import annotations + +# (required) topology-kind constants — referenced internally +SIP_TOPO_RING = 0 +SIP_TOPO_TORUS = 1 +SIP_TOPO_MESH = 2 + +# (optional) topology name → kind mapping. Used by the backend to +# translate ccl.yaml/topology's string SIP topology into an integer. +TOPO_NAME_TO_KIND = { + "ring_1d": SIP_TOPO_RING, + "torus_2d": SIP_TOPO_TORUS, + "mesh_2d_no_wrap": SIP_TOPO_MESH, +} + +# (required) kernel argument builder +def kernel_args(world_size: int, n_elem: int, *, cube_w: int = 4, cube_h: int = 4) -> tuple: + return (n_elem, cube_w, cube_h, world_size) + +# (required) kernel function (TLContext is injected via the `tl=...` +# keyword argument). +def my_allreduce_kernel(t_ptr, n_elem, cube_w, cube_h, n_sips, + sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h, *, tl): + ... + +# (required) kernel alias — the backend accesses `module.kernel` +kernel = my_allreduce_kernel +``` + +- The `kernel` alias is the entry point the backend invokes. Whatever + the function name is (e.g., `allreduce_intercube_multidevice`), it + must be exposed via `module.kernel = fn`. +- Without `kernel_args`, the backend has no way to build the + algorithm's argument list. See D2 for the signature. +- If `TOPO_NAME_TO_KIND` is absent, the backend falls back to + `sip_topo_kind = 0`. An algorithm supporting only a single topology + may omit it. + +### D2. `kernel_args` signature — `(world_size, n_elem, *, cube_w, cube_h)` + +```python +def kernel_args(world_size: int, n_elem: int, *, + cube_w: int = 4, cube_h: int = 4) -> tuple: + return (n_elem, cube_w, cube_h, world_size) +``` + +- **Positional arguments**: `world_size` (= number of ranks), `n_elem` + (= element count of a single shard, f16-based). +- **Keyword arguments**: `cube_w`, `cube_h` (= cube-mesh dimensions). + Default 4×4 — aligned with `topology.yaml`'s `sip.cube_mesh` default. +- **Return**: a tuple in the order the kernel's positional arguments + expect. + +When the backend calls `all_reduce`: + +```python +kernel_args_tuple = self._algo_module.kernel_args( + self._world_size, n_elem, cube_w=eff_cube_w, cube_h=eff_cube_h, +) +extra_args = (sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h) +pending = self.ctx.launch( + self._merged["algorithm"], kernel_fn, tensor, + *kernel_args_tuple, *extra_args, _defer_wait=True, +) +``` + +So the kernel's full positional argument list becomes: `(tensor_ptr, +*kernel_args_tuple, sip_rank, sip_topo_kind, sip_topo_w, +sip_topo_h)`, with `tl=...` injected as a keyword. The tuple length +and order returned by `kernel_args` must **match the kernel signature +1:1**. + +### D3. Kernel signature — standardized 9 + tl arguments + +Recommended signature: + +```python +def my_kernel( + t_ptr: int, # VA base of the row-wise-sharded tensor on this SIP + n_elem: int, # element count per cube tile (or per shard) + cube_w: int, # cube mesh width (from kernel_args) + cube_h: int, # cube mesh height (from kernel_args) + n_sips: int, # equal to world_size (rank = SIP, ADR-0024) + sip_rank: int, # this SIP's rank + sip_topo_kind: int, # result of TOPO_NAME_TO_KIND lookup + sip_topo_w: int, # SIP mesh width (0 for ring_1d) + sip_topo_h: int, # SIP mesh height (0 for ring_1d) + *, tl, # TLContext (auto-injected) +) -> None: +``` + +Even if `kernel_args` chose a different positional argument order, the +kernel's **last four positional arguments are always +`(sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)`** — the backend +appends them as `extra_args` (ADR-0047 D5). A custom algorithm must +accept these four, but a single-SIP algorithm may simply ignore them. + +`tl` is injected via keyword — `RuntimeContext.launch` adds `tl=tl_ctx` +just before invoking the kernel. The signature therefore exposes `tl` +as keyword-only (`*, tl`) or as the trailing keyword parameter. + +### D4. Kernel body — freedom and constraints + +Available inside the kernel: every `tl.*` primitive from ADR-0046 D3. + +Common patterns: + +- `cube_id = tl.program_id(axis=1)` — this PE's cube index. +- `pe_addr = t_ptr + cube_id * nbytes` — per-cube VA of the tile. +- `acc = tl.load(pe_addr, shape=(n_elem,), dtype="f16")` — load local + data. +- `tl.send(dir=...)` / `tl.recv(dir=..., shape=, dtype=)` — IPCQ + collective. +- `acc = acc + recv` — TensorHandle arithmetic operators (ADR-0046 D4). +- `tl.store(pe_addr, acc)` — store the result. + +The kernel body is plain Python — branching and loops are fine. But: + +- No SimPy `yield` or `async` (ADR-0046 D1). +- No direct access to TensorHandle `.data` — the Phase 1 timing model + doesn't see data dependencies (ADR-0020's 2-pass separation). +- Kernel execution must be deterministic — the same input must produce + the same op sequence. No random or external IO. + +### D5. SIP topology semantics — meaning of `sip_topo_kind` + +The backend looks up `topology.yaml`'s `system.sips.topology` string +in the algorithm module's `TOPO_NAME_TO_KIND` and passes the integer +as `sip_topo_kind`. The algorithm then branches: + +```python +if sip_topo_kind == SIP_TOPO_RING: + acc = _inter_sip_ring(...) +elif sip_topo_kind == SIP_TOPO_TORUS: + acc = _inter_sip_torus_2d(...) +elif sip_topo_kind == SIP_TOPO_MESH: + acc = _inter_sip_mesh_2d(...) +``` + +Each topology branch communicates with peers via IPCQ direction names +(`"global_E"`, `"W"`, `"S"`, `"N"` …). Direction semantics are defined +in ADR-0023/0025; `configure_sfr_intercube_multisip` installs the IPCQ +neighbor table accordingly. + +If a topology kind not supported by the algorithm appears, prefer an +explicit `raise ValueError(f"unsupported topology kind +{sip_topo_kind}")` over a silent no-op — fail fast on misconfiguration. + +### D6. The `ccl.yaml` algorithm entry + +The algorithm module is paired with a `ccl.yaml` entry (ADR-0023 D10 + +ADR-0047 D3): + +```yaml +defaults: + algorithm: lrab_hierarchical_allreduce + n_elem: 8 + +algorithms: + lrab_hierarchical_allreduce: + module: kernbench.ccl.algorithms.lrab_hierarchical_allreduce + # optional: world_size override + # optional: per-algorithm parameters consumed by configure_sfr_intercube_multisip +``` + +- `module`: the full Python module path; `importlib.import_module` + consumes this string as-is. +- `world_size` (optional): when set, overrides the topology fallback + (ADR-0047 D2). +- Algorithm-specific parameters are consumed by + `configure_sfr_intercube_multisip`. + +Workflow to add a new algorithm: + +1. Write `src/kernbench/ccl/algorithms/.py` following D1. +2. Add the entry under `algorithms` in `ccl.yaml`. +3. (If needed) extend `kernbench.ccl.sfr_config` with the SFR-install + branch. +4. Add tests (e.g., `tests/sccl/test_.py`, extending the + ADR-0043 eval harness). + +### D7. Legacy "rank = flat PE index" mode + +The `world_size` override in `ccl.yaml`, surfaced by ADR-0047 D2, is +used by legacy "rank = flat PE index" tests. The algorithm module can +assume `n_sips=world_size` ranks even in this mode — the backend +maintains the rank↔(SIP, cube, PE) mapping, so no modal branching is +needed inside the algorithm body. + +In single-cube workloads (where `cube_w=cube_h=1`), the algorithm must +skip mesh-based phases — see the +`single_cube = (cube_w == 1 and cube_h == 1)` pattern in +`lrab_hierarchical_allreduce.py`. + +## Alternatives Considered + +### A1. Organize the algorithm module as a class (`class Allreduce: kernel(...)`) + +Rejected. The Python module namespace already identifies an algorithm +(see ADR-0047 D3's `importlib.import_module`). A class wrapper adds +indirection without simplifying dispatch. Module-level free functions +plus a `kernel` alias are clean and obvious. + +### A2. Type `kernel_args` with an explicit dataclass + +Rejected (currently). Each algorithm normally has a different argument +count; forcing one dataclass would hurt cross-algorithm interchange. +The tuple return is simple and unpacks cleanly with the backend's +`*kernel_args_tuple`. If an algorithm wants stronger internal typing, +it may define its own NamedTuple. + +### A3. Move SFR installation inside the algorithm module + +Rejected. SFR installation +(`configure_sfr_intercube_multisip`) is a cross-module decision +combining topology + algorithm; `kernbench.ccl.sfr_config` is a more +natural home than the algorithm module itself. D6's "extend +sfr_config if needed" workflow keeps responsibility boundaries clear. + +### A4. Auto-register algorithm names via a decorator (analogous to ADR-0045's `@bench`) + +Rejected. Unlike benches, algorithms are already tied to `ccl.yaml` +entries; an additional registry would be redundant. The string mapping +in `module` is sufficient. + +## Consequences + +- ADR-0047 D3's one-line contract expands to a D1–D7 author-facing + guide; new algorithm signatures no longer need to be grep-derived. +- D3's standardized 9 + tl signature couples naturally with the + backend's `extra_args` append (ADR-0047 D5). It is explicit that + even single-SIP-only algorithms must accept the four `sip_*` trailing + arguments. +- D5's fail-loud recommendation means a `ccl.yaml` topology that the + algorithm doesn't support will surface as an explicit `ValueError` + rather than a silent wrong result. +- D6's step-by-step addition workflow makes clear how far a new + algorithm has to reach into sfr_config / tests / ccl.yaml. diff --git a/docs/adr/ADR-0051-lat-routing-helper-api.md b/docs/adr/ADR-0051-lat-routing-helper-api.md new file mode 100644 index 0000000..004ee9c --- /dev/null +++ b/docs/adr/ADR-0051-lat-routing-helper-api.md @@ -0,0 +1,288 @@ +# ADR-0051: Routing Helper API — `AddressResolver` + `PathRouter` + +## Status + +Accepted (2026-05-22). + +Pins down every public API, argument, return value, and adjacency-graph +selection of the two helper classes (`AddressResolver`, `PathRouter`) +exposed by `policy/routing/router.py`. ADR-0002 defines routing +distance, ordering, and bypass rules, but **the helper API surface +itself** has had no ADR-level coverage. + +## First action + +### `AddressResolver(graph)` + +On construction, caches two pieces of state: + +1. `self._node_ids = set(graph.nodes)` — a set of all node ids for + lookup. +2. `self._hbm_slice_bytes = hbm_total_gb * (1 << 30) // slices_per_cube` + — derived from `graph.spec.cube.memory_map` (default `48 GB / 8 + slices = 6 GB`). `resolve()` uses this value to decode `pe_id` from + an HBM PA's `hbm_offset`. + +In short, **AddressResolver's first act is "precompute the full set of +node ids and the HBM slice size"**. It does not retain the graph +itself. + +### `PathRouter(graph)` + +On construction, **builds four separate adjacency graphs in one pass**: + +1. `self._adj_all`: every edge (used for component-to-component + routing). +2. `self._adj`: edges with `kind != "command"` (PE DMA / generic data + paths). +3. `self._adj_mcpu_dma`: excludes + `_MCPU_DMA_EXCLUDE = {"pe_internal", "pe_to_router"}` (M_CPU DMA + must not pass through PE pipeline nodes). +4. `self._adj_local`: excludes the 8-element `_UCIE_KINDS` set (UCIe + would look like a zero-distance bus to Dijkstra, which would prefer + it over the mesh — for cube-local routing this must be avoided). + +Each graph is a `defaultdict(list)` of `(neighbor, weight)`. The +weight is `edge.routing_weight_mm or edge.distance_mm`. + +In short, **PathRouter's first act is "classify topology edges into +four policy-specific adjacency lists simultaneously"**. Each `find_*()` +call picks the appropriate graph and runs Dijkstra. + +## Context + +`policy/routing/router.py` performs two responsibilities together: + +- **Naming**: it is the sole owner of the topology naming convention + (`sip{S}.cube{C}.`, `sip{S}.io{I}.pcie_ep`, etc.). Components / + probe / IPCQ install / runtime API do not build node-id strings + themselves — they call helpers. +- **Path decisions**: policy separation by `edge.kind`. For the same + src→dst, different routing intents (PE DMA vs M_CPU DMA vs general + component routing) call for different adjacencies and so produce + different paths. + +This helper API is widely consumed (probe.py / distributed.py / +install.py / various components / tests), yet **the exact signatures / +return semantics / adjacency picks** are not gathered in any ADR. This +ADR closes that gap. + +## Decision + +### D1. `AddressResolver` exposes five public methods + +#### D1.1. `resolve(addr: PhysAddr) -> str` + +Translates a `PhysAddr` to a destination node id in the topology: + +``` +addr.kind == "hbm" → f"sip{s}.cube{d}.hbm_ctrl.pe{pe_id}" + where pe_id = addr.hbm_offset // self._hbm_slice_bytes (ADR-0017 D4/D9) + +addr.kind == "pe_resource": + addr.unit_type == PE → f"sip{s}.cube{d}.pe{addr.pe_id}.pe_tcm" + addr.unit_type == SRAM → f"sip{s}.cube{d}.sram" + addr.unit_type == MCPU → f"sip{s}.cube{d}.m_cpu" + others → RoutingError("unsupported unit_type") + +other kinds → RoutingError("unsupported address kind") +``` + +If the derived node id is not in `self._node_ids`, raises +`RoutingError(f"node {node_id} not found in topology")`. So even when +the address has valid syntax, an absent node in the topology +fails-loud. + +#### D1.2. `find_m_cpu(sip, cube) -> str` + +Returns `f"sip{sip}.cube{cube}.m_cpu"`; absent → `RoutingError`. + +#### D1.3. `find_pcie_ep(sip, io_id="io0") -> str` + +Returns `f"sip{sip}.{io_id}.pcie_ep"`; absent → `RoutingError`. + +#### D1.4. `find_io_cpu(sip, io_id="io0") -> str` + +Returns `f"sip{sip}.{io_id}.io_cpu"`; absent → `RoutingError`. + +#### D1.5. `find_all_pcie_eps() -> list[str]` + +All PCIE_EP node ids across all SIPs, sorted. Filtered by +`endswith(".pcie_ep")`. Cross-SIP IPCQ uses this when enumerating +PCIE_EPs. + +This class is the sole owner of the naming convention +(`sip{S}.cube{C}.`, `sip{S}.{io_id}.`) — ADR-0015 D4. +The topology builder produces nodes with the same naming convention; +components never build node-id strings directly — they go through +these helpers. + +### D2. `PathRouter`'s four adjacency graphs + +Constructed in one pass. `edge.kind` drives policy: + +| graph | excluded edge kinds | use case | +|-------------------|--------------------------------------------------------------------------------------------------------------------------------------|-------------------------------------------------| +| `_adj_all` | (none) | M_CPU↔NOC command included, IO_CPU/M_CPU routes | +| `_adj` | `"command"` | PE DMA / generic data paths | +| `_adj_mcpu_dma` | `"pe_internal"`, `"pe_to_router"` | M_CPU DMA (skips PE pipeline) | +| `_adj_local` | `_UCIE_KINDS` (`ucie_internal`, `ucie_conn_to_router`, `router_to_ucie_conn`, `ucie_conn_to_noc`, `noc_to_ucie_conn`, `ucie_mesh`, `io_to_cube`, `cube_to_io`) | same-cube routing (UCIe bus excluded) | + +Each graph is `dict[node_id, list[(neighbor, weight)]]` with weight = +`edge.routing_weight_mm or edge.distance_mm`. Excluding command edges +prevents them from influencing routing; isolating `_adj_local` keeps +UCIe's "zero-distance bus" from out-competing the mesh — consistent +with ADR-0017 D7's cross-PE-slice mesh-distance requirement. + +### D3. `PathRouter` exposes six public methods (+ two backward-compat shims) + +#### D3.1. `find_path(src_pe: str, dst_node: str) -> list[str]` + +**PE DMA routing**. `src_pe` is a PE prefix (e.g., +`"sip0.cube0.pe0"`); the function auto-prepends `.pe_dma`, making the +true start node `"sip0.cube0.pe0.pe_dma"`. + +Adjacency depends on cube-locality (`_same_cube`): + +- **Same-cube** (src and dst share `sip{S}.cube{C}.` prefix): uses + `_adj_local`. Excluding UCIe lets cross-PE-slice access pay accurate + mesh distance (ADR-0017 D7). +- **Cross-cube**: uses `_adj`. UCIe naturally becomes the right choice + for the cross-cube portion. + +#### D3.2. `find_path_with_distance(src_pe, dst_node) -> tuple[list[str], float]` + +Same adjacency policy as D3.1, but returns `(path, total_distance)`. +Used by probe and analysis tools that need the distance metric. + +#### D3.3. `find_mcpu_dma_path(m_cpu_id: str, dst_hbm_id: str) -> list[str]` + +**M_CPU DMA path**. Same cube → `_adj_local` (stay within the mesh); +different cube → `_adj_all` (cross via UCIe). The +`_MCPU_DMA_EXCLUDE` set ensures PE-pipeline nodes never appear on +M_CPU's routes. + +#### D3.4. `find_memory_path(src: str, dst: str) -> list[str]` + +Direct memory path like +`pcie_ep → io_noc → cube → router mesh → hbm_ctrl`. Uses +`_adj_mcpu_dma` to exclude `pe_internal` and `pe_to_router`, so +host-issued reads/writes never leak into the PE pipeline. Probe +(ADR-0049 D1's H2D/D2H cases) calls this directly. + +#### D3.5. `find_node_path(src: str, dst: str) -> list[str]` + +Generic routing between arbitrary nodes, **including command edges** +(via `_adj_all`). IoCpuComponent / MCpuComponent use this when they +need to route through M_CPU ↔ NOC command-kind links. + +#### D3.6. Backward-compat shims + +- `_dijkstra(start, goal) -> list[str]` — thin wrapper for + `_run_dijkstra(self._adj, …)`. +- `_dijkstra_with_dist(start, goal) -> tuple[list[str], float]` — + distance-aware variant. + +Despite the underscore prefixes (suggesting internal API), existing +tests call these directly. New code should prefer D3.1–D3.5; these two +shims are deprecation candidates. + +### D4. Dijkstra — single-source shortest path + +`_run_dijkstra_with_dist(adj, start, goal)`: + +- `heapq` priority queue. +- `best: dict[node, distance]` — best known distance to each node. +- `prev: dict[node, predecessor]` — for path reconstruction. +- Edge weight = `routing_weight_mm or distance_mm`. The separation + matters because UCIe (and a few others) declare an explicit + `routing_weight_mm` distinct from physical `distance_mm`. + +`start == goal` short-circuits to `([start], 0.0)`. Unreachable target +→ `RoutingError(f"no path from {start} to {goal}")`. + +The algorithm is **deterministic**: identical graph + start/goal gives +the same path, satisfying SPEC R1 ("routing MUST be deterministic"). +Tie-breaks follow `heapq`'s push order (Python list order is +deterministic). + +### D5. Single-owner principle for helper-API decisions + +The following decisions live only inside router.py: + +- Naming convention: `sip{S}.cube{C}.`, + `sip{S}.{io_id}.`, + `sip{S}.cube{C}.hbm_ctrl.pe{pe_id}`. +- Adjacency policy: which edge kinds belong to which graph. +- Algorithm for recovering PE id from an HBM slice size. +- Dijkstra weight selection + (`routing_weight_mm or distance_mm`). + +Breaking single ownership (e.g., a component starting to build +`f"sip{s}..."` itself) would explode the blast radius of naming- +convention changes. This aligns with ADR-0015 D4. + +### D6. Consumers of the helper API + +Methods listed in this ADR are called from (current corpus): + +- `probes/probe.py` (ADR-0049): `find_pcie_ep`, `find_io_cpu`, + `find_m_cpu`, `find_node_path`, `find_mcpu_dma_path`, + `find_memory_path`, `find_path`, `resolve`. +- `runtime_api/distributed.py` (ADR-0047): indirectly (engine-internal + routing). +- `ccl/install.py` (ADR-0023): `find_all_pcie_eps`, `resolve`. +- `sim_engine/event_log.py`: like probe — `find_pcie_ep`, + `find_memory_path`. +- `components/builtin/m_cpu.py`, `components/builtin/io_cpu.py`: + `find_node_path`, `find_mcpu_dma_path`. +- Tests (test_routing.py, test_cross_sip_routing.py, …): most of + D3.1–D3.5. + +When a new consumer arrives, D1/D3 act as a first-pass guide on +whether an existing method matches the intent or a new one is needed. + +## Alternatives Considered + +### A1. One adjacency graph + per-call edge-kind filtering + +Rejected. Re-filtering the graph on every `find_*()` call hurts +Dijkstra cache locality. Constructing four graphs in one pass (D2) +has modest memory cost (edges ≤ a few × 10⁴), and selection happens +in O(1) at call time. + +### A2. Drive adjacency separation by separate edge metadata rather than `kind` + +Rejected. `edge.kind` is already assigned by the topology builder +(ADR-0015 D4 + ADR-0017); a parallel metadata field would force +synchronization between two systems. + +### A3. Use BFS with uniform weights instead of Dijkstra + +Rejected. With per-edge `routing_weight_mm` (mesh link / UCIe / +IO-internal), BFS minimizes hop count rather than total +latency/distance. SPEC R1 + R2 require deterministic and accurate +routing, which BFS does not deliver. + +### A4. Express the helper API as module functions instead of classes + +Rejected. Each class +(`AddressResolver`, `PathRouter`) maintains caches +(`_node_ids`, `_hbm_slice_bytes`, four adjacency graphs) reused across +many routing queries on the same graph. Module functions would have +to rebuild state per call or go global, hurting safety and +performance. + +## Consequences + +- When components / probe / IPCQ install / runtime API all go through + router.py helpers, a naming-convention change (e.g., `.io0.` → + `.iochiplet0.`) is a one-file edit (D5). +- D2's four-graph split is now ADR-locked, so when a new edge kind is + added (e.g., a new inter-die UCIe-link kind), the right adjacency + category is decided explicitly rather than by default. +- D3.1's same-cube vs cross-cube branching (ADR-0017 D7) is explicit, + so anyone changing routing knows which adjacency to touch. +- D6's consumer list bounds PR-review scope for helper-API changes, + and the backward-compat shims (D3.6) are flagged as deprecation + candidates. diff --git a/docs/adr/ADR-0052-dev-oplog-memory-store-schemas.md b/docs/adr/ADR-0052-dev-oplog-memory-store-schemas.md new file mode 100644 index 0000000..dd76def --- /dev/null +++ b/docs/adr/ADR-0052-dev-oplog-memory-store-schemas.md @@ -0,0 +1,371 @@ +# ADR-0052: OpLog + MemoryStore Schemas — sim_engine internals + +## Status + +Accepted (2026-05-22). + +Pins down the `OpRecord` schema and the `record_start` / `record_end` / +`record_copy` behavior in `sim_engine/op_log.py`, plus the +(space, addr) namespace and read/write semantics of `MemoryStore` in +`sim_engine/memory_store.py`. ADR-0020 (2-pass data execution) declares +that these two facilities exist, but **the precise record fields and +semantics** had no ADR-level coverage, and several recent ADRs +(ADR-0046 D3.2's `tl.store` visibility, ADR-0023 D9's IPCQ copy +record) depend on these semantics. + +## First action + +### `OpLogger(memory_store=None)` + +On construction, initialize three fields: + +1. `self._records: list[OpRecord] = []` — accumulated records. +2. `self._pending: dict[int, dict] = {}` — partial records keyed by + `id(msg)` (created at `record_start`, completed at `record_end`). +3. `self._memory_store = memory_store` — optional MemoryStore + reference. Used to capture math-op input snapshots and dma_write + HBM-source snapshots. + +Records and pending are empty; the `record_*` calls accumulate data +over time. + +### `MemoryStore()` + +On construction, initialize a single field: +`self._storage: dict[str, dict[int, np.ndarray]] = {}` — a two-level +dict (`space → addr → ndarray`). Inner dicts are created lazily as new +spaces appear. + +In short, **both facilities' first act is "set up an empty accumulator +buffer plus a sparse, per-space dict"**. The first record / write +fills the fields when it arrives. + +## Context + +ADR-0020 D2/D5/D7 (2-pass data execution) declares: + +- During Phase 1 (timing), `ComponentBase._on_process_start/end` hooks + call `OpLogger.record_start/end`, recording the time and metadata of + every data op. +- Phase 2 (data) replays the op log in `t_start` order to compute real + data. +- Data payloads live in `MemoryStore`, keyed by (space, addr). + +Subsequent ADRs (ADR-0023 D9's IPCQ atomic write, ADR-0027's Megatron +TP scratch-overwrite avoidance, ADR-0046 D3.2's `tl.store` visibility) +depend on op_log and MemoryStore behavior, but **the exact record +fields / space names / snapshot timing** are only discoverable via +source grep. This ADR codifies them. + +## Decision + +### D1. `OpRecord` schema — seven fields + +```python +@dataclass +class OpRecord: + t_start: float + t_end: float + component_id: str + op_kind: str # "memory" | "gemm" | "math" | "unknown" + op_name: str # e.g. "dma_read", "gemm_f16", "exp", + # "TileToken/DMA_READ", "composite_gemm", + # "ipcq_copy" + params: dict[str, Any] + dependency_ids: list[int] = field(default_factory=list) +``` + +- **`t_start` / `t_end`**: SimPy time (float ns). `t_start` is when the + component begins the op; `t_end` is completion. Duration = + `t_end - t_start`. +- **`component_id`**: the node id where the op occurred (e.g., + `"sip0.cube0.pe0.pe_dma"`). +- **`op_kind`**: one of four. Phase 2 DataExecutor branches on this. +- **`op_name`**: a debug/analysis-friendly name. For a TileToken, + expands to `"TileToken/{stage_type}"` (e.g., + `"TileToken/DMA_READ"`) to disambiguate stages. +- **`params`**: op-specific metadata dict (see D3). +- **`dependency_ids`**: currently unused (default `[]`). Reserved for + future cross-op dependency tracking. + +### D2. `OpLogger.records` — guaranteed `t_start` sort + +```python +@property +def records(self) -> list[OpRecord]: + self._records.sort(key=lambda r: r.t_start) + return self._records +``` + +A stable sort by `t_start` runs on each access. Records with the same +`t_start` preserve insertion order. Aligns with ADR-0020 D5's +"t_start stable ordering" requirement. + +Phase 2 DataExecutor always accesses via the `records` property, so +even when `record_end` calls arrive out of `t_start` order (e.g., a +short op started later but finished earlier), the sequence handed to +Phase 2 is consistent. + +### D3. `params` schema per `op_name` (matrix from `_extract_op_info`) + +#### D3.1. `op_kind="memory", op_name="dma_read"` (DmaReadCmd) + +```python +{"src_addr": int, "nbytes": int, "handle_id": str} +``` + +#### D3.2. `op_kind="memory", op_name="dma_write"` (DmaWriteCmd) + +```python +{ + "src_space": str, # handle.space ("tcm"|"hbm"|"sram"), default "tcm" + "src_addr": int, # handle.addr + "shape": tuple, "dtype": str, + "dst_space": "hbm", # DmaWrite always targets HBM + "dst_addr": int, + "nbytes": int, + "handle_id": str, + # When src_space == "hbm" at record_end, a snapshot is added (D4) + "snapshot": np.ndarray | None, +} +``` + +#### D3.3. `op_kind="gemm", op_name=f"gemm_{dtype_a}"` (GemmCmd) + +```python +{ + "src_a_addr": int, "src_b_addr": int, "dst_addr": int, + "shape_a": tuple, "shape_b": tuple, "shape_out": tuple, + "dtype_in": str, "dtype_out": str, + "m": int, "k": int, "n": int, + # ADR-0027: per-operand + output spaces preserved + "src_a_space": str, "src_b_space": str, "dst_space": str, +} +``` + +#### D3.4. `op_kind="math", op_name=msg.op` (MathCmd; op = "exp", "sum", "add", "where", …) + +```python +{ + "input_addrs": list[int], # addrs of input handles + "input_shapes": list[tuple], + "input_spaces": list[str], + "input_dtypes": list[str], + "dst_addr": int, "dst_space": str, + "shape_out": tuple, "dtype": str, + "axis": int | None, # only meaningful for reductions + # All inputs get snapshots at record_end (D4) + "input_snapshots": list[np.ndarray | None], +} +``` + +#### D3.5. `op_kind="gemm" or "math", op_name=f"composite_{op}"` (CompositeCmd) + +```python +{ + "op": str, # "gemm" | "math" + "out_addr": int, "out_nbytes": int, + # If op == "gemm", same fields as GemmCmd are added: + "src_a_addr": int, "src_b_addr": int, + "shape_a": tuple, "shape_b": tuple, + "dtype_in": str, "dtype_out": str, + "src_a_space": str, "src_b_space": str, + "dst_space": "hbm", "dst_addr": int, # = out_addr +} +``` + +If `op == "gemm"`, `op_kind = "gemm"`; otherwise `"math"`. An alias so +Phase 2 replays composite-gemm on the same path as `GemmCmd`. + +#### D3.6. `op_kind="memory", op_name="ipcq_copy"` (record_copy path) + +```python +{ + "src_space": str, "src_addr": int, + "dst_space": str, "dst_addr": int, + "shape": tuple, "dtype": str, "nbytes": int, + "snapshot": np.ndarray | None, # passed by caller; if None, record_copy reads fresh +} +``` + +`PE_DMA._handle_ipcq_inbound` (ADR-0023 D9) emits this record so Phase +2 can replay the IPCQ slot's inbound copy. It bypasses +`record_start` / `record_end` and pushes directly via `record_copy()`. + +#### D3.7. `op_kind="unknown", op_name=type(msg).__name__` + +Fallback for messages `_extract_op_info` doesn't recognize. `params = +{}`. If DataExecutor encounters this kind, it skips — Phase 2 replay +is unaffected. + +### D4. Snapshot capture timing + +When `OpLogger._memory_store` is set, `record_end` performs: + +- **Math op**: read every input + (addr/shape/space/dtype) from `self._memory_store.read(...)` and + attach an ndarray copy to `params["input_snapshots"]`. Read failure + → `None`. +- **`dma_write` op**: snapshot the source **only if `src_space == + "hbm"`** and attach to `params["snapshot"]`. TCM (PE scratch) + sources are **deliberately skipped** — TCM is repopulated by Phase 2 + math/gemm replay, and a Phase-1-time snapshot would capture a + previous kernel's stale value (ADR-0027 postmortem: TP gemm → + all_reduce race). +- **`ipcq_copy`**: the caller passes the in-flight snapshot via + `snapshot=token.data`. If absent, `record_copy` attempts a fresh + read from MemoryStore. + +Snapshots are taken with `.copy()` (fresh allocation), making them +safe against later storage mutation. This is the foundation of +ADR-0027's "cross-PE Phase 2 ordering" race-avoidance. + +When `memory_store` is `None` (Phase 1 timing-only mode), all +snapshot steps are skipped. Only the timing portion of the record is +preserved; data replay is unavailable. + +### D5. TileToken handling — `record_start` captures stage info + +ADR-0014 D6's self-routing tile token (pipeline mode) may have already +advanced its `stage_idx` by the time `record_end` runs (the TileToken +caches the next stage's params as it moves to the next component). +Therefore: + +`record_start` pre-saves the following in `pending[id(msg)]["snap"]`: + +```python +snap["stage_type"] = stage.stage_type.name # "DMA_READ", "GEMM", ... +snap["stage_params"] = dict(stage.params) # copy of params at start time +``` + +`record_end` retrieves this snap and merges into params: + +- Adds `params["stage_type"]` to final params. +- Merges `stage_params` keys (keeps existing values if any). +- If `op_name == "TileToken"`, rewrites it to + `f"TileToken/{stage_type}"` (e.g., `"TileToken/DMA_READ"`), + disambiguating different stages emitted by the same component. + +Thanks to this, DMA_READ vs DMA_WRITE, FETCH vs STORE coming from the +same component (e.g., pe_dma) are distinguishable in reports. + +### D6. `MemoryStore` — two-level (space, addr) dict + +```python +class MemoryStore: + def __init__(self) -> None: + self._storage: dict[str, dict[int, np.ndarray]] = {} + + def write(self, space, addr, data): self._storage[space][addr] = data + def read(self, space, addr, shape=None, dtype=None) -> np.ndarray: ... + def has(self, space, addr) -> bool: ... + def snapshot(self) -> MemoryStore: ... +``` + +#### D6.1. Space namespace + +A string key. Standard values: + +- `"hbm"`: HBM data (deploy_tensor + Phase 2 dma_write results). +- `"tcm"`: PE-local TCM (Phase 2 math/gemm output). +- `"sram"`: cube-level SRAM (ADR-0023 D9.7's IPCQ slot tier). + +Other spaces (e.g., `"reg"`) are allowed — `_storage` is a lazy dict +that creates a new space when `write` first touches it. + +#### D6.2. Address keying + +`addr` is an integer. It may be a **physical address (PA) or a virtual +address (VA)** — `MemoryStore` itself doesn't know address-space +semantics; it just uses them as keys. Phase 1's `MemoryWriteMsg` +writes both PA and VA +(`_create_tensor` zero-inits at PA and at the VA base too); Phase 2 +reads/writes via the addresses captured by op_log. + +The caller decides `addr`'s meaning — `MemoryStore` provides only +lookup. + +#### D6.3. read/write semantics — reference store (no copy) + +`write(space, addr, data)`: stores the ndarray reference. **No copy.** +If the caller later mutates the same ndarray, the stored value +changes. + +`read(space, addr, shape=None, dtype=None)`: returns the stored +ndarray reference. If `shape`/`dtype` are provided: + +- `dtype != stored.dtype`: `arr.view(np_dtype)` reinterprets as a + view (no copy). +- `shape != stored.shape`: if `nbytes` matches, `arr.reshape(shape)` + is a view. +- `nbytes` mismatch → `ValueError`. + +To detach the data, the caller must call `arr.copy()`. ADR-0027's +race-avoidance requires explicit `.copy()` in op_log snapshot steps +for exactly this reason. + +#### D6.4. `has(space, addr) -> bool` + +Existence check; does not materialize data. + +#### D6.5. `snapshot() -> MemoryStore` + +Shallow copy. Creates a new instance of inner dicts but shares +ndarray references. Used at Phase 2 init to fork from Phase 1's +store, so Phase 2 mutations don't affect Phase 1's remaining +consumers. + +### D7. op_log assumes a single-threaded SimPy + +`OpLogger`'s `_records` and `_pending` are lock-free. SimPy is +single-threaded, so nothing else can intrude between `record_start` +and `record_end` for the same message. + +When multi-process kernbench (ADR-0047 D6) arrives, OpLogger must be +split per process — one OpLogger instance cannot receive records from +multiple processes. + +## Alternatives Considered + +### A1. Externalize op_log to SQLite / parquet + +Rejected (currently). The in-memory list minimizes Phase 1 → Phase 2 +hand-off latency. Externalization makes sense for long-running batch +runs but adds overhead for the current single-run workload. + +### A2. Capture snapshots at `record_start` + +Rejected. At `record_start`, inputs are often not yet populated (e.g., +a math op's input is the output of a just-issued previous op). +`record_end` is the correct point. + +### A3. Per-component MemoryStore + +Rejected. The (space, addr) key already disambiguates effectively, and +splitting per component would complicate cross-PE IPCQ copy (ADR-0023 +D9), which needs access to both source and destination stores. + +### A4. Explicit dependency edges in op_log + +Partially adopted. The `dependency_ids` field exists on `OpRecord` but +is currently unused (D1). Phase 2 DataExecutor orders via `t_start` + +a secondary sort (memory ops before math at the same `t_start`). When +an explicit dependency graph is required, this field is the home. +Current ordering rules are sufficient, so it remains unused. + +## Consequences + +- ADR-0020's op_log / MemoryStore declarations are expanded into the + concrete D1–D6 schemas, so writing/modifying Phase 2 DataExecutor + doesn't need source-grep to learn field semantics. +- D3's per-`op_name` params matrix makes adding new ops (e.g., a new + reduction type) a question of branching in `_extract_op_info`. +- D4's per-op snapshot policy (math = input snapshot, dma_write = + HBM-only snapshot) is ADR-locked, so ADR-0027's race-avoidance + decision won't silently regress on future refactors. +- D6.3's reference-store semantics are explicit, putting mutation + safety on the caller. ADR-0027's explicit `.copy()` pattern is + justified. +- D7's single-thread assumption is recorded, so multi-process + kernbench (ADR-0047 D6's supersession candidate) will need OpLogger + separation when introduced. diff --git a/docs/adr/ADR-0053-dev-topology-builder-algorithms.md b/docs/adr/ADR-0053-dev-topology-builder-algorithms.md new file mode 100644 index 0000000..f889192 --- /dev/null +++ b/docs/adr/ADR-0053-dev-topology-builder-algorithms.md @@ -0,0 +1,351 @@ +# ADR-0053: Topology Builder + Visualizer Algorithms + +## Status + +Accepted (2026-05-22). + +Pins down the key algorithmic choices of the topology compile and +visualization pipeline jointly implemented by `topology/builder.py`, +`topology/mesh_gen.py`, and `topology/visualizer.py` — +placement-driven router attachment, mesh auto-layout, the source_hash +cache, view projections, and SVG rendering. ADR-0006 defines the +high-level intent of topology compilation (compiled topology, distance +extraction, automatic diagram generation), but **which algorithms the +builder actually uses** was only discoverable via source grep. + +## First action + +When `resolve_topology(path_str)` is called, four steps run in order: + +1. **Path validation** (`builder.py::resolve_topology`): + `Path(path_str).expanduser().resolve()`, existence check, file + check. Failure → `FileNotFoundError` or `ValueError`. +2. **YAML parsing** (`_read_spec`): `yaml.safe_load`. Parse errors + yield a `ValueError` with line/column. Non-dict roots are + rejected. +3. **Auto-generate the mesh** (`mesh_gen.ensure_mesh_file`): create or + reuse a `cube_mesh.yaml` next to the topology file. Cache hit on + matching source_hash; miss triggers regeneration. This step decides + the cube NoC's router grid and attachment information. +4. **Compile the graph** (`_compile_graph`): system → IO chiplets → + cubes → inter-cube edges → IO↔cube edges → system↔IO edges, then + build four view projections (system, sip, cube, pe) and wrap into + a `TopologyGraph`. + +In short, **topology compilation's first act is "read topology.yaml as +a dict, create/validate cube_mesh.yaml in the same directory, then +build the flat graph + 4-view projection in system → sip → cube → pe +order"**. + +## Context + +`topology/` package responsibilities: + +- **builder.py** (1207 lines): turns topology.yaml into a + `TopologyGraph` (nodes + edges + 4 view projections). +- **mesh_gen.py** (305 lines): auto-decides the cube NoC's router + grid and PE/UCIe/M_CPU/SRAM attachment positions and caches them in + `cube_mesh.yaml`. +- **visualizer.py** (887 lines): generates four SVG diagrams (system / + sip / cube / pe) from a `TopologyGraph`. + +ADR-0006 makes the high-level decision that "the result of topology +compilation is the single source for distance metadata and diagram +generation", but specific algorithms (e.g., placement-driven nearest- +router attachment, the HBM exclusion zone, which fields in source_hash +trigger regeneration) are not in any ADR. + +In particular, these decisions are absent at ADR level: + +- Why is mesh_gen cached in a separate file (`cube_mesh.yaml`)? +- Which fields are in source_hash, and which changes force + regeneration? +- Why placement coordinates in mm rather than cube coordinates? +- How are the HBM exclusion zone and UCIe N/S/E/W distribution + decided inside the mesh? +- What is the abstraction-level difference among the four view + projections (system/sip/cube/pe)? + +This ADR captures these decisions in one place. + +## Decision + +### D1. Compile pipeline — six stages + +`_compile_graph(spec)`: + +1. **System nodes** (`_instantiate_system`): add system-level nodes + like `fabric.switch0` and the host CPU. +2. **Per-SIP loop** (`for sip_id in range(system.sips.count)`): + - **IO chiplets** (`_instantiate_io_chiplets`): create pcie_ep / + io_cpu / io_noc / io_ucie PHYs / conn nodes and their bidirectional + internal edges. + - **Cube instantiation** (`_instantiate_cube`): using + cube_mesh.yaml's router grid, instantiate cube routers, PE + sub-components (pe_cpu, pe_dma, pe_fetch_store, pe_gemm, pe_math, + pe_mmu, pe_tcm, pe_scheduler, pe_ipcq), m_cpu, sram, hbm_ctrl, + and their internal edges. + - **Inter-cube edges** (`_add_inter_cube_edges`): the UCIe + N/S/E/W mesh edges. + - **IO ↔ cube edges** (`_add_io_to_cube_edges`): connect io_noc to + each cube's edge UCIe phy. +3. **Switch ↔ IO edges** (`_add_system_to_io_edges`): bidirectional + edges between `fabric.switch0` and each SIP's `pcie_ep` (the + cross-SIP IPCQ path of ADR-0038 D3 + ADR-0010). +4. **Build four view projections**: + - `_build_system_view(spec)` — Tray level: SIPs and the system + switch. + - `_build_sip_view(spec)` — inside one SIP: cube mesh + IO + chiplet. + - `_build_cube_view(spec)` — inside one cube: router grid + PE / + M_CPU / SRAM / HBM_CTRL attachments. + - `_build_pe_view(spec)` — inside one PE: nine sub-components + + internal edges (pe_internal kind). +5. **Return `TopologyGraph`**: `TopologyGraph(spec, nodes, edges, + system_view, sip_view, cube_view, pe_view)`. + +The six stages are **ordered for a reason**: only after cubes exist +do inter-cube edges have valid src/dst, and IO chiplets must precede +the IO ↔ cube edges that reference them. New node types must slot in +the right spot. + +### D2. `cube_mesh.yaml` — a separate file with a source_hash cache + +`mesh_gen.ensure_mesh_file(cube_spec, mesh_path)`: + +1. Compute `source_hash = _compute_source_hash(cube_spec)` from these + input fields: + - `geometry` (cube_mm.w/h …). + - `pe_layout` (corners, pe_per_corner). + - `ucie.n_connections`. + - `memory_map.hbm_mapping_mode`. + - `placement` (m_cpu/sram pos_mm). +2. If `mesh_path` (= `cube_mesh.yaml` next to topology.yaml) exists + and `existing.source_hash == source_hash`, reuse it (cache hit). +3. Otherwise, generate a new mesh via + `_generate_mesh(cube_spec, source_hash)` and write to yaml. + +Caching as a separate file because: + +- Mesh generation involves nontrivial PE/UCIe/router attachment math + and is too expensive to redo every time. +- Multiple runs with the same cube spec must guarantee an identical + mesh. +- The resulting mesh is itself an inspectable / debuggable artifact. + +The five fields listed in source_hash are the ones that determine +mesh shape; other changes (e.g., bandwidth, overhead_ns) do not +trigger mesh regeneration. + +### D3. Cube NoC mesh auto-layout + +`_generate_mesh(cube_spec)`: + +#### D3.1. Rows / columns + +- `pe_positions = _corner_pe_positions(cube_w, cube_h)`: PE-center + coordinates (mm) per corner (NW/NE/SW/SE). Hardcoded patterns like + `(1.5, 1.5)` and `(cube_w-1.5, cube_h-1.5)`; with `pe_per_corner=2`, + each corner has two PE positions. +- `col_xs = _compute_col_positions(...)`: union of PE x-coordinates, + plus relay columns inserted when any gap exceeds + `max_spacing = 3.0 mm`. +- `row_ys, rows_per_half = _compute_row_positions(cube_h, + n_connections, pe_positions)`: + - `n_conn = max(n_connections, 2)` (hot-path minimum). + - `rows_per_half = ceil(n_conn / 2)`. + - Top half + two HBM rows + bottom half. HBM sits at + `(cube_h/2 - 1.5, cube_h/2 + 1.5)`. The gap between PE rows and + HBM rows is `hbm_gap = 1.5 mm`. + +#### D3.2. HBM exclusion zone + +`hbm_row_start = rows_per_half`, +`hbm_row_end = rows_per_half + 1`. +`hbm_col_start = n_cols // 2 - 1`, +`hbm_col_end = n_cols // 2`. + +Router slots inside this (row, col) rectangle are marked `None` (no +router). HBM controllers are added separately as +`hbm_ctrl.pe{X}` nodes following ADR-0017 D9's per-PE partition +pattern. + +#### D3.3. PE attachment + +Each corner's PEs map to a row: + +- Top half: NW → row 0, NE → row 1 (top_corners index). +- Bottom half: SW → row `hbm_row_end + 1`, SE → row + `hbm_row_end + 2`. + +Each PE's x-coordinate attaches to the nearest column's router +(`min(range(n_cols), key=lambda c: abs(col_xs[c] - pe_x))`). +Attachment items are `pe{pe_idx}.dma`, `pe{pe_idx}.cpu`, +`pe{pe_idx}.hbm` (pushed into the router's attach list). + +#### D3.4. M_CPU / SRAM attachment — nearest router by Euclidean distance + +For `placement.m_cpu.pos_mm` (default `[1.5, 5.5]`) and +`placement.sram.pos_mm` (default `[1.5, 8.5]`), find the router with +the smallest Euclidean distance and append `"m_cpu"` / `"sram"` to +its attach list. + +#### D3.5. UCIe N/S/E/W distribution + +`ucie_pe_rows = top_pe_rows + bot_pe_rows` (total +`2 * rows_per_half`). + +- UCIe-E: one PE row at a time, attach `ucie_e.c{i}` to the rightmost + column's router. +- UCIe-W: attach `ucie_w.c{i}` to the leftmost column's router (E's + mirror). +- UCIe-N/S: split PE columns into left and right halves; attach to + the top row's / bottom row's matching columns. + +Each UCIe connection is suffixed `c{i}`, distributing +ucie_n_connections PHYs (ADR-0017 D5+). + +### D4. Node naming convention — single ownership + +builder.py creates nodes with the following naming convention (the +single-owner principle from ADR-0051 D5): + +- `fabric.switch0` — system-level switch. +- `sip{S}.{io_id}.{pcie_ep|io_cpu|io_noc|io_ucie.{dir}|conn.{id}}` — + IO chiplet. +- `sip{S}.cube{C}.{m_cpu|sram|hbm_ctrl.pe{X}|noc.r{R}c{C}|...}` — + inside cube. +- `sip{S}.cube{C}.pe{P}.{pe_cpu|pe_dma|pe_fetch_store|pe_gemm|pe_math|pe_mmu|pe_tcm|pe_scheduler|pe_ipcq}` — + PE sub-components. + +Changing this convention requires updating both builder.py and +router.py's helpers (ADR-0051). Components never know the convention +directly — they only call the helpers. + +### D5. Edge `kind` classification + +Every edge gets a `kind`; routing policy (ADR-0051 D2) reads it. Major +kinds: + +- `"pe_internal"` — within a PE between sub-components. +- `"pe_to_router"` — PE_DMA ↔ cube NoC router. +- `"router_mesh"` — between cube NoC routers. +- `"router_to_hbm"`, `"router_to_mcpu"`, `"router_to_sram"`, + `"sram_to_router"`, etc. — between cube-attached components. +- `"ucie_internal"`, `"ucie_conn_to_router"`, + `"router_to_ucie_conn"`, `"ucie_conn_to_noc"`, + `"noc_to_ucie_conn"`, `"ucie_mesh"` — UCIe-related. +- `"io_internal"` — inside IO chiplet. +- `"io_to_cube"`, `"cube_to_io"` — at the IO ↔ cube boundary. +- `"pcie"` — switch ↔ pcie_ep. +- `"command"` — control-plane edges only (e.g., M_CPU ↔ NOC; excluded + from PE DMA paths). + +Adding a new edge kind requires picking a category in router.py's +four adjacency graphs (ADR-0051 D2). If you forget, it defaults to +`_adj_all` only, which can produce unintended routes. + +### D6. View projection — four abstraction levels + +`TopologyGraph` keeps four view projections alongside the flat +nodes+edges: + +- **system_view** (`_build_system_view`): Tray level. SIP blocks and + `fabric.switch0`. PCIe links shown. For external high-level + overview. +- **sip_view** (`_build_sip_view`): inside one SIP — cube mesh + IO + chiplet (pcie_ep + io_cpu + io_noc). UCIe N/S/E/W appear as + cube-cube links. +- **cube_view** (`_build_cube_view`): inside one cube — router grid + + PE / M_CPU / SRAM / HBM_CTRL attachments + UCIe PHY edges. For + intra-cube routing / placement debugging. +- **pe_view** (`_build_pe_view`): inside one PE — nine sub-components + + internal edges (pe_internal kind). For detailed PE-internal + dataflow review. + +Views are selectively rendered via the spec's +`visualization.emit_views: [system, sip, cube]` (ADR-0006). The pe +view is omitted from default output but the code is retained for +detailed debugging. + +### D7. visualizer.py — SVG diagram output + +`emit_diagrams(graph, out_dir)` renders every view as SVG. Key +functions: + +- `_render_view_svg(view)` — generic view render (no router grid). +- `_render_cube_view_svg(view, spec)` — cube-view specific (HBM block, + router grid layout, PE/M_CPU/SRAM/HBM placement). +- `_draw_node`, `_draw_edge` — node/edge visual representation. +- `_pick_scale`, `_compute_node_sizes` — auto-scaling. + +The visualizer is a **derived artifact** (ADR-0006); changes here do +not pass production checks. Aligns with CLAUDE.md's "Derived +Artifacts" guidance. + +### D8. Blast radius of spec changes + +| spec field | effect | mesh regenerated? | +|---------------------------------------|---------------------|-------------------| +| `system.sips.count` | SIP count, node count | No | +| `sip.cube_mesh.w/h` | cube mesh shape | No | +| `cube.geometry.cube_mm.w/h` | cube size (mm) | **Yes** | +| `cube.pe_layout.corners/pe_per_corner`| PE attachment positions | **Yes** | +| `cube.ucie.n_connections` | UCIe PHY distribution | **Yes** | +| `cube.memory_map.hbm_mapping_mode` | HBM distribution mode | **Yes** | +| `cube.placement` | M_CPU/SRAM positions | **Yes** | +| `cube.memory_map.*` (besides above) | HBM capacity / BW | No | +| `*.links.*.bw_gbs` | edge bandwidth | No | +| `*.attrs.overhead_ns` | component latency | No | + +The table mirrors D2's `_compute_source_hash` inputs. Changes that +require mesh regeneration automatically invalidate `cube_mesh.yaml`'s +source_hash. + +## Alternatives Considered + +### A1. Regenerate the mesh on every compile without a cache file + +Rejected. The cost of mesh generation would be paid repeatedly (CLI +runs, probe, tests) for the same spec, and the human-inspectable +artifact would disappear. + +### A2. Merge mesh generation into builder.py + +Rejected (currently). It is a 305-line algorithm of its own, and the +mesh-layout decisions (placement-driven router attachment, HBM +exclusion zone) are different from builder's general node/edge +emission. Keeping it separate respects single-responsibility. + +### A3. Express placement coordinates in cube coordinates (col/row) + +Rejected. mm coordinates flow consistently between the visualizer and +mesh layout (for nearest-router computation). Cube coordinates are +undefined until the router grid is fixed, so they are unsuitable as +placement input. + +### A4. Lazy view projection generation + +Rejected (currently). The four views are cheap to build (typically < +100 ms), and eager construction guarantees `TopologyGraph` as the +single source of truth. + +### A5. Visualizer output in formats besides SVG (PNG/PDF) + +Rejected. SVG is vector + text-searchable + directly renderable in +browsers. PNG conversion, when required, is downstream +post-processing (e.g., rsvg-convert). + +## Consequences + +- ADR-0006's high-level intent is fleshed out via D1–D7; topology + changes can be assessed quickly via D8's table. +- D3's mesh-layout algorithm is ADR-locked, so future PE attachment + patterns (e.g., a 6-zone HBM split) make clear which stage they + affect. +- D5's edge-kind list and D7's view structure are explicit, giving PR + reviewers a quick map of where (builder + router + visualizer) a + new component type ripples through. +- D2's source_hash invalidation rules are explicit, so a stale + `cube_mesh.yaml` (e.g., when only bandwidth changed) is recognized + as correct behavior.