Fix ADR-0025: IPCQ direction addressing via address-based matching

2-rank bidirectional ring deadlock: when E and W neighbors point to the
same peer, sender-coord matching in _handle_meta_arrival / _credit_worker
picked the first direction in dict order, landing data in the wrong rx
slot relative to what the kernel recv(W) was waiting on.

Fix (ADR-0025 D1/D2/D3):
- install.reverse_direction: prefer OPPOSITE direction (E↔W, N↔S) when
  peer has it pointing back to us; fallback to any matching for
  topologies without opposite convention (tree_binary parent/child).
- _handle_meta_arrival: match by token.dst_addr range against each qp's
  my_rx_base_pa + n_slots × slot_size window (unambiguous).
- _credit_worker: match by credit.dst_rx_base_pa == qp.peer.rx_base_pa.
- IpcqCreditMetadata: new dst_rx_base_pa field carrying receiver-side
  rx base; _delayed_credit_send fills it from the consuming qp.

Tests (Phase 1 → Phase 2):
- test_reverse_direction_opposite_preference_2rank_ring
- test_reverse_direction_opposite_preference_4rank_ring_sanity
- test_meta_arrival_matches_by_dst_addr_same_peer
- test_credit_matches_by_dst_rx_base_pa_same_peer
- Existing credit-return test updated with dst_rx_base_pa.

508 tests pass.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

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 단순화 |

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 |

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 |

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` | 신규 (선택): 사용자 가이드 |

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__` 등 추가 |

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 추가 |

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 |

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` 회귀 |

src/kernbench/ccl/install.py

@@ -219,9 +219,24 @@ def install_ipcq(
         "neighbor_table": neighbor_table,
     }
 
-    def reverse_direction(my_rank: int, peer_rank: int) -> str | None:
-        """Find which direction in peer's neighbor table points back to my_rank."""
-        for d, target in neighbor_table[peer_rank].items():
+    _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 (ADR-0025 D1). This matters in 2-rank
+        bidirectional rings where both E and W on one side point to the
+        same peer — without the preference, dict-order first-match would
+        route data into the wrong rx slot. Falls back to any direction
+        pointing back for topologies without an opposite convention
+        (e.g. 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
@@ -234,7 +249,7 @@ def install_ipcq(
             if peer_rank is None:
                 continue
             peer_s, peer_c, peer_p = rank_pe[peer_rank]
-            peer_dir = reverse_direction(r, peer_rank)
+            peer_dir = reverse_direction(r, peer_rank, d)
             if peer_dir is None:
                 # Peer doesn't have a reverse entry — skip (asymmetric topology)
                 continue

src/kernbench/common/ipcq_types.py

@@ -196,10 +196,17 @@ class IpcqCreditMetadata:
     Sent by ``PeIpcqComponent._delayed_credit_send`` after a
     bottleneck-BW based latency, putting the metadata directly into
     the peer's pre-wired credit store (no fabric routing).
+
+    ``dst_rx_base_pa`` is the receiver's ``my_rx_base_pa`` for the direction
+    whose slot was consumed. The original sender matches this against
+    ``qp.peer.rx_base_pa`` to find the correct direction (ADR-0025 D3) —
+    unambiguous even when multiple directions share the same peer (e.g.
+    2-rank bidirectional ring).
     """
 
     consumer_seq: int            # my_tail at recv side (new tail value)
-    src_sip: int                 # which peer is sending the credit
+    dst_rx_base_pa: int          # receiver-side my_rx_base_pa (ADR-0025 D3)
+    src_sip: int                 # which peer is sending the credit (diag)
     src_cube: int
     src_pe: int
     src_direction: str           # sender-side direction (peer maps to its own)

src/kernbench/components/builtin/pe_ipcq.py

@@ -370,11 +370,21 @@ class PeIpcqComponent(ComponentBase):
     # ── Metadata arrival from PE_DMA (D9) ──
 
     def _handle_meta_arrival(self, msg: IpcqMetaArrival) -> None:
+        """Match arrival to the correct direction by dst_addr range (ADR-0025 D2).
+
+        Each direction has a unique rx buffer address range
+        ([my_rx_base_pa, my_rx_base_pa + n_slots * slot_size)). The token's
+        dst_addr (set by the sender's IPCQ when computing the peer slot
+        address) falls within exactly one such range. Address-based matching
+        is unambiguous even when multiple directions share the same peer
+        (2-rank bidirectional ring).
+        """
         token = msg.token
-        sender_key = (token.src_sip, token.src_cube, token.src_pe)
+        dst_addr = token.dst_addr
         for d, qp in self._queue_pairs.items():
-            p = qp["peer"]
-            if (p.sip, p.cube, p.pe) == sender_key:
+            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)
                 # Track arrived token for strict-mode peek
                 self._arrived_tokens.setdefault(d, []).append(token)
@@ -391,19 +401,22 @@ class PeIpcqComponent(ComponentBase):
                     if not ev.triggered:
                         ev.succeed()
                 return
-        # Unknown sender — silently drop (could log)
+        # Unknown dst_addr — silently drop (could log)
 
     # ── Credit return (fast path) ──
 
     def _credit_worker(self, env: simpy.Environment) -> Generator:
-        """Process IpcqCreditMetadata from credit_inbox."""
+        """Process IpcqCreditMetadata from credit_inbox.
+
+        Matches credit to the correct direction by `credit.dst_rx_base_pa ==
+        qp.peer.rx_base_pa` (ADR-0025 D3). This is unambiguous even when
+        multiple directions share the same peer (2-rank bidirectional ring).
+        """
         assert self._credit_inbox is not None
         while True:
             credit: IpcqCreditMetadata = yield self._credit_inbox.get()
-            sender_key = (credit.src_sip, credit.src_cube, credit.src_pe)
             for d, qp in self._queue_pairs.items():
-                p = qp["peer"]
-                if (p.sip, p.cube, p.pe) == sender_key:
+                if qp["peer"].rx_base_pa == credit.dst_rx_base_pa:
                     qp["peer_tail_cache"] = max(qp["peer_tail_cache"], credit.consumer_seq)
                     # Wake any blocked send on this direction
                     waiters = self._send_waiters.get(d, [])
@@ -421,12 +434,19 @@ class PeIpcqComponent(ComponentBase):
         new_tail: int,
     ) -> Generator:
         """Wait bottleneck-BW latency, then put IpcqCreditMetadata into peer
-        credit store (D9 fast path)."""
+        credit store (D9 fast path).
+
+        Carries ``dst_rx_base_pa`` = this PE's my_rx_base_pa for the
+        consumed direction. The peer (original sender) matches this against
+        qp.peer.rx_base_pa to identify the correct qp (ADR-0025 D3).
+        """
         latency_ns = self._credit_latency_ns(direction)
         if latency_ns > 0:
             yield env.timeout(latency_ns)
+        qp = self._queue_pairs[direction]
         meta = IpcqCreditMetadata(
             consumer_seq=new_tail,
+            dst_rx_base_pa=qp["my_rx_base_pa"],
             src_sip=self._self_sip,
             src_cube=self._self_cube,
             src_pe=self._self_pe,

tests/test_ccl_install.py

@@ -98,3 +98,63 @@ def test_install_ipcq_credit_stores_wired():
 
     qp_e = pe0.queue_pairs["E"]
     assert qp_e["peer_credit_store"] is pe1.credit_inbox
+
+
+# ── ADR-0025 D1: reverse_direction opposite-preference ───────────────
+
+
+def test_reverse_direction_opposite_preference_2rank_ring():
+    """ADR-0025 D1: In a 2-rank bidirectional ring both E and W point to the
+    same peer; reverse_direction must pick the OPPOSITE direction (W for E,
+    E for W) so rx_base targets the semantically-correct slot.
+
+    Concretely: rank 0 sending via E to rank 1 must target rank 1's W-rx
+    buffer (not rank 1's E-rx), because rank 1's kernel recv(W) reads from
+    its W-rx.
+    """
+    engine, topo = _engine()
+    cfg = load_ccl_config()
+    merged = resolve_algorithm_config(cfg, name="ring_allreduce_tcm")
+    merged["world_size"] = 2
+    install_ipcq(engine, topo.spec, merged)
+
+    ipcq0 = engine._components["sip0.cube0.pe0.pe_ipcq"]
+    ipcq1 = engine._components["sip0.cube0.pe1.pe_ipcq"]
+
+    rank1_e_rx = ipcq1.queue_pairs["E"]["my_rx_base_pa"]
+    rank1_w_rx = ipcq1.queue_pairs["W"]["my_rx_base_pa"]
+
+    qp0_e = ipcq0.queue_pairs["E"]
+    qp0_w = ipcq0.queue_pairs["W"]
+
+    # rank 0's E entry should target rank 1's W-rx (opposite), NOT rank 1's E-rx.
+    assert qp0_e["peer"].rx_base_pa == rank1_w_rx, (
+        f"expected rank 0's E peer.rx_base_pa == rank 1's W-rx ({rank1_w_rx:#x}), "
+        f"got {qp0_e['peer'].rx_base_pa:#x} (matches E-rx: {rank1_e_rx:#x}) — "
+        f"reverse_direction picked same-label instead of opposite"
+    )
+    # rank 0's W entry should target rank 1's E-rx (opposite).
+    assert qp0_w["peer"].rx_base_pa == rank1_e_rx
+
+
+def test_reverse_direction_opposite_preference_4rank_ring_sanity():
+    """ADR-0025 D1 sanity: ws>=3 ring. E and W have distinct peers, so
+    opposite-preference produces same result as old dict-order first-match.
+    This test should PASS both under current and post-fix code.
+    """
+    engine, topo = _engine()
+    cfg = load_ccl_config()
+    merged = resolve_algorithm_config(cfg, name="ring_allreduce_tcm")
+    merged["world_size"] = 4
+    install_ipcq(engine, topo.spec, merged)
+
+    ipcq0 = engine._components["sip0.cube0.pe0.pe_ipcq"]
+    ipcq1 = engine._components["sip0.cube0.pe1.pe_ipcq"]
+    ipcq3 = engine._components["sip0.cube0.pe3.pe_ipcq"]
+
+    # rank 0 E → rank 1 → rank 1's W-rx
+    qp0_e = ipcq0.queue_pairs["E"]
+    assert qp0_e["peer"].rx_base_pa == ipcq1.queue_pairs["W"]["my_rx_base_pa"]
+    # rank 0 W → rank 3 (last in ring) → rank 3's E-rx
+    qp0_w = ipcq0.queue_pairs["W"]
+    assert qp0_w["peer"].rx_base_pa == ipcq3.queue_pairs["E"]["my_rx_base_pa"]

tests/test_ipcq_types.py

@@ -63,7 +63,8 @@ def test_ipcq_dma_token():
 
 def test_ipcq_credit_metadata():
     cm = IpcqCreditMetadata(
-        consumer_seq=3, src_sip=0, src_cube=0, src_pe=1, src_direction="W",
+        consumer_seq=3, dst_rx_base_pa=0x1000,
+        src_sip=0, src_cube=0, src_pe=1, src_direction="W",
     )
     assert cm.consumer_seq == 3
     assert cm.src_direction == "W"
@@ -71,7 +72,8 @@ def test_ipcq_credit_metadata():
 
 def test_ipcq_credit_metadata_frozen():
     cm = IpcqCreditMetadata(
-        consumer_seq=3, src_sip=0, src_cube=0, src_pe=1, src_direction="W",
+        consumer_seq=3, dst_rx_base_pa=0x1000,
+        src_sip=0, src_cube=0, src_pe=1, src_direction="W",
     )
     with pytest.raises(Exception):
         cm.consumer_seq = 99  # type: ignore

tests/test_pe_ipcq.py

@@ -291,9 +291,12 @@ def test_send_blocks_when_peer_slot_full():
     env.run(until=20)
     assert not req5.done.triggered
 
-    # Send a credit return: peer (E direction, pe=1) consumed slot 0
+    # Send a credit return: peer (E direction, pe=1) consumed slot 0.
+    # dst_rx_base_pa is the peer-side rx buffer — which equals my qp_E's
+    # peer.rx_base_pa (0x10_000 from _install_two_neighbors).
     credit = IpcqCreditMetadata(
         consumer_seq=1,  # peer consumed up to my_tail=1
+        dst_rx_base_pa=0x10_000,  # E's peer.rx_base_pa (ADR-0025 D3)
         src_sip=0, src_cube=0, src_pe=1, src_direction="W",  # peer's view
     )
     comp.credit_inbox.put(credit)
@@ -315,3 +318,156 @@ def test_init_installs_neighbors():
     assert comp._queue_pairs["W"]["peer"].pe == 2
     assert comp._queue_pairs["E"]["my_head"] == 0
     assert comp._queue_pairs["E"]["peer_tail_cache"] == 0
+
+
+# ── ADR-0025: address-based matching in meta arrival / credit ────────
+
+
+def _install_same_peer_neighbors(
+    env: simpy.Environment, comp: PeIpcqComponent,
+) -> tuple[simpy.Store, simpy.Store]:
+    """Install E and W neighbors BOTH pointing to the same peer (pe=1).
+
+    This mirrors the 2-rank bidirectional ring topology (ADR-0025 motivation):
+    rank 0's E and W neighbors are the same peer rank, but target different
+    rx slots on that peer (E→peer's W-rx, W→peer's E-rx).
+
+    - E's peer.rx_base_pa = 0x10_000 (peer's W-rx buffer)
+    - W's peer.rx_base_pa = 0x20_000 (peer's E-rx buffer)
+    - my_rx_base_pa: E=0x30_000, W=0x40_000 (local rx for each dir)
+    """
+    peer_e_credit = simpy.Store(env)
+    peer_w_credit = simpy.Store(env)
+
+    ep_e = IpcqEndpoint(
+        sip=0, cube=0, pe=1,
+        buffer_kind="tcm",
+        rx_base_pa=0x10_000, rx_base_va=0,
+        n_slots=4, slot_size=4096,
+    )
+    ep_w = IpcqEndpoint(
+        sip=0, cube=0, pe=1,  # SAME peer as ep_e
+        buffer_kind="tcm",
+        rx_base_pa=0x20_000, rx_base_va=0,  # different target slot
+        n_slots=4, slot_size=4096,
+    )
+    init_msg = IpcqInitMsg(
+        correlation_id="t", request_id="t",
+        target_sips=(0,), target_cubes=(0,), target_pe=0,
+        entries=(
+            IpcqInitEntry(
+                direction="E", peer=ep_e,
+                my_rx_base_pa=0x30_000, my_rx_base_va=0,
+                n_slots=4, slot_size=4096,
+                peer_credit_store=peer_e_credit,
+            ),
+            IpcqInitEntry(
+                direction="W", peer=ep_w,
+                my_rx_base_pa=0x40_000, my_rx_base_va=0,
+                n_slots=4, slot_size=4096,
+                peer_credit_store=peer_w_credit,
+            ),
+        ),
+        backpressure_mode="sleep",
+        buffer_kind="tcm",
+        credit_size_bytes=16,
+    )
+    done = env.event()
+    comp.in_ports["host"].put(_FakeTxn(request=init_msg, done=done))
+    env.run(until=done)
+    return peer_e_credit, peer_w_credit
+
+
+def test_meta_arrival_matches_by_dst_addr_same_peer():
+    """ADR-0025 D2: when E and W point to the same peer (2-rank ring),
+    dst_addr range must determine which qp's peer_head_cache updates.
+
+    Under the old sender-key matching, the first matching direction (E)
+    would win for any arrival, regardless of which rx slot was written.
+    Under D2 address-based matching, dst_addr within W's rx range
+    (my_rx_base_pa_W .. +n_slots*slot_size) must update W, and dst_addr
+    within E's rx range must update E.
+    """
+    env = simpy.Environment()
+    comp = _make_pe_ipcq(env)
+    _install_same_peer_neighbors(env, comp)
+
+    # Arrival into W's rx buffer (my_rx_base_pa=0x40_000)
+    token_into_w = IpcqDmaToken(
+        src_addr=0, src_space="tcm",
+        dst_addr=0x40_000, dst_endpoint=comp._queue_pairs["W"]["peer"],
+        nbytes=64, handle_id="w1",
+        shape=(8,), dtype="f16",
+        sender_seq=0,
+        src_sip=0, src_cube=0, src_pe=1, src_direction="E",
+    )
+    comp.in_ports["host"].put(IpcqMetaArrival(token=token_into_w))
+    env.run(until=5)
+
+    # W's peer_head_cache should increment; E's stays 0.
+    assert comp._queue_pairs["W"]["peer_head_cache"] == 1, (
+        "W qp should have been updated because dst_addr is in W's rx range"
+    )
+    assert comp._queue_pairs["E"]["peer_head_cache"] == 0, (
+        "E qp should NOT be updated; current sender-key matching wrongly "
+        "picks the first direction with a matching peer"
+    )
+
+    # Second arrival into E's rx buffer (my_rx_base_pa=0x30_000)
+    token_into_e = IpcqDmaToken(
+        src_addr=0, src_space="tcm",
+        dst_addr=0x30_000, dst_endpoint=comp._queue_pairs["E"]["peer"],
+        nbytes=64, handle_id="e1",
+        shape=(8,), dtype="f16",
+        sender_seq=0,
+        src_sip=0, src_cube=0, src_pe=1, src_direction="W",
+    )
+    comp.in_ports["host"].put(IpcqMetaArrival(token=token_into_e))
+    env.run(until=10)
+
+    assert comp._queue_pairs["E"]["peer_head_cache"] == 1
+    assert comp._queue_pairs["W"]["peer_head_cache"] == 1
+
+
+def test_credit_matches_by_dst_rx_base_pa_same_peer():
+    """ADR-0025 D3: credit must carry dst_rx_base_pa (the receiver-side
+    rx buffer base) so the original sender can match it against
+    qp.peer.rx_base_pa and find the correct direction. Under old
+    sender-key matching, first-match-wins would always pick E when
+    E and W share the same peer.
+    """
+    env = simpy.Environment()
+    comp = _make_pe_ipcq(env)
+    _install_same_peer_neighbors(env, comp)
+
+    # Credit corresponding to a send through W direction:
+    # - My W sent to peer's rx at 0x20_000 (qp_w["peer"].rx_base_pa)
+    # - Peer consumed it; sends credit back with dst_rx_base_pa=0x20_000
+    # - Receiver (me, the original sender) should update W's peer_tail_cache
+    credit_for_w = IpcqCreditMetadata(
+        consumer_seq=1,
+        dst_rx_base_pa=0x20_000,  # matches W's peer.rx_base_pa
+        src_sip=0, src_cube=0, src_pe=1, src_direction="E",
+    )
+    comp.credit_inbox.put(credit_for_w)
+    env.run(until=5)
+
+    assert comp._queue_pairs["W"]["peer_tail_cache"] == 1, (
+        "W's peer_tail_cache should update — credit.dst_rx_base_pa matches "
+        "W qp's peer.rx_base_pa"
+    )
+    assert comp._queue_pairs["E"]["peer_tail_cache"] == 0, (
+        "E's peer_tail_cache should NOT update"
+    )
+
+    # Second credit: for E direction
+    credit_for_e = IpcqCreditMetadata(
+        consumer_seq=2,
+        dst_rx_base_pa=0x10_000,  # matches E's peer.rx_base_pa
+        src_sip=0, src_cube=0, src_pe=1, src_direction="W",
+    )
+    comp.credit_inbox.put(credit_for_e)
+    env.run(until=10)
+
+    assert comp._queue_pairs["E"]["peer_tail_cache"] == 2
+    assert comp._queue_pairs["W"]["peer_tail_cache"] == 1