diff --git a/docs/adr/ADR-0024-sip-tp-launcher.md b/docs/adr/ADR-0024-sip-tp-launcher.md new file mode 100644 index 0000000..9e6ecbd --- /dev/null +++ b/docs/adr/ADR-0024-sip-tp-launcher.md @@ -0,0 +1,990 @@ +# ADR-0024: SIP-level TP Launcher — rank = SIP (host-driven dispatch) + +## Status + +Proposed (Revision 8 — Hierarchical content split out to ADR-0029) + +## Context + +### 목표 + +`torch.distributed` collective 호출의 참여 단위(rank)를 **SIP**(device) +경계에 맞춘다. 실제 PyTorch DDP/TP 스크립트와 **호스트 레벨에서 구분 없이** +읽히는 bench 코드를 목표로 한다. + +real PyTorch와 비교: + +| 차원 | real PyTorch | KernBench (이 ADR 이후) | +|---|---|---| +| 프로세스 모델 | N개 프로세스, 각 1 GPU | 1 프로세스, N greenlet, 각 1 SIP | +| `get_rank()` | `RANK` env var | greenlet-local 레지스트리 | +| `get_world_size()` | `WORLD_SIZE` env var | topology의 SIP 수 | +| `torch.cuda.set_device(r)` (real) / `torch.ahbm.set_device(r)` (KernBench) | rank → GPU | rank → SIP | +| `mp.spawn` | OS 프로세스 fork | greenlet fan-out | + +### 설계 원칙 — 공개 API의 추상화, 내부는 기존 path 활용 + +**공개 API (bench worker) 수준의 추상화**: +``` +rank = SIP +DPPolicy = intra-device (cube × PE) 분산만 +dist.all_reduce, torch.ahbm.set_device, mp.spawn 등 PyTorch-style 표면 +``` + +**Framework 내부 구현**: +``` +build_install_plans (host): topology + mapper + algorithm → SipInstallPlan + ↓ +backend (host): plan의 per-PE spec을 engine.submit으로 IpcqInitMsg 디스패치 + ↓ +engine: 기존 PE-scoped routing (MmuMapMsg 등과 동일 경로) + ↓ +PE_IPCQ: 자체 message loop에서 IpcqInitMsg 처리 (기존 capability) +``` + +**핵심**: 새 message 타입이나 IO_CPU 확장 없음. 기존 engine routing과 기존 +`IpcqInitMsg` 타입을 그대로 사용. 기존의 "sideband direct call" 우회만 +제거하여 convention 일원화. + +### 현재 상태 + +- `DistributedContext` facade 존재 +- `init_process_group("ahbm")` → `AhbmCCLBackend`가 `ctx.install_ipcq` 호출 + → `ccl/install.py`가 **sideband direct call** (`pe_ipcq._install_neighbors`)로 + PE_IPCQ에 neighbor table 설치 +- `get_rank()` 항상 `0` (single-driver) +- `get_world_size()` fallback: 총 PE 수 (rank = PE) +- `benches/ccl_allreduce.py`: `worker(rank=0, world_size=total_PEs)` 1회 호출 + +### 풀어야 할 문제 + +1. **공개 API에서 rank = SIP** — bench worker가 PE 개념을 알지 않도록. +2. **Multi-worker 실행** — N개 rank가 독립 worker 코드 실행. 1 프로세스 제약 + 하에서 greenlet + barrier 동기화. +3. **Cross-rank collective submit 동기화** — 첫 rank가 혼자 wait하면 peer 부재로 + SimPy deadlock. 모든 rank submit 후 drain 보장. +4. **기존 sideband install 제거** — IpcqInitMsg를 engine.submit으로 일원화. + MmuMapMsg 등 다른 control-plane 메시지와 동일 패턴. +5. **Algorithm / mapper / validator 분리** — 알고리즘 모듈은 kernel 코드만 + 담고, topology / mapping / validation은 registry + 선언. + +### Non-problem (이 ADR 밖) + +- IPCQ direction addressing fix → **ADR-0025** +- `DPPolicy.sip`/`num_sips` 제거 → **ADR-0026** +- Megatron-style TP → **ADR-0027** +- DTensor → **ADR-0028 (future)** +- **IO_CPU를 SIP-level control-plane 단일 endpoint로 승격**: 이 ADR에서는 + invariant으로 채택하지 않음. 현재 KernBench에 해당 원칙이 없고, 단독으로 + 도입하기엔 정당화가 약함. 미래에 control-plane latency 모델링 정밀도 요구가 + 생기면 별도 ADR. + +### TODO (이 ADR 구현 이후) + +- Tensor Parallelism (ADR-0027) +- Hierarchical all-reduce 알고리즘 설계 (ADR-0029) — 본 ADR의 mapper / + validator registry 인프라를 활용하는 첫 사례 + +--- + +## Decision + +### D1. rank = SIP (world_size 해석) + +```python +def _resolve_world_size(self) -> int: + if "world_size" in self._merged: + return int(self._merged["world_size"]) + defaults = self._cfg_all.get("defaults", {}) + if "world_size" in defaults: + return int(defaults["world_size"]) + spec = self.ctx.spec or {} + return int(spec.get("system", {}).get("sips", {}).get("count", 1)) +``` + +우선순위: 알고리즘 override > defaults override > SIP count. `ccl.yaml` +override는 legacy "rank = PE" 테스트 경로로 유지. + +### D2. Install 경로 — engine.submit 일원화 + +`ccl/install.py`의 sideband direct call을 제거하고, `IpcqInitMsg`를 +`engine.submit`으로 보낸다. MmuMapMsg / MemoryWriteMsg 등이 이미 동일 패턴. + +```python +# Backend (AhbmCCLBackend.__init__ 또는 init_process_group 시점) +from kernbench.ccl.install_plan import build_install_plans + +plans = build_install_plans( + world_size=self._world_size, + algorithm=self._merged["algorithm"], + algorithm_config=self._merged, + spec=self.ctx.spec, +) +self._plans = plans + +# Each PE_IPCQ가 자기 neighbor table을 받도록 engine 경유 submit +handles = [] +for plan in plans: + for pe_install in plan.pe_installs: + h = self.ctx.submit(IpcqInitMsg( + correlation_id=self.ctx.correlation_id, + request_id=f"ipcq_init_s{plan.sip}c{pe_install.cube}p{pe_install.pe}", + target_sips=(plan.sip,), + target_cubes=(pe_install.cube,), + target_pe=pe_install.pe, + entries=pe_install.neighbors, + buffer_kind=plan.buffer_kind, + n_slots=plan.n_slots, + slot_size=plan.slot_size, + # ... (기존 IpcqInitMsg 필드) + )) + handles.append(h) + +# Eager install — init_process_group이 반환하기 전에 완료 보장 +for h in handles: + self.ctx.wait(h) +``` + +**PE_IPCQ 컴포넌트**는 이미 `IpcqInitMsg`를 main loop에서 처리 (`pe_ipcq.py` +라인 145-147). 변경 불필요. 유일한 차이는 "message가 sideband Python call이 +아니라 engine queue를 거쳐 도착한다"는 점. + +**Correctness invariant (equivalence)**: `init_process_group()`은 모든 +install handle을 `wait()`한 후 반환하므로 launch-before-install 문제는 +구조적으로 없다. 남는 correctness 질문은 단 하나: + +> Engine-routed `IpcqInitMsg` 처리가 기존 sideband +> `pe_ipcq._install_neighbors(msg)` 호출과 **동일한 최종 PE_IPCQ 상태**를 +> 생성하는가. + +검증 포인트 (T3 참고): + +1. **State equivalence**: `_install_neighbors()` 내부 상태 전이가 engine + dispatch path에서도 동일하게 일어나 최종 PE_IPCQ state + (`_queue_pairs`, `_installed`, `_credit_inbox` 등)가 일치. + +2. **Sideband-only side effect 부재**: Sideband path에서만 있던 부수 효과가 + 없음 (예: engine.submit이 설정하는 request_id / correlation tracking 등이 + install semantics를 왜곡하지 않음). + +3. **Ordering independence**: 서로 다른 PE들의 install message가 engine + 큐에서 임의 순서로 처리되어도 최종 상태가 동일. 즉 install은 **PE별 + 독립 연산**이어야 하고, cross-PE 순서 의존성이 있으면 안 됨. + +4. **Idempotency**: 동일 PE에 대해 `IpcqInitMsg`가 두 번 도착하면? 현재 + 설계 전제는 "per-PE 단 한 번 install". 중복 install 시 동작은 정의되지 + 않음. 보수적 정책: + - 최초 install 시 `_installed = True`로 전이 + - 이후 중복 install msg는 **에러** (raise) 또는 **silent idempotent** + (no-op) 둘 중 하나로 명시 + - Recommend: **raise** (명시적 에러 → 버그 조기 검출). T3에 duplicate + install 케이스 추가. + +5. **Partial install visibility**: 일부 PE만 install 완료된 중간 상태가 + 외부에 observable한가? 현재 구조에서는 `init_process_group()`의 eager + wait-all이 barrier 역할을 하므로 partial state는 bench 코드에 노출되지 + 않음. 단, debugging / introspection API는 중간 상태를 볼 수 있음 (문제 + 아님, 문서화만). + +**Timing 영향**: Engine-routed install은 `init_process_group()`이 SimPy 시간을 +소비하게 만든다. 기존 sideband install은 사실상 zero-cost. ADR 계약: + +> Benchmarks must not rely on zero-cost initialization. +> `init_process_group()` consumes simulated time proportional to the number +> of participating PEs × per-PE install latency. First collective call +> starts at a well-defined but non-zero sim time. + +### D3. Launch 경로 — non-CCL 커널과 동일 primitive + +**CCL 커널은 non-CCL 커널과 동일한 `KernelLaunchMsg` submission path를 쓴다.** +Engine 내부의 IO_CPU/M_CPU transit 같은 것은 **기존 구현 세부이지 CCL-specific +장치가 아님**. Backend는 plan의 `participating_pes` 목록을 돌면서 `KernelLaunchMsg`를 +submit할 뿐이다. 새 메시지 타입 없음, 새 라우팅 경로 없음. + +```python +# AhbmCCLBackend.all_reduce +def all_reduce(self, tensor, op="sum"): + if op != "sum": + raise NotImplementedError(...) + if tensor._handle is None or not tensor._handle.shards: + raise RuntimeError(...) + + # Validator — global handle 기준 (D8) + validator_name = self._merged.get("validator") + if validator_name: + resolve_validator(validator_name)(tensor._handle, self._world_size, self.ctx.spec) + + rank = self.ctx.distributed.get_rank() + plan = self._plans[rank] + tensor_view = _tensor_slice_for_sip(tensor._handle, plan.sip) + + # Plan에서 kernel args 계산 (host-side) + import importlib + mod = importlib.import_module(plan.kernel_module) + n_elem = tensor_view.shards[0].nbytes // tensor.itemsize + kargs = mod.kernel_args(n_elem=n_elem, world_size=plan.world_size, + **plan.kernel_config) + + def _submit(): + out = [] + for (cube, pe) in plan.participating_pes: + h = self.ctx.submit(KernelLaunchMsg( + correlation_id=self.ctx.correlation_id, + request_id=f"allreduce_r{rank}_c{cube}p{pe}", + kernel_ref=KernelRef(name=plan.algorithm_name, kind="builtin"), + args=(_tensor_arg_for_pe(tensor_view, cube, pe), *kargs), + target_sips=(plan.sip,), + target_cubes=(cube,), + target_pe=pe, + )) + out.append(h) + return out + + self._barrier.submit_and_drain(self.ctx, rank, _submit) +``` + +### D4. Algorithm ABI — 얇게 + 명시적 arg 계약 + +각 알고리즘 모듈은 **kernel + kernel_args만 필수**. + +```python +# src/kernbench/ccl/algorithms/ring_allreduce.py +def kernel(t_ptr, n_elem, world_size, tl): + """PE-side kernel code. + + Signature convention: first positional arg is the tensor pointer + (per-PE slice), subsequent positional args are whatever + kernel_args() returns. `tl` is injected by the TLContext runtime. + """ + +def kernel_args(*, n_elem: int, world_size: int, **kw) -> tuple: + """Return the tuple of non-tensor positional args. + + Signature contract: + - Called keyword-only with n_elem and world_size plus kernel_config. + - Returns a tuple (possibly empty) of scalar / metadata args. + - The backend constructs the final KernelLaunchMsg.args as: + (per_pe_tensor_arg, *kernel_args(...)) + where per_pe_tensor_arg is a TensorArg containing only the shards + local to the receiving PE (derived from tensor_view). + """ + return (n_elem, world_size) +``` + +**Arg assembly in backend (reference)**: + +```python +# AhbmCCLBackend.all_reduce (D3에서 발췌) +kargs = mod.kernel_args(n_elem=n_elem, world_size=plan.world_size, + **plan.kernel_config) +for (cube, pe) in plan.participating_pes: + pe_tensor_arg = _tensor_arg_for_pe(tensor_view, cube, pe) + self.ctx.submit(KernelLaunchMsg( + args=(pe_tensor_arg, *kargs), # tensor first, then kernel_args return + target_sips=(plan.sip,), + target_cubes=(cube,), + target_pe=pe, + ... + )) +``` + +**ccl.yaml**에서 선언적 metadata: + +```yaml +algorithms: + ring_allreduce_tcm: + module: kernbench.ccl.algorithms.ring_allreduce + topology: ring_1d # kernbench/ccl/topologies.py + mapper: leader_only # kernbench/ccl/mappers.py (신규) + validator: single_shard_per_rank # kernbench/ccl/validators.py (신규) + buffer_kind: tcm + n_elem: 8 +``` + +- `topology` (필수) +- `mapper` (선택, default `"leader_only"`) +- `validator` (선택) + +알고리즘 모듈 자체에는 mapper/validator/participating_pes/neighbor +생성기가 **들어가지 않음**. + +### D5. Mapper + validator — registry key **또는** import path + +Host-side framework가 built-in registry 제공. 커스텀 확장은 dot-import path. + +```python +# src/kernbench/ccl/mappers.py (new) +Mapper = Callable[[dict, int], list[tuple[int, int]]] + +def leader_only(spec, rank): + """Single leader PE per SIP. Ring/tree/mesh용.""" + return [(0, 0)] + +def all_pes(spec, rank): + """Every PE in the SIP. 알고리즘이 intra-SIP 전체 PE를 참여시킬 때 사용 + (e.g. intra-SIP reduction, intra-SIP broadcast, hierarchical collective + 의 낮은 레벨 등).""" + cm = spec["sip"]["cube_mesh"] + pl = spec["cube"]["pe_layout"] + n_cubes = cm["w"] * cm["h"] + n_pes = pl["pe_per_corner"] * len(pl["corners"]) + return [(c, p) for c in range(n_cubes) for p in range(n_pes)] + +MAPPER_REGISTRY = {"leader_only": leader_only, "all_pes": all_pes} + +def resolve_mapper(key_or_path: str) -> Mapper: + if key_or_path in MAPPER_REGISTRY: + return MAPPER_REGISTRY[key_or_path] + if "." in key_or_path: + import importlib + mod_path, fn_name = key_or_path.rsplit(".", 1) + return getattr(importlib.import_module(mod_path), fn_name) + raise ValueError(f"unknown mapper: {key_or_path!r}") +``` + +Validator도 동일 패턴 (`src/kernbench/ccl/validators.py`). 입력은 **global +TensorHandle** (D8 참고). + +### D6. Host-side install plan builder + +```python +# src/kernbench/ccl/install_plan.py (new; 기존 install.py의 재구성) +from dataclasses import dataclass +from typing import Any, Mapping + +@dataclass(frozen=True) +class NeighborTableEntry: + direction: str + peer_direction: str # ADR-0025 + peer_sip: int + peer_cube: int + peer_pe: int + rx_base_pa: int + # ... 기타 IPCQ 설정 ... + +@dataclass(frozen=True) +class PeInstallSpec: + cube: int + pe: int + neighbors: tuple[NeighborTableEntry, ...] + +@dataclass(frozen=True) +class SipInstallPlan: + algorithm_name: str # human-readable ("ring_allreduce_tcm") + sip: int + rank: int + world_size: int + pe_installs: tuple[PeInstallSpec, ...] # per-PE neighbor tables + buffer_kind: str + n_slots: int + slot_size: int + kernel_module: str + participating_pes: tuple[tuple[int, int], ...] + kernel_config: Mapping[str, Any] + + +def build_install_plans( + world_size: int, + algorithm: str, + algorithm_config: dict, + spec: dict, +) -> list[SipInstallPlan]: + """Compose topology + mapper + algorithm into per-SIP plan list.""" + topo_fn = _resolve_topology(algorithm_config["topology"]) + mapper = resolve_mapper(algorithm_config.get("mapper", "leader_only")) + + # kernel_config: launch 시 kernel_args에 전달할 algorithm-specific params + kernel_config = { + k: v for k, v in algorithm_config.items() + if k in {"n_elem", "reduce_op", "chunk_size"} or k.startswith("kernel_") + } + + plans = [] + for rank in range(world_size): + sip = rank # identity mapping (non-identity는 open question) + pes = mapper(spec, rank) + pe_installs = _build_pe_installs( + rank=rank, world_size=world_size, sip=sip, + pes=pes, topo_fn=topo_fn, algorithm_config=algorithm_config, spec=spec, + ) + plans.append(SipInstallPlan( + algorithm_name=algorithm, + sip=sip, rank=rank, world_size=world_size, + pe_installs=pe_installs, + buffer_kind=algorithm_config["buffer_kind"], + n_slots=algorithm_config["n_slots"], + slot_size=algorithm_config["slot_size"], + kernel_module=algorithm_config["module"], + participating_pes=tuple(pes), + kernel_config=kernel_config, + )) + return plans +``` + +`_build_pe_installs`는 기존 `ccl/install.py`의 neighbor 계산 로직을 재활용 +(ADR-0025의 `reverse_direction` 개선 반영). + +**Multi-PE 매퍼와 neighbor 생성 책임**: mapper가 SIP 내 여러 PE를 반환하는 +경우 (`all_pes` 등), PE-level neighbor 그래프는 `_build_pe_installs` 내부에 +형성된다. 즉 topology 모듈은 rank-level 관계만 제공하고, PE-level 연결은 +builder에서 풀어낸다. 복잡한 multi-level 패턴을 쓰는 알고리즘은 이 책임 +분산이 관리 부담이 될 수 있음 — 관련 논의는 ADR-0029 참고. + +### D7. Epoch-based collective barrier + +Cross-rank submit 동기화. 각 collective 호출은 독립 epoch. 같은 rank의 +중복 join은 즉시 에러. + +```python +# src/kernbench/runtime_api/distributed.py +@dataclass +class _EpochState: + participants: set[int] = field(default_factory=set) + pending: list = field(default_factory=list) + drained: bool = False + returned: int = 0 + + +class _CollectiveBarrier: + """Epoch-based barrier. + + Contract: + - Each call joins the earliest non-drained epoch. + - Each rank may join a given epoch at most once. Duplicate join raises. + - Last arriver (participants == world_size) performs drain and advances + _next_epoch. Earlier arrivers yield and re-check drained on resume. + - Epoch state is GC'd when returned == world_size (success path). + - On failure paths, residual state is acceptable; reset() clears it. + """ + + def __init__(self, world_size: int): + self._world_size = world_size + self._next_epoch = 0 + self._state: dict[int, _EpochState] = {} + + def submit_and_drain(self, ctx, rank: int, submit_fn) -> None: + epoch = self._next_epoch + state = self._state.setdefault(epoch, _EpochState()) + + if rank in state.participants: + raise RuntimeError( + f"rank {rank} attempted duplicate join to epoch {epoch}" + ) + state.participants.add(rank) + + handles = submit_fn() + state.pending.extend(handles) + + is_last = len(state.participants) >= self._world_size + + if is_last: + for h in state.pending: + ctx.wait(h) + state.drained = True + self._next_epoch = epoch + 1 + else: + from greenlet import getcurrent + g = getcurrent() + if g.parent is None: + raise RuntimeError("barrier requires a bound worker greenlet") + while not state.drained: + g.parent.switch() + + state.returned += 1 + if state.returned >= self._world_size: + self._state.pop(epoch, None) + + def reset(self) -> None: + """Explicit cleanup on spawn exception unwinding.""" + self._state.clear() + self._next_epoch = 0 +``` + +### D8. Per-rank tensor view + validator contract + +**Validator** (host-side, pre-slice, global handle 기준): + +```python +# src/kernbench/ccl/validators.py +Validator = Callable[[TensorHandle, int, dict], None] + +def single_shard_per_rank(handle, world_size, spec): + """Ring 계열: 정확히 world_size개 shard, SIP당 1개.""" + if len(handle.shards) != world_size: + raise ValueError(...) + per_sip = {} + for s in handle.shards: + per_sip[s.sip] = per_sip.get(s.sip, 0) + 1 + if any(c != 1 for c in per_sip.values()): + raise ValueError(...) + +def multi_pe_sip_local(handle, world_size, spec): + """Multi-PE per SIP layout: 각 SIP에 intra-SIP PE 수만큼 shard 존재. + Intra-SIP 전체 PE를 참여시키는 알고리즘이 사용.""" + cm = spec["sip"]["cube_mesh"] + pl = spec["cube"]["pe_layout"] + per_sip = cm["w"] * cm["h"] * pl["pe_per_corner"] * len(pl["corners"]) + if len(handle.shards) != world_size * per_sip: + raise ValueError(...) + +VALIDATOR_REGISTRY = {...} +def resolve_validator(key_or_path): ... +``` + +Validator는 world 전체의 shard layout 불변량을 본다. Per-rank view는 +backend가 validator 호출 **후** `_tensor_slice_for_sip`로 생성. + +**Per-rank tensor view** — SIP-local slice: + +```python +def _tensor_slice_for_sip(handle, sip) -> TensorArg: + sip_shards = [s for s in handle.shards if s.sip == sip] + if not sip_shards: + raise RuntimeError(f"tensor has no shards on SIP {sip}") + # Deterministic ordering contract: (cube, pe, offset_bytes) ascending. + # Multi-PE mappers (hierarchical 등) rely on this ordering to align + # per-PE tensor arg construction with participating_pes enumeration. + sip_shards.sort(key=lambda s: (s.cube, s.pe, s.offset_bytes)) + min_offset = min(s.offset_bytes for s in sip_shards) + local_va_base = handle.va_base + min_offset if handle.va_base else 0 + return TensorArg( + shards=tuple(TensorArgShard(...) for s in sip_shards), + va_base=local_va_base, + ) +``` + +**Ordering invariant**: slice의 shard는 `(cube, pe, offset_bytes)` 오름차순. +Backend가 `participating_pes`를 iterate하며 `_tensor_arg_for_pe(view, cube, pe)`를 +구성할 때, 결정론적 ordering을 전제할 수 있다. 특히 `all_pes` mapper + +hierarchical 알고리즘이 per-PE slice 조합을 순서 의존적으로 해석하는 경우에 +중요. + +### D9. Greenlet-local rank registry (+ debug warning) + +```python +class DistributedContext: + def __init__(self): + self._backend = None + self._rank_by_greenlet: dict = {} + + def _bind_rank(self, g, rank: int) -> None: + self._rank_by_greenlet[g] = int(rank) + + def get_rank(self) -> int: + self._ensure_initialized() + from greenlet import getcurrent + g = getcurrent() + if g not in self._rank_by_greenlet: + if os.environ.get("KERNBENCH_DEBUG"): + warnings.warn( + "get_rank() called outside a bound greenlet — returning 0. " + "Likely a bug unless running single-driver." + ) + return 0 + return int(self._rank_by_greenlet[g]) +``` + +### D10. `torch.ahbm.set_device(rank)` — SIP 바인딩 + +KernBench 백엔드 이름은 `ahbm` (ADR-0023 D10). Real PyTorch는 +`torch.cuda.set_device(r)`이지만 우리는 CUDA가 아니므로 honestly-named +namespace를 사용한다. + +```python +class _AhbmNamespace: + """torch.ahbm — per-greenlet SIP device binding. + + Real-PyTorch parity idiom: ``torch.cuda.set_device(rank)``. Since + KernBench's backend is 'ahbm' (not CUDA), we expose the equivalent + API under ``torch.ahbm`` to avoid pretending to be a CUDA runtime. + """ + + def __init__(self): + self._device_by_greenlet: dict = {} + + def set_device(self, device: int) -> None: + from greenlet import getcurrent + self._device_by_greenlet[getcurrent()] = int(device) + + def current_device(self) -> int | None: + from greenlet import getcurrent + return self._device_by_greenlet.get(getcurrent()) + +# Attached to RuntimeContext as `self.ahbm = _AhbmNamespace()`. +# Bench code: `torch.ahbm.set_device(rank)` mirrors `torch.cuda.set_device`. +``` + +**PyTorch 2.x style 병행 지원**: 최신 PyTorch는 device-agnostic한 +`torch.accelerator` 네임스페이스를 지향 (`torch.accelerator.set_device_index(r)`, +`torch.accelerator.current_device_index()`). Device vendor에 종속되지 않는 +코드를 쓰려는 사용자를 위해 KernBench도 이 표면을 병행 지원한다. + +```python +class _AcceleratorNamespace: + """torch.accelerator — device-agnostic API (PyTorch 2.x style). + + Aliases torch.ahbm for bench code that prefers device-neutral idiom: + torch.accelerator.set_device_index(rank) + torch.accelerator.current_device_index() + """ + + def __init__(self, ahbm: _AhbmNamespace): + self._ahbm = ahbm + + def set_device_index(self, device: int) -> None: + self._ahbm.set_device(device) + + def current_device_index(self) -> int | None: + return self._ahbm.current_device() + +# RuntimeContext +self.ahbm = _AhbmNamespace() +self.accelerator = _AcceleratorNamespace(self.ahbm) # alias +``` + +Bench 작성자는 다음 중 하나를 선택 — 둘 다 내부적으로 같은 레지스트리를 보유: + +```python +torch.ahbm.set_device(rank) # KernBench-native, explicit backend +torch.accelerator.set_device_index(rank) # PyTorch 2.x device-agnostic +``` + +### D11. Tensor placement = structural (sip, cube, pe) 좌표 + +`resolve_dp_policy`가 `target_sip`을 직접 받아 구조적 좌표로 placement 생성. +세부는 ADR-0026. + +```python +# RuntimeContext._create_tensor +current_sip = self.ahbm.current_device() # (D10 naming) +if current_sip is None: + current_sip = 0 # single-driver fallback (D9와 일관) +placement = resolve_dp_policy( + dp, shape=shape_2d, itemsize=itemsize, + num_pe=eff_num_pe, num_cubes=eff_num_cubes, + target_sip=current_sip, +) +``` + +Post-hoc `pe_index` shifting 제거 — ShardSpec이 `(sip, cube, pe)` 구조적 +좌표 보유. + +### D12. `torch.multiprocessing.spawn`-compat surface + +Bench 작성자 표면은 real PyTorch `mp.spawn`과 동일: + +```python +# src/kernbench/runtime_api/multiprocessing.py (new) +def spawn(fn, args=(), nprocs=1, join=True, daemon=False, start_method="spawn"): + """Drop-in for torch.multiprocessing.spawn. + Internal: greenlet fan-out + epoch-barrier sync + exception propagation. + """ + ... + +# torch namespace에 부착 +torch.multiprocessing = SimpleNamespace(spawn=spawn) +``` + +Bench: + +```python +import torch.multiprocessing as mp +mp.spawn(worker, nprocs=world_size, args=(world_size, torch)) +``` + +### D13. Scheduler + exception handling + +```python +def spawn(fn, args, nprocs, ...): + dist = torch.distributed + gs: list[greenlet] = [] + errors: dict[int, Exception] = {} + + for rank in range(nprocs): + def _entry(r=rank): + try: + fn(r, *args) + except Exception as e: + errors[r] = e + raise + g = greenlet(_entry) + dist._bind_rank(g, rank) + gs.append(g) + + try: + while True: + alive = [g for g in gs if not g.dead] + if not alive: + break + for g in alive: + if not g.dead: + g.switch() + except Exception as outer: + for other in gs: + if not other.dead: + try: + other.throw(SystemExit) + except Exception: + pass + # Epoch barrier state 명시적 cleanup + backend = getattr(dist, "_backend", None) + if backend is not None and hasattr(backend, "_barrier"): + backend._barrier.reset() + raise SpawnException(errors) from outer +``` + +**Scheduler contract**: +- Deterministic round-robin over insertion order (rank 0, 1, ..., N-1). +- 동기화 지점은 epoch barrier (D7)만. Scheduler 순서에 의존하는 correctness 없음. +- 예외 발생 시 다른 greenlet 강제 종료 + `SpawnException` 전파. + +**Starvation guideline**: +- 일반적으로 collective barrier가 workers를 동기화. 큰 편차 없음. +- 극단적 non-collective 루프 대비 cooperative yield 제공: + `torch.distributed.cooperative_yield()`. + +### D14. Backward compatibility + +1. **Single-driver 호출**: `get_rank()` 0 반환 (D9). +2. **`ccl.yaml` world_size override**: D1 fallback 우회 — legacy "rank = PE" + 테스트 경로로 사용 가능. +3. **`DPPolicy.sip="column_wise"` 명시**: ADR-0026 scope. +4. **`install_ipcq()` compatibility wrapper**: + +기존 `ccl/install.py`의 `install_ipcq()` API는 곧바로 제거하지 않는다. +Thin compatibility wrapper로 남겨 기존 직접 호출자가 점진적으로 migration할 +수 있게 한다. + +```python +# src/kernbench/ccl/install.py (after this ADR) +def install_ipcq(engine, spec, merged, *, algo_module=None, rank_to_pe=None): + """DEPRECATED: legacy host-side PE installer. + + Internally delegates to build_install_plans + engine-routed IpcqInitMsg. + Use dist.init_process_group() instead. + """ + from kernbench.ccl.install_plan import build_install_plans + import warnings + warnings.warn( + "install_ipcq() is deprecated; use dist.init_process_group()", + DeprecationWarning, stacklevel=2, + ) + plans = build_install_plans( + world_size=merged.get("world_size", 1), + algorithm=merged["algorithm"], + algorithm_config=merged, + spec=spec, + ) + handles = [] + for plan in plans: + for pe_install in plan.pe_installs: + h = engine.submit(IpcqInitMsg( + target_sips=(plan.sip,), + target_cubes=(pe_install.cube,), + target_pe=pe_install.pe, + entries=pe_install.neighbors, + buffer_kind=plan.buffer_kind, + n_slots=plan.n_slots, + slot_size=plan.slot_size, + )) + handles.append(h) + for h in handles: + engine.wait(h) + return {"world_size": merged.get("world_size", 1), "plans": plans} +``` + +Migration 스케줄: +- Phase 1: wrapper로 유지 + DeprecationWarning +- Phase 2: 직접 호출자 grep-audit → 각각 `dist.init_process_group()` 또는 + `build_install_plans()` 직접 사용으로 이관 +- Phase 3: wrapper 제거 (별도 cleanup ADR 또는 PR) + +--- + +## Dependencies + +- **ADR-0023** (IPCQ): `IpcqInitMsg` 메시지 타입과 PE_IPCQ 핸들링을 그대로 + 활용. Engine-routed submit으로 전환하는 것이 유일한 변경. +- **ADR-0025** (IPCQ direction fix): `_build_pe_installs`의 neighbor 계산이 + 2-rank ring 등에서 정확히 동작하려면 필요. +- **ADR-0003 / 0016** (IO_CPU): IO_CPU는 기존 transit 역할 그대로. 본 ADR에서 + IO_CPU 역할 변경 없음. + +--- + +## Non-goals + +- **IPCQ protocol 수정**: ADR-0023 유지. +- **DPPolicy 필드 정리**: ADR-0026. +- **Megatron-style TP**: ADR-0027. +- **Multi-node (프로세스 간)**: 단일 프로세스. +- **IO_CPU SIP control-plane 단일 endpoint 원칙 채택**: 본 ADR 범위 밖. 현재 + KernBench에 이 원칙이 없고, 도입은 별도 ADR. +- **Hierarchical all-reduce 알고리즘 설계**: ADR-0029. 본 ADR은 그 알고리즘이 + 쓸 framework 인프라 (`all_pes` mapper, `multi_pe_sip_local` validator, + registry 확장점)만 제공. + +--- + +## Open questions + +### 🔴 Critical — 구현 blocker 가능성 (integration 전 반드시 검증) + +- **`IpcqInitMsg`의 engine routing — primary implementation risk**: 현재 + sideband만 쓰여서 engine routing path가 실사용 검증되지 않은 상태. **본 + ADR 전체가 "engine routing이 동작한다"는 가정 위에 서 있다**. 이것이 + 실제로 안 되면 D2, D14, T3 등이 전부 영향 받음. 반드시 **ADR 구현 착수 + 전 스파이크 검증**: + - `engine.submit(IpcqInitMsg(target_sips=..., target_cubes=..., target_pe=...))` + 가 PE_IPCQ로 정확히 배달되는지 (기존 `MmuMapMsg` / `MemoryWriteMsg` 라우팅 + 패턴과 비교) + - 미지원 시 minor hook: engine의 message-type → component-kind 매핑 테이블에 + `IpcqInitMsg → "pe_ipcq"` 등록 (localized change, topology builder / + message schema 영향 없음) + - 결과에 따라 D2 채택 여부가 달라질 수 있음 — 만약 routing 불가 시 sideband + path 유지로 fallback 후 본 ADR 범위 재조정 + +- **Engine-routed install vs sideband equivalence** (D2 검증점 1-5): T3의 + equivalence test가 실제 동작하는지 스파이크. 특히 ordering independence와 + idempotency는 기존 테스트에 없는 속성이라 신규 검증 필요. + +- **`install_ipcq()` 직접 호출자 audit** (구현 전 필수): deprecated wrapper + 전략은 적절하지만 실제 migration 리스크는 호출자 목록에 따라 다름. 착수 전 + grep audit: + - Pattern: `install_ipcq(` (cwd 전체) + - Scope: `src/`, `tests/`, `benches/`, `scripts/`, `src/kernbench/cli/` + - 각 호출자의 예상 migration path (→ `dist.init_process_group` vs + `build_install_plans` 직접)를 정리한 후 wrapper 도입 + +### 🟡 Nice-to-have — scope 경계 관련 + +- **Install timing 허용치**: SimPy 시간 상 install이 몇 ns~us 소모. 기존 + sideband는 0ns. 기존 테스트가 t=0 시작을 전제로 하는지 확인 (audit 결과에 + 따라 테스트 교정 필요). + +- **`IpcqInitMsg` 배치 가능성**: MmuMapMsg처럼 `target_pe="all"` 브로드캐스트 + 는 IPCQ에서는 부적합 (PE마다 neighbor가 다름). 현재는 per-PE 개별 submit. + Per-PE payload를 담는 batched IpcqInitMsg 타입은 future optimization. + +- **`_rank_to_sip` 매핑**: 현재 identity. Non-trivial mapping 요구 시 별도. + +- **Cooperative yield API 위치**: `torch.distributed.cooperative_yield()`로 + 노출 예정. 실제 필요성은 Phase 2 이후 벤치 추가 시 판단. + +(PE-level topology 일원화 관련 중장기 방향은 **ADR-0029** 참고 — 복잡한 +multi-level 알고리즘이 driving force가 되는 framework 진화 방향.) + +--- + +## Test strategy + +### T1. Launcher infrastructure + +`tests/test_ccl_ddp_launcher.py`: +- `test_world_size_equals_sip_count` — D1 +- `test_ahbm_set_device_binds_tensor_to_single_sip` — D10/D11 +- `test_get_rank_is_greenlet_local` — D9 +- `test_run_spawns_one_worker_per_rank` — D12/D13 +- `test_get_rank_debug_warning` — D9 warning path + +### T2. Install plan builder + +`tests/test_ccl_install_plan.py` (new): +- `build_install_plans` — ring_1d × leader_only 조합 (단일 PE per rank) +- `build_install_plans` — ring_1d × all_pes 조합 (multi-PE per rank; mapper + framework 동작 확인, 알고리즘-무관) +- Mapper / validator registry resolution (built-in key vs import path vs + unknown) +- Import path fallback (`"pkg.mod.fn"` 형식) 동작 검증 + +### T3. Engine-routed IpcqInitMsg (equivalence — 핵심 검증) + +`tests/test_ipcq_init_routing.py` (new): +- **Routing**: `engine.submit(IpcqInitMsg)` → 지정 PE_IPCQ가 실제 설치 수행 +- **Equivalence**: 동일한 IpcqInitMsg를 (a) sideband `_install_neighbors` + 직접 호출, (b) engine.submit 두 경로로 보낸 뒤 PE_IPCQ 최종 state + (`_queue_pairs`, `_installed` 등) 동일성 비교 +- **Ordering independence**: 서로 다른 PE의 install msg를 engine 큐에 임의 + 순서로 넣어도 최종 state가 동일 +- **Idempotency (duplicate install)**: 동일 PE에 두 번 install msg → 두 + 번째는 에러 raise (policy: explicit error; D2 검증점 4 참고) +- **Multi-PE 병렬 install**: per-PE submit이 interference 없이 완료 +- **Install 후 send 성공**: 설치 직후 `IpcqSendCmd` 실행해서 neighbor table + state가 실제로 유효한지 확인 + +### T4. Barrier correctness + +`tests/test_collective_barrier.py` (new): +- Single collective 정상 +- 다중 collective 연속 호출 (epoch 격리) +- 동일 rank의 duplicate join → RuntimeError +- Rank 1이 all_reduce 전 종료 → SpawnException + barrier.reset() +- Conditional branch 시 모든 rank 도달하면 정상 + +### T5. E2E + +`tests/test_ccl_allreduce_matrix.py`: +- `ring_tcm` / `ring_hbm` / `ring_sram` @ ws=SIP_count + +### T6. 회귀 + +기존 `test_ccl_framework`, `test_ccl_install`, `test_ccl_topologies`, +`test_ccl_mock_runtime`, `test_pe_ipcq`, `test_ipcq_e2e`, 기타 non-CCL +모두 통과. + +--- + +## Consequences + +### Positive + +- **새 message 타입 0개**: 기존 `IpcqInitMsg` + `KernelLaunchMsg`만으로 구현. +- **IO_CPU / engine 변경 없음**: 기존 routing 그대로. +- **Sideband install convention 제거**: MmuMapMsg 등과 동일 패턴으로 일원화. +- **Plan state stale 문제 소멸**: Plan은 host 단일 소유. +- **Bench = real PyTorch DDP** (공개 API 관점). +- **Algorithm ABI 경량**: `kernel` + `kernel_args`만 필수. +- **Epoch-based barrier**: interleaved collective 안전. +- **Control/data plane 분리**: data plane(PE_IPCQ)은 ADR-0023 유지, control + plane은 host-driven. +- 장기 확장성: Megatron TP, DTensor 기반. + +### Negative + +- 신규 모듈: `install_plan.py`, `mappers.py`, `validators.py`, + `multiprocessing.py`. +- Engine이 `IpcqInitMsg`를 엔진-path로 라우팅할 수 있는지 구현 시 확인 필요 + (minor hook 가능성). +- Install이 SimPy 시간을 소모 (positive로도 볼 수 있으나, 기존 sideband 시점 + 0ns 전제인 테스트가 있으면 교정 필요). + +### Neutral + +- IPCQ PE-level protocol (ADR-0023) 불변. +- `DPPolicy` 필드 변경은 ADR-0026. +- IO_CPU 역할 불변 (기존 transit 그대로). + +--- + +## Affected files + +| File | Change | +|------|--------| +| `src/kernbench/runtime_api/distributed.py` | D1/D2/D7/D9: world_size fallback, rank_to_sip, plan 소유, engine-routed install/launch, epoch barrier | +| `src/kernbench/runtime_api/context.py` | D10/D11: `_AhbmNamespace`, `ctx.ahbm`, `_create_tensor`가 `target_sip` 전달 | +| `src/kernbench/runtime_api/multiprocessing.py` (new) | D12/D13: `spawn` + scheduler + exception | +| `src/kernbench/ccl/install_plan.py` (new) | D6: `build_install_plans`, `SipInstallPlan`, `PeInstallSpec`, `NeighborTableEntry` | +| `src/kernbench/ccl/mappers.py` (new) | D5: `leader_only`, `all_pes`, registry + resolver | +| `src/kernbench/ccl/validators.py` (new) | D5: validator registry + resolver | +| `src/kernbench/ccl/install.py` | Thin deprecated compat wrapper (D14) | +| `src/kernbench/ccl/algorithms/ring_allreduce.py` | D4: `kernel` + `kernel_args` 유지 (큰 변화 없음) | +| `src/kernbench/ccl/algorithms/mesh_allreduce.py` | D4 동일 | +| `src/kernbench/ccl/algorithms/tree_allreduce.py` | D4 동일 | +| `ccl.yaml` | 각 알고리즘에 `mapper` / `validator` 선언 추가 | +| `src/kernbench/sim_engine/engine.py` | (If needed) `IpcqInitMsg` → PE_IPCQ 라우팅 확인 hook | +| `benches/ccl_allreduce.py` | 새 launcher 기반 rewrite | +| `tests/test_ccl_ddp_launcher.py` (new) | T1 | +| `tests/test_ccl_install_plan.py` (new) | T2 | +| `tests/test_ipcq_init_routing.py` (new) | T3 | +| `tests/test_collective_barrier.py` (new) | T4 | +| `tests/test_ccl_allreduce_matrix.py` | T5: ws=SIP_count 단순화 | diff --git a/docs/adr/ADR-0025-ipcq-direction-addressing.md b/docs/adr/ADR-0025-ipcq-direction-addressing.md new file mode 100644 index 0000000..5ef6e62 --- /dev/null +++ b/docs/adr/ADR-0025-ipcq-direction-addressing.md @@ -0,0 +1,365 @@ +# ADR-0025: IPCQ Direction Addressing — address-based matching + +## Status + +Proposed (Revision 2 — Address-based matching; peer_direction field dropped) + +## Context + +### 목표 + +ADR-0023의 IPCQ protocol에서 **"어느 direction pair를 통한 전송인가"의 식별**을 +topology / dict-order에 의존하지 않고 **주소 기반**으로 일관되게 한다. +2-rank bidirectional ring (또는 여러 direction이 동일 peer를 가리키는 +topology 일반)에서 정확히 동작하도록 한다. + +### 현재 상태 (ADR-0023 D9 구현) + +`src/kernbench/components/builtin/pe_ipcq.py` — `_handle_meta_arrival`: + +```python +def _handle_meta_arrival(self, msg: IpcqMetaArrival) -> None: + token = msg.token + sender_key = (token.src_sip, token.src_cube, token.src_pe) + for d, qp in self._queue_pairs.items(): + p = qp["peer"] + if (p.sip, p.cube, p.pe) == sender_key: + qp["peer_head_cache"] = max(qp["peer_head_cache"], token.sender_seq + 1) + # ... wake recv waiters ... + return +``` + +`_credit_worker`도 동일한 "sender-coord-first-match" 패턴. + +`src/kernbench/ccl/install.py` — `reverse_direction`: + +```python +def reverse_direction(my_rank: int, peer_rank: int) -> str | None: + for d, target in neighbor_table[peer_rank].items(): + if target == my_rank: + return d + return None +``` + +### 드러난 버그 — 2-rank bidirectional ring + +`ring_1d(rank, world_size=2)` → `{"E": 1, "W": 1}` (rank 0). 양쪽 방향이 같은 peer. + +**버그 1 (install)**: +- `reverse_direction(0, 1)` → dict order로 "E" 반환 (틀림, "W"가 맞음 — opposite + direction convention) +- rank 0의 E entry가 `peer.rx_base_pa = rx_base(sip1, cube0, pe0, d="E")`로 설정 +- tl.send(E) → data가 sip1의 E-rx buffer로 landing (should be W-rx) + +**버그 2 (runtime)**: +- 설령 install이 올바른 주소로 설정해도, receiver의 `_handle_meta_arrival`이 + sender 좌표만으로 direction 매칭 → 첫 direction (E) 승 +- peer_head_cache[E] 증가, peer_head_cache[W]는 불변 +- Kernel의 tl.recv(W)는 peer_head_cache[W] 대기 → 영원히 블록 → IpcqDeadlock + +### 근본 원인 + +두 축에서 동일 문제: +1. **Install-time pairing**: "내 direction과 peer의 어느 direction이 짝인가" + 결정이 dict-iteration-order에 의존 → 여러 direction이 같은 peer를 가리킬 때 + fragile +2. **Runtime identification**: "어느 qp를 업데이트해야 하는가" 결정이 sender + 좌표만으로 이루어짐 → direction 중복 시 ambiguous + +### 해결 방향 — address-based matching + +각 PE의 rx buffer는 **direction별로 고유한 주소 range**에 위치 (rx_base_pa + +direction_idx × bytes_per_direction). 따라서: + +- **Runtime**: sender coord 대신 **dst_addr 범위**로 매칭 → unambiguous +- **Install**: opposite-direction 우선 선택 heuristic (ring / mesh의 자연스러운 + 대칭성) +- `peer_direction` 같은 이중 메타데이터 불필요 — **주소가 single source of + truth** + +이 설계는 **PhysAddr 전환 (ADR-0030)과 독립적**으로 작동. 현재 synthetic +주소든 PhysAddr든 direction별 range 유일성만 지켜지면 동일하게 적용 가능. + +--- + +## Decision + +### D1. Install — `reverse_direction` opposite-preference + +`src/kernbench/ccl/install.py`: + +```python +_OPPOSITE_DIR = {"E": "W", "W": "E", "N": "S", "S": "N"} + +def reverse_direction(my_rank: int, peer_rank: int, my_dir: str) -> str | None: + """Find peer's direction that reciprocates my_dir→peer_rank. + + Prefer the OPPOSITE direction (E↔W, N↔S) when the peer has it + pointing back to us. This matters in 2-rank bidirectional rings + where both E and W on one side point to the same peer — without + the preference, the first-match-wins iteration would route data + into the wrong rx slot. Falls back to any direction pointing back + for topologies without an opposite convention (tree_binary's + parent/child). + """ + nt = neighbor_table[peer_rank] + opp = _OPPOSITE_DIR.get(my_dir) + if opp is not None and nt.get(opp) == my_rank: + return opp + for d, target in nt.items(): + if target == my_rank: + return d + return None +``` + +호출부: + +```python +for d, peer_rank in nbrs.items(): + peer_dir = reverse_direction(r, peer_rank, d) # my_dir 전달 + if peer_dir is None: + continue + ... +``` + +### D2. Runtime — `_handle_meta_arrival` dst_addr 매칭 + +`src/kernbench/components/builtin/pe_ipcq.py`: + +```python +def _handle_meta_arrival(self, msg: IpcqMetaArrival) -> None: + """Match incoming token to the receiver-side direction by dst_addr range. + + Each direction has a unique rx buffer address range + (my_rx_base_pa + n_slots * slot_size). The token's dst_addr (set by + the sender's IPCQ when computing peer's slot address) falls within + exactly one such range. This address-based matching is unambiguous + even when multiple directions have the same peer (2-rank ring). + """ + token = msg.token + dst_addr = token.dst_addr + for d, qp in self._queue_pairs.items(): + base = qp["my_rx_base_pa"] + size = qp["n_slots"] * qp["slot_size"] + if base <= dst_addr < base + size: + qp["peer_head_cache"] = max(qp["peer_head_cache"], + token.sender_seq + 1) + self._arrived_tokens.setdefault(d, []).append(token) + waiters = self._recv_waiters.get(d, []) + self._recv_waiters[d] = [] + for ev in waiters: + if not ev.triggered: + ev.succeed() + any_waiters = self._any_recv_waiters + self._any_recv_waiters = [] + for ev in any_waiters: + if not ev.triggered: + ev.succeed() + return + # Unknown dst_addr — diagnostic log (should not happen under correct install) +``` + +Sender 좌표 검사는 **제거**. `dst_addr`가 이미 direction을 결정. + +### D3. Credit — `dst_rx_base_pa` 필드 추가 + +`src/kernbench/common/ipcq_types.py`: + +```python +@dataclass(frozen=True) +class IpcqCreditMetadata: + consumer_seq: int + dst_rx_base_pa: int # NEW: 원 sender의 peer.rx_base_pa와 매칭용 + # 기존 필드 (diagnostic / log 용도로 유지) + src_sip: int + src_cube: int + src_pe: int + src_direction: str +``` + +Credit 생성 시 (`_delayed_credit_send`): 자기 direction의 `my_rx_base_pa`를 +`dst_rx_base_pa`로 실어 보냄 (이게 상대방이 sender 당시 썼던 `peer.rx_base_pa`). + +수신 측 (`_credit_worker`): + +```python +def _credit_worker(self, env): + while True: + credit = yield self._credit_inbox.get() + for d, qp in self._queue_pairs.items(): + # peer의 rx_base_pa와 credit의 dst_rx_base_pa가 일치하는 qp 찾기 + if qp["peer"].rx_base_pa == credit.dst_rx_base_pa: + qp["peer_tail_cache"] = max(qp["peer_tail_cache"], + credit.consumer_seq) + waiters = self._send_waiters.get(d, []) + self._send_waiters[d] = [] + for ev in waiters: + if not ev.triggered: + ev.succeed() + break +``` + +Sender 좌표 검사 제거. `dst_rx_base_pa` 매칭으로 unambiguous. + +### D4. `IpcqInitEntry`에 `peer_direction` 필드를 **추가하지 않음** + +ADR-0025 rev 1에서 제안했던 `IpcqInitEntry.peer_direction`은 **불필요**. +이유: +- Meta arrival은 dst_addr로 매칭 (D2) +- Credit은 dst_rx_base_pa로 매칭 (D3) +- qp에 peer_direction 저장 필요 없음 +- Install은 rx_base_pa 계산 시 내부적으로만 peer_dir 사용 (`reverse_direction`) + +IpcqInitEntry schema 변경 없음. Rev 1 대비 **단순화**. + +### D5. `IpcqDmaToken.src_direction` 유지 (diagnostic only) + +기존 `src_direction` 필드는 제거하지 않는다. 다음 용도로 유지: +- Logging / trace: `KERNBENCH_CCL_TRACE=1` 출력의 `(rank, t, dir, nbytes)` +- Diagnostics: pointer_dump 등에서 direction 표시 +- 미래 확장 여지 + +Runtime matching은 `dst_addr`만 사용. + +### D6. Invariants (ADR-0023 I3 강화) + +**I3 (엄격)**: 각 방향 pair `(my_direction, peer_direction)`에 대해 my +rx_base와 peer rx_base는 **별개의 direction slot**을 가리켜야 함. Install은 +이를 보장해야 한다 (reverse_direction opposite-preference). + +**I3.1 (신규)**: 모든 qp에 대해 `qp["my_rx_base_pa"]`와 `qp["peer"].rx_base_pa`는 +서로 disjoint한 주소 range를 점유한다 (다른 direction의 buffer는 절대 겹치지 +않음). 이것이 D2/D3의 주소-기반 매칭의 전제. + +Install time에 검증 가능: +```python +# ccl/install_plan.py: build_install_plans 끝에 assertion +all_rx_ranges = set() +for plan in plans: + for pe_install in plan.pe_installs: + for entry in pe_install.neighbors: + r = (entry.my_rx_base_pa, + entry.my_rx_base_pa + plan.n_slots * plan.slot_size) + overlap = any(_ranges_overlap(r, e) for e in all_rx_ranges) + assert not overlap + all_rx_ranges.add(r) +``` + +--- + +## Dependencies + +- **ADR-0023** (IPCQ protocol): 본 ADR은 ADR-0023의 runtime 매칭 로직 수정 + (D2, D3) + install heuristic 개선 (D1). IPCQ 프로토콜의 semantic layer + 변경은 없음. +- **ADR-0024** (launcher): 2-rank bidirectional ring이 실제 쓰이는 경우가 + ADR-0024의 ws=SIP_count 모델. 본 ADR이 그 케이스를 작동시킴. +- **ADR-0030** (PhysAddr transition, stub): **독립적** — ADR-0025의 + 주소-기반 매칭은 현재 synthetic 주소든 PhysAddr이든 동일하게 작동. + +--- + +## Non-goals + +- **IPCQ 주소 체계를 PhysAddr로 전환**: ADR-0030 scope. 본 ADR은 주소가 어떻게 + 인코딩되는가와 무관. +- **Multi-hop routing**: ADR-0023 D5의 single-hop DMA write 전제 유지. +- **Unidir ring 특수화**: `ring_1d_unidir`는 direction 하나만 있으므로 본 버그 + 무관. + +--- + +## Open questions + +- **주소 매칭 성능**: `_handle_meta_arrival`과 `_credit_worker`가 qp를 선형 + 순회 (max 4 direction). 성능 영향 무시 가능 수준. 문제 시 dict lookup으로 + 전환 가능 (`_qp_by_rx_base`). +- **`IpcqDmaToken.src_direction` 필요성 재평가**: diagnostic 용도로만 남긴 + 필드를 계속 유지할지, 또는 logging 외부로 분리할지. 현재는 유지. +- **Install-time invariant 검증 cost**: D6의 I3.1 검증은 O(N_PE × N_direction)^2. + 대형 topology에서 느려질 수 있음 → interval tree 등 자료구조로 개선 가능. + 단순 구현 먼저. + +--- + +## Test strategy + +### T1. Unit — `reverse_direction` opposite-preference + +`tests/test_ccl_install.py` (확장): +- Ring ws=2: `reverse_direction(0, 1, "E")` → "W", `reverse_direction(0, 1, "W")` → "E" +- Ring ws=4: `reverse_direction(0, 1, "E")` → "W" (자연스러운 opposite) +- Mesh 2×2: `reverse_direction(r, peer, "N")` → "S", "E" ↔ "W" +- Tree binary: opposite 없는 direction (parent) → fallback 경로 +- Non-symmetric topology: opposite가 peer에 없고 다른 direction만 있는 경우 + +### T2. Runtime — `_handle_meta_arrival` dst_addr 매칭 + +`tests/test_pe_ipcq.py` (확장): +- 2-rank pair install 후, E direction dst_addr로 meta arrival → E의 `peer_head_cache` + 증가 (W는 불변) +- W direction dst_addr로 meta arrival → W의 `peer_head_cache` 증가 +- 잘못된 dst_addr (어느 rx range에도 속하지 않음) → 에러 또는 silent drop + (결정 후 명시) + +### T3. Credit — `dst_rx_base_pa` 매칭 + +`tests/test_pe_ipcq.py` (확장): +- E direction send 후 peer가 consume → credit에 자기 W의 `my_rx_base_pa` + 담아 송신 → sender의 E direction `peer_tail_cache` 증가 +- W direction도 동일 + +### T4. E2E — 2-rank bidirectional ring + +`tests/test_ipcq_e2e.py`: +- 2-rank ring_1d로 tl.send(E) + tl.recv(W) pattern이 양방향으로 작동 +- ADR-0024의 `test_ccl_allreduce_matrix.py`에서 ring at ws=2가 통과 + +### T5. Install invariant — rx_base range disjointness + +`tests/test_ccl_install_plan.py` (확장): +- I3.1 검증: `build_install_plans` 결과에서 모든 qp의 rx_base range가 disjoint + +### T6. 회귀 + +- 기존 ws≥3 ring / mesh / tree 테스트 그대로 통과 +- `test_pe_ipcq`, `test_ipcq_e2e` 기존 케이스 회귀 + +--- + +## Consequences + +### Positive + +- **단순함**: `peer_direction` 이중 메타데이터 제거. 주소가 single source of truth. +- **Unambiguous matching**: 모든 topology (direction 중복 포함)에서 동작. +- **Schema 변경 최소**: `IpcqInitEntry` 불변, `IpcqCreditMetadata`에 1 필드 추가. +- **PhysAddr 전환 (ADR-0030) 독립**: 주소-기반 매칭은 주소 인코딩 방식과 무관. +- **Diagnostic 유지**: `IpcqDmaToken.src_direction`은 로깅 용도로 존치. + +### Negative + +- Runtime 매칭이 주소 비교로 바뀌어서 디버깅 시 "왜 peer_head_cache[E]가 아닌 + W가 업데이트됐나" 같은 질문에 address range를 추적해야 함 (기존엔 direction + 이름으로 충분). 해결: pointer_dump에 "direction ↔ rx_base_pa" 매핑 포함. + +### Neutral + +- IPCQ protocol의 semantic layer (sender가 dst_addr 계산, receiver가 수신)는 + 불변. + +--- + +## Affected files + +| File | Change | +|------|--------| +| `src/kernbench/ccl/install.py` | D1: `reverse_direction`에 `my_dir` 인자 추가, opposite-preference | +| `src/kernbench/components/builtin/pe_ipcq.py` | D2: `_handle_meta_arrival` dst_addr 매칭 / D3: `_credit_worker` dst_rx_base_pa 매칭 / `_delayed_credit_send`가 `dst_rx_base_pa` 필드 채움 | +| `src/kernbench/common/ipcq_types.py` | D3: `IpcqCreditMetadata`에 `dst_rx_base_pa` 필드 추가 | +| `src/kernbench/ccl/install_plan.py` (ADR-0024 신규) | D6: I3.1 invariant 검증 (optional) | +| `docs/adr/ADR-0023-ipcq-pe-collective.md` | Reference note: runtime 매칭 방식이 ADR-0025에서 바뀜 | +| `tests/test_ccl_install.py` | T1 | +| `tests/test_pe_ipcq.py` | T2, T3 | +| `tests/test_ipcq_e2e.py` | T4 | +| `tests/test_ccl_install_plan.py` | T5 | diff --git a/docs/adr/ADR-0026-dppolicy-intra-device.md b/docs/adr/ADR-0026-dppolicy-intra-device.md new file mode 100644 index 0000000..3d87aad --- /dev/null +++ b/docs/adr/ADR-0026-dppolicy-intra-device.md @@ -0,0 +1,476 @@ +# ADR-0026: DPPolicy = Intra-Device Only — sip/num_sips 필드 제거 + +## Status + +Proposed (Revision 4 — 문서 일관성 + grep audit 구체화) + +## Context + +### 목표 + +`DPPolicy`를 **한 device(SIP) 내부의 cube × PE 분산**만 표현하는 순수한 +intra-device 추상화로 명확화한다. SIP 간 분산(TP)은 별도 레이어로 분리 +(ADR-0024의 `torch.ahbm.set_device(rank)` 또는 ADR-0027의 Megatron parallel +layers가 담당). + +### 현재 상태 + +`src/kernbench/policy/placement/dp.py`: + +```python +@dataclass(frozen=True) +class DPPolicy: + sip: Literal["replicate", "column_wise", "row_wise"] = "replicate" + cube: Literal["replicate", "column_wise", "row_wise"] = "replicate" + pe: Literal["replicate", "column_wise", "row_wise"] = "replicate" + num_pes: int | None = None + num_cubes: int | None = None + num_sips: int | None = None # ← 제거 대상 +``` + +`sip` / `num_sips` 필드는 텐서를 SIP 경계 **너머**로 분산하는 경로를 제공함. +이는: + +- **ADR-0024의 launcher 모델과 충돌**: ADR-0024는 "rank = SIP = 1 worker per SIP" + 모델. 각 worker가 자기 SIP에 텐서를 생성. 텐서가 여러 SIP에 걸치는 경우는 + Megatron-style TP가 개별 primitive로 처리해야 함. +- **사용자 의도와 불일치**: "DPPolicy는 한 디바이스 내에서 PE들로 분산하는 방법" + (사용자 진술). +- **개념 혼동**: `DPPolicy.sip="column_wise"`는 실제로 **TP**. 이름이 DP인데 + 하는 일은 TP → 신규 사용자에게 혼란. + +### 영향받는 call site (rollback 시점 grep 결과) + +**생성 사이트** (`DPPolicy(sip=...` 또는 `num_sips=...`): +- `tests/test_runtime_api_tensor.py` +- `benches/ccl_allreduce.py` (ADR-0024 scope 내에서 이미 개편됨) +- `tests/test_va_offset.py` +- `benches/va_offset_verify.py` +- `tests/test_sip_parallel.py` + +**참조 사이트** (`dp.sip`, `policy.sip`, `num_sips` 등): +- `src/kernbench/runtime_api/context.py` (`_create_tensor`, `launch`) +- `src/kernbench/components/builtin/pe_cpu.py` +- `src/kernbench/components/legacy/builtin/pe_cpu.py` +- `src/kernbench/policy/placement/dp.py` (구현 자체) +- `tests/test_tensor.py`, `test_ipcq_types.py` + +**핵심 테스트**: `test_sip_parallel.py`는 이름 그대로 "SIP 병렬성을 DPPolicy로 +표현하는" 테스트. 이 ADR 이후 **새 launcher 모델로 재작성** 필요. + +--- + +## Decision + +### D1. `DPPolicy`에서 `sip` + `num_sips` 필드 제거 + +```python +@dataclass(frozen=True) +class DPPolicy: + """Intra-device (cube × PE) data-parallel policy. + + SIP-level placement is controlled by ``torch.cuda.set_device(rank)`` + (ADR-0024) and, for model-level TP, by Megatron-style parallel layers + (ADR-0027). DPPolicy does not cross SIP boundaries. + """ + cube: Literal["replicate", "column_wise", "row_wise"] = "replicate" + pe: Literal["replicate", "column_wise", "row_wise"] = "replicate" + num_pes: int | None = None + num_cubes: int | None = None +``` + +제거되는 필드: `sip`, `num_sips`. + +### D2. `ShardSpec` — structural (sip, cube, pe) 좌표, `pe_index` 완전 제거 + +현재 `ShardSpec.pe_index`는 **global flat index** (`sip × cubes × pes + cube × +pes + pe`). 이는 ADR-0024 D11이 "abstraction leakage"로 지적한 형태. + +본 ADR에서 ShardSpec을 **structural 좌표로 재정의**하고, `pe_index`는 +property로도 **남기지 않는다**: + +```python +# src/kernbench/policy/placement/dp.py (after) +@dataclass(frozen=True) +class ShardSpec: + """Structural shard placement — intra-SIP (cube × PE) coord. + + Global-flat `pe_index` was removed in ADR-0026. Callers must use + structural coords (sip, cube, pe) directly. If a flat integer key is + needed (e.g. dict lookup), compute it explicitly at the call site. + """ + sip: int # structural — which SIP this shard lives on + cube: int # local within SIP + pe: int # local within cube + offset_bytes: int + nbytes: int +``` + +**핵심 원칙**: +- ShardSpec의 정체성은 `(sip, cube, pe)` 3튜플. +- **`pe_index` property도 없음** — silent semantics drift 차단. +- Global flat을 기대한 기존 호출자는 `.pe_index` 접근 시 **즉시 + `AttributeError`** → 반드시 구조적 좌표로 migration. +- Flat integer key가 필요한 국소 문맥 (예: 내부 dict lookup)은 호출자가 + 명시적으로 `spec.sip * N_CUBES * N_PE + spec.cube * N_PE + spec.pe`를 계산. + +**Property 제거 정당화**: KernBench는 사내 프로젝트로 call site가 한정되어 +있음. Silent drift 위험 (의미만 바뀌고 타입은 같은 int) 대비 explicit breakage +(AttributeError)가 훨씬 안전. + +### D3. `resolve_dp_policy`가 `target_sip`을 받아 structural 좌표 생성 + +ADR-0024 D11의 계약 구현. Post-hoc shifting 없음. + +```python +# src/kernbench/policy/placement/dp.py (after) + +@dataclass(frozen=True) +class _LocalPeShard: + """Internal — PE resolver의 반환. Cube 내 local PE 식별자 + payload.""" + local_pe: int # cube-local PE index (0..num_pe-1) + offset_bytes: int + nbytes: int + + +def resolve_dp_policy( + policy: DPPolicy, + *, + shape: tuple[int, int], + itemsize: int, + num_pe: int, + num_cubes: int = 1, + target_sip: int, # NEW — 어느 SIP에 배치할지 명시 +) -> list[ShardSpec]: + """2-level resolution (cube × PE) on a specified SIP. + + Returns ShardSpecs with structural coords (sip=target_sip, cube, pe). + No SIP-level split — DPPolicy is intra-device only. + """ + resolver = _PE_RESOLVERS[policy.pe] + all_shards: list[ShardSpec] = [] + + # Level 1: cube within SIP + cube_splits = _split_shape(policy.cube, shape, num_cubes, itemsize) + + for cube_id, (cube_shape, cube_offset) in enumerate(cube_splits): + # Level 2: PE within cube — resolver returns _LocalPeShard (local_pe) + local_shards = resolver(shape=cube_shape, itemsize=itemsize, + num_pe=num_pe) + + for ls in local_shards: + all_shards.append(ShardSpec( + sip=target_sip, # from caller (current_device) + cube=cube_id, # local within SIP + pe=ls.local_pe, # local within cube (explicit name) + offset_bytes=cube_offset + ls.offset_bytes, + nbytes=ls.nbytes, + )) + + return all_shards +``` + +**내부 resolver** (`column_wise`, `row_wise`, `replicate`)는 `_LocalPeShard` +리스트 반환 — `local_pe` 필드명으로 **"cube-local PE identifier"임이 명시적**. +과거 `ShardSpec.pe_index`와 이름이 혼동되던 문제 해소. + +**이름 규약 정리** (전체 ADR): +- `ShardSpec.pe`: 최종 외부 API — cube-local PE (structural coord) +- `_LocalPeShard.local_pe`: 내부 resolver 단계의 동일 의미 +- `pe_index`: **제거**. 외부/내부 어디에도 남기지 않는다 (silent drift 차단의 + 부가 효과: 이름 재등장 없음). + +### D4. `_create_tensor` — 구조적 좌표로 직접 placement + +ADR-0024 D11 연속선. Post-hoc shifting 제거, 구조적 좌표를 `resolve_dp_policy` +호출 시점에 직접 지정. + +```python +# context.py _create_tensor (after) +current_sip = self.ahbm.current_device() +if current_sip is None: + # Single-driver fallback (ADR-0024 D9와 일관). + # Launcher 기반 코드가 set_device()를 빼먹으면 조용히 SIP 0에 박히는 + # 문제가 있음 → debug mode에서 경고. + if os.environ.get("KERNBENCH_DEBUG"): + import warnings + warnings.warn( + "torch.ahbm.current_device() is None; defaulting to SIP 0. " + "If this is a multi-rank launcher context, you likely forgot " + "torch.ahbm.set_device(rank) inside the worker.", + stacklevel=2, + ) + current_sip = 0 + +placement = resolve_dp_policy( + dp, + shape=shape_2d, + itemsize=itemsize, + num_pe=eff_num_pe, + num_cubes=eff_num_cubes, + target_sip=current_sip, # ← 구조적 좌표 일차 지정 +) + +# placement의 각 ShardSpec은 이미 (sip=current_sip, cube=local, pe=local) 포함. +# 과거의 post-hoc shifting 블록은 완전히 제거. +``` + +**모든** 텐서가 current device SIP에 배치됨. Multi-SIP 텐서를 만들고 싶으면 +ADR-0027의 TP primitive 사용. + +**Single-driver fallback의 trade-off**: set_device 없는 호출에서 SIP 0으로 +default는 기존 single-driver 테스트 호환을 위해 유지. `KERNBENCH_DEBUG=1` +환경에서는 launcher 컨텍스트의 실수로 set_device 누락 시 조용히 잘못된 SIP에 +배치되는 것을 감지할 수 있도록 warning. + +### D5. Downstream — allocator lookup은 구조적 tuple key로 + +기존 `deploy_tensor` (`src/kernbench/runtime_api/tensor.py`): + +```python +for spec in placement: + alloc = allocators[spec.pe_index] # ← AttributeError (property 제거됨) +``` + +`pe_index`가 없어졌으므로 구조적 좌표로 **강제** migration: + +```python +for spec in placement: + alloc = allocators[(spec.sip, spec.cube, spec.pe)] +``` + +`_ensure_allocators`의 dict population도 tuple key로: + +```python +# context.py _ensure_allocators (after) +for sip_id in sip_range: + for cube_id in range(cubes_per_sip): + for pe_id in range(pes_per_cube): + self._allocators[(sip_id, cube_id, pe_id)] = PEMemAllocator( + rack_id=0, sip_id=sip_id, cube_id=cube_id, pe_id=pe_id, cfg=cfg, + ) +``` + +`_free_tensor`도 동일: 기존 `flat_idx = sip * ... + cube * ... + pe` 계산 +블록 제거, `(shard.sip, shard.cube, shard.pe)` 직접 사용. + +**Tuple vs dataclass `PEIdentity`**: Tuple이 단순하고 hashable로 바로 써서 +권고. `PEIdentity` 값객체는 명시적 타입 장점은 있지만 boilerplate가 크고 현재 +allocator dict의 유일한 key라 오버엔지니어링. Tuple 유지. + +### D6. Migration — 기존 call site + +**(A) `DPPolicy(sip=..., num_sips=..., ...)` 사용하던 코드**: + +- `DPPolicy(sip="column_wise", cube=..., pe=...)` 패턴 → **해당 bench를 ADR-0024 + launcher로 재작성**. worker가 `set_device(rank)`로 SIP 선택, DPPolicy는 + cube/PE만. +- `DPPolicy(sip="replicate", num_sips=1, ...)` 패턴 → `DPPolicy(cube=..., pe=...)`로 + 축소 (필드가 사라지니 자연스럽게). + +**(B) `dp.sip`, `dp.num_sips` 읽던 코드**: + +- 제거. `launch()`의 `_compute_local_shape`에서 `dp.sip` 분기 삭제. +- `pe_cpu.py`가 `dp.sip`을 참조하던 곳도 정리. + +**(C) `ShardSpec.pe_index`를 사용하던 코드 — 전부 수정 필요**: + +- `.pe_index` 접근은 이제 `AttributeError` 발생 → 모든 call site 수정 필수. +- Allocator lookup: `allocators[spec.pe_index]` → + `allocators[(spec.sip, spec.cube, spec.pe)]` +- Flat integer가 꼭 필요한 국소 문맥: `spec.sip * N_CUBES * N_PE + spec.cube * + N_PE + spec.pe` 명시적 계산. **국소 변수로만 사용하고 공개 API에 노출하지 + 않는다**. + +**구현 착수 전 grep audit 체크리스트**: + +1. **Property 참조**: + - `\.pe_index\b` — 필드/property 접근 모두 (regex) + - `pe_index=` — 생성 시점의 키워드 인자 + - `pe_index:` — dataclass 필드 선언 +2. **Allocator / dict indexing**: + - `allocators\[` — dict lookup 패턴. `allocators[spec.pe_index]` 같은 + 것이 걸리는지 + - `_allocators\[` — 같은 패턴 (prefix _) +3. **Flat index 수동 계산 블록**: + - `flat_idx =` + - `pe_index =` (좌변) + - `* pes_per_cube +` (전형적 flat 계산 패턴) + - `* self._num_cubes \* self._pes_per_cube` (global flat 계산) +4. **Serialization / logging**: + - `asdict(.*shard` — dataclass 직렬화 시 `pe_index` 자동 포함 여부 + - `repr(.*ShardSpec` — 로그 포맷에서 의존하는지 + - JSON/YAML 저장 포맷에서 `pe_index` 키 사용 여부 +5. **Tests asserting integer PE identity**: + - `assert .*pe_index` — 정수 동일성 주장 + - `spec.pe_index ==` — 비교 (SIP-local 의미로 변하면 테스트가 깨질 수 있음) + +각 match마다 "이 호출자가 global flat / SIP-local / 내부 lookup 중 무엇을 +기대했나"를 판단한 뒤 구조적 좌표로 교체. + +**(D) `test_sip_parallel.py`**: + +- 이름 유지, 내용은 ADR-0024의 multi-greenlet launcher 기반 재작성. +- "SIP 병렬성 = rank 별 worker × 각자 DPPolicy" 로 검증. + +**(E) `test_va_offset.py`, `benches/va_offset_verify.py`**: + +- `num_sips=1`만 쓰는 경우가 대부분. 단순히 필드 제거. +- SIP offset 테스트가 핵심이면 `set_device(rank)` + 구조적 좌표 관찰로 이식. + +### D7. 하위 호환 — 불가 (cleanup ADR) + +이 ADR은 **breaking change**. + +1. `DPPolicy(sip=...)` 또는 `DPPolicy(num_sips=...)` 호출 → `TypeError` +2. `ShardSpec.pe_index` 접근 → `AttributeError` + +모두 **즉시 명시적 breakage**. Deprecation warning / fallback 경로 없음. +KernBench는 사내 프로젝트로 call site가 한정되어 있어 한 번에 migration. + +**Silent drift 차단**이 property 완전 제거의 주된 이점: global flat을 기대한 +코드가 SIP-local 결과를 받아 조용히 잘못된 인덱싱을 할 가능성 제거. + +### D8. 문서 업데이트 + +- `ADR-0008` (tensor deploy) — DPPolicy 의미 갱신 note, ShardSpec 구조적 좌표 + 전환 명시 +- DPPolicy docstring에 "intra-device only" 명시 (D1 코드 스니펫의 docstring) +- ShardSpec docstring에 **structural coordinates `(sip, cube, pe)`를 직접 + 사용하며, `pe_index`는 더 이상 제공되지 않음**을 명시 (D2) +- `docs/ccl-author-guide` 등 튜토리얼에서 `sip=...` 예시 제거 + +--- + +## Dependencies + +- **ADR-0024** (launcher): `set_device(rank)` 및 current-device scoping이 + SIP 배치 메커니즘 제공. 본 ADR은 그 위에 서서 DPPolicy를 순수 intra-device로 + 좁힘. +- **ADR-0027** (Megatron TP): 다중 SIP에 걸친 텐서가 필요한 경우의 대안 경로. + 이 ADR 적용 후 multi-SIP use case는 ADR-0027로 이관. + +--- + +## Non-goals + +- **`DPPolicy.cube` / `pe` 재설계**: 기존 replicate/column_wise/row_wise 의미 + 유지. +- **Tiling 정책 통합**: `tiled_column_major` / `tiled_row_major`는 그대로. +- **Multi-device 텐서 추상화 신규**: DTensor-like는 ADR-0028. + +--- + +## Open questions + +- **`_create_tensor`의 current_sip 기본값**: set_device 없는 호출에서 rank=0 + (SIP 0)로 fallback할지, 아니면 error 낼지. 권고는 fallback (기존 single-driver + 테스트와의 호환). +- **`test_sip_parallel.py` 재작성 범위**: 기존 단위 테스트의 의도를 유지하며 + launcher 기반으로 옮기려면 추가 fixture 필요. 별도 작업으로 scope. +- **`DPPolicy`의 `num_sips=None` 의미**: 필드가 없어지면 `num_sips` 개념 자체가 + 사라짐. Multi-SIP을 표현하고 싶으면 ADR-0027의 TP primitive를 쓰라는 것이 + 명시적 답. + +**Resolved (이전 rev에서 open이었던 것들)**: +- ~~`ShardSpec.pe_index` property 존치 여부~~ → **완전 제거** (D2) +- ~~`_ensure_allocators` dict key 형식~~ → **tuple `(sip, cube, pe)`** (D5) + +--- + +## Test strategy + +### T1. 단위 테스트 갱신 + +- `tests/test_tensor.py`, `tests/test_ipcq_types.py`, `tests/test_runtime_api_tensor.py` + — DPPolicy 생성자 인자 정리, ShardSpec 구조적 좌표 검증 +- `tests/test_va_offset.py` — `num_sips=1` 제거 후 동작 유지 + +### T2. `resolve_dp_policy` 구조적 좌표 반환 + +`tests/test_dp_policy.py` (new 또는 확장): +- `resolve_dp_policy(dp, ..., target_sip=1)` 결과의 모든 ShardSpec이 `sip=1` +- 각 spec의 `(cube, pe)`가 local (0..num_cubes-1, 0..num_pe-1) +- 같은 topology에서 `target_sip=0`과 `target_sip=1` 결과가 sip 필드만 다름 + +### T3. `test_sip_parallel.py` 재작성 + +SIP 병렬성 검증을 launcher 기반으로: + +```python +def test_sip_parallel_via_launcher(topology): + ... + def worker(rank, ws, torch): + torch.ahbm.set_device(rank) + t = torch.zeros((1, 128), dtype="f16", + dp=DPPolicy(cube="column_wise", pe="column_wise")) + # verify shard.sip == rank (structural coord) + + spawn(worker, nprocs=n_sips, ...) +``` + +### T4. Allocator key migration + +`tests/test_allocator_structural_key.py` (new 또는 기존 확장): +- `PEMemAllocator` dict이 `(sip, cube, pe)` tuple key로 작동 +- `deploy_tensor`가 구조적 좌표로 allocator lookup +- `_free_tensor`도 동일 + +### T5. E2E 회귀 + +ADR-0024의 `test_ccl_allreduce_matrix.py` 그대로 통과. + +### T6. 오류 검증 + +- `DPPolicy(sip="column_wise")` 호출 → `TypeError`. 테스트로 명시. +- `DPPolicy(num_sips=2)` 호출 → `TypeError`. +- `spec.pe_index` 접근 → `AttributeError` (property 완전 제거 검증). + +--- + +## Consequences + +### Positive + +- **개념 분리 명확**: DPPolicy = intra-device, TP = inter-device. +- **API 단순화**: DPPolicy 생성자 필드 ~33% 축소. +- **Structural 좌표 일관성**: ShardSpec이 `(sip, cube, pe)` 튜플로 표현 → + abstraction leakage 해소 (ADR-0024 D11 계약 충족). +- **`pe_index` 의미 명확**: SIP-local이 단일 해석. Global flat이 필요하면 명시. +- **Launcher 모델 일관성**: ADR-0024의 "1 worker per SIP" 모델이 유일한 SIP + 경계 제어 메커니즘. + +### Negative + +- **Breaking change (explicit)**: `DPPolicy(sip=...)` → `TypeError`, + `spec.pe_index` → `AttributeError`. 모든 호출자 한 번에 수정 필요. +- **ShardSpec schema 변경**: `pe_index` 단일 필드 → `sip`/`cube`/`pe` 세 필드. + Downstream (`deploy_tensor`, `_free_tensor`, `_ensure_allocators`, + `allocators` dict key 등) 연쇄 수정. +- **Silent drift 없음**: property 완전 제거로 runtime에서 즉시 실패 → + migration leakage 원천 차단. (Negative가 아니라 explicit tradeoff) +- `test_sip_parallel.py` 재작성 비용. + +### Neutral + +- 기존 `cube` / `pe` 필드 의미 불변. + +--- + +## Affected files + +| File | Change | +|------|--------| +| `src/kernbench/policy/placement/dp.py` | D1: `sip`/`num_sips` 제거 / D2: `ShardSpec`에 `sip`/`cube`/`pe` structural fields 추가, **`pe_index` property 제거** / D3: `resolve_dp_policy`에 `target_sip`, SIP-level 루프 제거 / 내부 resolver가 반환하는 shard 타입 이름도 `local_pe`로 명확화 (이름 충돌 방지) | +| `src/kernbench/runtime_api/context.py` | D4: `_create_tensor` `target_sip` 전달 / D5: `_ensure_allocators` dict key → `(sip, cube, pe)` tuple / `launch`의 `dp.sip` 분기 제거 | +| `src/kernbench/runtime_api/tensor.py` | D5: `deploy_tensor`가 구조적 좌표로 allocator lookup | +| `src/kernbench/components/builtin/pe_cpu.py` | D6: `dp.sip` 참조 제거 | +| `src/kernbench/components/legacy/builtin/pe_cpu.py` | D6: 동일 | +| `benches/ccl_allreduce.py` | ADR-0024 scope에서 이미 처리 | +| `benches/va_offset_verify.py` | D6: `num_sips=1` 제거 | +| `tests/test_runtime_api_tensor.py` | D6 | +| `tests/test_va_offset.py` | D6 | +| `tests/test_tensor.py`, `test_ipcq_types.py` | D6 | +| `tests/test_sip_parallel.py` | T3: launcher 기반 재작성 | +| `tests/test_dp_policy.py` (new 또는 확장) | T2 | +| `tests/test_allocator_structural_key.py` (new) | T4 | diff --git a/docs/adr/ADR-0027-megatron-tp.md b/docs/adr/ADR-0027-megatron-tp.md new file mode 100644 index 0000000..a6c8ab3 --- /dev/null +++ b/docs/adr/ADR-0027-megatron-tp.md @@ -0,0 +1,341 @@ +# ADR-0027: Megatron-style Tensor Parallelism API + +## Status + +Proposed + +## Context + +### 목표 + +SIP 간 tensor parallelism(TP)을 **Megatron-LM 스타일의 명시적 parallel layer** +API로 지원한다. DTensor 같은 선언적 추상화는 별도 ADR(0028) future work. + +Megatron-style을 선택한 이유 (사용자 지시): +- TP는 일반적으로 **model의 특정 layer 경계**에서 발생. 명시적 primitive가 + mental model에 자연스러움. +- NVIDIA Megatron / DeepSpeed가 확립한 인더스트리 표준 패턴. +- DTensor는 선언적이라 **디자인 공간이 더 크다** → 단계적으로 접근. + +### 현재 상태 + +- KernBench는 TP가 없음. 기존 `DPPolicy.sip="column_wise"` 경로가 "SIP 간 + column sharding"을 흉내 냈으나 DP와 TP가 섞인 상태 (ADR-0026에서 정리). +- ADR-0024가 launcher 인프라 (rank = SIP, `set_device`, greenlet-local) 제공. +- 이 인프라 위에 TP primitive를 얹는다. + +### TP primitive 스펙 (Megatron-LM 참조) + +- **ColumnParallelLinear**: weight의 **column** 축을 TP ranks에 분산. 입력 + full-replicated, 출력 column-sharded. 후속에 row-parallel이 올 때 all-reduce + 없음. +- **RowParallelLinear**: weight의 **row** 축을 TP ranks에 분산. 입력이 이미 + column-sharded (ColumnParallel의 출력). forward 끝에 **all-reduce** 필요. +- **VocabParallelEmbedding**: embedding을 vocab 축에 분산. forward 끝에 all-reduce. +- **`copy_to_tp_region`**, **`reduce_from_tp_region`**, **`scatter_to_tp_region`**, + **`gather_from_tp_region`** — 기본 primitive (`identity` forward, `all-reduce` + backward 등). + +### 풀어야 할 문제 + +1. **Per-rank weight 분산 표현**: 각 worker(rank)가 weight tensor의 자기 + slice를 소유. ADR-0024의 `set_device(rank)` + ADR-0026의 intra-device + DPPolicy로 자연스러운 표현. + +2. **Forward / backward activation 흐름**: 현재 KernBench는 backward가 없음 + (simulation 목적). 본 ADR은 **forward만** 우선 지원. Training simulation이 + 추가되면 확장. + +3. **Collective 호출 지점**: RowParallelLinear가 forward 끝에 `all_reduce`를 + 호출. ADR-0024의 multi-greenlet 구조에서 자연스럽게 동작 (각 rank가 동시에 + 호출). + +4. **TP group 개념**: Megatron은 일반적으로 data_parallel × tensor_parallel × + pipeline_parallel group을 교차 사용. 초기 scope는 **TP group = 전체 SIP**로 + 단순화. Mixed DP+TP는 future. + +--- + +## Decision + +### D1. 새 패키지 `kernbench.tp` + +``` +src/kernbench/tp/ + __init__.py — public API re-exports + parallel_state.py — TP group 관리 (현재 단일 global group) + layers.py — ColumnParallelLinear, RowParallelLinear, VocabParallelEmbedding + primitives.py — copy/reduce/scatter/gather_to/from_tp_region + mappings.py — identity/all_reduce forward, all_reduce/identity backward (stub) +``` + +### D2. `parallel_state` — TP group + +```python +# parallel_state.py +_TP_WORLD_SIZE = None +_TP_RANK = None # greenlet-local via dist.get_rank() + +def initialize_model_parallel(tensor_model_parallel_size: int) -> None: + """Initialize TP group. Must be called after dist.init_process_group.""" + global _TP_WORLD_SIZE + from kernbench.runtime_api.distributed import get_dist + dist = get_dist() + total = dist.get_world_size() + if tensor_model_parallel_size != total: + raise NotImplementedError( + "Only TP == world_size supported in initial scope" + ) + _TP_WORLD_SIZE = tensor_model_parallel_size + +def get_tensor_model_parallel_world_size() -> int: + return _TP_WORLD_SIZE + +def get_tensor_model_parallel_rank() -> int: + from kernbench.runtime_api.distributed import get_dist + return get_dist().get_rank() # ADR-0024의 greenlet-local rank +``` + +초기 scope: **TP 사이즈 = world_size = topology SIP 수**. Pure TP 모델. + +### D3. `ColumnParallelLinear` + +```python +# layers.py +class ColumnParallelLinear: + """Weight의 K(out_features) 축을 TP rank에 분산. + + forward(x): + x: (M, N) — full-replicated across ranks + W_k: (N, K / world_size) — rank-local slice + y_k = x @ W_k → (M, K / world_size) — rank-local output + + 출력은 sharded. 후속 RowParallelLinear가 기대하는 입력 형태. + """ + + def __init__(self, in_features: int, out_features: int, bias: bool = False, + dtype: str = "f16", torch=None): + ws = get_tensor_model_parallel_world_size() + assert out_features % ws == 0 + k_local = out_features // ws + # 각 rank가 자기 slice 소유 (ADR-0024 set_device + ADR-0026 DPPolicy) + self.weight = torch.zeros( + (in_features, k_local), dtype=dtype, + dp=DPPolicy(cube="column_wise", pe="column_wise"), + name="col_parallel_w", + ) + # init with something sensible — TODO + if bias: + self.bias = torch.zeros((k_local,), ...) + else: + self.bias = None + + def forward(self, x): + # x는 full-replicated (caller 보장). 단순 local matmul. + y = torch.matmul(x, self.weight) + if self.bias is not None: + y = y + self.bias + return y +``` + +### D4. `RowParallelLinear` + +```python +class RowParallelLinear: + """Weight의 N(in_features) 축을 TP rank에 분산. + + forward(x): + x: (M, N / world_size) — rank-local slice (ColumnParallel의 출력) + W_k: (N / world_size, K) — rank-local slice + y_k = x @ W_k → (M, K) — partial sum on each rank + y = all_reduce(y_k, op="sum") → (M, K) on every rank + """ + + def __init__(self, in_features: int, out_features: int, bias: bool = False, + dtype: str = "f16", torch=None): + ws = get_tensor_model_parallel_world_size() + assert in_features % ws == 0 + n_local = in_features // ws + self.weight = torch.zeros( + (n_local, out_features), dtype=dtype, + dp=DPPolicy(cube="column_wise", pe="column_wise"), + name="row_parallel_w", + ) + # bias는 rank 0에만 (Megatron convention) + self.bias = torch.zeros(...) if bias else None + self._torch = torch + + def forward(self, x): + y_partial = torch.matmul(x, self.weight) + # Final all-reduce sums partial products across ranks + self._torch.distributed.all_reduce(y_partial, op="sum") + if self.bias is not None: + # bias는 reduce 이후에만 추가 (rank 0 보유) + rank = get_tensor_model_parallel_rank() + if rank == 0: + y_partial = y_partial + self.bias + return y_partial +``` + +### D5. Primitive 함수 + +```python +# primitives.py +def copy_to_tp_region(x): + """Forward: identity. Backward: all-reduce. (Training 추가 시 구현).""" + return x + +def reduce_from_tp_region(x): + """Forward: all-reduce. Backward: identity.""" + torch.distributed.all_reduce(x, op="sum") + return x + +def scatter_to_tp_region(x): + """x를 K 축으로 scatter. Forward: split. Backward: all-gather.""" + # 초기 scope에서는 사용자가 이미 sharded tensor를 만들었다고 가정 → + # no-op 또는 metadata 추가 + raise NotImplementedError("Phase 2 feature") + +def gather_from_tp_region(x): + """x를 K 축으로 all-gather. Forward: all-gather. Backward: split.""" + raise NotImplementedError("all-gather kernel이 먼저 필요 (future)") +``` + +### D6. 샘플 bench — 2-layer MLP with TP + +```python +# benches/tp_mlp.py (새 파일) +def worker(rank, world_size, torch): + torch.cuda.set_device(rank) + tp.initialize_model_parallel(world_size) + + B, D_in, D_hidden, D_out = 1, 512, 2048, 512 + fc1 = tp.ColumnParallelLinear(D_in, D_hidden, torch=torch) + fc2 = tp.RowParallelLinear(D_hidden, D_out, torch=torch) + + x = torch.zeros((B, D_in), dtype="f16", + dp=DPPolicy(cube="column_wise", pe="column_wise"), + name="x") + # ... init x ... + + h = fc1.forward(x) # column-sharded output + y = fc2.forward(h) # all-reduced output, full on every rank + + # verify... + +def run(torch): + torch.distributed.init_process_group(backend="ahbm") + torch.distributed.spawn(worker, nprocs=torch.distributed.get_world_size(), + args=(...,)) +``` + +### D7. Non-functional — training 미지원 + +본 ADR은 **inference/forward only**. Backward / gradient / optimizer는 future. +기존 KernBench가 training이 아니므로 자연스러움. + +### D8. 초기 scope 제약 + +- TP 사이즈 = world_size (mixed DP+TP 없음) +- `scatter_to_tp_region`, `gather_from_tp_region`은 unimplemented (별도 kernel + 필요) +- Weight init은 단순 zero (적절한 init은 future) +- Pipeline parallelism은 scope 밖 + +### D9. `distributed.all_reduce` 기반 + +RowParallelLinear의 모든 collective는 ADR-0024의 `dist.all_reduce`를 사용. +별도 TP-전용 collective 엔진 불필요. + +--- + +## Dependencies + +- **ADR-0024** (launcher): rank = SIP, greenlet-local rank, `dist.all_reduce`, + `torch.cuda.set_device(rank)`, `spawn_workers` 제공. +- **ADR-0026** (DPPolicy intra-device): weight tensor의 per-rank slice 표현. +- **ADR-0023 / ADR-0025** (IPCQ): `dist.all_reduce` 구현의 기반. + +--- + +## Non-goals + +- **Backward pass / training**: inference only. Training simulation은 별도 ADR. +- **Mixed parallelism (DP + TP + PP)**: 초기엔 pure TP만. +- **Weight init schemes**: 단순 zero / debug pattern. 실제 training init는 future. +- **Fused ops**: Megatron의 fused matmul+bias+gelu 등은 KernBench kernel 수준 + 문제. 본 ADR은 host-side API만. +- **DTensor 통합**: ADR-0028 future. + +--- + +## Open questions + +- **`initialize_model_parallel` 위치**: `kernbench.tp.initialize_model_parallel` + vs `torch.distributed.init_tp(...)` 확장. real PyTorch는 `torch.distributed. + init_device_mesh` 등을 권장. 우리는 당분간 TP-전용 모듈. +- **Weight의 DP 전략**: 본 ADR은 `DPPolicy(cube="column_wise", pe="column_wise")` + 를 가정. Intra-SIP DP를 다르게 주면? 성능 벤치마크로 결정. +- **Bias 배치 정책**: Megatron은 bias를 split하지 않음. RowParallelLinear는 + rank 0에만. 이게 항상 맞는가? 대안: replicate across ranks. +- **`VocabParallelEmbedding`**: 처음 몇 벤치엔 불필요할 수도. 샘플 구현은 넣되 + scope에서 제외할 수도. + +--- + +## Test strategy + +### T1. Unit — `tests/test_tp_layers.py` (신규) + +- `ColumnParallelLinear` forward: rank별 weight slice, 출력이 `(M, K / ws)`. +- `RowParallelLinear` forward: 입력이 sharded, all_reduce 후 `(M, K)` 일치. +- `VocabParallelEmbedding` forward (if implemented). +- `parallel_state` 초기화 / rank 조회. + +### T2. E2E — `tests/test_tp_mlp.py` (신규) + +- 2-layer MLP (ColumnParallel → RowParallel) forward가 single-driver reference + 와 일치 (numerical check, rtol/atol). +- ws = SIP count (current: 2). + +### T3. 회귀 + +- ADR-0024의 `test_ccl_allreduce_matrix` 그대로 통과 (TP가 호출하는 + `dist.all_reduce`의 기반). + +--- + +## Consequences + +### Positive + +- **Megatron 코드 이식 용이**: real training code와 API 일치. +- **TP 벤치마크 가능**: scaling, communication-compute overlap 등 HW 특성 + 연구. +- **DPPolicy 의미 명확화** (ADR-0026과 시너지). + +### Negative + +- 새 모듈 (`kernbench.tp`) 유지보수 비용. +- 초기 scope가 제한적 (pure TP only). + +### Neutral + +- ADR-0024/0026 기반 위에 순수한 상위 레이어 추가. Hardware simulation + stack에 영향 없음. + +--- + +## Affected files + +| File | Change | +|------|--------| +| `src/kernbench/tp/__init__.py` | 신규: public API re-export | +| `src/kernbench/tp/parallel_state.py` | 신규: D2 | +| `src/kernbench/tp/layers.py` | 신규: D3/D4 | +| `src/kernbench/tp/primitives.py` | 신규: D5 | +| `src/kernbench/tp/mappings.py` | 신규 stub (backward TODO) | +| `benches/tp_mlp.py` | 신규: D6 샘플 | +| `tests/test_tp_layers.py` | 신규: T1 | +| `tests/test_tp_mlp.py` | 신규: T2 | +| `docs/tp-author-guide.md` | 신규 (선택): 사용자 가이드 | diff --git a/docs/adr/ADR-0028-dtensor-support.md b/docs/adr/ADR-0028-dtensor-support.md new file mode 100644 index 0000000..7df566f --- /dev/null +++ b/docs/adr/ADR-0028-dtensor-support.md @@ -0,0 +1,171 @@ +# ADR-0028: DTensor Support — 선언적 분산 텐서 (Stub / Future) + +## Status + +Stub (Future Work) + +## Context + +### 목표 + +**선언적 분산 텐서 추상화**(PyTorch 2.x `DTensor` 스타일)를 KernBench에 +도입하기 위한 **디자인 공간 preliminary exploration**. 본 ADR은 **구현 계획이 +아닌 future 작업의 파일 플레이스홀더 + 초기 질문 목록**이다. + +### Megatron-style TP와의 차이 (Why DTensor) + +| 관점 | Megatron (ADR-0027) | DTensor (이 ADR) | +|---|---|---| +| 표현 | 명시적 parallel layer | 텐서 + placement spec | +| 호출 형태 | `ColumnParallelLinear(...)` | `distribute_tensor(x, mesh, [Shard(1)])` | +| Collective 삽입 | 레이어 내부 명시 | 연산 dispatch가 자동 | +| Learning curve | 낮음 (명시적) | 중~높음 (선언적 의미 이해) | +| 유연성 | 레이어 단위로 고정 | 레이어 경계 무관, 어디서나 | +| KernBench에 선행 필요한 것 | launcher (ADR-0024) + TP (0027) | 그 + operator dispatch overhaul | + +DTensor는 operator-level에서 "텐서의 placement를 보고 자동으로 collective +삽입". KernBench가 이를 지원하려면 **operator dispatch layer에 placement-aware +rewriting**이 들어가야 한다. 이는 비-trivial. + +### 현재 상태 + +- KernBench는 operator dispatch 레이어가 없음 (`torch.matmul`은 없음; kernel + launch로 대체). +- DPPolicy는 정적 placement metadata를 보유 (ADR-0026 후: intra-device only). +- ADR-0024 launcher가 rank / device 개념 제공. +- Megatron-style TP (ADR-0027)가 명시적 대안으로 기능할 것. + +--- + +## Preliminary decision space + +### DQ1. PyTorch DTensor API 수용 범위 + +- `DeviceMesh`: rank들의 논리적 grid. +- `Placements`: `Shard(dim)`, `Replicate()`, `Partial(reduce_op)`. +- `distribute_tensor(tensor, device_mesh, placements)`: local tensor → DTensor. +- Redistribute: `dt.redistribute(new_placements)`로 collective 자동 삽입. +- Operator forward: `dt @ dt`, `dt + dt` 등 → 적절한 collective 자동 dispatch. + +KernBench가 어느 수준까지 지원할지 결정 필요. 최소: `distribute_tensor` + +`redistribute`. 최대: 모든 operator overloading. + +### DQ2. Operator dispatch 레이어 + +KernBench에서 `dt @ dt`를 정의하려면 Tensor의 `__matmul__`이 placement를 +보고 적절한 action 수행: + +- 둘 다 replicated → local matmul +- A column-sharded, B row-sharded → local matmul + all-reduce (RowParallel) +- A replicated, B column-sharded → local matmul (ColumnParallel) +- etc. + +이는 Megatron-style의 **자동화된 버전**. Kernel은 기존 matmul kernel 사용. + +### DQ3. DeviceMesh와 기존 topology + +KernBench topology는 이미 SIP/cube/PE 계층. DTensor의 DeviceMesh는 추상 +`(tp_size, dp_size, ...)` grid. 매핑: + +- 1D mesh of size = SIP count → rank = SIP +- 2D mesh (tp × dp) → SIP을 그룹 분할 (pure TP 대신 mixed parallelism) + +초기엔 1D mesh만, DP × TP 2D는 future. + +### DQ4. Placement의 intra-device (DP) 통합 + +KernBench 특이점: 한 rank 내부에서 DPPolicy로 cube/PE에 분산. DTensor는 +device 내부를 보지 않음. 통합: + +- DTensor placement = rank (SIP) 간 분산 +- 각 rank의 local tensor는 여전히 DPPolicy로 cube/PE 배치 +- → DTensor wrapper가 local tensor의 DPPolicy도 보관 + +### DQ5. Collective 자동 삽입 지점 + +`redistribute` 또는 operator forward 시. ADR-0024의 submit+yield+wait 패턴을 +자동으로 호출하는 형태. `_launch_submit` 내부화. + +### DQ6. Autograd + +DTensor는 autograd와 상호작용 (backward에서 reverse collective). KernBench가 +backward 지원하기 전까지는 **forward-only DTensor**. + +--- + +## Open questions (to resolve before real design) + +1. **우선순위**: Megatron-style(ADR-0027)이 먼저 안착한 후 DTensor를 위에 + 얹는가, 아니면 공통 lower-layer를 먼저 설계하는가? +2. **호환성 목표**: PyTorch DTensor API와 몇 %까지 일치시키는가? 독자 API vs + 거의 동일? +3. **Operator dispatch**: KernBench `Tensor` 클래스에 `__matmul__` 등 연산자 + overloading을 도입하는가? (현재는 kernel launch만) +4. **Redistribute 정책**: `Shard(0) → Replicate()` 변환 시 어떤 collective + 사용? `all_gather`가 없으면 구현 전까지 제약. +5. **Mesh × DPPolicy interaction**: 하나의 DTensor가 2개 layer 분산을 갖는 + 경우의 metadata 표현. +6. **Partial placement의 reduce 시점**: 자동 vs 명시 `redistribute` 호출. +7. **Bench authoring impact**: 기존 Megatron-style bench가 DTensor 기반으로 + 얼마나 쉽게 포팅되는가? + +--- + +## Non-goals (for future real ADR) + +- 이번 stub에서 API 확정. Future ADR에서 구체화. +- Implementation timeline. 이번 round에서는 **설계 공간 매핑만**. + +--- + +## Dependencies (potential) + +- **ADR-0024** (launcher): rank / device 기반 +- **ADR-0026** (DPPolicy cleanup): DTensor placement와의 분리 명확화 +- **ADR-0027** (Megatron TP): 실용 TP 패턴 경험을 DTensor 설계로 환류 +- **Future ADR** (operator dispatch layer): KernBench Tensor에 operator + overloading 도입 + +--- + +## Expected consequences (hypothetical) + +### Positive + +- PyTorch training code 이식이 **매우 쉬워짐** (DTensor 코드 그대로). +- TP + DP + 더 복잡한 parallelism을 **하나의 추상화**로 표현. +- Collective 삽입이 자동 → bench 작성자 부담 감소. + +### Negative + +- Operator dispatch layer 신규 구축 → 상당한 엔지니어링. +- Implicit behavior 증가 → 디버깅 / 성능 분석 복잡. +- KernBench의 "명시적 kernel launch" 철학과 tension. + +--- + +## Action + +- **Phase 1 (현재)**: 본 stub 유지. Megatron-style (ADR-0027) 먼저 구현 + + 사용 경험 축적. +- **Phase 2 (future)**: 사용 경험을 바탕으로 본 ADR을 real design으로 승격. + 위 Open questions에 대한 답을 제시. +- **Phase 3 (future)**: Implementation. + +현재 구현 작업은 **없음**. 디자인 공간 매핑만. + +--- + +## Affected files + +본 ADR은 **stub**이므로 production 변경 없음. Future real ADR에서 갱신될 +파일 후보: + +| File | 예상 변경 (future) | +|------|---| +| `src/kernbench/dtensor/__init__.py` | 신규 패키지 | +| `src/kernbench/dtensor/device_mesh.py` | DeviceMesh | +| `src/kernbench/dtensor/placements.py` | Shard/Replicate/Partial | +| `src/kernbench/dtensor/api.py` | distribute_tensor, redistribute | +| `src/kernbench/dtensor/ops/*.py` | Operator dispatch (matmul 등) | +| `src/kernbench/runtime_api/tensor.py` | Tensor에 `__matmul__` 등 추가 | diff --git a/docs/adr/ADR-0029-hierarchical-allreduce.md b/docs/adr/ADR-0029-hierarchical-allreduce.md new file mode 100644 index 0000000..604abeb --- /dev/null +++ b/docs/adr/ADR-0029-hierarchical-allreduce.md @@ -0,0 +1,419 @@ +# ADR-0029: Hierarchical All-Reduce — 3-level intra/inter-SIP 알고리즘 + +## Status + +Proposed + +## Context + +### 목표 + +"Rank = SIP" 모델 (ADR-0024) 위에서 각 SIP 내부의 모든 PE를 참여시키는 +**3-level 계층 all-reduce** 알고리즘을 정의한다. 각 레벨이 서로 다른 물리 +연결(intra-cube ring, inter-cube NoC, inter-SIP UCIe)을 활용해 대역폭을 +극대화한다. + +### 왜 hierarchical인가 + +단순 ring/mesh/tree all-reduce는 SIP당 1 PE만 참여 (ADR-0024의 `leader_only` +mapper). 이는 inter-SIP 단계는 잘 모델링하지만: + +- **Intra-SIP PE가 노는 시간이 발생**. Leader PE가 inter-SIP 통신 중이면 + 나머지 7 PE / 16 cube는 유휴. +- **Intra-cube/inter-cube 연결 대역폭 미활용**. Cube NoC는 매우 빠르지만 + 단일 leader 사용 시 이 자원이 노출되지 않음. +- **실제 NCCL 등은 hierarchical**: NVLink(intra-node) + InfiniBand(inter-node) + 의 bandwidth 차이를 활용. KernBench 토폴로지도 동일 구조 + (intra-cube / inter-cube / inter-SIP의 bandwidth·latency 차이). + +### 현재 상태 + +- `src/kernbench/ccl/algorithms/hierarchical_allreduce.py` 이미 존재 + (git log `10b33b4` — "Tensor indexing + hierarchical 3-level all-reduce + kernel"). PE-level로 world_size = total PE를 가정하는 옛 모델 기반 구현. +- ADR-0024에 의해 launcher는 rank = SIP로 바뀜. +- Hierarchical 커널은 **재해석 필요**: 이제 각 worker(1 per SIP)가 자기 SIP의 + 모든 PE를 참여시키고, kernel은 intra-cube → inter-cube → inter-SIP 순으로 + 3-level reduce + broadcast. + +### 풀어야 할 문제 + +1. **ADR-0024 framework 위에 hierarchical 알고리즘 맞추기** + - Mapper: `all_pes` (ADR-0024 D5 제공) + - Validator: `multi_pe_sip_local` (ADR-0024 D8 제공) + - Kernel: 기존 `hierarchical_allreduce.py` 수정 — rank 계산 방식을 SIP 내 + local (cube, pe)로 바꿈 +2. **PE-level neighbor graph 생성** + - Intra-cube: `(sip, cube, pe) ↔ (sip, cube, pe±1 mod N_PE)` (ring 내부) + - Inter-cube: `(sip, cube, 0) ↔ (sip, cube±1 mod N_CUBE, 0)` (cube leader만) + - Inter-SIP: `(sip, 0, 0) ↔ (sip±1 mod N_SIP, 0, 0)` (SIP leader만) +3. **Tensor layout**: 각 PE가 1 tile을 소유하고 시작 (`multi_pe_sip_local` + validator가 이 layout 강제). DPPolicy(cube="column_wise", + pe="column_wise")로 달성 가능. +4. **PE-level topology 표현 부족** (ADR-0024 D6의 "책임 분산" 이슈 구체화) + - Ring/mesh/tree 같은 단순 패턴은 rank-level topology_fn + mapper 조합으로 + 충분. + - Hierarchical은 레벨마다 다른 peer 매핑이라 `_build_pe_installs`에서 + multi-level 해석을 해야 함. + - 장기적으로는 topology 모듈이 PE-level을 직접 표현하는 편이 명시적. + +### Non-problem (이 ADR 밖) + +- Launcher / barrier / rank-to-SIP / mapper-validator registry → ADR-0024 +- IPCQ direction addressing → ADR-0025 +- DPPolicy 필드 정리 → ADR-0026 +- Megatron TP → ADR-0027 + +--- + +## Decision + +### D1. 알고리즘 구조 — 3-level reduce + 역순 broadcast + +``` +Level 1 (intra-cube, E/W ring): + 각 cube의 N_PE개 PE가 bidirectional ring reduce → cube 내 PE 0에 부분합 집중 +Level 2 (inter-cube within SIP, N/S ring, PE 0만 참여): + N_CUBE개 cube-leader가 bidirectional ring reduce → SIP 내 (cube 0, PE 0)에 + SIP 전체 부분합 집중 +Level 3 (inter-SIP, N_SIP peers, (cube 0, PE 0)만 참여): + Ring 또는 pair exchange로 전역 합산 완료 +Broadcast: + 역순 — Level 3 결과를 (cube 0, PE 0)에서 SIP 내 모든 cube-leader로, 다시 + 각 cube 내 모든 PE로 전파 +``` + +세부는 기존 `hierarchical_allreduce.py`의 커널 구현과 일치. ADR-0024 이후 +변경점은 **rank 계산 방식**과 **n_elem 해석**뿐: + +- 기존 (rank=PE 모델): `rank = cube_id * pes_per_cube + local_pe`, `pe_addr = + t_ptr + rank * nbytes` +- 신규 (rank=SIP 모델): 커널은 SIP-local 좌표 `(cube_id, local_pe)`로만 동작. + 텐서의 per-PE slice는 backend가 per-PE `TensorArg`로 전달 (ADR-0024 D3). + 커널 내부 rank 계산 자체가 불필요해짐 — `tl.program_id(0/1)`로 충분. + +### D2. Framework integration — ADR-0024 infrastructure 재활용 + +`ccl.yaml`: + +```yaml +algorithms: + hierarchical_allreduce: + module: kernbench.ccl.algorithms.hierarchical_allreduce + topology: hierarchical_3level # NEW — D3 참고 + mapper: all_pes # ADR-0024 D5 built-in + validator: multi_pe_sip_local # ADR-0024 D8 built-in + buffer_kind: tcm + n_elem: 128 +``` + +Framework 관점에서 hierarchical은 **특별한 알고리즘이 아니라, 특정 +topology / mapper / validator 조합**. 본 ADR은 그 조합과 topology 패턴을 +정의. + +### D3. `hierarchical_3level` topology (신규) + +`kernbench/ccl/topologies.py`에 신규 추가: + +```python +def hierarchical_3level(rank: int, world_size: int, spec: dict) -> dict: + """3-level hierarchical neighbor pattern. + + Returns a nested structure describing intra-cube + inter-cube + inter-SIP + neighbors. Unlike ring_1d / mesh_2d which are rank → {dir: peer_rank}, + hierarchical is PE-level and requires spec for cube_mesh / pe_layout. + """ +``` + +반환 스키마 (초안): + +```python +{ + "intra_cube": { + # 각 cube 내 ring neighbors: (cube, pe) → {"E": (cube, pe_e), "W": (cube, pe_w)} + ... + }, + "inter_cube": { + # cube-leader 간 ring: (cube, 0) → {"N": (cube_n, 0), "S": (cube_s, 0)} + ... + }, + "inter_sip": { + # SIP-leader 간: rank → {"parent": peer_rank} (또는 ring 방식) + ... + }, +} +``` + +이 구조는 `_build_pe_installs`가 해석하여 각 PE의 neighbor table 엔트리 +(4-direction)에 대응시킨다. + +**Rank-level `topologies.py` 현 API와의 관계**: 기존 단순 패턴은 +`(rank → {dir: peer_rank})` 단일 레벨. Hierarchical은 multi-level이므로 +기존 API와 schema가 다름. `_resolve_topology`는 **알고리즘이 어떤 schema를 +쓰는지 선언**하고, builder가 그에 맞춰 해석하도록 확장 필요 (open question). + +### D4. PE-level neighbor graph — `_build_pe_installs` 확장 + +기존 (ring/mesh/tree): topology_fn이 반환한 `(rank → {dir: peer_rank})`를 +각 참여 PE에 그대로 매핑 (leader_only일 경우 peer PE도 leader). + +신규 (hierarchical): `hierarchical_3level`의 3단 구조를 per-PE neighbor +table로 펼침: + +```python +def _build_pe_installs_hierarchical(rank, world_size, sip, pes, topo, spec): + """Hierarchical 전용 PE neighbor table 빌더.""" + result = [] + for (cube, pe) in pes: + entries = [] + # Level 1: intra-cube ring (E/W) + for d, peer in topo["intra_cube"][(cube, pe)].items(): + entries.append(NeighborTableEntry(direction=d, ...)) + # Level 2: inter-cube ring (N/S) — cube leader (pe == 0)만 + if pe == 0: + for d, peer in topo["inter_cube"][(cube, 0)].items(): + entries.append(NeighborTableEntry(direction=d, ...)) + # Level 3: inter-SIP — SIP leader (cube == 0 and pe == 0)만 + if cube == 0 and pe == 0: + for d, peer_rank in topo["inter_sip"][rank].items(): + # peer_rank → peer SIP의 (0, 0) + entries.append(NeighborTableEntry( + direction=d, peer_sip=peer_rank, peer_cube=0, peer_pe=0, ...)) + result.append(PeInstallSpec(cube=cube, pe=pe, neighbors=tuple(entries))) + return tuple(result) +``` + +`build_install_plans`에서 algorithm_config의 `topology`에 따라 적절한 builder +선택 (기존 simple builder vs hierarchical builder). + +### D5. Kernel 재해석 — SIP-local 좌표로 + +`src/kernbench/ccl/algorithms/hierarchical_allreduce.py`를 ADR-0024 D3에 +맞춰 수정: + +```python +def kernel_args(*, n_elem: int, world_size: int, pes_per_cube: int, + cubes_per_sip: int, num_sips: int, **kw) -> tuple: + """world_size (= num_sips), pes_per_cube, cubes_per_sip를 스칼라로.""" + return (n_elem, pes_per_cube, cubes_per_sip, num_sips) + +def kernel(t_ptr, n_elem, pes_per_cube, cubes_per_sip, num_sips, tl): + """SIP-local 좌표 기반. + + 이전 (rank=PE 모델): + rank = cube_id * pes_per_cube + local_pe + pe_addr = t_ptr + rank * nbytes + 현재 (rank=SIP 모델): + per-PE tensor slice는 backend가 TensorArg로 전달 → t_ptr은 이미 local. + intra-cube ring은 tl.program_id(0) 사용. + inter-cube ring은 pe_id == 0 조건으로 제한. + inter-SIP reduce는 cube_id == 0 and pe_id == 0 조건으로 제한. + """ + local_pe = tl.program_id(axis=0) + cube_id = tl.program_id(axis=1) + + # Level 1: intra-cube ring + for _ in range(intra_rounds(pes_per_cube)): + tl.send(dir="E", src=acc) + recv = tl.recv(dir="W", shape=(n_elem,), dtype="f16") + acc = acc + recv + + # Level 2: inter-cube (cube leader only) + if local_pe == 0: + for _ in range(inter_cube_rounds(cubes_per_sip)): + tl.send(dir="N", src=acc) + recv = tl.recv(dir="S", shape=(n_elem,), dtype="f16") + acc = acc + recv + + # Level 3: inter-SIP (SIP leader only) + if local_pe == 0 and cube_id == 0: + for _ in range(inter_sip_rounds(num_sips)): + tl.send(dir="parent", src=acc) + recv = tl.recv(dir="parent", shape=(n_elem,), dtype="f16") + acc = acc + recv + + # Broadcast (reverse chain) + # ... + tl.store(t_ptr, acc) +``` + +`kernel_args`는 ADR-0024 D4의 keyword-only signature 계약을 따른다. + +### D6. Validator — `multi_pe_sip_local` + +ADR-0024 D8의 built-in 그대로 활용. `ccl.yaml`에서 `validator: +multi_pe_sip_local` 지정 시 backend가 각 SIP에 `cubes × pes_per_cube`개 +shard가 있는지 검증. + +### D7. Bench — 기본 all-reduce bench 확장 + +`benches/ccl_allreduce.py`의 worker는 `ccl.yaml`이 `hierarchical_allreduce`를 +선택하면 자동으로: + +```python +# Worker 예 +dp = DPPolicy(cube="column_wise", pe="column_wise") +tensor = torch.zeros((1, intra_sip_pes * n_elem), dp=dp, name="in") +# tensor는 각 SIP의 모든 PE에 1 tile씩 분산 (multi_pe_sip_local validator 통과) +dist.all_reduce(tensor, op="sum") +``` + +Worker 코드 자체는 알고리즘 종류를 모름 (`ccl.yaml` 선택에 의존). 단, +**DPPolicy가 hierarchical 요구와 일치해야** 함 — `cube/pe="column_wise"` +같은 SIP-내 분산을 하는 DPPolicy여야 `multi_pe_sip_local` 검증 통과. 이 +DPPolicy 선택은 bench 설정 또는 sample bench에서 결정. + +--- + +## Dependencies + +- **ADR-0024**: Launcher, `all_pes` mapper, `multi_pe_sip_local` validator, + registry + import path. 본 ADR 구현의 전제. +- **ADR-0025**: IPCQ direction addressing — cube/pe/SIP 간 다중 direction을 + 동시 사용하므로 정확한 direction 매칭 필수. +- **ADR-0023**: IPCQ protocol (neighbor table, send/recv, credit return). +- **기존 `hierarchical_allreduce.py`**: 본 ADR은 그 커널의 재해석 + 주변 + framework integration. + +--- + +## Non-goals + +- **ADR-0024 framework 변경**: 재활용만. +- **Alternative reduce topology (tree-in-tree 등)**: 3-level ring이 첫 구현. +- **Dynamic level count**: 현재 SIP/cube/PE 3단 고정. 2단 (SIP + PE, cube + skip) 또는 4단 이상은 future. +- **Bandwidth-optimal schedule tuning**: reduce round 수 / chunk size 조정 + 같은 tuning은 별도. +- **Pipelined hierarchical**: 여러 chunk를 파이프라인으로 겹쳐서 돌리는 + NCCL-style 최적화는 future. + +--- + +## Open questions + +### 🟠 중간 영향 — 구현 시 결정 필요 + +- **`topologies.py` 스키마 확장**: 기존 `ring_1d` 등은 단일 레벨 `(rank → + {dir: peer})`. `hierarchical_3level`은 multi-level. `_resolve_topology`가 + 둘을 모두 반환할 수 있도록 schema를 일반화할지, 아니면 hierarchical 전용 + return type을 두고 builder가 분기할지. + - Option A: 모든 topology를 neighbor-list 형태로 단일화 + (`[{direction, peer_sip, peer_cube, peer_pe}, ...]`) + - Option B: topology 모듈이 `kind` 필드 제공, builder가 분기 + - 권장: Option A (single source of truth, ADR-0024 Open Q의 + "PE-level topology 일원화" 방향과 일치) + +- **`hierarchical_3level` vs algorithm별 topology 모듈**: 향후 mesh-based + hierarchical 등 variant이 생기면? `hierarchical_3level` 같은 이름이 이미 + topology-specific. 변형은 새 key 추가 (`hierarchical_mesh_3level` 등) 또는 + 알고리즘 모듈에서 topology 생성 override. + +### 🟡 Nice-to-have + +- **Reduce round 수 최적화**: Bidirectional ring은 `ceil((N-1)/2)` round. + Non-power-of-2 group size에서 idle PE 발생 가능. +- **Non-uniform topology 대응**: cube_mesh가 w != h일 때 inter-cube ring + balance. +- **Single SIP 케이스**: world_size = 1 (SIP 1개)일 때 Level 3 skip. Degenerate + case 검증. + +### 🟢 Framework evolution 시사점 (ADR-0024로부터 이관) + +- **PE-level topology 일원화 (중장기)**: 현 설계는 + - topology (rank graph 또는 level-separated) + - mapper (per-SIP PE set) + - `_build_pe_installs` (actual edges) + + 의 3단 분산. Hierarchical이 이 분산을 가장 스트레스 받는 케이스. 중장기로는 + `topologies.py`가 PE-level neighbor list를 직접 반환하고 mapper는 단순히 + "어느 PE가 참여하느냐"만 결정, `_build_pe_installs`는 flat + mapping으로 단순화되는 방향이 자연스러움. **본 ADR에서 Option A를 채택**하면 + 이 방향으로 이미 정합. + +--- + +## Test strategy + +### T1. Topology generator + +`tests/test_hierarchical_topology.py` (new): +- `hierarchical_3level(rank, world_size, spec)` → 각 level의 neighbor set이 + 예상 구조인지 (intra-cube는 ring, inter-cube는 cube-leader만 참여, inter-SIP은 + SIP-leader만 참여) +- 2 SIP × 4 cubes × 4 PEs 같은 작은 토폴로지로 수작업 검증 가능 +- Symmetry: rank r의 E neighbor가 peer에서 W로 역포인팅 + +### T2. Install plan — hierarchical × all_pes + +`tests/test_ccl_install_plan.py` (확장): +- `build_install_plans(algorithm="hierarchical_allreduce", mapper="all_pes", + validator="multi_pe_sip_local")` 호출 시 + - 각 SIP의 모든 PE가 `participating_pes`에 포함 + - PE 0 (cube leader)만 inter-cube neighbor를 가짐 + - (cube 0, pe 0) (SIP leader)만 inter-SIP neighbor를 가짐 + - Non-leader PE는 intra-cube neighbor만 + +### T3. Kernel unit — mock runtime + +`tests/test_hierarchical_mock_runtime.py` (new): +- `run_kernel_in_mock` (kernbench.ccl.testing)을 확장해 multi-level 지원 +- 2 SIP × 2 cubes × 4 PEs (총 16 PE) 토폴로지에서 초기 tile을 rank+1로 채우고 + hierarchical all-reduce 실행 +- 모든 PE의 최종 결과가 `sum(1..16)`인지 + +### T4. E2E — 실제 SimPy backend + +`tests/test_ccl_allreduce_matrix.py` (확장): +- `hierarchical @ ws=SIP_count`: multi_pe_sip_local layout + 3-level 알고리즘 + 전체 stack 통과 검증 + +### T5. Validator enforcement + +- `multi_pe_sip_local` validator가 wrong layout (예: leader_only 스타일 1 + shard per rank) 입력에 raise + +### T6. 회귀 + +기존 ring/mesh/tree 알고리즘 모두 그대로 통과. 본 ADR은 그들을 건드리지 않음. + +--- + +## Consequences + +### Positive + +- **Intra-SIP PE 활용도 증가**: Inter-SIP 통신 중에도 intra-cube / inter-cube + reduce가 진행되어 전체 PE 가동률 향상. +- **Multi-level bandwidth 활용**: cube NoC, UCIe 모두 작동 → 더 정확한 HW 모델. +- **ADR-0024 framework 검증**: `all_pes` mapper + `multi_pe_sip_local` + validator의 첫 non-trivial use case. Framework 설계 타당성 확인. +- **기존 커널 재활용**: `hierarchical_allreduce.py` 큰 구조 유지, SIP-local + 좌표만 재해석. + +### Negative + +- **`topologies.py` schema 확장 필요**: Single-level vs multi-level 표현. + 해결안(Option A)은 기존 ring/mesh/tree의 마이그레이션 비용 유발. +- **Validator / mapper 조합 요구**: 사용자가 DPPolicy를 + `multi_pe_sip_local`에 맞춰 선택해야 함 (bench 설정 복잡도 증가). + +### Neutral + +- 본 ADR 구현 전까지 `hierarchical_allreduce.py`는 deprecated 상태 유지 또는 + ADR-0024 matrix test에서 제외. 현재 파일을 곧바로 삭제하지는 않음. + +--- + +## Affected files + +| File | Change | +|------|--------| +| `src/kernbench/ccl/topologies.py` | D3: `hierarchical_3level` topology 함수 추가. (Option A 채택 시) 기존 topology 출력 format 통일 | +| `src/kernbench/ccl/install_plan.py` | D4: hierarchical builder 분기 (또는 단일 builder가 level 개수로 dispatch) | +| `src/kernbench/ccl/algorithms/hierarchical_allreduce.py` | D5: SIP-local 좌표로 kernel 재작성, `kernel_args` keyword-only signature | +| `ccl.yaml` | D2: `hierarchical_allreduce` 엔트리 추가 (`mapper: all_pes`, `validator: multi_pe_sip_local`, `topology: hierarchical_3level`) | +| `tests/test_hierarchical_topology.py` (new) | T1 | +| `tests/test_ccl_install_plan.py` | T2 확장 | +| `tests/test_hierarchical_mock_runtime.py` (new) | T3 | +| `tests/test_ccl_allreduce_matrix.py` | T4: hierarchical row 추가 | diff --git a/docs/adr/ADR-0030-ipcq-physaddr.md b/docs/adr/ADR-0030-ipcq-physaddr.md new file mode 100644 index 0000000..e2ea903 --- /dev/null +++ b/docs/adr/ADR-0030-ipcq-physaddr.md @@ -0,0 +1,347 @@ +# ADR-0030: IPCQ Physical Addressing — PhysAddr integration + +## Status + +Proposed (Blocked on ADR-0031 — PhysAddr PE-resource extension) + +## Context + +### 목표 + +IPCQ ring buffer의 주소 체계를 ADR-0023의 **synthetic parallel namespace** +(`_IPCQ_BASE = 1<<60`)에서 **ADR-0001의 PhysAddr**로 이관한다. Routing / +allocator / MemoryStore의 정합성을 회복하고, buffer_kind (tcm/hbm/sram)별 +physical backing을 구조적 좌표로 표현한다. + +### 현재 상태 (ADR-0023 D2.5) + +`src/kernbench/ccl/install.py:52-56`: + +```python +_IPCQ_BASE = 1 << 60 +def _ipcq_base_for_pe(sip, cube, pe): + return _IPCQ_BASE | (sip << 40) | (cube << 32) | (pe << 24) + +def rx_base(s, c, p, d): + return _ipcq_base_for_pe(s, c, p) + direction_idx[d] * bytes_per_direction +``` + +- **bit 60** 사용 → ADR-0001의 51-bit PhysAddr 공간 밖 (`MAX_51 = (1 << 51) - 1`) +- `PhysAddr.decode(addr)` → `PhysAddrError("addr must be a 51-bit value")` +- `IpcqEndpoint.rx_base_pa: int` — 타입이 raw int, 구조 없음 +- `buffer_kind` (tcm/hbm/sram)와 synthetic 주소의 관계가 coupling 없음 +- Allocator (`PEMemAllocator`) 우회 — synthetic unique id per (sip, cube, pe, + direction). 진짜 physical allocation이 아님 + +ADR-0023 D2.5 원문: + +> This bypasses the topology's address resolver / PhysAddr encoding and +> treats IPCQ buffers as a separate, parallel address namespace. Real PA +> encoding can be plugged in later without changing the rest of the design. + +"later"가 이 ADR. + +### 왜 지금 다루는가 + +- ADR-0025 (direction addressing)은 주소-기반 매칭으로 전환. 주소가 correctness에 + 직접 기여 → 주소 체계가 설계 관점에서 더 중요해짐 +- ADR-0001의 "Routing consumes decoded domains, not raw bit-fields" 계약 위반 + 지속 → 기술 부채 +- Routing fabric (cube_noc / UCIe)은 PhysAddr.decode()로 destination을 정함. + IPCQ의 synthetic 주소가 fabric routing에서 실제로 어떻게 처리되는지 **검증되지 + 않음** (별도 경로로 배달되는 것으로 추정) +- TCM / HBM / SRAM의 실제 memory layout과 IPCQ ring buffer 위치가 **disjoint** + → allocator가 IPCQ 영역을 모르므로 실수로 겹칠 가능성 (현재는 bit 60로 완전 + 분리되어 문제 없지만 설계 원칙상 건강하지 않음) + +### 풀어야 할 문제 + +1. **IPCQ ring buffer의 PhysAddr 표현**: buffer_kind별로 어떤 PhysAddr factory를 + 쓸지. +2. **PhysAddr 공간 부족 가능성**: 51-bit 공간에 IPCQ 버퍼를 담을 여유가 있는지. +3. **Allocator 통합**: `PEMemAllocator`에 IPCQ buffer 영역 예약 기능 추가, 또는 + 기존 pool에서 정상 allocation. +4. **MemoryStore space naming 정리**: 현재는 `{"tcm", "hbm", "sram"}` 문자열로 + space 구분. IPCQ buffer도 이 space에 속하면 일반 data와 주소 겹침 방지 필요. +5. **Routing fabric 통합**: PhysAddr 기반 routing이 IPCQ 토큰을 올바른 SIP의 + 올바른 메모리로 배달. +6. **ADR-0025와의 정합**: 주소-기반 매칭이 PhysAddr에서도 동일하게 작동. + +--- + +## Decision + +### D1. IPCQ ring buffer = PhysAddr factory 사용 + +각 `buffer_kind`가 해당하는 PhysAddr factory를 호출: + +| buffer_kind | PhysAddr factory | 필요한 인자 | +|---|---|---| +| `tcm` | `PhysAddr.pe_tcm_addr(rack_id, sip_id, cube_id, pe_id, tcm_offset)` | PE-local TCM | +| `hbm` | `PhysAddr.pe_hbm_addr(rack_id, sip_id, cube_id, pe_id, pe_local_hbm_offset, slice_size_bytes)` | PE-local HBM slice | +| `sram` | `PhysAddr.cube_sram_addr(rack_id, sip_id, cube_id, sram_offset)` | Cube-shared SRAM | + +Install plan builder (`build_install_plans` in ADR-0024)가 각 PE의 rx_base를 +계산할 때: + +```python +# ADR-0030 후 install_plan.py (pseudocode) +def _compute_rx_base(sip, cube, pe, direction_idx, buffer_kind, n_slots, slot_size, + allocator_pool, rack_id=0) -> PhysAddr: + bytes_per_direction = n_slots * slot_size + offset = direction_idx * bytes_per_direction + + if buffer_kind == "tcm": + # TCM base (per-PE) + direction offset + tcm_base = allocator_pool.reserve_pe_tcm_for_ipcq(sip, cube, pe, + total_bytes=N_DIR * bytes_per_direction) + return PhysAddr.pe_tcm_addr(rack_id=rack_id, sip_id=sip, cube_id=cube, + pe_id=pe, tcm_offset=tcm_base + offset) + elif buffer_kind == "hbm": + hbm_base = allocator_pool.reserve_pe_hbm_for_ipcq(sip, cube, pe, + total_bytes=...) + return PhysAddr.pe_hbm_addr(rack_id=rack_id, sip_id=sip, cube_id=cube, + pe_id=pe, pe_local_hbm_offset=hbm_base + offset, + slice_size_bytes=slice_size) + elif buffer_kind == "sram": + sram_base = allocator_pool.reserve_cube_sram_for_ipcq(sip, cube, + total_bytes=...) + return PhysAddr.cube_sram_addr(rack_id=rack_id, sip_id=sip, cube_id=cube, + sram_offset=sram_base + offset) +``` + +`IpcqEndpoint.rx_base_pa`의 타입을 `PhysAddr` (또는 encoded `int`)로 변경: + +```python +@dataclass(frozen=True) +class IpcqEndpoint: + sip: int + cube: int + pe: int + buffer_kind: str + rx_base_pa: int # PhysAddr.encode() 결과 (51-bit) + rx_base_va: int + n_slots: int + slot_size: int +``` + +타입은 int 유지 (encoded form), 단 **반드시 PhysAddr.decode()로 복원 가능**한 +값임을 invariant으로 둔다. 디코더 호출자는 `PhysAddr.decode(rx_base_pa)`로 +구조적 좌표 획득. + +### D2. Allocator 확장 — IPCQ 예약 API + +`PEMemAllocator`에 IPCQ 전용 예약 기능 추가: + +```python +class PEMemAllocator: + def reserve_ipcq_tcm(self, total_bytes: int) -> int: + """Reserve TCM region for IPCQ ring buffers at this PE. + Returns tcm_offset (to be used in PhysAddr.pe_tcm_addr).""" + # TCM에서 `total_bytes` 연속 영역 예약. + # Tensor allocation과 겹치지 않도록. + + def reserve_ipcq_hbm(self, total_bytes: int) -> int: ... + # cube-level allocator도 유사 +``` + +Install plan 빌더가 각 PE allocator에서 예약. 예약 결과(offset)를 PhysAddr +factory에 전달. + +**기존 `_ipcq_base_for_pe` / `_IPCQ_BASE` 제거**. + +### D3. MemoryStore space 통합 + +현재 `MemoryStore`는 `{space_name: {addr: ndarray}}` 구조. IPCQ buffer는 일반 +tensor 데이터와 같은 space (tcm/hbm/sram)를 공유하게 됨. 주소 유일성은 ADR-0001의 +PhysAddr 계층 보장. + +Backward compatibility: 기존 IPCQ address (synthetic)을 쓰는 code path는 +**제거**하고, 모두 PhysAddr.encode() 결과만 사용. 이 자체는 API 변경이 아니라 +값 변경. + +### D4. Routing fabric 통합 + +IPCQ DMA write (`IpcqDmaToken`의 `src_addr → dst_addr`)이 PhysAddr encoding을 +사용하므로 **routing fabric이 `PhysAddr.decode(dst_addr)`로 destination +SIP/cube/PE를 정확히 찾을 수 있음**. Fabric routing 로직 변경 없음 (기존에도 +PhysAddr.decode를 쓰는 것으로 추정). + +**검증 필요**: 현재 fabric이 bit 60 synthetic 주소를 어떻게 라우팅하는지 확인. +별도 경로가 있다면 제거, PhysAddr 경로로 통합. + +### D5. ADR-0025와의 정합 + +ADR-0025의 주소-기반 매칭 (dst_addr로 direction 식별)은 PhysAddr.encode() +결과를 비교하는 것으로 자연스럽게 호환. 변경 없음. + +다만 debug / diagnostic 향상 가능: + +```python +# pointer_dump 등에서 +print(f"E: rx_base_pa={PhysAddr.decode(qp.peer.rx_base_pa)}") +# 출력 예: PhysAddr(sip=1, cube=0, pe=0, kind="pe_resource", unit_type=PE, ...) +``` + +이전 synthetic 주소는 decode 불가 → diagnostic 질 저하. PhysAddr 전환으로 개선. + +### D6. ADR-0023 D2.5 amendment + +ADR-0023의 "bypasses PhysAddr encoding" 문구를 **Accepted fallback → now +replaced by ADR-0030**으로 수정. 본 ADR이 적용되면 ADR-0023 D2.5의 "Real PA +encoding can be plugged in later" 약속이 이행된 것. + +--- + +## Migration strategy + +단계적 전환 (한 PR로 하지 않는다): + +### Phase 1: PhysAddr 공간 재검토 +- 51-bit PhysAddr 공간에 IPCQ ring buffer가 실제로 들어갈 수 있는지 확인. +- 각 buffer_kind (tcm/hbm/sram)별 factory가 제공하는 `local_offset` 범위가 + IPCQ 요구 (4 direction × n_slots × slot_size)를 수용 가능한지. +- 부족하면 PhysAddr layout 자체 확장 (ADR-0001 amendment 별도 필요). + +### Phase 2: Allocator API 확장 +- `PEMemAllocator.reserve_ipcq_*` 메소드 추가. +- 기존 tensor allocation과 영역 충돌 방지. + +### Phase 3: Install plan builder 전환 +- `_ipcq_base_for_pe` 제거, PhysAddr factory 호출로 대체. +- `IpcqEndpoint.rx_base_pa`가 PhysAddr.encode() 결과 (51-bit). + +### Phase 4: Routing fabric 검증 +- IPCQ DMA token이 fabric 정상 경로로 배달되는지 확인. +- 별도 fast-path가 있다면 제거, 통합. + +### Phase 5: MemoryStore space 검증 +- IPCQ buffer 주소가 기존 tensor 주소와 겹치지 않는지. +- Allocator 레벨에서 이미 예약했으므로 정상적으로 분리되어야 함. + +### Phase 6: ADR-0023 D2.5 업데이트 + 기존 sideband path 제거 (완료) + +--- + +## Dependencies + +- **ADR-0031** (PhysAddr PE-resource extension) — **Blocker**: PhysAddr가 PE + resource (특히 IPCQ ring buffer)를 충분히 표현할 수 있도록 schema 확장이 + 선행되어야 함. 본 ADR은 ADR-0031 완료 후에만 실행 가능. +- **ADR-0001** (PhysAddr layout): 본 ADR의 기반. 51-bit 공간 / factory API의 + ADR-0031 확장본을 사용. +- **ADR-0023** (IPCQ protocol): 본 ADR은 ADR-0023 D2.5의 "later" 약속 이행. + D9 piggyback / credit return 프로토콜 자체는 불변. +- **ADR-0024** (launcher + install_plan.py): `build_install_plans`가 PhysAddr + factory를 호출하게 됨. +- **ADR-0025** (direction addressing): 주소-기반 매칭이 PhysAddr에서도 동일하게 + 작동. 변경 없음. + +--- + +## Non-goals + +- **ADR-0001 PhysAddr layout 자체 변경**: 51-bit 공간과 segment 구조는 유지. + 부족 시 별도 ADR. +- **IPCQ protocol semantic 변경**: ADR-0023 D9 piggyback 등 프로토콜 로직 유지. +- **Allocator 전반 재설계**: IPCQ 예약 API 추가만. + +--- + +## Open questions + +### 🔴 Critical — Migration 전 반드시 검증 + +- **PhysAddr 51-bit 공간에 IPCQ 버퍼가 실제로 들어가는가**: 각 PE의 TCM + 영역에서 `4 direction × n_slots (default 4) × slot_size (default 4KB)` = + 64KB가 PE TCM 공간에 수용 가능. TCM size (e.g., 16MB) 대비 충분. HBM도 여유 + 많음. SRAM은 cube 공유라 direction × PE 곱이 있음 — 별도 검증 필요. +- **Routing fabric의 현재 IPCQ 주소 처리**: 현재 synthetic 주소가 fabric에서 + 어떻게 routing되는지 trace 필요. `PhysAddr.decode()`로 판독 불가한 값이 + fabric에서 정상 배달된다면 어떤 경로를 쓰는지 조사. + +### 🟡 Nice-to-have + +- **IPCQ 전용 kind / sub_offset 인코딩**: `UnitType.PE`의 sub_offset 공간을 + IPCQ와 공유. 충돌 방지를 위해 IPCQ 전용 sub-space 정의할지 여부. +- **Debug tool**: `pointer_dump`를 PhysAddr 포매팅으로 개선. + +--- + +## Test strategy + +### T1. PhysAddr round-trip + +`tests/test_ipcq_physaddr.py` (new): +- `PhysAddr.pe_tcm_addr(...)` → encode → decode → 동일 필드 복원 +- TCM / HBM / SRAM 각 factory에 대해 + +### T2. Allocator 예약 + +`tests/test_ipcq_alloc.py` (new): +- `PEMemAllocator.reserve_ipcq_tcm` → 반환된 offset이 valid TCM 영역 +- 중복 예약 → 에러 또는 non-overlapping offset +- Tensor allocation과 충돌 없음 + +### T3. Install plan PhysAddr integration + +`tests/test_ccl_install_plan.py` (확장): +- `build_install_plans` 결과의 `rx_base_pa`가 PhysAddr.decode() 가능 +- Decoded 좌표가 plan의 (sip, cube, pe)와 일치 +- I3.1 invariant (ADR-0025 D6) — rx_base range disjointness가 PhysAddr에서도 성립 + +### T4. Routing — IPCQ DMA fabric traversal + +`tests/test_ipcq_routing.py` (new): +- Cross-SIP IPCQ send → fabric이 `PhysAddr.decode(dst_addr)`로 destination SIP + 정확히 판단 → 올바른 MemoryStore에 write +- UCIe 경로 / cube_noc 경로 모두 검증 + +### T5. 회귀 + +- 기존 IPCQ E2E 테스트 (ring, mesh, tree) 모두 통과 +- ADR-0024, ADR-0025 통합 테스트 통과 + +--- + +## Consequences + +### Positive + +- **ADR-0001 정합성 회복**: routing과 addressing이 단일 체계. +- **buffer_kind 명확**: TCM/HBM/SRAM이 구조적 좌표로 구분. +- **Debug 향상**: PhysAddr.decode()로 사람이 읽을 수 있는 좌표. +- **Allocator 통합**: IPCQ 영역이 정상 예약 → tensor와의 충돌 리스크 사전 차단. +- **Fabric routing 일원화**: 별도 경로 없이 기존 PhysAddr-based routing 재활용. + +### Negative + +- **Migration 복잡도**: 6 Phase 단계적 전환 필요. 각 Phase마다 regression 리스크. +- **PhysAddr 공간 검증 부담**: Phase 1에서 TCM/HBM/SRAM 공간이 IPCQ 요구를 + 수용하는지 실측 필요. +- **Routing fabric 검증**: 현재 fabric이 synthetic 주소를 어떻게 처리하는지 + 조사 필요. + +### Neutral + +- IPCQ protocol semantic (ADR-0023 D9 등) 불변. +- ADR-0025의 direction addressing 로직 불변. + +--- + +## Affected files + +| File | Change | +|------|--------| +| `src/kernbench/ccl/install.py` | `_IPCQ_BASE`, `_ipcq_base_for_pe` 제거 | +| `src/kernbench/ccl/install_plan.py` (ADR-0024) | D1: PhysAddr factory 호출로 rx_base 계산 | +| `src/kernbench/policy/address/allocator.py` (or similar) | D2: IPCQ 예약 API (`reserve_ipcq_tcm` 등) | +| `src/kernbench/common/ipcq_types.py` | D1: `IpcqEndpoint.rx_base_pa` 문서화 — PhysAddr.encode 결과 | +| `src/kernbench/sim_engine/memory_store.py` | D3: IPCQ buffer가 기존 space와 공유되는지 검증 | +| `src/kernbench/sim_engine/engine.py` | D4: IPCQ token routing이 PhysAddr-based fabric 경로 사용 | +| `src/kernbench/ccl/diagnostics.py` | D5: pointer_dump를 PhysAddr 포매팅으로 개선 | +| `docs/adr/ADR-0023-ipcq-pe-collective.md` | D6: D2.5 amendment note | +| `tests/test_ipcq_physaddr.py` (new) | T1 | +| `tests/test_ipcq_alloc.py` (new) | T2 | +| `tests/test_ccl_install_plan.py` | T3 확장 | +| `tests/test_ipcq_routing.py` (new) | T4 | diff --git a/docs/adr/ADR-0031-physaddr-pe-resource-extension.md b/docs/adr/ADR-0031-physaddr-pe-resource-extension.md new file mode 100644 index 0000000..8633b2c --- /dev/null +++ b/docs/adr/ADR-0031-physaddr-pe-resource-extension.md @@ -0,0 +1,257 @@ +# ADR-0031: PhysAddr PE-Resource Extension + +## Status + +Stub (Blocker for ADR-0030 — specific range allocations TBD) + +## Context + +### 목표 + +ADR-0001의 `PhysAddr` schema를 **PE 내부의 다양한 resource**를 체계적으로 +표현할 수 있도록 확장한다. ADR-0030 (IPCQ PhysAddr integration) 및 향후의 +PE-local resource 추가 (scratchpad, register file, status register, 등)의 +기반을 제공한다. + +### 현재 상태 (ADR-0001) + +51-bit PhysAddr layout: + +``` +[50:47] rack_id (4) +[46:43] sip_id (4) +[42:38] sip_seg (5) # cube_id +[37:0] local_offset (38) +``` + +`local_offset` (38 bits) 내부: + +- `[37]` selector: 1 = HBM window (128GB), 0 = PE resource window +- PE resource window는 `unit_type` (3 bits: PE | MCPU | SRAM) + + `pe_id` (4 bits) + `ext` (1 bit) + `sub_offset` (29 bits) + +Factory API: +- `PhysAddr.hbm_addr(...)` — HBM generic +- `PhysAddr.pe_hbm_addr(...)` — PE-local HBM slice +- `PhysAddr.pe_tcm_addr(...)` — PE TCM (via `UnitType.PE` + `sub_offset`) +- `PhysAddr.cube_sram_addr(...)` — Cube-shared SRAM + +### 풀어야 할 문제 + +1. **PE 내부 resource 구분의 명시적 체계 부재**: 현재 `local_offset` (38 bits) + 이 평면 공간으로 취급되고, PE TCM / IPCQ ring / scratchpad / 향후 register + file 등이 관습적 offset 범위로만 구분됨. Schema 레벨에서 명확하지 않음. +2. **IPCQ 주소의 PhysAddr 표현 부재**: ADR-0030이 IPCQ ring buffer를 PhysAddr로 + 표현하려면 "이 주소가 IPCQ 영역"을 decode 가능해야 함. 현재는 불가. +3. **향후 PE resource 확장 경로**: register file, performance counter 등 + 추가 시 일관된 위치 할당 규칙 필요. + +### 설계 방향 — local_offset을 PE 컴포넌트별 range로 분할 + +`local_offset` (38 bits = 256GB per PE segment)을 **PE 컴포넌트마다 고정 +range**로 나누어 할당한다. 각 range는 해당 컴포넌트 전용 주소 공간이며, +`PhysAddr.decode()`가 주소가 어느 range에 속하는지 판별해 해당하는 `kind` / +`unit_type` / `sub_type` 필드를 채운다. + +개념적 구조 (구체적 bit 할당은 **TBD**): + +``` +local_offset [37:0] (38 bits total) +├── HBM window [37] = 1 (기존 128GB) +├── PE component ranges [37] = 0 +│ ├── TCM [range_1] +│ ├── IPCQ rings [range_2] +│ ├── Scratchpad [range_3] +│ ├── Register file [range_4] +│ ├── (reserved) ... +│ └── Sideband / status [range_N] +``` + +### 왜 range-based partition인가 + +- **Schema-level 명시성**: 주소 하나 보고 어느 컴포넌트의 자원인지 decode 가능. + "Routing consumes decoded domains" (ADR-0001 D5) 계약 충족. +- **Unit type enum 확장보다 유연**: 3-bit `UnitType` 공간을 고갈시키지 않고 + 세분화 가능. 미래 추가 컴포넌트도 빈 range 할당. +- **Allocator 통합 자연**: 각 PE-level allocator가 관리하는 하위 pool을 + address range와 1:1 매칭 (e.g., `reserve_ipcq_tcm()` → IPCQ range 안에서만 + 할당). +- **Decode routing 단순**: `PhysAddr.decode(addr)`가 range table을 참조해 + `kind` + sub-field를 채움. 기존 HBM selector bit 패턴의 일반화. + +### 왜 지금 다루는가 + +- ADR-0030 (IPCQ PhysAddr 통합)이 이 확장에 **의존**. ADR-0030 단독 진행 시 + `sub_offset` 공간을 불투명하게 재사용하게 되어 ADR-0001 계약 미충족. +- PE 내부 자원이 더 추가될 가능성 — 지금 구조를 정리해두면 일관된 확장 경로 확보. + +--- + +## Decision (pending specific range allocation) + +### D1. Range-based local_offset partition — approach + +`local_offset`을 고정 byte range로 분할하고, 각 range를 PE 컴포넌트에 할당한다. +주소의 어느 range에 속하는가로 `kind` / component type을 결정. + +```python +# src/kernbench/policy/address/phyaddr.py (conceptual, post-extension) +@dataclass(frozen=True) +class PeResourceRange: + name: str # e.g. "tcm", "ipcq", "scratchpad", "regfile" + start_offset: int # local_offset 내 시작 + end_offset: int # exclusive + byte_size: int # end - start + +PE_RESOURCE_MAP: tuple[PeResourceRange, ...] = ( + # TBD — 구체적 range 할당은 사용자가 별도 업데이트 +) +``` + +`PhysAddr.decode(addr)`의 PE resource 경로는: + +```python +def decode_pe_resource(local_offset: int) -> dict: + for r in PE_RESOURCE_MAP: + if r.start_offset <= local_offset < r.end_offset: + return { + "kind": "pe_resource", + "component": r.name, # NEW: "tcm"/"ipcq"/... + "component_offset": local_offset - r.start_offset, # within range + } + raise PhysAddrError(f"local_offset {local_offset} not in any PE range") +``` + +### D2. Specific range allocations — **TBD** + +> 사용자가 구체적 byte 할당을 별도로 정의한 뒤 본 ADR에 업데이트. +> +> 필요 정보: +> - 각 컴포넌트 (TCM, IPCQ, scratchpad, regfile, ...)의 이름 / byte size +> - `local_offset` 내 시작 offset (align 고려) +> - 현재 하드웨어 사양 / 시뮬레이션 요구 반영 + +이 섹션이 채워진 뒤 ADR status: **Stub → Proposed → Accepted** 승격. + +### D3. Factory API — per-component 함수 + +기존 `PhysAddr.pe_tcm_addr(...)` 패턴을 일반화: + +```python +# 기존 (이미 존재) +PhysAddr.pe_tcm_addr(rack_id, sip_id, cube_id, pe_id, tcm_offset) + +# 신규 (ADR-0031 후 추가) +PhysAddr.pe_ipcq_addr(rack_id, sip_id, cube_id, pe_id, ipcq_offset) +PhysAddr.pe_scratchpad_addr(...) +PhysAddr.pe_regfile_addr(...) +# ... +``` + +각 factory는 해당 컴포넌트의 range 내에서 `component_offset`만 받아 최종 +PhysAddr encoding. 호출자는 어느 range인지 몰라도 됨. + +### D4. Backward compatibility + +- 기존 `pe_tcm_addr()` signature / semantic 유지. +- 내부 인코딩만 신규 range table을 참조하도록 변경. +- 기존 `UnitType.PE` decoding 경로는 `PE_RESOURCE_MAP`에서 "tcm" range를 + 대응하도록 매핑 → 기존 코드 transparent. +- 기존 코드가 `PhysAddr.decode(addr).unit_type == UnitType.PE`를 체크하는 + 경우는 여전히 유효 (TCM 주소는 계속 PE unit_type). + +--- + +## Open questions + +### 🔴 Pending user input (ADR 승격 blocker) + +- **D2의 specific range allocation**: 사용자가 구체적 byte 할당 테이블을 + 제공해야 Stub → Proposed 승격 가능. 필요 정보: + - 컴포넌트 목록 (TCM, IPCQ, scratchpad, regfile 등) + - 각 컴포넌트의 byte size / 시작 offset + - Alignment 요구사항 (4KB / page-aligned 등) + +### 🟡 설계 세부 — range allocation 결정 과정에서 함께 결정 + +- **총 local_offset space 배분**: HBM window (bit 37 = 1, 128GB)을 유지할지, + 아니면 PE resource space를 확장하기 위해 HBM window 축소할지. +- **Range padding / reserved space**: 미래 컴포넌트 추가를 위한 "reserved" + range 몇 개를 미리 확보할지. +- **Address alignment**: 각 range의 시작 offset이 특정 alignment (page / + cache line) 만족해야 하는지. +- **Diagnostic / debug 포맷**: `PhysAddr.decode()` 출력에서 component 이름 + + component_offset을 사람이 읽기 좋게 표시 (e.g., "IPCQ ring sip=0 cube=0 pe=3 + offset=0x1234"). +- **기존 `UnitType` enum의 role**: Range-based 접근 후에도 `unit_type` 필드 + 유지할지 (decode 결과에 `component` 추가), 또는 enum 대체할지. + +### 🟢 ADR-0030 연동 질문 + +- **IPCQ range 내 direction/slot 표현**: PhysAddr는 `component_offset` 단위 + 까지만 표현. "direction=E, slot=2"는 IPCQ range 내 offset 계산으로 도출 + (`direction_idx * slot_region_size + slot_idx * slot_size`) — 이 공식은 + ADR-0030 scope에서 구체화. +- **Allocator pool 구조**: `PEMemAllocator`가 여러 range (TCM, IPCQ, + scratchpad)를 개별 pool로 관리할지, 단일 pool에서 kind별 reserved만 관리 + 할지. Range-based schema면 개별 pool이 자연스러움. + +--- + +## Non-goals (this ADR) + +- **51-bit 전체 layout 재작성**: 본 ADR은 `local_offset` (38 bits) 내부의 + subdivision만 다룬다. Rack / SIP / cube segment 같은 상위 bit 구조는 + 불변. +- **`UnitType` enum 재설계**: range-based 접근으로 대체 가능하지만, 기존 enum + (PE / MCPU / SRAM)은 backward compat 위해 유지. +- **Dynamic range allocation**: runtime에 range 크기 바꾸는 기능 불필요. 모든 + range는 컴파일 / 설정 시점에 고정. +- **Multi-process / multi-rack partitioning**: PE 내부 resource만 다룸. + +--- + +## Action + +### Phase 1 — User 입력: specific range allocation (**Blocker**) +- 사용자가 정의한 PE 컴포넌트별 byte range를 D2에 기입: + - `PE_RESOURCE_MAP` 테이블 내용 (name, start_offset, byte_size per 컴포넌트) + - 각 컴포넌트의 hardware spec 근거 note + +### Phase 2 — ADR Stub → Proposed 승격 +- D2 채워지면 status 변경. +- Open questions의 "🔴 Pending user input" 블록 제거. +- ADR-0001에 amendment note 초안 작성. + +### Phase 3 — 구현 +- `PhysAddr` range-based decode 구현. +- 신규 factory 함수 (`pe_ipcq_addr`, `pe_scratchpad_addr` 등 컴포넌트별) + 추가. +- 기존 `pe_tcm_addr` 내부 인코딩만 신규 range table 참조하도록 수정 + (signature 불변). +- 기존 코드 경로 회귀 확인. + +### Phase 4 — ADR-0030 unblock +- ADR-0030 "Blocked" 상태 해제. +- Install_plan builder가 `pe_ipcq_addr(...)` 등 확장된 factory 호출하도록 + 수정. + +--- + +## Dependencies + +- **ADR-0001** (PhysAddr layout): 본 ADR은 ADR-0001의 확장. +- **ADR-0023** (IPCQ protocol): IPCQ ring buffer의 주소 체계를 PhysAddr로 + 통합할 수 있게 하는 기반. +- **ADR-0030** (IPCQ PhysAddr integration): 본 ADR에 blocked. + +--- + +## Affected files (future, after promotion to Proposed) + +| File | Change | +|------|--------| +| `src/kernbench/policy/address/phyaddr.py` | Range table (`PE_RESOURCE_MAP`), range-based decode, 신규 component-specific factory들 (`pe_ipcq_addr` 등), 기존 `pe_tcm_addr` 내부 인코딩 갱신 | +| `src/kernbench/policy/address/allocator.py` | Range-aware pool 분리 (TCM pool / IPCQ pool / scratchpad pool 등 per-PE) | +| `docs/adr/ADR-0001-physaddr-layout.md` | Amendment note: range-based PE resource partition | +| `tests/test_phyaddr.py` | Range table 검증, 각 factory의 encode/decode round-trip, 기존 `pe_tcm_addr` 회귀 |