paper(figures): GQA long-context decode 4-cases comparative figures

Generated by scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py
from the bench's sweep.json (S_kv=8192, the headline 128K run was
sidestepped — Case 2's single-PE 128-tile sweep is the long pole
and shape/ordering is preserved across scales).

  gqa_decode_long_ctx_4cases_latency.png   end-to-end latency / case
  gqa_decode_long_ctx_4cases_traffic.png   ipcq / dma op-count breakdown
  gqa_decode_long_ctx_4cases_memory.png    KV bytes per cube (slide-11)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

.claude/commands/paper.md

@@ -0,0 +1,417 @@
+---
+description: Generate the 1H HW-SW codesign technical paper (PDF) — propose/confirm a TOC, then synthesize ADRs + live KernBench benchmark results into a sectioned LaTeX paper compiled with Tectonic.
+---
+
+# `/paper` — HW-SW Codesign Technical Report Generator
+
+Produces a **technical-paper-style PDF** at
+`docs/report/1H-codesign-paper/build/main.pdf`, synthesizing the
+ADR corpus, SPEC.md, and **live KernBench benchmark results** into a
+publication-shaped report on the three 1H optimizations:
+
+1. **GEMM Acceleration** using the composite command
+2. **All-Reduce Acceleration** using PE_IPCQ
+3. **Fused Grouped Query Attention** using both composite command and PE_IPCQ
+
+plus an explanation of the simulation platform itself (why KernBench
+exists, its execution model, and how it computes latency / how accurate
+that is).
+
+Audience is the **internal team** (engineers + architects), but the
+report reads like a real technical paper: **do not cite `ADR-NNNN` or
+`SPEC §X.Y` anywhere** — not in prose, not in a reference list. ADRs and
+SPEC are *grounding sources only* (use their content; never name them).
+The **References** section, if used, lists **external literature only**
+(e.g. FlashAttention, Megatron-LM, GPT-3, Llama 3) — actual papers, not
+internal docs.
+
+---
+
+## Operating principles (read first)
+
+- **Confirm the TOC before writing.** The first thing `/paper` does is
+  present the outline (saved `toc.md` if present, else the default
+  below) and **stop for the user's confirmation / adjustments**. Do not
+  generate sections until the user approves the TOC.
+- **Ground every number.** Every quantitative result in prose, tables,
+  or captions MUST trace to (a) a committed bench artifact
+  (`*.csv` / `*.json` / `*.png` under `1H_milestone_output/` or
+  `docs/diagrams/`), (b) a freshly-run bench output, or (c) a specific
+  ADR/SPEC quotation. **Never invent latencies, speedups, or
+  utilizations.** If a needed number does not exist, generate it via a
+  bench run (below) or state the gap — do not guess.
+- **Reuse existing results by default.** The three milestone benches
+  already have committed outputs. Do **not** re-run a bench unless the
+  user asks to regenerate, or the artifact is missing/stale.
+- **Design / proposal status honesty.** Some designs are `Proposed`,
+  not yet implemented (e.g. ADR-0061, ADR-0065). Present these as
+  proposed and clearly separate "modeled & measured" from
+  "designed, not yet measured." Do not report measurements for
+  unimplemented paths.
+- **English prose.** Architectural terms (composite command, PE_IPCQ,
+  GQA, TCM, SIP, CUBE, PE) used as-is.
+
+---
+
+## Invocation
+
+- `/paper` — **default**. Resolve the period, present the TOC
+  (`toc.md` if it exists, else the default TOC), and **stop** for
+  confirmation. Once the user approves (possibly after editing the TOC
+  in conversation), persist `toc.md` and proceed to full build.
+- `/paper toc` — only (re)negotiate and persist the TOC. No sections,
+  no build.
+- `/paper build` — assume the saved `toc.md` is approved; run the full
+  pipeline (sections → figures → compile PDF) without re-confirming.
+- `/paper section <id>` — regenerate a single section (`<id>` = a
+  filename stem under `sections/`, e.g. `04-allreduce`) and recompile.
+- `/paper figures` — (re)generate only the benchmark figures/data
+  (runs/refreshes benches), no prose changes, then recompile.
+- `/paper pdf` — only recompile the existing `.tex` to PDF (Tectonic).
+
+If `toc.md` does not yet exist, any mode except `/paper`/`/paper toc`
+first falls back to presenting the default TOC for confirmation.
+
+---
+
+## Output layout
+
+```
+docs/report/1H-codesign-paper/
+  toc.md                      # the agreed outline (persisted; human-editable)
+  main.tex                    # preamble + \input of each section
+  refs.bib                    # ADR/SPEC bibliography (optional; or inline refs)
+  sections/
+    00-exec-summary.tex
+    01-introduction.tex
+    02-platform.tex           # includes 3.4 Modeled Hardware Configuration
+    03-gemm.tex
+    04-allreduce.tex
+    05-gqa.tex
+    06-discussion.tex
+    07-conclusion.tex
+    08-future-work.tex
+  figures/                    # PNGs copied in from bench outputs (portable)
+  build/                      # Tectonic output; main.pdf lives here
+```
+
+Report-specific benchmark/plot harnesses (when a figure must be
+generated, see *Generating missing results*) live **isolated** under:
+
+```
+scripts/paper/                # report-only harnesses — NOT registered benches
+```
+
+This directory is intentionally separate from `src/kernbench/benches/`
+so the report's harnesses never collide with other people's benches
+(the bench registry only audits `src/kernbench/benches/`). Treat
+`scripts/paper/*` and everything under `docs/report/1H-codesign-paper/`
+as **derived artifacts** (like `docs/diagrams/`): creating/updating them
+does not require Phase 2 approval, but they MUST stay consistent with
+SPEC.md and ADRs.
+
+---
+
+## Default TOC
+
+If `toc.md` is absent, propose this and ask the user to confirm or edit:
+
+1. **Executive Summary** — the attention-optimization goal, the two
+   enabling optimizations (GEMM composite command + PE_IPCQ communication),
+   the fused GQA capstone, the headline results, and the bottom-line
+   recommendation on which HW changes are worth keeping. May run a
+   paragraph or two longer than a terse abstract; front-loads conclusions.
+2. **Introduction** — the 1H focus is **attention-kernel optimization**:
+   FlashAttention-style tiling and Grouped Query Attention (GQA), building
+   on prior work on Multi-Head Attention (MHA, studied earlier — reference
+   it as the established baseline, not re-derived here). Optimizing a fused
+   attention kernel requires two enabling optimizations, each studied in
+   its own right and then **combined inside the fused kernel**:
+   (a) **GEMM optimization** (the composite command) and
+   (b) **communication optimization** (PE_IPCQ). Frame the three result
+   sections as *two enablers → the fused GQA capstone that uses both*.
+   Motivate why HW-SW codesign (not software alone) is required to realize
+   the gains.
+3. **The KernBench Platform**
+   1. *Why KernBench* — source-level kernel execution with **no compiler /
+      software-stack dependency**; isolate algorithm-level optimization from
+      other layers so a kernel can fully exploit the hardware without E2E
+      cross-layer entanglement.
+   2. *Execution model* — discrete-event graph; `runtime API → sim_engine →
+      components`; PE pipeline + composite commands; 2-pass data/timing model.
+   3. *Latency model & accuracy* — traversal-based golden invariants
+      (every hop pays > 0, deterministic, explicit connectivity); the
+      structural CPU-dispatch cost model; per-edge BW occupancy, HBM
+      pseudo-channel parallelism, flit streaming; known simplifications &
+      what calibration anchors the constants.
+   4. *Modeled Hardware Configuration* — the shared, platform-wide
+      quantitative config used across **all** experiments, read from
+      `topology.yaml`: SIP/CUBE/PE hierarchy (counts & nesting), PE clock &
+      compute throughput, TCM/SRAM/HBM capacities & bandwidths, NOC /
+      inter-CUBE / inter-SIP link bandwidths & latencies, and the cost-model
+      constants. Per-experiment workload parameters (GEMM shapes, All-Reduce
+      topologies, GQA seq/head/user counts) stay in each evaluation section,
+      not here.
+4. **GEMM Acceleration via Composite Command** — necessity · design · results · analysis.
+5. **All-Reduce Acceleration via PE_IPCQ** — necessity · design · results · analysis.
+6. **Fused Grouped Query Attention** (composite + PE_IPCQ) — necessity · design · results · analysis.
+7. **Discussion** — which HW changes are meaningful, and why (cross-cutting).
+8. **Conclusion**.
+9. **Future Work — 2H** — add the FFN/MoE layer toward full LLM decoding;
+   analyze how compute & data should be distributed for agentic / MoE
+   workloads.
+10. **References** *(optional)* — **external literature only** (e.g.
+    FlashAttention, Megatron-LM, GPT-3, Llama 3). **No ADR/SPEC entries.**
+
+Each of §4/§5/§6 follows the same four-beat structure:
+**(a) why the optimization is needed**, **(b) the design (how it works)**,
+**(c) experimental results** (run/plot/measure), **(d) analysis & meaning**.
+
+---
+
+## Source map (ADR / SPEC → section)
+
+Use these as **grounding sources only** — read them for content, but
+**never name an ADR or SPEC in the paper**. Verify each quantitative
+claim against the actual file before asserting it (CLAUDE.md: no
+asserting architectural numbers from memory).
+
+- **Modeled Hardware Configuration (§3.4)**: `topology.yaml`
+  (authoritative for counts, clocks, BW, link params, cost-model
+  constants) is the **primary** source — read it directly and build the
+  HW table from it. Cross-check against ADR-0003 (hierarchy/scope),
+  ADR-0004 (memory BW guarantee), ADR-0033 (latency model), ADR-0064
+  (cost-model defaults). If `topology.yaml` and an ADR disagree, the
+  loaded `topology.yaml` is what the sim actually runs — report that
+  value and flag the mismatch.
+
+- **Platform — why/execution/latency**: `SPEC §0`, `SPEC §0.1`
+  (golden invariants), `SPEC §1 R1/R2/R3`, `SPEC §2.1` (graph execution);
+  ADR-0003 (hierarchy/scope), ADR-0007 (runtime↔sim_engine boundary),
+  ADR-0014 (PE pipeline / composite), ADR-0020 (2-pass data/timing),
+  ADR-0033 (latency model assumptions & simplifications),
+  ADR-0064 (structural CPU-dispatch cost model: `FIXED + logical_bytes×R`).
+- **GEMM / composite**: ADR-0014 (composite command D3.2/D3.3),
+  ADR-0042 (tile-plan generators), ADR-0044 (GEMM eval harness),
+  ADR-0064 (cost model — composite's per-command BW reduction),
+  ADR-0065 (flat-ops `CompositeCmd`, *Proposed*).
+- **All-Reduce / PE_IPCQ**: ADR-0023 (PE_IPCQ — control plane, ring
+  buffers, VC split), ADR-0025 (direction addressing), ADR-0032
+  (intercube all-reduce), ADR-0043 (allreduce eval harness),
+  ADR-0047 (AHBM CCL backend), ADR-0050 (CCL algorithm contract).
+- **GQA / fused attention**: ADR-0060 (fused GQA kernel),
+  ADR-0061 (`tl.broadcast`/`tl.repeat`, *Proposed*),
+  ADR-0062 (lazy `tl.load`), ADR-0063 (`tl.scratch_scope`),
+  ADR-0064 (cost model), ADR-0065 (flat-ops composite + `softmax_merge`,
+  *Proposed*). KO mirrors in `docs/adr-ko/`; DDD-0060/0065 design docs if present.
+- **Eval/milestone pattern**: ADR-0054 (self-contained milestone benches);
+  ADR-0045 (bench contract).
+
+---
+
+## Result map (bench → artifact → section) + run commands
+
+All runs use the repo's `topology.yaml` unless a section needs a
+specific topology. Reuse committed artifacts unless told to regenerate.
+
+### §4 GEMM — `milestone-1h-gemm`
+- Artifacts (committed): `src/kernbench/benches/1H_milestone_output/gemm/`
+  → `gemm_sweep.json`, `gemm_stage_breakdown.png`,
+  `gemm_mac_utilization_measured.png`,
+  `gemm_mac_utilization_theoretical_vs_measured.png`.
+- Regenerate: `kernbench run --topology topology.yaml --bench milestone-1h-gemm`
+  (full 24-shape sweep ≈ minutes; `MILESTONE_FAST=1` reuses committed JSON).
+
+### §5 All-Reduce — `milestone-1h-ccl`
+- Artifacts (committed): `src/kernbench/benches/1H_milestone_output/ccl/`
+  → `summary.csv`, `topology.png`, per-topology
+  `AllReduce_LRAB_*` PNGs, buffer-kind sweep
+  `AllReduce_LRAB_2Dtorus_..._with_TCM_SRAM_HBM.{png,csv}`,
+  `comparison_mesh_vs_ring_vs_2DTorus_vs_theoretical_vs_fsim.png`.
+- Regenerate: `kernbench run --topology topology.yaml --bench milestone-1h-ccl`.
+
+### §6 GQA — `milestone-gqa-headline`
+- Artifact (committed): `src/kernbench/benches/1H_milestone_output/gqa_headline/sweep.json`.
+  **Caveat:** this currently holds only `op_log_summary` **op-counts**
+  (`gemm_count`, `ipcq_copy_count`, `dma_read_count`, `dma_write_count`)
+  for 4 panels (single/multi-user × prefill/decode) — **no latency,
+  no figures**.
+- Regenerate counts:
+  `GQA_HEADLINE_RUN=1 kernbench run --topology topology.yaml --bench milestone-gqa-headline`.
+- For a real GQA **results** section, generate latency + figures via an
+  isolated report harness — see next.
+
+---
+
+## Generating missing results (isolation rule)
+
+When a section needs an artifact that does not yet exist (today: GQA
+latency and any GQA figure), create a **report-dedicated harness** under
+`scripts/paper/` — never under `src/kernbench/benches/`. Rules:
+
+- One file per concern, prefixed `paper_`, e.g.
+  `scripts/paper/paper_gqa_latency.py` (drives the engine, emits
+  latency + a baseline-vs-optimized comparison),
+  `scripts/paper/paper_plot_gqa.py` (renders figures from that data).
+- Follow the **ADR-0054 self-contained pattern**: the harness builds its
+  own `GraphEngine`(s), imports the existing GQA kernel modules
+  (`_gqa_attention_{prefill,decode}_{short,long}.py`,
+  `milestone_gqa_headline.py`), runs them, and reads latency from the
+  engine's completion timestamps / op log. It does **not** register a
+  bench and does **not** import or mutate other people's benches.
+- Output figures land in `docs/report/1H-codesign-paper/figures/` and,
+  if useful as a shared diagram, mirror into `docs/diagrams/gqa_plots/`.
+- **Only measure what the current code actually runs.** ADR-0065's GQA
+  opt2 (`softmax_merge`, two-composite decode) is *Proposed*; if it is
+  not wired in production, present it as designed-not-measured and base
+  the measured results on the implemented path. Verify implementation
+  status by reading the code before claiming a measurement.
+- These harnesses are derived-artifact tooling (no Phase 2 needed). **But**
+  if the report ever requires changing a *registered production bench*
+  under `src/kernbench/benches/` or any production module, that is a
+  non-trivial production change — STOP and follow CLAUDE.md Phase 1 →
+  approval → Phase 2.
+
+---
+
+## Procedure
+
+### Step 0 — Toolchain (Tectonic)
+
+Ensure a working Tectonic before any compile step. In order:
+
+1. If `tectonic` is on PATH (`tectonic --version` succeeds), use it.
+2. Else try, in order, whichever succeeds (capture output, don't hang):
+   - `winget install --id TectonicProject.Tectonic --silent --accept-package-agreements --accept-source-agreements`
+   - `scoop install tectonic`
+   - Direct download (no admin): fetch the latest Windows `x86_64-pc-windows-msvc`
+     release zip from `https://github.com/tectonic-typesetting/tectonic/releases`
+     via PowerShell `Invoke-WebRequest`, expand it, and place `tectonic.exe`
+     under `docs/report/1H-codesign-paper/build/.tools/`. Use that absolute
+     path for compilation.
+3. Re-verify with `tectonic --version`. If all methods fail, write the
+   `.tex` files anyway, **report the failure**, and tell the user the
+   exact command to compile manually (or to use Overleaf). Do not block
+   the rest of the pipeline.
+
+Tectonic fetches LaTeX packages on first compile (needs network once).
+
+### Step 1 — Period + TOC
+
+- Period from system date: month 1–6 → `{YYYY}-1H`, else `{YYYY}-2H`.
+  (Report dir is fixed at `1H-codesign-paper` for this half; adjust the
+  title page text with the period.)
+- If `toc.md` exists, present it; else present the **Default TOC**.
+- **Stop and ask the user to confirm or edit.** Incorporate edits.
+  Persist the agreed outline to `toc.md`. (In `/paper build` mode, skip
+  the stop and trust the saved `toc.md`.)
+
+### Step 2 — Platform section (§3, incl. §3.4 HW config)
+
+Write `sections/02-platform.tex` from the platform source map, covering
+all four sub-parts:
+
+- §3.1–3.3 (why / execution model / latency model + accuracy). Include
+  the cost-model formula `dispatch_cycles = FIXED + logical_bytes×R` with
+  the committed defaults (verify against ADR-0064 before quoting), and
+  explicitly address **how accurate** the latency is: what is modeled
+  precisely (per-edge BW occupancy, HBM pseudo-channels, flit streaming,
+  per-component overhead) and the known simplifications.
+- §3.4 *Modeled Hardware Configuration*: read `topology.yaml` directly
+  and build a compact, table-driven summary — SIP/CUBE/PE hierarchy
+  (counts & nesting), PE clock & compute throughput, TCM/SRAM/HBM
+  capacity & bandwidth, NOC / inter-CUBE / inter-SIP link bandwidth &
+  latency, and the cost-model constants (`FIXED_PER_CMD`, `R`, composite
+  cap). Every number from `topology.yaml`. This is the shared config;
+  later sections say "the modeled hardware" and only add their own
+  workload params.
+
+Do not name the source ADRs/SPEC in the prose.
+
+### Step 3 — The three result sections (§4 GEMM, §5 All-Reduce, §6 GQA)
+
+For each, in order, produce the four beats:
+
+- **(a) Necessity** — from the ADR Context: the bottleneck the
+  optimization targets (e.g. per-command dispatch overhead for GEMM;
+  collective comm cost / HoL blocking for All-Reduce; KV-load-bound
+  decode + softmax merge for GQA).
+- **(b) Design** — from the ADR Decisions: composite command tile
+  pipeline (ADR-0014); PE_IPCQ control/data split + ring buffers + VC
+  (ADR-0023); fused attention with online-softmax, lazy load, scratch
+  recycling (ADR-0060/0062/0063). Note Proposed-vs-Accepted status.
+- **(c) Results** — reuse committed artifacts (GEMM, All-Reduce). For
+  GQA, generate latency + figures via the isolated `scripts/paper/`
+  harness (above). Copy chosen PNGs into `figures/`, `\includegraphics`
+  them, and build result tables from the CSV/JSON — quoting actual
+  numbers, with a one-line note of the run command and source artifact.
+- **(d) Analysis** — interpret: what the curve/table shows, why the HW
+  feature produces it, where it saturates or fails to help, and the
+  cost (area/complexity) implied. Tie back to the cost model.
+
+### Step 4 — Discussion, Conclusion, Future Work (§6/§7/§8)
+
+- **Discussion**: synthesize across the three — which HW changes
+  (composite command, PE_IPCQ, the cost-model-exposed BW reduction)
+  carry their weight, and under what regimes.
+- **Conclusion**: the codesign thesis, supported by the measured
+  results; state plainly which HW changes are meaningful.
+- **Future Work (2H)**: FFN / MoE layer toward full LLM decoding; how
+  compute & data should be distributed for agentic / MoE workloads
+  (expert routing, token dispersion, all-to-all vs all-reduce, capacity).
+  Frame as the next codesign questions, not claims.
+
+### Step 5 — Assemble & compile
+
+- `main.tex`: article/IEEEtran-style preamble, title with the period,
+  author, executive-summary input, then `\input{sections/...}` in TOC
+  order, then references. Keep the preamble minimal (graphicx, booktabs,
+  hyperref, amsmath) so Tectonic compiles cleanly.
+- Compile: `tectonic main.tex --outdir build` (or the resolved Tectonic
+  path) from `docs/report/1H-codesign-paper/`. Re-run if references need
+  a second pass. Confirm `build/main.pdf` exists and report its size.
+
+### Step 6 — Chat report
+
+Emit to chat (not to any file):
+
+```
+## /paper — Build Summary
+**PDF:** docs/report/1H-codesign-paper/build/main.pdf  (<N> pages, <size>)
+**Toolchain:** tectonic <version | how installed>
+**Sections written:** <list>
+**Benches: reused vs regenerated**
+  - GEMM: <reused committed | regenerated via milestone-1h-gemm>
+  - All-Reduce: <reused | regenerated>
+  - GQA: <generated via scripts/paper/paper_gqa_latency.py | counts-only>
+**Figures included:** <list of figures/*.png>
+**Grounding check:** every result number traced to <artifact list>
+**Proposed-not-measured items flagged:** <e.g. ADR-0065 opt2, ADR-0061>
+**Open gaps / suggested next runs:** <…>
+```
+
+---
+
+## Constraints (do not violate)
+
+1. **TOC first.** Never write sections before the TOC is confirmed
+   (default mode). `toc.md` is the contract.
+2. **No fabricated numbers.** Every quantitative claim traces to a
+   committed artifact, a fresh bench run, or an ADR/SPEC quote. Verify
+   architectural constants against the source file, not memory.
+3. **Reuse by default.** Do not re-run a bench unless asked or the
+   artifact is missing/stale.
+4. **Isolation.** Report harnesses live only in `scripts/paper/`; never
+   register a bench, never touch other people's benches.
+5. **Production-change gate.** Any change to a registered bench or
+   production module triggers CLAUDE.md Phase 1 → approval → Phase 2.
+   Report tooling and `docs/report/...` are derived artifacts (exempt).
+6. **Status honesty.** Mark Proposed designs as proposed; never present
+   an unimplemented path as measured.
+7. **No internal references.** Never name `ADR-NNNN` or `SPEC §X.Y` in
+   prose, captions, or the reference list — they are grounding sources
+   only. A References section, if present, holds **external literature
+   only**. Hardware numbers come from `topology.yaml`.
+8. **Determinism.** Sections in TOC order; same inputs → same paper.
+9. **English prose**, internal-audience tone.
+```

CLAUDE.md

@@ -371,6 +371,13 @@ Concrete forms that Part 1's *Verification Plan* MUST take in this repo:
 - `kernbench run --device <id>` runs the benchmark on a single device.
 - Omitting `--device` runs the benchmark on all devices discovered in the topology (logically parallel).
 - Device enumeration is handled by the CLI only; benchmarks MUST remain single-device.
+- **Eval-bench exception (ADR-0054)**: a *milestone / eval bench*
+  (`milestone-1h-*`) may drive many configurations and build its own
+  per-config engines to regenerate a domain's full result + figure set; it
+  ignores `--device` and submits a sentinel tensor to satisfy the
+  "must submit ≥1 request" contract (ADR-0045 D4). This is the eval-harness
+  carve-out to the single-device rule, alongside the ADR-0024 multi-SIP CCL
+  exception.
 
 ## Derived Artifacts (Clarification)
 

docs/adr-ko/ADR-0026-par-dppolicy-intra-device.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Accepted (Revision 5 — Phase 2 landed 2026-04-14, 523 passed + 1 strict xfail)
+Accepted (Revision 6 — cube_start 추가 2026-06-04; Revision 5 landed 2026-04-14)
 
 ## Context
 
@@ -30,9 +30,11 @@ class DPPolicy:
     pe: Literal["replicate", "column_wise", "row_wise"] = "replicate"
     num_pes: int | None = None
     num_cubes: int | None = None
+    cube_start: int = 0   # Revision 6 — SIP 내 첫 cube 인덱스; D8 참조
 ```
 
 제거되는 필드: `sip`, `num_sips`.
+추가되는 필드 (Revision 6): `cube_start` — D8 참조.
 
 ### D2. `ShardSpec` — structural (sip, cube, pe) 좌표, `pe_index` 완전 제거
 
@@ -114,7 +116,7 @@ def resolve_dp_policy(
         for ls in local_shards:
             all_shards.append(ShardSpec(
                 sip=target_sip,                              # from caller (current_device)
-                cube=cube_id,                     # local within SIP
+                cube=policy.cube_start + cube_id,            # Rev 6: D8 cube_start만큼 shift; 기본 0 = local within SIP
                 pe=ls.local_pe,                              # local within cube (explicit name)
                 offset_bytes=cube_offset + ls.offset_bytes,
                 nbytes=ls.nbytes,
@@ -224,6 +226,47 @@ KernBench는 사내 프로젝트로 call site가 한정되어 있어 한 번에
 **Silent drift 차단**이 property 완전 제거의 주된 이점: global flat을 기대한
 코드가 SIP-local 결과를 받아 조용히 잘못된 인덱싱을 할 가능성 제거.
 
+### D8. SIP 내 분리된 cube sub-mesh를 위한 `cube_start: int = 0` 추가
+
+**Revision 6 추가 사항 (2026-06-04)**.
+
+**목적**: 한 SIP 내에서 분리된 cube sub-mesh 지정 (예: 0..7과 8..15을 함께
+사용). GQA Llama-70B 8-KV-group 헤드라인 타겟 — SIP당 2×4 KV-group 2개 ×
+4 SIP = 64 cubes — 의 두 번째 KV-group을 기본값 0..7이 아닌 cubes 8..15에
+배치하기 위해 필요.
+
+**Semantics**: `resolve_dp_policy`가
+`cube = policy.cube_start + cube_id`로 ShardSpec 생성 (cube_id는 launch 내에서
+`0..num_cubes-1` 순회). 선택된 cube는 target SIP 내의
+`[cube_start, cube_start + num_cubes)` 범위.
+
+**기본값**: `cube_start = 0`은 기존 모든 호출 사이트 동작을 bit-for-bit
+보존 (이전처럼 `ShardSpec.cube ∈ [0, num_cubes)`). CCL milestone bench의
+full-SIP `DPPolicy(num_cubes=16)`은 그대로 cubes 0..15 생성.
+
+**제약**: intra-device 불변성 보존.
+`cube_start ∈ [0, cubes_per_sip)`이며
+`cube_start + num_cubes ≤ cubes_per_sip`. SIP 경계 교차는 여전히
+`ahbm.set_device(rank)` (ADR-0024)의 책임.
+
+**왜 scalar인가** (2D `cube_mesh_origin`이나 임의 `cube_ids` list가 아닌):
+consumer 커널 (예: `_attention_mesh_mlo_2d`)이 row-major contiguous cube를
+가정; scalar `cube_start`는 `num_cubes`와 자연스럽게 쌍을 이룸 (범위 =
+`[start, start + count)`); 기본값 0에서 마이그레이션 부담 없음. 비연속
+케이스가 필요해지면 향후 더 일반적인 설계를 위에 얹을 수 있음.
+
+**하위 호환성**: 기본값이 있는 가산적 변경. 기존 호출 사이트의 변경을
+강제하지 않음. D7의 "breaking change" 입장은 원래 sip/num_sips 제거에만
+해당하며, `cube_start`는 기존 호출자를 깨지 않음.
+
+**커널 측 참고**: kernbench의 `tl.program_id(axis=1)`은 launch-local rank가
+아닌 물리 cube id를 반환 (ADR-0022). `program_id(axis=1)`에서 ring 위치를
+도출하는 커널은 `cube_start > 0`일 때 launch-local rank를 복원하기 위해
+`cube_start`를 빼야 함 — 그렇지 않으면 out-of-bounds sub-mesh 위치를
+계산하여 deadlock 발생. 참고 패턴은 `_attention_mesh_mlo_2d.py` 참조.
+ADR-0022의 향후 개정에서 `program_id(axis=1)`을 launch-local rank semantics로
+이동하면 커널 측 명시적 차감이 제거될 수 있음.
+
 ## Dependencies
 
 - **ADR-0024** (launcher): `set_device(rank)` 및 current-device scoping이

docs/adr-ko/ADR-0043-eval-allreduce-harness.md

@@ -7,6 +7,11 @@ Accepted
 `tests/sccl/` 평가 하니스를 문서화한다; 구현과 대조 검증 완료
 (상수, 파일 집합, 스윕 차원을 교차 확인).
 
+**ADR-0054로 개정됨**: 드라이버 코어, sweep, renderer가 `milestone-1h-ccl`
+bench(단일 home)로 이동했다; `tests/sccl/_allreduce_helpers.py`는 이제 거기서
+re-export한다(pytest 전용 param 빌더 + `_run_distributed` wrapper는 로컬
+유지). figure 테스트는 변경 없음.
+
 ## Context
 
 ADR-0032는 intercube all-reduce *알고리즘*을 정의하고, ADR-0023/0024/0027은

docs/adr-ko/ADR-0044-eval-gemm-harness.md

@@ -8,6 +8,12 @@ GEMM 평가/특성화 하니스를 문서화한다; 구현과 대조 검증 완
 (상수, tile 크기, figure 집합, script↔test 분할을 교차 확인). D5/D6
 caveat은 부정확이 아니라 기록된 한계다.
 
+**ADR-0054로 개정됨**: sweep + renderer가 `milestone-1h-gemm` bench(단일
+home)로 이동했다; `scripts/gemm_sweep.py`와 `tests/gemm/`는 이제 거기서
+re-export한다. D1/D2의 "데이터 생성은 수동 script / 무거운 작업은 opt-in"은
+평가-bench 패턴으로 대체된다(하나의 bench가 전부 재생성;
+`MILESTONE_FAST=1`은 committed JSON 재사용).
+
 ## Context
 
 ADR-0014(PE pipeline)와 ADR-0042(tile-plan generator)는 GEMM *구현*을

docs/adr-ko/ADR-0045-prog-bench-module-contract.md

@@ -0,0 +1,265 @@
+# ADR-0045: Bench Module Contract — registration, dispatch, and authoring
+
+## Status
+
+Accepted (2026-05-21).
+
+`src/kernbench/benches/` 패키지의 등록 메커니즘(@bench), CLI 디스패치 경로
+(`kernbench run/list`), 그리고 새 bench 모듈 작성 시 따라야 할 계약을 통합
+정의한다. ADR-0010 (CLI surface)이 `kernbench list/run` 인터페이스를 명세하나,
+**bench가 어떻게 등록되고 어떤 함수 시그너처를 따라야 하는가**는 ADR 레벨에
+없었음.
+
+**ADR-0054로 확장됨**: D5의 단일 구성 규칙에 세 번째 패턴이 추가된다 —
+*평가 bench*(예: `milestone-1h-*`)는 여러 구성을 구동하고, 구성별 자체 엔진을
+빌드하며, D4를 만족시키기 위해 sentinel 텐서를 제출한다.
+
+## First action (제일 처음에 하는 일)
+
+`kernbench.benches` 패키지가 임포트되면 `__init__.py` 가 즉시
+`_eager_import_and_audit(__path__, __name__)` 를 호출한다. 이 함수의 첫 일은
+패키지 디렉터리 안의 모든 형제 모듈을 `pkgutil.iter_modules(__path__)`로 나열한
+뒤, 다음 두 조건을 만족하지 않는 모듈을 모두 `importlib.import_module(...)`로
+**즉시 로드**하는 것이다:
+
+- 이름이 `registry` 인 경우 (인프라 자체)
+- 이름이 `_` 로 시작하는 경우 (helper 모듈)
+
+임포트 시점에 각 모듈 안의 `@bench(name=..., description=...)` 데코레이터가
+실행되어 `_PENDING` 리스트에 `(name, description, fn)` 튜플이 append 되고,
+`_REGISTERED_MODULES` 셋에 `fn.__module__` 가 추가된다.
+
+전체 임포트가 끝나면 `_audit_modules(imported, _REGISTERED_MODULES)` 가 호출되어,
+**임포트는 되었지만 @bench를 한 번도 호출하지 않은 모듈**이 있으면
+`RuntimeError("Bench module(s) missing @bench decorator: ...")` 가 즉시 발생한다.
+이 audit이 통과한 시점에 인덱스 할당은 아직 일어나지 않은 상태이며, 첫
+`list_all()` / `resolve(...)` 호출 시 `_finalize()` 가 이름 알파벳 정렬 순으로
+1-based index를 부여한다.
+
+즉, **bench 인프라의 첫 일은 "패키지 디렉터리의 모든 비-helper 모듈을 임포트
+하고, 각 모듈이 최소 한 번 @bench를 호출했는지 감사하는 것"** 이다.
+
+## Context
+
+`src/kernbench/benches/` 는 현재 8개의 bench 모듈을 보유한다 (`ccl_allreduce`,
+`gemm_single_pe`, `gpt3_qkv`, `ipcq_allreduce`, `matmul_composite`, `qkv_gemm`,
+`qkv_gemm_multi_pe`, `va_offset_verify`). 모든 bench는 다음 통합 흐름을 따른다:
+
+```
+kernbench run --topology <T> --bench <N>
+   ↓
+cli/main.py::cmd_run
+   ↓  resolve_topology(T)  + resolve(N)  + resolve_device(device_arg)
+   ↓
+runtime_api/bench_runner.py::run_bench(topology, bench_fn, device, engine_factory)
+   ↓  engine_factory(topology, device) → GraphEngine
+   ↓  RuntimeContext(engine, target_device, correlation_id, spec)
+   ↓
+bench_fn(ctx)        ← bench가 정의한 run(torch) 가 호출됨
+   ↓  ctx.empty/zeros/from_numpy/launch/distributed.* 등을 통해 submit
+   ↓
+ctx.wait_all()       ← 미완료 핸들이 있으면 drain
+   ↓
+BenchResult(completion, correlation_id, trace, traces, engine)
+```
+
+ADR-0010 은 CLI 표면만 다루고 (`run/list/probe/web`), ADR-0007 은 runtime API ↔
+sim_engine 책임 경계만 다룬다. 정작 "새 bench 파일을 추가하려면 어떤 모양으로
+써야 하는가"는 코드 컨벤션만으로 추적해야 한다. 결과적으로:
+
+- @bench 데코레이터의 호출 규약 (kebab-case 이름, non-empty description)이
+  코드에만 존재.
+- bench 함수 시그너처 (`def run(torch)`) 가 사실상 컨벤션인데, CLI 디스패치 측이
+  `spec.run` 을 호출한다는 사실로 강제되고 있음.
+- 신규 bench 추가자가 "helper 모듈은 `_` 접두로 분리해야 한다"는 것을 audit
+  RuntimeError를 받아본 뒤에야 학습.
+- single-device 컨벤션 (CLAUDE.md Part 2 CLI Semantics)이 bench 작성자 관점에서
+  어디까지 적용되는지 (CCL 멀티-SIP bench는 예외인가?) 명확하지 않음.
+
+이 ADR이 이런 모호함을 한 곳에 정리한다.
+
+## Decision
+
+### D1. @bench 데코레이터 계약
+
+```python
+from kernbench.benches.registry import bench
+
+@bench(name="my-bench", description="Short, complete-sentence description.")
+def run(torch):
+    ...
+```
+
+- `name`: kebab-case 문자열. 정규식 `^[a-z][a-z0-9]*(-[a-z0-9]+)*$` 통과 필요.
+  소문자/숫자/대시만 허용; 밑줄(`_`) 금지; 알파벳으로 시작.
+- `description`: non-empty 문자열 (strip 후 길이 > 0). CLI `list` 출력에 그대로
+  표시됨.
+- 데코레이터는 **fn을 변형 없이 반환**한다 — 즉 직접 호출도 가능. 부수효과로
+  `_PENDING` 에 등록만 추가한다.
+
+위 두 규칙 위반은 즉시 `ValueError`. duplicate name은 `_finalize()` 시점에
+`RuntimeError("duplicate bench name: ...")` 로 잡힌다.
+
+### D2. 모듈 파일 컨벤션
+
+`src/kernbench/benches/<slug>.py` 는 다음 중 하나여야 한다:
+
+- **bench 모듈**: 최상위 임포트 경로에서 적어도 한 번 `@bench(...)` 가 실행되어
+  최소 하나의 bench를 등록한다.
+- **helper 모듈**: 파일명이 `_` 로 시작 (예: `_shared_helpers.py`). `iter_modules`
+  순회에서 스킵된다.
+
+audit (`_audit_modules`) 는 helper가 아닌데도 @bench를 호출하지 않은 모듈을
+허용하지 않는다. 의도된 결과: 새 파일을 `benches/` 에 추가하기만 하면 자동
+등록되며, helper와의 구분은 **파일명 접두사** 하나로 명확하게 표시된다.
+
+### D3. bench 함수 시그너처는 `def run(torch)` 다
+
+데코레이터는 함수 이름을 강제하지 않지만, **CLI 디스패치는 `spec_entry.run`
+(즉 데코레이트된 callable) 을 호출**한다. 따라서 컨벤션은:
+
+- 함수 이름: `run`. 다른 이름으로 데코레이트해도 동작은 하지만 readability /
+  grep-ability 측면에서 항상 `run`.
+- 인자: 단일 위치 인자 `torch`. 실제로는 `RuntimeContext` 인스턴스이며 PyTorch
+  스타일의 namespace (zeros/empty/launch/distributed/...)를 노출한다 (ADR-0024 D3).
+- 반환값: 임의 (`Any`). 현재 `run_bench` 는 반환값을 무시하고 `ctx.handles()` /
+  `engine.get_completion()` 로 완료를 추적한다.
+
+`torch` 이름은 PyTorch 호환 idiom을 흉내내기 위함이며, 실제로 PyTorch 모듈이
+들어오는 것은 아니다 (ADR-0024 의 "rank = SIP" launcher 컨벤션과 정렬).
+
+### D4. bench는 최소 한 번의 submit을 수행해야 한다
+
+`run_bench` 는 `ctx.handles()` 가 비어 있는 경우 BenchResult.completion 을
+`ok=False, error_code="NO_REQUESTS"` 로 반환한다. 따라서 의미 있는 bench는
+다음 중 하나 이상을 호출해야 한다:
+
+- 텐서 생성 API: `torch.zeros(...)`, `torch.empty(...)` — 내부적으로
+  `MmuMapMsg` 와 (zeros 의 경우) `MemoryWriteMsg` 가 submit 됨.
+- 커널 실행 API: `torch.launch(name, fn, *args)` — `KernelLaunchMsg` 를 SIP 별로
+  submit.
+- (예외) 빈 placeholder bench: `ipcq_allreduce.py` 처럼 `print(...)` 만 하는
+  스텁은 NO_REQUESTS 결과를 받게 됨. CI 측에서 placeholder임을 인지하고 별도
+  처리하는 것을 가정한다.
+
+### D5. 단일-디바이스 컨벤션 + 멀티-SIP 예외 (ADR-0024/0027)
+
+CLAUDE.md Part 2 CLI Semantics 가 명시하는 **"benchmarks MUST remain
+single-device"** 컨벤션은 다음과 같이 해석된다:
+
+- **일반 bench (single-SIP 사용)**: `dp = DPPolicy(...)` 로 텐서 placement를
+  정의하고 `torch.launch(...)` 로 커널 발사. SIP 인덱스는 `--device` 가
+  결정한다 (CLI 측 책임).
+- **CCL bench (멀티-SIP 사용)**: 예외적으로 `torch.distributed.init_process_group
+  (backend="ahbm")` + `torch.multiprocessing.spawn(_worker, ..., nprocs=ws)` 로
+  rank = SIP 패턴 (ADR-0024 D3) 을 따른다. `--device` 는 무시되며 (또는
+  `all` 로 가정), 각 spawned worker가 `torch.ahbm.set_device(rank)` 로 자신의
+  SIP를 바인딩한다.
+
+이 두 패턴 외의 멀티-디바이스 호출 (예: 한 bench 함수가 동일 process에서 여러
+SIP을 직접 launch) 은 본 ADR이 금지한다. CLI 가 `--device all` 로 호출되어도
+bench는 한 번만 실행되며, 그 안에서 멀티-SIP을 다루려면 D5의 두 번째 패턴을
+사용한다.
+
+### D6. 이름·인덱스 해석 (`resolve`)
+
+`resolve(identifier: str)` 는 다음 순서로 BenchSpec을 반환한다:
+
+1. `identifier.isdigit()` → 정수 변환 후 `_REGISTRY` 의 entries에서 `index ==`
+   인 spec 반환. 없으면 `ValueError("No bench with index ..."`)`.
+2. `identifier in _REGISTRY` → 직접 lookup.
+3. 그 외 → `ValueError("Unknown bench ...")`.
+
+빈/공백 identifier 는 `ValueError("bench identifier must be a non-empty string.")`.
+
+CLI 는 `--bench` 의 인자를 그대로 `resolve` 에 넘긴다. 따라서 사용자는
+`kernbench run --bench gemm-single-pe` 또는 `kernbench run --bench 2` 형식 모두
+사용 가능.
+
+### D7. 인덱스는 안정 API가 아니다
+
+`_finalize()` 가 `_PENDING` 을 **이름 알파벳 정렬** 후 1-based index를 부여하므로,
+새 bench 가 추가되면 기존 bench의 index가 밀릴 수 있다. 따라서:
+
+- 사람-친화적 인터랙티브 사용: 인덱스 OK.
+- 스크립트 / CI 자동화: 반드시 이름을 사용한다.
+
+이 사실은 `registry.py` 모듈 docstring 에 명시되어 있다.
+
+### D8. RuntimeContext 가 bench에 노출하는 표면
+
+bench 함수가 `torch` 파라미터를 통해 정상적으로 사용할 수 있는 표면:
+
+- **텐서 생성**: `torch.empty(shape, dtype=..., dp=DPPolicy(...), name=...)`,
+  `torch.zeros(...)`, `torch.from_numpy(arr)`. 모두 host-side 메타 + 디바이스
+  배포 (MmuMap + MemoryWrite) 를 submit 한다.
+- **커널 발사**: `torch.launch(kernel_name, kernel_fn, *args)` —
+  `(Tensor, int, float)` 위치 인자를 `TensorArg` / `ScalarArg` 로 변환하여
+  SIP 별 `KernelLaunchMsg` 발행 후 drain.
+- **동기화**: `torch.wait(handle)`, `torch.wait_all()` (run_bench 가 자동 호출).
+- **분산**: `torch.distributed.init_process_group(backend="ahbm")`,
+  `torch.distributed.get_world_size()`, `torch.distributed.all_reduce(t, op=...)`
+  (ADR-0024/0027).
+- **멀티-프로세스 (rank=SIP)**: `torch.multiprocessing.spawn(_worker, ..., nprocs=ws)`
+  (ADR-0024 D3 / ADR-0027).
+- **디바이스 바인딩**: `torch.ahbm.set_device(rank)` 또는
+  `torch.accelerator.set_device_index(rank)` (둘 다 같은 namespace를 가리킴).
+- **IPCQ 설치**: `torch.install_ipcq(algorithm=..., ccl_yaml=...)` (ADR-0023 D10).
+- **스펙 조회**: `torch.spec` — 토폴로지 빌더가 만든 dict (시스템·cube_mesh·HBM
+  파라미터 등). bench가 toplogy.yaml 파라미터에 의존하지 않게 짜기 위함.
+
+bench는 위에 열거되지 않은 RuntimeContext 의 private 멤버 (`_handles`, `_traces`,
+`_allocators` 등) 에 직접 접근해선 안 된다. ADR-0007 의 layer boundary 정신과
+정렬: bench → runtime API → sim_engine 한 방향만 허용.
+
+### D9. 환경 변수로 파라미터화는 허용된다
+
+`matmul_composite.py` 처럼 `os.environ.get("MATMUL_M", ...)` 등으로 bench
+파라미터를 외부에서 주입하는 패턴은 허용한다. 이유:
+
+- bench 함수 시그너처는 D3 에 의해 `def run(torch)` 로 고정되어 있어 위치/키워드
+  인자로 파라미터를 받기 곤란.
+- 환경 변수 패턴은 `MATMUL_VARIANT` 같은 운영-시 스윕을 위한 자연스러운 hook.
+- `scripts/gemm_sweep.py` 같은 외부 드라이버 (ADR-0044) 가 이 hook을 사용한다.
+
+단, 환경 변수가 bench의 동작을 바꾼다면 모듈 docstring 에 모든 변수를 명시할 것
+(matmul_composite.py 가 그 예시).
+
+## Alternatives Considered
+
+### A1. 명시적 manifest 파일 (YAML)에 bench 목록 두기
+
+기각. @bench 데코레이터 + audit 패턴은 "파일 추가 = 자동 등록" 을 보장하여 신규
+bench 작성자의 인지 비용을 한 곳 (파일 작성)으로 집중시킨다. 별도 manifest는
+유지보수 측에서 drift 위험이 크고, helper 분리는 이미 `_` 접두로 명확하다.
+
+### A2. bench 함수 이름을 데코레이터 인자로 받기 (`@bench(name=..., entry="run_xxx")`)
+
+기각. 디스패치 측에서 `spec.run` 하나만 호출하면 되는 단순함을 깬다. `run` 컨벤션
+하나로 충분하며, 변종이 필요하면 같은 모듈에 여러 함수를 등록하면 된다 (각각
+@bench 데코레이트).
+
+### A3. CCL bench를 위한 별도 `@multi_device_bench` 데코레이터
+
+기각. D5에서 명시한 두 패턴 (single + ADR-0024 멀티-SIP) 만으로 현재 8개 bench가
+모두 표현 가능. 별도 데코레이터는 디스패치 측에서 분기를 강제하여 복잡도를 늘리며,
+멀티-SIP 사용 의도는 bench 함수 본문의 `init_process_group(...)` 호출로 충분히
+드러난다.
+
+### A4. 인덱스를 안정 API로 만들기 (등록 순서 / explicit index= 인자)
+
+기각. D7에서 명시한 trade-off — 사용자 친화성 (알파벳 정렬된 인덱스가 list 출력
+에서 자연스럽게 1, 2, 3...) 우선. 스크립트는 이름으로 지정하면 충분.
+
+## Consequences
+
+- "bench 추가 방법" 이 한 ADR로 정리됨 → 신규 작성자가 코드 grep 없이 D1-D3,
+  D8 만 따르면 됨.
+- helper 모듈을 `_` 접두로 분리하는 패턴이 ADR-level에서 정당화되어, 향후
+  `benches/_*.py` 식의 공유 helper 작성이 자유로워짐.
+- CLAUDE.md Part 2 CLI Semantics 의 single-device 컨벤션이 멀티-SIP CCL bench
+  와 모순되지 않음을 D5 가 명시 — 둘은 직교한다.
+- ADR-0044 (GEMM eval harness) 의 `scripts/gemm_sweep.py` 가 환경 변수 hook을
+  사용하는 근거 (D9) 가 본 ADR에 굳어짐.
+- 인덱스가 불안정함 (D7) 이 명시되어, CI 측 `kernbench run --bench 3` 같은
+  코드는 본 ADR 수락 직후 점검 대상.

docs/adr-ko/ADR-0046-prog-tl-context-contract.md

@@ -0,0 +1,307 @@
+# ADR-0046: TLContext — Kernel-side `tl.*` API Contract
+
+## Status
+
+Accepted (2026-05-22).
+
+`src/kernbench/triton_emu/` 의 `TLContext` 가 노출하는 `tl.*` primitive
+집합과 그 의미, 그리고 두 실행 모드 (command-list / greenlet runner) 의
+계약을 명시한다. ADR-0014/0020 가 PE 파이프라인과 2-pass 실행 모델을
+정의하나, **bench 의 kernel 함수가 호출하는 `tl.*` 표면 자체**는 ADR-level
+에 정리되어 있지 않았다.
+
+## First action (제일 처음에 하는 일)
+
+`TLContext(pe_id, num_programs, dispatch_cycles, runner, cube_id, num_cubes,
+scratch_base, scratch_size)` 생성 시 가장 먼저 다음 6개 필드를 초기화한다:
+
+- `self._pe_id`, `self._num_programs`, `self._cube_id`, `self._num_cubes` —
+  `tl.program_id` / `tl.num_programs` 가 반환할 값.
+- `self._dispatch_cycles` — 모든 `tl.*` API 호출 시작에서 자동으로 발행될
+  `PeCpuOverheadCmd(cycles)` 의 cycle 수.
+- `self._runner` — `KernelRunner` 인스턴스 (있으면 greenlet 모드, 없으면
+  command-list 모드).
+- `self._commands: list[PeCommand] = []` — command-list 모드에서 누적할
+  command 시퀀스.
+- `self._handle_counter = 0`, `self._completion_counter = 0` — 새 TensorHandle /
+  CompletionHandle id 생성용.
+- `self._scratch_base`, `self._scratch_size`, `self._scratch_cursor = 0` —
+  PE-로컬 scratch 영역 (math/dot/composite 의 output handle 주소 할당용).
+
+즉, **TLContext 의 첫 일은 "이 kernel 인스턴스가 어디서 (sip/cube/pe) 어떤
+규모 (num_programs/num_cubes) 로 실행되며, 어느 모드 (runner 유무) 로
+명령을 발사할지 메타데이터를 채우는 것"** 이다. 이 시점에 SimPy event 는
+없으며 command 도 발사되지 않는다.
+
+런타임 첫 동작은 kernel 함수가 `tl.<api>()` 를 처음 호출할 때 발생한다.
+모든 `tl.*` API 의 표준 entry 동작은:
+
+1. `self._emit_dispatch_overhead()` 호출 — `dispatch_cycles > 0` 인 경우
+   `PeCpuOverheadCmd(dispatch_cycles)` 를 즉시 `_emit`.
+2. API 별 처리 (TensorHandle 생성, command 구성).
+3. `self._emit(cmd)` — runner 모드면 greenlet.switch 로 SimPy 측에 cmd 전달,
+   아니면 `self._commands` 에 append.
+
+## Context
+
+`tl.*` 표면은 `TLContext` 가 노출하는 메소드들로 구성되며, kernel 함수가
+받는 `tl` 매개변수가 이 객체다. 사용자(bench 작성자) 입장에서 보이는
+contract:
+
+- 어떤 primitive 가 있는가
+- 각 primitive 가 어떤 데이터 흐름을 발생시키는가 (DMA / compute / IPCQ /
+  metadata-only)
+- TensorHandle 의 `space` 와 `addr` 가 어떻게 결정되는가
+- command-list 모드와 greenlet 모드의 차이
+
+ADR-0014 (PE pipeline) 가 PE_SCHEDULER 가 받는 PeCommand 들을 정의하나,
+`tl.*` 가 이들을 어떻게 emit 하는지는 코드 컨벤션에만 존재한다. 또한
+ADR-0020 (2-pass data execution) 가 greenlet 모드의 존재를 D3 에서
+언급하나, runner / non-runner 두 경로의 시그너처 차이 (return value 처리)
+는 ADR-level 에 명시되어 있지 않다. 이 ADR 이 그 빈자리를 채운다.
+
+## Decision
+
+### D1. `tl` 매개변수는 `TLContext` 인스턴스다
+
+bench 의 kernel 함수는 다음 시그너처를 따른다:
+
+```python
+def _kernel(arg1, arg2, ..., tl, **kwargs):
+    ...
+```
+
+`tl` 의 정체는 `kernbench.triton_emu.tl_context.TLContext` 인스턴스이다.
+real Triton 의 `triton.language` 모듈을 흉내내기 위한 이름이며, real
+Triton 모듈이 들어오는 것은 아니다.
+
+kernel 함수는 일반 Python 함수이며 `yield` / `async` 가 없다. `tl.*`
+호출이 SimPy event 를 발생시키지만, 호출자(kernel) 쪽에서는 동기 호출처럼
+보인다 — greenlet 모드에서 KernelRunner 가 SimPy ↔ kernel 사이를 중계
+하기 때문 (ADR-0020 D3).
+
+### D2. 두 실행 모드 — command-list / greenlet runner
+
+- **command-list 모드 (`runner is None`)**: `tl.*` 호출이 `self._commands`
+  리스트에 PeCommand 를 누적. DMA / GEMM / Math 가 실제 SimPy 시간을
+  소비하지 않으며, return value 가 metadata-only TensorHandle (data=None) 다.
+  이후 PE_SCHEDULER / sim_engine 가 command 시퀀스를 시간상 재생.
+
+- **greenlet runner 모드 (`runner is not None`)**: `tl.*` 호출이
+  `self._emit(cmd)` 를 통해 `runner.switch_to_simpy(cmd)` 로 부모 greenlet
+  (SimPy) 으로 컨트롤을 넘김. 부모는 cmd 를 컴포넌트에 분배하여 SimPy 시간을
+  소비한 뒤, DMA read 의 경우 실제 numpy 데이터를 반환. kernel 은 그
+  결과를 받아 다음 line 으로 진행 (ADR-0020 D3 의 데이터 인지 실행 모델).
+
+mode 선택은 KernelRunner 인스턴스를 TLContext 에 주입하는지 여부로 결정
+되며, `tl.*` 메소드들은 이 차이를 인지하지 않고 `_emit()` 헬퍼를 통해
+일관되게 동작한다.
+
+### D3. Primitive 카테고리
+
+#### D3.1. Reference (no DMA, metadata only)
+
+- `tl.ref(ptr, shape, dtype="f16") -> TensorHandle`: HBM 데이터를 참조하는
+  핸들만 만들고 DMA 는 발행하지 않음. composite scheduler 가 per-tile 로
+  스트리밍할 때 사용 (예: GEMM 의 b 피연산자).
+
+#### D3.2. Data movement (blocking, DMA engine)
+
+- `tl.load(ptr, shape, dtype="f16") -> TensorHandle`: HBM → 결과 핸들.
+  `DmaReadCmd` 발행. greenlet 모드에서는 결과 핸들의 `.data` 에 실제
+  numpy 배열 첨부; command-list 모드에서는 placeholder. 반환 핸들의
+  `space="hbm"`, `pinned=True`.
+- `tl.store(ptr, handle) -> None`: TCM → HBM. `DmaWriteCmd` 발행. greenlet
+  모드에서는 `handle.data` 가 있을 때만 `_store.write("hbm", ptr, data)` 를
+  먼저 호출 (visibility = issue time, ADR-0020 D3).
+
+#### D3.3. GEMM / compute (blocking)
+
+- `tl.dot(a, b) -> TensorHandle`: `a @ b`. 두 피연산자는 TCM 이어야 하며,
+  shape (M,K) × (K,N) → (M,N). `GemmCmd` 발행, output handle 은
+  `_make_compute_out(shape, dtype)` 로 PE-로컬 scratch 에 할당.
+- `tl.composite(op, a, b=None, out_ptr=0, math_op=None, epilogue=None,
+  acc_dtype=None, tile_shape=None) -> CompletionHandle`: 비차단(non-blocking)
+  tiled pipeline. `CompositeCmd` 발행. `epilogue` 는 dict list, 각 dict 는
+  `"op"` 키 + op-specific 필드 + 옵션 `"scope"` (k_tile / output_tile);
+  unknown op 나 missing field 는 즉시 ValueError. 반환된 CompletionHandle 은
+  `tl.wait(h)` 로 동기화.
+
+#### D3.4. Math: unary (blocking)
+
+- `tl.exp(x)`, `tl.log(x)`, `tl.sqrt(x)`, `tl.abs(x)`, `tl.sigmoid(x)`,
+  `tl.cos(x)`, `tl.sin(x)` — 모두 `MathCmd(op=<name>, inputs=(x,), out=)`
+  발행. `out` 은 동일 shape/dtype 의 scratch 할당.
+
+#### D3.5. Math: binary (blocking)
+
+- `tl.maximum(a, b)`, `tl.minimum(a, b)` — `_binary_math`.
+- `tl.fma(a, b, c)` — `a*b + c`. inputs 3개.
+- `tl.clamp(x, min, max)` — `MathCmd(op="clamp", inputs=(x, min, max))`.
+- `tl.where(cond, a, b)` — `MathCmd(op="where", inputs=(cond, a, b))`.
+- `tl.softmax(x, axis=-1)` — 단일 MathCmd(op="softmax") 로 시간 회계는
+  한 번에. Phase 2 DataExecutor 가 canonical (x-max → exp → sum → div) 로
+  expand 한다.
+
+#### D3.6. Reduction (blocking)
+
+- `tl.sum(x, axis)`, `tl.max(x, axis)`, `tl.min(x, axis)` — 해당 axis 의
+  크기를 1 로 줄인 output handle 을 반환. `MathCmd(op=<name>, inputs=(x,),
+  out=, axis=axis)` 발행.
+
+#### D3.7. Index / scalar (PE_CPU, no engine)
+
+- `tl.program_id(axis=0) -> int`: `axis==0` → pe_id (cube-local PE 인덱스),
+  `axis==1` → cube_id (ADR-0022).
+- `tl.num_programs(axis=0) -> int`: `axis==0` → num_programs (cube 당
+  PE 수), `axis==1` → num_cubes.
+- `tl.arange(start, end, dtype="i32") -> TensorHandle`: TCM 의 인덱스
+  range. command 발사 없이 metadata 만.
+- `tl.zeros(shape, dtype="f16") -> TensorHandle`, `tl.full(shape, value,
+  dtype="f16") -> TensorHandle`: TCM 에 placeholder. command 발사 없음.
+
+#### D3.8. Scalar helpers (no command, no engine)
+
+- `TLContext.cdiv(a, b) -> int` (static): ceiling division
+  `-(-a // b)`. real Triton 의 `tl.cdiv` 모방.
+
+#### D3.9. Metadata-only (no compute, no DMA)
+
+- `tl.trans(x) -> TensorHandle`: shape 의 마지막 두 dim 을 swap 한 새
+  핸들. 같은 addr/data 를 공유, command 발사 없음.
+
+#### D3.10. IPCQ (CCL) primitives (ADR-0023 D4)
+
+- `tl.send(dir, src=None, *, src_addr=None, nbytes=None, shape=None,
+  dtype="f16", space="tcm") -> None`: blocking send. handle 형태 또는
+  raw 주소 형태 둘 다 허용. `IpcqSendCmd` 발행. handle 의 `.data` 스냅샷이
+  명령에 실리는 경우, recv 측에서 받은 데이터의 race 회피.
+- `tl.recv(dir=None, shape=(), dtype="f16", space="tcm", dst_addr=None,
+  dst_space=None) -> TensorHandle`: blocking recv. `dst_addr/dst_space`
+  둘 다 주면 "copy_to_dst" 모드, 아니면 "return_slot" 모드. greenlet
+  모드에서 핸들의 `.data` 에 실제 데이터 첨부.
+- `tl.recv_no_consume(dir=None, shape=(), dtype="f16") -> TensorHandle`:
+  **DIAGNOSTIC ONLY**. recv blocking 동기화는 그대로 적용되나 slot-read
+  latency (slot-IO + PE↔bank fabric drain) 는 건너뛴다. pe2pe overview
+  플롯에서 `tl.store` 와의 apples-to-apples 비교용. production kernel 은
+  사용 금지 — `consume=False` 라는 별도 명령 분기로 격리되어 있어 실수
+  flag 가 작동하지 않는다.
+- `tl.recv_async(dir, shape=(), dtype="f16") -> RecvFuture`: non-blocking
+  recv. `RecvFuture` 를 반환; 이후 `tl.wait(future)` 로 결과 수령.
+
+#### D3.11. Composite + control
+
+- `tl.composite(...)`: D3.3 에서 설명.
+- `tl.wait(handle=None)`: `CompletionHandle` (composite) 또는 `RecvFuture`
+  (async recv) 또는 `None` (모든 pending composite) 대기.
+- `tl.cycles(n)`: PE_CPU scalar 실행 overhead 를 명시적으로 선언.
+  `PeCpuOverheadCmd(cycles=n)` 발행.
+
+### D4. TensorHandle 산술 연산자 — thread-local TLContext
+
+`tl_context.py` 모듈 로드 시점에 `_enable_tensor_ops()` 가 호출되어
+`TensorHandle.__add__`, `__sub__`, `__mul__`, `__truediv__` 를 patch한다.
+각 연산자는 thread-local `_ctx` (모듈 변수) 에 저장된 active TLContext 의
+`_binary_math` 를 호출한다.
+
+따라서 kernel 안에서 `c = a + b` 는 `MathCmd(op="add", inputs=(a,b),
+out=)` 발행 + new TensorHandle 반환 패턴과 동일하다.
+
+active TLContext 관리:
+
+- `TLContext._set_active(ctx)`: 현재 thread/greenlet 의 active ctx 설정.
+- `TLContext._get_active()`: 조회 (없으면 RuntimeError).
+- `run_kernel(kernel_fn, tl_ctx, *args, **kwargs)`: helper. 진입 시
+  active 설정, kernel 실행, 종료 시 None 으로 복원.
+
+`KernelRunner` 는 매 cmd 분배 시 `_switch_kernel` 가 직접 `_set_active(tl)`
+를 호출하여, 같은 thread 안의 다른 PE runner 가 active 를 덮어쓴 경우에도
+복원되도록 한다.
+
+### D5. Scratch allocator — compute output handles
+
+`tl.dot`, `tl.exp`, `tl.add` (TensorHandle `__add__`) 등 결과를 만드는 op 는
+`_make_compute_out(shape, dtype)` 를 호출하여 16-byte aligned scratch
+주소를 할당한다. 이 주소는 `space="tcm"` 로 발행되며, 이후 `tl.send` /
+`tl.store` 가 이 handle 을 source 로 사용할 수 있다.
+
+`_scratch_base == 0` (command-list 모드 등) 이면 할당 주소가 0으로
+반환되어 handle 은 send/store 의 source 로 사용 불가 (이 경우 `tl.load`
+로 받은 핸들만 source 가 될 수 있다).
+
+cursor 가 `_scratch_size` (default 1 MiB) 를 초과하면 RuntimeError.
+cursor 는 매 kernel invocation 시작 시 0 으로 리셋되어야 하나 (현재 코드는
+KernelRunner 가 새 TLContext 를 매번 생성하여 자연스럽게 리셋됨).
+
+### D6. Dispatch overhead — `PeCpuOverheadCmd(dispatch_cycles)`
+
+모든 non-metadata `tl.*` 호출의 entry 에서 `_emit_dispatch_overhead()` 가
+호출되며 `dispatch_cycles > 0` 일 때 `PeCpuOverheadCmd(dispatch_cycles)`
+를 발행한다. PE_CPU 가 명령 dispatch 자체에 소비하는 cycle 비용을
+모델링하기 위함이다.
+
+기본값:
+
+- `TLContext.__init__` 의 `dispatch_cycles` 매개변수 기본값: 1 cycle.
+- `KernelRunner` 가 만드는 TLContext: 0 cycles (greenlet 모드는 cycle
+  회계가 별도, ADR-0020 D3 정신).
+
+### D7. Kernel registry (`triton_emu/registry.py`)
+
+별도의 `_kernels: dict[str, Callable]` 가 kernel 이름 → 함수 매핑을 보유:
+
+- `register_kernel(name, fn)`: duplicate 등록 시 ValueError.
+- `get_kernel(name)`: 미등록 시 KeyError.
+- `clear_registry()`: 테스트 전용.
+
+`RuntimeContext.launch(kernel_name, kernel_fn, *args)` 가 매 호출마다
+`_kernels[kernel_name] = kernel_fn` 으로 idempotent 덮어쓴다 (last call
+wins). 이는 ADR-0045 D8 의 launch 동작과 정합된다.
+
+PE_CPU 는 `KernelRef.name` 으로 registry 에서 kernel 함수를 lookup 한 뒤
+KernelRunner 로 실행한다.
+
+## Alternatives Considered
+
+### A1. tl.* 를 ADR-0014 / ADR-0020 안으로 통합
+
+기각. ADR-0014 는 PE pipeline (PeCommand 의 sim_engine 측 소비) 를, ADR-0020
+은 2-pass 실행 (Phase 1 timing / Phase 2 data) 을 다룬다. `tl.*` 는 kernel
+작성자가 만나는 API 표면이라 독립 분리하는 것이 검색성·온보딩 측면에서
+낫다.
+
+### A2. command-list 모드 deprecation
+
+기각 (현재). 단순한 unit test 와 kernel verification 에서 command-list
+모드가 가볍게 동작한다. greenlet 의존성 없이 PeCommand 시퀀스를 검사할 수
+있는 출입구로 유지한다. greenlet 모드만의 의미 (실데이터, Phase 2) 가
+필요하면 D2 의 mode 선택으로 명시적으로 들어간다.
+
+### A3. TensorHandle 산술 연산자 제거
+
+기각. real Triton 의 kernel 코드 가독성을 흉내내기 위함이며 (예: `c = a +
+b`), thread-local active ctx 패턴이 깔끔하게 작동 중. 명시적 `tl.add(a, b)`
+도 D3.5 에 노출되어 있어, 연산자가 헷갈리면 함수형 호출로 대체 가능.
+
+### A4. softmax 를 명시적 시퀀스 (max → exp → sum → div) 로 expand
+
+부분 채택. `tl.softmax` 는 단일 `MathCmd(op="softmax")` 로 timing 회계는
+한 번에 처리한다 (D3.5). 실 데이터 expansion 은 Phase 2 DataExecutor 가
+canonical 시퀀스로 풀어준다. 즉, 시간 모델은 atomic, 데이터 모델은
+expansion — 두 마리 토끼를 의도적으로 분리.
+
+## Consequences
+
+- bench 작성자가 만나는 모든 `tl.*` primitive 가 한 ADR 에 분류·정의됨.
+  ADR-0045 D8 의 host-side surface (torch.empty 등) 와 짝을 이루어 "kernel
+  안 / 밖" 양쪽 작성 가이드가 완성.
+- command-list / greenlet 두 모드의 차이가 D2 에 명시되어, 새로운 `tl.*`
+  primitive 추가 시 `_emit()` 패턴만 따르면 양쪽 자동 호환됨.
+- thread-local active ctx 패턴 (D4) 이 ADR-level 에서 정당화되어, 향후
+  multi-PE 동일-thread 실행 시 reset 책임이 어디인지 명확해짐
+  (`_switch_kernel` 가 cmd 분배 시 active 복원 — KernelRunner.run 의
+  contract).
+- `tl.recv_no_consume` 의 진단 전용 격리(D3.10) 가 ADR 에 굳어져, 실수로
+  production kernel 에서 사용되는 것을 막는 layer 가 명확.
+- registry (D7) 가 별도 D 항목으로 분리되어, kernel 이름 충돌 / 동적
+  재등록 동작의 사양이 명시.

docs/adr-ko/ADR-0047-par-ahbm-ccl-backend.md

@@ -0,0 +1,243 @@
+# ADR-0047: AHBM CCL Backend — `torch.distributed`-compat shim
+
+## Status
+
+Accepted (2026-05-22).
+
+`runtime_api/distributed.py` 의 `AhbmCCLBackend` + `DistributedContext` —
+즉 `torch.distributed.init_process_group(backend="ahbm")` 진입점이 실제로
+무엇을 설치하고 어떤 의미로 `all_reduce`/`barrier`/`get_rank` 등을
+구현하는지를 명시한다. ADR-0023 D11 이 "torch.distributed compatibility"
+의도를 언급하나, **backend 자체의 동작 모델**은 ADR-level 에 없었다.
+
+## First action (제일 처음에 하는 일)
+
+`RuntimeContext.__post_init__` 가 자동으로 `DistributedContext()` 인스턴스를
+만들어 `self.distributed` 에 attach 한다. 그 시점의 첫 일은:
+
+1. `self._backend: AhbmCCLBackend | None = None` 으로 초기화 (아직 init
+   되지 않은 상태).
+2. `self._rank_by_greenlet: dict = {}` 로 greenlet-local rank 레지스트리
+   초기화 (ADR-0024 D2).
+3. 호출자(RuntimeContext) 측에서 `dc._ctx_ref = self` 로 back-reference 를
+   심어, 이후 `init_process_group` 가 `ctx.engine` / `ctx.spec` / `ctx.launch`
+   에 도달할 수 있게 한다.
+
+즉, **DistributedContext 의 첫 일은 "RuntimeContext 에 자기 자신을
+back-reference 와 함께 부착하고 backend 슬롯을 비워두는 것"**. 실제 backend
+설치(IPCQ install, world_size 산출, 알고리즘 모듈 로드)는 사용자 코드의
+`torch.distributed.init_process_group(backend="ahbm")` 호출 시점에 비로소
+일어난다.
+
+해당 시점의 `init_process_group` 의 첫 일은:
+
+1. `backend != "ahbm"` 이면 즉시 `ValueError("Unsupported backend ...")`.
+2. `getattr(self, "_ctx_ref", None)` 가 None 이면
+   `RuntimeError("DistributedContext not bound to a RuntimeContext")`.
+3. `self._backend = AhbmCCLBackend(torch_ctx=ctx)` — 이 생성자 안에서
+   ccl.yaml load + 알고리즘 모듈 import + world_size 산출 + SFR 설정 +
+   IPCQ install 이 모두 일어난다.
+4. `self._backend._dist_ctx = self` — backend 가 거꾸로
+   `_rank_by_greenlet` 에 접근할 수 있게 함.
+
+## Context
+
+PyTorch DDP 의 collective 호출 (`init_process_group`, `all_reduce` 등) 을
+그대로 사용할 수 있게 만들어, bench 코드가 "진짜 DDP training script" 와
+동일한 모습이 되도록 하는 것이 `AhbmCCLBackend` 의 목적이다 (ADR-0024 +
+ADR-0027 의 launcher 모델과 정렬).
+
+이 backend 가 책임지는 것:
+
+- `init_process_group` 시점에 **IPCQ neighbor table 을 한 번 설치** (real
+  NCCL communicator creation 과 유사).
+- `all_reduce(tensor, op="sum")` 호출 시 **설정된 algorithm 의 kernel 함수
+  를 `ctx.launch(...)` 로 발사**.
+- `get_world_size` / `get_rank` 를 greenlet-local rank 레지스트리와
+  ccl.yaml/topology 로부터 일관되게 답함.
+
+ADR-0023 D10 (IPCQ install plan), ADR-0024 (SIP launcher) 가 부분적으로
+이를 다루나, **`AhbmCCLBackend` 자체의 책임 범위와 의사결정 순서**는
+어디에도 명시되어 있지 않다. 본 ADR 이 채운다.
+
+## Decision
+
+### D1. backend 는 `init_process_group(backend="ahbm")` 시점에만 생성된다
+
+`DistributedContext` 는 `__init__` 시점에 `_backend = None` 으로 시작한다.
+backend 객체는 사용자가 `dist.init_process_group(backend="ahbm")` 를
+호출하기 전까지 존재하지 않으며, 그 외 API (`is_initialized`,
+`get_world_size`, `all_reduce`, `barrier`) 가 backend 가 None 인 채로
+호출되면 `RuntimeError("Default process group has not been initialized...")`
+를 던진다 (`_ensure_initialized` 헬퍼).
+
+`backend != "ahbm"` 은 즉시 `ValueError`. 다른 backend 명 (nccl, gloo
+등) 은 인식하지 않는다.
+
+### D2. world_size 산출 우선순위 — algorithm > defaults > topology
+
+`AhbmCCLBackend._resolve_world_size` (ADR-0024 D1) 의 결정 순서:
+
+1. `ccl.yaml` 의 algorithm entry 에 `world_size` 가 있으면 그 값.
+2. `defaults.world_size` 가 있으면 그 값.
+3. 둘 다 없으면 `spec.system.sips.count` (=topology 의 SIP 개수).
+
+기본 의미는 **rank = SIP** (ADR-0024). cube/PE-level parallelism 은 각
+rank 안에서 DPPolicy 로 표현되며 world_size 에 영향을 주지 않는다. 명시적
+`ccl.yaml` 의 world_size override 가 있으면 legacy "rank = flat PE 인덱스"
+테스트 경로를 위해 그대로 존중된다.
+
+`init_process_group(world_size=..., rank=...)` 의 사용자 인자는 **수신하나
+무시**된다 (real PyTorch 의 `RANK` / `WORLD_SIZE` env var 와 같은 의미).
+
+### D3. `init_process_group` 가 즉시 하는 4가지 설치 작업
+
+`AhbmCCLBackend.__init__` 안에서 다음이 순차 실행된다:
+
+1. **ccl.yaml 로딩**: `kernbench.ccl.install.load_ccl_config()` →
+   `resolve_algorithm_config(_cfg_all)` 로 `defaults.algorithm` (또는
+   사용자가 지정한 알고리즘) 의 merged config 산출.
+2. **알고리즘 모듈 import**: `importlib.import_module(self._merged["module"])`.
+   이 모듈은 `kernel` 함수, `kernel_args(world_size, n_elem, cube_w, cube_h)`
+   helper, optional `TOPO_NAME_TO_KIND` 매핑을 노출해야 한다.
+3. **world_size 산출** (D2).
+4. **topology 메타 수집**: `spec` 으로부터 `n_sips`, `sip_topo` (`ring_1d`
+   기본), `cube_w`/`cube_h`, `sips.w`/`sips.h`. SIP topology 가 ring_1d 가
+   아니면 explicit `w`/`h` 또는 square root 로 (`w*h == n_sips` 보장)
+   `_sip_topo_w/h` 산출. 불일치 시 `ValueError`.
+5. **SFR + IPCQ 설치**: `kernbench.ccl.sfr_config.configure_sfr_intercube_multisip
+   (engine, spec, self._merged)` 를 호출. 이 함수가 모든 SIP/cube 의 pe0 에
+   IPCQ neighbor table 을 푸시 (real NCCL communicator 의 일회성 설정에
+   해당).
+
+이 순서가 변하면 (예: SFR 전에 algorithm 모듈 load 가 실패하면) 부분 초기화
+상태가 발생할 수 있다. 따라서 D3 는 atomic 한 4-단계로 본다 — 실패 시
+backend 는 미설치 상태로 남는다.
+
+### D4. greenlet-local rank 등록 (ADR-0024 D2)
+
+`DistributedContext._rank_by_greenlet: dict[greenlet, int]` 은 spawn 된
+worker greenlet 각각에 rank 를 매핑한다. bench launcher (예:
+`torch.multiprocessing.spawn`) 가 worker 를 띄울 때
+`dc._bind_rank(g, rank)` 를 호출하여 등록한다.
+
+`get_rank()` 는 `getcurrent()` 의 greenlet 을 lookup. 미등록 greenlet은
+fallback 으로 0 을 반환 — single-driver / 테스트 호환성 유지.
+
+backend 는 `_dist_ctx._rank_by_greenlet` 를 통해 `all_reduce` 시 현재
+greenlet 의 rank 를 가져온다 (D5).
+
+### D5. `all_reduce(tensor, op="sum")` 동작
+
+검증 단계:
+
+- `op != "sum"` → `NotImplementedError`. 현재 kernel 들은 add reduction만 구현.
+- `tensor._handle is None` → `RuntimeError("not deployed")`.
+- `tensor._handle.shards` 가 비면 `RuntimeError("no shards")`.
+
+준비 단계:
+
+- `n_elem = shards[0].nbytes // tensor.itemsize` — 단일 shard 의 element 수.
+- `kernel_fn = self._algo_module.kernel` — D3 에서 import 된 알고리즘 모듈의
+  진입 함수.
+- effective cube dims 결정: 첫 번째 SIP 의 cube 갯수가 1 이면 (1,1) 으로
+  scalar 처리, 아니면 토폴로지의 `cube_w`/`cube_h` 사용. TP 가 일부 cube
+  만 쓰는 경우를 자연스럽게 흡수.
+- `kernel_args = self._algo_module.kernel_args(world_size, n_elem, cube_w,
+  cube_h)` — 알고리즘이 자기 kernel 에 넘길 인자 셋을 결정.
+
+dispatch:
+
+- 현재 greenlet 의 rank 를 `_rank_by_greenlet.get(g, 0)` 로 lookup.
+- `extra_args = (sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)` 를 append.
+- `pending = self.ctx.launch(algorithm_name, kernel_fn, tensor, *kernel_args,
+  *extra_args, _defer_wait=True)` — `_defer_wait=True` 로 collective drain
+  을 메인 scheduler 에 위임 (ADR-0027 D0.4).
+
+drain:
+
+- 부모 greenlet 이 살아있으면 (multi-greenlet 모드) `_pending_collective_handles`
+  에 enqueue 한 뒤 부모로 switch. 메인 scheduler 가 모든 rank 의 launch 후
+  일괄 drain.
+- 단일-driver 모드면 inline 으로 `for h, _sip_id, meta in pending:
+  self.ctx.wait(h, _meta=meta)` 즉시 drain.
+
+### D6. `barrier()` 는 no-op 이다 (single-driver 모델)
+
+kernbench 는 하나의 Python process 안에서 모든 rank 를 greenlet 으로 다룬다.
+process 간 동기화가 필요한 상황이 없으므로 `barrier()` 는 호출 가능하지만
+실제 어떤 동기화도 수행하지 않는다. real PyTorch DDP 와의 API 호환성을
+위해 유지 (호출자가 NotImplementedError 를 받지 않도록).
+
+장래에 multi-process kernbench (예: SimPy event loop 가 process 별로
+독립) 가 도입되면 D6 를 supersede 하는 새 ADR 이 필요.
+
+### D7. `get_rank` / `get_world_size` / `get_backend` 의 의미
+
+- `get_rank()` (D4): 현재 greenlet 의 bound rank. 미등록은 0.
+- `get_world_size()` (D2): backend 가 D3 에서 산출한 world_size.
+- `get_backend()`: 항상 `"ahbm"` 문자열. backend 객체가 존재하지 않으면
+  `_ensure_initialized` 에서 RuntimeError.
+
+real PyTorch 와의 차이:
+
+- real PyTorch `get_rank()` 는 process global 값이지만, kernbench 는
+  greenlet-local. spawn 된 worker 안에서 호출하면 rank, main thread 에서
+  호출하면 0. bench 작성자는 worker 함수 안에서만 의미 있는 rank 를 기대해야
+  한다.
+
+### D8. 지원하는 API 표면 (final)
+
+`DistributedContext` 가 노출하는 API:
+
+- `init_process_group(backend="ahbm", world_size=None, rank=None, **kwargs)`
+- `is_initialized() -> bool`
+- `get_world_size() -> int`
+- `get_rank() -> int`
+- `get_backend() -> str`
+- `all_reduce(tensor, op="sum") -> None`
+- `barrier() -> None`
+- (internal) `_bind_rank(g, rank)`
+
+이외의 PyTorch distributed API (broadcast, reduce, all_gather, gather,
+scatter, send/recv 등) 는 **아직 구현되어 있지 않다**. kernel 레벨에서는
+`tl.send`/`tl.recv` (ADR-0046 D3.10) 로 직접 표현 가능하나, dist.* surface
+로는 노출되지 않는다. 추가 collective 가 필요해질 시 별도 알고리즘 모듈
++ `DistributedContext` 메소드 한 쌍을 추가하여 D8 를 확장한다.
+
+## Alternatives Considered
+
+### A1. backend 를 `RuntimeContext.__init__` 에서 즉시 생성
+
+기각. ccl.yaml 이 없거나 알고리즘 모듈을 import 할 수 없는 경우, bench 가
+distributed 기능을 안 쓰는데도 RuntimeContext 생성 자체가 실패하게 된다.
+"호출 시점에 비로소 설치" (D1) 가 lazy 의미상 옳다.
+
+### A2. world_size 를 항상 topology 로부터 자동 산출 (override 금지)
+
+기각. ADR-0024 D1 의 "explicit override" 경로가 legacy 테스트에서 사용 중.
+한 SIP 안에서 PE-level rank 를 따로 정의해야 하는 진단 시나리오를 위해
+유지.
+
+### A3. `op != "sum"` 을 silent fallback 으로 처리
+
+기각. 사용자가 `op="prod"` / `"max"` / `"avg"` 를 의도했는데 silently sum
+이 실행되면 결과 검증이 매우 어렵다. 명시적 `NotImplementedError` 가 안전.
+
+### A4. `barrier` 를 SimPy event 로 구현
+
+기각 (현재). single-driver 모델에서 cross-process 동기화 의미가 없으므로
+no-op 가 의미적으로 정확. SimPy fake-barrier 는 의미 없이 코드 복잡도만
+높임. multi-process kernbench 도입 시 재평가.
+
+## Consequences
+
+- `torch.distributed.init_process_group(backend="ahbm")` 의 4-단계 설치
+  (D3) 가 ADR-level 에서 굳어져, 향후 새 collective 알고리즘이 어디에
+  훅을 걸어야 하는지 명확.
+- D2 의 우선순위 (algorithm > defaults > topology) 가 명시되어, ccl.yaml
+  변경 시 영향 범위를 빠르게 가늠 가능.
+- D6 의 barrier no-op 결정이 ADR-level 에 굳어져, multi-process kernbench
+  도입 시 별도 ADR 로 supersede 해야 함이 분명.
+- D8 의 미지원 API 목록이 명시되어, 사용자가 `dist.broadcast(...)` 를
+  호출하려 할 때의 명확한 거절 근거 제공.

docs/adr-ko/ADR-0048-mem-allocator-algorithms.md

@@ -0,0 +1,262 @@
+# ADR-0048: Memory Allocator Algorithms — VirtualAllocator + PEMemAllocator
+
+## Status
+
+Accepted (2026-05-22).
+
+`policy/address/allocator.py` 의 `_FreeList` / `PEMemAllocator` 와
+`va_allocator.py` 의 `VirtualAllocator` 가 사용하는 free-list 알고리즘,
+페이지 정렬, coalescing 규칙을 명시한다. ADR-0001 (PhysAddr 레이아웃) 과
+ADR-0011 (PA/VA/LA 모델) 이 주소 스킴을 정의하나, **할당 알고리즘**은 별도
+ADR 이 없었다.
+
+## First action (제일 처음에 하는 일)
+
+### `_FreeList(capacity)`
+
+생성 즉시 `self._capacity = capacity`, `self._used = 0`, `self._free =
+[(0, capacity)]` 로 초기화. 첫 일은 **전 영역을 single free block 으로
+세우는 것** — 즉 `(offset=0, size=capacity)` 한 튜플이 free list 의 유일한
+원소다.
+
+### `PEMemAllocator(sip_id, die_id, pe_id, cfg)`
+
+생성 즉시 두 개의 `_FreeList` 를 만든다:
+
+- `self._hbm = _FreeList(cfg.hbm_slice_bytes)` — 이 PE 가 소유한 HBM
+  slice 의 바이트 크기 (`hbm_bytes_per_cube // hbm_slices_per_cube`) 만큼.
+- `self._tcm = _FreeList(cfg.tcm_allocatable_bytes)` — `tcm_bytes_per_pe -
+  tcm_scheduler_reserved_bytes` 만큼 (scheduler 예약분은 사전 분리).
+
+따라서 PEMemAllocator 의 첫 일은 **이 PE 의 HBM slice 와 사용자
+TCM 영역을 각각 단일 free block 으로 세우는 것**.
+
+### `VirtualAllocator(va_base, va_size, page_size=2*1024*1024)`
+
+생성 즉시 `self._va_base = va_base`, `self._va_size = va_size`,
+`self._page_size = page_size`, `self._used = 0`, `self._free = [(va_base,
+va_size)]`. 첫 일은 **VA base 부터 size 까지 single block 으로 세우고
+page_size 를 회수**.
+
+## Context
+
+`runtime_api/context.py::_ensure_allocators` 는 다음 단계로 allocator 세트를
+구성한다:
+
+1. spec 으로부터 `hbm_total_gb_per_cube`, `hbm_slices_per_cube`,
+   `tcm_size_mb`, target_device 별 SIP 범위 등을 읽음.
+2. `AddressConfig` 로 모든 파라미터를 frozen 하게 패킹.
+3. target SIP 범위 × cube × PE 의 모든 조합에 대해
+   `PEMemAllocator(sip, cube, pe, cfg)` 인스턴스를 1개씩 생성.
+4. `VirtualAllocator(va_base=0x1_0000_0000, va_size=64 GiB,
+   page_size=pe_mmu.page_size)` 를 1개 생성.
+
+allocator 들의 책임:
+
+- **PEMemAllocator**: PE-로컬 HBM slice / TCM 의 PA-공간 할당 (PhysAddr
+  encoding 까지 포함).
+- **VirtualAllocator**: device-wide VA 공간을 페이지 정렬로 할당. 이후
+  `RuntimeContext._create_tensor` 가 VA → PA 매핑을 `MmuMapMsg` 로 fabric
+  에 push.
+
+이 알고리즘들은:
+
+- **first-fit** 으로 단순.
+- 자유 블록 리스트는 **offset 정렬 (sorted by start)** 유지.
+- `free()` 시 **양쪽 인접 블록과 coalesce**.
+
+이런 결정의 근거가 어디에도 없으므로, 향후 누군가 "왜 best-fit 이 아닌가",
+"왜 buddy allocator 가 아닌가", "왜 partial overlap free 가 silently
+허용되는가" 라는 질문에 답할 기준이 필요. 본 ADR 이 그 기준을 마련한다.
+
+## Decision
+
+### D1. `_FreeList` — offset-기반 first-fit + coalescing
+
+`policy/address/allocator.py::_FreeList`:
+
+- 내부 표현: `list[tuple[int, int]]` = `[(start_offset, size), ...]` —
+  start offset 으로 정렬된 자유 블록의 sorted list.
+- `alloc(nbytes)`:
+  1. free list 를 앞에서부터 순회 (first-fit).
+  2. 처음 만나는 `size >= nbytes` 인 블록에서 앞부분을 잘라 사용.
+  3. 정확히 일치하면 블록 통째로 제거; 아니면 `(start+nbytes, size-nbytes)`
+     로 축소.
+  4. `_used += nbytes`, 잘라낸 `start` 반환.
+  5. 맞는 블록이 없으면 `AllocationError("overflow ... largest free block
+     ...")`.
+- `free(offset, nbytes)`:
+  1. `_used -= nbytes`.
+  2. `bisect_left(self._free, (offset,))` 로 삽입 위치 결정.
+  3. 직전 블록과 인접 (`prev_start + prev_size == offset`) 하면 흡수.
+  4. 직후 블록과 인접 (`offset+nbytes == next_start`) 하면 흡수.
+  5. coalesced range 를 정렬 위치에 insert.
+
+이 알고리즘은 fragmentation 에 약점이 있으나 (best-fit / buddy 대비), 본
+시뮬레이터의 워크로드 특성상 (deploy/free 패턴이 거의 stack-like) 충분
+하다는 것이 디자인 가정이다. 워크로드가 변하면 D1 supersede 후보.
+
+### D2. partial overlap free 는 **검사하지 않는다**
+
+`_FreeList.free(offset, nbytes)` 는 호출자가 정확한 (offset, nbytes) 를
+넘긴다고 신뢰한다. 다음을 검증하지 않는다:
+
+- 그 range 가 실제로 alloc 된 것인지.
+- 그 range 가 다른 alloc 된 영역과 겹치지 않는지.
+
+이유: 시뮬레이터 컨텍스트에서 호출자는 항상 `alloc()` 의 반환값을 그대로
+저장했다가 `free()` 에 넘기는 패턴이며, 외부 사용자 입력이 아니다. 안전성
+검사를 추가하면 매 free 마다 O(N) 비용이 들어 시뮬 wall-clock 에 영향.
+
+이 신뢰 모델이 깨지면 (예: 두 텐서가 같은 PA 를 가리키는 코드 경로 도입)
+즉시 ADR-level 으로 재검토.
+
+### D3. `PEMemAllocator` — HBM/TCM 두 채널 분리
+
+`PEMemAllocator(sip_id, die_id, pe_id, cfg)` 는 두 `_FreeList` 를 보유:
+
+- `_hbm`: `cfg.hbm_slice_bytes` 크기.
+- `_tcm`: `cfg.tcm_allocatable_bytes` (= `tcm_bytes_per_pe -
+  tcm_scheduler_reserved_bytes`) 크기.
+
+`alloc_hbm(nbytes) -> PhysAddr`:
+
+- `_hbm.alloc(nbytes)` 로 offset 획득.
+- `PhysAddr.pe_hbm_addr(sip_id, die_id, pe_id, pe_local_hbm_offset=offset,
+  slice_size_bytes=cfg.hbm_slice_bytes)` 로 PA 인코딩.
+- 실패 시 `AllocationError("HBM overflow ...")`.
+
+`free_hbm(pa, nbytes)`:
+
+- `pa.hbm_offset - pe_id * cfg.hbm_slice_bytes` 로 PE-local offset 복원.
+- `_hbm.free(offset, nbytes)`.
+
+`alloc_tcm(nbytes) -> PhysAddr`: 유사하게 `PhysAddr.pe_tcm_addr` 로 인코딩.
+
+`free_tcm(pa, nbytes)`: `pa.sub_offset` 을 그대로 사용 (TCM 은 PE-local
+offset 이 곧 sub_offset).
+
+scheduler-reserved TCM 영역 (`cfg.tcm_scheduler_reserved_bytes`) 은
+allocator 가 인지하지 않는다 (`_tcm` 의 capacity 에서 사전 차감되어 있음).
+이는 ADR-0014 의 PE_SCHEDULER 내부 buffer 예약과 정합된다.
+
+### D4. `VirtualAllocator` — 페이지 정렬 first-fit + coalescing
+
+`policy/address/va_allocator.py::VirtualAllocator`:
+
+- 내부 표현: `_FreeList` 와 동일한 sorted `list[tuple[int, int]]`.
+  최초: `[(va_base, va_size)]`.
+- `_align_up(nbytes) = ceil(nbytes / page_size) * page_size`.
+- `alloc(nbytes) -> int`:
+  1. `aligned = _align_up(nbytes)`.
+  2. first-fit 으로 `size >= aligned` 인 블록 탐색.
+  3. 블록 앞부분 `aligned` 만큼 잘라 사용. 정확히 일치하면 제거.
+  4. `_used += aligned`. 블록 `start` (= aligned 된 VA) 반환.
+  5. 실패 시 `VaAllocationError`.
+- `free(va, nbytes)`: `_align_up(nbytes)` 단위로 free. _FreeList 와 동일한
+  coalesce 알고리즘.
+
+`page_size` 의 실제 값은 두 곳에서 다른 기본을 갖는다:
+
+- `VirtualAllocator.__init__` 의 매개변수 기본값: `2 MiB`. 직접 호출하는
+  테스트가 그대로 받는다.
+- `RuntimeContext._ensure_allocators` 가 인스턴스화할 때:
+  `pe_mmu.attrs.get("page_size", 4096)` — `topology.yaml` 의
+  `pe_mmu.attrs.page_size` 가 있으면 그 값, 없으면 fallback 4 KiB.
+
+두 기본이 다른 이유: VirtualAllocator 의 standalone 기본은 ADR-0039 의
+PE_MMU stopgap 기본 (2 MiB) 과 정합되어 직접 테스트가 자연스럽고, context
+fallback 의 4 KiB 는 topology 미설정 시 안전한 minimum page 다. 실제 사용
+경로는 항상 후자이며 (`_ensure_allocators` 가 인스턴스화하므로),
+`topology.yaml` 에서 `page_size` 가 명시되면 그 값이 양쪽 (MMU + VA
+allocator) 으로 일관되게 흐른다.
+
+만약 이 일치가 깨지면 (예: VirtualAllocator 의 page_size 를 PE_MMU 와
+다르게 인스턴스화) MMU `map()` 가 서브-페이지 region 모드 (ADR-0039 D3) 로
+흐른다.
+
+VA 기본 범위: `va_base = 0x1_0000_0000` (= 4 GiB), `va_size = 64 GiB`. 이
+값은 `_ensure_allocators` 에 하드코딩되어 있으며 ADR-0011 의 VA 모델에서
+직접적인 의미를 갖지는 않는다 — 단지 host 코드와 충돌하지 않을 만큼 큰
+주소 공간을 device-wide 로 잡아둔 것.
+
+### D5. allocator 인스턴스의 lifecycle
+
+- `RuntimeContext._ensure_allocators` 가 lazy 하게 호출됨 (`_create_tensor`
+  의 첫 호출 시점).
+- 한 번 생성된 allocator dict (`self._allocators`) 는 RuntimeContext 의
+  lifetime 동안 재사용. 같은 process 안의 두 번째 deploy 는 새 객체를
+  만들지 않는다.
+- `RuntimeContext.cleanup()` 이 모든 living tensor 의 `_free_tensor()` 를
+  호출 → MMU unmap + `va_allocator.free` + `pemem_allocator.free_hbm` 으로
+  free list 가 원상복구. 다음 RuntimeContext 가 다시 만들면 초기 상태부터.
+
+allocator 상태가 RuntimeContext 간에 공유되지 않는 점이 단일 process 안의
+연속 실행에서 deploy → cleanup → deploy 의 결정성을 보장한다.
+
+### D6. Allocator 실패는 raise 한다 (silent OOM 금지)
+
+`_FreeList.alloc` / `VirtualAllocator.alloc` 모두 충분한 free block 이
+없으면 `AllocationError` / `VaAllocationError` 를 던진다. 메시지에는
+"required size + largest available block" 가 포함되어, fragmentation
+인지 진짜 OOM 인지 진단 가능.
+
+silent fallback (예: 가장 큰 블록만큼만 alloc) 는 절대 금지 — 부분 할당된
+텐서가 SimPy 단계에 들어가면 라우팅·DMA 가 잘못된 PA 를 인지하여 시뮬
+정확도가 깨진다.
+
+### D7. address space 와 allocator 의 1:1 대응
+
+물리 주소 공간 분리는 PhysAddr 의 sub-unit (ADR-0001 D2.3) 으로 표현되며,
+각 sub-unit 마다 별도 allocator 인스턴스를 둔다:
+
+- HBM slice → `PEMemAllocator._hbm`.
+- PE TCM → `PEMemAllocator._tcm`.
+- (현재 미사용) M_CPU local memory, CUBE SRAM → 별도 allocator 필요. 현재
+  구현은 아직 IPCQ-only slot 으로 처리 (ADR-0023 D9.7) 하며 PA 공간을
+  share 하지 않으므로 별도 free-list 가 없음.
+
+cube-level SRAM allocator 가 필요해지면 `_FreeList(cfg.sram_bytes_per_cube)`
+인스턴스를 cube 단위로 추가한다 (`cfg.sram_bytes_per_cube` 는 이미
+`AddressConfig` 에 정의되어 있어 데이터 모델은 준비됨).
+
+## Alternatives Considered
+
+### A1. best-fit / buddy allocator
+
+기각 (현재). 워크로드의 alloc/free 패턴이 stack-like (deploy 순서 = free
+순서) 라 first-fit + coalescing 으로 fragmentation 이 충분히 통제된다.
+LLM kernel sweep 에서 long-running fragmentation 이 관찰되면 buddy 로
+교체하는 ADR 을 별도로 만든다.
+
+### A2. partial overlap free 검증 추가
+
+기각. D2 의 신뢰 모델 + O(N) 검사 비용. 단, 디버그 모드 (`KERNBENCH_DEBUG`
+env var 등) 에서 활성화하는 옵션은 후속 작업으로 가능.
+
+### A3. VA 와 PA 의 통합 allocator
+
+기각. VA 공간 (64 GiB device-wide) 과 PA 공간 (slice 별 ~6 GiB) 는 의미
+차원이 다르다. VA 는 host kernel 의 view, PA 는 device sub-unit 의 view.
+ADR-0011 의 VA 모델 정신 (MMU 가 둘 사이를 매핑) 과 정합하기 위해
+allocator 도 분리.
+
+### A4. page_size 의 multi-tier 지원 (large page + small page)
+
+기각 (현재). 단일 page_size (현재 2 MiB) 가 LLM kernel 의 텐서 단위 (수
+MiB~수 GiB) 에 맞고, ADR-0039 D3 의 서브-페이지 region 으로 작은 매핑이
+필요할 때 흡수된다. multi-tier page 는 MMU 자체 모델을 확장해야 하므로
+별도 ADR 후보.
+
+## Consequences
+
+- allocator 알고리즘이 ADR-level 에서 굳어져 (D1·D3·D4), 새로운 시뮬
+  시나리오에서 fragmentation 이슈가 발생할 때 "여기서 first-fit + coalesce
+  를 쓰고 있다" 가 명확.
+- D2 의 신뢰 모델이 명시되어, 향후 사용자 입력으로부터 직접 alloc/free 를
+  받는 경로가 도입되면 본 ADR supersede 가 필요함을 일찍 인지 가능.
+- D7 의 sub-unit별 allocator 1:1 대응이 명시되어, M_CPU/SRAM 별도 영역이
+  필요해질 때 어디에 free-list 를 추가해야 하는지 명확.
+- `VirtualAllocator` 의 page_size 가 PE_MMU 설정과 일치해야 함이 D4 에
+  적혀 있어, 향후 topology.yaml 의 page_size 변경 시 ADR-0039 stopgap 동작
+  과의 상호작용을 빠르게 가늠 가능.

docs/adr-ko/ADR-0049-ver-probe-subcommand.md

@@ -0,0 +1,231 @@
+# ADR-0049: `kernbench probe` Subcommand — Traffic-Pattern Verification Harness
+
+## Status
+
+Accepted (2026-05-22).
+
+`probes/probe.py` 의 `run_probe(...)` 가 노출하는 traffic-pattern catalog,
+formula vs actual 비교, 그리고 monotonicity / D2H≥H2D 같은 invariant
+체크의 의미를 명시한다. ADR-0010 (CLI surface) 가 `kernbench probe`
+subcommand 를 enumerate 하나, **probe 가 실제로 측정하는 것**과 **어떤
+invariant 를 PASS/FAIL 로 판정하는가**는 ADR-level 에 없었다.
+
+## First action (제일 처음에 하는 일)
+
+`run_probe(topology_path, case_filter=None)` 의 첫 4가지 작업:
+
+1. `Path(topology_path).expanduser().resolve()` 로 절대 경로 산출.
+2. `load_topology(path)` → `TopologyGraph` 인스턴스 (그래프 + spec).
+3. `_build_edge_map(graph)` → `{(src, dst): Edge}` 빠른 lookup 테이블.
+4. `AddressResolver(graph)` + `PathRouter(graph)` 인스턴스화.
+
+그 다음 `nbytes = 32768` (= 32 KiB, summary table 의 기준 데이터 크기) 와
+`show_all = (case_filter is None or case_filter == "all")` 를 설정.
+
+즉, **probe 의 첫 일은 "토폴로지를 한 번 로드하여 edge map / resolver /
+router 를 준비하고, 32 KiB 라는 표준 측정 크기를 픽스하는 것"**. 그 이후
+H2D → D2H → PE DMA 세 카테고리의 case 들이 각각 별도의 `GraphEngine`
+인스턴스에서 실행된다 (case 간 cross-talk 차단).
+
+## Context
+
+`kernbench probe` 는 다음 의도로 도입된 verification 도구다:
+
+- **수동 분석 ground truth**: 실 시뮬레이션 (`kernbench run --bench ...`)
+  결과의 latency 가 비정상으로 보일 때, 단순 traffic pattern 의 정답을 별도
+  로 얻어 비교.
+- **formula vs actual 비교**: 분석 모델 (wire latency + overhead + drain)
+  과 시뮬레이션 결과 (`total_ns`) 가 일치하는지 확인. 일치하지 않으면 모델
+  단순화 가정 (ADR-0033) 어디가 빠진 것인지 단서.
+- **monotonicity check**: hop 수가 늘면 latency 가 단조 증가해야 한다는
+  invariant 의 자동 확인.
+- **utilization sweep**: 데이터 크기 (4 KiB ~ 1 MiB) 별 BW 활용률 표.
+
+이 도구의 동작 사양이 ADR-level 에 없으면:
+
+- 다른 형식의 traffic pattern (예: MCpuDma, IPCQ) 을 추가하려는 사람이 기존
+  카테고리의 표 포맷 / 측정 단위를 일관되게 따르기 어렵다.
+- monotonicity 가 무엇을 기준으로 검사되는지 (hop 수? cube 거리? wire
+  길이?) 모호.
+- 32 KiB 라는 기준 크기와 `[4 KiB, 16 KiB, 64 KiB, 256 KiB, 1 MiB]` sweep
+  의 의미가 코드 grep 으로만 확인 가능.
+
+## Decision
+
+### D1. 세 가지 case category — H2D / D2H / PE DMA
+
+각 category 는 토폴로지 상 별개의 데이터 경로를 가지며, 별도의 summary
+table + sweep table + route detail block 으로 출력된다.
+
+- **H2D (Host→Device Write)**: `MemoryWriteMsg(dst_sip=0, dst_cube,
+  dst_pe=0, pattern="zero")` 가 `pcie_ep → io_cpu → m_cpu → hbm_ctrl` 경로
+  를 흐른다. cube 인덱스로 hop 수가 증가:
+  - h2d-1hop: cube=0, hops=1
+  - h2d-2hop: cube=4, hops=2
+  - h2d-3hop: cube=8, hops=3
+  - h2d-4hop: cube=12, hops=4
+- **D2H (Device→Host Read)**: `MemoryReadMsg(src_sip=0, src_cube, src_pe=0)`.
+  forward command path + reverse data path 의 합 latency. 같은 4 hops
+  카테고리.
+- **PE DMA (PE-initiated)**: `PeDmaMsg(src_sip, src_cube, src_pe, dst_pa)`.
+  5 가지 케이스로 cube/PE 위치 변화:
+  - pe-local-hbm: same cube, same PE
+  - pe-same-half-hbm: same cube, different PE (PE 1)
+  - pe-cross-half-hbm: same cube, far PE (PE 4)
+  - pe-cross-cube-hbm-best: adjacent cube (cube 1)
+  - pe-cross-cube-hbm-worst: diagonal far cube (cube 15)
+
+cube 인덱스가 4/8/12 (H2D), 1/4/15 (PE DMA) 같이 의미 있는 이유는
+4x4 cube mesh (sip.cube_mesh.w=4, h=4) 에서의 거리 정의 — 추후 cube_mesh
+크기 변경 시 이 값들이 같이 갱신되어야 한다.
+
+### D2. 표준 측정 크기 — `nbytes = 32768` (32 KiB)
+
+모든 case 의 summary table 은 `nbytes=32768` 로 한 번 실행한 결과를
+보여준다. 32 KiB 가 선택된 이유:
+
+- DMA overhead 와 BW drain 이 한쪽으로 치우치지 않는 적당한 크기.
+- 다수 sub-unit (TCM, register file) 의 1회 transfer 단위와 비교 가능.
+
+크기별 utilization 변화는 별도 sweep table 이 보여준다 (D3).
+
+### D3. Utilization sweep — `[4 KiB, 16 KiB, 64 KiB, 256 KiB, 1 MiB]`
+
+`SWEEP_SIZES = [4096, 16384, 65536, 262144, 1048576]`, `SWEEP_LABELS =
+["4KB", "16KB", "64KB", "256KB", "1MB"]`. 매 size 마다 다음 공식:
+
+```
+drain   = nbytes / bottleneck_bw
+total   = overhead + wire + drain
+eff_bw  = nbytes / total
+util%   = eff_bw / bottleneck_bw × 100
+```
+
+`bn_bw is None or <= 0` 이면 그 컬럼은 0.0 % 로 출력. 의미: hop 수가 늘
+수록 작은 transfer 는 overhead-bound, 큰 transfer 는 drain-bound 가 되는
+패턴을 한 표에서 확인.
+
+### D4. 측정 항목 — actual / formula / breakdown
+
+각 case 행에 표시되는 컬럼:
+
+- `Actual` (total_ns): SimPy 실행 결과의 `trace["total_ns"]`.
+- `Ovhd`: 경로상 모든 node 의 `node.attrs["overhead_ns"]` 합 (formula
+  breakdown).
+- `Drain`: `nbytes / min(edge.bw_gbs over path)` (formula).
+- `Wire`: `Σ edge.distance_mm * (ns_per_mm from spec)`.
+- `Ovhd%` / `Drain%`: Ovhd/Drain 이 Actual 에서 차지하는 비율 (formula 의
+  Wire 는 통상 매우 작아 표시하지 않음).
+- `Eff.BW`: `nbytes / total_ns` (실 측정 BW).
+- `BN.BW`: bottleneck bandwidth (formula). path 상 모든 edge 의 BW 중 최소.
+  edge BW 가 없으면 "-".
+- `Util%`: `Eff.BW / BN.BW × 100`. 100% 면 single-stream BW upper bound 에
+  도달.
+
+formula 의 합 (`wire + ovhd + drain`) 과 actual 의 차이가 크면 모델
+단순화가 잡지 못하는 요소가 있다는 신호 (ADR-0033 의 가정 점검).
+
+### D5. Invariant 자동 체크 — PASS/FAIL
+
+다음 invariant 들이 자동으로 확인되어 `[v] PASS` / `[x] FAIL` 로 출력:
+
+- **H2D / D2H monotonic increase**: hop 수가 늘면 actual latency 가
+  단조 증가해야 함. `all(lats[i] < lats[i+1] for ...)`.
+- **D2H ≥ H2D**: 같은 hop 인덱스에서 D2H ≥ H2D (D2H 는 forward command
+  + reverse data 두 leg 이므로). `all(d2h[i].total >= h2d[i].total)`.
+- **PE DMA best < worst**: cross-cube best (adjacent) latency < cross-cube
+  worst (diagonal) latency.
+- **PE DMA local vs remote**: local BN BW vs remote BN BW 의 비교 출력
+  (PASS/FAIL 이 아닌 정보성).
+
+체크가 FAIL 이면 사람이 즉시 모델/토폴로지 회귀를 인지할 수 있도록 한
+줄로 분명하게 출력.
+
+### D6. Route detail — per-hop timestamp trace
+
+summary 와 sweep 표 이후 각 case 의 path 와 per-hop 누적 시간 (
+`_hop_timestamps`) 가 별도 섹션에서 출력된다:
+
+- H2D: leg1 (`pcie_ep → io_cpu`) + leg2 (`io_cpu → m_cpu`) + leg3
+  (`m_cpu → hbm_ctrl`) + per-hop trace.
+- D2H: forward (cmd, no data) + reverse (data) trace 분리 표시.
+- PE DMA: `pe_dma → router → hbm_ctrl` path + per-hop trace.
+
+각 hop 의 timestamp 는 cumulative `wire_ns + overhead_ns` 누적. terminal
+hop 의 annotation 에 `drain:Xns` 가 붙는다. bottleneck edge 는
+`<BN:XXGB/s>` 로 표시되어 시각적으로 식별 가능.
+
+### D7. case_filter 인자의 의미
+
+- `None` 또는 `"all"`: 모든 case 실행 (default).
+- 다른 문자열: 그 이름과 정확히 일치하는 case 만 실행. 예: `kernbench
+  probe --case h2d-2hop`.
+
+각 카테고리 안에서 `name != case_filter` 면 skip 되며, 그 카테고리의
+monotonicity / D2H≥H2D 비교는 데이터가 1개일 때 자연히 skip 된다.
+
+CLI parser 의 `--case` 기본값은 `"all"`이라 인자 생략 시 전체 실행.
+
+### D8. 매 case 별 fresh GraphEngine
+
+H2D 4개, D2H 4개, PE DMA 5개의 case 가 각각 **새로운 GraphEngine**
+인스턴스에서 실행된다 (`engine = GraphEngine(graph)`). 이유:
+
+- case 간 누적 상태 (op_log, completion 추적, allocator 등) 가 cross-talk
+  하지 않도록 격리.
+- 한 case 의 traffic 이 다른 case 의 BW 측정에 영향을 주지 않도록 보장.
+
+이 격리는 probe 의 측정 결과를 **각 case 단독 single-flow** 의 latency 로
+해석할 수 있게 한다. multi-flow contention 측정은 별도 도구 (예:
+`pe2pe_overview` 플롯, ADR-0033 의 multi-flow merging 모델) 책임.
+
+### D9. 출력 포맷의 안정성
+
+probe 의 stdout 출력은 사람이 읽기 위함이며, 정확한 컬럼 폭/구분자/공백 은
+machine-readable contract 가 아니다. 자동화된 도구가 probe 결과를 파싱
+하려면 별도 JSON 출력 모드를 추가해야 한다 (현재 미구현).
+
+PASS/FAIL 줄의 `[v]` / `[x]` 접두사는 CI grep 용 anchor 로 안정 보장.
+
+## Alternatives Considered
+
+### A1. Probe 를 별도 bench 로 등록 (`@bench(name="probe")`)
+
+기각. probe 는 bench 가 아니라 verification 도구로 의도된다 — sweep / 분석
+용 multi-engine 실행과 invariant PASS/FAIL 출력이 본질이며, ADR-0045 의
+"단일 디바이스 + 단일 RuntimeContext" bench 모델과 맞지 않는다.
+
+### A2. monotonicity 위반 시 exit code 1
+
+기각 (현재). 인간 검사 도구 위주로 의도되어 있어 PASS/FAIL 줄을 출력하고
+exit 0 로 종료. CI 가 violation 으로 fail 하길 원하면 별도 wrapper 가
+`grep "\[x\]"` 결과로 판단하면 됨. 후속으로 strict-mode flag (`--strict`)
+도입 가능.
+
+### A3. probe 의 case 정의를 외부 YAML 로
+
+기각 (현재). 8개 case (4 H2D + 4 D2H + 5 PE DMA — 합 13개) 는 코드에
+하드코딩되어 있고 의미가 토폴로지 mesh 구조에 단단히 묶여 있다. 외부
+YAML 로 옮기면 cube 인덱스의 의미 (4, 8, 12 / 1, 4, 15) 를 별도로 문서화
+해야 하므로 응집도 손실. 케이스 추가가 잦아지면 그때 별도 ADR 로 도입.
+
+### A4. multi-flow contention 측정 추가
+
+기각 (probe 범위 밖). D8 에서 명시한 single-flow 격리 모델이 probe 의 핵심
+의도. multi-flow contention 은 ADR-0033 latency model 의 다른 영역으로,
+별도 도구 또는 별도 case category 로 처리.
+
+## Consequences
+
+- probe 의 case catalog (D1) 와 측정 단위 (D2/D3) 가 ADR-level 에서 명시
+  되어, 새 traffic 카테고리 추가 시 어떤 표 포맷을 따라야 하는지 분명.
+- formula vs actual 의 컬럼 의미 (D4) 가 굳어져, probe 결과를 보고 "왜
+  Drain% 가 5% 인가 / 70% 인가" 같은 질문을 빠르게 ADR-0033 가정 점검으로
+  연결 가능.
+- invariant 자동 체크 (D5) 가 ADR 에 굳어져, 향후 latency 모델 변경 시
+  monotonicity / D2H≥H2D 회귀를 probe 가 즉시 잡아낸다는 안전망 정착.
+- D8 의 case 간 격리가 명시되어, probe 결과를 single-flow 측정으로 안전
+  하게 해석 가능. multi-flow 측정이 필요해지면 별도 도구 트랙이 필요함이
+  분명.
+- A2 의 strict-mode flag 가 후속 작업 후보로 기록되어, CI 통합 요구 시
+  최소 추가 작업으로 도입 가능.

docs/adr-ko/ADR-0050-par-ccl-algorithm-module-contract.md

@@ -0,0 +1,308 @@
+# ADR-0050: CCL Algorithm Module Contract — `ccl/algorithms/*.py`
+
+## Status
+
+Accepted (2026-05-22).
+
+`src/kernbench/ccl/algorithms/` 디렉터리 안의 모듈이 AHBM CCL backend
+(ADR-0047) 에서 collective algorithm 으로 사용되려면 갖춰야 할 인터페이스,
+kernel 시그너처, 그리고 새 알고리즘 추가 절차를 명시한다. ADR-0047 D3 가
+"algorithm 모듈은 `kernel`, `kernel_args`, optional `TOPO_NAME_TO_KIND` 를
+expose 해야 한다" 라고만 한 줄로 언급하나, **algorithm 모듈 작성자가 따라야
+할 contract** 는 ADR-level 에서 정리된 적이 없다. ADR-0045 가 bench 모듈
+contract 를 다루는 것과 짝을 이룬다.
+
+## First action (제일 처음에 하는 일)
+
+알고리즘 모듈이 import 되는 시점은 두 가지다:
+
+1. **AHBM backend 진입**: 사용자 코드가 `dist.init_process_group(backend="ahbm")`
+   를 호출하면, `AhbmCCLBackend.__init__` 안에서 `self._algo_module =
+   importlib.import_module(self._merged["module"])` 가 실행된다. 이때 모듈
+   레벨에서 가장 먼저 일어나는 일:
+   - `SIP_TOPO_RING/TORUS/MESH` 같은 정수 상수가 모듈 namespace 에 노출.
+   - `TOPO_NAME_TO_KIND` 사전이 모듈 namespace 에 노출 — backend 가
+     `topo_map = getattr(self._algo_module, "TOPO_NAME_TO_KIND", None)` 로
+     조회.
+   - `kernel_args` 함수 정의 — 호출 시 호출자가 사용.
+   - `allreduce_intercube_multidevice` 같은 알고리즘 함수 정의.
+   - 모듈 마지막 줄에서 `kernel = allreduce_intercube_multidevice` 로
+     alias 가 노출.
+
+2. **ccl.yaml install 단계**: `kernbench.ccl.install.install_ipcq` 가 호출
+   되어 IPCQ neighbor table 을 푸시할 때 같은 알고리즘 모듈이 import 됨.
+
+즉, **algorithm 모듈의 첫 일은 "topology-kind 상수, `TOPO_NAME_TO_KIND`
+사전, `kernel_args` 함수, 그리고 `kernel` alias 를 모듈 namespace 에 노출
+하는 것"** 이다. 모든 노출은 import-time 부수효과로 충분하며 별도 초기화
+함수 호출이 필요하지 않다.
+
+## Context
+
+`AhbmCCLBackend` (ADR-0047) 는 process group 초기화 시점에 `ccl.yaml` 의
+`defaults.algorithm` (또는 사용자가 지정한 알고리즘 이름) 으로부터 모듈
+경로를 얻어 dynamic import 한다. backend 는 그 모듈로부터 다음 4 가지를
+기대한다:
+
+- `kernel`: collective 의 진입 함수.
+- `kernel_args(world_size, n_elem, cube_w=, cube_h=) -> tuple`: kernel 에
+  넘길 위치 인자 묶음.
+- `TOPO_NAME_TO_KIND` (optional): `topology.yaml` 의 `sips.topology`
+  문자열 (예: `"ring_1d"`, `"torus_2d"`, `"mesh_2d_no_wrap"`) 을 정수
+  상수로 매핑하는 dict.
+- (간접) IPCQ neighbor table 설치: `configure_sfr_intercube_multisip` 가
+  알고리즘 모듈의 `TOPO_NAME_TO_KIND` 와 `cube_w/h` 를 보고 SFR 을 결정.
+
+현재 코퍼스의 유일한 algorithm 모듈은 `lrab_hierarchical_allreduce.py`
+(248 줄) 이다. 이름은 "**l**eft-**r**ight **a**lternating **b**roadcast
+**hierarchical allreduce**". 향후 `ring_allreduce`, `tree_allreduce`,
+`broadcast` 같은 모듈이 추가될 때마다 이 contract 를 따라야 일관된
+디스패치가 가능하다.
+
+이 contract 가 ADR-level 에 없으면:
+
+- 새 algorithm 작성자가 ADR-0047 D3 의 한 줄 만으로 시그너처를 추론해야.
+- kernel 함수 인자 순서 (특히 `t_ptr, n_elem, cube_w, cube_h, n_sips,
+  sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h, tl`) 의 의미가 코드
+  grep 없이는 명확하지 않다.
+- `kernel_args` 가 어떤 인자를 받고 어떤 tuple 을 돌려줘야 하는지 관례
+  로만 굳어진다.
+
+## Decision
+
+### D1. algorithm 모듈은 4 가지 public symbol 을 노출한다
+
+```python
+# src/kernbench/ccl/algorithms/<name>.py
+
+from __future__ import annotations
+
+# (필수) topology-kind 상수 — 알고리즘 내부에서 사용
+SIP_TOPO_RING = 0
+SIP_TOPO_TORUS = 1
+SIP_TOPO_MESH = 2
+
+# (선택) topology 이름 → kind 매핑. backend 가 ccl.yaml/topology 의
+# 문자열 SIP topology 를 정수로 변환하는 데 사용.
+TOPO_NAME_TO_KIND = {
+    "ring_1d": SIP_TOPO_RING,
+    "torus_2d": SIP_TOPO_TORUS,
+    "mesh_2d_no_wrap": SIP_TOPO_MESH,
+}
+
+# (필수) kernel 인자 빌더
+def kernel_args(world_size: int, n_elem: int, *, cube_w: int = 4, cube_h: int = 4) -> tuple:
+    return (n_elem, cube_w, cube_h, world_size)
+
+# (필수) kernel 함수 (`tl=...` 키워드를 통해 TLContext 가 주입됨)
+def my_allreduce_kernel(t_ptr, n_elem, cube_w, cube_h, n_sips,
+                         sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h, *, tl):
+    ...
+
+# (필수) kernel alias — backend 가 `module.kernel` 로 접근
+kernel = my_allreduce_kernel
+```
+
+- `kernel` alias 는 backend 가 직접 호출하는 entry point 다. 함수 이름이
+  무엇이든 (`allreduce_intercube_multidevice` 처럼) `module.kernel = fn`
+  으로 노출해야 한다.
+- `kernel_args` 가 없으면 backend 가 알고리즘 인자를 만들 방법이 없다.
+  signature 는 D2 참고.
+- `TOPO_NAME_TO_KIND` 가 없으면 backend 는 `sip_topo_kind = 0` 으로
+  fallback 한다. 단일 topology 만 지원하는 알고리즘이라면 생략 가능.
+
+### D2. `kernel_args` 시그너처 — `(world_size, n_elem, *, cube_w, cube_h)`
+
+```python
+def kernel_args(world_size: int, n_elem: int, *,
+                cube_w: int = 4, cube_h: int = 4) -> tuple:
+    return (n_elem, cube_w, cube_h, world_size)
+```
+
+- **위치 인자**: `world_size` (= rank 수), `n_elem` (= 단일 shard 의
+  element 수, f16 기준).
+- **키워드 인자**: `cube_w`, `cube_h` (= cube mesh 크기). default 는
+  4×4 — `topology.yaml` 의 `sip.cube_mesh` 기본값과 정합.
+- **반환**: kernel 의 위치 인자 순서대로 묶은 tuple.
+
+backend 의 `all_reduce` 가 호출 시:
+
+```python
+kernel_args_tuple = self._algo_module.kernel_args(
+    self._world_size, n_elem, cube_w=eff_cube_w, cube_h=eff_cube_h,
+)
+extra_args = (sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)
+pending = self.ctx.launch(
+    self._merged["algorithm"], kernel_fn, tensor,
+    *kernel_args_tuple, *extra_args, _defer_wait=True,
+)
+```
+
+즉 kernel 의 최종 위치 인자는: `(tensor_ptr, *kernel_args_tuple,
+sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)` 이며, 거기에 `tl=...` 가
+키워드로 자동 주입된다. `kernel_args` 가 돌려주는 tuple 의 길이/순서는
+**kernel signature 와 1:1 일치** 해야 한다.
+
+### D3. `kernel` 함수 시그너처 — 정형화된 9 + tl 인자
+
+권장 시그너처:
+
+```python
+def my_kernel(
+    t_ptr: int,             # VA base of the row-wise-sharded tensor on this SIP
+    n_elem: int,            # element count per cube tile (or per shard)
+    cube_w: int,            # cube mesh width  (kernel_args 에서 옴)
+    cube_h: int,            # cube mesh height (kernel_args 에서 옴)
+    n_sips: int,            # world_size 와 동일 (rank = SIP, ADR-0024)
+    sip_rank: int,          # 이 SIP 의 rank
+    sip_topo_kind: int,     # TOPO_NAME_TO_KIND lookup 결과
+    sip_topo_w: int,        # SIP mesh width (ring_1d 면 0)
+    sip_topo_h: int,        # SIP mesh height (ring_1d 면 0)
+    *, tl,                  # TLContext (auto-injected)
+) -> None:
+```
+
+`kernel_args` 가 다른 위치 인자 순서를 채택하더라도, kernel 의 **마지막
+4 개 위치 인자는 항상 `(sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)`**
+이며 backend 가 `extra_args` 로 append 한다 (ADR-0047 D5). 이 4 개 인자는
+사용자 정의 algorithm 도 받아야 하지만, 알고리즘이 single-SIP 이라면
+그냥 무시하면 된다.
+
+`tl` 은 위치 인자가 아닌 키워드로 주입된다 — `RuntimeContext.launch` 가
+kernel 호출 직전에 `tl=tl_ctx` 를 추가한다. 따라서 kernel signature 의
+`tl` 은 keyword-only (`*, tl`) 또는 마지막 키워드 매개변수 형태여야
+한다.
+
+### D4. kernel body 의 자유도와 제약
+
+kernel body 안에서 사용 가능한 표면: ADR-0046 D3 의 모든 `tl.*` primitive.
+
+특히 자주 쓰이는 패턴:
+
+- `cube_id = tl.program_id(axis=1)` — 이 PE 가 속한 cube 인덱스.
+- `pe_addr = t_ptr + cube_id * nbytes` — cube-별 tile 의 VA 계산.
+- `acc = tl.load(pe_addr, shape=(n_elem,), dtype="f16")` — local 데이터
+  로드.
+- `tl.send(dir=...)` / `tl.recv(dir=..., shape=, dtype=)` — IPCQ
+  collective.
+- `acc = acc + recv` — TensorHandle 산술 연산자 (ADR-0046 D4).
+- `tl.store(pe_addr, acc)` — 결과 저장.
+
+kernel body 는 일반 Python 함수이며, branching/looping 자유. 단:
+
+- SimPy `yield` 또는 `async` 금지 (ADR-0046 D1).
+- TensorHandle 의 `.data` 직접 접근 금지 — phase 1 timing 모델은
+  데이터 의존을 모른다 (ADR-0020 의 2-pass 분리).
+- kernel 실행은 deterministic 해야 한다 — 같은 입력으로 두 번 실행하면
+  같은 op 시퀀스 발사. random / external IO 금지.
+
+### D5. SIP topology semantics — `sip_topo_kind` 의 의미
+
+backend 가 `topology.yaml` 의 `system.sips.topology` 문자열을 algorithm
+모듈의 `TOPO_NAME_TO_KIND` 로 lookup 하여 `sip_topo_kind` 정수로 변환.
+algorithm 은 이 정수를 보고 분기:
+
+```python
+if sip_topo_kind == SIP_TOPO_RING:
+    acc = _inter_sip_ring(...)
+elif sip_topo_kind == SIP_TOPO_TORUS:
+    acc = _inter_sip_torus_2d(...)
+elif sip_topo_kind == SIP_TOPO_MESH:
+    acc = _inter_sip_mesh_2d(...)
+```
+
+각 topology branch 는 IPCQ direction 이름 (예: `"global_E"`, `"W"`, `"S"`,
+`"N"`) 을 통해 peer 와 통신. direction 의 의미는 ADR-0023/0025 가 정의
+하며, `configure_sfr_intercube_multisip` 가 IPCQ neighbor table 을 그에
+맞춰 설치한다.
+
+algorithm 모듈은 자기가 지원하지 않는 topology kind 가 들어오면 silent
+no-op 으로 두기보다 명시적으로 `raise ValueError(f"unsupported topology
+kind {sip_topo_kind}")` 하는 것을 권장 — 실수로 backend 에 잘못 dispatch
+된 경우 빠르게 fail.
+
+### D6. ccl.yaml 의 algorithm entry 구조
+
+algorithm 모듈은 `ccl.yaml` 의 entry 와 짝을 이룬다 (ADR-0023 D10 +
+ADR-0047 D3):
+
+```yaml
+defaults:
+  algorithm: lrab_hierarchical_allreduce
+  n_elem: 8
+
+algorithms:
+  lrab_hierarchical_allreduce:
+    module: kernbench.ccl.algorithms.lrab_hierarchical_allreduce
+    # optional: world_size override
+    # optional: per-algorithm parameters consumed by configure_sfr_intercube_multisip
+```
+
+- `module`: full Python module path. backend 의 `importlib.import_module`
+  가 이 문자열을 그대로 사용.
+- `world_size` (optional): 명시되면 topology fallback 을 override
+  (ADR-0047 D2).
+- algorithm-specific parameters 는 `configure_sfr_intercube_multisip` 가
+  소비.
+
+새 algorithm 추가 시:
+
+1. `src/kernbench/ccl/algorithms/<name>.py` 작성 (D1 컨벤션).
+2. `ccl.yaml` 의 `algorithms` 섹션에 entry 추가.
+3. (필요 시) `kernbench.ccl.sfr_config` 에 SFR 설치 분기 추가.
+4. test 추가 (예: `tests/sccl/test_<name>.py`, ADR-0043 의 eval harness
+   확장).
+
+### D7. legacy "rank = flat PE index" 모드
+
+ADR-0047 D2 가 명시한 `ccl.yaml` 의 `world_size` override 경로는 legacy
+"rank = flat PE index" 테스트가 사용한다. algorithm 모듈은 이 모드 에서도
+`n_sips=world_size` 만큼의 rank 가 들어옴을 가정하면 된다 — backend 가
+rank↔(SIP, cube, PE) 매핑을 사전에 분리해 두므로 algorithm 본체에서는
+modal 분기가 필요 없다.
+
+단, single-cube workload 에서는 `cube_w=cube_h=1` 이 들어와 mesh-기반
+phase 들이 skip 되도록 작성해야 한다 (`lrab_hierarchical_allreduce.py`
+의 `single_cube = (cube_w == 1 and cube_h == 1)` 패턴 참고).
+
+## Alternatives Considered
+
+### A1. algorithm 모듈을 class 로 구조화 (`class Allreduce: kernel(...)` 등)
+
+기각. Python 모듈 namespace 자체가 algorithm 의 identity 로 사용 중이며
+(ADR-0047 D3 의 `importlib.import_module`), class 한 겹은 추가 indirection
+만 늘리고 dispatch 측 코드를 두텁게 만든다. 모듈-레벨 free function
++ `kernel` alias 패턴이 충분히 명확.
+
+### A2. `kernel_args` 를 명시적 dataclass 로 typing
+
+기각 (현재). algorithm 마다 인자 갯수가 다른 것이 정상이며, dataclass 한
+종류를 강제하면 다양한 algorithm 간 호환이 어려워진다. tuple 반환은 simple
+하고 backend 측 `*kernel_args_tuple` unpacking 과 깨끗이 맞물린다.
+algorithm 별 자체 타입 강도가 필요해지면 그 algorithm 모듈 안에서 NamedTuple
+사용은 자유.
+
+### A3. SFR 설치를 algorithm 모듈 안으로
+
+기각. SFR 설치 (`configure_sfr_intercube_multisip`) 는 topology + algorithm
+모두를 보고 IPCQ neighbor table 을 설치하는 cross-module 결정이라, algorithm
+모듈 내부보다 `kernbench.ccl.sfr_config` 같은 전용 위치가 자연스럽다. D6 의
+"필요 시 sfr_config 분기 추가" 워크플로우가 책임 분리 측면에서 더 명확.
+
+### A4. algorithm name 을 모듈 namespace 에 자동 등록 (decorator)
+
+기각. ADR-0045 (bench) 와 달리 algorithm 은 ccl.yaml entry 와 직접 묶여
+있어 추가 등록 레지스트리가 중복이다. `module` 문자열 매핑 하나면 충분.
+
+## Consequences
+
+- ADR-0047 D3 의 한 줄 contract 가 D1–D7 의 작성자-친화적 가이드로 확장
+  되어, 새 algorithm 추가 시 시그너처를 grep 으로 추론할 필요 없음.
+- D3 의 9 + tl 인자 시그너처가 표준화되어, backend 의 `extra_args` append
+  (ADR-0047 D5) 와 자연스럽게 맞물림. 향후 single-SIP-only algorithm 도
+  4 개의 sip_* 인자를 받아야 함이 명시.
+- D5 의 fail-loud 권장으로, ccl.yaml 의 topology 가 algorithm 미지원
+  topology 로 잘못 설정되면 backend 가 silent wrong-result 가 아닌
+  ValueError 로 fail.
+- D6 의 단계별 추가 절차가 명시되어, 새 algorithm 추가가 sfr_config /
+  test / ccl.yaml 어디까지 손대야 하는지 분명.

docs/adr-ko/ADR-0051-lat-routing-helper-api.md

@@ -0,0 +1,267 @@
+# ADR-0051: Routing Helper API — `AddressResolver` + `PathRouter`
+
+## Status
+
+Accepted (2026-05-22).
+
+`policy/routing/router.py` 가 노출하는 두 helper 클래스
+(`AddressResolver`, `PathRouter`) 의 모든 public API, 인자, 반환 값,
+그리고 네 가지 다른 adjacency graph 의 사용처를 명시한다. ADR-0002 가
+routing distance 와 ordering, bypass 규칙을 정의하나, **helper API 표면
+자체** 는 ADR-level 에 정리된 적이 없다.
+
+## First action (제일 처음에 하는 일)
+
+### `AddressResolver(graph)`
+
+생성 즉시 다음 두 가지를 캐시한다:
+
+1. `self._node_ids = set(graph.nodes)` — 모든 node id 의 set (lookup 용).
+2. `self._hbm_slice_bytes = hbm_total_gb * (1 << 30) // slices_per_cube` —
+   `graph.spec.cube.memory_map` 으로부터 산출 (기본 `48 GB / 8 slices = 6
+   GB`). 이 값이 `resolve()` 가 HBM PA 의 `hbm_offset` 에서 `pe_id` 를
+   복원하는 데 쓰인다.
+
+즉, **AddressResolver 의 첫 일은 "전체 node id 집합과 HBM slice 크기를
+미리 계산해 두는 것"** 이다. graph 자체는 보유하지 않는다.
+
+### `PathRouter(graph)`
+
+생성 즉시 **네 개의 별도 adjacency graph 를 동시 구축**한다:
+
+1. `self._adj_all`: 모든 edge 포함 (component-to-component routing 용).
+2. `self._adj`: `kind != "command"` 인 edge 만 (PE DMA / 일반 data path).
+3. `self._adj_mcpu_dma`: `_MCPU_DMA_EXCLUDE = {"pe_internal",
+   "pe_to_router"}` 를 제외 (M_CPU DMA 가 PE pipeline 노드로 잘못 라우팅
+   되지 않게).
+4. `self._adj_local`: `_UCIE_KINDS` 8 종을 제외 (cube-local routing 용 —
+   UCIe 가 zero-distance bus 처럼 보여 Dijkstra 가 mesh 보다 선호하는
+   것을 막음).
+
+각 그래프는 `defaultdict(list)` of `(neighbor, weight)` 형태이며,
+`edge.routing_weight_mm or edge.distance_mm` 이 weight 로 쓰인다.
+
+즉, **PathRouter 의 첫 일은 "topology edge 들을 4개의 다른 정책으로 동시
+분류하여 4 개의 인접 리스트로 구축하는 것"**. 매 `find_*()` 호출 시 적절
+한 그래프를 골라 Dijkstra 를 돌린다.
+
+## Context
+
+`policy/routing/router.py` 는 다음 두 책임을 함께 수행한다:
+
+- **이름 매핑**: 토폴로지 명명 규칙 (`sip{S}.cube{C}.<comp>`,
+  `sip{S}.io{I}.pcie_ep` 등) 의 단일 소유자. 컴포넌트 / probe / IPCQ
+  install / runtime API 가 이름 문자열을 직접 만들지 않고 helper 를 호출.
+- **경로 결정**: edge 의 `kind` 에 따른 정책 분리. 같은 src→dst 라도
+  routing 의도 (PE DMA vs M_CPU DMA vs general component routing) 에 따라
+  다른 adjacency 를 사용해야 결과가 달라진다.
+
+이 helper API 가 코드 전반에서 광범위하게 소비되는데도 (probe.py /
+distributed.py / install.py / 각종 component / tests), ADR-level 에서
+**정확한 시그너처 / 반환 의미 / 어떤 adjacency 를 쓰는지** 가 한 곳에
+정리되어 있지 않다. 본 ADR 이 그 빈자리를 채운다.
+
+## Decision
+
+### D1. `AddressResolver` 의 5 개 public API
+
+#### D1.1. `resolve(addr: PhysAddr) -> str`
+
+`PhysAddr` 인스턴스를 토폴로지의 destination node id 로 변환.
+
+```
+addr.kind == "hbm"             → f"sip{s}.cube{d}.hbm_ctrl.pe{pe_id}"
+  where pe_id = addr.hbm_offset // self._hbm_slice_bytes  (ADR-0017 D4/D9)
+
+addr.kind == "pe_resource":
+  addr.unit_type == PE         → f"sip{s}.cube{d}.pe{addr.pe_id}.pe_tcm"
+  addr.unit_type == SRAM       → f"sip{s}.cube{d}.sram"
+  addr.unit_type == MCPU       → f"sip{s}.cube{d}.m_cpu"
+  그 외                          → RoutingError("unsupported unit_type")
+
+다른 kind                       → RoutingError("unsupported address kind")
+```
+
+산출된 node id 가 `self._node_ids` 에 없으면 `RoutingError(f"node {node_id}
+not found in topology")`. 즉, address 의 syntax 가 valid 해도 topology 에
+실제로 매핑되는 노드가 없으면 fail-loud.
+
+#### D1.2. `find_m_cpu(sip, cube) -> str`
+
+`f"sip{sip}.cube{cube}.m_cpu"`. 없으면 `RoutingError`.
+
+#### D1.3. `find_pcie_ep(sip, io_id="io0") -> str`
+
+`f"sip{sip}.{io_id}.pcie_ep"`. 없으면 `RoutingError`.
+
+#### D1.4. `find_io_cpu(sip, io_id="io0") -> str`
+
+`f"sip{sip}.{io_id}.io_cpu"`. 없으면 `RoutingError`.
+
+#### D1.5. `find_all_pcie_eps() -> list[str]`
+
+전 SIP 의 PCIE_EP node id 를 정렬된 리스트로 반환. `endswith(".pcie_ep")`
+필터링. cross-SIP IPCQ 가 모든 PCIE_EP 를 enumerate 할 때 사용.
+
+명명 규칙 (`sip{S}.cube{C}.<comp>`, `sip{S}.{io_id}.<comp>`) 의 단일
+소유자가 이 클래스다 (ADR-0015 D4). 토폴로지 빌더가 같은 명명 규칙으로
+노드를 만들고, 컴포넌트는 이름 문자열을 절대 직접 구성하지 않는다 —
+모두 helper 를 거친다.
+
+### D2. `PathRouter` 의 4 개 adjacency graph
+
+생성자가 한 번에 구축. edge `kind` 가 정책을 결정:
+
+| graph             | 제외 edge kinds                               | 용도                                       |
+|-------------------|-----------------------------------------------|--------------------------------------------|
+| `_adj_all`        | (none)                                        | M_CPU↔NOC command 포함, IO_CPU/M_CPU routing |
+| `_adj`            | `"command"`                                   | PE DMA / 일반 data path                    |
+| `_adj_mcpu_dma`   | `"pe_internal"`, `"pe_to_router"`            | M_CPU DMA (PE pipeline 우회)               |
+| `_adj_local`      | `_UCIE_KINDS` (`ucie_internal`, `ucie_conn_to_router`, `router_to_ucie_conn`, `ucie_conn_to_noc`, `noc_to_ucie_conn`, `ucie_mesh`, `io_to_cube`, `cube_to_io`) | same-cube routing (UCIe bus 우회) |
+
+각 그래프는 `dict[node_id, list[(neighbor, weight)]]` 이며, weight 는
+`edge.routing_weight_mm or edge.distance_mm`. command edge 의 routing
+영향력을 명시적으로 가르고, UCIe 의 "0-distance bus" 가 mesh 보다 선호
+되는 것을 막기 위한 `_adj_local` 분리가 ADR-0017 D7 의 cross-PE-slice
+mesh-distance 요구와 정합.
+
+### D3. `PathRouter` 의 6 개 public API (+ 2 backward-compat)
+
+#### D3.1. `find_path(src_pe: str, dst_node: str) -> list[str]`
+
+**PE DMA routing**. `src_pe` 는 PE prefix (예: `"sip0.cube0.pe0"`) 이며,
+함수가 `.pe_dma` 를 자동으로 prepend 하여 실제 시작 노드를
+`"sip0.cube0.pe0.pe_dma"` 로 설정.
+
+cube-local 여부 (`_same_cube`) 에 따라 adjacency 선택:
+
+- **same-cube** (src 와 dst 가 `sip{S}.cube{C}.` prefix 공유):
+  `_adj_local` 사용. UCIe 우회를 막아 cross-PE-slice 가 mesh 거리를 정확
+  히 지불 (ADR-0017 D7).
+- **cross-cube**: `_adj` 사용. UCIe 가 자연스럽게 cross-cube path 의
+  최적 선택지로 포함됨.
+
+#### D3.2. `find_path_with_distance(src_pe, dst_node) -> tuple[list[str], float]`
+
+D3.1 과 동일한 adjacency 정책을 사용하나, 결과로 `(path, total_distance)`
+를 함께 반환. probe / 분석 도구에서 distance 메트릭이 필요할 때 사용.
+
+#### D3.3. `find_mcpu_dma_path(m_cpu_id: str, dst_hbm_id: str) -> list[str]`
+
+**M_CPU DMA path**. cube 가 같으면 `_adj_local` (mesh 안에서 마무리), 다르
+면 `_adj_all` (UCIe 경유). `_MCPU_DMA_EXCLUDE` 가 PE pipeline 노드를 자동
+배제하므로, M_CPU 가 PE 의 내부 stage 를 거쳐 routing 되는 잘못된 경로가
+나오지 않는다.
+
+#### D3.4. `find_memory_path(src: str, dst: str) -> list[str]`
+
+`pcie_ep → io_noc → cube → router mesh → hbm_ctrl` 같은 직접 메모리
+경로. `_adj_mcpu_dma` 를 사용하여 `pe_internal` 및 `pe_to_router` edge
+를 제외 — host-issued read/write 가 PE pipeline 으로 새지 않게 보장.
+probe (ADR-0049 D1 의 H2D/D2H case) 에서 직접 호출.
+
+#### D3.5. `find_node_path(src: str, dst: str) -> list[str]`
+
+임의의 두 node 사이의 path. **command edge 포함** (`_adj_all` 사용). M_CPU
+↔ NOC 같은 command-kind link 를 거쳐야 하는 IoCpuComponent /
+MCpuComponent 등이 호출.
+
+#### D3.6. backward-compat shims
+
+- `_dijkstra(start, goal) -> list[str]` — `_run_dijkstra(self._adj, …)`
+  의 thin wrapper.
+- `_dijkstra_with_dist(start, goal) -> tuple[list[str], float]` — distance
+  포함 버전.
+
+언더스코어 prefix 에서 보듯이 내부 API 인 척이지만 기존 테스트가 직접
+호출. 새 코드는 D3.1–D3.5 를 사용하고, 이 두 shim 은 deprecation 후보.
+
+### D4. Dijkstra 알고리즘 — single-source shortest path
+
+`_run_dijkstra_with_dist(adj, start, goal)`:
+
+- `heapq` priority queue.
+- `best: dict[node, distance]` — 노드별 최단 거리 캐시.
+- `prev: dict[node, predecessor]` — path reconstruction.
+- weight 는 `routing_weight_mm or distance_mm`. UCIe 처럼 routing_weight 가
+  명시되어 distance 와 다른 edge 가 있으므로 weight 분리가 의도된 것.
+
+`start == goal` 은 빠른 path `([start], 0.0)` 반환. 도달 불가는
+`RoutingError(f"no path from {start} to {goal}")`.
+
+이 알고리즘은 **deterministic** 하다 — 같은 graph + start/goal 이면 같은
+경로. 이는 SPEC R1 의 "Routing MUST be deterministic" 요구와 정합. tie-
+break 는 `heapq` 의 push 순서를 따른다 (Python list 순서가 deterministic).
+
+### D5. helper API 의 단일 소유자 원칙
+
+다음 정보는 오직 router.py 안에서만 결정된다:
+
+- 명명 규칙: `sip{S}.cube{C}.<comp>`, `sip{S}.{io_id}.<comp>`,
+  `sip{S}.cube{C}.hbm_ctrl.pe{pe_id}`.
+- adjacency 정책: 어떤 edge kind 가 어떤 그래프에 포함되는가.
+- HBM slice 크기로부터 PE id 복원 방법.
+- Dijkstra의 weight 결정 (`routing_weight_mm or distance_mm`).
+
+이 단일 소유자 원칙이 깨지면 (예: 컴포넌트가 자체적으로 `f"sip{s}..."` 를
+구성하기 시작하면) 명명 규칙 변경 시 영향 범위가 폭발한다. ADR-0015 D4 의
+정신과 정렬.
+
+### D6. helper API consumer 의 목록
+
+본 helper 가 노출하는 메소드를 호출하는 곳을 명시 (현재 코퍼스 기준):
+
+- `probes/probe.py` (ADR-0049): `find_pcie_ep`, `find_io_cpu`,
+  `find_m_cpu`, `find_node_path`, `find_mcpu_dma_path`,
+  `find_memory_path`, `find_path`, `resolve`.
+- `runtime_api/distributed.py` (ADR-0047): 간접 (engine 내부 routing).
+- `ccl/install.py` (ADR-0023): `find_all_pcie_eps`, `resolve`.
+- `sim_engine/event_log.py`: probe 와 유사하게 `find_pcie_ep`,
+  `find_memory_path`.
+- `components/builtin/m_cpu.py`, `components/builtin/io_cpu.py`:
+  `find_node_path`, `find_mcpu_dma_path`.
+- 각종 tests (test_routing.py, test_cross_sip_routing.py 등): D3.1–D3.5
+  대부분.
+
+새 consumer 가 추가될 때 본 ADR 의 D1/D3 가 그 의도에 맞는 메소드가
+이미 있는지 / 새 메소드를 추가해야 하는지 1차 판단의 기준이 된다.
+
+## Alternatives Considered
+
+### A1. 단일 adjacency graph + edge-kind filter 동적 적용
+
+기각. 매 `find_*()` 마다 graph filtering 을 다시 하면 Dijkstra 의 cache
+locality 와 성능이 떨어진다. 4 개 그래프 동시 구축 (D2) 은 메모리 비용
+이 작고 (edge ≤ 수만 건 규모), 호출 시점에 정책 선택이 O(1) 로 결정.
+
+### A2. adjacency 분리를 edge 의 `kind` 가 아닌 별도 metadata 로
+
+기각. edge `kind` 는 이미 topology builder 가 부여하며 (ADR-0015 D4 +
+ADR-0017), 별도 metadata 를 도입하면 두 시스템이 동기화되어야 하는
+중복이 생긴다.
+
+### A3. Dijkstra 대신 BFS + uniform weight
+
+기각. routing_weight_mm 이 edge 별로 다른 (mesh link / UCIe / IO-internal)
+현실에서 BFS 는 hop 수 최소화일 뿐 latency / distance 최단을 보장하지
+않는다. SPEC R1 + R2 의 결정적·정확한 routing 요구에 어긋남.
+
+### A4. helper API 를 클래스 메서드가 아닌 모듈 함수로
+
+기각. 두 클래스 (`AddressResolver`, `PathRouter`) 가 각각 cache 상태
+(`_node_ids`, `_hbm_slice_bytes`, 4 adjacency graphs) 를 보유해야 하며,
+같은 graph 인스턴스에 여러 routing 질의가 발생한다. 모듈 함수는 매 호출
+시 state 를 다시 만들거나 global 로 두어야 해서 안전성/성능 저하.
+
+## Consequences
+
+- 컴포넌트 / probe / IPCQ install / runtime API 가 모두 router.py 의
+  helper 만 호출하면 명명 규칙 변경 (예: `.io0.` → `.iochiplet0.`) 이
+  단 한 파일 수정으로 끝남 (D5).
+- D2 의 4 그래프 분리가 ADR 에 굳어져, 새 edge kind 가 추가될 때 (예:
+  Inter-die UCIe link 의 새 kind) 어느 그래프에 포함시킬지 결정의 명확
+  한 기준 제공.
+- D3.1 의 cube-local vs cross-cube 분기 (ADR-0017 D7) 가 명시되어, 향후
+  routing 동작을 변경하려는 사람이 어느 adjacency 를 건드려야 할지 안다.
+- D6 의 consumer 목록이 명시되어, helper API 변경 시 PR review 범위가
+  분명. backward-compat shim (D3.6) 의 deprecation 후보가 식별됨.

docs/adr-ko/ADR-0052-dev-oplog-memory-store-schemas.md

@@ -0,0 +1,352 @@
+# ADR-0052: OpLog + MemoryStore Schemas — sim_engine internals
+
+## Status
+
+Accepted (2026-05-22).
+
+`sim_engine/op_log.py` 의 `OpRecord` 스키마와 `OpLogger` 의 record_start /
+record_end / record_copy 동작, 그리고 `sim_engine/memory_store.py` 의
+`MemoryStore` 가 사용하는 (space, addr) 주소공간 namespace 와 read/write
+의미를 명시한다. ADR-0020 (2-pass data execution) 가 두 인프라의 존재를
+선언하나, **레코드의 정확한 필드와 의미** 는 ADR-level 에서 정리되지
+않았고 ADR-0046 D3.2 (`tl.store` visibility), ADR-0023 D9 (IPCQ copy
+record) 등 여러 ADR 이 이들의 동작에 의존하고 있다.
+
+## First action (제일 처음에 하는 일)
+
+### `OpLogger(memory_store=None)`
+
+생성 즉시 다음 3 가지 필드 초기화:
+
+1. `self._records: list[OpRecord] = []` — 누적된 op record.
+2. `self._pending: dict[int, dict] = {}` — `id(msg)` 키로 partial record
+   (record_start 시점에 만들어졌고 record_end 가 아직 안 온 것).
+3. `self._memory_store = memory_store` — 옵션 MemoryStore reference.
+   math op 의 input 스냅샷 + dma_write 의 HBM source 스냅샷 캡처에 사용.
+
+생성 시점에는 records / pending 모두 비어 있으며, `record_*` 호출이
+순차적으로 데이터를 누적한다.
+
+### `MemoryStore()`
+
+생성 즉시 `self._storage: dict[str, dict[int, np.ndarray]] = {}` 단 하나
+의 필드 초기화. 두 단계 dict (`space → addr → ndarray`) 이며 lazy 하게
+필요한 space 가 생길 때마다 inner dict 가 채워진다.
+
+즉, **두 인프라의 첫 일은 "비어 있는 누적 buffer + space-별 sparse dict
+를 만들어 두는 것"** 이다. 첫 record / write 가 실제로 도착하면 그때
+필드가 채워지기 시작한다.
+
+## Context
+
+ADR-0020 (2-pass data execution) 의 D2/D5/D7 가 다음을 선언:
+
+- Phase 1 (timing) 동안 `ComponentBase._on_process_start/end` hook 이
+  `OpLogger.record_start/end` 를 호출하여 모든 data op 의 시간 + 메타
+  데이터를 기록.
+- Phase 2 (data) 가 op_log 를 t_start 순으로 재생하여 실 데이터 결과를
+  계산.
+- 데이터 페이로드 자체는 `MemoryStore` 에 (space, addr) 키로 보관.
+
+ADR-0023 D9 (IPCQ atomic write), ADR-0027 (Megatron TP scratch
+overwrite 회피), ADR-0046 D3.2 (`tl.store` visibility) 등 후속 ADR 들이
+op_log 와 MemoryStore 의 동작에 의존하지만, **정확한 record 필드 / space
+이름 / 스냅샷 시점** 은 코드 grep 으로만 확인 가능하다. 본 ADR 이 이를
+정리한다.
+
+## Decision
+
+### D1. `OpRecord` 스키마 — 7 개 필드
+
+```python
+@dataclass
+class OpRecord:
+    t_start: float
+    t_end: float
+    component_id: str
+    op_kind: str               # "memory" | "gemm" | "math" | "unknown"
+    op_name: str               # e.g. "dma_read", "gemm_f16", "exp",
+                               #     "TileToken/DMA_READ", "composite_gemm",
+                               #     "ipcq_copy"
+    params: dict[str, Any]
+    dependency_ids: list[int] = field(default_factory=list)
+```
+
+- **`t_start` / `t_end`**: SimPy 시간 (float ns). `t_start` 는 component
+  가 op 를 시작한 시점, `t_end` 는 완료 시점. duration = `t_end - t_start`.
+- **`component_id`**: op 가 발생한 node id (예:
+  `"sip0.cube0.pe0.pe_dma"`).
+- **`op_kind`**: 4 가지 중 하나. Phase 2 DataExecutor 가 이 값으로 분기.
+- **`op_name`**: 디버깅 / 분석용 사람-친화 이름. TileToken 일 경우
+  `"TileToken/{stage_type}"` (예: `"TileToken/DMA_READ"`) 로 stage 를
+  구분.
+- **`params`**: op-종속 메타데이터 dict (D3 참고).
+- **`dependency_ids`**: 현재 사용되지 않음 (default `[]`). 향후 cross-op
+  dependency 추적이 필요해질 때를 위한 자리.
+
+### D2. `OpLogger.records` — t_start 정렬 보장
+
+```python
+@property
+def records(self) -> list[OpRecord]:
+    self._records.sort(key=lambda r: r.t_start)
+    return self._records
+```
+
+매 접근 시 `t_start` 로 stable sort. 즉 같은 t_start 인 record 들은 insertion
+순서를 유지. ADR-0020 D5 의 "t_start stable ordering" 요구와 정합.
+
+Phase 2 DataExecutor 는 항상 `records` property 를 통해 접근하므로,
+record_end 호출이 t_start 와 다른 순서로 도착해도 (예: 짧은 op 가 긴
+op 보다 늦게 시작했으나 먼저 끝남) 재정렬되어 일관된 시퀀스를 받는다.
+
+### D3. op_name 별 `params` 스키마 (`_extract_op_info` 매핑)
+
+#### D3.1. `op_kind="memory", op_name="dma_read"` (DmaReadCmd)
+
+```python
+{"src_addr": int, "nbytes": int, "handle_id": str}
+```
+
+#### D3.2. `op_kind="memory", op_name="dma_write"` (DmaWriteCmd)
+
+```python
+{
+    "src_space": str,   # handle.space ("tcm"|"hbm"|"sram"), default "tcm"
+    "src_addr": int,    # handle.addr
+    "shape": tuple, "dtype": str,
+    "dst_space": "hbm", # DmaWrite 는 항상 HBM 으로
+    "dst_addr": int,
+    "nbytes": int,
+    "handle_id": str,
+    # record_end 시점에 src_space == "hbm" 이면 snapshot 추가 (D4)
+    "snapshot": np.ndarray | None,
+}
+```
+
+#### D3.3. `op_kind="gemm", op_name=f"gemm_{dtype_a}"` (GemmCmd)
+
+```python
+{
+    "src_a_addr": int, "src_b_addr": int, "dst_addr": int,
+    "shape_a": tuple, "shape_b": tuple, "shape_out": tuple,
+    "dtype_in": str, "dtype_out": str,
+    "m": int, "k": int, "n": int,
+    # ADR-0027: per-operand + output spaces 보존
+    "src_a_space": str, "src_b_space": str, "dst_space": str,
+}
+```
+
+#### D3.4. `op_kind="math", op_name=msg.op` (MathCmd; op = "exp", "sum", "add", "where" 등)
+
+```python
+{
+    "input_addrs": list[int],   # 입력 핸들들의 addr
+    "input_shapes": list[tuple],
+    "input_spaces": list[str],
+    "input_dtypes": list[str],
+    "dst_addr": int, "dst_space": str,
+    "shape_out": tuple, "dtype": str,
+    "axis": int | None,         # reduction 인 경우만 의미 있음
+    # record_end 시점에 모든 input 의 스냅샷이 채워짐 (D4)
+    "input_snapshots": list[np.ndarray | None],
+}
+```
+
+#### D3.5. `op_kind="gemm" or "math", op_name=f"composite_{op}"` (CompositeCmd)
+
+```python
+{
+    "op": str,              # "gemm" | "math"
+    "out_addr": int, "out_nbytes": int,
+    # op == "gemm" 인 경우 GemmCmd 와 같은 필드 추가:
+    "src_a_addr": int, "src_b_addr": int,
+    "shape_a": tuple, "shape_b": tuple,
+    "dtype_in": str, "dtype_out": str,
+    "src_a_space": str, "src_b_space": str,
+    "dst_space": "hbm", "dst_addr": int,  # = out_addr
+}
+```
+
+`op == "gemm"` 이면 `op_kind = "gemm"`, 아니면 `"math"`. Phase 2 측에서
+GemmCmd 와 동일 path 로 재생되도록 alias.
+
+#### D3.6. `op_kind="memory", op_name="ipcq_copy"` (record_copy 전용 경로)
+
+```python
+{
+    "src_space": str, "src_addr": int,
+    "dst_space": str, "dst_addr": int,
+    "shape": tuple, "dtype": str, "nbytes": int,
+    "snapshot": np.ndarray | None,   # 호출자가 전달, 없으면 record_copy 가 fresh read
+}
+```
+
+`PE_DMA._handle_ipcq_inbound` (ADR-0023 D9) 가 이 record 를 발사하여 IPCQ
+slot 의 inbound copy 를 Phase 2 가 재생 가능하게 한다. 이 record 는
+`record_start` / `record_end` 를 거치지 않고 직접 `record_copy()` 로 push.
+
+#### D3.7. `op_kind="unknown", op_name=type(msg).__name__`
+
+`_extract_op_info` 가 인식 못 한 message 의 fallback. params = `{}`.
+DataExecutor 가 이 op_kind 를 만나면 skip — Phase 2 replay 에 영향 없음.
+
+### D4. snapshot 캡처 시점
+
+`OpLogger._memory_store` 가 set 되어 있을 때 record_end 가 다음을 수행:
+
+- **math op**: 모든 input addr/shape/space/dtype 으로
+  `self._memory_store.read(...)` 를 호출하여 `params["input_snapshots"]` 에
+  ndarray copy 첨부. read 실패 시 None.
+- **dma_write op**: `src_space == "hbm"` 인 경우에만 source HBM 의
+  스냅샷을 `params["snapshot"]` 에 첨부. TCM source 는 **명시적으로
+  스킵** — TCM (PE scratch) 은 Phase 2 math/gemm 재생이 다시 채우므로,
+  Phase-1-time snapshot 을 잡으면 이전 kernel 의 stale 데이터를 잡을 위험
+  (ADR-0027 postmortem: TP gemm → all_reduce race).
+- **ipcq_copy**: `record_copy` 호출자가 `snapshot=token.data` 같이 in-flight
+  스냅샷을 전달. 없으면 record_copy 가 fresh read 로 대체 시도.
+
+스냅샷은 `.copy()` 가 호출되어 (`ndarray.copy()` 가 fresh allocation) 이후
+storage mutation 으로부터 안전. ADR-0027 의 "cross-PE Phase 2 ordering"
+race 회피의 근간.
+
+`memory_store` 가 None 인 경우 (Phase 1 timing-only 모드) 스냅샷 단계는
+전부 skip. record 의 timing 정보만 보존되며 데이터 replay 는 불가능.
+
+### D5. TileToken 처리 — record_start 가 stage 정보를 캡처
+
+ADR-0014 D6 의 self-routing tile token (pipeline 모드) 은 stage_idx 가
+record_end 시점에 이미 advance 되어 있을 수 있다 (TileToken 이 다음
+component 로 이동하면서 next stage 의 params 를 캐시). 따라서:
+
+`record_start` 가 다음을 `pending[id(msg)]["snap"]` 에 미리 저장:
+
+```python
+snap["stage_type"] = stage.stage_type.name        # "DMA_READ", "GEMM", 등
+snap["stage_params"] = dict(stage.params)         # 시점의 params 복사본
+```
+
+`record_end` 에서 이 snap 을 꺼내 params 에 merge:
+
+- `params["stage_type"]` 가 final params 에 추가.
+- `stage_params` 의 key 들이 (이미 있으면 보존) merge.
+- `op_name == "TileToken"` 이면 `op_name = f"TileToken/{stage_type}"` 로
+  rewrite (예: `"TileToken/DMA_READ"`) — 같은 component 에서 발생한 서로
+  다른 stage 의 record 를 disambiguate.
+
+이 메커니즘 덕분에 DMA_READ vs DMA_WRITE, FETCH vs STORE 가 같은 component
+(예: pe_dma) 에서 발생하더라도 reporting 측에서 구분 가능.
+
+### D6. `MemoryStore` — (space, addr) 두 단계 dict
+
+```python
+class MemoryStore:
+    def __init__(self) -> None:
+        self._storage: dict[str, dict[int, np.ndarray]] = {}
+
+    def write(self, space, addr, data): self._storage[space][addr] = data
+    def read(self, space, addr, shape=None, dtype=None) -> np.ndarray: ...
+    def has(self, space, addr) -> bool: ...
+    def snapshot(self) -> MemoryStore: ...
+```
+
+#### D6.1. space namespace
+
+문자열 키. 표준 값:
+
+- `"hbm"`: HBM 데이터 (deploy_tensor + Phase 2 dma_write 결과).
+- `"tcm"`: PE-로컬 TCM (Phase 2 math/gemm 결과).
+- `"sram"`: cube-level SRAM (ADR-0023 D9.7 IPCQ slot tier).
+
+다른 space (예: `"reg"`) 도 자유롭게 허용 — `_storage` 가 lazy dict 라
+새 space 가 write 호출과 함께 자동 생성.
+
+#### D6.2. address keying
+
+`addr` 는 정수. **physical address (PA) 또는 virtual address (VA)** 일 수
+있다 — MemoryStore 자체는 address space 의 의미를 모르고 그저 키로 쓴다.
+Phase 1 의 `MemoryWriteMsg` 는 PA + VA 둘 다 write (`_create_tensor` 에서
+PA 로 zero-init, VA base 로도 zero-init), Phase 2 는 op_log 가 captured
+한 address 로 read/write.
+
+`addr` 의 의미는 호출자가 결정한다 — `MemoryStore` 는 lookup 만 제공.
+
+#### D6.3. read/write 의미 — reference store (no copy)
+
+`write(space, addr, data)`: `data` ndarray 의 reference 를 저장. **copy
+하지 않음**. 호출자가 같은 ndarray 를 이후 mutate 하면 stored value 도
+변경된다.
+
+`read(space, addr, shape=None, dtype=None)`: 저장된 ndarray 의 reference
+반환. `shape` 또는 `dtype` 이 제공되면:
+
+- `dtype != stored.dtype`: `arr.view(np_dtype)` 로 reinterpret cast (no
+  copy).
+- `shape != stored.shape`: `nbytes` 가 일치하면 `arr.reshape(shape)` (view).
+- `nbytes` 불일치: `ValueError`.
+
+데이터를 안전하게 분리하려면 호출자가 `arr.copy()` 호출. ADR-0027 의
+race 회피가 op_log snapshot 단계에서 명시적 copy 를 강제하는 이유.
+
+#### D6.4. `has(space, addr) -> bool`
+
+해당 키의 존재 여부만 확인. 데이터 인스턴스화는 안 함.
+
+#### D6.5. `snapshot() -> MemoryStore`
+
+shallow copy. inner dict 의 새 인스턴스를 만들되 ndarray reference 는
+공유. Phase 2 초기화 시점에 Phase 1 의 store 를 fork 하여 Phase 2 의
+mutation 이 Phase 1 의 다른 사용처에 영향을 주지 않게 분리하는 데 사용.
+
+### D7. op_log 가 SimPy 단일-스레드를 가정한다
+
+`OpLogger` 의 `_records`, `_pending` 은 lock 없이 사용. SimPy 가 single-
+threaded 라 `record_start` → `record_end` 사이에 다른 thread 가 끼어들
+수 없다는 가정.
+
+향후 multi-process kernbench (ADR-0047 D6) 가 도입되면 OpLogger 도 process
+별로 분리되어야 함이 명시. 단일 OpLogger 인스턴스가 multiple process 의
+record 를 받지 못한다.
+
+## Alternatives Considered
+
+### A1. op_log 를 SQLite / parquet 같은 외부 store 로
+
+기각 (현재). in-memory list 가 Phase 1 → Phase 2 의 핸드오프 latency 를
+최소화한다. 외부화는 long-running batch run 에서 의미가 있겠으나, 현재
+single-run 워크로드 에서는 overhead 만 추가.
+
+### A2. snapshot 을 record_start 시점에 캡처
+
+기각. record_start 시점은 input 이 아직 채워지지 않은 상황 (예: math
+op 의 input 이 직전 op 의 output 일 때) 이 흔하다. record_end 가 정확한
+시점.
+
+### A3. MemoryStore 를 component-별 store 로 분리
+
+기각. (space, addr) 키가 이미 충분히 disambiguation 을 제공하며, component
+별 분리는 cross-PE IPCQ copy (ADR-0023 D9) 가 source/destination 양쪽
+store 를 접근해야 하는 케이스를 복잡하게 만든다.
+
+### A4. op_log 에 cross-op dependency edge 명시
+
+부분 채택. `dependency_ids` 필드가 OpRecord 에 자리 잡고 있지만 현재
+사용되지 않음 (D1). Phase 2 DataExecutor 가 t_start 정렬 + secondary sort
+(memory ops before math at same t_start) 로 ordering 을 결정하며, 명시적
+dependency graph 가 필요해지면 이 필드가 채워질 자리. 현재는 ordering rule
+이 충분하므로 미사용.
+
+## Consequences
+
+- ADR-0020 의 op_log / MemoryStore 선언이 D1–D6 의 구체 schema 로 확장
+  되어, Phase 2 DataExecutor 작성/수정 시 정확한 필드 의미를 grep 없이
+  ADR 에서 확인 가능.
+- D3 의 op_name 별 params 스키마가 명시되어, 새 op (예: 새 reduction
+  type) 추가 시 `_extract_op_info` 분기 어디에 끼울지 명확.
+- D4 의 snapshot 시점 차이 (math = input snapshot, dma_write = HBM-only
+  snapshot) 가 ADR 에 굳어져, ADR-0027 의 cross-PE race 회피 결정이 향후
+  refactor 에서 silently 깨지지 않음.
+- D6.3 의 reference-store 의미가 명시되어, 호출자가 mutation safety 책임
+  을 인지. ADR-0027 의 explicit `.copy()` 패턴이 정당화됨.
+- D7 의 single-thread 가정이 명시되어, multi-process kernbench (ADR-0047
+  D6 supersession 후보) 도입 시 OpLogger 분리가 필요함이 분명.

docs/adr-ko/ADR-0053-dev-topology-builder-algorithms.md

@@ -0,0 +1,307 @@
+# ADR-0053: Topology Builder + Visualizer Algorithms
+
+## Status
+
+Accepted (2026-05-22).
+
+`topology/builder.py`, `topology/mesh_gen.py`, `topology/visualizer.py` 가
+함께 수행하는 토폴로지 컴파일·시각화 파이프라인의 핵심 알고리즘 선택
+(placement-driven router attachment, mesh auto-layout, source_hash 캐시,
+view projection, SVG rendering) 을 명시한다. ADR-0006 가 topology
+compilation 의 high-level intent (compiled topology, distance extraction,
+automatic diagram generation) 를 정의하나, **builder 가 실제로 어떤
+알고리즘을 사용하는지** 는 코드 grep 으로만 확인 가능했다.
+
+## First action (제일 처음에 하는 일)
+
+`resolve_topology(path_str)` 가 호출되면 다음 4 단계가 순서대로 일어난다:
+
+1. **경로 검증** (`builder.py::resolve_topology`):
+   `Path(path_str).expanduser().resolve()`, 존재 확인, file 여부 확인.
+   실패 시 `FileNotFoundError` 또는 `ValueError`.
+2. **YAML 파싱** (`_read_spec`): `yaml.safe_load`. parse error 면 line/
+   column 정보 포함한 `ValueError`. dict 가 아니면 reject.
+3. **mesh 자동 생성** (`mesh_gen.ensure_mesh_file`): topology yaml 과
+   같은 디렉터리에 `cube_mesh.yaml` 을 만들거나 (캐시 invalid 시) 재사용
+   (캐시 hit 시). 이 단계가 cube NoC 의 라우터 grid 와 부착 정보를 결정.
+4. **graph 컴파일** (`_compile_graph`): system → IO chiplets → cubes →
+   inter-cube edges → IO↔cube edges → system↔IO edges 순으로 nodes/edges
+   를 누적, 그 다음 4 개의 view projection (system, sip, cube, pe) 을
+   생성하여 `TopologyGraph` 로 묶음.
+
+즉, **topology compile 의 첫 일은 "topology.yaml 을 dict 로 읽고, 동일
+디렉터리에 cube_mesh.yaml 을 생성/검증한 뒤, system→sip→cube→pe 순으로
+flat graph + 4-view projection 을 만드는 것"** 이다.
+
+## Context
+
+`topology/` 패키지의 책임:
+
+- **builder.py** (1207 줄): topology.yaml 을 받아 `TopologyGraph` (nodes
+  + edges + 4 view projections) 를 컴파일.
+- **mesh_gen.py** (305 줄): cube NoC 의 라우터 grid 와 PE/UCIe/M_CPU/SRAM
+  부착 위치를 자동 결정하여 `cube_mesh.yaml` 로 캐시.
+- **visualizer.py** (887 줄): `TopologyGraph` 로부터 SVG 다이어그램 4종
+  (system / sip / cube / pe) 을 생성.
+
+ADR-0006 가 "topology compilation 의 결과는 distance metadata 와 diagram
+generation 의 single source" 라는 high-level 결정을 정의하나, 구체 알고리즘
+(예: placement-driven nearest-router attachment, HBM 제외 zone 산출,
+source_hash 의 어떤 필드가 invalidation 을 트리거하는가) 은 ADR 에 없다.
+
+특히 다음 결정들이 ADR-level 에 부재:
+
+- 왜 mesh_gen 이 별도 파일 (`cube_mesh.yaml`) 로 캐시되는가?
+- source_hash 가 어떤 필드를 포함하며, 어떤 변경이 재생성을 강제하는가?
+- placement coordinate 가 cube 좌표가 아닌 mm 단위인 이유?
+- HBM zone 제외와 UCIe N/S/E/W 분배가 mesh 안에서 어떻게 결정되는가?
+- view projection 4 개 (system/sip/cube/pe) 의 추상화 레벨 차이?
+
+이 ADR 이 이 결정들을 한 곳에 정리한다.
+
+## Decision
+
+### D1. compile 파이프라인 — 6 단계
+
+`_compile_graph(spec)`:
+
+1. **시스템 노드 생성** (`_instantiate_system`): `fabric.switch0`, host CPU
+   등 system-level 노드 추가.
+2. **per-SIP loop** (`for sip_id in range(system.sips.count)`):
+   - **IO chiplets** (`_instantiate_io_chiplets`): pcie_ep / io_cpu /
+     io_noc / io_ucie PHY / conn 노드 + 내부 양방향 edge 생성.
+   - **cube instantiation** (`_instantiate_cube`): cube_mesh.yaml 의 router
+     grid 를 토대로 cube-별 라우터, PE sub-components (pe_cpu, pe_dma,
+     pe_fetch_store, pe_gemm, pe_math, pe_mmu, pe_tcm, pe_scheduler,
+     pe_ipcq), m_cpu, sram, hbm_ctrl 인스턴스화 + 내부 edge 깔기.
+   - **inter-cube edges** (`_add_inter_cube_edges`): UCIe N/S/E/W mesh
+     edge.
+   - **IO ↔ cube edges** (`_add_io_to_cube_edges`): io_noc 와 cube 의
+     edge UCIe phy 사이 연결.
+3. **switch ↔ IO edges** (`_add_system_to_io_edges`): `fabric.switch0`
+   와 각 SIP 의 `pcie_ep` 사이 양방향 edge (ADR-0038 D3 + ADR-0010 의
+   cross-SIP IPCQ 경로).
+4. **view projections** 4 종 build:
+   - `_build_system_view(spec)` — Tray 레벨, SIP 들과 system switch.
+   - `_build_sip_view(spec)` — SIP 안의 cube mesh + IO chiplet.
+   - `_build_cube_view(spec)` — 단일 cube 안의 router grid + PE/M_CPU/SRAM/
+     HBM_CTRL 부착.
+   - `_build_pe_view(spec)` — 단일 PE 안의 9 sub-components + 내부 edge.
+5. **TopologyGraph 리턴**: `TopologyGraph(spec, nodes, edges, system_view,
+   sip_view, cube_view, pe_view)`.
+
+이 6 단계는 **순서가 의미를 가진다**: cubes 가 만들어진 후에야 inter-cube
+edges 가 valid 한 src/dst 를 갖고, IO chiplet 이 먼저 만들어져야 IO ↔ cube
+edge 가 그를 참조할 수 있다. 새 노드 종류를 끼울 때는 의존 관계를 보고
+적절한 위치에 삽입해야 한다.
+
+### D2. `cube_mesh.yaml` — 별도 파일 + source_hash 캐시
+
+`mesh_gen.ensure_mesh_file(cube_spec, mesh_path)`:
+
+1. `source_hash = _compute_source_hash(cube_spec)` 산출. 입력 필드:
+   - `geometry` (cube_mm.w/h 등).
+   - `pe_layout` (corners, pe_per_corner).
+   - `ucie.n_connections`.
+   - `memory_map.hbm_mapping_mode`.
+   - `placement` (m_cpu/sram pos_mm).
+2. `mesh_path` (= `topology.yaml` 와 같은 디렉터리의 `cube_mesh.yaml`) 이
+   존재하고 `existing.source_hash == source_hash` 면 재사용 (캐시 hit).
+3. 아니면 `_generate_mesh(cube_spec, source_hash)` 로 새 mesh 생성 후
+   yaml 로 저장.
+
+별도 파일로 캐시하는 이유:
+
+- mesh 생성은 PE/UCIe/router 부착 계산이 들어가 매번 다시 하기 무거움.
+- 같은 cube spec 으로 여러 번 실행 시 동일 mesh 가 보장되어야 함.
+- 사람이 직접 mesh 를 inspect / debug 할 수 있는 artifact 가 됨.
+
+`source_hash` 가 list 한 5 개 필드가 mesh 형상을 결정하는 핵심이며, 그
+외 (예: bandwidth, overhead_ns) 변경은 mesh 재생성을 트리거하지 않는다.
+
+### D3. cube NoC mesh auto-layout 알고리즘
+
+`_generate_mesh(cube_spec)`:
+
+#### D3.1. 행/열 결정
+
+- `pe_positions = _corner_pe_positions(cube_w, cube_h)`: 4 corner (NW/NE/
+  SW/SE) 마다 PE center 좌표 (mm). hardcoded `(1.5, 1.5)` / `(cube_w-1.5,
+  cube_h-1.5)` 패턴 + `pe_per_corner=2` 면 각 corner 에 2 PE 위치.
+- `col_xs = _compute_col_positions(...)`: PE 들의 x 좌표 union + `max_spacing
+  = 3.0 mm` 보다 큰 gap 에 relay 컬럼 삽입.
+- `row_ys, rows_per_half = _compute_row_positions(cube_h, n_connections,
+  pe_positions)`:
+  - `n_conn = max(n_connections, 2)` (hot path minimum).
+  - `rows_per_half = ceil(n_conn / 2)`.
+  - top 절반 + HBM 두 row + bottom 절반. HBM 은 `(cube_h/2 - 1.5, cube_h/2
+    + 1.5)` 에 위치. PE rows 와 HBM rows 사이 `hbm_gap = 1.5 mm`.
+
+#### D3.2. HBM 제외 zone
+
+`hbm_row_start = rows_per_half`, `hbm_row_end = rows_per_half + 1`.
+`hbm_col_start = n_cols // 2 - 1`, `hbm_col_end = n_cols // 2`.
+
+이 (row, col) 사각형 안의 router 슬롯은 `None` 으로 마킹 (라우터 없음).
+실제 HBM 컨트롤러는 별도 `hbm_ctrl.pe{X}` 노드로 ADR-0017 D9 의 per-PE
+파티션 패턴을 따라 부착.
+
+#### D3.3. PE 부착
+
+각 corner 의 PE 들은 다음 row 에 매핑:
+
+- Top half: NW → row 0, NE → row 1 (top_corners 안의 index).
+- Bottom half: SW → row `hbm_row_end + 1`, SE → row `hbm_row_end + 2`.
+
+각 PE 의 x 좌표가 가장 가까운 col 의 router 에 부착 (`min(range(n_cols),
+key=lambda c: abs(col_xs[c] - pe_x))`). 부착 항목은 `pe{pe_idx}.dma`,
+`pe{pe_idx}.cpu`, `pe{pe_idx}.hbm` 세 가지 (router 별 attach list 에 push).
+
+#### D3.4. M_CPU / SRAM 부착 — nearest router by Euclidean distance
+
+`placement.m_cpu.pos_mm` (default `[1.5, 5.5]`) 와 `placement.sram.pos_mm`
+(default `[1.5, 8.5]`) 의 좌표에서 가장 가까운 router 를 Euclidean
+distance 로 찾아 attach list 에 `"m_cpu"` / `"sram"` 추가.
+
+#### D3.5. UCIe N/S/E/W 분배
+
+`ucie_pe_rows = top_pe_rows + bot_pe_rows` (총 `2 * rows_per_half` 개).
+
+- UCIe-E: 매 PE row 마다 rightmost col 의 router 에 `ucie_e.c{i}`.
+- UCIe-W: leftmost col 의 router 에 `ucie_w.c{i}` (E 의 mirror).
+- UCIe-N/S: PE column 들 중 절반을 좌측, 절반을 우측으로 나눠 top row /
+  bottom row 의 해당 col 에 부착.
+
+각 UCIe connection 은 `c{i}` index 가 붙어 ucie_n_connections 만큼의 PHY
+가 분산된다 (ADR-0017 D5+).
+
+### D4. node 명명 규칙 — 단일 소유자
+
+builder.py 는 다음 명명 규칙으로 노드를 만든다 (ADR-0051 D5 의 단일
+소유자 원칙):
+
+- `fabric.switch0` — system-level switch.
+- `sip{S}.{io_id}.{pcie_ep|io_cpu|io_noc|io_ucie.{dir}|conn.{id}}` — IO
+  chiplet.
+- `sip{S}.cube{C}.{m_cpu|sram|hbm_ctrl.pe{X}|noc.r{R}c{C}|...}` — cube 내부.
+- `sip{S}.cube{C}.pe{P}.{pe_cpu|pe_dma|pe_fetch_store|pe_gemm|pe_math|pe_mmu|pe_tcm|pe_scheduler|pe_ipcq}` — PE sub-components.
+
+이 명명 규칙을 변경하려면 builder.py 와 router.py (ADR-0051) 의 helper
+양쪽이 함께 갱신되어야 한다. 컴포넌트는 명명 규칙을 직접 알지 못하고
+helper 만 호출한다.
+
+### D5. edge `kind` 분류
+
+각 edge 가 부여받는 `kind` 가 라우팅 정책 (ADR-0051 D2) 의 입력. 주요
+kind 값:
+
+- `"pe_internal"` — PE 내부 sub-component 간.
+- `"pe_to_router"` — PE_DMA ↔ cube NoC router.
+- `"router_mesh"` — cube NoC router 간.
+- `"router_to_hbm"`, `"router_to_mcpu"`, `"router_to_sram"`,
+  `"sram_to_router"` 등 — cube-attached component 간.
+- `"ucie_internal"`, `"ucie_conn_to_router"`, `"router_to_ucie_conn"`,
+  `"ucie_conn_to_noc"`, `"noc_to_ucie_conn"`, `"ucie_mesh"` — UCIe 관련.
+- `"io_internal"` — IO chiplet 내부.
+- `"io_to_cube"`, `"cube_to_io"` — IO ↔ cube 경계.
+- `"pcie"` — switch ↔ pcie_ep.
+- `"command"` — control-plane only edges (M_CPU ↔ NOC 등; PE DMA path 에서
+  제외).
+
+새 edge kind 를 추가하면 router.py 의 4 adjacency graph (ADR-0051 D2) 의
+어느 카테고리에 속할지 결정해야 한다 — 그렇지 않으면 default 로 `_adj_all`
+에만 포함되어 의도와 다른 routing 발생 가능.
+
+### D6. view projection — 4 추상화 레벨
+
+`TopologyGraph` 는 flat (nodes + edges) 외에 4 개의 view projection 을
+보유:
+
+- **system_view** (`_build_system_view`): Tray 레벨. SIP 박스들 + `fabric.
+  switch0`. PCIE 링크 표시. 외부 발표용 high-level overview.
+- **sip_view** (`_build_sip_view`): 한 SIP 안. cube mesh + IO chiplet
+  (pcie_ep + io_cpu + io_noc). UCIe N/S/E/W 가 cube 간 연결로 보임.
+- **cube_view** (`_build_cube_view`): 한 cube 안. router grid + PE/M_CPU/
+  SRAM/HBM_CTRL 부착 + UCIe PHY edge 부분. cube 내부 라우팅 / placement
+  진단용.
+- **pe_view** (`_build_pe_view`): 한 PE 안. 9 sub-components + 내부 edge
+  (pe_internal kind). 자세한 PE 내부 dataflow 검토용.
+
+view 는 spec 에서 `visualization.emit_views: [system, sip, cube]` 같이
+선택적으로 출력 (ADR-0006). pe view 는 기본 출력에서 빠져 있으나 코드는
+유지 (자세한 디버그용).
+
+### D7. visualizer.py — SVG 다이어그램 출력
+
+`emit_diagrams(graph, out_dir)` 가 모든 view 를 SVG 로 렌더. 핵심 함수:
+
+- `_render_view_svg(view)` — 일반적인 view 렌더 (router grid 가 없는
+  경우).
+- `_render_cube_view_svg(view, spec)` — cube view 전용 (HBM block 그리기,
+  router grid layout, PE/M_CPU/SRAM/HBM positioning).
+- `_draw_node`, `_draw_edge` — 노드 / edge 의 시각적 표현.
+- `_pick_scale`, `_compute_node_sizes` — 자동 스케일링.
+
+visualizer 는 **derived artifact** (ADR-0006) 로 분류되며, 코드 변경 시
+production check 대상이 아니다. CLAUDE.md 의 "Derived Artifacts" 항목과
+정합.
+
+### D8. spec 변경의 영향 범위
+
+| spec 필드                             | 영향              | mesh 재생성 |
+|---------------------------------------|-------------------|-------------|
+| `system.sips.count`                   | SIP 갯수, node 수 | No          |
+| `sip.cube_mesh.w/h`                   | cube mesh 형상    | No          |
+| `cube.geometry.cube_mm.w/h`           | cube 크기 (mm)    | **Yes**     |
+| `cube.pe_layout.corners/pe_per_corner`| PE 부착 위치       | **Yes**     |
+| `cube.ucie.n_connections`             | UCIe PHY 분배     | **Yes**     |
+| `cube.memory_map.hbm_mapping_mode`    | HBM 분배 모드     | **Yes**     |
+| `cube.placement`                      | M_CPU/SRAM 위치   | **Yes**     |
+| `cube.memory_map.*` (위 제외)         | HBM 용량 / BW     | No          |
+| `*.links.*.bw_gbs`                    | edge bandwidth    | No          |
+| `*.attrs.overhead_ns`                 | 컴포넌트 latency  | No          |
+
+위 표가 D2 의 `_compute_source_hash` 입력과 일치. mesh 재생성이 필요한
+변경은 `cube_mesh.yaml` 의 source_hash 가 자동 invalidate.
+
+## Alternatives Considered
+
+### A1. mesh 를 별도 캐시 파일 없이 매 compile 시 재생성
+
+기각. 같은 spec 으로 여러 번 호출되는 케이스 (CLI run, probe, test) 마다
+mesh 생성 비용을 다시 지불. 또한 사람이 mesh 를 inspect 할 수 있는 artifact
+가 사라짐.
+
+### A2. mesh 생성을 builder.py 에 합치기
+
+기각 (현재). 305 줄 짜리 자체 알고리즘이며, mesh layout 의 결정 (placement-
+driven router attachment, HBM exclusion zone) 이 builder 의 일반적인
+node/edge 생성 책임과 다르다. 분리 유지가 단일 책임 원칙에 더 부합.
+
+### A3. placement coordinate 를 cube 좌표 (col/row) 로 표현
+
+기각. mm 단위 좌표가 시각화 측 (visualizer) 과 mesh layout 측 (nearest-
+router 산출) 양쪽에서 일관되게 쓰인다. cube 좌표는 router grid 가 결정
+되기 전까지는 정의되지 않으므로 placement 입력에 부적절.
+
+### A4. view projection 을 lazy 하게 생성
+
+기각 (현재). 4 개 view 의 생성 비용이 작고 (보통 < 100 ms), eager 생성이
+`TopologyGraph` 를 통한 single source of truth 를 보장.
+
+### A5. visualizer 출력 형식을 SVG 외 (PNG/PDF) 도
+
+기각. SVG 가 vector + 텍스트 검색 가능 + 브라우저 직접 렌더가 가능한 가장
+유연한 형식. PNG 변환이 필요하면 별도 도구 (rsvg-convert 등) 로 후처리.
+
+## Consequences
+
+- ADR-0006 의 high-level intent 가 D1–D7 로 구체화되어, topology 변경
+  영향을 D8 표로 빠르게 가늠 가능.
+- D3 의 mesh auto-layout 알고리즘이 ADR-level 에서 굳어져, 추후 새 PE
+  부착 패턴 (예: HBM 의 6-zone 분할) 도입 시 어느 단계가 영향받는지 명확.
+- D5 의 edge kind 목록과 D7 의 view 구조가 명시되어, 새 component 종류
+  추가 시 (builder + router + visualizer) 어디까지 손대야 하는지 PR
+  reviewer 가 한눈에 파악 가능.
+- D2 의 source_hash invalidation 규칙이 명시되어, cube_mesh.yaml 이 stale
+  하게 남는 경우 (예: bw 값만 바꿨을 때) 가 정상 동작임이 분명.

docs/adr-ko/ADR-0054-eval-milestone-benches.md

@@ -0,0 +1,138 @@
+# ADR-0054: 마일스톤 평가 bench — 자기완결적 sweep + figure bench
+
+## Status
+
+Accepted (2026-05-22).
+
+ADR-0044(D1/D2)와 ADR-0045(D5)를 개정하고, ADR-0043/0044의 "로직이
+`scripts/` + `tests/`에 산다" 배치를 대체한다: GEMM/allreduce 평가
+하니스가 이제 사용자가 실행하여 모든 결과 + figure를 재생성하는
+자기완결적 **bench**가 된다.
+
+## Context
+
+ADR-0043(allreduce 평가)과 ADR-0044(GEMM 평가)는 각 하니스를 **sweep**
+(수동 `scripts/` 드라이버, 또는 allreduce의 경우 parametrized 테스트
+자체) + committed 데이터를 렌더링하는 **figure 테스트**로 분리했다.
+따라서 sweep/render 로직은 `scripts/gemm_sweep.py`,
+`tests/gemm/_gemm_plot_helpers.py`, `tests/sccl/_allreduce_helpers.py`에
+존재했다.
+
+마일스톤 요구사항("사용자가 *하나의 bench*를 실행해 모든 결과와 플롯을
+생성하도록 allreduce + GEMM 평가를 리팩터")은 그 배치로는 충족 불가다:
+bench는 production 코드이며 **`tests/`를 import할 수 없다**(ADR-0007 레이어
+방향). 평가 로직은 bench에서 닿을 수 있도록 production으로 이동해야 했다.
+
+선택한 home은 별도 `kernbench.eval` 패키지가 아니라 bench 모듈 자체다.
+bench 파일은 임의의 모듈 레벨 코드를 가질 수 있으며, 하니스를 bench로
+합치면 도메인당 파일 하나가 유지되고 패키지 레이어가 하나 줄어든다.
+
+## Decision
+
+### D1. 두 마일스톤 bench가 평가 로직을 보유
+
+- `src/kernbench/benches/milestone_1h_gemm.py` — GEMM shape×variant sweep
+  + 세 figure renderer(`scripts/gemm_sweep.py` +
+  `tests/gemm/_gemm_plot_helpers.py`에서 이동).
+- `src/kernbench/benches/milestone_1h_ccl.py` — distributed allreduce
+  드라이버, latency + buffer-kind sweep, topology diagram, FSIM 비교, 그리고
+  direct-launch 패리티 레퍼런스(`tests/sccl/_allreduce_helpers.py`에서 이동).
+
+각 파일은 해당 도메인 평가 로직의 **단일 home**이다.
+
+### D2. "평가 bench" 패턴 (ADR-0045 D5 확장)
+
+ADR-0045 D5는 bench를 단일 구성(single-SIP, 또는 ADR-0024 multi-SIP CCL
+예외)으로 고정했다. 본 ADR은 세 번째 패턴을 추가한다:
+
+- **평가 bench**는 *여러* 구성을 구동하고 figure를 렌더링할 수 있다. 외부
+  `run_bench` 엔진 대신 sweep 지점마다 자체 `GraphEngine` /
+  `RuntimeContext`를 빌드한다.
+- 그러면 외부 ctx에 제출된 handle이 없으므로, bench는 마지막에
+  **sentinel 텐서**(`torch.zeros((1, 1), …)`)를 제출하여 `run_bench`의
+  "최소 한 번 제출" 계약(ADR-0045 D4)을 만족시키고 CLI가 0으로 종료되게
+  한다.
+
+### D3. 출력 위치
+
+두 bench 모두 `src/kernbench/benches/1H_milestone_output/{gemm,ccl}/`에
+쓴다(사용자 요청 — bench 옆 아티팩트). 디렉터리는 생성된 PNG/CSV/JSON만
+보유하며(`.py`/`__init__.py` 없음), 따라서 eager-import audit(ADR-0045
+첫 동작)이 무시한다 — `pkgutil.iter_modules`는 비-패키지 하위 디렉터리를
+yield하지 않는다. `docs/diagrams/` 아티팩트처럼 **커밋된다**(원격에서
+figure를 볼 수 있도록); bench 재실행 시 제자리에서 재생성된다.
+
+### D4. GEMM 무거운 sweep — 기본은 fresh, `MILESTONE_FAST`로 재사용
+
+`milestone-1h-gemm`은 기본적으로 전체 24-sim sweep을 실행한다(분 단위;
+한 shape는 2048 tile). `MILESTONE_FAST=1`은 committed
+`docs/diagrams/gemm_sweep.json`을 재사용하고 렌더링만 한다(초 단위). 이는
+ADR-0044 D1/D2의 "무거운 sweep은 수동/`slow` 단계로 유지"를 뒤집는다:
+bench 실행이 곧 재생성이다. slow 경로는 `@pytest.mark.slow` bench
+테스트로 행사되고, fast 경로는 기본 실행된다.
+
+### D5. 테스트 + 스크립트는 thin re-export shim으로 재사용 (단일 home 유지)
+
+기존 figure 테스트와 `scripts/gemm_sweep.py` 진입점은 유지되며 이제 bench
+모듈을 재사용한다:
+
+- `tests/gemm/_gemm_plot_helpers.py` → renderer +
+  `GEMM_SWEEP_JSON`/`GEMM_PLOTS_DIR`/`ROOT`를
+  `kernbench.benches.milestone_1h_gemm`에서 re-export.
+- `tests/sccl/_allreduce_helpers.py` → 드라이버 코어, config writer, sweep
+  상수, renderer, disk aggregator를 `kernbench.benches.milestone_1h_ccl`에서
+  re-export하고, **pytest 전용** 조각은 로컬 유지: `pytest.param` 행렬
+  (`CONFIGS` / `_sweep_params` / `_bk_params`)과 fixture 결합
+  `_run_distributed`(`monkeypatch.chdir` + `_drive_distributed`) wrapper.
+- `scripts/gemm_sweep.py` → bench의 `run_sweep` 위 thin wrapper.
+
+테스트가 bench 모듈을 import하는 것은 허용된다(테스트는 production 위에
+위치, ADR-0007); 이는 전체 패키지 eager audit을 유발하며, 그것은 이미 매
+`kernbench` 실행 시 동작한다. matplotlib는 renderer 내부에서 lazy import로
+유지되어 audit의 startup 비용은 불변이다.
+
+### D6. 평면 모듈 네이밍 (`benches/` 하위 폴더 없음)
+
+`1H_milestone…`로 명명된 `benches/` 하위 패키지는 불가능하다 — Python
+패키지 이름은 숫자로 시작할 수 없다. 따라서 bench는 평면 모듈
+`milestone_1h_gemm.py` / `milestone_1h_ccl.py`이며 bench 이름은
+`milestone-1h-gemm` / `milestone-1h-ccl`(kebab-case, ADR-0045 D1에 따라
+글자로 시작)이다.
+
+## Consequences
+
+### Positive
+
+- `kernbench run --bench milestone-1h-gemm`(또는 `…-ccl`)이 도메인의 모든
+  결과 + figure를 한 명령으로 재생성한다 — 마일스톤 요구사항.
+- 평가 로직의 단일 소스(bench), shim을 통해 테스트와 스크립트가 재사용;
+  중복 없음.
+- figure 테스트와 `scripts/gemm_sweep.py`는 변경 없이 계속 동작.
+
+### Negative / limitations
+
+- 두 bench 파일이 크다(CCL 쪽은 distributed 드라이버, sweep, matplotlib
+  드로잉을 섞는다). 대부분 평가 하니스인 "bench"는 이례적이며, 본 ADR이
+  이를 정당화한다.
+- 생성 아티팩트가 명시적 요청에 의해 source tree(`src/kernbench/benches/`)
+  안에 살며 커밋된다(원격에서 figure를 볼 수 있도록); bench 재실행 시
+  재생성된다.
+- `milestone-1h-ccl`(및 기본 `milestone-1h-gemm`)은 분 단위 소요 —
+  on-demand 마일스톤 아티팩트에는 수용 가능, 일상 실행에는 아님.
+
+## Dependencies
+
+- **ADR-0007**: 레이어 방향(테스트는 production을 import할 수 있으나 bench는
+  테스트를 import할 수 없는 이유).
+- **ADR-0043 / ADR-0044**: 본 ADR이 bench로 이전하는 allreduce / GEMM 평가
+  하니스.
+- **ADR-0045**: bench 모듈 계약; 여기 D2가 그 D5(single-device 규칙)를
+  평가-bench 패턴으로 확장하고, sentinel을 위해 D4(NO_REQUESTS)에 의존.
+- **ADR-0024**: allreduce sweep이 구동하는 rank = SIP launcher.
+
+## Open questions
+
+- GEMM theoretical 모델 상수(ADR-0044 D5)를 복사 대신 ADR-0033/0014에서
+  소싱해야 하는가? 본 ADR로는 불변.
+- `build_overview_slides.py`가 GEMM 막대를 네이티브로 그리는 대신 마일스톤
+  출력 PNG를 소비해야 하는가? 여전히 open(ADR-0044 D6 / Negative).

docs/adr-ko/ADR-0062-prog-tl-async-load.md

@@ -0,0 +1,148 @@
+# ADR-0062: Lazy `tl.load` — 첫 사용 시점 auto-wait를 갖는 non-blocking HBM load
+
+## Status
+
+Accepted
+
+> **ADR-0060**(AHBM GQA Fused Attention)의 보조 ADR. Decode와 long-context
+> 어텐션은 **KV-load-bound**이며, load/compute 오버랩이 지배적 레버다. 본
+> ADR은 별도의 `load_async` op을 추가하는 대신 `tl.load` 자체를 **lazy**
+> (non-blocking, wait를 첫 사용 시점에 자동 삽입)로 만든다. 오늘날 유일한
+> async primitive는 IPCQ 통신용이지 HBM load용이 아니다.
+
+## Context
+
+### 오버랩이 요구하는 것
+
+FlashAttention은 operand를 스트리밍한다: GEMM/MATH 엔진이 현재 타일을 처리하는
+동안 DMA 엔진은 이미 다음 operand를 당기고 있어야 한다. 그 오버랩이 있으면
+대역폭-bound 커널은 타일당 `compute + dma` 대신 대략 `max(compute, dma)`로
+돈다.
+
+ADR-0060의 하이브리드 설계에서 두 GEMM은 `tl.composite(op="gemm")`로 발행되고,
+그 K/V operand는 `tl.ref`(HBM 상주, PE_SCHEDULER가 타일당 스트리밍)이므로 **타일당
+K/V prefetch는 composite scheduler가 처리**한다. lazy `tl.load`는 나머지 명시적
+load(Q group, 그리고 비-composite 커널)를 담당하여 그것들도 greenlet을 멈추는
+대신 뒤따르는 compute와 오버랩되게 한다.
+
+### 현재 존재하는 것
+
+- `tl.load(ptr, shape, dtype)`은 **blocking**이다: `DmaReadCmd`를 emit하고
+  greenlet 커널은 PE_DMA가 완료를 알릴 때까지 suspend된다
+  (`tl_context.py:177-203`; greenlet 구동 `kernel_runner.py:146-153`). 한
+  커널에서 두 `tl.load`가 in-flight일 수 없다.
+- async 패턴이 **존재**하긴 하나 IPCQ 전용이다:
+  `tl.recv_async(dir, ...) -> RecvFuture` + 지연 wait
+  (`kernel_runner.py:248-285`). 이는 메커니즘을 입증한다 — future를 반환하고
+  나중에 wait 체크(`if not future.event.triggered: yield future.event`)로
+  resolve되는 non-blocking 커맨드가 greenlet 모델에서 동작한다.
+- DMA 엔진은 read 채널을 write 채널과 분리된 SimPy resource(capacity 1,
+  `pe_dma.py:45`)로 모델한다. 따라서 in-flight read는 표현 가능하며, 단일 read
+  채널에서 직렬화되면서 PE_GEMM/PE_MATH의 compute와 오버랩된다. 오직
+  *커널-대면 API*만이 현재 load-vs-compute를 직렬화한다.
+
+`tl.load`는 오늘날 뒤따르는 compute와 오버랩될 수 없다.
+
+## Decision
+
+**`tl.load`를 lazy로 만든다: `DmaReadCmd`를 발행하고 핸들을 즉시 반환
+(non-blocking); 런타임은 load된 데이터가 실제로 처음 소비되는 지점에 wait를
+자동 삽입한다.** 커널-대면 API는 불변이다 — 작성자는 계속 `tl.load`를 쓴다 —
+따라서 이는 새 op이 아니라 *시맨틱* 변경이다. 기존 `recv_async`/wait 기계장치를
+(1) HBM-load 경로로, (2) 명시적 `tl.wait` 호출에서 암묵적·의존성 기반 first-use
+wait로 일반화한다.
+
+### D1. `tl` 표면 — 불변
+
+`tl.load(ptr, shape, dtype)`은 시그니처와 `TensorHandle` 반환을 유지한다.
+바뀌는 것은 *언제* blocking하느냐다: 발행 시점에는 결코 아니고, 결과가 소비 op에
+의해 처음 읽힐 때만 암묵적으로.
+
+### D2. 메커니즘 — non-blocking 발행 + 사용 시점 auto-wait
+
+- `tl.load`는 완료 event를 yield하지 않고 `DmaReadCmd`를 PE_DMA에 post하며,
+  반환 핸들에 pending event를 기록한다(`recv_async` 패턴을 load에 적용).
+- 소비 op(`tl.dot`, MATH op, `tl.store`, `tl.composite` operand, …)이
+  dispatch될 때, 런타임은 각 입력 핸들의 pending load event를 확인하고 아직
+  triggered가 아니면 먼저 yield한다 — 즉 wait가 가장 늦은 올바른 지점(첫 사용)에
+  자동 삽입된다.
+- op_log 엔트리는 불변(`memory/dma_read`): asynchrony는 스케줄링 속성이지 새 op
+  종류가 아니므로 `dma_read_count` 및 기존 op_log 소비자가 계속 동작한다.
+
+### D3. 범위 — 전역
+
+`tl.load`는 opt-in 플래그 뒤가 아니라 **모든 곳에서** lazy다. 이것이 충실한
+모델이다: 매 load마다 blocking하는 것은 *naive* 커널의 속성이고, 효율 커널(그리고
+실제 컴파일러)은 load를 hoist하고 사용 시점에만 wait한다. 결과: `tl.load`와 첫
+사용 사이에 독립 작업이 있는 기존 커널은 **더 낮은(빠른) latency**를 본다 —
+동작의 회귀가 아니라 모델의 정확도 개선. 바뀌는 골든 latency는 **재생성**해야
+하며; load 직후 곧바로 사용하는 커널은 변화가 없다(auto-wait이 즉시 발동하여
+blocking과 동일).
+
+### D4. Latency / 오버랩 시맨틱
+
+- `tl.load`는 **발행**(descriptor push)만 청구하고 커널은 진행한다. (per-op
+  발행 비용은 `dispatch_cycles`, 현재 0; op 종류별 차등 발행 비용은 별도로
+  추적 — ADR-0060 §1, §9, ADR-0064.)
+- DMA 전송은 read 채널(capacity 1)을 모델 지속시간만큼 점유하며 커널이 다음에
+  내는 compute와 병렬; 다중 in-flight load는 채널에서 직렬화되나 compute와
+  오버랩.
+- 자동 삽입된 wait는 전송이 끝나지 않았을 때만 blocking.
+- 결정성 보존: wait 지점은 프로그램 순서(첫 사용)로 고정되고, 완료는 modeled DMA
+  채널의 scheduled event다(SPEC §0.1, R8). latency 빼기 없음 — 오버랩은 실제
+  모델 동시성.
+
+> **모델링 가정.** auto-wait-at-first-use는 *잘 스케줄된* 커널(컴파일러가
+> wait를 가장 늦은 올바른 지점에 둠)을 모델한다. 실제 컴파일러는 이를 근사하며,
+> 일부 load는 hoist 불가(register pressure, aliasing). 성능 시뮬레이터(SPEC
+> §0)에서 잘 스케줄된 경우를 모델하는 것이 의도된 동작이다.
+
+## Alternatives
+
+### A1. 별도 `tl.load_async` op (이전 제안)
+
+커널이 손으로 호출하는 명시적 `load_async`/`wait` 쌍 추가(더블버퍼 댄스). 기각:
+커널-대면 API를 키우고 버퍼 수명 부기를 모든 커널 작성자에게 떠넘긴다. lazy
+`tl.load`는 **API 표면 변경 없이** 동일 오버랩과 컴파일러식 auto-wait를 준다.
+
+### A2. 커널별 opt-in lazy load
+
+`tl.load`를 기본 blocking으로 두고 새 커널만 opt-in. 기각: `tl.load` 시맨틱을 두
+변종으로 쪼개고 기존 bench로부터 모델 개선을 숨긴다; 전역 lazy가 더 깔끔하고
+골든 재생성 비용은 1회성(D3).
+
+### A3. composite 스트리밍에만 의존 (lazy load 없이)
+
+ADR-0060의 composite는 `tl.ref` K/V operand를 스트리밍하므로 GEMM operand DMA는
+이미 오버랩된다. 그러나 명시적 load(Q group, 비-composite 커널)는 lazy `tl.load`
+없이 여전히 stall. 단독으로는 불충분.
+
+## Consequences
+
+### Positive
+- 커널-대면 API 변경 **제로**로 load/compute 오버랩; 작성자는 계속 `tl.load`.
+- `recv_async`/wait와 대칭 — 개념적 표면적 작음.
+- 일반적: 어떤 대역폭-bound 커널도 자동 prefetch.
+
+### Negative
+- load와 사용 사이에 독립 작업이 있는 커널의 기존 골든 latency가 이동(빨라짐)
+  → 1회성 재생성(D3).
+- auto-wait는 런타임이 핸들별 pending event를 추적하고 소비-op dispatch에서
+  확인하도록 요구(데이터 의존성 추적).
+- scratch 수명(ADR-0063)과 상호작용: in-flight load의 대상 버퍼는 auto-wait이
+  발동하기 전에 recycle되면 안 된다. recycling scope는 살아있는(미-wait된) load
+  버퍼를 제외해야 한다.
+
+## Test Requirements
+
+1. **오버랩이 실재**: `tl.load` 후 비슷한 지속시간의 독립 GEMM을 내는 커널이
+   `load+gemm`이 아니라 ≈`max(load, gemm)`으로 완료(end-to-end latency가 직렬
+   합보다 엄격히 작음).
+2. **auto-wait 정확성**: load된 `TensorHandle`이 처음 소비될 때, 같은 주소에
+   대해 오늘날 blocking `tl.load`와 동일한 바이트를 보유(Phase 2).
+3. **둘 in-flight**: 서로 다른 주소로의 두 `tl.load`가 나중에 소비되어 둘 다
+   올바른 독립 텐서로 resolve; 그 DMA는 read 채널에서 직렬화.
+4. **op_log 호환성**: 각 `tl.load`은 여전히 정확히 하나의 `memory/dma_read`를
+   로그; `dma_read_count`는 blocking 버전 대비 불변.
+5. **무-오버랩 무-변화**: load 직후 곧바로 사용하는 커널은 blocking 모델과 동일한
+   latency(auto-wait 즉시 발동).

docs/adr-ko/ADR-0063-prog-tl-scratch-scope.md

@@ -0,0 +1,222 @@
+# ADR-0063: `tl.scratch_scope` — long-context를 위한 tile 단위 scratch 재활용
+
+## Status
+
+Accepted
+
+> **ADR-0060**(AHBM GQA Fused Attention)의 보조 ADR. Long-context attention
+> 은 수많은 K/V tile을 sweep하며, 매 tile의 중간 결과
+> (`scores`, `P`, `exp`, partial `O`, …) 는 현재 kernel 호출 동안 회수되지
+> 않는 fresh scratch를 새로 할당한다. PE당 1 MiB scratch 예산은 실제로
+> 의미 있는 context length에 도달하기 훨씬 전에 소진된다.
+
+## Context
+
+### Bump allocator
+
+`TLContext` 는 모든 math/compute output handle을 PE-local scratch pool
+에서 **linear bump cursor** 로 할당한다
+(`src/kernbench/triton_emu/tl_context.py`; `_scratch_alloc`):
+
+```python
+def _scratch_alloc(self, nbytes):
+    aligned = (nbytes + 15) & ~15
+    addr = self._scratch_base + self._scratch_cursor
+    self._scratch_cursor += aligned
+    if self._scratch_cursor > self._scratch_size:      # default 1 << 20 = 1 MiB
+        raise RuntimeError("TLContext scratch overflow: ...")
+    return addr
+```
+
+Docstring은 cursor가 **"매 kernel 호출마다 reset된다"** 고 명시한다 —
+즉 kernel entry에서만 reset, kernel body 내에서는 절대 reset되지 않는다.
+모든 `tl.dot`, `tl.softmax`, `tl.exp`, `tl.sum`, `a - b`, `a * b`, … 는
+fresh slice를 가져가고 kernel이 return될 때까지 아무것도 회수되지 않는다.
+
+### 왜 GQA kernel에서 문제가 되는가
+
+현재 milestone bench는 이 한계 **때문에** `S_q = S_kv_per_rank = 16` 으로
+명시적으로 고정해 두었다
+(`tests/attention/test_milestone_gqa_llama70b.py:123-148`):
+
+> "S_q_prefill and S_kv_per_rank are deliberately small (16 each) so the
+> simulator's 1 MB per-PE TCM kernel scratch is not exhausted by the
+> bump-allocated handle outputs of softmax/exp/dot/sum chains over
+> n_ranks ring steps."
+
+`n_tiles`개의 tile sweep에 대해 FlashAttention은 O(`n_tiles` ×
+tile당 중간값) 만큼 할당한다. 의미 있는 context에서는 overflow된다.
+Flash attention의 *수학* 자체는 **O(1)** 의 live scratch만 필요로 한다 —
+running `(m, l, O)` + 현재 tile의 working set — 매 tile의 임시값은
+running state에 fold된 직후 dead가 되기 때문이다. 다만 allocator가 그
+사실을 모를 뿐이다.
+
+## Decision
+
+Kernel이 할당 영역을 **재활용 가능(reclaimable)** 으로 표시할 수 있는
+**scratch scope** 을 추가한다. 이를 통해 tile당 임시값은 재활용되고,
+live running state(및 in-flight prefetch buffer)는 보존된다.
+
+### D1. `tl` 표면 — context manager
+
+```python
+with tl.scratch_scope():
+    s   = tl.dot(q, k_t)        # `with` 내부의 모든 handle은
+    p   = tl.softmax(s)         # 종료 시 rewind되는 영역을 공유
+    o_j = tl.dot(p, v)
+    # ... o_j를 외부 영역에 사는 running (m,l,O) 에 fold ...
+# __exit__: cursor가 __enter__ 시점 값으로 rewind
+```
+
+Semantics:
+- `__enter__` 는 현재 `_scratch_cursor` 를 save-point로 기록.
+- `__exit__` 는 cursor를 save-point로 복원, 내부에서 할당된 모든 것을
+  free.
+- Scope **외부**에서 할당된 handle (running `m,l,O`, ADR-0062의
+  `LoadFuture` 가 보유한 prefetch buffer) 은 주소를 유지 — save-point
+  이전에 또는 enclosing scope에서 할당되었기 때문.
+
+### D2. 안전 계약
+
+Scope 내부에서 할당된 handle은 scope exit 이후 **읽으면 안 된다** —
+그 바이트는 다음 scope의 할당으로 덮어쓰여질 수 있다. Flash loop은
+이를 자연스럽게 지킨다: tile iteration 사이에 살아남는 값은 running
+accumulator 뿐이고, 이들은 per-tile scope **외부**에서 할당되며
+*내부*의 op들이 외부 주소로 쓰면서 갱신된다(merge가 새 running state
+를 쓰는 방식 — D3 참조).
+
+### D3. Running accumulator와의 상호작용
+
+Online-softmax merge는 이전 running `(m, l, O)` 와 현재 tile의
+`(m_j, l_j, O_j)` 를 읽어 새 running 값을 만든다. 새 running 값을
+재활용 영역 바깥에 두기 위해, merge는 결과를 loop 시작 전 1회 할당된
+**stable scratch** (per-tile scope과는 별개의 작은 고정 "running-state"
+arena) 로 쓴다. 구체적으로 kernel은 두 개의 arena를 둔다:
+
+- **persistent arena** (1회 할당): `m, l, O` (merge용 double-buffer가
+  필요한 경우 그것까지).
+- **scoped arena** (매 tile rewind): `scores, P, exp, O_j, scale_*`.
+
+이는 실제 flash-attention SRAM budgeting의 모습과 같다: 작은 persistent
+accumulator 영역 + 재활용되는 tile working set.
+
+#### D3.1 Persistent-arena 쓰기 mechanism: `tl.copy_to(dst, src)`
+
+Merge op (`tl.maximum`, `tl.exp`, binary `*` / `+`) 모두 D1의 bump
+cursor로부터 할당하는 `_make_compute_out(...)` 을 호출한다. 따라서
+`scratch_scope` 내부에서 그 결과 handle들은 scope **내부**에 살고
+`__exit__` 시점에 사라진다. D3의 two-arena 분리를 실현하려면 kernel은
+**scoped 결과의 바이트를 persistent 주소로 쓰는** 방법이 필요하다.
+
+이 gap을 메우는 primitive:
+
+```python
+def copy_to(self, dst: TensorHandle, src: TensorHandle) -> None:
+    """``src`` 의 바이트를 ``dst`` 의 주소로 복사한다 (양쪽 모두 TCM).
+
+    Shape, dtype 일치 필요. ``dst`` 는 통상 어떤 활성 ``scratch_scope`` 도
+    바깥에서 할당된 handle (persistent arena); ``src`` 는 scope ``__exit__``
+    이후에도 바이트가 살아 있어야 하는 scoped handle.
+    """
+```
+
+`tl.store` (HBM 측 바이트 복사) 와 대칭이지만, running-state writeback이
+op_log를 불필요한 DMA로 오염시키지 않도록 TCM 전용으로 유지한다.
+
+**Mechanics:**
+- 신규 `CopyCmd(src, dst, nbytes, data_op=True)` command.
+- op_log: `op_kind="math"`, `op_name="copy"` — vector engine에서 실행.
+- Latency: `pe_math._compute_ns(prod(shape))` — HBM 전송이 아니라
+  on-chip register writeback을 모델링.
+- Emit-time 검증: `dst.shape == src.shape`, `dst.dtype == src.dtype`,
+  `dst.space == "tcm"`, `src.space == "tcm"`. 작성자 오류는 Phase 2
+  data execution 깊숙한 곳이 아니라 Phase 1에서 노출.
+
+**호출 패턴:**
+
+```python
+m, l, O = init_running(...)                  # persistent (scope 외부)
+for j in range(n_tiles):
+    with tl.scratch_scope():
+        ...                                  # tile 작업 (재활용)
+        m_new = tl.maximum(m, mj)            # scoped scratch
+        l_new = l * scale_old + l_step * scale_step
+        O_new = O * scale_old + O_step * scale_step
+
+        tl.copy_to(m, m_new)                 # ← 새 running state를 persist
+        tl.copy_to(l, l_new)
+        tl.copy_to(O, O_new)
+    # exit: scoped m_new/l_new/O_new 는 사라지고, 그 바이트는
+    #       persistent m/l/O 에 살아있다
+```
+
+`copy_to` 는 `__exit__` **이전** 에 실행되므로 `src` (scoped) 의 read는
+유효하며, exit 이후에는 `dst` (persistent) 만 read되어 D2의
+"exit 이후에는 읽지 말 것" 안전 계약을 만족한다.
+
+**왜 모든 math op에 `dst=` kwarg를 추가하지 않고 전용 primitive 인가**
+(검토 후 기각): `tl.maximum`, `tl.exp`, `_binary_math`, `_unary_math`,
+`_reduction` 에 `dst=` 를 추가하는 것은 5개 op-family 에 걸쳐 ~25 LOC
+이며, "호출은 fresh handle을 return한다" 라는 일관된 패턴을 깬다.
+`tl.copy_to` 는 단일 primitive, 단일 command, 단일 executor handler —
+같은 효과에 비해 최소한의 surface area.
+
+### D4. Nesting
+
+Scope는 nest 가능 (save-point의 stack). Inner scope exit은 inner
+save-point로 rewind, outer exit은 더 멀리 rewind한다. Prefill의
+"per-query-block" 외부 scope을 둘러싼 "per-KV-tile" 내부 scope을
+지원한다.
+
+## Alternatives
+
+### A1. Temporary를 HBM에 round-trip
+
+중간값을 HBM에 store하고 reload하여 TCM scratch를 "free" 한다. 기각:
+TCM-resident streaming kernel을 HBM-bandwidth-bound로 만들어 목적에
+정반대이며, op_log를 불필요한 DMA로 오염시킨다.
+
+### A2. Tiled `tl.composite` (scheduler-관리 scratch)
+
+`tl.composite` 는 PE_SCHEDULER 내부에서 tile당 scratch를 자동 재활용
+한다. ADR-0062 A1과 마찬가지로 매력적이지만 flash-capable composite
+kind (두 GEMM + carry되는 `(m,l,O)`) 가 필요하므로 훨씬 큰 변경에
+종속된다. `scratch_scope` 는 작고 일반적인 primitive로 greenlet kernel
+에 동일한 memory 거동을 제공한다. 두 방식은 공존 가능.
+
+### A3. Scratch 예산 증가
+
+`pe_tcm.kernel_scratch_mb` 를 증가. 해결책으로는 기각: O(`n_tiles`)
+scratch leak이 여전히 존재하면서 context-length ceiling만 선형적으로
+밀어내며, hardware 실제 모습 (실 TCM은 작고, flash attention의 핵심은
+O(1) working set) 을 잘못 표현한다. 거친 조절 knob으로만 유용, 재활용의
+대체재는 아니다.
+
+## Consequences
+
+### Positive
+- 인위적인 `S = 16` validation-scale ceiling 제거; 시간/데이터 mode
+  양쪽에서 의미 있는 context length 가능.
+- Flash attention의 O(1) working-set 특성을 충실히 모델링.
+- 작고 일반적인 primitive (모든 tiled kernel이 이점).
+
+### Negative
+- Use-after-scope 버그는 stale 바이트를 조용히 읽는다. D2 계약, attention
+  helper 안에서 scope discipline 공유, (선택적으로) rewound 영역을
+  poison시키는 debug build로 완화.
+- ADR-0062와 조정 필요: prefetch buffer는 tile iteration 사이에 살아
+  있어야 하므로 scoped arena가 아니라 persistent arena에 속한다.
+
+## Test Requirements
+
+1. **Recycling**: `tl.scratch_scope()` 내부의 `N` tile loop에서 peak
+   `_scratch_cursor` 가 `N`과 무관하게 한 tile의 footprint로 bound됨
+   (오늘날은 선형 증가 후 overflow).
+2. **Correctness**: scope이 있는 flash sweep이 scope이 없는 동일 sweep
+   과 (재활용 없이도 수렴하는 작은 `N` 에서) 동일한 `O` 를 산출
+   (Phase 2).
+3. **Long context**: scope 없이 1 MiB를 초과하는 `N` 의 sweep이 완료
+   (오늘날 `S=16` cap이 회피하는 그 실패).
+4. **Persistent-vs-scoped isolation**: scope 외부에서 할당된 running
+   `(m,l,O)` 가 `__exit__` 이후에도 올바른 값을 유지.
+5. **Nesting**: 중첩된 scope이 올바른 save-point로 rewind.

docs/adr-ko/ADR-0064-perf-cpu-issue-cost-model.md

@@ -0,0 +1,366 @@
+# ADR-0064: 구조적 CPU dispatch cost 모델 (`logical_bytes` + FIXED + R)
+
+## Status
+
+Accepted (Revision 2)
+
+> **ADR-0060** (AHBM GQA Fused Attention) 와 **ADR-0065** (flat-ops
+> composite + 첫 stateful recipe) 의 보조 ADR. 그 hybrid 의 핵심 — "GEMM 은
+> `tl.composite` 로, softmax merge 는 커널에서" — 은 **PE_SCHEDULER 로
+> tiling 을 offload 해서 CPU 가 굵은 descriptor 만 issue 하고 앞서가며
+> 엔진을 포화시킴** 으로 이깁니다. 그 win 이 현재 시뮬레이터에선 보이지
+> 않습니다 (per-op CPU issue cost = 0).
+>
+> Revision 2 는 원안의 **op-type calibration 표** 를 **logical_bytes 기반
+> 구조 공식** 으로 대체합니다 — op-type 별 calibration 불필요, 새 op_kind
+> 도 자동 적용.
+
+## Context
+
+### 현재 코드
+
+- 모든 `tl.*` op 가 cmd emit 전에 `_emit_dispatch_overhead()` 호출
+  (`tl_context.py:196-212`) → `dispatch_cycles > 0` 일 때만
+  `PeCpuOverheadCmd(cycles=dispatch_cycles)` 발행. knob 은 op kind 무관
+  **uniform** 이며 두 실행 경로 모두 **0** 하드코딩
+  (`pe_cpu.py:101` greenlet, `:195` replay).
+- ⇒ 명령 issue (descriptor 구성 + scheduler 큐 push) 가 PE_CPU 에서
+  **0 ns** 비용.
+- `PeCpuOverheadCmd` 는 PE_CPU 에서 `yield env.timeout(cmd.cycles)` 로
+  소비 (`kernel_runner.py:131-132`).
+
+### uniform-zero 가 hybrid 에 부적합한 이유
+
+ADR-0060 §1 의 핵심: composite 1 개가 `N_tiles` 분량의 GEMM tiling 을
+offload → CPU 는 `O(1)` 개의 굵은 cmd 만 issue, primitive 경로는
+`O(N_tiles × ops/tile)`. issue cost = 0 이면 모델은:
+- primitive 경로의 *CPU 가 push 를 못 따라가 엔진 idle* 을 못 보임
+- composite 의 *구성 비용 ≫ primitive 1 개이지만 ≪ 대체된 primitive 들의 합* 을 못 보임
+
+### Revision 1 의 op-type calibration 이 과한 이유
+
+원안은 `cost_table[kind]` (kind = `composite`, `load`, `dot`, `math`, …)
+형태였습니다. 비용:
+- kind 마다 값 필요 (calibration ≥ |kinds|)
+- 새 kind 추가 시마다 entry 추가
+- 그런데 잡으려던 *ratio* — "composite ≫ primitive, ≪ 대체된 primitive 들의 합" —
+  는 *op kind* 가 아니라 *cmd 가 들고 있는 필드 수* 의 함수.
+
+cmd 의 *byte 발자국* 이 자연 proxy: N OpSpec composite 는 1 OpSpec primitive 의 ~N 배 bytes.
+*고정* 부분 (큐 tail 업데이트, completion 등록, MMIO-급 latency) 은 cmd 마다.
+둘 합치면: `FIXED + bytes × R`.
+
+### 이 모델이 실제로 노출하는 것
+
+**1차** 신호는 per-cmd FIXED 비용에 의한 **command-count 감소**.
+**byte** 항은 *비정상적으로 큰* composite 가 공짜로 보이지 않도록 막는
+2차 보정. 현실적 on-die 큐 대역폭 (16 B/cycle, D3) 에서 FIXED 가
+decode opt2 의 opt3 (≈10 cmds/tile) vs opt2 (≈2 cmds/tile) 의 dispatch
+cost 차이 중 ≥85% 차지.
+
+이 framing 이 모델의 한계도 정의: 단일 composite 가 무한히 커지면 byte
+항만으로는 HW 가 받지 못할 정도의 fused command 를 모델이 보상하는 것을
+막지 못함 — 그래서 **D7** 의 descriptor 크기 cap 이 필요.
+
+## Decision
+
+### D1. 구조적 dispatch cost 공식
+
+PE_SCHEDULER 로 가는 모든 PE command 가 PE_CPU 에서 dispatch cycles 소모:
+
+```
+dispatch_cycles(cmd) = FIXED_PER_CMD + cmd.logical_bytes × R
+```
+
+- `FIXED_PER_CMD` (cycles/cmd): 큐 tail 업데이트, MMIO-급 RTT,
+  completion-event 등록 — cmd 크기와 무관 고정 비용.
+- `R` (cycles/byte): scheduler 큐에 cmd 를 직렬화하는 큐-write 대역폭.
+- `cmd.logical_bytes` (int): cmd 의 *HW-논리* byte 수 — D2 룰로 계산,
+  Python `sys.getsizeof` 가 아님.
+
+PE_CPU 는 기존 hook 그대로 `PeCpuOverheadCmd(cycles=dispatch_cycles(cmd))`
+를 dispatch 전에 발행 — 사이클 값만 변경.
+
+### D2. `logical_bytes` 룰
+
+각 PE command dataclass 가 `logical_bytes: int` (property) 노출. 룰
+(HW 친화적, Python overhead 무시):
+
+| 필드 종류 | bytes |
+|---|---|
+| cmd framing (cmd 종류 discriminator + completion id 참조) | 4 |
+| Opcode (op kind enum) | 1 |
+| Enum (scope 등) | 1 |
+| `TensorHandle` 참조 (address only — shape/dtype 은 descriptor table 가정) | 8 |
+| Scalar (int/float) | 4 |
+| Tuple length marker | 1 |
+
+`CompositeCmd` 는 `ops` 와 `rw_handles` 재귀 합산:
+
+```python
+@property
+def logical_bytes(self) -> int:
+    return (
+        4                                # framing
+        + 1 + sum(op.logical_bytes for op in self.ops)
+        + 1 + 8 * len(self.rw_handles)
+    )
+```
+
+`OpSpec`:
+
+```python
+@property
+def logical_bytes(self) -> int:
+    return (
+        1 + 1                                        # opcode + scope
+        + 1 + 8 * len(self.operands)                 # named operand handles
+        + (8 if self.out is not None else 0)         # out handle
+        + 1 + sum(_extra_bytes(v) for v in self.extra.values())
+    )
+
+
+def _extra_bytes(v) -> int:
+    """OpSpec.extra 값의 타입-aware byte 카운트."""
+    if isinstance(v, bool):              return 1
+    if isinstance(v, (int, float)):      return 4
+    if isinstance(v, (tuple, list)):     return 1 + 4 * len(v)   # shape, axes, …
+    if isinstance(v, str):               return 1                # opcode-like tag
+    return 4                                                     # default scalar
+
+
+# extra 의 예시 타입:
+#   m, k, n            int   → 4 each
+#   reduce_axis        int   → 4
+#   shape=(64, 64)     tuple → 1 + 8 = 9
+#   factor=1.0         float → 4
+```
+
+(`DmaReadCmd`, `MathCmd` 등도 같은 룰 — dataclass 당 property 1 개,
+약 3 줄.)
+
+**카운팅 룰 — per-op 합산, dedup 없음.** `OpSpec` 안의 각 `TensorHandle`
+참조는 독립적으로 카운트 — 같은 핸들이 여러 OpSpec 에 등장하거나
+`rw_handles` 에도 있어도 마찬가지. `rw_handles` block 은 cross-composite
+해저드 추적 (ADR-0065 D6.3) 의 메타데이터로 `ops` 의 operand 참조와
+**별도** 카운트. 중복 제거 없음. HW 와 일치: descriptor 가 각 operand
+slot 을 독립 address field 로 인코딩하고 dispatcher 가 `rw_handles` 를
+별도 metadata block 으로 추적. 예: `OpSpec(kind="mul_bcast", operands=
+{"src_a": O, "src_b": corr}, out=O)` 는 핸들 `O` 를 *두 번* 카운트
+(`src_a` 한 번, `out` 한 번); 만약 `O` 가 둘러싼 CompositeCmd 의
+`rw_handles` 에도 있으면 세 번째 카운트.
+
+### D3. Default — 일반 composite ≈ 43 ns 기준
+
+Anchor: **DMA-staged GEMM 경로용 단일-OpSpec composite** —
+`OpSpec(kind="gemm", ...)` 1 개; DMA stage 는 PE_SCHEDULER 가 operand
+`space` 에서 자동 삽입 (ADR-0065 D4), **`logical_bytes` 에 등장하지 않음**
+(커널이 별도 cmd 로 emit 안 함). 분해:
+
+```
+framing                  4
+ops tuple length         1
+GEMM OpSpec             40   (opcode 1 + scope 1 + 1 + 2 handles 16
+                              + out 8 + 1 + extra m/k/n 12)
+rw_handles tuple length  1
+rw_handles content       8   (출력 RW 핸들 1 개)
+─────────────────────────────
+total                  ~54 bytes
+```
+
+Dispatch 목표 ≈43 ns. on-die producer→consumer 큐 16 bytes/cycle
+(일반 on-die descriptor queue 폭) 가정.
+
+```
+FIXED_PER_CMD        = 40 cycles
+R                    = 0.0625 cycles/byte  (= 16 bytes/cycle)
+```
+
+검증: `40 + 54 × 0.0625 = 43.375 cycles ≈ 43 ns` ✓
+
+Cycle→ns 변환은 PE 노드의 기존 `clock_freq_ghz` attr (PE_MATH `_compute_ns`
+가 쓰는 같은 것) 사용. cost-model knob 은 **cycle-domain 만** — clock
+설정 중복 안 함.
+
+### D4. topology config override
+
+default 는 `pe_cpu.py` 에 내장. topology yaml 의 PE 노드 attrs 아래
+`pe_cost_model:` 절로 override (cycle-domain knob 만; clock 은 PE 의
+기존 `clock_freq_ghz` 사용):
+
+```yaml
+pe:
+  attrs:
+    clock_freq_ghz: 1.0                    # 기존, cycle→ns 변환용
+    pe_cost_model:
+      fixed_per_cmd_cycles: 40
+      byte_cycles_recip: 0.0625            # = 16 bytes/cycle
+      max_composite_logical_bytes: 1024    # D7 — descriptor 크기 cap
+```
+
+누락된 키는 default. dispatch 공식은 PE_CPU init 시
+`node.attrs["pe_cost_model"]` 에서 읽음.
+
+### D5. Scope — 무엇이 cost 를 내고 무엇이 안 내나
+
+| 경로 | dispatch cost? |
+|---|---|
+| PE_CPU → PE_SCHEDULER 의 모든 `PeCommand` | **예** |
+| `PeCpuOverheadCmd` 자체 (이미 cycles 명시) | **아니오** (공식 우회) |
+| PE_SCHEDULER 가 자동 생성한 Stage (DMA_READ/WRITE/FETCH/STORE) | **아니오** (PE_SCHEDULER 내부) |
+| 엔진 compute latency (PE_DMA `drain_ns`, GEMM/MATH `_compute_ns`) | **변화 없음** — 엔진에 그대로 (SPEC §0.1) |
+
+"latency 는 모델링된 컴포넌트에서" invariant 유지 — dispatch cost 는
+*추가* CPU-side 시간, 엔진 시간에 fold 안 함.
+
+### D6. 설정 가능 값; goldens 재생성
+
+issue cost 를 0 → non-zero 로 바꾸면 **모든** bench 의 latency 변화.
+이 ADR land 시 골든 latency **한 번** 재생성 — ADR-0062 D3 lazy-load 와
+같은 패턴. 재생성 후 동일 calibration 이 ADR-0065 opt2 측정에 적용.
+
+### D7. Composite 크기 cap (하드 제한 — validation error)
+
+각 `CompositeCmd` 의 `logical_bytes` 가 **`max_composite_logical_bytes`**
+(default **1024 bytes**, D4 로 override) 로 제한. `logical_bytes` 가 cap
+을 초과하는 composite 는 host-side TLContext 가 emit 시점에 `ValueError`
+로 **거부** — **자동 segmentation 없음**. 커널 작성자가 각 `CompositeCmd`
+가 descriptor capacity 에 맞도록 recipe 를 재구성(예: 여러 개의 더 작은
+composite 로 명시적으로 분할)해야 함.
+
+**cap 이 필요한 이유.** 실제 하드웨어는 descriptor queue entry 크기,
+scheduler parser buffer, command SRAM, firmware 입력의 하드 제한 보유.
+cap 없으면 `FIXED + bytes × R` 모델이 하드웨어가 받지 못할 정도로 큰
+fused composite 를 보상 (예: primitive 100 개를 composite 1 개로 fuse,
+FIXED 1 회만 지불).
+
+**1024 bytes 라는 *특정* 숫자의 근거.** 이는 측정된 HW 숫자가 아니라
+**safe engineering limit** — kernbench 코드베이스의 현재 모든 알려진
+composite (가장 큰 것이 decode opt2 의 `#2` ~322 bytes) 보다 훨씬 위이면서
+미래 recipe 도 지켜야 할 *유한* descriptor capacity 를 의도적으로 표현.
+값은 topology 별 override (D4); 실제 HW reference 등장 시 재캘리브레이션.
+이 default 의 역할은 cap 을 *원칙으로 존재시키는 것*, 특정 HW 에 맞추는
+것이 아님.
+
+**자동 segmentation 대신 하드 에러를 택한 근거.** 초과 composite 를 자동
+분할(원래 Revision 2 제안)하는 것은 inter-segment 순서, `rw_handles`
+공유, completion 체이닝 등 emitter 복잡도를 추가하면서, 실질적으로는
+하드웨어가 허용하는 것보다 큰 descriptor 를 요청한 커널을 덮어주는
+것일 뿐. 명시적 에러로 표면화하면 emitter 가 단순해지고, HW 제약이
+작업을 어떻게 분할할지 가장 잘 아는 커널 작성자에게 보임.
+
+Decode opt2 의 `#2` composite (10 ops, ~322 bytes) 는 1024 cap 안에
+편안히 — GQA workload 에 에러 없음.
+
+## Alternatives
+
+### A1. Revision 1 의 op-type calibration 표 유지
+
+기각: calibration 비용이 |kinds| 에 비례, 그리고 표가 잡으려던 *ratio* 는
+구조적으로 cmd 크기의 함수. 구조 공식은 같은 정성적 동작을 N 개가 아닌
+2 개의 calibratable 숫자로 달성.
+
+### A2. byte-only 공식 (FIXED 항 없음)
+
+기각. FIXED = 0 이면 opt2 (ADR-0065 Option Y) 가 opt3 를 **이기지 못함**
+— per-tile dispatch *총 bytes* 가 비슷 (opt3 ≈ 232, opt2 ≈ 380); win 은
+전적으로 *per-cmd fixed cost 횟수 감소* 에서 옴. byte-only 는 모델이
+보여야 할 신호를 지움.
+
+### A3. PE_SCHEDULER 에 dispatch 비용 부과
+
+기각: saturation 질문은 *"CPU 가 descriptor 를 push 하는 속도가 엔진을
+계속 바쁘게 할 수 있는가?"* — **PE_CPU** 의 issue-bandwidth 특성.
+scheduler 에 부과하면 CPU back-pressure 모델링 안 됨.
+
+### A4. DMA program/setup 시간을 별도 fixed per-descriptor 비용으로 모델
+
+연기: 처음엔 descriptor-program 비용을 **issuing op 의** dispatch cost 에
+fold. calibration 이 분리 필요성을 보이면 PE_DMA fixed setup 으로 분리.
+
+## Consequences
+
+### Positive
+- hybrid 의 CPU-offload / saturation win (ADR-0060 §1) 이 **측정 가능**,
+  구조적으로 정직한 모델 (calibration 표 없음).
+- 새 op_kind 추가 (예: ADR-0065 의 `softmax_merge` 8-step recipe) 가
+  *zero cost* — 같은 공식 자동 적용.
+- 하드웨어에 더 충실 (queue-head MMIO RTT + queue-write 대역폭).
+
+### Negative
+- **모든** bench goldens 변화 → 일회성 재생성 (D6); CI 골든 fixtures 업데이트.
+- 두 calibration knob (FIXED, R) 필요; default 는 문서화된 가정에 기반 —
+  HW reference 없는 동안엔 절대 latency 를 잠정 취급, **ratio** 만 방어.
+- 각 PE command dataclass 에 작은 `logical_bytes` property 추가.
+
+## Open review items
+
+1. **FIXED 와 R 의 calibration 출처.** default 는 "일반 composite ≈ 43 ns
+   + on-die 큐 16 bytes/cycle"; on-die descriptor 큐로 합리.
+   HW reference 등장 시 재방문.
+2. **대형 composite 에서 Scheduler plan-gen 비용.** 0 유지 — D5 가
+   PE_SCHEDULER 의 plan 생성을 dispatch 공식 밖에 둠. D7 cap (1024 bytes
+   ≈ 30–35 OpSpec) 이 최악을 묶지만, cap 근처 composite 도 현재
+   zero-cost 모델에서 1-op composite 와 같은 scheduler 비용. stress test
+   (대형 composite microbench) 가 scheduler-bound 행동을 보이면
+   per-op-count `overhead_ns` 노출.
+3. **Override 위치.** topology yaml 의 PE 노드 attrs 아래 `pe_cost_model:`
+   block — 모든 knob 을 한 곳에, 리뷰 가능. clock 은 PE 의 기존
+   `clock_freq_ghz` 에서 — 여기 중복 안 함.
+4. **경로 parity.** greenlet (`_execute_legacy`, `kernel_runner`) 와
+   replay 둘 다 같은 cost model 읽어야 함. 검증.
+5. **R 에 대한 결론의 민감도.** opt2 < opt3 가 합리적 `R` 범위 (큐 대역폭
+   가정) 에서 유지되어야 함. 민감도 sweep 이 Test Requirements (#9) 의
+   일부.
+
+## Test Requirements
+
+테스트는 specific 숫자 anchor 가 아닌 **공식 (D1)** 에 대해 작성 —
+calibration 이 바뀌거나 OpSpec/CompositeCmd 필드가 추가되어도 유효.
+
+1. **공식 보존.** 임의 `CompositeCmd` `c` 에 대해 PE_CPU 의 기록된
+   dispatch overhead 가 `FIXED_PER_CMD + c.logical_bytes × R`
+   (floor/round-off 의 ±1 cycle 내). 여러 composite parametrize:
+   1-OpSpec GEMM composite, 5-op MATH chain, 10-op recipe composite.
+   Default-calibration 숫자 (1-OpSpec composite 의 anchor ≈43 ns) 는
+   *informative* 참고, 테스트 gate 가 아님 — gate 는 공식 등가.
+2. **정성적 ratio (robust).** opt3 per-tile PE_CPU dispatch 가 opt2
+   per-tile dispatch 를 최소 2× 차이로 엄격히 초과 — `opt3 > 2 × opt2`.
+   Default-calibration 모델은 ≈4× 예측; gate 는 느슨한 2× 한계라
+   calibration 이동 (예: HW reference 가 default 대체) 시 깨지지 않음.
+   Informative 숫자 — ADR-0065 §verification 와 DDD-0065 §11 의 모델
+   기대치 참조.
+3. **Override 경로.** topology yaml 의 `pe_cost_model:` block 이 per-PE
+   dispatch cost 변경; block 누락 시 default 복귀. #1 의 공식 등가가
+   override 값에서도 성립.
+4. **`PeCpuOverheadCmd` 우회.** 수동 `tl.cycles(n)` 는 정확히 `n` cycles,
+   `n + dispatch_cycles(...)` 아님.
+5. **double-count 없음.** PE_DMA `drain_ns`, PE_GEMM/MATH `_compute_ns` 가
+   pre-ADR 값과 동일.
+6. **결정성.** 동일 입력 → 동일 op_log + latency (SPEC §0.1).
+7. **경로 parity.** greenlet 과 replay 가 동일 커널에 대해 동일
+   dispatch-cycle 회계.
+8. **Composite 크기 cap (D7).** `logical_bytes >
+   max_composite_logical_bytes` 가 될 composite 는 emit 시점에 `ValueError`
+   (segmentation 없음). cap 이내 composite 는 정상 emit.
+9. **민감도 (정성적).** `R ∈ {0.25, 0.0625, 0.03125}` cycles/byte 에서
+   3 지점 모두 `opt3 > opt2`. 방향 (R 감소 시 ratio 단조 증가) 도 단언,
+   단 절대 ratio 값은 *불필요*.
+
+## Migration
+
+ADR-0064 Revision 2 는 단일 PR 로 land:
+- 각 `PeCommand` dataclass 의 `logical_bytes` property (D2 의 타입-aware
+  extra 카운트)
+- `pe_cpu.py` dispatch 경로의 공식 적용
+- PE_CPU init 의 `pe_cost_model:` override read (cycle-domain knob +
+  `max_composite_logical_bytes`)
+- **composite 크기 cap (D7)** — TLContext-side 하드 cap,
+  `max_composite_logical_bytes` default 1024: cap 초과 composite 는 emit
+  시점에 `ValueError` (자동 segmentation 없음). 기존 bench 는 트리거
+  안 됨 (현재 최대 composite 가 ~322 bytes), 하지만 recipe 가 커질 때
+  descriptor-capacity 제한이 강제되도록 체크 land.
+- 일회성 골든 재생성
+
+land 후 ADR-0065 가 위에 쌓임 — 기존 bench 추가 골든 churn 없음
+(ADR-0065 는 기존 경로의 `CompositeCmd` 의미 보존 refactor; opt2 만 새 bench).

docs/adr-ko/ADR-0065-prog-flat-ops-composite-softmax-merge.md

@@ -0,0 +1,440 @@
+# ADR-0065: 평평한 ops `CompositeCmd` + 첫 stateful recipe (`softmax_merge`)
+
+## Status
+
+Accepted
+
+> **ADR-0060** (AHBM GQA Fused Attention) 의 보조 ADR. §5.6 / §8 item 4 의
+> carve-out 을 정확히 구현: decode opt2 = (#1 기존 Q·Kᵀ GEMM composite) +
+> (#2 softmax + P·V + online-softmax merge 를 담은 단일 composite). ADR-0060
+> §5.6 의 "ex_composite" / "flash_pv_merge" 명칭은 *retire* — 본 ADR 이
+> 정식 모양을 확정: 기존 `tl.composite` 입구를 그대로 사용하면서 두 가지
+> 구조적 추가 — (a) 평평한 `ops` 튜플, (b) MATH micro-op 시퀀스로
+> 펼쳐지는 첫 stateful recipe `softmax_merge`.
+
+## Context
+
+### ADR-0060 가 이 ADR 에 남긴 것
+
+ADR-0060 §5.6 는 opt3 (software pipelining) 를 *지금 ship* 으로 권고,
+opt2 는 — per-tile dispatch 적음, K-before-V DMA priority,
+**ADR-0064 의 cost model 이후에만 측정 가능** — 후속으로 carve out.
+
+§8 item 4 가 carve-out 의 sizing: "재방문 시 **두 개** composite 로 분할 —
+`#1` = Q·Kᵀ (기존 composite + `scale`), `#2` = softmax + P·V + 온라인-softmax
+누산기 merge". 본 ADR 이 `#2` 를 구현.
+
+### composite 가 구조 변경이 필요한 이유
+
+현재 `CompositeCmd` 는 `(op, a, b, out_addr, ops=(head, *epi))` 모양 + 암묵 컨벤션:
+- `op` 가 head 엔진 선택 (gemm | math)
+- `ops[0]` 는 head OpSpec, `ops[1:]` 는 epilogue OpSpec 으로 head *이후*
+  per-output-tile 또는 per-K-tile 에 실행 (scope 결정)
+
+decode opt2 는 MATH op 이 head GEMM **이전** 에 실행 (`m, l, O` 의
+online-softmax 갱신 + GEMM 입력 `P` emit) 이 필요. 현 모양엔 자리 없음.
+
+이전에 검토 후 기각된 경로:
+- softmax_merge 를 #1 (Q·Kᵀ) 의 *epilogue* 로 두기 — `Sj` 를 읽고 *다음*
+  GEMM (#2 P·V) 가 소비할 `P` 를 쓰는 구조라 순환.
+- 단일 mixed-engine recipe `flash_pv_merge` 가 MATH + GEMM + MATH 흡수 —
+  engine boundary 위반, PE_SCHEDULER 복잡화.
+
+깔끔한 모양: **command 포맷에서 head/epilogue 구분 제거**. `ops` 를 평평한
+순서 튜플로, 각 op 의 `scope` + position 이 tile-loop 배치 결정. "prologue"
+개념은 *사용자 API* 의 `tl.composite(...)` 에 ergonomic 으로 남되 —
+**HW command 는 보지 않음**.
+
+### recipe 가 (generic op list 가 아닌) 이유
+
+decode opt2 의 `#2` MATH chain 은 8 단계 고정 구조 (online softmax merge).
+커널 작성자에게 8 단계 직접 와이어링은 verbose + correctness hazard
+(단계 순서 의존). *Recipe* — TLContext 가 8 단계로 (주소 미리 채워)
+펼치는 단일 명명 op (`softmax_merge`) — 가 커널을 간결히 + 구조를 단일
+source 에 보존. recipe 표는 TLContext (compiler analog) 에 위치;
+PE_SCHEDULER 는 안 읽음. D5 boundary 참조.
+
+## Decision
+
+### D1. `CompositeCmd` 는 평평한 ordered op 리스트
+
+```python
+@dataclass(frozen=True)
+class CompositeCmd:
+    completion: CompletionHandle
+    ops: tuple[OpSpec, ...]                       # ordered MATH/GEMM ops
+    rw_handles: tuple[TensorHandle, ...] = ()     # cross-composite hazard
+    data_op: bool = True
+
+    @property
+    def logical_bytes(self) -> int: ...           # ADR-0064 D2
+```
+
+이전의 `(op, a, b, out_addr, out_nbytes, math_op)` 필드는 `ops` 에 흡수.
+`ops` 의 첫 OpSpec 이 `kind == "gemm"` 이면 head GEMM; 앞의 OpSpec 들이
+pre-loop, 뒤의 OpSpec 들이 `scope` 에 따라 post-loop / per-tile epilogue.
+
+### D2. `OpSpec` 이 named operand 보유
+
+```python
+@dataclass(frozen=True)
+class OpSpec:
+    kind: str                                     # "gemm" | "rmax" | ...
+    scope: Scope                                  # KERNEL | K_TILE | OUTPUT_TILE
+    operands: dict[str, TensorHandle]             # named — D5 참조
+    extra: dict[str, Any]                         # scalars, axes, m/k/n 등
+    out: TensorHandle | None = None               # 명시적 write-back handle
+
+    @property
+    def logical_bytes(self) -> int: ...
+```
+
+named operand 로 PE_SCHEDULER 가 (positional 컨벤션 없이) 엔진 port 에
+입력 라우팅 (예: GEMM 은 `operands["a"]`, `operands["b"]`).
+
+### D3. Position + scope 가 tile-loop 배치 결정
+
+**Semantics 분리.** *Position* 이 **phase** (tile-loop 수명주기 어디서
+실행되는지) 결정; *scope* 가 그 phase 안에서의 **repetition** 결정.
+`Scope.KERNEL` 의미는 "**`CompositeCmd` invocation 당 1 회**" — kernel
+launch 당 1 회 아님. 그 1 회 실행이 op 의 시퀀스 내 *position* 에서
+발생. 같은 KERNEL 값이 OpSpec 이 GEMM op 의 앞/위치/뒤에 있느냐에 따라
+*phase* 가 달라짐.
+
+PE_SCHEDULER 가 `cmd.ops` 에서 GEMM op 검색 (composite 당 ≤ 1, D6 참조).
+인덱스 `g`:
+
+| OpSpec position | scope | Phase | plan 내 배치 |
+|---|---|---|---|
+| `0 .. g-1` | KERNEL | pre-tile-loop | tile loop 진입 전 1 회 |
+| `g` | KERNEL | head GEMM | `extra["m"], extra["k"], extra["n"]` 로 tile loop drive |
+| `g+1 .. ` | K_TILE | in-accumulation | per K-tile epilogue (K 누산 loop 안) |
+| `g+1 .. ` | OUTPUT_TILE | per-output-tile | per (m, n) tile epilogue (K 누산 후) |
+| `g+1 .. ` | KERNEL | post-tile-loop | tile loop 종료 후 1 회 |
+
+GEMM op 가 없는 composite (예: MATH-only) 는 모든 op 를 KERNEL-scope
+직렬로 처리 — phase 가 "순서대로 single-shot" 으로 collapse.
+
+### D4. PE_SCHEDULER 가 operand `pinned` 플래그 보고 DMA 자동 삽입
+
+PE_SCHEDULER 가 GEMM 의 `operands` 검사. 각각:
+- **not `pinned`** (데이터가 HBM 에 있어 스트리밍 필요) → FETCH/GEMM Stage
+  앞에 DMA_READ Stage 삽입
+- **`pinned`** (이전 `tl.load` 로 이미 TCM 에 staged, 또는 recipe 의 TCM
+  scratch / primary-out `P`) → DMA Stage 없음, in-place 소비
+
+커널은 composite operand 의 명시적 DMA cmd 를 절대 emit 안 함;
+PE_SCHEDULER 가 핸들에서 추론. `tl.ref` operand 는 not pinned (DMA_READ);
+`tl.load` 결과와 recipe scratch / primary-out 은 pinned (in-place).
+
+> **as-built 노트 (P3).** D4 의 초안은 DMA 결정을 `space` 태그
+> (`hbm`→DMA, `tcm`→in-place) 에서 도출하며 현재 `space` 값
+> (오늘 `tl.load`=`hbm`, `tl.ref`=`tcm`) 을 뒤집어야 했음. 구현은 기존
+> **`pinned` 플래그**를 DMA 신호로 그대로 유지 — `pinned` 이 이미
+> "이미 TCM 에 있음, DMA skip" 을 인코딩하고, `space` 뒤집기는 기존 bench
+> Stage 시퀀스를 바꾸면서 기능적 이득이 없음. `space` 통일 + `pinned`
+> 폐기는 후순위 cleanup 으로 연기.
+
+**Prologue recipe op 은 TCM-only (Phase 1 한계, TLContext emit 시 강제).**
+*prologue recipe* (비 GEMM) OpSpec 의 모든 operand 는 **반드시** TCM 상주;
+HBM 핸들을 recipe operand 로 전달하면 emit 시 validation error. Phase 1
+에서 DMA staging 을 head op 의 operand 로 제한. **head op 자체** (gemm
+또는 math) 는 기존 DMA-staged-from-HBM 동작 유지 — D4 의 TCM-only 룰은
+head 에 적용 **안 됨**. 이 invariant 는 D6 #7 에 재기술.
+
+### D5. RECIPE_DESCRIPTORS — TLContext 내부, `pe_commands.py` 가 아님
+
+Recipe 는 새 TLContext-adjacent 모듈
+(`src/kernbench/triton_emu/tl_recipes.py`) 에 위치. PE_SCHEDULER 는
+import **안 함**. 첫 recipe:
+
+```python
+@dataclass(frozen=True)
+class RecipeDescriptor:
+    operands: dict[str, str]                       # name → "R" | "RW"
+    primary_out: PrimaryOutSpec | None             # 암묵 primary output 의 type/shape 룰
+    tile_alignment: Literal["single_shot", "tile_aligned"]
+    internal_scratch_bytes_fn: Callable[..., int]
+    engine_seq: tuple[EngineOp, ...]               # TLContext 가 평평한 OpSpec 으로 펼침
+
+
+RECIPE_DESCRIPTORS["softmax_merge"] = RecipeDescriptor(
+    operands={"s": "R", "m": "RW", "l": "RW", "O": "RW"},
+    primary_out=PrimaryOutSpec(
+        from_shape="s", from_dtype="s", transform="identity",
+    ),
+    tile_alignment="single_shot",
+    internal_scratch_bytes_fn=lambda G, TILE, d, bpe: bpe * (
+        G + G + G + G * TILE + G  # m_loc + m_new + corr + P + l_loc
+    ),
+    engine_seq=(
+        EngineOp("MATH", "rmax",       src="s",         dst="m_loc",  reduce_axis=-1),
+        EngineOp("MATH", "max_elem",   src_a="m",       src_b="m_loc", dst="m_new"),
+        EngineOp("MATH", "exp_diff",   src_a="m",       src_b="m_new", dst="corr"),
+        EngineOp("MATH", "exp_diff",   src_a="s",       src_b="m_new", dst="P",     bcast_axis=0),
+        EngineOp("MATH", "rsum",       src="P",         dst="l_loc",   reduce_axis=-1),
+        EngineOp("MATH", "fma",        src_a="l",       src_b="corr",  src_c="l_loc", dst="l"),
+        EngineOp("MATH", "mul_bcast",  src_a="O",       src_b="corr",  dst="O",     bcast_axis=1),
+        EngineOp("MATH", "copy",       src="m_new",     dst="m"),
+    ),
+)
+```
+
+TLContext 가 `tl.composite(prologue=[{"op": "softmax_merge", ...}],
+op="gemm", ...)` 에서:
+1. operand R/RW 를 `RECIPE_DESCRIPTORS["softmax_merge"]` 와 검증.
+2. scratch (m_loc, m_new, corr, P, l_loc) 를 scratch_scope helper 로 할당 (ADR-0063).
+3. primary output P 의 shape 을 `s.shape` 에서 derive (identity).
+4. `engine_seq` 를 8 개의 평평한 MATH OpSpec 으로 펼침 (모든 주소/크기 채움).
+   각 OpSpec 의 `scope = KERNEL`.
+5. **Auto-bind (충돌 검사, D6.6).** head GEMM 이 `operands["a"]` 를
+   *명시하지 않은* 경우에만 `operands["a"] = P_handle` auto-bind.
+   커널이 `a` 를 명시적으로 제공했다면 validation error — 모호한
+   primary-output 대체.
+6. `CompositeCmd(ops=(8 MATH + 1 GEMM + epilogue), rw_handles=(m, l, O))` emit.
+
+**RECIPE_DESCRIPTORS 는 HW 경로 어디에도 안 나타남**. PE_SCHEDULER 는
+평평한 ops 리스트만 봄.
+
+### D6. Invariants
+
+1. **GEMM 수.** `0 ≤ count(op.kind == "gemm" in cmd.ops) ≤ 1`. GEMM 0 개면
+   MATH-only composite.
+2. **softmax_merge 가 V operand 없음.** `RECIPE_DESCRIPTORS[
+   "softmax_merge"].operands` 에 HBM-resident handle 없음 → prologue 동안
+   DMA 삽입 없음 → V DMA 는 GEMM head 시작 시에만 발생 → decode opt2 의
+   #1 / #2 간 K-before-V DMA priority 자연 강제.
+3. **Cross-composite strict FIFO.** PE_SCHEDULER 가 in-flight `rw_handles`
+   추적. 신규 composite 의 `rw_handles` (또는 op 입력) 가 in-flight 와
+   교차하면 모든 이전 composite 완료까지 대기 — out-of-order reorder 없음.
+4. **legacy caller 후방 호환.** 사용자 API `tl.composite(op="gemm", a, b,
+   epilogue=[...])` 보존. TLContext 가 내부적으로 평평한 ops `CompositeCmd`
+   로 lowering.
+5. **PE_SCHEDULER recipe-free.** PE_SCHEDULER 는 RECIPE_DESCRIPTORS import
+   안 함, "softmax_merge" 의 존재 모름, recipe 별 코드 분기 없음.
+6. **암묵 operand 대체 없음 (auto-bind 충돌).** prologue recipe 가
+   `primary_out` 선언하고 head op 의 auto-bind 대상 operand (예: GEMM
+   `a`) 도 *커널이 명시적* 으로 제공하면, TLContext 가 emit 시 validation
+   error. 커널은 recipe 가 바인딩하게 두거나 (operand 생략) **혹은**
+   명시적으로 지정 (prologue 생략, 또는 `primary_out` 없는 recipe 사용)
+   해야 함. "어느 값이 이기느냐" 의 모호함 방지.
+7. **Prologue MATH operand TCM-only.** D4 의 재기술: *prologue recipe*
+   (비 GEMM) OpSpec 의 모든 operand 는 TCM 상주. head op (gemm 또는 math)
+   은 예외 — HBM 스트리밍 허용. Phase 1 에서 TLContext emit 시 강제.
+
+### D7. Boundary 요약 (compiler vs scheduler vs engine)
+
+```
+HOST                                       DEVICE
+TLContext (compiler analog)          PE_SCHEDULER (scheduler/dispatcher)
+  - RECIPE_DESC 와 검증                 - 평평한 ops 리스트 스캔
+  - primary_out shape derive            - GEMM 식별 (≤ 1)
+  - scratch 할당                        - position-기반 stage 배치
+  - recipe → 평평한 OpSpec 펼침        - space 보고 DMA stage 자동 삽입
+  - rw_handles 계산                     - strict-FIFO RW 해저드 추적
+  - CompositeCmd emit                   - Stage 리스트 emit
+  imports: RECIPE_DESCRIPTORS            imports: recipe 관련 없음
+
+                ENGINES (PE_MATH / PE_GEMM / PE_DMA)
+                  - Stage.params 읽음 (op_kind, addresses, n_elements)
+                  - 기존 latency 모델 변화 없음
+                  imports: recipe 관련 없음
+```
+
+### D8. Composite 출력 handle + 출력-space DMA
+
+`tl.composite(...)` 는 **출력 `TensorHandle`** 을 반환(단순 `CompletionHandle`
+아님) — composite 결과를 `tl.dot` 결과처럼 downstream op 에 먹일 수 있음.
+handle 의 `pending` 필드가 composite 완료를 참조; downstream consumer 가
+auto-await(`_await_pending`, ADR-0062 lazy 패턴), `tl.wait(handle)` 동작.
+
+**`out` 은 항상 `TensorHandle`** (raw 주소 아님) — 위치는 커널의 선택, 추측 안 함:
+
+- **HBM 출력** → `out=tl.ref(addr, shape)`. `tl.ref` 가 `space="hbm"` handle
+  반환 → composite 가 DMA_WRITE 로 write-back. STORE + DMA_WRITE 는 composite
+  파이프라인 **내부**에 유지 — fused op 의 일부이지 별도 `tl.store` 아님.
+- **In-place TCM** → `out=<기존 TCM handle>` (예: recipe 누적기 `O`).
+  `space="tcm"` → STORE 만, 결과 TCM 에 남음 (chainable).
+- **미지정** → TLContext 가 **TCM scratch 자동 할당** (`tl.dot` 과 동일) 후 반환.
+
+편의상 `out_ptr: int` 는 **HBM 출력 shorthand** 로 허용 —
+`out=tl.ref(out_ptr, <출력 shape>)` 와 동치(컴파일러가 `space="hbm"` handle 로
+감쌈). 출력 미지정 시 scratch 주소는 커널이 아니라 컴파일러가 소유.
+(`tl.ref` 는 `space="hbm"` 반환 — HBM 데이터 참조이므로. operand-**입력** DMA
+결정은 D4 대로 `pinned` 기반이라 이 라벨이 입력 스트리밍에 영향 없음.)
+
+출력 handle 의 **`space` 가 타일 루프의 write-back 을 결정** (D4 의 operand
+`pinned` 규칙의 출력판):
+
+| `out.space` | 타일 루프 write-back |
+| --- | --- |
+| `"tcm"` | STORE 만 — TCM 에 남음 (chainable; **DMA_WRITE 없음**) |
+| `"hbm"` | STORE + DMA_WRITE — 최종 HBM 출력 |
+
+이로써 출력이 TCM 에 남는 composite(예: recipe 누적기)가 합법이 되고,
+TCM-scratch 주소가 DMA 엔진의 PA 디코더(고비트 scratch 주소를 거부)로 가는
+일을 막음. 기존 벤치는 HBM 출력을 `tl.ref` 로 감싸므로 write-back
+불변(byte-equal).
+
+## Alternatives
+
+### A1. tile 당 3 composite (prologue 개념 없이)
+
+#2 를 `softmax_merge` + `gemm(P·V)` + `add(O)` 의 3 개로 분할. 장점:
+flat-ops 재구조화 없음 (`add` 가 기존 epilogue 경로 재사용). 단점: tile
+당 dispatch 3 회 (대신 2) — ADR-0060 §8 item 4 sizing ("**두 개** 로
+분할") 과 어긋남; ~32 ns × N_tiles 의 fixed-cost 절감 소실. **기각.**
+
+### A2. Mixed-engine recipe (`flash_pv_merge`)
+
+MATH + GEMM + MATH 흡수하는 단일 recipe (ADR-0060 §5.6 원안 표현). 장점:
+가장 적은 dispatch (1 composite). 단점: PE_SCHEDULER 가 engine-경계
+recipe 알아야 함, "엔진은 Stage.op_kind 만" boundary 위반, recipe 표가
+ISA-급이 되어 HW 인터페이스 오염. **기각.**
+
+### A3. Generic op-list DAG (`tl.ex_composite([...])`)
+
+사용자 정의 임의 op 시퀀스 + 명명 슬롯 + 데이터플로 분석. 장점: 최대
+유연성. 단점: 단일 user (softmax_merge) 가 generic DAG 언어를 정당화 못함;
+PE_SCHEDULER 에 데이터플로 분석 필요; ADR-0060 §8 item 4 의 "inner loop
+흡수" 트랩 재등장. **기각 (Simplicity First).**
+
+### A4. RW-aware scheduler reorder (strict FIFO 가 아님)
+
+PE_SCHEDULER 가 `rw_handles` 교차 없으면 후속 composite 가 in-flight 를
+추월 허용 (예: 다음 tile 의 #1 가 현재 tile 의 #2 와 겹침 — #1 은
+m/l/O 안 만짐). 장점: 더 많은 overlap. 단점: scheduler state 증가, 이득
+한정 (다음 #1 은 어차피 K DMA 대기). **연기** — strict FIFO 먼저;
+RW-aware reorder 는 후속 ADR 가능.
+
+### A5. softmax_merge 를 #1 (Q·Kᵀ) 의 epilogue 로 embedding
+
+오답: softmax_merge 의 출력 `P` 는 #2 의 P·V GEMM 의 입력, #1 의
+epilogue *이후* 실행. embedding 하면 #2 가 자기 입력을 #1 의 epilogue
+scratch 에서 읽는 형태 — 순환 tile-loop 의존성. **기각 (incorrect).**
+
+## Consequences
+
+### Positive
+- ADR-0060 §5.6 / §8 item 4 carve-out 정확 구현: opt2 = tile 당 2
+  composite, K-before-V DMA priority 자연.
+- composite contract 균일화: scope-기반 배치의 평평한 ops 리스트. HW cmd
+  에서 "head" 와 "epilogue" 구조 구분 없음.
+- DMA 가 operand `space` 에서 자동 추론; 커널 표면 단순화.
+- HW 인터페이스 (`CompositeCmd` 모양) 가 recipe 증가에도 작게 유지 —
+  recipe 펼침이 host-side.
+- 새 fused 패턴 (linear+gelu+norm, rmsnorm+linear 등) 은 `RECIPE_DESCRIPTORS`
+  entry 추가만으로 가능; PE_SCHEDULER 변경 없음.
+
+### Negative
+- CompositeCmd 구조 변경 — 모든 현재 caller 의 refactor. 의미 보존
+  (D6.4); ADR-0064 Revision 2 land 후 기존 골든 불변.
+- `OpSpec.operands` 가 positional `tuple[Any, ...]` 에서
+  `dict[str, TensorHandle]` 로 변경 — 기존 epilogue lowering 접촉.
+- 새 TLContext 코드: recipe 펼침 + scratch slot 할당 + primary_out
+  auto-bind. ~80–120 LOC.
+- PE_SCHEDULER 의 `_generate_plan` 에 position-기반 스캔 + DMA 자동 삽입
+  + strict-FIFO RW tracker. ~60–100 LOC.
+- **strict-FIFO RW 해저드 추적이 안전하지만 보수적.** 두 composite 가
+  `rw_handle` 을 공유할 때 신규는 모든 이전 겹치는 composite 가 *완전히
+  완료* 될 때까지 대기 — 원칙적으로 overlap 가능한 경우에도
+  (예: 다음-tile #1 = Q·Kᵀ 가 m/l/O 안 만지므로 현재-tile 의 #2 와
+  overlap 가능하지만, 같은 핸들을 만지는 후속 path 도 FIFO 로 직렬화).
+  더 똑똑한 scheduler 대비 **composite-level overlap 을 과소 노출**.
+  RW-aware reorder (A4) 가 연기된 개선.
+
+## Open review items
+
+1. **identity 를 넘는 recipe shape derivation 룰.** 미래 recipe 가
+   non-identity transform (예: 전치된 primary_out) 필요할 수 있음.
+   `PrimaryOutSpec.transform` 필드가 forward-compat; 필요 시 새 transform 추가.
+2. **Recipe scratch allocator** 의 ADR-0063 `scratch_scope` 통합. 초기
+   배선: TLContext 가 활성 `scratch_scope` 가 있으면 그 안에서 할당,
+   없으면 kernel-scope persistent slot.
+3. **`tl_recipes.py` 위치.** `triton_emu/` 아래 새 모듈. recipe 지식을
+   compiler analog (TLContext) 와 함께 위치.
+4. **GEMM `out=TensorHandle`.** 기존 `out_ptr: int` 와 함께 새 형태.
+   둘 다 허용; handle 형태가 신규 코드에 권장 (handle identity 로
+   strict-FIFO RW 추적 가능).
+
+## Test Requirements
+
+1. **CompositeCmd 평평화 refactor — 의미 보존.** 기존 모든 bench 의
+   op_log 가 refactor 전후로 byte-equal (legacy API `tl.composite(
+   op="gemm", a, b, epilogue=[...])` 가 같은 Stage 시퀀스로 lowering).
+2. **softmax_merge recipe lowering.** TLContext 호출 `tl.composite(
+   prologue=[{"op": "softmax_merge", "s": Sj, "m": m, "l": l, "O": O}],
+   op="gemm", b=V_ref, out=O, epilogue=[{"op": "add", "other": O}])` 가
+   정확히 10 ops (8 MATH + 1 GEMM + 1 MATH(add)) + `rw_handles == (m, l, O)`
+   의 `CompositeCmd` 생성.
+3. **K-before-V DMA priority invariant.** decode opt2 의 op_log 에서
+   #2 의 MATH prologue 동안 V 관련 DMA 없음 (#2 의 V DMA 는 GEMM head
+   시작 시에만).
+4. **Strict-FIFO RW 직렬화.** `rw_handle` 공유하는 두 연속 composite 가
+   dispatch 순서대로 완료; stage 가 interleave 안 함.
+5. **`pinned` 에서 DMA 자동 삽입.** not-`pinned` operand (예: `b=tl.ref(...)`)
+   의 GEMM composite 가 DMA_READ Stage emit; `pinned` operand (예: recipe
+   의 TCM scratch / primary-out, 또는 `tl.load` 결과) 는 emit 안 함.
+   (as-built D4 노트: DMA 결정은 `space` 태그가 아니라 기존 `pinned`
+   플래그 사용.)
+6. **opt2 가 opt3 와 수치 동등.** data mode 에서 opt2 의 최종
+   `(m, l, O)` 가 opt3 와 fp tolerance 안.
+7. **opt2 dispatch ratio (ADR-0064 Rev2 이후) — robust.** opt3 vs opt2
+   per-tile PE_CPU dispatch cycles 가 `opt3 > 2 × opt2` 만족
+   (정성적 gate, calibration-독립). Default-calibration 모델 기대치는
+   ≈ 4.0× (FIXED=40 cycles, R=0.0625 cycles/byte) — DDD-0065 §11 의
+   모델-유도 숫자 참조; 테스트 gate 는 느슨한 `> 2×` 한계라
+   calibration 변경에도 살아남음. Ratio 가 FIXED-dominated —
+   command-count 감소가 1차 신호임을 반영.
+8. **GEMM-count invariant.** GEMM OpSpec 두 개 가진 composite 가
+   TLContext emit 시 validation error.
+9. **Composite 크기 cap 안에 (ADR-0064 D7).** Decode opt2 의 `#2`
+   composite (10 ops, ~322 logical bytes — operand 핸들의 per-op 합산)
+   가 default `MAX_COMPOSITE_LOGICAL_BYTES=1024` 안에 편안히 —
+   segmentation 없음. recipe lowering 테스트가 `cmd.logical_bytes
+   < 1024` 단언.
+10. **Auto-bind 충돌.** `tl.composite(op="gemm", a=A_handle,
+    prologue=[{"op": "softmax_merge", ...}], ...)` 처럼 softmax_merge 가
+    `primary_out` 선언하는데 `a` 도 명시되면 emit 시 validation error
+    (D6.6).
+11. **MATH operand TCM invariant.** `tl.composite(op="math",
+    a=tl.ref(addr, shape), ...)` (HBM-resident 핸들을 MATH operand 로
+    전달) 가 emit 시 validation error (D6.7).
+
+## Dependencies
+
+- **ADR-0060** §5.6, §8 item 4 — 본 ADR 이 구현하는 carve-out.
+- **ADR-0064 Revision 2** — dispatch cost model; 먼저 land, opt2 의
+  fewer-issues win 을 측정 가능하게 함.
+- **ADR-0063** `scratch_scope` — recipe intermediate 가 안에 할당;
+  `m, l, O` 는 persistent arena 로 밖에 할당 (DDD-0060 §6.2 참조).
+- **ADR-0046** `tl_context` contract — `prologue=[...]` kwarg,
+  `out=TensorHandle`, RECIPE_DESCRIPTORS lookup 으로 확장.
+- **ADR-0042** tile plan generators — (재설계 아닌) 확장 — 평평한 ops
+  리스트 처리 + DMA 자동 삽입.
+- **ADR-0014** PE pipeline — boundary 보존: scheduler 가 recipe-free,
+  엔진이 op_kind-opaque.
+
+## Migration
+
+> **구현 순서 노트.** 아래의 ADR-0064-우선 순서가 원래 계획이었음.
+> 실제 구현은 **1↔2 단계를 뒤집어** — ADR-0065 Phase 1 (평평화
+> `CompositeCmd`) 이 ADR-0064 Rev2 *보다 먼저* land. 이유: `logical_bytes`
+> (ADR-0064 D2) 가 평평화 형태 위에서 자연스럽게 정의되므로, refactor 를
+> 먼저 하면 버려지는 transitional `logical_bytes` 를 피함. 각 단계는
+> 독립적으로 안전: P1 은 기존 cost 경로 하에서 의미 보존 refactor
+> (op_log byte-equal); ADR-0064 Rev2 가 이미 평평화된 형태 위에 구조적
+> 공식을 교체. 최종 상태는 어느 쪽이든 동일.
+
+Land 순서 (원래 계획; as-built 1↔2 뒤집힘은 위 노트 참조):
+1. **ADR-0064 Revision 2** (별도 PR): 구조적 dispatch cost +
+   `logical_bytes` property + topology config override + 일회성 골든 재생성.
+2. **ADR-0065 Phase 1** (본 ADR, PR a): `CompositeCmd` 평평화 refactor +
+   legacy-API lowering. refactor only; 골든 불변.
+3. **ADR-0065 Phase 2** (본 ADR, PR b): `tl_recipes.py` + `softmax_merge`
+   recipe + PE_SCHEDULER position-스캔 + DMA 자동 삽입 + strict-FIFO
+   RW tracker.
+4. **ADR-0065 Phase 3** (본 ADR, PR c): `_gqa_decode_long.py` opt2
+   variant + 수치 동등 테스트 + opt3 대비 dispatch-ratio 측정.
+
+상세 구현 계획: **DDD-0065** 참조.

docs/adr-ko/INDEX.md

@@ -0,0 +1,175 @@
+# ADR Index
+
+Auto-generated by `tools/generate_adr_index.py`. Total ADRs: **47**.
+
+Classification mirrors the `/report` skill's section assignment. When adding a new ADR, also add an entry to the `CLASSIFICATION` table in `tools/generate_adr_index.py`.
+
+## Design Principles
+
+- [ADR-0013](./ADR-0013-ver-verification-strategy.md) — 검증 전략 및 Phase 1 테스트 계획
+- [ADR-0033](./ADR-0033-lat-latency-model-assumptions.md) — 레이턴시 모델: 가정 및 알려진 단순화
+
+## High-level Architecture
+
+- [ADR-0003](./ADR-0003-dev-target-system-hierarchy.md) — 타겟 시스템 계층 및 모델링 범위  _(System hierarchy (Tray / SIP / CUBE / PE))_
+- [ADR-0007](./ADR-0007-api-runtime-api-boundaries.md) — 런타임 API 및 시뮬레이션 엔진 경계  _(Runtime API ↔ sim_engine boundaries)_
+- [ADR-0016](./ADR-0016-dev-iochiplet-noc-and-memory-path.md) — IOChiplet NoC와 메모리 데이터 경로  _(IOChiplet NOC and memory data path)_
+- [ADR-0017](./ADR-0017-dev-cube-noc-and-hbm-connectivity.md) — 큐브 NoC와 HBM 연결성  _(Cube NOC and HBM connectivity)_
+
+## Detailed Architecture
+
+One subsection per component file under `src/kernbench/components/builtin/`.
+
+### forwarding
+
+- [ADR-0037](./ADR-0037-dev-forwarding-component.md) — Forwarding 컴포넌트 (forwarding_v1)
+
+### hbm_ctrl
+
+- [ADR-0034](./ADR-0034-dev-hbm-controller-internal-design.md) — HBM 컨트롤러 내부 설계
+
+### io_cpu
+
+- [ADR-0036](./ADR-0036-dev-io-cpu-component-model.md) — IO_CPU 컴포넌트 모델
+
+### m_cpu
+
+- [ADR-0035](./ADR-0035-dev-m-cpu-and-m-cpu-dma-component-model.md) — M_CPU 및 M_CPU.DMA 컴포넌트 모델
+
+### pcie_ep
+
+- [ADR-0038](./ADR-0038-dev-pcie-ep-component-model.md) — PCIE_EP Component Model
+
+### pe_cpu
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_dma
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+- [ADR-0023](./ADR-0023-dev-ipcq-pe-collective.md) — PE-level IPCQ — Inter-PE Collective Communication
+
+### pe_fetch_store
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_gemm
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_ipcq
+
+- [ADR-0023](./ADR-0023-dev-ipcq-pe-collective.md) — PE-level IPCQ — Inter-PE Collective Communication
+
+### pe_math
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_mmu
+
+- [ADR-0039](./ADR-0039-dev-pe-mmu-component-model.md) — PE_MMU Component Model — 컴포넌트 + 유틸리티 이중 역할
+
+### pe_scheduler
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_tcm
+
+- [ADR-0040](./ADR-0040-dev-pe-tcm-component-model.md) — PE_TCM Component Model — 듀얼 채널 BW 직렬화
+
+### sram
+
+- [ADR-0041](./ADR-0041-dev-cube-sram-component-model.md) — Cube SRAM Component Model — terminal scratchpad on cube NoC
+
+### tiling
+
+- [ADR-0042](./ADR-0042-prog-tile-plan-generators.md) — Tile Plan Generators — GEMM/Math 파이프라인 plan 빌더
+
+## Implementation Decisions
+
+### Address Scheme
+
+- [ADR-0001](./ADR-0001-mem-physaddr-layout.md) — 51비트 물리 주소 레이아웃 및 디코딩 계약
+- [ADR-0011](./ADR-0011-mem-memory-addressing-simplification.md) — 메모리 주소 지정 — PA / VA / LA 주소 모델
+
+### Routing & Helper API
+
+- [ADR-0002](./ADR-0002-lat-routing-distance.md) — 라우팅 거리, 순서 및 우회 규칙
+- [ADR-0051](./ADR-0051-lat-routing-helper-api.md) — Routing Helper API — `AddressResolver` + `PathRouter`
+
+### Memory Semantics & Local-HBM Bandwidth
+
+- [ADR-0004](./ADR-0004-mem-memory-semantics-local-hbm.md) — 메모리 시맨틱 및 로컬 HBM 대역폭 보장
+
+### Topology Compilation, Diagrams & Builder Algorithms
+
+- [ADR-0005](./ADR-0005-dev-diagram-views-distance-layout.md) — 다이어그램 뷰 및 거리 기반 레이아웃 규칙
+- [ADR-0006](./ADR-0006-dev-topology-compilation-distance-diagram.md) — 토폴로지 컴파일, 거리 추출, 그리고 자동 다이어그램 생성
+- [ADR-0053](./ADR-0053-dev-topology-builder-algorithms.md) — Topology Builder + Visualizer Algorithms
+
+### Tensor Deployment and Allocation
+
+- [ADR-0008](./ADR-0008-api-tensor-deploy-and-allocation.md) — 텐서 배포 및 할당 (호스트 할당기, PA 우선)
+
+### Kernel Execution and Host-Device Messaging
+
+- [ADR-0009](./ADR-0009-api-kernel-execution-messaging.md) — 커널 실행 메시징 및 완료 시맨틱
+- [ADR-0012](./ADR-0012-api-host-io-message-schema.md) — Host ↔ IO_CPU 메시지 스키마 (PA-우선, PE-태깅)
+
+### CLI Surface and Semantics
+
+- [ADR-0010](./ADR-0010-api-cli-surface-and-semantics.md) — 명령줄 인터페이스 및 실행 시맨틱
+
+### Component Port/Wire Fabric Model
+
+- [ADR-0015](./ADR-0015-dev-component-port-wire-model.md) — 컴포넌트 포트/와이어 모델과 패브릭 라우팅
+
+### Two-Pass Data Execution
+
+- [ADR-0020](./ADR-0020-prog-data-execution-two-pass.md) — 2-Pass 데이터 실행 모델 (타이밍 / 데이터 분리)
+
+### 2D Grid Program Identity
+
+- [ADR-0022](./ADR-0022-prog-program-id-2d-grid.md) — 2D 그리드 program_id 시맨틱
+
+### Parallelism (Launcher, DP, TP, AHBM backend, CCL algorithm)
+
+- [ADR-0024](./ADR-0024-par-sip-tp-launcher.md) — SIP-level Launcher — rank = SIP
+- [ADR-0026](./ADR-0026-par-dppolicy-intra-device.md) — DPPolicy = Intra-Device Only — sip/num_sips 필드 제거
+- [ADR-0027](./ADR-0027-par-megatron-tp.md) — Megatron-style Tensor Parallelism API
+- [ADR-0047](./ADR-0047-par-ahbm-ccl-backend.md) — AHBM CCL Backend — `torch.distributed`-compat shim
+- [ADR-0050](./ADR-0050-par-ccl-algorithm-module-contract.md) — CCL Algorithm Module Contract — `ccl/algorithms/*.py`
+
+### IPCQ Direction Addressing
+
+- [ADR-0025](./ADR-0025-algo-ipcq-direction-addressing.md) — IPCQ Direction Addressing — address-based matching
+
+### Intercube All-Reduce
+
+- [ADR-0032](./ADR-0032-algo-intercube-allreduce.md) — 큐브 간 All-Reduce — pe0 큐브-메시 리듀스 + 다중-SIP 교환
+
+### Evaluation Harnesses
+
+- [ADR-0043](./ADR-0043-eval-allreduce-harness.md) — Allreduce 평가 하니스 — `tests/sccl/`
+- [ADR-0044](./ADR-0044-eval-gemm-harness.md) — GEMM 평가 하니스 — `scripts/gemm_sweep.py` + `tests/gemm/`
+- [ADR-0054](./ADR-0054-eval-milestone-benches.md) — 마일스톤 평가 bench — 자기완결적 sweep + figure bench
+
+### Bench Module Contract
+
+- [ADR-0045](./ADR-0045-prog-bench-module-contract.md) — Bench Module Contract — registration, dispatch, and authoring
+
+### Kernel-side tl.* API (TLContext)
+
+- [ADR-0046](./ADR-0046-prog-tl-context-contract.md) — TLContext — Kernel-side `tl.*` API Contract
+
+### Memory Allocator Algorithms
+
+- [ADR-0048](./ADR-0048-mem-allocator-algorithms.md) — Memory Allocator Algorithms — VirtualAllocator + PEMemAllocator
+
+### Probe Subcommand
+
+- [ADR-0049](./ADR-0049-ver-probe-subcommand.md) — `kernbench probe` Subcommand — Traffic-Pattern Verification Harness
+
+### Sim-engine Op Log and Memory Store Schemas
+
+- [ADR-0052](./ADR-0052-dev-oplog-memory-store-schemas.md) — OpLog + MemoryStore Schemas — sim_engine internals

docs/adr-proposed/ADR-0061-prog-tl-broadcast.md

@@ -0,0 +1,201 @@
+# ADR-0061: `tl.broadcast` / `tl.repeat` — data-faithful GQA head reuse
+
+## Status
+
+Proposed
+
+> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention).
+>
+> **⚠ Re-evaluated during design iteration — this is now OPTIONAL, not a
+> GQA blocker.** The original framing (below) held that GQA `h_q > h_kv`
+> is blocked by the MemoryStore byte-conservation check on a broadcast
+> view. That is true *only of the baseline's head-packing hack*
+> (`_view(K, (h_q·d, S_kv))`). The correct GQA design (ADR-0060 §5.2)
+> processes **one KV head at a time** and folds the `G` group rows into
+> the matmul **M** dimension — using only byte-conserving reshapes — so it
+> needs **no broadcast op at all**, and `np.matmul` already broadcasts a
+> shared operand in Phase-2 data mode. This ADR therefore drops to a
+> *convenience* primitive: it cleans up additive-mask construction across
+> the `G·T_q` rows and serves general kernels. **Lowest priority of the
+> three supporting ADRs.** The rest of this document is retained as the
+> design for that convenience op.
+
+## Context
+
+### The GQA reuse lever
+
+Grouped-Query Attention amortises K/V loads: `G = H_q / H_kv` query heads
+share one KV head. For Llama3-70B `G = 8`. In a kernel this means a KV
+tile of shape `[S_kv, d]` (for one KV head) must be matmul'd against the
+`G` query rows of its group. Loading K/V **once** and reusing it across
+the group is the decode-time efficiency win (ADR-0060 §5.2).
+
+### Why it does not work today
+
+Pre-ADR-0060 mesh kernels reshaped tensors with a metadata-only `_view`
+helper that rewrote `shape` but kept the original `nbytes`:
+
+```python
+def _view(handle, new_shape):
+    return TensorHandle(..., shape=new_shape, nbytes=handle.nbytes, ...)
+    #                          ^^^ new shape          ^^^ OLD nbytes (unchanged)
+```
+
+This rewrites `shape` but keeps the original `nbytes`. It is purely
+symbolic. It survives **Phase 1 (timing)** because Phase 1 never reads
+tensor bytes. It fails **Phase 2 (`enable_data=True`)** because the
+data executor materialises every op through `MemoryStore.read`, which
+enforces byte conservation
+(`src/kernbench/sim_engine/memory_store.py:67-73`):
+
+```python
+if shape is not None and arr.shape != shape:
+    if arr.nbytes != np.prod(shape) * arr.dtype.itemsize:
+        raise ValueError(f"Shape mismatch: stored {arr.shape} ({arr.nbytes}B) "
+                         f"vs requested {shape} ({...}B)")
+    arr = arr.reshape(shape)
+```
+
+A reshape conserves bytes; a **broadcast does not** (`[S_kv, 1, d]` →
+`[S_kv, G, d]` multiplies the element count by `G`). So the moment a
+kernel tries to view a `h_kv`-headed K as if it had `h_q` heads, the
+stored array has `1/G` of the requested bytes and `read` raises.
+
+The pre-ADR-0060 baseline documented this as a deliberate limitation:
+
+> "v1 uses `h_q == h_kv == 1` to avoid … GQA broadcast view (which is
+> symbolic and does not survive MemoryStore's nbytes check under
+> simulator data execution). Real GQA (`h_q > h_kv`) is deferred."
+
+There is **no** broadcast / repeat / expand primitive in the `tl` API
+today (`src/kernbench/triton_emu/tl_context.py` exposes `trans`,
+`zeros`, `full`, `arange` as metadata-only ops; none replicate data).
+
+## Decision
+
+Add a **data-faithful broadcast op** to the `tl` API and the data
+executor. Unlike `_view`, it updates `nbytes` and, in Phase 2, actually
+replicates the underlying ndarray. The latency it charges models the
+on-chip register/SRAM fan-out, not an HBM re-load (the KV tile is read
+from HBM once; broadcast is an on-chip replication).
+
+### D1. `tl` surface
+
+```python
+def broadcast(self, x: TensorHandle, new_shape: tuple[int, ...]) -> TensorHandle:
+    """Replicate x along size-1 axes to new_shape (numpy broadcast rules).
+
+    new_shape must be broadcast-compatible with x.shape: every axis is
+    either equal to x's, or x's is 1 (the axis being expanded). nbytes of
+    the result reflect new_shape (unlike the metadata-only _view).
+    """
+```
+
+`tl.repeat(x, axis, n)` is offered as a thin convenience wrapper over
+`broadcast` for the common "insert a group axis and expand it to G"
+pattern, but `broadcast` is the primitive.
+
+### D2. Command + op_log
+
+Emit a new `BroadcastCmd` (sibling of `MathCmd`):
+
+```python
+@dataclass(frozen=True)
+class BroadcastCmd:
+    src: TensorHandle
+    out: TensorHandle          # out.nbytes reflects new_shape
+    new_shape: tuple[int, ...]
+    data_op: bool = True
+```
+
+op_log entry: `op_kind="math"`, `op_name="broadcast"` (it runs on the
+MATH/vector engine, like the other element-wise ops). The output handle
+is allocated from PE scratch via the existing `_make_compute_out`
+(`tl_context.py:148-156`), so `out.nbytes` is computed from `new_shape`
+and the byte-conservation check passes in Phase 2.
+
+### D3. Phase 2 data execution
+
+Add `_execute_broadcast` to the data executor
+(`src/kernbench/sim_engine/data_executor.py`):
+
+```python
+def _execute_broadcast(self, op):
+    p = op.params
+    src = self._resolve_read(p["src_space"], p["src_addr"],
+                             p["src_shape"], p["dtype"], op.t_start)
+    result = np.broadcast_to(src, p["new_shape"]).copy()   # .copy() → writable, owns bytes
+    self.store.write(p["dst_space"], p["dst_addr"], result)
+```
+
+`.copy()` is required: `np.broadcast_to` returns a zero-stride view, and
+MemoryStore stores references — a later consumer's nbytes check must see
+a real `prod(new_shape)`-element array.
+
+### D4. Latency model
+
+Broadcast latency = vector-engine cost of writing `prod(new_shape)`
+elements (reuse `pe_math._compute_ns(num_elements)`,
+`src/kernbench/components/builtin/pe_math.py:46-51`). This models the
+on-chip replication into the register file / SRAM that real GQA hardware
+performs. It is **not** an HBM transaction — no `dma_read` is logged,
+consistent with "K/V tile loaded once" (ADR-0060 §5.2).
+
+## Alternatives
+
+### A1. Indexed per-head dot (no broadcast op)
+
+Loop the `G` query heads, each doing its own `tl.dot(Q_g, Kᵀ)` against
+the *same* (re-`ref`'d, not re-loaded) KV tile:
+
+```python
+for g in range(G):
+    s_g = tl.dot(Q[g], K_T)     # G separate GEMMs, KV reused via tl.ref
+```
+
+- Pro: no new primitive; works in Phase 2 today.
+- Con: `G×` more GEMM commands (op_log inflation), no batched-GEMM
+  modelling, and the per-call overhead is charged `G` times. It also
+  obscures the "single batched matmul over the group" the hardware does.
+  Rejected as the primary path; acceptable as a fallback if the batched
+  2D dot's shape handling proves awkward.
+
+### A2. Keep `_view`, relax the nbytes check
+
+Make `MemoryStore.read` tolerate broadcast (stride-0) shapes. Rejected:
+it silently breaks byte conservation for *all* readers and would mask
+genuine shape bugs elsewhere. The check is a correctness invariant
+(ADR-0052), not an obstacle to route around.
+
+### A3. Materialise the broadcast at deploy time
+
+Store K/V already expanded to `h_q` heads in HBM. Rejected: defeats the
+entire point of GQA (it `G×`'s KV-cache HBM footprint and bandwidth).
+
+## Consequences
+
+### Positive
+- Unblocks real GQA (`h_q > h_kv`) in Phase 2 — the headline blocker for
+  ADR-0060 and the deferred milestone work.
+- A general, reusable primitive (any kernel needing numpy-style
+  broadcast benefits, not just attention).
+- Byte conservation stays intact; no weakening of MemoryStore.
+
+### Negative
+- One new command type + one data-executor handler + one `tl` method.
+- Phase 2 `.copy()` allocates the expanded array in the store; for very
+  large groups this is real memory. Bounded by tile size in the tiled
+  design (ADR-0060), so acceptable.
+
+## Test Requirements
+
+1. **Phase 1 timing**: `tl.broadcast([S,1,d] → [S,G,d])` logs one
+   `math/broadcast` op with `num_elements == S*G*d`; no `dma_read`.
+2. **Phase 2 data**: broadcasting a known array yields the numpy
+   `broadcast_to(...).copy()` result; a downstream `tl.dot` consuming it
+   passes the MemoryStore nbytes check (the exact case that fails today).
+3. **GQA end-to-end**: re-run a reduced `milestone-gqa-headline` panel
+   with `h_q=8, h_kv=1` under `enable_data=True`; it must complete (today
+   it raises `ValueError: Shape mismatch`).
+4. **Broadcast-incompatible shape** (e.g. axis size 3 → 8) raises a
+   clear error at emit time, not deep in Phase 2.

docs/adr-proposed/DDD-0060-gqa-fused-attention-detailed-design.md

@@ -0,0 +1,355 @@
+# Detailed Design Document — AHBM GQA Fused Attention (ADR-0060)
+
+**Status:** Draft (companion to ADR-0060, Proposed)
+**Scope:** implementation-ready plan for the **two** GQA FlashAttention
+kernels (decode=reduce, prefill=ring) on AHBM in kernbench.
+**Audience:** reviewer resuming with *"GQA 검토 시작하자"*. Read **ADR-0060
+first** — it is now the authoritative design record (TL;DR has full
+pseudocode for both kernels). This DDD is the *how to build it*: file plan,
+phase plan, helper signatures, tests. It does **not** re-derive the design;
+it points to ADR-0060 sections.
+
+> **Alignment note (this revision).** ADR-0060 moved to a **composite
+> hybrid + hierarchical CUBE-Group SP** design with **two kernels**. This
+> DDD is rewritten to match. The earlier DDD (greenlet-primitive,
+> `tl.load_async`, single unified kernel) is superseded.
+
+---
+
+## 1. Goal and success criteria
+
+**Goal:** two *efficient* GQA fused-attention kernels on AHBM in kernbench —
+**decode+SP** (head-replicated, KV static shard, 2-level reduce) and
+**prefill+SP** (1 Q head per CUBE, Ring KV) — at realistic scale, in both
+timing (`enable_data=False`) and data (`enable_data=True`) modes.
+
+**Success criteria:**
+1. Correct: kernel output matches a numpy FlashAttention reference within fp
+   tolerance, with **real GQA** (`H_q=64, H_kv=8, G=8`).
+2. Efficient — levers observable in op_log (ADR-0060 §11):
+   GQA K/V-load amortisation; decode 2-level reduce (`⌈log₂P⌉` + center-mesh
+   over `C`, not `C·P−1`); prefill ring (no reduce, KV rotates); load/compute
+   overlap (lazy `tl.load` + composite streaming); composite GEMM offload.
+3. Scales: context length not capped by the 1 MiB scratch bump allocator.
+4. Deterministic and SPEC-compliant (all latency from modelled events).
+
+**Non-goals:** RoPE / QKV projection / KV-cache writes (upstream
+`qkv_rope`); output projection (downstream); cross-**SIP** head split
+(a head stays in one SIP, ADR-0060 §0); cycle-accurate microarchitecture
+(SPEC §5); a bespoke "flash-composite" command kind (ADR-0060 §8 item 4).
+
+---
+
+## 2. Where the design sits in the current code
+
+```
+runtime_api (host)
+  RuntimeContext.zeros/empty/launch    context.py  (tensor deploy, DPPolicy(+cube_start), launch)
+        │  KernelLaunchMsg
+        ▼
+sim_engine
+  GraphEngine(enable_data=…)           engine.py
+  DataExecutor (Phase 2)               data_executor.py     ← _execute_broadcast (ADR-0061, optional)
+  MemoryStore                          memory_store.py
+        │
+        ▼
+components / PE pipeline (per PE)
+  PE_CPU → PE_SCHEDULER → PE_{DMA,GEMM,MATH,IPCQ,TCM}
+  greenlet kernel ↔ SimPy              kernel_runner.py, pe_cpu.py
+  CompositeCmd → tile plan → engines   pe_scheduler.py (ADR-0014 D6)
+        │
+        ▼
+tl programming model                   tl_context.py
+  load (→ lazy, ADR-0062) / store / composite / dot / max/sum/exp/maximum
+  send/recv/recv_async / scratch_scope (ADR-0063) / broadcast (ADR-0061, opt)
+```
+
+Both kernels are **bench-side Python** (`src/kernbench/benches/`) run as
+greenlet kernels emitting `tl.*` ops; the **GEMMs are `tl.composite`**
+(scheduler-managed tiling + K/V DMA streaming, existing `CompositeCmd` — no
+new command kind), the softmax merge + reduction/ring stay kernel-level
+(ADR-0060 §1). The remote work already added building blocks:
+`DPPolicy.cube_start`, `_attention_mesh_mlo_2d.py` (2D cube reduce —
+currently AllReduce, to become reduce-to-root), `topologies/llama70b_4sip.yaml`.
+
+---
+
+## 3. File plan
+
+### 3.1 New / modified production files
+
+| File | Change | ADR |
+|---|---|---|
+| `src/kernbench/triton_emu/tl_context.py` | `tl.load` → **lazy** (non-blocking + auto-wait on first use); `scratch_scope` ctx-mgr; `broadcast` (optional) | 0062/0063/0061 |
+| `src/kernbench/components/builtin/pe_dma.py` | non-blocking read path: post DMA, resolve on a scheduled event (the `recv_async` pattern for loads) | 0062 |
+| `src/kernbench/triton_emu/kernel_runner.py` | auto-wait: yield a handle's pending load event at first consume; lazy-load dispatch arm | 0062 |
+| `src/kernbench/sim_engine/data_executor.py` | `_execute_broadcast` (optional) | 0061 |
+| `src/kernbench/components/builtin/pe_cpu.py` + cost table | per-op-type CPU issue cost (replaces uniform `dispatch_cycles`) | **0064** (separate) |
+
+No new GEMM command kind: the two GEMMs use the **existing** `CompositeCmd`.
+
+### 3.2 Kernel + bench files
+
+| File | Role |
+|---|---|
+| `src/kernbench/benches/_gqa_decode.py` | decode+SP kernel — evolve `_attention_mesh_mlo`/`_mlo_2d` to composite-hybrid + **2-level reduce-to-root** (Level-2 PE tree + Level-1 center-mesh; currently 2D AllReduce) |
+| `src/kernbench/benches/_gqa_prefill.py` | prefill+SP kernel — evolve `_attention_mesh_kv` to **head-parallel (1 Q head/CUBE) Ring KV** + composite-hybrid |
+| `src/kernbench/benches/_gqa_helpers.py` | `head_of_group`, `valid_len_2level`, `init_running`, tree/center-mesh topology, `hierarchical_reduce_and_store`, ring helpers, causal `block_*` predicates + mask builders |
+| extend `src/kernbench/benches/milestone_gqa_llama70b.py` | headline panels: real `G`, `C=G=8`, contiguous KV, 4-SIP |
+| `topologies/llama70b_4sip.yaml` | **exists** (remote) — 4 SIP × 16-CUBE 4×4 × 8 PE |
+
+### 3.3 New tests
+
+| File | Covers |
+|---|---|
+| `tests/attention/test_gqa_decode.py` | decode reduce: GQA reuse (dma_read ⟂ `G`), 2-level reduce step count, root-only output, long context |
+| `tests/attention/test_gqa_prefill.py` | prefill ring: no `(m,ℓ,O)` reduce, `C` KV rotations, causal step-skip, per-CUBE distributed output |
+| `tests/test_tl_lazy_load.py` | ADR-0062 unit (overlap, auto-wait correctness, op_log parity) |
+| `tests/test_tl_scratch_scope.py` | ADR-0063 unit |
+| `tests/test_tl_broadcast.py` | ADR-0061 unit (optional) |
+| `tests/test_dppolicy_cube_start.py` | **exists** (remote) — sub-mesh placement |
+
+---
+
+## 4. Data model and launch contract
+
+### 4.1 Tensor placement — CUBE Group, 2-level SP, contiguous KV
+
+A `CUBE Group` is `C` CUBEs in one SIP owning one KV head (ADR-0060 §0).
+Placement via `DPPolicy(..., cube_start=...)`:
+
+| Tensor | Decode+SP | Prefill+SP (`C=G`) |
+|---|---|---|
+| `Q` | `replicate` (all `G` heads, M-folded) over `C×P` ranks | **head-parallel**: CUBE `i` holds Q head `i` |
+| `K`,`V` | sequence-sharded **two axes** (`cube` row_wise over `C`, `pe` row_wise over `P`) | sequence-sharded over `C` CUBEs (the ring blocks) |
+| `O` | reduced to CUBE-Group root (all `G` heads) | per-CUBE (one head each), distributed |
+
+**KV cache is *contiguous* `C×P` position blocks** (rank `r` = `[r·B,(r+1)·B)`),
+**shared by both kernels** — prefill writes it (via `qkv_rope`), decode reads
++ extends it, no reshard. Contiguous is required for prefill causal skip
+(ADR-0060 §2.1). `cube_start` offsets the group's CUBE sub-mesh within the
+4×4 SIP mesh.
+
+**Driver SP-enable threshold (fallback).** Contiguous sharding under-utilizes
+ranks for short/early decode (only frontier ranks hold data). So the driver
+**enables SP only past a context-length threshold** where KV-sweep time
+dominates reduction + placement imbalance; below it, fall back to a smaller
+`C` (or `C=P=1`, single rank, no reduce). The threshold is a sweep item
+(P7 / §9), not a fixed constant.
+
+`--device` enumerates **SIPs**; one SIP-device runs its 2 CUBE Groups
+(2 KV heads); the head is picked by CUBE coordinate (`head_of_group`), not an
+in-kernel loop (ADR-0060 §0).
+
+### 4.2 Launch signatures
+
+```python
+# decode+SP
+ctx.launch(f"{panel}_gqa_decode", gqa_decode_sp,
+           q, k, v, o, counter, start_pe, start_cube, C, P, softmax_scale)
+# prefill+SP
+ctx.launch(f"{panel}_gqa_prefill", gqa_prefill_sp,
+           q, k, v, o, T_q, S_kv_local, d, C, softmax_scale, q_block, cube_start)
+```
+
+Full per-kernel I/O contract: ADR-0060 §0.5.3/§0.5.4 (note the **output head
+distribution differs** — decode root-gathered, prefill distributed).
+
+### 4.3 GQA reuse — fold `G` into the matmul M dim (decode)
+
+Decode replicates Q and **M-folds** the `G` query heads into the GEMM row
+dim so one `Q·Kᵀ` serves all `G` heads sharing one `K` (ADR-0060 §0, §5.2):
+`Q` group `[G, T_q, d]` → `[G·T_q, d]` (byte-conserving `_view`); the
+composite carries `m = G·T_q` (timing counts all `G` rows). Prefill does
+**not** M-fold (1 head/CUBE). Caveat: `tl.trans` is reshape-not-transpose →
+store `K` pre-transposed `[d, S/(C·P)]` (ADR-0060 §3, §B).
+
+---
+
+## 5. Kernel design
+
+**The full pseudocode for both kernels (and the 3 decode CPU-pipelining
+variants) lives in ADR-0060 TL;DR + §5.** Implementation notes only here.
+
+### 5.1 Decode (reduce) — `_gqa_decode.py`
+- Inner tile = §3 composite-hybrid: `Sj = composite(q_g, Kⱼ)·scale` → softmax
+  MATH → `Oj = composite(P, Vⱼ)` → running merge (ADR-0060 §3).
+- **Ship opt3** (software pipelining: issue next tile's `Q·Kᵀ` before this
+  tile's softmax; `Sj` in a persistent double buffer) — removes the
+  GEMM-engine bubble, no new command kind (ADR-0060 §5.6). opt1 is the naïve
+  baseline; opt2 (`ex_composite`) waits on ADR-0064.
+- Combine = **`hierarchical_reduce_and_store`** (ADR-0060 §4): Level-2 PE
+  reduce-to-root tree (intra-CUBE, the cube's KV slice further split across
+  its `P` PEs — decision (a)) → Level-1 center-root CUBE-mesh reduce
+  (intra-CUBE-Group). Data-driven, level-pipelined, root-only.
+
+### 5.2 Prefill (ring) — `_gqa_prefill.py`
+- Head-parallel: CUBE `i` owns Q head `i` + KV slice `i`. **No reduce** (each
+  CUBE outputs a distinct head). Ring rotates KV blocks; GQA reuse via the
+  rotation (ADR-0060 §5.5).
+- **K/V are `tl.load`'d TCM handles** (not `tl.ref`) so the ring can
+  `tl.send`/`recv_async` them; `K` pre-transposed `[d, S/C]`. Receive buffers
+  **ping-pong** (persistent arena) — recv into a *separate* buffer while
+  computing/sending the current one (do not clobber).
+- Causal step-skip + boundary mask; `recv_async` overlaps next block's
+  receive with current compute.
+- Within a CUBE: tile the Q head's `T_q` rows across `P` PEs (disjoint output
+  rows → no intra-CUBE reduce); fall back to KV-block split only if `T_q < P`
+  (ADR-0060 §B decode-split item 3).
+
+### 5.3 Shared helpers — `_gqa_helpers.py`
+`head_of_group(cube_id) = cube_id // C`; `init_running(...)` → persistent
+`(m=-inf, ℓ=0, O=0)`; `valid_len_2level(...)` (contiguous block extent);
+`hierarchical_reduce_and_store(...)`; tree/center-mesh child/parent dirs;
+`block_all_future`/`block_partial` + `causal_mask`.
+
+---
+
+## 6. Supporting primitives — integration points
+
+### 6.1 ADR-0062 lazy `tl.load` (efficiency — overlap)
+- `tl.load` issues `DmaReadCmd` non-blocking, returns a handle with a pending
+  event; the runtime **auto-inserts the wait at first consume** (generalises
+  the `recv_async`/wait pattern, `kernel_runner.py:248-285`).
+- `pe_dma.py`: non-blocking read; DMA occupies the (capacity-1) read channel,
+  overlaps compute on PE_GEMM/PE_MATH.
+- **Global** semantics change → existing goldens regenerate (ADR-0062 D3).
+- op_log unchanged (`memory/dma_read`) ⇒ `dma_read_count` stable.
+
+### 6.2 ADR-0063 `tl.scratch_scope` (scale)
+- ctx-mgr saving/restoring `_scratch_cursor`. Persistent arena for `(m,ℓ,O)`
+  (+ decode opt3 `Sj` double buffer, + prefill ring ping-pong buffers, +
+  in-flight lazy-load buffers) lives **outside** the scope. Removes `S=16`.
+
+### 6.3 ADR-0061 `tl.broadcast` (optional)
+- Not on the GQA critical path (M-fold gives reuse). Convenience for additive
+  mask construction. Lowest priority.
+
+### 6.4 ADR-0064 per-op-type CPU issue cost (separate ADR)
+- Replaces uniform `dispatch_cycles=0`. Makes the hybrid's CPU-saturation
+  win (and the decode opt2 `ex_composite` fewer-issues win) **measurable**.
+  Land before claiming hybrid latency wins (ADR-0060 §9).
+
+---
+
+## 7. Phased implementation plan (test-first per CLAUDE.md)
+
+Each phase is an independent Phase-1→Phase-2 cycle; each keeps the baseline
+green.
+
+| Phase | Deliverable | Gate |
+|---|---|---|
+| **P1** | Real GQA, decode, composite-hybrid M-fold; one-shot, no tiling. **No new primitive** (existing `CompositeCmd`). | data-mode completes `h_q=G·h_kv`; K/V `dma_read_count` ⟂ `G`; composite tile plan `m=G·T_q` |
+| **P2** | Decode **2-level reduce-to-root** (Level-2 PE tree + Level-1 center-mesh), replacing the 2D AllReduce | reduce rounds = `⌈log₂P⌉`+center-mesh; **root-only** output; Level-1 on CUBE NOC, Level-2 on PE IPCQ |
+| **P3** | **ADR-0063** `scratch_scope` + tiled sweep | long-`S` decode completes (today caps S=16); scratch O(1) |
+| **P4** | **ADR-0062** lazy `tl.load` + composite K/V streaming | tiled-sweep latency < serial `Σ(load+compute)`; goldens regenerated |
+| **P5** | Decode **opt3** software pipelining | GEMM engine non-idle across tiles (structural / op_log) |
+| **P6** | Prefill **head-parallel Ring KV** + causal step-skip | `C` KV rotations, no `(m,ℓ,O)` reduce, per-CUBE distributed `O`, causal skip ≈½ |
+| **P7** | Contiguous shared KV layout + 4-SIP topology + headline milestone panels (real `G`, `C=G`) | shared layout (no reshard); headline sweep.json rows |
+| **P8** *(opt)* | **ADR-0064** cost model + decode **opt2** `ex_composite`; **ADR-0061** broadcast | cost-model goldens; opt2 fewer-issues measurable |
+
+P1 delivers "GQA runs" on the composite path. P3 is the key *scale* feature;
+P4 the overlap lever; P2/P6 the SP algorithms.
+
+---
+
+## 8. Verification plan (concrete)
+
+Mirrors ADR-0060 §11; grounded in SPEC R2/R5, ADR-0023/0025, ADR-0046, ADR-0054.
+
+**Runs + numeric (`enable_data=True`):**
+- decode & prefill complete for `(G∈{1,8}, T_q∈{1,16}, S, C∈{1,8}, P∈{1,8})`
+  without byte-conservation error (the `G=8` cases baseline cannot express).
+- *Numeric parity (secondary):* `O ≈ numpy FlashAttention` for
+  symmetric/identity inputs (reshape-as-transpose exact); full asymmetric
+  parity gated on a real `tl.transpose` (deferred).
+
+**Levers (op_log):**
+- GQA amortisation: K/V `dma_read_count` ⟂ `G`; compute scales with `G`.
+- Decode reduce: rounds = `⌈log₂P⌉` + center-mesh (not `C·P−1`); result at
+  one rank; Level-1=CUBE NOC, Level-2=PE IPCQ.
+- Prefill ring: `C` KV rotations, **no** reduce, each CUBE writes a distinct
+  head.
+- Causal skip: GEMM/step count = lower-triangular.
+- Overlap: tiled-sweep latency < `Σ(load+compute)`.
+- Scratch: peak cursor bounded by one tile.
+
+**Invariants:** determinism (identical op_log + latency); every routed
+request latency > 0 (SPEC §0.1).
+
+---
+
+## 9. Performance model (expected)
+
+Per-rank decode (KV-load-bound), composite-streamed + overlapped:
+```
+t_decode_rank ≈ max( DMA(K+V over S/(C·P) tiles),  compute(QKᵀ+PV, ×G amortised) )
+              + t_reduce( ⌈log₂P⌉ intra-CUBE  +  center-mesh over C inter-CUBE )
+```
+Prefill (ring): `t ≈ C · max( recv(KV block), compute(QKᵀ+PV over block) )`
+with causal step-skip removing ≈half the steps; no reduce term.
+
+Levers: GQA reuse → KV bytes `H_kv·S·d` not `H_q·S·d`; overlap → `max` not
+sum; decode reduce `⌈log₂P⌉+center-mesh` vs baseline `C·P−1`; prefill ring
+trades the reduce for KV rotation (good when `O` is big). The CPU-saturation
+win (composite offload) is *measurable* only with ADR-0064.
+
+---
+
+## 10. Open items (status)
+
+Most original "Open Decisions" are now **resolved** in ADR-0060; the live
+review items live in **ADR-0060 §B** (three groups: hybrid pivot,
+hierarchical CUBE-Group SP, decode/prefill split). Key resolved choices:
+
+| Was open | Now |
+|---|---|
+| greenlet vs composite | **composite hybrid** (GEMMs→composite, merge→kernel) |
+| one kernel vs two | **two kernels** (decode reduce / prefill ring) |
+| GQA matmul shape | **M-fold** (decode); head-parallel (prefill) |
+| reduction topology | decode **2-level reduce-to-root**; prefill **no reduce (ring)** |
+| `tl.load_async` | **lazy `tl.load`** (ADR-0062, redefined) |
+| KV placement | **contiguous `C×P`**, shared prefill/decode (ADR-0060 §2.1) |
+
+Still to confirm (ADR-0060 §B): DDD↔impl reconcile (`_attention_mesh_mlo_2d`
+AllReduce → reduce-to-root); `head_of_group` sub-mesh partition; `C=G`
+coupling; short-context KV balance; ADR-0064 calibration. Pre-existing:
+ghost ADRs 0055–0059 (separate backfill); `bf16→f16` proxy.
+
+---
+
+## 11. Risks
+
+| Risk | Likelihood | Mitigation |
+|---|---|---|
+| Numeric parity blocked by trans-as-reshape | med | structural-first verify; pre-transpose K (ADR-0060 §B item 2) |
+| 2-level reduce pairs not 1-hop on the mesh | med | verify SFR neighbour table; center-root sub-mesh partition (§B) |
+| scratch_scope use-after-scope (Sj / ring / lazy-load buffers) | med | persistent-arena discipline; ping-pong buffers (§5.2) |
+| prefill ring buffer clobbered while in flight | med | recv into the *other* ping-pong buffer (§5.2) |
+| decode opt2 `acc` read before composites drain | med | `tl.wait()` after the loop (ADR-0060 TL;DR opt2) |
+| impl drift (`_attention_mesh_mlo_2d` is AllReduce, Q-replicated) | med | reconcile to reduce-to-root; add prefill ring (§B) |
+| CPU-saturation win invisible | exp | ADR-0064 cost model (P8) |
+
+---
+
+## 12. Glossary & references
+
+- **GQA / G** — `G = H_q/H_kv` query heads share one KV head. Llama3-70B: `G=8`.
+- **CUBE Group** — `C` CUBEs in one SIP owning one KV head (ADR-0060 §0).
+- **2-level SP** — Level-1 inter-CUBE (over `C`) × Level-2 intra-CUBE PE
+  (over `P`); ranks = `C·P`.
+- **Composite hybrid** — GEMMs via `tl.composite` (scheduler), softmax merge
+  + reduction in the kernel (ADR-0060 §1).
+- **M-fold** — stack the `G` query heads into the GEMM M (row) dim so one
+  `Q·Kᵀ` does all heads sharing one `K`.
+- **FlashAttention / FlashDecoding / Ring Attention** — ADR-0060 lineage.
+
+Source anchors: `tl_context.py` (tl API), `pe_scheduler.py:104-143`
+(composite tile plan), `pe_dma.py:45,89` (read channel, drain), `pe_cpu.py`
+(dispatch_cycles), `lrab_hierarchical_allreduce.py` (center-root pattern),
+`_attention_mesh_mlo_2d.py` / `_attention_mesh_kv.py` (impl to evolve),
+`DPPolicy.cube_start`, `topologies/llama70b_4sip.yaml`.
+
+ADRs: **0060** (this design), **0062** (lazy load), **0063** (scratch scope),
+**0061** (broadcast, optional), **0064** (CPU issue cost model); related
+accepted **0014/0017/0020/0023/0025/0046/0054**.

docs/adr/ADR-0026-par-dppolicy-intra-device.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Accepted (Revision 5 — Phase 2 landed 2026-04-14, 523 passed + 1 strict xfail)
+Accepted (Revision 6 — cube_start added 2026-06-04; Revision 5 landed 2026-04-14)
 
 ## Context
 
@@ -31,9 +31,11 @@ class DPPolicy:
     pe: Literal["replicate", "column_wise", "row_wise"] = "replicate"
     num_pes: int | None = None
     num_cubes: int | None = None
+    cube_start: int = 0   # Revision 6 — first cube in SIP; see D8
 ```
 
 Removed fields: `sip`, `num_sips`.
+Added fields (Revision 6): `cube_start` — see D8.
 
 ### D2. `ShardSpec` — structural (sip, cube, pe) coordinates, `pe_index` fully removed
 
@@ -120,7 +122,7 @@ def resolve_dp_policy(
         for ls in local_shards:
             all_shards.append(ShardSpec(
                 sip=target_sip,                              # from caller (current_device)
-                cube=cube_id,                     # local within SIP
+                cube=policy.cube_start + cube_id,            # Rev 6: shifted by D8 cube_start; default 0 = local within SIP
                 pe=ls.local_pe,                              # local within cube (explicit name)
                 offset_bytes=cube_offset + ls.offset_bytes,
                 nbytes=ls.nbytes,
@@ -239,6 +241,53 @@ property: code that expected a global flat could otherwise silently
 receive a SIP-local result and index incorrectly — that possibility is
 eliminated.
 
+### D8. Add `cube_start: int = 0` for disjoint cube sub-meshes within a SIP
+
+**Revision 6 addition (2026-06-04)**.
+
+**Purpose**: address a disjoint cube sub-mesh within one SIP (e.g.
+cubes 8..15 alongside cubes 0..7). Required by the GQA Llama-70B
+8-KV-group headline target — two 2×4 KV-groups per SIP × 4 SIPs = 64
+cubes — where the second KV-group on each SIP must land on cubes 8..15
+instead of the default 0..7.
+
+**Semantics**: `resolve_dp_policy` returns `ShardSpec` with
+`cube = policy.cube_start + cube_id`, where `cube_id` iterates
+`0..num_cubes-1` within the launch. Selected cubes lie in
+`[cube_start, cube_start + num_cubes)` within the target SIP.
+
+**Default**: `cube_start = 0` preserves every existing call site
+bit-for-bit (`ShardSpec.cube ∈ [0, num_cubes)` as before). The CCL
+milestone bench's full-SIP `DPPolicy(num_cubes=16)` continues to
+produce cubes 0..15.
+
+**Constraint**: intra-device invariant preserved.
+`cube_start ∈ [0, cubes_per_sip)` and
+`cube_start + num_cubes ≤ cubes_per_sip`. SIP boundary crossing remains
+the job of `ahbm.set_device(rank)` (ADR-0024).
+
+**Why scalar** (not 2D `cube_mesh_origin` or arbitrary `cube_ids` list):
+consumer kernels (e.g. `_attention_mesh_mlo_2d`) assume row-major
+contiguous cubes; scalar `cube_start` pairs naturally with `num_cubes`
+(range = `[start, start + count)`); zero migration churn at default 0.
+More general designs can be added on top later if a non-contiguous use
+case appears.
+
+**Backward compatibility**: additive change with default value. No
+existing call site is forced to change. The "breaking change" stance
+in D7 applies only to the original sip/num_sips removal — `cube_start`
+does NOT break existing callers.
+
+**Kernel-side note**: kernbench's `tl.program_id(axis=1)` returns the
+physical cube id (ADR-0022), not a launch-local rank. Kernels that
+derive ring positions from `program_id(axis=1)` must subtract
+`cube_start` to recover launch-local rank when `cube_start > 0` —
+otherwise they compute out-of-bounds sub-mesh positions and deadlock.
+See `_attention_mesh_mlo_2d.py` for the reference pattern. A future
+revision of ADR-0022 could move `program_id(axis=1)` to launch-local
+rank semantics; that change would let kernels drop the explicit
+subtraction.
+
 ## Dependencies
 
 - **ADR-0024** (launcher): `set_device(rank)` and current-device scoping

docs/adr/ADR-0043-eval-allreduce-harness.md

@@ -7,6 +7,11 @@ Accepted
 Documents the `tests/sccl/` evaluation harness; verified against the
 implementation (constants, file set, and sweep dimensions cross-checked).
 
+**Amended by ADR-0054**: the driver core, sweeps, and renderers moved into
+the `milestone-1h-ccl` bench (single home); `tests/sccl/_allreduce_helpers.py`
+now re-exports from it (keeping the pytest-only param builders +
+`_run_distributed` wrapper local). The figure tests are unchanged.
+
 ## Context
 
 ADR-0032 defines the intercube all-reduce *algorithm*; ADR-0023/0024/0027

docs/adr/ADR-0044-eval-gemm-harness.md

@@ -9,6 +9,12 @@ implementation (constants, tile sizes, figure set, and the script↔test
 split cross-checked). The D5/D6 caveats are recorded limitations, not
 inaccuracies.
 
+**Amended by ADR-0054**: the sweep + renderers moved into the
+`milestone-1h-gemm` bench (single home); `scripts/gemm_sweep.py` and
+`tests/gemm/` now re-export from it. D1/D2's "data generation stays a manual
+script / heavy work is opt-in" is superseded by the eval-bench pattern (one
+bench regenerates everything; `MILESTONE_FAST=1` reuses the committed JSON).
+
 ## Context
 
 ADR-0014 (PE pipeline) and ADR-0042 (tile-plan generators) define the GEMM

docs/adr/ADR-0045-prog-bench-module-contract.md

@@ -0,0 +1,295 @@
+# ADR-0045: Bench Module Contract — registration, dispatch, and authoring
+
+## Status
+
+Accepted (2026-05-21).
+
+Unifies the `src/kernbench/benches/` registration mechanism (@bench), the
+CLI dispatch path (`kernbench run/list`), and the contract a new bench
+module must follow. ADR-0010 (CLI surface) specifies the `kernbench
+list/run` interface, but **how benches are registered and what signature
+they must follow** had no ADR-level coverage.
+
+**Extended by ADR-0054**: D5's single-config rule gains a third pattern —
+the *eval bench* (e.g. `milestone-1h-*`) drives many configs, builds its
+own per-config engines, and submits a sentinel tensor to satisfy D4.
+
+## First action
+
+When `kernbench.benches` is imported, `__init__.py` immediately calls
+`_eager_import_and_audit(__path__, __name__)`. Its first action is to
+enumerate every sibling module in the package directory via
+`pkgutil.iter_modules(__path__)` and **eagerly import** each one via
+`importlib.import_module(...)` — except modules matching either:
+
+- name `registry` (the infrastructure module itself), or
+- name starting with `_` (helper modules).
+
+At import time, each `@bench(name=..., description=...)` decorator inside
+the imported module runs, appending `(name, description, fn)` to
+`_PENDING` and adding `fn.__module__` to `_REGISTERED_MODULES`.
+
+Once imports finish, `_audit_modules(imported, _REGISTERED_MODULES)`
+runs; if any imported module did not invoke `@bench` at least once, it
+raises `RuntimeError("Bench module(s) missing @bench decorator: ...")`
+immediately. At this point indices are still unassigned — the first call
+to `list_all()` / `resolve(...)` triggers `_finalize()`, which sorts
+`_PENDING` alphabetically by name and assigns 1-based indices.
+
+In short, **the bench infrastructure's first act is "eagerly import
+every non-helper module in the package and audit that each one
+registered at least one bench"**.
+
+## Context
+
+`src/kernbench/benches/` currently holds 8 bench modules (`ccl_allreduce`,
+`gemm_single_pe`, `gpt3_qkv`, `ipcq_allreduce`, `matmul_composite`,
+`qkv_gemm`, `qkv_gemm_multi_pe`, `va_offset_verify`). Every bench follows
+the same unified flow:
+
+```
+kernbench run --topology <T> --bench <N>
+   ↓
+cli/main.py::cmd_run
+   ↓  resolve_topology(T)  + resolve(N)  + resolve_device(device_arg)
+   ↓
+runtime_api/bench_runner.py::run_bench(topology, bench_fn, device, engine_factory)
+   ↓  engine_factory(topology, device) → GraphEngine
+   ↓  RuntimeContext(engine, target_device, correlation_id, spec)
+   ↓
+bench_fn(ctx)        ← invokes the bench's run(torch)
+   ↓  ctx.empty/zeros/from_numpy/launch/distributed.* etc. submit work
+   ↓
+ctx.wait_all()       ← drains any outstanding handles
+   ↓
+BenchResult(completion, correlation_id, trace, traces, engine)
+```
+
+ADR-0010 covers only the CLI surface (`run/list/probe/web`); ADR-0007
+covers only the runtime API ↔ sim_engine boundary. The question "what
+shape must a new bench file take?" had to be answered by grepping the
+codebase. As a result:
+
+- The @bench decorator contract (kebab-case name, non-empty description)
+  lived only in the source.
+- The bench function signature (`def run(torch)`) was a de-facto
+  convention enforced by the CLI dispatcher calling `spec.run`.
+- New bench authors learned the "helpers must use `_` prefix" rule only
+  after seeing the audit's RuntimeError.
+- The single-device convention (CLAUDE.md Part 2 CLI Semantics) and its
+  interaction with multi-SIP CCL benches was ambiguous for bench
+  authors.
+
+This ADR consolidates all of it in one place.
+
+## Decision
+
+### D1. @bench decorator contract
+
+```python
+from kernbench.benches.registry import bench
+
+@bench(name="my-bench", description="Short, complete-sentence description.")
+def run(torch):
+    ...
+```
+
+- `name`: kebab-case string matching `^[a-z][a-z0-9]*(-[a-z0-9]+)*$`.
+  Lowercase letters, digits, and dashes only; underscores forbidden;
+  must start with a letter.
+- `description`: non-empty string (stripped length > 0). Displayed
+  verbatim by `kernbench list`.
+- The decorator **returns the function unchanged** — direct invocation
+  is fine. Its only side effect is appending to `_PENDING`.
+
+Violations of the first two rules raise `ValueError` at decoration time.
+Duplicate names are caught at `_finalize()` with
+`RuntimeError("duplicate bench name: ...")`.
+
+### D2. Module-file convention
+
+Every `src/kernbench/benches/<slug>.py` must be one of:
+
+- **A bench module**: at top-level import, `@bench(...)` runs at least
+  once to register at least one bench.
+- **A helper module**: the filename starts with `_` (e.g.,
+  `_shared_helpers.py`). `iter_modules` skips it.
+
+The audit (`_audit_modules`) rejects any non-helper that fails to call
+`@bench`. Intended consequence: dropping a new file into `benches/`
+automatically registers its benches, and helper modules are clearly
+flagged by their filename prefix alone.
+
+### D3. The bench function signature is `def run(torch)`
+
+The decorator does not enforce a function name, but **CLI dispatch calls
+`spec_entry.run`** (the decorated callable). The convention is therefore:
+
+- Function name: `run`. Other names work, but always use `run` for
+  readability and grep-ability.
+- Argument: a single positional `torch`. In practice this is a
+  `RuntimeContext` instance exposing PyTorch-style namespaces
+  (zeros/empty/launch/distributed/...) — see ADR-0024 D3.
+- Return value: any (`Any`). `run_bench` ignores it and tracks
+  completion via `ctx.handles()` / `engine.get_completion()`.
+
+The `torch` name imitates a PyTorch-compatible idiom; the actual PyTorch
+module is not passed in (aligned with ADR-0024's "rank = SIP" launcher
+convention).
+
+### D4. A bench must submit at least once
+
+If `ctx.handles()` is empty after the bench returns, `run_bench` reports
+`BenchResult.completion = ok=False, error_code="NO_REQUESTS"`. So a
+meaningful bench must invoke at least one of:
+
+- Tensor-creation APIs: `torch.zeros(...)`, `torch.empty(...)` — these
+  internally submit `MmuMapMsg` and (for zeros) `MemoryWriteMsg`.
+- Kernel-launch API: `torch.launch(name, fn, *args)` — submits per-SIP
+  `KernelLaunchMsg`.
+- (Exception) Empty placeholder benches: e.g.,
+  `ipcq_allreduce.py`'s `print(...)`-only stub will receive a
+  NO_REQUESTS result. CI is expected to recognize and handle placeholder
+  benches specially.
+
+### D5. Single-device convention + multi-SIP exception (ADR-0024/0027)
+
+CLAUDE.md Part 2 CLI Semantics' **"benchmarks MUST remain
+single-device"** rule is interpreted as follows:
+
+- **Standard bench (single-SIP use)**: define tensor placement with
+  `dp = DPPolicy(...)` and launch with `torch.launch(...)`. The SIP
+  index is chosen by `--device` (CLI's responsibility).
+- **CCL bench (multi-SIP use)**: as an exception, use
+  `torch.distributed.init_process_group(backend="ahbm")` plus
+  `torch.multiprocessing.spawn(_worker, ..., nprocs=ws)` for the
+  rank = SIP pattern (ADR-0024 D3). `--device` is ignored (or treated
+  as `all`); each spawned worker calls `torch.ahbm.set_device(rank)` to
+  bind to its SIP.
+
+Multi-device patterns outside these two (e.g., one bench function
+launching across multiple SIPs in the same process) are forbidden by
+this ADR. Even with `--device all`, the bench runs once; multi-SIP use
+inside that single run must follow D5's second pattern.
+
+### D6. Name/index resolution (`resolve`)
+
+`resolve(identifier: str)` returns a BenchSpec via:
+
+1. If `identifier.isdigit()`: convert to int and find the spec where
+   `index ==` that value. If none, `ValueError("No bench with index
+   ...")`.
+2. If `identifier in _REGISTRY`: direct lookup.
+3. Otherwise: `ValueError("Unknown bench ...")`.
+
+Empty or whitespace-only identifiers raise `ValueError("bench
+identifier must be a non-empty string.")`.
+
+The CLI passes `--bench` directly to `resolve`, so users can use either
+`kernbench run --bench gemm-single-pe` or `kernbench run --bench 2`.
+
+### D7. Indices are not a stable API
+
+`_finalize()` sorts `_PENDING` alphabetically by name and assigns
+1-based indices. Adding a new bench can shift existing benches'
+indices. Therefore:
+
+- Human-interactive use: indices are fine.
+- Scripts / CI automation: always use the name.
+
+This caveat is documented in `registry.py`'s module docstring.
+
+### D8. Surface RuntimeContext exposes to benches
+
+A bench's `torch` parameter may legitimately use:
+
+- **Tensor creation**: `torch.empty(shape, dtype=..., dp=DPPolicy(...),
+  name=...)`, `torch.zeros(...)`, `torch.from_numpy(arr)`. All submit
+  host-side metadata plus device deployment (`MmuMapMsg` +
+  `MemoryWriteMsg`).
+- **Kernel launch**: `torch.launch(kernel_name, kernel_fn, *args)` —
+  converts `(Tensor, int, float)` positional args to `TensorArg` /
+  `ScalarArg`, submits per-SIP `KernelLaunchMsg`, and drains.
+- **Synchronization**: `torch.wait(handle)`, `torch.wait_all()`
+  (`run_bench` calls the latter automatically).
+- **Distributed**: `torch.distributed.init_process_group(backend="ahbm")`,
+  `torch.distributed.get_world_size()`,
+  `torch.distributed.all_reduce(t, op=...)` (ADR-0024/0027).
+- **Multi-process (rank = SIP)**:
+  `torch.multiprocessing.spawn(_worker, ..., nprocs=ws)` (ADR-0024 D3 /
+  ADR-0027).
+- **Device binding**: `torch.ahbm.set_device(rank)` or
+  `torch.accelerator.set_device_index(rank)` (both point to the same
+  namespace).
+- **IPCQ install**: `torch.install_ipcq(algorithm=..., ccl_yaml=...)`
+  (ADR-0023 D10).
+- **Spec lookup**: `torch.spec` — the dict produced by the topology
+  builder (system / cube_mesh / HBM parameters etc.). Use it so the
+  bench does not hardcode topology.yaml values.
+
+Benches must not access RuntimeContext private members (`_handles`,
+`_traces`, `_allocators`, etc.) directly. This aligns with ADR-0007's
+layer-boundary spirit: bench → runtime API → sim_engine flows in one
+direction.
+
+### D9. Environment-variable parameterization is allowed
+
+Benches may parameterize themselves via `os.environ.get(...)`, as
+`matmul_composite.py` does for `MATMUL_M`, `MATMUL_K`, `MATMUL_N`,
+`MATMUL_DTYPE`, `MATMUL_VARIANT`. Rationale:
+
+- The bench function signature is fixed by D3 to `def run(torch)`, so
+  positional/keyword arguments cannot carry parameters.
+- The env-var pattern is a natural hook for operational sweeps (e.g.,
+  `MATMUL_VARIANT`).
+- External drivers such as `scripts/gemm_sweep.py` (ADR-0044) consume
+  this hook (it sets `MATMUL_M/K/N/VARIANT` at
+  `scripts/gemm_sweep.py:115-118`).
+
+When environment variables alter bench behavior, the module docstring
+must list every variable used (`matmul_composite.py` is the canonical
+example).
+
+## Alternatives Considered
+
+### A1. An explicit manifest file (YAML) listing benches
+
+Rejected. The `@bench` + audit pattern guarantees "drop in file → auto-
+register", concentrating cognitive cost in one place (the file itself).
+A separate manifest is prone to drift, and helper separation is already
+clear via the `_` prefix.
+
+### A2. Allowing the bench's entry-point name in the decorator
+(`@bench(name=..., entry="run_xxx")`)
+
+Rejected. Breaks the simplicity of dispatch (`spec.run` is a single
+callable). The `run` convention is sufficient; variants can register
+multiple `@bench`-decorated functions in the same module.
+
+### A3. A separate `@multi_device_bench` decorator for CCL
+
+Rejected. The two patterns named in D5 (single + ADR-0024 multi-SIP)
+cover all 8 current benches. A separate decorator would force dispatch
+to branch and add complexity; the multi-SIP intent is already obvious
+from the bench's `init_process_group(...)` call.
+
+### A4. Make indices a stable API (registration order or explicit
+`index=` argument)
+
+Rejected. D7's trade-off favors user-friendliness — alphabetically
+sorted 1-based indices read naturally in the `list` output. Scripts can
+use names.
+
+## Consequences
+
+- "How to add a bench" is consolidated in one ADR — new authors only
+  need to read D1-D3 and D8 without grepping source.
+- The `_`-prefixed helper-module pattern is legitimized at ADR level,
+  so future `benches/_*.py` shared helpers are free to be added.
+- The CLI's single-device convention and CCL's multi-SIP exception are
+  shown to be consistent (D5) — they are orthogonal.
+- The rationale for ADR-0044's GEMM eval harness using env-var hooks
+  (D9) is now ADR-pinned.
+- Indices are explicitly unstable (D7), so any CI code calling
+  `kernbench run --bench 3` is flagged for review after this ADR is
+  accepted.

docs/adr/ADR-0046-prog-tl-context-contract.md

@@ -0,0 +1,327 @@
+# ADR-0046: TLContext — Kernel-side `tl.*` API Contract
+
+## Status
+
+Accepted (2026-05-22).
+
+Documents the set of `tl.*` primitives exposed by
+`src/kernbench/triton_emu/`'s `TLContext`, their semantics, and the two
+execution-mode contracts (command-list / greenlet runner). ADR-0014/0020
+defines the PE pipeline and the 2-pass execution model, but **the `tl.*`
+surface that bench kernel functions call** had no ADR-level coverage.
+
+## First action
+
+When `TLContext(pe_id, num_programs, dispatch_cycles, runner, cube_id,
+num_cubes, scratch_base, scratch_size)` is instantiated, the first action
+is to initialize six categories of state:
+
+- `self._pe_id`, `self._num_programs`, `self._cube_id`, `self._num_cubes` —
+  values that `tl.program_id` / `tl.num_programs` will return.
+- `self._dispatch_cycles` — cycle count emitted as `PeCpuOverheadCmd(cycles)`
+  at the start of every `tl.*` API call.
+- `self._runner` — `KernelRunner` instance (present → greenlet mode;
+  absent → command-list mode).
+- `self._commands: list[PeCommand] = []` — command-list accumulator
+  (command-list mode only).
+- `self._handle_counter = 0`, `self._completion_counter = 0` — counters
+  for generating TensorHandle / CompletionHandle ids.
+- `self._scratch_base`, `self._scratch_size`, `self._scratch_cursor = 0` —
+  PE-local scratch region (used for math/dot/composite output handle
+  addresses).
+
+In short, **TLContext's first act is "record where (sip/cube/pe) and at
+what scale (num_programs/num_cubes) this kernel instance runs, and pick
+its dispatch mode (runner present or not)"**. No SimPy event is created
+and no command is emitted at this moment.
+
+The runtime first action happens when the kernel function first calls a
+`tl.<api>()`. The standard entry for every `tl.*` API is:
+
+1. Call `self._emit_dispatch_overhead()` — if `dispatch_cycles > 0`,
+   immediately `_emit` a `PeCpuOverheadCmd(dispatch_cycles)`.
+2. Per-API processing (TensorHandle creation, command construction).
+3. `self._emit(cmd)` — in runner mode this `greenlet.switch()`es the cmd
+   to SimPy; in command-list mode it appends to `self._commands`.
+
+## Context
+
+The `tl.*` surface consists of `TLContext`'s methods, and the `tl`
+parameter received by a kernel function is one of these objects. The
+contract the user (bench author) sees:
+
+- Which primitives exist.
+- What data flow each primitive triggers (DMA / compute / IPCQ /
+  metadata-only).
+- How a TensorHandle's `space` and `addr` are decided.
+- The difference between command-list and greenlet modes.
+
+ADR-0014 (PE pipeline) defines the PeCommands consumed by PE_SCHEDULER,
+but how `tl.*` emits them is a code-only convention. ADR-0020 (2-pass
+data execution) mentions greenlet mode in D3 but does not pin down the
+signature difference (return-value handling) between the runner /
+non-runner paths. This ADR fills the gap.
+
+## Decision
+
+### D1. The `tl` parameter is a `TLContext` instance
+
+A bench kernel function has the signature:
+
+```python
+def _kernel(arg1, arg2, ..., tl, **kwargs):
+    ...
+```
+
+`tl` is a `kernbench.triton_emu.tl_context.TLContext` instance. The name
+imitates real Triton's `triton.language` module; the actual Triton
+module is **not** passed in.
+
+The kernel is plain Python — no `yield` or `async`. `tl.*` calls produce
+SimPy events, but to the caller they appear synchronous because in
+greenlet mode the KernelRunner relays between SimPy and the kernel
+(ADR-0020 D3).
+
+### D2. Two execution modes — command-list / greenlet runner
+
+- **Command-list mode (`runner is None`)**: `tl.*` calls append PeCommand
+  to `self._commands`. DMA / GEMM / Math consume no SimPy time and return
+  metadata-only TensorHandles (`data=None`). PE_SCHEDULER / sim_engine
+  later replays the command sequence in time.
+
+- **Greenlet runner mode (`runner is not None`)**: `tl.*` calls
+  `self._emit(cmd)` → `runner.switch_to_simpy(cmd)`, handing control to
+  the parent greenlet (SimPy). The parent distributes the cmd to
+  components, consumes SimPy time, and (for DMA reads) returns real numpy
+  data. The kernel receives the result and continues to the next line
+  (the data-aware execution model from ADR-0020 D3).
+
+The choice of mode is decided by whether a KernelRunner is injected into
+the TLContext. The `tl.*` methods themselves are mode-blind — they go
+through `_emit()` uniformly.
+
+### D3. Primitive categories
+
+#### D3.1. Reference (no DMA, metadata only)
+
+- `tl.ref(ptr, shape, dtype="f16") -> TensorHandle`: create a handle
+  referencing HBM data without issuing DMA. Used when the scheduler
+  streams the data per-tile (e.g., the b operand of a composite GEMM).
+
+#### D3.2. Data movement (blocking, DMA engine)
+
+- `tl.load(ptr, shape, dtype="f16") -> TensorHandle`: HBM → handle.
+  Emits `DmaReadCmd`. In greenlet mode the returned handle's `.data`
+  carries real numpy data; in command-list mode it is a placeholder.
+  The handle has `space="hbm"`, `pinned=True`.
+- `tl.store(ptr, handle) -> None`: TCM → HBM. Emits `DmaWriteCmd`. In
+  greenlet mode, when `handle.data` is present, `_store.write("hbm",
+  ptr, data)` runs first (visibility = issue time, ADR-0020 D3).
+
+#### D3.3. GEMM / compute (blocking)
+
+- `tl.dot(a, b) -> TensorHandle`: `a @ b`. Both operands must live in
+  TCM; shapes `(M,K) × (K,N) → (M,N)`. Emits `GemmCmd`; the output
+  handle is allocated from PE-local scratch via
+  `_make_compute_out(shape, dtype)`.
+- `tl.composite(op, a, b=None, out_ptr=0, math_op=None, epilogue=None,
+  acc_dtype=None, tile_shape=None) -> CompletionHandle`: non-blocking
+  tiled pipeline. Emits `CompositeCmd`. `epilogue` is a list of dicts,
+  each with `"op"` plus op-specific fields and an optional `"scope"`
+  (k_tile / output_tile). Unknown ops or missing fields raise
+  ValueError immediately. The returned CompletionHandle synchronizes
+  via `tl.wait(h)`.
+
+#### D3.4. Math: unary (blocking)
+
+- `tl.exp(x)`, `tl.log(x)`, `tl.sqrt(x)`, `tl.abs(x)`, `tl.sigmoid(x)`,
+  `tl.cos(x)`, `tl.sin(x)` — each emits `MathCmd(op=<name>,
+  inputs=(x,), out=)`. `out` is scratch-allocated with the same
+  shape/dtype as `x`.
+
+#### D3.5. Math: binary (blocking)
+
+- `tl.maximum(a, b)`, `tl.minimum(a, b)` — `_binary_math`.
+- `tl.fma(a, b, c)` — `a*b + c`. Three inputs.
+- `tl.clamp(x, min, max)` — `MathCmd(op="clamp", inputs=(x, min, max))`.
+- `tl.where(cond, a, b)` — `MathCmd(op="where", inputs=(cond, a, b))`.
+- `tl.softmax(x, axis=-1)` — a single `MathCmd(op="softmax")` so timing
+  accounts at one dispatch. Phase 2 DataExecutor expands it to the
+  canonical (x-max → exp → sum → div) sequence.
+
+#### D3.6. Reduction (blocking)
+
+- `tl.sum(x, axis)`, `tl.max(x, axis)`, `tl.min(x, axis)` — return an
+  output handle with the axis size collapsed to 1. Emit
+  `MathCmd(op=<name>, inputs=(x,), out=, axis=axis)`.
+
+#### D3.7. Index / scalar (PE_CPU, no engine)
+
+- `tl.program_id(axis=0) -> int`: `axis==0` → pe_id (cube-local PE
+  index), `axis==1` → cube_id (ADR-0022).
+- `tl.num_programs(axis=0) -> int`: `axis==0` → num_programs (PEs per
+  cube), `axis==1` → num_cubes.
+- `tl.arange(start, end, dtype="i32") -> TensorHandle`: an index range
+  in TCM. No command emitted.
+- `tl.zeros(shape, dtype="f16") -> TensorHandle`, `tl.full(shape,
+  value, dtype="f16") -> TensorHandle`: TCM placeholder. No command
+  emitted.
+
+#### D3.8. Scalar helpers (no command, no engine)
+
+- `TLContext.cdiv(a, b) -> int` (static): ceiling division
+  `-(-a // b)`. Mirrors real Triton's `tl.cdiv`.
+
+#### D3.9. Metadata-only (no compute, no DMA)
+
+- `tl.trans(x) -> TensorHandle`: a new handle with the last two dims
+  swapped. Shares `addr` and `data`; no command emitted.
+
+#### D3.10. IPCQ (CCL) primitives (ADR-0023 D4)
+
+- `tl.send(dir, src=None, *, src_addr=None, nbytes=None, shape=None,
+  dtype="f16", space="tcm") -> None`: blocking send. Accepts either
+  handle form or raw-address form. Emits `IpcqSendCmd`. The handle's
+  `.data` snapshot rides along on the command — avoiding the race
+  where a later inbound IPCQ overwrites the slot before the outbound
+  PE_DMA reads it.
+- `tl.recv(dir=None, shape=(), dtype="f16", space="tcm", dst_addr=None,
+  dst_space=None) -> TensorHandle`: blocking recv. Providing both
+  `dst_addr` and `dst_space` enters "copy_to_dst" mode; otherwise
+  "return_slot" mode. In greenlet mode the handle's `.data` carries
+  the real data.
+- `tl.recv_no_consume(dir=None, shape=(), dtype="f16") -> TensorHandle`:
+  **DIAGNOSTIC ONLY**. Has the same blocking-arrival semantics as
+  `tl.recv` but skips the slot-read latency charge (slot-IO + PE↔bank
+  fabric drain). Used in the pe2pe overview plot for an apples-to-apples
+  comparison against `tl.store`. Production kernels MUST NOT use it —
+  the diagnostic flag is isolated in its own command branch
+  (`consume=False`) so it cannot be accidentally enabled.
+- `tl.recv_async(dir, shape=(), dtype="f16") -> RecvFuture`: non-blocking
+  recv. Returns a `RecvFuture`; resolved later by `tl.wait(future)`.
+
+#### D3.11. Composite + control
+
+- `tl.composite(...)`: see D3.3.
+- `tl.wait(handle=None)`: wait on a `CompletionHandle` (composite), a
+  `RecvFuture` (async recv), or `None` (all pending composites).
+- `tl.cycles(n)`: declare a scalar PE_CPU overhead. Emits
+  `PeCpuOverheadCmd(cycles=n)`.
+
+### D4. TensorHandle arithmetic operators — thread-local TLContext
+
+At module load, `tl_context.py::_enable_tensor_ops()` runs and patches
+`TensorHandle.__add__`, `__sub__`, `__mul__`, `__truediv__`. Each
+operator calls `_binary_math` on the active TLContext stored in a
+module-level thread-local `_ctx`.
+
+So inside a kernel, `c = a + b` is equivalent to emitting
+`MathCmd(op="add", inputs=(a, b), out=)` and returning a new
+TensorHandle.
+
+Active-TLContext management:
+
+- `TLContext._set_active(ctx)`: set the active ctx for the current
+  thread/greenlet.
+- `TLContext._get_active()`: read it (RuntimeError if unset).
+- `run_kernel(kernel_fn, tl_ctx, *args, **kwargs)`: helper. Sets active
+  on entry, runs the kernel, restores `None` on exit.
+
+`KernelRunner` re-asserts `_set_active(tl)` inside its `_switch_kernel`
+just before resuming the kernel, so a sibling PE runner that overwrote
+the thread-local context is correctly recovered.
+
+### D5. Scratch allocator — compute output handles
+
+Ops that produce a result — `tl.dot`, `tl.exp`, `tl.add` (via
+TensorHandle `__add__`), etc. — call `_make_compute_out(shape, dtype)`
+to obtain a 16-byte-aligned scratch address. The address is published
+with `space="tcm"`, so the handle can later be the source of a
+`tl.send` / `tl.store`.
+
+When `_scratch_base == 0` (e.g., command-list mode), the address is 0
+and the handle cannot be a send/store source (in that case, only
+`tl.load`-returned handles are valid sources).
+
+When the cursor exceeds `_scratch_size` (default 1 MiB), a
+RuntimeError is raised. The cursor must reset between kernel
+invocations (current code naturally satisfies this: KernelRunner
+creates a fresh TLContext each time).
+
+### D6. Dispatch overhead — `PeCpuOverheadCmd(dispatch_cycles)`
+
+Every non-metadata `tl.*` call starts with `_emit_dispatch_overhead()`,
+which — when `dispatch_cycles > 0` — emits
+`PeCpuOverheadCmd(dispatch_cycles)`. This models the cycles PE_CPU
+spends dispatching the command.
+
+Defaults:
+
+- `TLContext.__init__`'s `dispatch_cycles` parameter default: `1` cycle.
+- TLContext built by `KernelRunner`: `0` cycles (greenlet mode handles
+  cycle accounting differently — aligned with ADR-0020 D3 intent).
+
+### D7. Kernel registry (`triton_emu/registry.py`)
+
+A separate `_kernels: dict[str, Callable]` holds the name → function
+mapping:
+
+- `register_kernel(name, fn)`: ValueError on duplicate.
+- `get_kernel(name)`: KeyError if missing.
+- `clear_registry()`: test-only.
+
+`RuntimeContext.launch(kernel_name, kernel_fn, *args)` overwrites
+`_kernels[kernel_name] = kernel_fn` on every call (last-call-wins,
+idempotent) — consistent with ADR-0045 D8's `launch` behavior.
+
+PE_CPU looks up `KernelRef.name` in the registry and runs the function
+through KernelRunner.
+
+## Alternatives Considered
+
+### A1. Fold `tl.*` into ADR-0014 / ADR-0020
+
+Rejected. ADR-0014 covers the PE pipeline (sim_engine-side consumption
+of PeCommands); ADR-0020 covers 2-pass execution (Phase 1 timing /
+Phase 2 data). The `tl.*` surface is what the kernel author touches; a
+dedicated ADR improves findability and onboarding.
+
+### A2. Deprecate command-list mode
+
+Rejected (currently). Simple unit tests and kernel verification benefit
+from the lighter command-list path — it exposes a PeCommand sequence
+inspector without requiring greenlet machinery. When greenlet-mode
+semantics (real data, Phase 2) are needed, D2 explicitly selects them.
+
+### A3. Remove TensorHandle arithmetic operators
+
+Rejected. They mimic real Triton kernel ergonomics (e.g., `c = a + b`),
+and the thread-local active-ctx pattern works cleanly. The explicit
+function-form (`tl.add(a, b)`) is also exposed in D3.5, so the
+operators are syntactic sugar.
+
+### A4. Expand softmax into the explicit sequence (max → exp → sum → div)
+
+Partially adopted. `tl.softmax` is a single `MathCmd(op="softmax")` for
+timing accounting (D3.5), but Phase 2 DataExecutor expands it to the
+canonical sequence for real-data computation. Timing model atomic,
+data model expanded — the two split intentionally.
+
+## Consequences
+
+- Every `tl.*` primitive a bench author meets is classified and defined
+  in a single ADR. Paired with ADR-0045 D8's host-side surface
+  (`torch.empty` etc.), the inside-kernel and outside-kernel authoring
+  guides are now complete.
+- The command-list / greenlet difference is pinned in D2, so any new
+  `tl.*` primitive that follows the `_emit()` pattern auto-supports
+  both modes.
+- The thread-local active-ctx pattern (D4) is justified at ADR level,
+  clarifying who owns the reset responsibility when multiple PE
+  runners share a thread (KernelRunner.run's contract restores active
+  inside `_switch_kernel`).
+- `tl.recv_no_consume`'s diagnostic isolation (D3.10) is hardened in
+  ADR form — accidental production use is blocked by a separate
+  command branch.
+- The registry (D7) gets its own D-section, formalizing the
+  name-collision and dynamic-re-registration semantics.

docs/adr/ADR-0047-par-ahbm-ccl-backend.md

@@ -0,0 +1,259 @@
+# ADR-0047: AHBM CCL Backend — `torch.distributed`-compat shim
+
+## Status
+
+Accepted (2026-05-22).
+
+Pins down what `runtime_api/distributed.py`'s `AhbmCCLBackend` +
+`DistributedContext` actually install — i.e., the entry point
+`torch.distributed.init_process_group(backend="ahbm")` — and how
+`all_reduce`/`barrier`/`get_rank` etc. are implemented. ADR-0023 D11
+mentions the "torch.distributed compatibility" intent, but **the backend
+itself** had no ADR-level coverage.
+
+## First action
+
+`RuntimeContext.__post_init__` automatically constructs a
+`DistributedContext()` and attaches it to `self.distributed`. The first
+action at that moment:
+
+1. `self._backend: AhbmCCLBackend | None = None` — uninitialized.
+2. `self._rank_by_greenlet: dict = {}` — greenlet-local rank registry
+   (ADR-0024 D2).
+3. The caller (RuntimeContext) sets `dc._ctx_ref = self` so subsequent
+   `init_process_group` can reach `ctx.engine` / `ctx.spec` / `ctx.launch`.
+
+In short, **DistributedContext's first act is "attach to RuntimeContext
+with a back-reference and leave the backend slot empty"**. Actual
+backend installation (IPCQ install, world_size derivation, algorithm
+module import) happens only when user code calls
+`torch.distributed.init_process_group(backend="ahbm")`.
+
+At that moment, `init_process_group`'s first action is:
+
+1. If `backend != "ahbm"`, raise `ValueError("Unsupported backend ...")`
+   immediately.
+2. If `getattr(self, "_ctx_ref", None)` is None,
+   `RuntimeError("DistributedContext not bound to a RuntimeContext")`.
+3. `self._backend = AhbmCCLBackend(torch_ctx=ctx)` — inside this
+   constructor, ccl.yaml is loaded, the algorithm module is imported,
+   world_size is derived, SFR is configured, and IPCQ is installed.
+4. `self._backend._dist_ctx = self` — the backend gets a back-reference
+   so it can read `_rank_by_greenlet`.
+
+## Context
+
+The `AhbmCCLBackend` exists so that PyTorch DDP collective calls
+(`init_process_group`, `all_reduce`, etc.) work unchanged and bench code
+reads identically to a real DDP training script (in line with
+ADR-0024 + ADR-0027's launcher model).
+
+The backend's responsibilities:
+
+- At `init_process_group` time, install the **IPCQ neighbor table once**
+  (analogous to NCCL communicator creation).
+- For each `all_reduce(tensor, op="sum")`, dispatch the configured
+  algorithm's kernel function via `ctx.launch(...)`.
+- Answer `get_world_size` / `get_rank` consistently from the
+  greenlet-local rank registry plus ccl.yaml/topology.
+
+ADR-0023 D10 (IPCQ install plan) and ADR-0024 (SIP launcher) touch
+parts of this, but **the backend's own responsibility scope and decision
+order** are not pinned anywhere. This ADR fills that gap.
+
+## Decision
+
+### D1. The backend is created only at `init_process_group(backend="ahbm")` time
+
+`DistributedContext` starts with `_backend = None`. The backend object
+does not exist until the user calls
+`dist.init_process_group(backend="ahbm")`. Any other API
+(`is_initialized`, `get_world_size`, `all_reduce`, `barrier`) called
+while `_backend` is None raises
+`RuntimeError("Default process group has not been initialized...")` via
+the `_ensure_initialized` helper.
+
+`backend != "ahbm"` raises `ValueError` immediately. Other backend names
+(`nccl`, `gloo`, etc.) are not recognized.
+
+### D2. world_size resolution priority — algorithm > defaults > topology
+
+`AhbmCCLBackend._resolve_world_size` (ADR-0024 D1):
+
+1. If `ccl.yaml`'s algorithm entry has `world_size`, use it.
+2. Else if `defaults.world_size` is set, use it.
+3. Else fall back to `spec.system.sips.count` (the topology's SIP count).
+
+The default interpretation is **rank = SIP** (ADR-0024). Cube/PE-level
+parallelism is expressed inside each rank via DPPolicy and does not
+affect world_size. An explicit `ccl.yaml` override is preserved for the
+legacy "rank = flat PE index" test path.
+
+User arguments to `init_process_group(world_size=..., rank=...)` are
+**accepted but ignored** (same as real PyTorch's `RANK` / `WORLD_SIZE`
+env vars).
+
+### D3. `init_process_group` performs four installation steps
+
+Inside `AhbmCCLBackend.__init__`, in order:
+
+1. **Load ccl.yaml**: `kernbench.ccl.install.load_ccl_config()` →
+   `resolve_algorithm_config(_cfg_all)` produces the merged config for
+   `defaults.algorithm` (or the user-specified algorithm).
+2. **Import algorithm module**:
+   `importlib.import_module(self._merged["module"])`. The module must
+   expose a `kernel` function, a `kernel_args(world_size, n_elem,
+   cube_w, cube_h)` helper, and optionally a `TOPO_NAME_TO_KIND` map.
+3. **Resolve world_size** (D2).
+4. **Collect topology metadata** from `spec`: `n_sips`, `sip_topo`
+   (`ring_1d` default), `cube_w`/`cube_h`, `sips.w`/`sips.h`. When the
+   SIP topology is not `ring_1d`, derive `_sip_topo_w/h` from explicit
+   `w`/`h` or via square-root (require `w*h == n_sips`). Mismatch raises
+   `ValueError`.
+5. **Install SFR + IPCQ**:
+   `kernbench.ccl.sfr_config.configure_sfr_intercube_multisip(engine,
+   spec, self._merged)`. This pushes IPCQ neighbor tables to every
+   SIP/cube's pe0 (one-time setup analogous to NCCL communicator
+   creation).
+
+If the order changes (e.g., SFR runs before the algorithm module
+loads), partial initialization can result. So D3 is treated as an
+atomic 4-step block — on failure the backend remains uninstalled.
+
+### D4. Greenlet-local rank binding (ADR-0024 D2)
+
+`DistributedContext._rank_by_greenlet: dict[greenlet, int]` maps spawned
+worker greenlets to their ranks. When the bench launcher (e.g.,
+`torch.multiprocessing.spawn`) spawns a worker, it registers via
+`dc._bind_rank(g, rank)`.
+
+`get_rank()` looks up `getcurrent()`'s greenlet. Unregistered greenlets
+fall back to 0 — preserves single-driver / test compatibility.
+
+The backend reads the current greenlet's rank from
+`_dist_ctx._rank_by_greenlet` during `all_reduce` (D5).
+
+### D5. `all_reduce(tensor, op="sum")` behavior
+
+Validation:
+
+- `op != "sum"` → `NotImplementedError`. Current kernels only
+  implement add reduction.
+- `tensor._handle is None` → `RuntimeError("not deployed")`.
+- `tensor._handle.shards` empty → `RuntimeError("no shards")`.
+
+Preparation:
+
+- `n_elem = shards[0].nbytes // tensor.itemsize` — element count of a
+  single shard.
+- `kernel_fn = self._algo_module.kernel` — the algorithm module's entry
+  function (imported in D3).
+- Decide effective cube dims: if the first SIP has just 1 cube, use
+  `(1, 1)`; otherwise use the topology's `cube_w`/`cube_h`. This
+  naturally absorbs TP runs that use only a subset of cubes.
+- `kernel_args = self._algo_module.kernel_args(world_size, n_elem,
+  cube_w, cube_h)` — the algorithm decides which arguments to pass to
+  its kernel.
+
+Dispatch:
+
+- Resolve the current greenlet's rank via
+  `_rank_by_greenlet.get(g, 0)`.
+- Append `extra_args = (sip_rank, sip_topo_kind, sip_topo_w,
+  sip_topo_h)`.
+- `pending = self.ctx.launch(algorithm_name, kernel_fn, tensor,
+  *kernel_args, *extra_args, _defer_wait=True)` — `_defer_wait=True`
+  delegates collective drain to the main scheduler (ADR-0027 D0.4).
+
+Drain:
+
+- If the parent greenlet is alive (multi-greenlet mode), enqueue
+  `_pending_collective_handles` and switch to parent. The main
+  scheduler drains after all ranks have launched.
+- If single-driver mode, drain inline:
+  `for h, _sip_id, meta in pending: self.ctx.wait(h, _meta=meta)`.
+
+### D6. `barrier()` is a no-op (single-driver model)
+
+kernbench runs all ranks as greenlets inside a single Python process,
+so no cross-process synchronization is needed. `barrier()` is callable
+but does no synchronization. Kept for real-PyTorch API compatibility so
+callers don't get `NotImplementedError`.
+
+If multi-process kernbench (SimPy event loop per process) is introduced
+in the future, D6 needs a superseding ADR.
+
+### D7. Semantics of `get_rank` / `get_world_size` / `get_backend`
+
+- `get_rank()` (D4): the current greenlet's bound rank; unregistered → 0.
+- `get_world_size()` (D2): the world_size resolved by the backend in D3.
+- `get_backend()`: always the literal string `"ahbm"`. Calling before
+  backend exists triggers `_ensure_initialized`'s RuntimeError.
+
+Differences vs. real PyTorch:
+
+- Real PyTorch `get_rank()` is a process-global value; here it is
+  greenlet-local. Inside a spawned worker → the worker's rank; in the
+  main thread → 0. Bench authors should expect meaningful ranks only
+  inside worker functions.
+
+### D8. Supported API surface (final)
+
+`DistributedContext` exposes:
+
+- `init_process_group(backend="ahbm", world_size=None, rank=None,
+  **kwargs)`
+- `is_initialized() -> bool`
+- `get_world_size() -> int`
+- `get_rank() -> int`
+- `get_backend() -> str`
+- `all_reduce(tensor, op="sum") -> None`
+- `barrier() -> None`
+- (internal) `_bind_rank(g, rank)`
+
+Other PyTorch distributed APIs (`broadcast`, `reduce`, `all_gather`,
+`gather`, `scatter`, point-to-point `send/recv`, etc.) are **not
+implemented**. Kernel-level expression is available via
+`tl.send`/`tl.recv` (ADR-0046 D3.10), but the `dist.*` surface does not
+expose them. If additional collectives are needed, add a paired
+(algorithm module, `DistributedContext` method) and extend D8.
+
+## Alternatives Considered
+
+### A1. Create the backend in `RuntimeContext.__init__`
+
+Rejected. If `ccl.yaml` is missing or the algorithm module can't be
+imported, RuntimeContext construction would fail even when the bench
+does not use distributed features. Lazy creation at call time (D1) is
+the right semantics.
+
+### A2. Always derive world_size from topology (no override)
+
+Rejected. ADR-0024 D1's "explicit override" path is used by legacy
+tests. Diagnostic scenarios that define PE-level ranks within a single
+SIP also need this escape hatch.
+
+### A3. Silent fallback for unsupported `op`
+
+Rejected. If the user intends `op="prod"` / `"max"` / `"avg"` and silent
+`sum` runs instead, result validation gets very hard. Explicit
+`NotImplementedError` is safer.
+
+### A4. Implement `barrier` as a SimPy event
+
+Rejected (currently). With single-driver semantics there is no
+cross-process synchronization to express, so a no-op is meaningfully
+correct. A fake-barrier SimPy event would add code complexity for no
+semantic gain. Revisit when multi-process kernbench arrives.
+
+## Consequences
+
+- The 4-step installation (D3) for
+  `torch.distributed.init_process_group(backend="ahbm")` is locked in,
+  making clear where future collective algorithms must hook.
+- The priority order in D2 (algorithm > defaults > topology) makes the
+  blast radius of ccl.yaml changes quickly knowable.
+- The no-op `barrier` (D6) is recorded so multi-process kernbench, if
+  introduced, must explicitly supersede this ADR.
+- D8's list of unsupported APIs explicitly grounds the rejection
+  message when users call, e.g., `dist.broadcast(...)`.

docs/adr/ADR-0048-mem-allocator-algorithms.md

@@ -0,0 +1,278 @@
+# ADR-0048: Memory Allocator Algorithms — VirtualAllocator + PEMemAllocator
+
+## Status
+
+Accepted (2026-05-22).
+
+Pins down the free-list algorithm, page alignment, and coalescing rules
+used by `policy/address/allocator.py`'s `_FreeList` / `PEMemAllocator`
+and `va_allocator.py`'s `VirtualAllocator`. ADR-0001 (PhysAddr layout)
+and ADR-0011 (PA/VA/LA models) define the address schemes; the
+**allocation algorithms** had no ADR-level coverage.
+
+## First action
+
+### `_FreeList(capacity)`
+
+On construction: `self._capacity = capacity`, `self._used = 0`,
+`self._free = [(0, capacity)]`. The first act is **establishing the
+entire region as one free block** — the tuple `(offset=0,
+size=capacity)` is the sole entry in the free list.
+
+### `PEMemAllocator(sip_id, die_id, pe_id, cfg)`
+
+On construction, builds two `_FreeList`s:
+
+- `self._hbm = _FreeList(cfg.hbm_slice_bytes)` — the size of this PE's
+  HBM slice (`hbm_bytes_per_cube // hbm_slices_per_cube`).
+- `self._tcm = _FreeList(cfg.tcm_allocatable_bytes)` — equals
+  `tcm_bytes_per_pe - tcm_scheduler_reserved_bytes` (the scheduler
+  reservation is pre-deducted).
+
+So PEMemAllocator's first act is **constructing single-free-block
+HBM-slice and TCM regions for this PE**.
+
+### `VirtualAllocator(va_base, va_size, page_size=2*1024*1024)`
+
+On construction: `self._va_base = va_base`, `self._va_size = va_size`,
+`self._page_size = page_size`, `self._used = 0`, `self._free =
+[(va_base, va_size)]`. The first act is **establishing one block from
+va_base to va_size and stashing page_size**.
+
+## Context
+
+`runtime_api/context.py::_ensure_allocators` builds the allocator set
+in these stages:
+
+1. Read `hbm_total_gb_per_cube`, `hbm_slices_per_cube`, `tcm_size_mb`,
+   per-target_device SIP range, etc. from `spec`.
+2. Pack everything into a frozen `AddressConfig`.
+3. For every combination in the target SIP range × cubes × PEs,
+   construct one `PEMemAllocator(sip, cube, pe, cfg)` instance.
+4. Construct one `VirtualAllocator(va_base=0x1_0000_0000, va_size=64
+   GiB, page_size=pe_mmu.page_size)`.
+
+Allocator responsibilities:
+
+- **PEMemAllocator**: PA-space allocation in the PE-local HBM slice /
+  TCM (including PhysAddr encoding).
+- **VirtualAllocator**: device-wide VA allocation, page-aligned.
+  `RuntimeContext._create_tensor` then pushes VA → PA mappings to
+  components via `MmuMapMsg`.
+
+These algorithms are:
+
+- **First-fit**, kept simple.
+- The free-block list is **sorted by start offset**.
+- On `free()`, **adjacent blocks coalesce**.
+
+The rationale was not documented anywhere, so when someone asks "why
+not best-fit?", "why not a buddy allocator?", "why does partial-overlap
+free pass silently?", there was no anchor to answer from. This ADR
+provides it.
+
+## Decision
+
+### D1. `_FreeList` — offset-keyed first-fit + coalescing
+
+`policy/address/allocator.py::_FreeList`:
+
+- Internal representation: `list[tuple[int, int]] = [(start_offset,
+  size), ...]` — sorted by start offset.
+- `alloc(nbytes)`:
+  1. Iterate the free list from the front (first-fit).
+  2. Take from the first block with `size >= nbytes`.
+  3. Exact match → drop the block; otherwise shrink it to `(start +
+     nbytes, size - nbytes)`.
+  4. `_used += nbytes`; return the taken `start`.
+  5. If no block fits, `AllocationError("overflow ... largest free
+     block ...")`.
+- `free(offset, nbytes)`:
+  1. `_used -= nbytes`.
+  2. `bisect_left(self._free, (offset,))` finds the insertion index.
+  3. If adjacent to the previous block (`prev_start + prev_size ==
+     offset`), merge.
+  4. If adjacent to the next block (`offset + nbytes == next_start`),
+     merge.
+  5. Insert the coalesced range at the right sorted position.
+
+This algorithm is weaker than best-fit / buddy on fragmentation, but
+the simulator's workload (mostly stack-like deploy/free) tolerates it.
+If the workload shape changes, D1 is a supersession candidate.
+
+### D2. Partial-overlap free is **not** validated
+
+`_FreeList.free(offset, nbytes)` trusts the caller to pass the exact
+`(offset, nbytes)`. It does **not** verify:
+
+- That the range was actually allocated.
+- That the range does not overlap another allocated region.
+
+Reason: in a simulator context, callers always store the return value
+of `alloc()` and pass it back to `free()` — there is no external user
+input. Adding a safety check would cost O(N) per free and impact
+simulation wall-clock.
+
+If this trust model breaks (e.g., a code path lets two tensors point
+at the same PA), this ADR must be revisited.
+
+### D3. `PEMemAllocator` — two channels for HBM/TCM
+
+`PEMemAllocator(sip_id, die_id, pe_id, cfg)` holds two `_FreeList`s:
+
+- `_hbm`: size `cfg.hbm_slice_bytes`.
+- `_tcm`: size `cfg.tcm_allocatable_bytes` (= `tcm_bytes_per_pe -
+  tcm_scheduler_reserved_bytes`).
+
+`alloc_hbm(nbytes) -> PhysAddr`:
+
+- `_hbm.alloc(nbytes)` → offset.
+- `PhysAddr.pe_hbm_addr(sip_id, die_id, pe_id,
+  pe_local_hbm_offset=offset, slice_size_bytes=cfg.hbm_slice_bytes)`.
+- Failure raises `AllocationError("HBM overflow ...")`.
+
+`free_hbm(pa, nbytes)`:
+
+- Recover PE-local offset via `pa.hbm_offset - pe_id *
+  cfg.hbm_slice_bytes`.
+- `_hbm.free(offset, nbytes)`.
+
+`alloc_tcm(nbytes) -> PhysAddr`: similar; uses `PhysAddr.pe_tcm_addr`.
+
+`free_tcm(pa, nbytes)`: uses `pa.sub_offset` directly (TCM's PE-local
+offset equals its sub_offset).
+
+The allocator does not see the scheduler-reserved TCM region
+(`cfg.tcm_scheduler_reserved_bytes`) — it is pre-subtracted from the
+`_tcm` capacity. This is consistent with ADR-0014's PE_SCHEDULER
+internal-buffer reservation.
+
+### D4. `VirtualAllocator` — page-aligned first-fit + coalescing
+
+`policy/address/va_allocator.py::VirtualAllocator`:
+
+- Internal representation: same sorted `list[tuple[int, int]]` as
+  `_FreeList`. Initially `[(va_base, va_size)]`.
+- `_align_up(nbytes) = ceil(nbytes / page_size) * page_size`.
+- `alloc(nbytes) -> int`:
+  1. `aligned = _align_up(nbytes)`.
+  2. First-fit a block with `size >= aligned`.
+  3. Take `aligned` from the block's front; remove if exact.
+  4. `_used += aligned`. Return the block's `start` (which is page-
+     aligned).
+  5. Failure → `VaAllocationError`.
+- `free(va, nbytes)`: free `_align_up(nbytes)` worth. Coalesces with
+  the same algorithm as `_FreeList`.
+
+`page_size` has different defaults in two places:
+
+- `VirtualAllocator.__init__`'s parameter default: `2 MiB`. Direct-call
+  tests receive this.
+- `RuntimeContext._ensure_allocators` when constructing the instance:
+  `pe_mmu.attrs.get("page_size", 4096)` — uses
+  `topology.yaml`'s `pe_mmu.attrs.page_size` if set, else falls back
+  to `4 KiB`.
+
+The two defaults differ on purpose: `VirtualAllocator`'s standalone
+default (`2 MiB`) aligns with ADR-0039's PE_MMU stopgap default for
+direct-test ergonomics; the context fallback (`4 KiB`) is the safe
+minimum when `topology.yaml` doesn't specify a page size. The
+production path is always the latter (via `_ensure_allocators`), and
+when `topology.yaml` sets `page_size`, that value flows consistently
+into both the MMU and the VA allocator.
+
+If consistency breaks (e.g., VirtualAllocator instantiated with a
+page_size different from PE_MMU's), MMU `map()` falls into the
+sub-page region mode (ADR-0039 D3).
+
+VA range defaults: `va_base = 0x1_0000_0000` (= 4 GiB), `va_size = 64
+GiB`. These are hardcoded in `_ensure_allocators` and have no
+semantic meaning in ADR-0011's VA model — they simply reserve enough
+device-wide space without colliding with host code.
+
+### D5. Lifecycle of allocator instances
+
+- `RuntimeContext._ensure_allocators` is lazy — called on the first
+  `_create_tensor`.
+- The allocator dict (`self._allocators`) lives for the
+  RuntimeContext's lifetime. A second deploy in the same process
+  does not construct new objects.
+- `RuntimeContext.cleanup()` walks living tensors and calls
+  `_free_tensor()`, which issues MMU unmaps + `va_allocator.free` +
+  `pemem_allocator.free_hbm` — restoring the free lists. A subsequent
+  RuntimeContext starts fresh.
+
+This per-RuntimeContext isolation guarantees deterministic deploy →
+cleanup → deploy sequences within a single process.
+
+### D6. Allocator failure raises (no silent OOM)
+
+Both `_FreeList.alloc` and `VirtualAllocator.alloc` raise
+`AllocationError` / `VaAllocationError` when no block fits. The message
+includes "required size + largest available block" to distinguish
+fragmentation from true OOM.
+
+A silent fallback (e.g., allocating only as much as the largest free
+block) is strictly forbidden — a partially-allocated tensor reaching
+SimPy would cause routing / DMA to see incorrect PAs and silently
+corrupt simulation results.
+
+### D7. One allocator per address space
+
+Physical address spaces are separated by PhysAddr sub-units (ADR-0001
+D2.3); each sub-unit gets its own allocator instance:
+
+- HBM slice → `PEMemAllocator._hbm`.
+- PE TCM → `PEMemAllocator._tcm`.
+- (Currently unused) M_CPU local memory, CUBE SRAM → would need their
+  own allocators. Today these are handled as IPCQ-only slots (ADR-0023
+  D9.7) and do not share PA space, so no free-list exists for them.
+
+When a cube-level SRAM allocator is needed,
+`_FreeList(cfg.sram_bytes_per_cube)` is added per-cube
+(`cfg.sram_bytes_per_cube` is already defined in `AddressConfig` —
+the data model is ready).
+
+## Alternatives Considered
+
+### A1. Best-fit / buddy allocator
+
+Rejected (currently). The workload's alloc/free pattern is stack-like
+(deploy order ≈ free order), so first-fit + coalescing controls
+fragmentation well enough. If long-running fragmentation appears in LLM
+kernel sweeps, a buddy-allocator ADR will replace D1.
+
+### A2. Add partial-overlap free validation
+
+Rejected. D2's trust model plus the O(N) per-free cost makes this
+unattractive. A debug mode (e.g., `KERNBENCH_DEBUG` env var) that
+enables the check could be added later.
+
+### A3. A unified allocator for VA and PA
+
+Rejected. VA space (64 GiB device-wide) and PA space (per-slice ~6
+GiB) have different semantic dimensions — VA is the kernel's view, PA
+is the device sub-unit's view. ADR-0011's VA model (MMU maps between
+the two) calls for separated allocators.
+
+### A4. Multi-tier page sizes (large pages + small pages)
+
+Rejected (currently). A single page size (2 MiB) matches LLM kernel
+tensor sizes (a few MiB to GiB); smaller mappings are absorbed by
+ADR-0039 D3's sub-page region mode. Multi-tier paging would require
+extending the MMU model itself — a separate ADR candidate.
+
+## Consequences
+
+- The allocator algorithm is pinned at ADR level (D1, D3, D4), so any
+  future simulation scenario hitting fragmentation has a clear "we're
+  using first-fit + coalescing" anchor to inspect.
+- D2's trust model is explicit, so any future code path that exposes
+  alloc/free to direct user input will trigger this ADR's supersession
+  early.
+- D7's one-allocator-per-sub-unit mapping is on record, so when M_CPU
+  or SRAM need their own free-list, the addition point is obvious.
+- D4 captures the page_size dual-default and its production path
+  (`_ensure_allocators` always wins), letting future `topology.yaml`
+  `page_size` changes be assessed against ADR-0039's stopgap
+  interaction quickly.

docs/adr/ADR-0049-ver-probe-subcommand.md

@@ -0,0 +1,247 @@
+# ADR-0049: `kernbench probe` Subcommand — Traffic-Pattern Verification Harness
+
+## Status
+
+Accepted (2026-05-22).
+
+Pins down the traffic-pattern catalog, formula-vs-actual comparison, and
+invariant checks (monotonicity, D2H ≥ H2D, etc.) exposed by
+`probes/probe.py::run_probe(...)`. ADR-0010 (CLI surface) enumerates the
+`kernbench probe` subcommand, but **what probe actually measures** and
+**which invariants it judges PASS/FAIL** had no ADR-level coverage.
+
+## First action
+
+`run_probe(topology_path, case_filter=None)` performs four startup steps:
+
+1. `Path(topology_path).expanduser().resolve()` → absolute path.
+2. `load_topology(path)` → `TopologyGraph` (graph + spec).
+3. `_build_edge_map(graph)` → a `{(src, dst): Edge}` lookup table.
+4. Instantiate `AddressResolver(graph)` + `PathRouter(graph)`.
+
+Then it sets `nbytes = 32768` (= 32 KiB, the summary-table reference
+size) and `show_all = (case_filter is None or case_filter == "all")`.
+
+In short, **probe's first act is "load the topology once and prepare
+edge map / resolver / router, plus pin 32 KiB as the standard measurement
+size"**. After that, the H2D → D2H → PE DMA categories execute in
+separate `GraphEngine` instances (no cross-talk between cases).
+
+## Context
+
+`kernbench probe` was introduced as a verification tool for these
+purposes:
+
+- **Manual ground truth**: when a real-simulation result (`kernbench run
+  --bench ...`) shows abnormal latency, derive the answer for a simple
+  traffic pattern in isolation and compare.
+- **Formula vs actual**: check whether the analytical model
+  (wire latency + overhead + drain) matches the simulator's
+  `total_ns`. A mismatch points to which simplifying assumption in
+  ADR-0033 is missing.
+- **Monotonicity check**: latency should grow monotonically with hop
+  count.
+- **Utilization sweep**: a BW-utilization table across data sizes
+  (4 KiB ~ 1 MiB).
+
+Without an ADR for this tool:
+
+- Adding a new traffic-pattern category (e.g., MCpuDma, IPCQ) is hard
+  because the table format / measurement units of existing categories
+  aren't documented at the ADR level.
+- The basis for the monotonicity check (hop count? cube distance? wire
+  length?) is ambiguous.
+- The reference size 32 KiB and the sweep `[4 KiB, 16 KiB, 64 KiB, 256
+  KiB, 1 MiB]` are only discoverable by reading source.
+
+## Decision
+
+### D1. Three case categories — H2D / D2H / PE DMA
+
+Each category has a distinct data path in the topology and gets its own
+summary table + sweep table + route-detail block.
+
+- **H2D (Host → Device Write)**: `MemoryWriteMsg(dst_sip=0, dst_cube,
+  dst_pe=0, pattern="zero")` flows along `pcie_ep → io_cpu → m_cpu →
+  hbm_ctrl`. The cube index varies the hop count:
+  - h2d-1hop: cube=0, hops=1
+  - h2d-2hop: cube=4, hops=2
+  - h2d-3hop: cube=8, hops=3
+  - h2d-4hop: cube=12, hops=4
+- **D2H (Device → Host Read)**: `MemoryReadMsg(src_sip=0, src_cube,
+  src_pe=0)`. Total latency = forward command path + reverse data path.
+  Same 4-hops category as H2D.
+- **PE DMA (PE-initiated)**: `PeDmaMsg(src_sip, src_cube, src_pe,
+  dst_pa)`. Five cases cover varying cube/PE positions:
+  - pe-local-hbm: same cube, same PE
+  - pe-same-half-hbm: same cube, different PE (PE 1)
+  - pe-cross-half-hbm: same cube, far PE (PE 4)
+  - pe-cross-cube-hbm-best: adjacent cube (cube 1)
+  - pe-cross-cube-hbm-worst: diagonal far cube (cube 15)
+
+The cube indices 4/8/12 (H2D) and 1/4/15 (PE DMA) are meaningful for a
+4 × 4 cube mesh (`sip.cube_mesh.w=4, h=4`); changes to the mesh size
+require these to be updated in lockstep.
+
+### D2. Standard measurement size — `nbytes = 32768` (32 KiB)
+
+Every case in the summary table runs once with `nbytes=32768`. 32 KiB
+was chosen because:
+
+- DMA overhead and BW drain are balanced — neither dominates.
+- It compares cleanly against the one-shot transfer size of several
+  sub-units (TCM, register file).
+
+Per-size utilization variations are shown in a separate sweep table
+(D3).
+
+### D3. Utilization sweep — `[4 KiB, 16 KiB, 64 KiB, 256 KiB, 1 MiB]`
+
+`SWEEP_SIZES = [4096, 16384, 65536, 262144, 1048576]`,
+`SWEEP_LABELS = ["4KB", "16KB", "64KB", "256KB", "1MB"]`. Per size:
+
+```
+drain   = nbytes / bottleneck_bw
+total   = overhead + wire + drain
+eff_bw  = nbytes / total
+util%   = eff_bw / bottleneck_bw × 100
+```
+
+When `bn_bw is None or <= 0`, the column shows 0.0 %. The intent: the
+table shows in one view how small transfers become overhead-bound and
+large transfers become drain-bound as hop count rises.
+
+### D4. Measured columns — actual / formula / breakdown
+
+Per-case columns:
+
+- `Actual` (total_ns): the SimPy run's `trace["total_ns"]`.
+- `Ovhd`: sum of `node.attrs["overhead_ns"]` along the path (formula).
+- `Drain`: `nbytes / min(edge.bw_gbs over path)` (formula).
+- `Wire`: `Σ edge.distance_mm * (ns_per_mm from spec)`.
+- `Ovhd%` / `Drain%`: each portion as a percentage of Actual. Wire is
+  usually too small to display.
+- `Eff.BW`: `nbytes / total_ns` (measured BW).
+- `BN.BW`: bottleneck bandwidth (formula). The minimum edge BW along
+  the path. Missing edge BW shows "-".
+- `Util%`: `Eff.BW / BN.BW × 100`. 100 % means the single-stream BW
+  upper bound is reached.
+
+A large gap between the formula sum (`wire + ovhd + drain`) and Actual
+signals a factor the simplified model misses (a place to inspect
+ADR-0033's assumptions).
+
+### D5. Automatic invariant checks — PASS/FAIL
+
+The following invariants are reported with `[v] PASS` / `[x] FAIL`:
+
+- **H2D / D2H monotonic increase**: as hop count rises, actual latency
+  must grow monotonically. `all(lats[i] < lats[i+1] for ...)`.
+- **D2H ≥ H2D**: for the same hop index, D2H ≥ H2D (D2H has both
+  forward command and reverse data legs). `all(d2h[i].total >=
+  h2d[i].total)`.
+- **PE DMA best < worst**: cross-cube best (adjacent) latency must be
+  less than cross-cube worst (diagonal).
+- **PE DMA local vs remote**: prints the local BN BW vs remote BN BW
+  side-by-side (informational, not PASS/FAIL).
+
+When a check fails, a single clear line surfaces the regression for
+human review.
+
+### D6. Route detail — per-hop timestamp trace
+
+After the summary and sweep tables, each case's path and cumulative
+per-hop timestamps (`_hop_timestamps`) appear in a separate section:
+
+- H2D: leg1 (`pcie_ep → io_cpu`) + leg2 (`io_cpu → m_cpu`) + leg3
+  (`m_cpu → hbm_ctrl`) + per-hop trace.
+- D2H: forward (cmd, no data) and reverse (data) traces shown
+  separately.
+- PE DMA: `pe_dma → router → hbm_ctrl` path + per-hop trace.
+
+Each hop's timestamp is cumulative `wire_ns + overhead_ns`. The
+terminal hop's annotation appends `drain:Xns`. Bottleneck edges are
+marked `<BN:XXGB/s>` so they are visually identifiable.
+
+### D7. Semantics of the `case_filter` argument
+
+- `None` or `"all"`: run all cases (default).
+- Other strings: run only the case whose name matches exactly. Example:
+  `kernbench probe --case h2d-2hop`.
+
+Within a category, cases with `name != case_filter` are skipped; if
+only one data point remains, the category's monotonicity / D2H ≥ H2D
+comparisons are naturally skipped.
+
+The CLI parser's `--case` default is `"all"`, so omitting it runs
+everything.
+
+### D8. Fresh GraphEngine per case
+
+Each of the 4 H2D, 4 D2H, and 5 PE DMA cases runs in **its own
+GraphEngine** (`engine = GraphEngine(graph)`). Reasons:
+
+- Isolate accumulated state (op_log, completion tracking, allocators)
+  so cases do not cross-talk.
+- Guarantee one case's traffic does not perturb another case's BW
+  measurement.
+
+This isolation lets probe results be interpreted as **single-flow**
+per-case latency. Multi-flow contention measurement is handled by
+separate tooling (e.g., the `pe2pe_overview` plot or ADR-0033's
+multi-flow merging model).
+
+### D9. Output-format stability
+
+probe's stdout is meant for humans; precise column widths, separators,
+and whitespace are **not** a machine-readable contract. Automated tools
+that wish to parse probe output should use a separate JSON-output mode
+(not yet implemented).
+
+The `[v]` / `[x]` prefix on PASS/FAIL lines is a stable CI grep anchor.
+
+## Alternatives Considered
+
+### A1. Register probe as another bench (`@bench(name="probe")`)
+
+Rejected. probe is a verification tool, not a bench — multi-engine
+execution for sweeps/analysis and PASS/FAIL invariant output are
+essential, none of which fits ADR-0045's "single device + single
+RuntimeContext" bench model.
+
+### A2. Exit code 1 on monotonicity violation
+
+Rejected (currently). probe is positioned as a human inspection tool —
+PASS/FAIL is printed and exit is 0. A wrapper can `grep "\[x\]"` to
+decide. A future `--strict` flag could opt into non-zero exits.
+
+### A3. Externalize the case catalog to YAML
+
+Rejected (currently). The 8 cases (4 H2D + 4 D2H + 5 PE DMA = 13 total)
+are hardcoded and their semantics are tightly bound to the mesh
+topology. Moving cube-index meaning (4, 8, 12 / 1, 4, 15) into YAML
+would require separate documentation and lose cohesion. Externalize
+only when case additions become frequent.
+
+### A4. Add multi-flow contention measurement
+
+Rejected (out of probe scope). D8's single-flow isolation is probe's
+core intent. Multi-flow contention belongs in a different area of the
+ADR-0033 latency model — either a separate tool or a new case
+category.
+
+## Consequences
+
+- probe's case catalog (D1) and measurement units (D2/D3) are pinned at
+  ADR level, so new traffic categories know which table format to
+  follow.
+- The semantics of the formula-vs-actual columns (D4) are locked in, so
+  questions like "why is Drain% 5 % or 70 %?" can quickly be linked to
+  ADR-0033 assumption checks.
+- Automatic invariant checks (D5) are pinned, so latency-model changes
+  immediately catch monotonicity / D2H ≥ H2D regressions.
+- D8's case-isolation is explicit, so probe results are safe to read as
+  single-flow measurements. If multi-flow is needed, a separate tool
+  track is clearly required.
+- A2's strict-mode flag is recorded as a follow-up so CI integration
+  has a minimal change path when requested.

docs/adr/ADR-0050-par-ccl-algorithm-module-contract.md

@@ -0,0 +1,322 @@
+# ADR-0050: CCL Algorithm Module Contract — `ccl/algorithms/*.py`
+
+## Status
+
+Accepted (2026-05-22).
+
+Pins down the interface, kernel signature, and addition workflow that a
+module under `src/kernbench/ccl/algorithms/` must satisfy in order to be
+used as a collective algorithm by the AHBM CCL backend (ADR-0047).
+ADR-0047 D3 states only that "the algorithm module must expose `kernel`,
+`kernel_args`, optionally `TOPO_NAME_TO_KIND`"; **the contract an
+algorithm-module author needs to follow** has had no ADR-level coverage.
+This ADR pairs with ADR-0045's bench-module contract.
+
+## First action
+
+An algorithm module is imported at two moments:
+
+1. **AHBM backend entry**: when user code calls
+   `dist.init_process_group(backend="ahbm")`,
+   `AhbmCCLBackend.__init__` runs
+   `self._algo_module = importlib.import_module(self._merged["module"])`.
+   At module level, the following occur first:
+   - Topology-kind integer constants like `SIP_TOPO_RING/TORUS/MESH`
+     are bound in the module namespace.
+   - The `TOPO_NAME_TO_KIND` dict is bound; the backend reads it via
+     `getattr(self._algo_module, "TOPO_NAME_TO_KIND", None)`.
+   - `kernel_args` function is defined for the caller.
+   - The actual algorithm function (e.g.,
+     `allreduce_intercube_multidevice`) is defined.
+   - At the bottom of the module, `kernel = allreduce_intercube_multidevice`
+     publishes the alias.
+
+2. **ccl.yaml install stage**:
+   `kernbench.ccl.install.install_ipcq` imports the same algorithm
+   module while pushing the IPCQ neighbor table.
+
+In short, **the algorithm module's first act is "publish topology-kind
+constants, the `TOPO_NAME_TO_KIND` dict, the `kernel_args` function, and
+the `kernel` alias into the module namespace"** — all as import-time
+side effects, no separate initialization call.
+
+## Context
+
+`AhbmCCLBackend` (ADR-0047), at process-group creation, dynamically
+imports a module path obtained from `ccl.yaml`'s `defaults.algorithm` (or
+a user-specified algorithm). The backend expects four things from the
+module:
+
+- `kernel`: the collective's entry function.
+- `kernel_args(world_size, n_elem, cube_w=, cube_h=) -> tuple`: a tuple
+  packing the kernel's positional arguments.
+- `TOPO_NAME_TO_KIND` (optional): a dict mapping `topology.yaml`'s
+  `sips.topology` string (e.g., `"ring_1d"`, `"torus_2d"`,
+  `"mesh_2d_no_wrap"`) to the integer kind constants.
+- (Indirectly) IPCQ neighbor-table install:
+  `configure_sfr_intercube_multisip` reads
+  the module's `TOPO_NAME_TO_KIND` plus cube dimensions to decide the
+  SFR.
+
+The current corpus has one algorithm module:
+`lrab_hierarchical_allreduce.py` (248 lines). The name expands to
+"**l**eft-**r**ight **a**lternating **b**roadcast hierarchical allreduce".
+When future modules like `ring_allreduce`, `tree_allreduce`, or
+`broadcast` are added, they must follow this contract for the backend's
+dispatch path to keep working.
+
+Without an ADR-level contract:
+
+- A new algorithm author has to infer the signature from ADR-0047 D3's
+  one-liner.
+- The kernel-function argument order (especially `t_ptr, n_elem,
+  cube_w, cube_h, n_sips, sip_rank, sip_topo_kind, sip_topo_w,
+  sip_topo_h, tl`) is unclear without grep.
+- It is conventional, but not documented, what `kernel_args` takes as
+  inputs and what tuple it must return.
+
+## Decision
+
+### D1. The algorithm module exposes four public symbols
+
+```python
+# src/kernbench/ccl/algorithms/<name>.py
+
+from __future__ import annotations
+
+# (required) topology-kind constants — referenced internally
+SIP_TOPO_RING = 0
+SIP_TOPO_TORUS = 1
+SIP_TOPO_MESH = 2
+
+# (optional) topology name → kind mapping. Used by the backend to
+# translate ccl.yaml/topology's string SIP topology into an integer.
+TOPO_NAME_TO_KIND = {
+    "ring_1d": SIP_TOPO_RING,
+    "torus_2d": SIP_TOPO_TORUS,
+    "mesh_2d_no_wrap": SIP_TOPO_MESH,
+}
+
+# (required) kernel argument builder
+def kernel_args(world_size: int, n_elem: int, *, cube_w: int = 4, cube_h: int = 4) -> tuple:
+    return (n_elem, cube_w, cube_h, world_size)
+
+# (required) kernel function (TLContext is injected via the `tl=...`
+# keyword argument).
+def my_allreduce_kernel(t_ptr, n_elem, cube_w, cube_h, n_sips,
+                         sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h, *, tl):
+    ...
+
+# (required) kernel alias — the backend accesses `module.kernel`
+kernel = my_allreduce_kernel
+```
+
+- The `kernel` alias is the entry point the backend invokes. Whatever
+  the function name is (e.g., `allreduce_intercube_multidevice`), it
+  must be exposed via `module.kernel = fn`.
+- Without `kernel_args`, the backend has no way to build the
+  algorithm's argument list. See D2 for the signature.
+- If `TOPO_NAME_TO_KIND` is absent, the backend falls back to
+  `sip_topo_kind = 0`. An algorithm supporting only a single topology
+  may omit it.
+
+### D2. `kernel_args` signature — `(world_size, n_elem, *, cube_w, cube_h)`
+
+```python
+def kernel_args(world_size: int, n_elem: int, *,
+                cube_w: int = 4, cube_h: int = 4) -> tuple:
+    return (n_elem, cube_w, cube_h, world_size)
+```
+
+- **Positional arguments**: `world_size` (= number of ranks), `n_elem`
+  (= element count of a single shard, f16-based).
+- **Keyword arguments**: `cube_w`, `cube_h` (= cube-mesh dimensions).
+  Default 4×4 — aligned with `topology.yaml`'s `sip.cube_mesh` default.
+- **Return**: a tuple in the order the kernel's positional arguments
+  expect.
+
+When the backend calls `all_reduce`:
+
+```python
+kernel_args_tuple = self._algo_module.kernel_args(
+    self._world_size, n_elem, cube_w=eff_cube_w, cube_h=eff_cube_h,
+)
+extra_args = (sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)
+pending = self.ctx.launch(
+    self._merged["algorithm"], kernel_fn, tensor,
+    *kernel_args_tuple, *extra_args, _defer_wait=True,
+)
+```
+
+So the kernel's full positional argument list becomes: `(tensor_ptr,
+*kernel_args_tuple, sip_rank, sip_topo_kind, sip_topo_w,
+sip_topo_h)`, with `tl=...` injected as a keyword. The tuple length
+and order returned by `kernel_args` must **match the kernel signature
+1:1**.
+
+### D3. Kernel signature — standardized 9 + tl arguments
+
+Recommended signature:
+
+```python
+def my_kernel(
+    t_ptr: int,             # VA base of the row-wise-sharded tensor on this SIP
+    n_elem: int,            # element count per cube tile (or per shard)
+    cube_w: int,            # cube mesh width  (from kernel_args)
+    cube_h: int,            # cube mesh height (from kernel_args)
+    n_sips: int,            # equal to world_size (rank = SIP, ADR-0024)
+    sip_rank: int,          # this SIP's rank
+    sip_topo_kind: int,     # result of TOPO_NAME_TO_KIND lookup
+    sip_topo_w: int,        # SIP mesh width (0 for ring_1d)
+    sip_topo_h: int,        # SIP mesh height (0 for ring_1d)
+    *, tl,                  # TLContext (auto-injected)
+) -> None:
+```
+
+Even if `kernel_args` chose a different positional argument order, the
+kernel's **last four positional arguments are always
+`(sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)`** — the backend
+appends them as `extra_args` (ADR-0047 D5). A custom algorithm must
+accept these four, but a single-SIP algorithm may simply ignore them.
+
+`tl` is injected via keyword — `RuntimeContext.launch` adds `tl=tl_ctx`
+just before invoking the kernel. The signature therefore exposes `tl`
+as keyword-only (`*, tl`) or as the trailing keyword parameter.
+
+### D4. Kernel body — freedom and constraints
+
+Available inside the kernel: every `tl.*` primitive from ADR-0046 D3.
+
+Common patterns:
+
+- `cube_id = tl.program_id(axis=1)` — this PE's cube index.
+- `pe_addr = t_ptr + cube_id * nbytes` — per-cube VA of the tile.
+- `acc = tl.load(pe_addr, shape=(n_elem,), dtype="f16")` — load local
+  data.
+- `tl.send(dir=...)` / `tl.recv(dir=..., shape=, dtype=)` — IPCQ
+  collective.
+- `acc = acc + recv` — TensorHandle arithmetic operators (ADR-0046 D4).
+- `tl.store(pe_addr, acc)` — store the result.
+
+The kernel body is plain Python — branching and loops are fine. But:
+
+- No SimPy `yield` or `async` (ADR-0046 D1).
+- No direct access to TensorHandle `.data` — the Phase 1 timing model
+  doesn't see data dependencies (ADR-0020's 2-pass separation).
+- Kernel execution must be deterministic — the same input must produce
+  the same op sequence. No random or external IO.
+
+### D5. SIP topology semantics — meaning of `sip_topo_kind`
+
+The backend looks up `topology.yaml`'s `system.sips.topology` string
+in the algorithm module's `TOPO_NAME_TO_KIND` and passes the integer
+as `sip_topo_kind`. The algorithm then branches:
+
+```python
+if sip_topo_kind == SIP_TOPO_RING:
+    acc = _inter_sip_ring(...)
+elif sip_topo_kind == SIP_TOPO_TORUS:
+    acc = _inter_sip_torus_2d(...)
+elif sip_topo_kind == SIP_TOPO_MESH:
+    acc = _inter_sip_mesh_2d(...)
+```
+
+Each topology branch communicates with peers via IPCQ direction names
+(`"global_E"`, `"W"`, `"S"`, `"N"` …). Direction semantics are defined
+in ADR-0023/0025; `configure_sfr_intercube_multisip` installs the IPCQ
+neighbor table accordingly.
+
+If a topology kind not supported by the algorithm appears, prefer an
+explicit `raise ValueError(f"unsupported topology kind
+{sip_topo_kind}")` over a silent no-op — fail fast on misconfiguration.
+
+### D6. The `ccl.yaml` algorithm entry
+
+The algorithm module is paired with a `ccl.yaml` entry (ADR-0023 D10 +
+ADR-0047 D3):
+
+```yaml
+defaults:
+  algorithm: lrab_hierarchical_allreduce
+  n_elem: 8
+
+algorithms:
+  lrab_hierarchical_allreduce:
+    module: kernbench.ccl.algorithms.lrab_hierarchical_allreduce
+    # optional: world_size override
+    # optional: per-algorithm parameters consumed by configure_sfr_intercube_multisip
+```
+
+- `module`: the full Python module path; `importlib.import_module`
+  consumes this string as-is.
+- `world_size` (optional): when set, overrides the topology fallback
+  (ADR-0047 D2).
+- Algorithm-specific parameters are consumed by
+  `configure_sfr_intercube_multisip`.
+
+Workflow to add a new algorithm:
+
+1. Write `src/kernbench/ccl/algorithms/<name>.py` following D1.
+2. Add the entry under `algorithms` in `ccl.yaml`.
+3. (If needed) extend `kernbench.ccl.sfr_config` with the SFR-install
+   branch.
+4. Add tests (e.g., `tests/sccl/test_<name>.py`, extending the
+   ADR-0043 eval harness).
+
+### D7. Legacy "rank = flat PE index" mode
+
+The `world_size` override in `ccl.yaml`, surfaced by ADR-0047 D2, is
+used by legacy "rank = flat PE index" tests. The algorithm module can
+assume `n_sips=world_size` ranks even in this mode — the backend
+maintains the rank↔(SIP, cube, PE) mapping, so no modal branching is
+needed inside the algorithm body.
+
+In single-cube workloads (where `cube_w=cube_h=1`), the algorithm must
+skip mesh-based phases — see the
+`single_cube = (cube_w == 1 and cube_h == 1)` pattern in
+`lrab_hierarchical_allreduce.py`.
+
+## Alternatives Considered
+
+### A1. Organize the algorithm module as a class (`class Allreduce: kernel(...)`)
+
+Rejected. The Python module namespace already identifies an algorithm
+(see ADR-0047 D3's `importlib.import_module`). A class wrapper adds
+indirection without simplifying dispatch. Module-level free functions
+plus a `kernel` alias are clean and obvious.
+
+### A2. Type `kernel_args` with an explicit dataclass
+
+Rejected (currently). Each algorithm normally has a different argument
+count; forcing one dataclass would hurt cross-algorithm interchange.
+The tuple return is simple and unpacks cleanly with the backend's
+`*kernel_args_tuple`. If an algorithm wants stronger internal typing,
+it may define its own NamedTuple.
+
+### A3. Move SFR installation inside the algorithm module
+
+Rejected. SFR installation
+(`configure_sfr_intercube_multisip`) is a cross-module decision
+combining topology + algorithm; `kernbench.ccl.sfr_config` is a more
+natural home than the algorithm module itself. D6's "extend
+sfr_config if needed" workflow keeps responsibility boundaries clear.
+
+### A4. Auto-register algorithm names via a decorator (analogous to ADR-0045's `@bench`)
+
+Rejected. Unlike benches, algorithms are already tied to `ccl.yaml`
+entries; an additional registry would be redundant. The string mapping
+in `module` is sufficient.
+
+## Consequences
+
+- ADR-0047 D3's one-line contract expands to a D1–D7 author-facing
+  guide; new algorithm signatures no longer need to be grep-derived.
+- D3's standardized 9 + tl signature couples naturally with the
+  backend's `extra_args` append (ADR-0047 D5). It is explicit that
+  even single-SIP-only algorithms must accept the four `sip_*` trailing
+  arguments.
+- D5's fail-loud recommendation means a `ccl.yaml` topology that the
+  algorithm doesn't support will surface as an explicit `ValueError`
+  rather than a silent wrong result.
+- D6's step-by-step addition workflow makes clear how far a new
+  algorithm has to reach into sfr_config / tests / ccl.yaml.

docs/adr/ADR-0051-lat-routing-helper-api.md

@@ -0,0 +1,288 @@
+# ADR-0051: Routing Helper API — `AddressResolver` + `PathRouter`
+
+## Status
+
+Accepted (2026-05-22).
+
+Pins down every public API, argument, return value, and adjacency-graph
+selection of the two helper classes (`AddressResolver`, `PathRouter`)
+exposed by `policy/routing/router.py`. ADR-0002 defines routing
+distance, ordering, and bypass rules, but **the helper API surface
+itself** has had no ADR-level coverage.
+
+## First action
+
+### `AddressResolver(graph)`
+
+On construction, caches two pieces of state:
+
+1. `self._node_ids = set(graph.nodes)` — a set of all node ids for
+   lookup.
+2. `self._hbm_slice_bytes = hbm_total_gb * (1 << 30) // slices_per_cube`
+   — derived from `graph.spec.cube.memory_map` (default `48 GB / 8
+   slices = 6 GB`). `resolve()` uses this value to decode `pe_id` from
+   an HBM PA's `hbm_offset`.
+
+In short, **AddressResolver's first act is "precompute the full set of
+node ids and the HBM slice size"**. It does not retain the graph
+itself.
+
+### `PathRouter(graph)`
+
+On construction, **builds four separate adjacency graphs in one pass**:
+
+1. `self._adj_all`: every edge (used for component-to-component
+   routing).
+2. `self._adj`: edges with `kind != "command"` (PE DMA / generic data
+   paths).
+3. `self._adj_mcpu_dma`: excludes
+   `_MCPU_DMA_EXCLUDE = {"pe_internal", "pe_to_router"}` (M_CPU DMA
+   must not pass through PE pipeline nodes).
+4. `self._adj_local`: excludes the 8-element `_UCIE_KINDS` set (UCIe
+   would look like a zero-distance bus to Dijkstra, which would prefer
+   it over the mesh — for cube-local routing this must be avoided).
+
+Each graph is a `defaultdict(list)` of `(neighbor, weight)`. The
+weight is `edge.routing_weight_mm or edge.distance_mm`.
+
+In short, **PathRouter's first act is "classify topology edges into
+four policy-specific adjacency lists simultaneously"**. Each `find_*()`
+call picks the appropriate graph and runs Dijkstra.
+
+## Context
+
+`policy/routing/router.py` performs two responsibilities together:
+
+- **Naming**: it is the sole owner of the topology naming convention
+  (`sip{S}.cube{C}.<comp>`, `sip{S}.io{I}.pcie_ep`, etc.). Components /
+  probe / IPCQ install / runtime API do not build node-id strings
+  themselves — they call helpers.
+- **Path decisions**: policy separation by `edge.kind`. For the same
+  src→dst, different routing intents (PE DMA vs M_CPU DMA vs general
+  component routing) call for different adjacencies and so produce
+  different paths.
+
+This helper API is widely consumed (probe.py / distributed.py /
+install.py / various components / tests), yet **the exact signatures /
+return semantics / adjacency picks** are not gathered in any ADR. This
+ADR closes that gap.
+
+## Decision
+
+### D1. `AddressResolver` exposes five public methods
+
+#### D1.1. `resolve(addr: PhysAddr) -> str`
+
+Translates a `PhysAddr` to a destination node id in the topology:
+
+```
+addr.kind == "hbm"             → f"sip{s}.cube{d}.hbm_ctrl.pe{pe_id}"
+  where pe_id = addr.hbm_offset // self._hbm_slice_bytes  (ADR-0017 D4/D9)
+
+addr.kind == "pe_resource":
+  addr.unit_type == PE         → f"sip{s}.cube{d}.pe{addr.pe_id}.pe_tcm"
+  addr.unit_type == SRAM       → f"sip{s}.cube{d}.sram"
+  addr.unit_type == MCPU       → f"sip{s}.cube{d}.m_cpu"
+  others                       → RoutingError("unsupported unit_type")
+
+other kinds                    → RoutingError("unsupported address kind")
+```
+
+If the derived node id is not in `self._node_ids`, raises
+`RoutingError(f"node {node_id} not found in topology")`. So even when
+the address has valid syntax, an absent node in the topology
+fails-loud.
+
+#### D1.2. `find_m_cpu(sip, cube) -> str`
+
+Returns `f"sip{sip}.cube{cube}.m_cpu"`; absent → `RoutingError`.
+
+#### D1.3. `find_pcie_ep(sip, io_id="io0") -> str`
+
+Returns `f"sip{sip}.{io_id}.pcie_ep"`; absent → `RoutingError`.
+
+#### D1.4. `find_io_cpu(sip, io_id="io0") -> str`
+
+Returns `f"sip{sip}.{io_id}.io_cpu"`; absent → `RoutingError`.
+
+#### D1.5. `find_all_pcie_eps() -> list[str]`
+
+All PCIE_EP node ids across all SIPs, sorted. Filtered by
+`endswith(".pcie_ep")`. Cross-SIP IPCQ uses this when enumerating
+PCIE_EPs.
+
+This class is the sole owner of the naming convention
+(`sip{S}.cube{C}.<comp>`, `sip{S}.{io_id}.<comp>`) — ADR-0015 D4.
+The topology builder produces nodes with the same naming convention;
+components never build node-id strings directly — they go through
+these helpers.
+
+### D2. `PathRouter`'s four adjacency graphs
+
+Constructed in one pass. `edge.kind` drives policy:
+
+| graph             | excluded edge kinds                                                                                                                   | use case                                       |
+|-------------------|--------------------------------------------------------------------------------------------------------------------------------------|-------------------------------------------------|
+| `_adj_all`        | (none)                                                                                                                               | M_CPU↔NOC command included, IO_CPU/M_CPU routes |
+| `_adj`            | `"command"`                                                                                                                          | PE DMA / generic data paths                     |
+| `_adj_mcpu_dma`   | `"pe_internal"`, `"pe_to_router"`                                                                                                    | M_CPU DMA (skips PE pipeline)                   |
+| `_adj_local`      | `_UCIE_KINDS` (`ucie_internal`, `ucie_conn_to_router`, `router_to_ucie_conn`, `ucie_conn_to_noc`, `noc_to_ucie_conn`, `ucie_mesh`, `io_to_cube`, `cube_to_io`) | same-cube routing (UCIe bus excluded) |
+
+Each graph is `dict[node_id, list[(neighbor, weight)]]` with weight =
+`edge.routing_weight_mm or edge.distance_mm`. Excluding command edges
+prevents them from influencing routing; isolating `_adj_local` keeps
+UCIe's "zero-distance bus" from out-competing the mesh — consistent
+with ADR-0017 D7's cross-PE-slice mesh-distance requirement.
+
+### D3. `PathRouter` exposes six public methods (+ two backward-compat shims)
+
+#### D3.1. `find_path(src_pe: str, dst_node: str) -> list[str]`
+
+**PE DMA routing**. `src_pe` is a PE prefix (e.g.,
+`"sip0.cube0.pe0"`); the function auto-prepends `.pe_dma`, making the
+true start node `"sip0.cube0.pe0.pe_dma"`.
+
+Adjacency depends on cube-locality (`_same_cube`):
+
+- **Same-cube** (src and dst share `sip{S}.cube{C}.` prefix): uses
+  `_adj_local`. Excluding UCIe lets cross-PE-slice access pay accurate
+  mesh distance (ADR-0017 D7).
+- **Cross-cube**: uses `_adj`. UCIe naturally becomes the right choice
+  for the cross-cube portion.
+
+#### D3.2. `find_path_with_distance(src_pe, dst_node) -> tuple[list[str], float]`
+
+Same adjacency policy as D3.1, but returns `(path, total_distance)`.
+Used by probe and analysis tools that need the distance metric.
+
+#### D3.3. `find_mcpu_dma_path(m_cpu_id: str, dst_hbm_id: str) -> list[str]`
+
+**M_CPU DMA path**. Same cube → `_adj_local` (stay within the mesh);
+different cube → `_adj_all` (cross via UCIe). The
+`_MCPU_DMA_EXCLUDE` set ensures PE-pipeline nodes never appear on
+M_CPU's routes.
+
+#### D3.4. `find_memory_path(src: str, dst: str) -> list[str]`
+
+Direct memory path like
+`pcie_ep → io_noc → cube → router mesh → hbm_ctrl`. Uses
+`_adj_mcpu_dma` to exclude `pe_internal` and `pe_to_router`, so
+host-issued reads/writes never leak into the PE pipeline. Probe
+(ADR-0049 D1's H2D/D2H cases) calls this directly.
+
+#### D3.5. `find_node_path(src: str, dst: str) -> list[str]`
+
+Generic routing between arbitrary nodes, **including command edges**
+(via `_adj_all`). IoCpuComponent / MCpuComponent use this when they
+need to route through M_CPU ↔ NOC command-kind links.
+
+#### D3.6. Backward-compat shims
+
+- `_dijkstra(start, goal) -> list[str]` — thin wrapper for
+  `_run_dijkstra(self._adj, …)`.
+- `_dijkstra_with_dist(start, goal) -> tuple[list[str], float]` —
+  distance-aware variant.
+
+Despite the underscore prefixes (suggesting internal API), existing
+tests call these directly. New code should prefer D3.1–D3.5; these two
+shims are deprecation candidates.
+
+### D4. Dijkstra — single-source shortest path
+
+`_run_dijkstra_with_dist(adj, start, goal)`:
+
+- `heapq` priority queue.
+- `best: dict[node, distance]` — best known distance to each node.
+- `prev: dict[node, predecessor]` — for path reconstruction.
+- Edge weight = `routing_weight_mm or distance_mm`. The separation
+  matters because UCIe (and a few others) declare an explicit
+  `routing_weight_mm` distinct from physical `distance_mm`.
+
+`start == goal` short-circuits to `([start], 0.0)`. Unreachable target
+→ `RoutingError(f"no path from {start} to {goal}")`.
+
+The algorithm is **deterministic**: identical graph + start/goal gives
+the same path, satisfying SPEC R1 ("routing MUST be deterministic").
+Tie-breaks follow `heapq`'s push order (Python list order is
+deterministic).
+
+### D5. Single-owner principle for helper-API decisions
+
+The following decisions live only inside router.py:
+
+- Naming convention: `sip{S}.cube{C}.<comp>`,
+  `sip{S}.{io_id}.<comp>`,
+  `sip{S}.cube{C}.hbm_ctrl.pe{pe_id}`.
+- Adjacency policy: which edge kinds belong to which graph.
+- Algorithm for recovering PE id from an HBM slice size.
+- Dijkstra weight selection
+  (`routing_weight_mm or distance_mm`).
+
+Breaking single ownership (e.g., a component starting to build
+`f"sip{s}..."` itself) would explode the blast radius of naming-
+convention changes. This aligns with ADR-0015 D4.
+
+### D6. Consumers of the helper API
+
+Methods listed in this ADR are called from (current corpus):
+
+- `probes/probe.py` (ADR-0049): `find_pcie_ep`, `find_io_cpu`,
+  `find_m_cpu`, `find_node_path`, `find_mcpu_dma_path`,
+  `find_memory_path`, `find_path`, `resolve`.
+- `runtime_api/distributed.py` (ADR-0047): indirectly (engine-internal
+  routing).
+- `ccl/install.py` (ADR-0023): `find_all_pcie_eps`, `resolve`.
+- `sim_engine/event_log.py`: like probe — `find_pcie_ep`,
+  `find_memory_path`.
+- `components/builtin/m_cpu.py`, `components/builtin/io_cpu.py`:
+  `find_node_path`, `find_mcpu_dma_path`.
+- Tests (test_routing.py, test_cross_sip_routing.py, …): most of
+  D3.1–D3.5.
+
+When a new consumer arrives, D1/D3 act as a first-pass guide on
+whether an existing method matches the intent or a new one is needed.
+
+## Alternatives Considered
+
+### A1. One adjacency graph + per-call edge-kind filtering
+
+Rejected. Re-filtering the graph on every `find_*()` call hurts
+Dijkstra cache locality. Constructing four graphs in one pass (D2)
+has modest memory cost (edges ≤ a few × 10⁴), and selection happens
+in O(1) at call time.
+
+### A2. Drive adjacency separation by separate edge metadata rather than `kind`
+
+Rejected. `edge.kind` is already assigned by the topology builder
+(ADR-0015 D4 + ADR-0017); a parallel metadata field would force
+synchronization between two systems.
+
+### A3. Use BFS with uniform weights instead of Dijkstra
+
+Rejected. With per-edge `routing_weight_mm` (mesh link / UCIe /
+IO-internal), BFS minimizes hop count rather than total
+latency/distance. SPEC R1 + R2 require deterministic and accurate
+routing, which BFS does not deliver.
+
+### A4. Express the helper API as module functions instead of classes
+
+Rejected. Each class
+(`AddressResolver`, `PathRouter`) maintains caches
+(`_node_ids`, `_hbm_slice_bytes`, four adjacency graphs) reused across
+many routing queries on the same graph. Module functions would have
+to rebuild state per call or go global, hurting safety and
+performance.
+
+## Consequences
+
+- When components / probe / IPCQ install / runtime API all go through
+  router.py helpers, a naming-convention change (e.g., `.io0.` →
+  `.iochiplet0.`) is a one-file edit (D5).
+- D2's four-graph split is now ADR-locked, so when a new edge kind is
+  added (e.g., a new inter-die UCIe-link kind), the right adjacency
+  category is decided explicitly rather than by default.
+- D3.1's same-cube vs cross-cube branching (ADR-0017 D7) is explicit,
+  so anyone changing routing knows which adjacency to touch.
+- D6's consumer list bounds PR-review scope for helper-API changes,
+  and the backward-compat shims (D3.6) are flagged as deprecation
+  candidates.

docs/adr/ADR-0052-dev-oplog-memory-store-schemas.md

@@ -0,0 +1,371 @@
+# ADR-0052: OpLog + MemoryStore Schemas — sim_engine internals
+
+## Status
+
+Accepted (2026-05-22).
+
+Pins down the `OpRecord` schema and the `record_start` / `record_end` /
+`record_copy` behavior in `sim_engine/op_log.py`, plus the
+(space, addr) namespace and read/write semantics of `MemoryStore` in
+`sim_engine/memory_store.py`. ADR-0020 (2-pass data execution) declares
+that these two facilities exist, but **the precise record fields and
+semantics** had no ADR-level coverage, and several recent ADRs
+(ADR-0046 D3.2's `tl.store` visibility, ADR-0023 D9's IPCQ copy
+record) depend on these semantics.
+
+## First action
+
+### `OpLogger(memory_store=None)`
+
+On construction, initialize three fields:
+
+1. `self._records: list[OpRecord] = []` — accumulated records.
+2. `self._pending: dict[int, dict] = {}` — partial records keyed by
+   `id(msg)` (created at `record_start`, completed at `record_end`).
+3. `self._memory_store = memory_store` — optional MemoryStore
+   reference. Used to capture math-op input snapshots and dma_write
+   HBM-source snapshots.
+
+Records and pending are empty; the `record_*` calls accumulate data
+over time.
+
+### `MemoryStore()`
+
+On construction, initialize a single field:
+`self._storage: dict[str, dict[int, np.ndarray]] = {}` — a two-level
+dict (`space → addr → ndarray`). Inner dicts are created lazily as new
+spaces appear.
+
+In short, **both facilities' first act is "set up an empty accumulator
+buffer plus a sparse, per-space dict"**. The first record / write
+fills the fields when it arrives.
+
+## Context
+
+ADR-0020 D2/D5/D7 (2-pass data execution) declares:
+
+- During Phase 1 (timing), `ComponentBase._on_process_start/end` hooks
+  call `OpLogger.record_start/end`, recording the time and metadata of
+  every data op.
+- Phase 2 (data) replays the op log in `t_start` order to compute real
+  data.
+- Data payloads live in `MemoryStore`, keyed by (space, addr).
+
+Subsequent ADRs (ADR-0023 D9's IPCQ atomic write, ADR-0027's Megatron
+TP scratch-overwrite avoidance, ADR-0046 D3.2's `tl.store` visibility)
+depend on op_log and MemoryStore behavior, but **the exact record
+fields / space names / snapshot timing** are only discoverable via
+source grep. This ADR codifies them.
+
+## Decision
+
+### D1. `OpRecord` schema — seven fields
+
+```python
+@dataclass
+class OpRecord:
+    t_start: float
+    t_end: float
+    component_id: str
+    op_kind: str               # "memory" | "gemm" | "math" | "unknown"
+    op_name: str               # e.g. "dma_read", "gemm_f16", "exp",
+                               #     "TileToken/DMA_READ", "composite_gemm",
+                               #     "ipcq_copy"
+    params: dict[str, Any]
+    dependency_ids: list[int] = field(default_factory=list)
+```
+
+- **`t_start` / `t_end`**: SimPy time (float ns). `t_start` is when the
+  component begins the op; `t_end` is completion. Duration =
+  `t_end - t_start`.
+- **`component_id`**: the node id where the op occurred (e.g.,
+  `"sip0.cube0.pe0.pe_dma"`).
+- **`op_kind`**: one of four. Phase 2 DataExecutor branches on this.
+- **`op_name`**: a debug/analysis-friendly name. For a TileToken,
+  expands to `"TileToken/{stage_type}"` (e.g.,
+  `"TileToken/DMA_READ"`) to disambiguate stages.
+- **`params`**: op-specific metadata dict (see D3).
+- **`dependency_ids`**: currently unused (default `[]`). Reserved for
+  future cross-op dependency tracking.
+
+### D2. `OpLogger.records` — guaranteed `t_start` sort
+
+```python
+@property
+def records(self) -> list[OpRecord]:
+    self._records.sort(key=lambda r: r.t_start)
+    return self._records
+```
+
+A stable sort by `t_start` runs on each access. Records with the same
+`t_start` preserve insertion order. Aligns with ADR-0020 D5's
+"t_start stable ordering" requirement.
+
+Phase 2 DataExecutor always accesses via the `records` property, so
+even when `record_end` calls arrive out of `t_start` order (e.g., a
+short op started later but finished earlier), the sequence handed to
+Phase 2 is consistent.
+
+### D3. `params` schema per `op_name` (matrix from `_extract_op_info`)
+
+#### D3.1. `op_kind="memory", op_name="dma_read"` (DmaReadCmd)
+
+```python
+{"src_addr": int, "nbytes": int, "handle_id": str}
+```
+
+#### D3.2. `op_kind="memory", op_name="dma_write"` (DmaWriteCmd)
+
+```python
+{
+    "src_space": str,   # handle.space ("tcm"|"hbm"|"sram"), default "tcm"
+    "src_addr": int,    # handle.addr
+    "shape": tuple, "dtype": str,
+    "dst_space": "hbm", # DmaWrite always targets HBM
+    "dst_addr": int,
+    "nbytes": int,
+    "handle_id": str,
+    # When src_space == "hbm" at record_end, a snapshot is added (D4)
+    "snapshot": np.ndarray | None,
+}
+```
+
+#### D3.3. `op_kind="gemm", op_name=f"gemm_{dtype_a}"` (GemmCmd)
+
+```python
+{
+    "src_a_addr": int, "src_b_addr": int, "dst_addr": int,
+    "shape_a": tuple, "shape_b": tuple, "shape_out": tuple,
+    "dtype_in": str, "dtype_out": str,
+    "m": int, "k": int, "n": int,
+    # ADR-0027: per-operand + output spaces preserved
+    "src_a_space": str, "src_b_space": str, "dst_space": str,
+}
+```
+
+#### D3.4. `op_kind="math", op_name=msg.op` (MathCmd; op = "exp", "sum", "add", "where", …)
+
+```python
+{
+    "input_addrs": list[int],   # addrs of input handles
+    "input_shapes": list[tuple],
+    "input_spaces": list[str],
+    "input_dtypes": list[str],
+    "dst_addr": int, "dst_space": str,
+    "shape_out": tuple, "dtype": str,
+    "axis": int | None,         # only meaningful for reductions
+    # All inputs get snapshots at record_end (D4)
+    "input_snapshots": list[np.ndarray | None],
+}
+```
+
+#### D3.5. `op_kind="gemm" or "math", op_name=f"composite_{op}"` (CompositeCmd)
+
+```python
+{
+    "op": str,              # "gemm" | "math"
+    "out_addr": int, "out_nbytes": int,
+    # If op == "gemm", same fields as GemmCmd are added:
+    "src_a_addr": int, "src_b_addr": int,
+    "shape_a": tuple, "shape_b": tuple,
+    "dtype_in": str, "dtype_out": str,
+    "src_a_space": str, "src_b_space": str,
+    "dst_space": "hbm", "dst_addr": int,  # = out_addr
+}
+```
+
+If `op == "gemm"`, `op_kind = "gemm"`; otherwise `"math"`. An alias so
+Phase 2 replays composite-gemm on the same path as `GemmCmd`.
+
+#### D3.6. `op_kind="memory", op_name="ipcq_copy"` (record_copy path)
+
+```python
+{
+    "src_space": str, "src_addr": int,
+    "dst_space": str, "dst_addr": int,
+    "shape": tuple, "dtype": str, "nbytes": int,
+    "snapshot": np.ndarray | None,   # passed by caller; if None, record_copy reads fresh
+}
+```
+
+`PE_DMA._handle_ipcq_inbound` (ADR-0023 D9) emits this record so Phase
+2 can replay the IPCQ slot's inbound copy. It bypasses
+`record_start` / `record_end` and pushes directly via `record_copy()`.
+
+#### D3.7. `op_kind="unknown", op_name=type(msg).__name__`
+
+Fallback for messages `_extract_op_info` doesn't recognize. `params =
+{}`. If DataExecutor encounters this kind, it skips — Phase 2 replay
+is unaffected.
+
+### D4. Snapshot capture timing
+
+When `OpLogger._memory_store` is set, `record_end` performs:
+
+- **Math op**: read every input
+  (addr/shape/space/dtype) from `self._memory_store.read(...)` and
+  attach an ndarray copy to `params["input_snapshots"]`. Read failure
+  → `None`.
+- **`dma_write` op**: snapshot the source **only if `src_space ==
+  "hbm"`** and attach to `params["snapshot"]`. TCM (PE scratch)
+  sources are **deliberately skipped** — TCM is repopulated by Phase 2
+  math/gemm replay, and a Phase-1-time snapshot would capture a
+  previous kernel's stale value (ADR-0027 postmortem: TP gemm →
+  all_reduce race).
+- **`ipcq_copy`**: the caller passes the in-flight snapshot via
+  `snapshot=token.data`. If absent, `record_copy` attempts a fresh
+  read from MemoryStore.
+
+Snapshots are taken with `.copy()` (fresh allocation), making them
+safe against later storage mutation. This is the foundation of
+ADR-0027's "cross-PE Phase 2 ordering" race-avoidance.
+
+When `memory_store` is `None` (Phase 1 timing-only mode), all
+snapshot steps are skipped. Only the timing portion of the record is
+preserved; data replay is unavailable.
+
+### D5. TileToken handling — `record_start` captures stage info
+
+ADR-0014 D6's self-routing tile token (pipeline mode) may have already
+advanced its `stage_idx` by the time `record_end` runs (the TileToken
+caches the next stage's params as it moves to the next component).
+Therefore:
+
+`record_start` pre-saves the following in `pending[id(msg)]["snap"]`:
+
+```python
+snap["stage_type"] = stage.stage_type.name        # "DMA_READ", "GEMM", ...
+snap["stage_params"] = dict(stage.params)         # copy of params at start time
+```
+
+`record_end` retrieves this snap and merges into params:
+
+- Adds `params["stage_type"]` to final params.
+- Merges `stage_params` keys (keeps existing values if any).
+- If `op_name == "TileToken"`, rewrites it to
+  `f"TileToken/{stage_type}"` (e.g., `"TileToken/DMA_READ"`),
+  disambiguating different stages emitted by the same component.
+
+Thanks to this, DMA_READ vs DMA_WRITE, FETCH vs STORE coming from the
+same component (e.g., pe_dma) are distinguishable in reports.
+
+### D6. `MemoryStore` — two-level (space, addr) dict
+
+```python
+class MemoryStore:
+    def __init__(self) -> None:
+        self._storage: dict[str, dict[int, np.ndarray]] = {}
+
+    def write(self, space, addr, data): self._storage[space][addr] = data
+    def read(self, space, addr, shape=None, dtype=None) -> np.ndarray: ...
+    def has(self, space, addr) -> bool: ...
+    def snapshot(self) -> MemoryStore: ...
+```
+
+#### D6.1. Space namespace
+
+A string key. Standard values:
+
+- `"hbm"`: HBM data (deploy_tensor + Phase 2 dma_write results).
+- `"tcm"`: PE-local TCM (Phase 2 math/gemm output).
+- `"sram"`: cube-level SRAM (ADR-0023 D9.7's IPCQ slot tier).
+
+Other spaces (e.g., `"reg"`) are allowed — `_storage` is a lazy dict
+that creates a new space when `write` first touches it.
+
+#### D6.2. Address keying
+
+`addr` is an integer. It may be a **physical address (PA) or a virtual
+address (VA)** — `MemoryStore` itself doesn't know address-space
+semantics; it just uses them as keys. Phase 1's `MemoryWriteMsg`
+writes both PA and VA
+(`_create_tensor` zero-inits at PA and at the VA base too); Phase 2
+reads/writes via the addresses captured by op_log.
+
+The caller decides `addr`'s meaning — `MemoryStore` provides only
+lookup.
+
+#### D6.3. read/write semantics — reference store (no copy)
+
+`write(space, addr, data)`: stores the ndarray reference. **No copy.**
+If the caller later mutates the same ndarray, the stored value
+changes.
+
+`read(space, addr, shape=None, dtype=None)`: returns the stored
+ndarray reference. If `shape`/`dtype` are provided:
+
+- `dtype != stored.dtype`: `arr.view(np_dtype)` reinterprets as a
+  view (no copy).
+- `shape != stored.shape`: if `nbytes` matches, `arr.reshape(shape)`
+  is a view.
+- `nbytes` mismatch → `ValueError`.
+
+To detach the data, the caller must call `arr.copy()`. ADR-0027's
+race-avoidance requires explicit `.copy()` in op_log snapshot steps
+for exactly this reason.
+
+#### D6.4. `has(space, addr) -> bool`
+
+Existence check; does not materialize data.
+
+#### D6.5. `snapshot() -> MemoryStore`
+
+Shallow copy. Creates a new instance of inner dicts but shares
+ndarray references. Used at Phase 2 init to fork from Phase 1's
+store, so Phase 2 mutations don't affect Phase 1's remaining
+consumers.
+
+### D7. op_log assumes a single-threaded SimPy
+
+`OpLogger`'s `_records` and `_pending` are lock-free. SimPy is
+single-threaded, so nothing else can intrude between `record_start`
+and `record_end` for the same message.
+
+When multi-process kernbench (ADR-0047 D6) arrives, OpLogger must be
+split per process — one OpLogger instance cannot receive records from
+multiple processes.
+
+## Alternatives Considered
+
+### A1. Externalize op_log to SQLite / parquet
+
+Rejected (currently). The in-memory list minimizes Phase 1 → Phase 2
+hand-off latency. Externalization makes sense for long-running batch
+runs but adds overhead for the current single-run workload.
+
+### A2. Capture snapshots at `record_start`
+
+Rejected. At `record_start`, inputs are often not yet populated (e.g.,
+a math op's input is the output of a just-issued previous op).
+`record_end` is the correct point.
+
+### A3. Per-component MemoryStore
+
+Rejected. The (space, addr) key already disambiguates effectively, and
+splitting per component would complicate cross-PE IPCQ copy (ADR-0023
+D9), which needs access to both source and destination stores.
+
+### A4. Explicit dependency edges in op_log
+
+Partially adopted. The `dependency_ids` field exists on `OpRecord` but
+is currently unused (D1). Phase 2 DataExecutor orders via `t_start` +
+a secondary sort (memory ops before math at the same `t_start`). When
+an explicit dependency graph is required, this field is the home.
+Current ordering rules are sufficient, so it remains unused.
+
+## Consequences
+
+- ADR-0020's op_log / MemoryStore declarations are expanded into the
+  concrete D1–D6 schemas, so writing/modifying Phase 2 DataExecutor
+  doesn't need source-grep to learn field semantics.
+- D3's per-`op_name` params matrix makes adding new ops (e.g., a new
+  reduction type) a question of branching in `_extract_op_info`.
+- D4's per-op snapshot policy (math = input snapshot, dma_write =
+  HBM-only snapshot) is ADR-locked, so ADR-0027's race-avoidance
+  decision won't silently regress on future refactors.
+- D6.3's reference-store semantics are explicit, putting mutation
+  safety on the caller. ADR-0027's explicit `.copy()` pattern is
+  justified.
+- D7's single-thread assumption is recorded, so multi-process
+  kernbench (ADR-0047 D6's supersession candidate) will need OpLogger
+  separation when introduced.

docs/adr/ADR-0053-dev-topology-builder-algorithms.md

@@ -0,0 +1,351 @@
+# ADR-0053: Topology Builder + Visualizer Algorithms
+
+## Status
+
+Accepted (2026-05-22).
+
+Pins down the key algorithmic choices of the topology compile and
+visualization pipeline jointly implemented by `topology/builder.py`,
+`topology/mesh_gen.py`, and `topology/visualizer.py` —
+placement-driven router attachment, mesh auto-layout, the source_hash
+cache, view projections, and SVG rendering. ADR-0006 defines the
+high-level intent of topology compilation (compiled topology, distance
+extraction, automatic diagram generation), but **which algorithms the
+builder actually uses** was only discoverable via source grep.
+
+## First action
+
+When `resolve_topology(path_str)` is called, four steps run in order:
+
+1. **Path validation** (`builder.py::resolve_topology`):
+   `Path(path_str).expanduser().resolve()`, existence check, file
+   check. Failure → `FileNotFoundError` or `ValueError`.
+2. **YAML parsing** (`_read_spec`): `yaml.safe_load`. Parse errors
+   yield a `ValueError` with line/column. Non-dict roots are
+   rejected.
+3. **Auto-generate the mesh** (`mesh_gen.ensure_mesh_file`): create or
+   reuse a `cube_mesh.yaml` next to the topology file. Cache hit on
+   matching source_hash; miss triggers regeneration. This step decides
+   the cube NoC's router grid and attachment information.
+4. **Compile the graph** (`_compile_graph`): system → IO chiplets →
+   cubes → inter-cube edges → IO↔cube edges → system↔IO edges, then
+   build four view projections (system, sip, cube, pe) and wrap into
+   a `TopologyGraph`.
+
+In short, **topology compilation's first act is "read topology.yaml as
+a dict, create/validate cube_mesh.yaml in the same directory, then
+build the flat graph + 4-view projection in system → sip → cube → pe
+order"**.
+
+## Context
+
+`topology/` package responsibilities:
+
+- **builder.py** (1207 lines): turns topology.yaml into a
+  `TopologyGraph` (nodes + edges + 4 view projections).
+- **mesh_gen.py** (305 lines): auto-decides the cube NoC's router
+  grid and PE/UCIe/M_CPU/SRAM attachment positions and caches them in
+  `cube_mesh.yaml`.
+- **visualizer.py** (887 lines): generates four SVG diagrams (system /
+  sip / cube / pe) from a `TopologyGraph`.
+
+ADR-0006 makes the high-level decision that "the result of topology
+compilation is the single source for distance metadata and diagram
+generation", but specific algorithms (e.g., placement-driven nearest-
+router attachment, the HBM exclusion zone, which fields in source_hash
+trigger regeneration) are not in any ADR.
+
+In particular, these decisions are absent at ADR level:
+
+- Why is mesh_gen cached in a separate file (`cube_mesh.yaml`)?
+- Which fields are in source_hash, and which changes force
+  regeneration?
+- Why placement coordinates in mm rather than cube coordinates?
+- How are the HBM exclusion zone and UCIe N/S/E/W distribution
+  decided inside the mesh?
+- What is the abstraction-level difference among the four view
+  projections (system/sip/cube/pe)?
+
+This ADR captures these decisions in one place.
+
+## Decision
+
+### D1. Compile pipeline — six stages
+
+`_compile_graph(spec)`:
+
+1. **System nodes** (`_instantiate_system`): add system-level nodes
+   like `fabric.switch0` and the host CPU.
+2. **Per-SIP loop** (`for sip_id in range(system.sips.count)`):
+   - **IO chiplets** (`_instantiate_io_chiplets`): create pcie_ep /
+     io_cpu / io_noc / io_ucie PHYs / conn nodes and their bidirectional
+     internal edges.
+   - **Cube instantiation** (`_instantiate_cube`): using
+     cube_mesh.yaml's router grid, instantiate cube routers, PE
+     sub-components (pe_cpu, pe_dma, pe_fetch_store, pe_gemm, pe_math,
+     pe_mmu, pe_tcm, pe_scheduler, pe_ipcq), m_cpu, sram, hbm_ctrl,
+     and their internal edges.
+   - **Inter-cube edges** (`_add_inter_cube_edges`): the UCIe
+     N/S/E/W mesh edges.
+   - **IO ↔ cube edges** (`_add_io_to_cube_edges`): connect io_noc to
+     each cube's edge UCIe phy.
+3. **Switch ↔ IO edges** (`_add_system_to_io_edges`): bidirectional
+   edges between `fabric.switch0` and each SIP's `pcie_ep` (the
+   cross-SIP IPCQ path of ADR-0038 D3 + ADR-0010).
+4. **Build four view projections**:
+   - `_build_system_view(spec)` — Tray level: SIPs and the system
+     switch.
+   - `_build_sip_view(spec)` — inside one SIP: cube mesh + IO
+     chiplet.
+   - `_build_cube_view(spec)` — inside one cube: router grid + PE /
+     M_CPU / SRAM / HBM_CTRL attachments.
+   - `_build_pe_view(spec)` — inside one PE: nine sub-components +
+     internal edges (pe_internal kind).
+5. **Return `TopologyGraph`**: `TopologyGraph(spec, nodes, edges,
+   system_view, sip_view, cube_view, pe_view)`.
+
+The six stages are **ordered for a reason**: only after cubes exist
+do inter-cube edges have valid src/dst, and IO chiplets must precede
+the IO ↔ cube edges that reference them. New node types must slot in
+the right spot.
+
+### D2. `cube_mesh.yaml` — a separate file with a source_hash cache
+
+`mesh_gen.ensure_mesh_file(cube_spec, mesh_path)`:
+
+1. Compute `source_hash = _compute_source_hash(cube_spec)` from these
+   input fields:
+   - `geometry` (cube_mm.w/h …).
+   - `pe_layout` (corners, pe_per_corner).
+   - `ucie.n_connections`.
+   - `memory_map.hbm_mapping_mode`.
+   - `placement` (m_cpu/sram pos_mm).
+2. If `mesh_path` (= `cube_mesh.yaml` next to topology.yaml) exists
+   and `existing.source_hash == source_hash`, reuse it (cache hit).
+3. Otherwise, generate a new mesh via
+   `_generate_mesh(cube_spec, source_hash)` and write to yaml.
+
+Caching as a separate file because:
+
+- Mesh generation involves nontrivial PE/UCIe/router attachment math
+  and is too expensive to redo every time.
+- Multiple runs with the same cube spec must guarantee an identical
+  mesh.
+- The resulting mesh is itself an inspectable / debuggable artifact.
+
+The five fields listed in source_hash are the ones that determine
+mesh shape; other changes (e.g., bandwidth, overhead_ns) do not
+trigger mesh regeneration.
+
+### D3. Cube NoC mesh auto-layout
+
+`_generate_mesh(cube_spec)`:
+
+#### D3.1. Rows / columns
+
+- `pe_positions = _corner_pe_positions(cube_w, cube_h)`: PE-center
+  coordinates (mm) per corner (NW/NE/SW/SE). Hardcoded patterns like
+  `(1.5, 1.5)` and `(cube_w-1.5, cube_h-1.5)`; with `pe_per_corner=2`,
+  each corner has two PE positions.
+- `col_xs = _compute_col_positions(...)`: union of PE x-coordinates,
+  plus relay columns inserted when any gap exceeds
+  `max_spacing = 3.0 mm`.
+- `row_ys, rows_per_half = _compute_row_positions(cube_h,
+  n_connections, pe_positions)`:
+  - `n_conn = max(n_connections, 2)` (hot-path minimum).
+  - `rows_per_half = ceil(n_conn / 2)`.
+  - Top half + two HBM rows + bottom half. HBM sits at
+    `(cube_h/2 - 1.5, cube_h/2 + 1.5)`. The gap between PE rows and
+    HBM rows is `hbm_gap = 1.5 mm`.
+
+#### D3.2. HBM exclusion zone
+
+`hbm_row_start = rows_per_half`,
+`hbm_row_end = rows_per_half + 1`.
+`hbm_col_start = n_cols // 2 - 1`,
+`hbm_col_end = n_cols // 2`.
+
+Router slots inside this (row, col) rectangle are marked `None` (no
+router). HBM controllers are added separately as
+`hbm_ctrl.pe{X}` nodes following ADR-0017 D9's per-PE partition
+pattern.
+
+#### D3.3. PE attachment
+
+Each corner's PEs map to a row:
+
+- Top half: NW → row 0, NE → row 1 (top_corners index).
+- Bottom half: SW → row `hbm_row_end + 1`, SE → row
+  `hbm_row_end + 2`.
+
+Each PE's x-coordinate attaches to the nearest column's router
+(`min(range(n_cols), key=lambda c: abs(col_xs[c] - pe_x))`).
+Attachment items are `pe{pe_idx}.dma`, `pe{pe_idx}.cpu`,
+`pe{pe_idx}.hbm` (pushed into the router's attach list).
+
+#### D3.4. M_CPU / SRAM attachment — nearest router by Euclidean distance
+
+For `placement.m_cpu.pos_mm` (default `[1.5, 5.5]`) and
+`placement.sram.pos_mm` (default `[1.5, 8.5]`), find the router with
+the smallest Euclidean distance and append `"m_cpu"` / `"sram"` to
+its attach list.
+
+#### D3.5. UCIe N/S/E/W distribution
+
+`ucie_pe_rows = top_pe_rows + bot_pe_rows` (total
+`2 * rows_per_half`).
+
+- UCIe-E: one PE row at a time, attach `ucie_e.c{i}` to the rightmost
+  column's router.
+- UCIe-W: attach `ucie_w.c{i}` to the leftmost column's router (E's
+  mirror).
+- UCIe-N/S: split PE columns into left and right halves; attach to
+  the top row's / bottom row's matching columns.
+
+Each UCIe connection is suffixed `c{i}`, distributing
+ucie_n_connections PHYs (ADR-0017 D5+).
+
+### D4. Node naming convention — single ownership
+
+builder.py creates nodes with the following naming convention (the
+single-owner principle from ADR-0051 D5):
+
+- `fabric.switch0` — system-level switch.
+- `sip{S}.{io_id}.{pcie_ep|io_cpu|io_noc|io_ucie.{dir}|conn.{id}}` —
+  IO chiplet.
+- `sip{S}.cube{C}.{m_cpu|sram|hbm_ctrl.pe{X}|noc.r{R}c{C}|...}` —
+  inside cube.
+- `sip{S}.cube{C}.pe{P}.{pe_cpu|pe_dma|pe_fetch_store|pe_gemm|pe_math|pe_mmu|pe_tcm|pe_scheduler|pe_ipcq}` —
+  PE sub-components.
+
+Changing this convention requires updating both builder.py and
+router.py's helpers (ADR-0051). Components never know the convention
+directly — they only call the helpers.
+
+### D5. Edge `kind` classification
+
+Every edge gets a `kind`; routing policy (ADR-0051 D2) reads it. Major
+kinds:
+
+- `"pe_internal"` — within a PE between sub-components.
+- `"pe_to_router"` — PE_DMA ↔ cube NoC router.
+- `"router_mesh"` — between cube NoC routers.
+- `"router_to_hbm"`, `"router_to_mcpu"`, `"router_to_sram"`,
+  `"sram_to_router"`, etc. — between cube-attached components.
+- `"ucie_internal"`, `"ucie_conn_to_router"`,
+  `"router_to_ucie_conn"`, `"ucie_conn_to_noc"`,
+  `"noc_to_ucie_conn"`, `"ucie_mesh"` — UCIe-related.
+- `"io_internal"` — inside IO chiplet.
+- `"io_to_cube"`, `"cube_to_io"` — at the IO ↔ cube boundary.
+- `"pcie"` — switch ↔ pcie_ep.
+- `"command"` — control-plane edges only (e.g., M_CPU ↔ NOC; excluded
+  from PE DMA paths).
+
+Adding a new edge kind requires picking a category in router.py's
+four adjacency graphs (ADR-0051 D2). If you forget, it defaults to
+`_adj_all` only, which can produce unintended routes.
+
+### D6. View projection — four abstraction levels
+
+`TopologyGraph` keeps four view projections alongside the flat
+nodes+edges:
+
+- **system_view** (`_build_system_view`): Tray level. SIP blocks and
+  `fabric.switch0`. PCIe links shown. For external high-level
+  overview.
+- **sip_view** (`_build_sip_view`): inside one SIP — cube mesh + IO
+  chiplet (pcie_ep + io_cpu + io_noc). UCIe N/S/E/W appear as
+  cube-cube links.
+- **cube_view** (`_build_cube_view`): inside one cube — router grid +
+  PE / M_CPU / SRAM / HBM_CTRL attachments + UCIe PHY edges. For
+  intra-cube routing / placement debugging.
+- **pe_view** (`_build_pe_view`): inside one PE — nine sub-components
+  + internal edges (pe_internal kind). For detailed PE-internal
+  dataflow review.
+
+Views are selectively rendered via the spec's
+`visualization.emit_views: [system, sip, cube]` (ADR-0006). The pe
+view is omitted from default output but the code is retained for
+detailed debugging.
+
+### D7. visualizer.py — SVG diagram output
+
+`emit_diagrams(graph, out_dir)` renders every view as SVG. Key
+functions:
+
+- `_render_view_svg(view)` — generic view render (no router grid).
+- `_render_cube_view_svg(view, spec)` — cube-view specific (HBM block,
+  router grid layout, PE/M_CPU/SRAM/HBM placement).
+- `_draw_node`, `_draw_edge` — node/edge visual representation.
+- `_pick_scale`, `_compute_node_sizes` — auto-scaling.
+
+The visualizer is a **derived artifact** (ADR-0006); changes here do
+not pass production checks. Aligns with CLAUDE.md's "Derived
+Artifacts" guidance.
+
+### D8. Blast radius of spec changes
+
+| spec field                            | effect              | mesh regenerated? |
+|---------------------------------------|---------------------|-------------------|
+| `system.sips.count`                   | SIP count, node count | No                |
+| `sip.cube_mesh.w/h`                   | cube mesh shape     | No                |
+| `cube.geometry.cube_mm.w/h`           | cube size (mm)      | **Yes**           |
+| `cube.pe_layout.corners/pe_per_corner`| PE attachment positions | **Yes**       |
+| `cube.ucie.n_connections`             | UCIe PHY distribution | **Yes**         |
+| `cube.memory_map.hbm_mapping_mode`    | HBM distribution mode | **Yes**         |
+| `cube.placement`                      | M_CPU/SRAM positions | **Yes**          |
+| `cube.memory_map.*` (besides above)   | HBM capacity / BW   | No                |
+| `*.links.*.bw_gbs`                    | edge bandwidth      | No                |
+| `*.attrs.overhead_ns`                 | component latency   | No                |
+
+The table mirrors D2's `_compute_source_hash` inputs. Changes that
+require mesh regeneration automatically invalidate `cube_mesh.yaml`'s
+source_hash.
+
+## Alternatives Considered
+
+### A1. Regenerate the mesh on every compile without a cache file
+
+Rejected. The cost of mesh generation would be paid repeatedly (CLI
+runs, probe, tests) for the same spec, and the human-inspectable
+artifact would disappear.
+
+### A2. Merge mesh generation into builder.py
+
+Rejected (currently). It is a 305-line algorithm of its own, and the
+mesh-layout decisions (placement-driven router attachment, HBM
+exclusion zone) are different from builder's general node/edge
+emission. Keeping it separate respects single-responsibility.
+
+### A3. Express placement coordinates in cube coordinates (col/row)
+
+Rejected. mm coordinates flow consistently between the visualizer and
+mesh layout (for nearest-router computation). Cube coordinates are
+undefined until the router grid is fixed, so they are unsuitable as
+placement input.
+
+### A4. Lazy view projection generation
+
+Rejected (currently). The four views are cheap to build (typically <
+100 ms), and eager construction guarantees `TopologyGraph` as the
+single source of truth.
+
+### A5. Visualizer output in formats besides SVG (PNG/PDF)
+
+Rejected. SVG is vector + text-searchable + directly renderable in
+browsers. PNG conversion, when required, is downstream
+post-processing (e.g., rsvg-convert).
+
+## Consequences
+
+- ADR-0006's high-level intent is fleshed out via D1–D7; topology
+  changes can be assessed quickly via D8's table.
+- D3's mesh-layout algorithm is ADR-locked, so future PE attachment
+  patterns (e.g., a 6-zone HBM split) make clear which stage they
+  affect.
+- D5's edge-kind list and D7's view structure are explicit, giving PR
+  reviewers a quick map of where (builder + router + visualizer) a
+  new component type ripples through.
+- D2's source_hash invalidation rules are explicit, so a stale
+  `cube_mesh.yaml` (e.g., when only bandwidth changed) is recognized
+  as correct behavior.

docs/adr/ADR-0054-eval-milestone-benches.md

@@ -0,0 +1,143 @@
+# ADR-0054: Milestone Eval Benches — self-contained sweep + figure benches
+
+## Status
+
+Accepted (2026-05-22).
+
+Amends ADR-0044 (D1/D2) and ADR-0045 (D5) and supersedes the "logic lives
+in `scripts/` + `tests/`" arrangement of ADR-0043/0044: the GEMM and
+allreduce evaluation harnesses are now self-contained **benches** that a
+user runs to regenerate every result + figure.
+
+## Context
+
+ADR-0043 (allreduce eval) and ADR-0044 (GEMM eval) split each harness into
+a **sweep** (a manual `scripts/` driver, or — for allreduce — the
+parametrized tests themselves) plus **figure tests** that render committed
+data. The sweep/render logic therefore lived under `scripts/gemm_sweep.py`,
+`tests/gemm/_gemm_plot_helpers.py`, and `tests/sccl/_allreduce_helpers.py`.
+
+A milestone requirement ("refactor allreduce + GEMM evaluation so a user
+can run *one bench* to generate all the results and plots") cannot be met
+by that layout: a bench is production code and **must not import from
+`tests/`** (ADR-0007 layer direction). The eval logic had to move into
+production, reachable from a bench.
+
+The chosen home is the bench module itself — not a separate
+`kernbench.eval` package. A bench file may contain arbitrary module-level
+code; collapsing the harness into the bench keeps one file per domain and
+avoids an extra package layer.
+
+## Decision
+
+### D1. Two milestone benches own the eval logic
+
+- `src/kernbench/benches/milestone_1h_gemm.py` — GEMM shape×variant sweep +
+  the three figure renderers (moved from `scripts/gemm_sweep.py` +
+  `tests/gemm/_gemm_plot_helpers.py`).
+- `src/kernbench/benches/milestone_1h_ccl.py` — the distributed allreduce
+  driver, latency + buffer-kind sweeps, topology diagram, FSIM comparison,
+  and the direct-launch parity reference (moved from
+  `tests/sccl/_allreduce_helpers.py`).
+
+Each file is the **single home** for its domain's eval logic.
+
+### D2. The "eval bench" pattern (extends ADR-0045 D5)
+
+ADR-0045 D5 fixed a bench to a single configuration (single-SIP, or the
+ADR-0024 multi-SIP CCL exception). This ADR adds a third pattern:
+
+- An **eval bench** may drive *many* configurations and render figures. It
+  builds its own per-config `GraphEngine` / `RuntimeContext` instances
+  (one per sweep point) rather than using the outer `run_bench` engine.
+- Because the outer ctx then has no submitted handles, the bench submits a
+  **sentinel tensor** (`torch.zeros((1, 1), …)`) at the end to satisfy
+  `run_bench`'s "must submit at least one request" contract (ADR-0045 D4),
+  so the CLI exits 0.
+
+### D3. Output location
+
+Both benches write to `src/kernbench/benches/1H_milestone_output/{gemm,ccl}/`
+(per user request — artifacts beside the bench). The directory holds only
+generated PNG/CSV/JSON (never a `.py`/`__init__.py`), so the eager-import
+audit (ADR-0045 first action) ignores it — `pkgutil.iter_modules` does not
+yield non-package subdirectories. It is **committed** (like the
+`docs/diagrams/` artifacts) so the figures are viewable on the remote;
+rerunning the bench regenerates it in place.
+
+### D4. GEMM heavy sweep — fresh by default, `MILESTONE_FAST` to reuse
+
+`milestone-1h-gemm` runs the full 24-sim sweep by default (minutes; one
+shape is 2048 tiles). `MILESTONE_FAST=1` reuses the committed
+`docs/diagrams/gemm_sweep.json` and only re-renders (seconds). This
+reverses ADR-0044 D1/D2's "heavy sweep stays a manual/`slow`-marked step":
+running the bench *is* the regeneration. The slow path is exercised by a
+`@pytest.mark.slow` bench test; the fast path runs by default.
+
+### D5. Tests + script reuse via thin re-export shims (single home kept)
+
+The pre-existing figure tests and the `scripts/gemm_sweep.py` entry point
+are retained and now reuse the bench modules:
+
+- `tests/gemm/_gemm_plot_helpers.py` → re-exports the renderers +
+  `GEMM_SWEEP_JSON`/`GEMM_PLOTS_DIR`/`ROOT` from
+  `kernbench.benches.milestone_1h_gemm`.
+- `tests/sccl/_allreduce_helpers.py` → re-exports the driver core, config
+  writers, sweep constants, renderers, and disk aggregators from
+  `kernbench.benches.milestone_1h_ccl`, and keeps the **pytest-only** pieces
+  local: the `pytest.param` matrices (`CONFIGS` / `_sweep_params` /
+  `_bk_params`) and the fixture-coupled `_run_distributed`
+  (`monkeypatch.chdir` + `_drive_distributed`) wrapper.
+- `scripts/gemm_sweep.py` → thin wrapper over the bench's `run_sweep`.
+
+Tests importing a bench module is permitted (tests sit above production,
+ADR-0007); it triggers the whole-package eager audit, which already runs on
+every `kernbench` invocation. matplotlib stays lazily imported inside the
+renderers, so the audit's startup cost is unchanged.
+
+### D6. Flat module naming (no `benches/` subfolder)
+
+A `benches/` subpackage named `1H_milestone…` is impossible — a Python
+package name cannot start with a digit. The benches are therefore flat
+modules `milestone_1h_gemm.py` / `milestone_1h_ccl.py` with bench names
+`milestone-1h-gemm` / `milestone-1h-ccl` (kebab-case, letter-first per
+ADR-0045 D1).
+
+## Consequences
+
+### Positive
+
+- `kernbench run --bench milestone-1h-gemm` (or `…-ccl`) regenerates all of
+  a domain's results + figures in one command — the milestone requirement.
+- Single source for the eval logic (the bench), reused by tests and the
+  script via shims; no duplication.
+- The figure tests and `scripts/gemm_sweep.py` keep working unchanged.
+
+### Negative / limitations
+
+- The two bench files are large (the CCL one mixes the distributed driver,
+  sweeps, and matplotlib drawing). A "bench" that is mostly an eval harness
+  is unusual; this ADR legitimizes it.
+- Generated artifacts live inside the source tree (`src/kernbench/benches/`)
+  by explicit request and are committed (so the figures are viewable on the
+  remote); rerunning the bench regenerates them.
+- `milestone-1h-ccl` (and the default `milestone-1h-gemm`) take minutes —
+  acceptable for an on-demand milestone artifact, not for routine runs.
+
+## Dependencies
+
+- **ADR-0007**: layer direction (why tests may import production but a bench
+  may not import tests).
+- **ADR-0043 / ADR-0044**: the allreduce / GEMM eval harnesses this ADR
+  relocates into benches.
+- **ADR-0045**: bench module contract; D2 here extends its D5 (single-device
+  rule) with the eval-bench pattern, and relies on D4 (NO_REQUESTS) for the
+  sentinel.
+- **ADR-0024**: rank = SIP launcher driven by the allreduce sweeps.
+
+## Open questions
+
+- Should the GEMM theoretical-model constants (ADR-0044 D5) be sourced from
+  ADR-0033/0014 rather than copied? Unchanged by this ADR.
+- Should `build_overview_slides.py` consume the milestone output PNGs
+  instead of drawing GEMM bars natively? Still open (ADR-0044 D6 / Negative).

docs/adr/ADR-0062-prog-tl-async-load.md

@@ -0,0 +1,169 @@
+# ADR-0062: Lazy `tl.load` — non-blocking HBM load with auto-wait on first use
+
+## Status
+
+Accepted
+
+> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention). Decode and
+> long-context attention are **KV-load-bound**; load/compute overlap is a
+> dominant lever. This ADR makes `tl.load` itself **lazy** (non-blocking,
+> with the wait automatically inserted at first use) rather than adding a
+> separate `load_async` op. Today the only async primitive is for IPCQ
+> comms, not for HBM loads.
+
+## Context
+
+### What overlap requires
+
+FlashAttention streams operands: while the GEMM/MATH engine works on the
+current tile, the DMA engine should already be pulling the next operand.
+With that overlap a bandwidth-bound kernel runs at roughly
+`max(compute, dma)` per tile instead of `compute + dma`.
+
+In ADR-0060's hybrid design the two GEMMs are issued as
+`tl.composite(op="gemm")` whose K/V operands are `tl.ref` (HBM-resident,
+streamed per tile by PE_SCHEDULER), so the **per-tile K/V prefetch is
+handled by the composite scheduler**. Lazy `tl.load` covers the remaining
+explicit loads (the Q group, and any non-composite kernel) so those also
+overlap the compute that follows instead of stalling the greenlet.
+
+### What exists
+
+- `tl.load(ptr, shape, dtype)` is **blocking**: it emits `DmaReadCmd` and
+  the greenlet kernel suspends until PE_DMA signals completion
+  (`tl_context.py:177-203`; greenlet drive `kernel_runner.py:146-153`).
+  No two `tl.load`s can be in flight from one kernel.
+- An async pattern **does** exist, but only for IPCQ:
+  `tl.recv_async(dir, ...) -> RecvFuture` + a deferred wait
+  (`kernel_runner.py:248-285`). It proves the machinery — a non-blocking
+  command that returns a future, resolved later by a wait check
+  (`if not future.event.triggered: yield future.event`) — works in the
+  greenlet model.
+- The DMA engine models a read channel as a SimPy resource (capacity 1,
+  `pe_dma.py:45`) separate from the write channel, so in-flight reads are
+  representable; they serialise on the single read channel while
+  overlapping compute on PE_GEMM/PE_MATH. Only the *kernel-facing API*
+  currently serialises load-vs-compute.
+
+`tl.load` cannot today overlap with the compute that follows it.
+
+## Decision
+
+**Make `tl.load` lazy: it issues the `DmaReadCmd` and returns a handle
+immediately (non-blocking); the runtime auto-inserts the wait at the
+first point the loaded data is actually consumed.** The kernel-facing API
+is unchanged — authors keep writing `tl.load` — so this is a *semantics*
+change, not a new op. It generalises the existing `recv_async`/wait
+machinery (1) to the HBM-load path and (2) from an explicit `tl.wait`
+call to an implicit, dependency-driven wait at first use.
+
+### D1. `tl` surface — unchanged
+
+`tl.load(ptr, shape, dtype)` keeps its signature and `TensorHandle`
+return. What changes is *when* it blocks: never at issue, only implicitly
+when its result is first read by a consuming op.
+
+### D2. Mechanism — non-blocking issue + auto-wait on use
+
+- `tl.load` posts the `DmaReadCmd` to PE_DMA without yielding its
+  completion event, and records the pending event on the returned handle
+  (the `recv_async` pattern, applied to loads).
+- When a consuming op (`tl.dot`, a MATH op, `tl.store`, a `tl.composite`
+  operand, …) is dispatched, the runtime checks each input handle for a
+  pending load event and yields it first if not yet triggered — i.e. the
+  wait is inserted automatically at the latest correct point (first use).
+- The op_log entry is unchanged (`memory/dma_read`): asynchrony is a
+  scheduling property, not a new op kind, so `dma_read_count` and existing
+  op_log consumers keep working.
+
+### D3. Scope — global
+
+`tl.load` is lazy **everywhere**, not behind an opt-in flag. This is the
+faithful model: blocking on every load is a property of a *naive* kernel;
+an efficient kernel (and a real compiler) hoists the load and waits only
+at use. Consequence: existing kernels that have independent work between a
+`tl.load` and its first use see **lower (faster) latency** — a
+correctness improvement of the model, not a behaviour regression. Golden
+latencies that change must be **regenerated**; kernels that load then
+immediately use see no change (the auto-wait fires at once, identical to
+blocking).
+
+### D4. Latency / overlap semantics
+
+- `tl.load` charges the **issue** (descriptor push) only; the kernel
+  proceeds. (Per-op issue cost is `dispatch_cycles`, currently 0; an
+  op-type-differentiated issue cost is tracked separately — ADR-0060 §1,
+  §9.)
+- The DMA transfer occupies the read channel (capacity 1) for its
+  modelled duration in parallel with whatever compute the kernel issues
+  next; multiple in-flight loads serialise on the channel but overlap
+  compute.
+- The auto-inserted wait blocks only if the transfer has not finished.
+- Determinism is preserved: the wait point is fixed by program order
+  (first use), and completion is a scheduled event on the modelled DMA
+  channel (SPEC §0.1, R8). No latency subtraction — the overlap is real
+  modelled concurrency.
+
+> **Modelling assumption.** Auto-wait-at-first-use models a *well-
+> scheduled* kernel (the compiler places the wait at the latest correct
+> point). Real compilers approximate this; some loads cannot be hoisted
+> (register pressure, aliasing). For a performance simulator (SPEC §0)
+> modelling the well-scheduled case is the intended behaviour.
+
+## Alternatives
+
+### A1. Separate `tl.load_async` op (earlier proposal)
+
+Add an explicit `load_async`/`wait` pair the kernel calls by hand (the
+double-buffer dance). Rejected: it enlarges the kernel-facing API and
+pushes buffer-lifetime bookkeeping onto every kernel author, when lazy
+`tl.load` gives the same overlap with **no** API-surface change and a
+compiler-style auto-wait.
+
+### A2. Per-kernel opt-in lazy load
+
+Keep `tl.load` blocking by default; make new kernels opt in. Rejected:
+splits `tl.load` semantics into two variants and hides the model
+improvement from existing benches; global lazy is cleaner and the
+golden-regeneration cost is one-time (D3).
+
+### A3. Rely on composite streaming only (no lazy load)
+
+ADR-0060's composites stream their `tl.ref` K/V operands, so the GEMM
+operand DMA already overlaps. But explicit loads (the Q group, and any
+non-composite kernel) still stall without lazy `tl.load`. Insufficient on
+its own.
+
+## Consequences
+
+### Positive
+- Load/compute overlap with **zero** kernel-facing API change; authors
+  keep writing `tl.load`.
+- Symmetric with `recv_async`/wait — low conceptual surface area.
+- General: any bandwidth-bound kernel prefetches automatically.
+
+### Negative
+- Existing golden latencies shift (faster) for kernels with independent
+  work between load and use → one-time regeneration (D3).
+- Auto-wait requires the runtime to track per-handle pending events and
+  check them at consuming-op dispatch (data-dependency tracking).
+- Interacts with scratch lifetime (ADR-0063): an in-flight load's target
+  buffer must not be recycled before its auto-wait fires. The recycling
+  scope must exclude live (un-waited) load buffers.
+
+## Test Requirements
+
+1. **Overlap is real**: a kernel that issues `tl.load` then an
+   independent GEMM of comparable duration completes in ≈`max(load,
+   gemm)`, not `load+gemm` (end-to-end latency strictly below the serial
+   sum).
+2. **Auto-wait correctness**: the loaded `TensorHandle`, when first
+   consumed, carries the same bytes as today's blocking `tl.load` for the
+   same address (Phase 2).
+3. **Two in flight**: two `tl.load`s to distinct addresses, consumed
+   later, both resolve to correct, independent tensors; their DMAs
+   serialise on the read channel.
+4. **op_log compatibility**: each `tl.load` still logs exactly one
+   `memory/dma_read`; `dma_read_count` unchanged vs the blocking version.
+5. **No-overlap no-change**: a kernel that loads then immediately uses has
+   identical latency to the blocking model (auto-wait fires at once).

docs/adr/ADR-0063-prog-tl-scratch-scope.md

@@ -0,0 +1,224 @@
+# ADR-0063: `tl.scratch_scope` — per-tile scratch recycling for long context
+
+## Status
+
+Accepted
+
+> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention). Long-context
+> attention sweeps many K/V tiles; each tile's intermediates
+> (`scores`, `P`, `exp`, partial `O`, …) currently allocate fresh
+> scratch that is never freed within a kernel invocation. The 1 MiB
+> per-PE scratch budget is exhausted long before a realistic context
+> length.
+
+## Context
+
+### The bump allocator
+
+`TLContext` allocates every math/compute output handle from a per-PE
+scratch pool with a **linear bump cursor**
+(`src/kernbench/triton_emu/tl_context.py`; `_scratch_alloc`):
+
+```python
+def _scratch_alloc(self, nbytes):
+    aligned = (nbytes + 15) & ~15
+    addr = self._scratch_base + self._scratch_cursor
+    self._scratch_cursor += aligned
+    if self._scratch_cursor > self._scratch_size:      # default 1 << 20 = 1 MiB
+        raise RuntimeError("TLContext scratch overflow: ...")
+    return addr
+```
+
+The docstring states the cursor **"resets on every kernel invocation"** —
+i.e. only at kernel entry, never *within* the kernel body. Every
+`tl.dot`, `tl.softmax`, `tl.exp`, `tl.sum`, `a - b`, `a * b`, … grabs a
+fresh slice and nothing is reclaimed until the kernel returns.
+
+### Why this bites the GQA kernel
+
+The current milestone bench keeps `S_q = S_kv_per_rank = 16` **explicitly
+because** of this limit
+(`tests/attention/test_milestone_gqa_llama70b.py:123-148`):
+
+> "S_q_prefill and S_kv_per_rank are deliberately small (16 each) so the
+> simulator's 1 MB per-PE TCM kernel scratch is not exhausted by the
+> bump-allocated handle outputs of softmax/exp/dot/sum chains over
+> n_ranks ring steps."
+
+A FlashAttention sweep over `n_tiles` tiles allocates O(`n_tiles` ×
+per-tile-intermediates). For any realistic context this overflows. The
+*math* of flash attention needs only **O(1)** live scratch — the running
+`(m, l, O)` plus the current tile's working set — because each tile's
+temporaries are dead once that tile is folded into the running state. The
+allocator just doesn't know they're dead.
+
+## Decision
+
+Add a **scratch scope** that lets a kernel mark a region of allocations
+as reclaimable, so per-tile temporaries are recycled while live running
+state (and in-flight prefetch buffers) are preserved.
+
+### D1. `tl` surface — context manager
+
+```python
+with tl.scratch_scope():
+    s   = tl.dot(q, k_t)        # all handles allocated inside the `with`
+    p   = tl.softmax(s)         # share a region that is rewound on exit
+    o_j = tl.dot(p, v)
+    # ... fold o_j into running (m,l,O) which live OUTSIDE the scope ...
+# on __exit__: cursor rewinds to its value at __enter__
+```
+
+Semantics:
+- `__enter__` records the current `_scratch_cursor` as a save-point.
+- `__exit__` restores the cursor to the save-point, freeing everything
+  allocated inside.
+- Handles allocated **outside** the scope (running `m,l,O`, prefetch
+  buffers held by `LoadFuture` from ADR-0062) keep their addresses —
+  they were allocated before the save-point or in an enclosing scope.
+
+### D2. Safety contract
+
+A handle allocated inside a scope **must not** be read after the scope
+exits — its bytes may be overwritten by the next scope's allocations.
+The flash loop respects this naturally: the only values that survive a
+tile iteration are the running accumulators, which are allocated outside
+the per-tile scope and updated by ops *inside* it writing to outside
+addresses (the merge writes new running state — see D3).
+
+### D3. Interaction with running accumulators
+
+The online-softmax merge reads the old running `(m, l, O)` and the
+current tile's `(m_j, l_j, O_j)`, producing new running values. To keep
+the new running values outside the recycled region, the merge writes them
+to **stable scratch** allocated once before the loop (a small fixed
+"running-state" arena, distinct from the per-tile scope). Concretely the
+kernel keeps two arenas:
+
+- **persistent arena** (allocated once): `m, l, O` (and their
+  double-buffer if needed for the merge).
+- **scoped arena** (rewound each tile): `scores, P, exp, O_j, scale_*`.
+
+This mirrors how real flash-attention SRAM budgeting works: a small
+persistent accumulator region + a recycled tile working set.
+
+#### D3.1 The persistent-arena write mechanism: `tl.copy_to(dst, src)`
+
+The merge ops (`tl.maximum`, `tl.exp`, binary `*` / `+`) all call
+`_make_compute_out(...)` which allocates from the bump cursor (D1).
+Inside a `scratch_scope`, their result handles therefore live **inside**
+the scope and vanish on `__exit__`. To realise D3's two-arena split, the
+kernel needs a way to **write a scoped result's bytes back to a
+persistent address**.
+
+The primitive that closes this gap:
+
+```python
+def copy_to(self, dst: TensorHandle, src: TensorHandle) -> None:
+    """Copy ``src``'s bytes into ``dst``'s address (both TCM).
+
+    Shapes and dtypes must match. ``dst`` is typically a handle
+    allocated outside any active ``scratch_scope`` (the persistent
+    arena); ``src`` is a scoped handle whose bytes must outlive
+    scope ``__exit__``.
+    """
+```
+
+Symmetric to `tl.store` (the HBM-side byte copy), kept TCM-only here so
+the running-state writeback doesn't pollute op_log with spurious DMA.
+
+**Mechanics:**
+- New `CopyCmd(src, dst, nbytes, data_op=True)` command.
+- op_log: `op_kind="math"`, `op_name="copy"` — runs on the vector engine.
+- Latency: `pe_math._compute_ns(prod(shape))` — models on-chip register
+  writeback, not HBM transfer.
+- Emit-time validation: `dst.shape == src.shape`, `dst.dtype == src.dtype`,
+  `dst.space == "tcm"`, `src.space == "tcm"`. Authoring errors surface in
+  Phase 1, not deep in Phase 2 data execution.
+
+**Call-site pattern:**
+
+```python
+m, l, O = init_running(...)                  # persistent (outside scope)
+for j in range(n_tiles):
+    with tl.scratch_scope():
+        ...                                  # per-tile work (recycled)
+        m_new = tl.maximum(m, mj)            # scoped scratch
+        l_new = l * scale_old + l_step * scale_step
+        O_new = O * scale_old + O_step * scale_step
+
+        tl.copy_to(m, m_new)                 # ← persist new running state
+        tl.copy_to(l, l_new)
+        tl.copy_to(O, O_new)
+    # exit: scoped m_new/l_new/O_new gone; their bytes live in persistent m/l/O
+```
+
+The copy happens **before** `__exit__`, so the read of `src` (scoped) is
+valid; after exit only `dst` (persistent) is read, satisfying D2's
+no-read-after-exit safety contract.
+
+**Why a dedicated primitive rather than `dst=` kwargs on every math op**
+(considered, rejected): adding `dst=` to `tl.maximum`, `tl.exp`,
+`_binary_math`, `_unary_math`, `_reduction` is ~25 LOC across 5
+op-families and breaks the uniform "call returns a fresh handle"
+pattern. `tl.copy_to` is one primitive, one command, one executor
+handler — minimal surface area for the same effect.
+
+### D4. Nesting
+
+Scopes nest (stack of save-points). Inner scope exit rewinds to the inner
+save-point; outer exit rewinds further. This supports an outer
+"per-query-block" scope around an inner "per-KV-tile" scope for prefill.
+
+## Alternatives
+
+### A1. Round-trip temporaries through HBM
+
+Store intermediates to HBM and reload to "free" TCM scratch. Rejected:
+turns a TCM-resident streaming kernel into an HBM-bandwidth-bound one —
+the opposite of the goal, and it pollutes op_log with spurious DMA.
+
+### A2. Tiled `tl.composite` (scheduler-managed scratch)
+
+`tl.composite` recycles per-tile scratch inside PE_SCHEDULER
+automatically. As in ADR-0062 A1, this is attractive but blocked on a
+flash-capable composite kind (two GEMMs + carried `(m,l,O)`), a much
+larger change. `scratch_scope` gives the same memory behaviour for the
+greenlet kernel with a tiny, general primitive. The two can coexist.
+
+### A3. Grow the scratch budget
+
+Bump `pe_tcm.kernel_scratch_mb`. Rejected as a fix: it only pushes the
+context-length ceiling out linearly while still leaking O(`n_tiles`)
+scratch, and it misrepresents the hardware (real TCM is small; the point
+of flash attention is O(1) working set). Useful only as a coarse knob,
+not a substitute for recycling.
+
+## Consequences
+
+### Positive
+- Removes the artificial `S = 16` validation-scale ceiling; enables
+  realistic context lengths in both timing and data modes.
+- Faithfully models flash attention's O(1) working-set property.
+- Small, general primitive (any tiled kernel benefits).
+
+### Negative
+- A use-after-scope bug silently reads stale bytes. Mitigated by the D2
+  contract, by keeping the scope discipline inside a shared attention
+  helper, and (optionally) by a debug build that poisons rewound regions.
+- Must coordinate with ADR-0062: prefetch buffers are live across tile
+  iterations, so they belong to the persistent arena, not the scoped one.
+
+## Test Requirements
+
+1. **Recycling**: a loop of `N` tiles inside `tl.scratch_scope()` keeps
+   peak `_scratch_cursor` bounded by one tile's footprint, independent of
+   `N` (today it grows linearly and overflows).
+2. **Correctness**: a flash sweep with scopes produces the same `O` as
+   the same sweep without scopes at a small `N` that fits without
+   recycling (Phase 2).
+3. **Long context**: a sweep at an `N` that would overflow 1 MiB without
+   scopes completes (the exact failure the `S=16` cap avoids today).
+4. **Persistent-vs-scoped isolation**: running `(m,l,O)` allocated
+   outside the scope retains correct values across `__exit__`.
+5. **Nesting**: nested scopes rewind to the correct save-points.

docs/adr/ADR-0064-perf-cpu-issue-cost-model.md

@@ -0,0 +1,397 @@
+# ADR-0064: Structural CPU dispatch cost model (`logical_bytes` + FIXED + R)
+
+## Status
+
+Accepted (Revision 2)
+
+> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention) and **ADR-0065**
+> (flat-ops composite + first stateful recipe). The hybrid decision there
+> (GEMMs via `tl.composite`, softmax merge in the kernel) wins by
+> **offloading tiling to PE_SCHEDULER so the CPU issues coarse descriptors
+> and runs ahead, keeping the engines saturated**. That win is currently
+> **invisible in the simulator** because per-op CPU issue cost is zero.
+>
+> Revision 2 replaces the **op-type calibration table** (the original
+> proposal) with a **structural formula** derived from each command's
+> `logical_bytes` — no per-op-type calibration needed; new op kinds are
+> covered automatically.
+
+## Context
+
+### What exists today
+
+- Every `tl.*` op calls `_emit_dispatch_overhead()` before emitting its
+  command (`tl_context.py:196-212`), which emits
+  `PeCpuOverheadCmd(cycles=dispatch_cycles)` **only if** `dispatch_cycles
+  > 0`. The knob is **uniform** across op kinds and hardcoded to **0**
+  in both live execution paths (`pe_cpu.py:101` greenlet, `:195` replay).
+- ⇒ issuing a command — *constructing the descriptor and pushing it to
+  the scheduler queue* — currently costs **0 ns** on PE_CPU.
+- `PeCpuOverheadCmd` is consumed as `yield env.timeout(cmd.cycles)` on
+  PE_CPU (`kernel_runner.py:131-132`).
+
+### Why uniform-and-zero is wrong for the hybrid
+
+ADR-0060 §1's argument is that **one** composite descriptor offloads
+`N_tiles` worth of GEMM tiling, so the CPU issues `O(1)` coarse commands
+instead of `O(N_tiles × ops/tile)` fine ones. With issue cost = 0, the
+model cannot show:
+- that the primitive path may **fail to saturate** the engines when the
+  CPU cannot push fast enough, nor
+- that a composite **costs more to construct** than a single primitive
+  but far less than the many primitives it replaces.
+
+### Why per-op-type calibration (Revision 1) was over-shaped
+
+The original proposal had a `cost_table[kind]` keyed by op kind
+(`composite`, `load`, `dot`, `math`, …). That required:
+- a value per kind (calibration cost ≥ |kinds|),
+- a new entry every time a new kind appears,
+- and yet the *ratio* it tried to capture — "composite ≫ primitive,
+  but ≪ the primitives it replaces" — is structurally a function of
+  **how many fields the command carries**, not of the op kind.
+
+A command's *byte footprint* is the natural proxy: a composite carrying
+N OpSpecs has ~N× the bytes of a primitive op with one OpSpec. The
+*fixed* part (queue head update, completion register, MMIO-class
+latency) is per-command. The two together compose: `FIXED + bytes × R`.
+
+### What this model actually exposes
+
+The **primary** signal is **command-count reduction** through the
+per-command FIXED cost. The **byte** term is a secondary refinement
+that prevents pathologically-large composites from looking free. With
+realistic on-die queue bandwidth (16 B/cycle, D3), FIXED accounts for
+≥85% of the dispatch cost differential between opt3 (≈10 cmds/tile)
+and opt2 (≈2 cmds/tile) for decode opt2.
+
+This framing also bounds the model: if a single composite were allowed
+to grow unboundedly large, the byte term alone would not stop the
+formula from rewarding ever-bigger fused commands beyond what real HW
+supports — hence the descriptor-size cap in **D7**.
+
+## Decision
+
+### D1. Structural dispatch cost formula
+
+Each PE command going to PE_SCHEDULER incurs PE_CPU dispatch cycles:
+
+```
+dispatch_cycles(cmd) = FIXED_PER_CMD + cmd.logical_bytes × R
+```
+
+where:
+- `FIXED_PER_CMD` (cycles per command) models queue-tail update, MMIO-class
+  RTT, completion-event registration — fixed per command regardless of size.
+- `R` (cycles per byte) models the queue-write bandwidth — bytes of the
+  command serialized into the scheduler queue.
+- `cmd.logical_bytes` (int) is each command's *HW-logical* byte size,
+  computed from D2 below — not Python's `sys.getsizeof`.
+
+PE_CPU emits `PeCpuOverheadCmd(cycles=dispatch_cycles(cmd))` before
+dispatching, exactly as the existing hook (`tl_context.py:_emit_dispatch_
+overhead`) — only the cycle value changes.
+
+### D2. `logical_bytes` rule
+
+Each PE command dataclass exposes `logical_bytes: int` (property). The
+counting rule (HW-friendly, ignores Python overhead):
+
+| Field kind | Bytes |
+|---|---|
+| Command framing (cmd-type discriminator + completion id ref) | 4 |
+| Opcode (op kind enum) | 1 |
+| Enum (scope, etc.) | 1 |
+| `TensorHandle` reference (address only — shape/dtype assumed in descriptor table) | 8 |
+| Scalar (int/float) | 4 |
+| Tuple length marker | 1 |
+
+`CompositeCmd` recursively sums its `ops` and `rw_handles`:
+
+```python
+@property
+def logical_bytes(self) -> int:
+    return (
+        4                               # framing
+        + 1 + sum(op.logical_bytes for op in self.ops)
+        + 1 + 8 * len(self.rw_handles)
+    )
+```
+
+`OpSpec`:
+
+```python
+@property
+def logical_bytes(self) -> int:
+    return (
+        1 + 1                                       # opcode + scope
+        + 1 + 8 * len(self.operands)                # named operand handles
+        + (8 if self.out is not None else 0)        # out handle
+        + 1 + sum(_extra_bytes(v) for v in self.extra.values())
+    )
+
+
+def _extra_bytes(v) -> int:
+    """Type-aware byte count for OpSpec.extra values."""
+    if isinstance(v, bool):              return 1
+    if isinstance(v, (int, float)):      return 4
+    if isinstance(v, (tuple, list)):     return 1 + 4 * len(v)   # shape, axes, …
+    if isinstance(v, str):               return 1                # opcode-like tag
+    return 4                                                     # default scalar
+
+
+# Example types in extra:
+#   m, k, n            int   → 4 each
+#   reduce_axis        int   → 4
+#   shape=(64, 64)     tuple → 1 + 8 = 9
+#   factor=1.0         float → 4
+```
+
+(Identical rule for `DmaReadCmd`, `MathCmd`, etc. — one property per
+dataclass, ~3 lines each.)
+
+**Counting rule — per-op summation, no deduplication.** Each
+`TensorHandle` reference inside an `OpSpec` is counted independently,
+even when the same handle appears in multiple OpSpecs or in
+`rw_handles`. The `rw_handles` block is metadata for cross-composite
+hazard tracking (ADR-0065 D6.3) and is counted separately from
+operand references in `ops`. There is no deduplication. This matches
+HW reality: the descriptor encodes each operand slot as an independent
+address field, and the dispatcher tracks `rw_handles` as a distinct
+metadata block. Example: `OpSpec(kind="mul_bcast", operands={"src_a":
+O, "src_b": corr}, out=O)` counts handle `O` *twice* (once for
+`src_a`, once for `out`); if `O` is also in the enclosing
+CompositeCmd's `rw_handles`, it is counted a third time.
+
+### D3. Defaults — anchored on a typical composite ≈ 43 ns
+
+Anchor: a **single-OpSpec composite for a DMA-staged GEMM path** —
+one `OpSpec(kind="gemm", ...)`; DMA stages are auto-inserted by
+PE_SCHEDULER from operand `space` (ADR-0065 D4) and **do not appear in
+`logical_bytes`** (the kernel does not issue them as separate cmds).
+Breakdown:
+
+```
+framing                  4
+ops tuple length         1
+GEMM OpSpec             40   (opcode 1 + scope 1 + 1 + 2 handles 16
+                              + out 8 + 1 + extra m/k/n 12)
+rw_handles tuple length  1
+rw_handles content       8   (one RW handle for the output)
+─────────────────────────────
+total                  ~54 bytes
+```
+
+Target dispatch = ~43 ns. On-die producer→consumer queue at 16 bytes/cycle
+(typical on-die descriptor queue width).
+
+```
+FIXED_PER_CMD        = 40 cycles
+R                    = 0.0625 cycles/byte  (= 16 bytes/cycle)
+```
+
+Verification: `40 + 54 × 0.0625 = 43.375 cycles ≈ 43 ns` ✓
+
+Cycle→ns conversion uses the PE node's existing `clock_freq_ghz` attr
+(the same one used by PE_MATH `_compute_ns`). The cost-model knobs are
+**cycle-domain only** — they do not duplicate the clock setting.
+
+### D4. Topology config override
+
+Defaults are baked into `pe_cpu.py`. Topology yaml may override under a
+`pe_cost_model:` section at the PE node attrs (cycle-domain knobs only;
+clock comes from the PE's existing `clock_freq_ghz`):
+
+```yaml
+pe:
+  attrs:
+    clock_freq_ghz: 1.0                    # existing, used for cycle→ns
+    pe_cost_model:
+      fixed_per_cmd_cycles: 40
+      byte_cycles_recip: 0.0625            # = 16 bytes/cycle
+      max_composite_logical_bytes: 1024    # D7 — descriptor size cap
+```
+
+Missing keys fall back to defaults. The dispatch formula reads from
+`node.attrs["pe_cost_model"]` at PE_CPU init.
+
+### D5. Scope — what does and does not pay
+
+| Path | Pays dispatch cost? |
+|---|---|
+| PE_CPU → PE_SCHEDULER for any `PeCommand` | **Yes** |
+| `PeCpuOverheadCmd` itself (already cycles-explicit) | **No** (formula bypass) |
+| Stages auto-generated by PE_SCHEDULER (DMA_READ/WRITE/FETCH/STORE) | **No** (PE_SCHEDULER-internal) |
+| Engine compute latency (DMA `drain_ns`, GEMM/MATH `_compute_ns`) | **No change** — stays on engines (SPEC §0.1) |
+
+This preserves the "latency on modelled components" invariant — dispatch
+cost is *additional* CPU-side time, not folded into engine times.
+
+### D6. Configurable values; goldens regenerate
+
+Turning issue cost non-zero changes **every** bench's latency. Golden
+latencies are **regenerated once** when this ADR lands — same posture as
+ADR-0062 D3 lazy-load. After regeneration, the same calibration is in
+effect for ADR-0065 opt2 measurement.
+
+### D7. Composite size cap (hard limit — validation error)
+
+Each `CompositeCmd`'s `logical_bytes` is capped at
+**`max_composite_logical_bytes`** (default **1024 bytes**, overridable
+per D4). A composite whose `logical_bytes` exceeds the cap is **rejected
+at emit time** by the host-side TLContext with a `ValueError` — there is
+**no automatic segmentation**. The kernel author must restructure the
+recipe (e.g., split it into multiple smaller composites explicitly) so
+each `CompositeCmd` fits within the descriptor capacity.
+
+**Why a cap is needed.** Real hardware imposes hard limits — descriptor
+queue entry size, scheduler parser buffer, command SRAM, firmware
+input. Without a cap, the `FIXED + bytes × R` model would reward
+arbitrarily large fused composites beyond what hardware accepts (e.g.,
+fusing 100 primitive ops into one composite, paying one `FIXED`).
+
+**Why 1024 bytes specifically.** This is a **safe engineering limit**,
+not a measured HW number — intentionally chosen to be well above all
+currently known composites (decode opt2's `#2` at ~322 bytes is the
+largest in the kernbench codebase) while still representing a *finite*
+descriptor capacity that future recipes must respect. The number is
+overridable per topology (D4); when a real HW reference appears, the
+value should be recalibrated. The role of this default is to make the
+cap *exist as a discipline*, not to fit a specific HW.
+
+**Rationale for a hard error over auto-segmentation.** Auto-splitting an
+oversized composite (the original Revision 2 proposal) added emitter
+complexity — inter-segment ordering, shared `rw_handles`, completion
+chaining — to paper over what is, in practice, a kernel that asked for a
+descriptor larger than the hardware allows. Surfacing it as an explicit
+error keeps the emitter simple and makes the HW constraint visible to the
+kernel author, who is best placed to decide how to split the work.
+
+Decode opt2's `#2` composite (10 ops, ~322 bytes) sits comfortably
+inside the 1024 cap — no error for the GQA workload.
+
+## Alternatives
+
+### A1. Keep Revision 1's op-type calibration table
+
+Rejected: calibration cost scales with |kinds|, and the *ratio* the
+table tried to capture is structurally a function of cmd size. The
+structural formula reaches the same qualitative behaviour with two
+calibratable numbers instead of N.
+
+### A2. Byte-only formula (no FIXED term)
+
+Rejected. With FIXED = 0, opt2 (Option Y per ADR-0065) does **not** win
+over opt3 — the total *bytes* dispatched per tile are similar (opt3 ≈ 232,
+opt2 ≈ 380); the win is entirely in *fewer per-cmd fixed costs*. A
+byte-only formula erases the very signal the model needs to expose.
+
+### A3. Charge dispatch on PE_SCHEDULER instead of PE_CPU
+
+Rejected: the saturation question is *"can the CPU push descriptors fast
+enough to keep the engines busy?"* — that is a **PE_CPU** issue-bandwidth
+property. Charging on the scheduler would not model CPU back-pressure.
+
+### A4. Model DMA program/setup time as a separate fixed per-descriptor cost
+
+Deferred: initially fold the descriptor-program cost into the **issuing
+op's** dispatch cost. Split it out to a PE_DMA fixed setup only if
+calibration shows it matters.
+
+## Consequences
+
+### Positive
+- Hybrid's CPU-offload / saturation win (ADR-0060 §1) becomes
+  **measurable**, with a structurally honest model (no calibration table).
+- Adding new op kinds (e.g., ADR-0065's `softmax_merge` 8-step recipe)
+  costs zero — they fit the same formula automatically.
+- More faithful to hardware (queue-head MMIO RTT + queue-write bandwidth).
+
+### Negative
+- **All** bench goldens shift → one-time regeneration (D6); CI golden
+  fixtures update.
+- Two calibration knobs (FIXED, R) need values; defaults are anchored on
+  a documented assumption — treat absolute latencies as provisional until
+  a reference exists; keep the **ratios** defensible.
+- Adds a small `logical_bytes` property to each PE command dataclass.
+
+## Open review items
+
+1. **Calibration source for FIXED and R.** Defaults from "typical
+   composite ≈ 43 ns + on-die queue 16 bytes/cycle"; reasonable for an
+   on-die descriptor queue. Revisit when a HW reference appears.
+2. **Scheduler plan-gen cost vs large composites.** Stays 0 — D5 keeps
+   PE_SCHEDULER's plan-generation outside the dispatch formula. The
+   D7 cap (1024 bytes ≈ 30–35 OpSpecs) bounds the worst case, but a
+   composite near the cap still costs the scheduler the same as a 1-op
+   composite under the current zero-cost model. If a stress test
+   (large-composite microbench) shows scheduler-bound behaviour,
+   expose `overhead_ns` per-op-count.
+3. **Where the override lives.** `pe_cost_model:` block under PE node
+   attrs in topology yaml — keeps all knobs in one place, reviewable.
+   Clock comes from the PE's existing `clock_freq_ghz` attr, not
+   duplicated here.
+4. **Path parity.** Both greenlet (`_execute_legacy` and `kernel_runner`)
+   and replay paths must read the same cost model. Verify.
+5. **Sensitivity of conclusions to R.** opt2 < opt3 must hold across a
+   reasonable range of `R` (queue-bandwidth assumption). Sensitivity
+   sweep is part of Test Requirements (#9).
+
+## Test Requirements
+
+Tests are written against the **formula** (D1), not against specific
+numeric anchors, so they remain valid when calibration changes or when
+OpSpec/CompositeCmd fields are added.
+
+1. **Formula preservation.** For any `CompositeCmd` `c`, PE_CPU's
+   recorded dispatch overhead equals `FIXED_PER_CMD + c.logical_bytes
+   × R` (within ±1 cycle for floor/round-off). Parametrized over
+   several composites: a 1-OpSpec GEMM composite, a 5-op MATH chain,
+   and a 10-op recipe composite. Default-calibration numbers (anchor
+   ≈43 ns for the 1-OpSpec composite) are informative reference, not
+   the test gate — the test gate is the formula equality.
+2. **Qualitative ratio (robust).** opt3 per-tile PE_CPU dispatch
+   strictly exceeds opt2 per-tile dispatch by at least a 2× margin —
+   `opt3 > 2 × opt2`. The default-calibration model predicts ≈4×;
+   the gate is the loose 2× bound so the test does not break when
+   calibration is moved (e.g., when a HW reference replaces the
+   default). Informative numbers — see ADR-0065 §verification and
+   DDD-0065 §11 for the model expectation.
+3. **Override path.** Topology yaml `pe_cost_model:` block changes the
+   per-PE dispatch cost; default is recovered when block is missing.
+   The formula identity from #1 must hold with the override values.
+4. **`PeCpuOverheadCmd` bypass.** Manual `tl.cycles(n)` issues exactly `n`
+   cycles, not `n + dispatch_cycles(...)`.
+5. **No double-count.** PE_DMA `drain_ns`, PE_GEMM/MATH `_compute_ns`
+   identical to pre-ADR values.
+6. **Determinism.** Identical inputs → identical op_log + latency
+   (SPEC §0.1).
+7. **Path parity.** Greenlet and replay paths produce identical
+   dispatch-cycle accounting for the same kernel.
+8. **Composite size cap (D7).** A composite that would emit `logical_bytes
+   > max_composite_logical_bytes` raises a `ValueError` at emit time (no
+   segmentation). A composite within the cap emits normally.
+9. **Sensitivity (qualitative).** At `R ∈ {0.25, 0.0625, 0.03125}`
+   cycles/byte, `opt3 > opt2` at *all* three points. Direction (ratio
+   monotonically increases as R decreases) is also asserted, but
+   absolute ratio values are *not* required.
+
+## Migration
+
+ADR-0064 Revision 2 lands as a single PR with:
+- `logical_bytes` property on each `PeCommand` dataclass (type-aware
+  extra-field counting per D2)
+- formula application in `pe_cpu.py` dispatch path
+- `pe_cost_model:` override read at PE_CPU init (cycle-domain knobs +
+  `max_composite_logical_bytes`)
+- **composite size cap (D7)** — TLContext-side hard cap with
+  `max_composite_logical_bytes` default 1024: a composite over the cap
+  raises a `ValueError` at emit time (no auto-segmentation). Not
+  triggered by any existing bench (largest current composite is ~322
+  bytes), but the check lands so the descriptor-capacity limit is
+  enforced as recipes grow.
+- one-time goldens regeneration
+
+After this lands, ADR-0065 builds on top with no further goldens churn
+in existing benches (ADR-0065 is a meaning-preserving refactor of
+`CompositeCmd` for the existing path; only opt2 is a new bench).

docs/adr/ADR-0065-prog-flat-ops-composite-softmax-merge.md

@@ -0,0 +1,488 @@
+# ADR-0065: Flat-ops `CompositeCmd` + first stateful recipe (`softmax_merge`)
+
+## Status
+
+Accepted
+
+> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention). Implements
+> the §5.6 / §8 item 4 carve-out exactly: opt2 of decode = (#1 existing
+> GEMM composite for Q·Kᵀ) + (#2 a single composite carrying softmax +
+> P·V + online-softmax merge). The "ex_composite" / "flash_pv_merge" name
+> referenced in ADR-0060 §5.6 is *retired* — this ADR fixes the canonical
+> shape using the existing `tl.composite` entry with two structural
+> additions: (a) a flat `ops` tuple, and (b) a first stateful recipe
+> `softmax_merge` that lowers to a sequence of MATH micro-ops.
+
+## Context
+
+### What ADR-0060 left to this ADR
+
+ADR-0060 §5.6 shipped opt3 (software pipelining) as the recommended now,
+and explicitly carved out opt2 — fewer per-tile dispatches, K-before-V
+DMA priority, **measurable only after ADR-0064**'s cost model — as the
+follow-up.
+
+§8 item 4 sized the carve-out: "if revisited, split it into **two**
+composites — `#1` = Q·Kᵀ (existing composite + `scale`), `#2` = softmax +
+P·V + online-softmax accumulator merge". This ADR implements `#2`.
+
+### Why composite needs a structural change
+
+Today's `CompositeCmd` has the shape `(op, a, b, out_addr, ops=(head,
+*epi))` with implicit conventions:
+- `op` selects the head engine (gemm | math)
+- `ops[0]` is the head OpSpec, `ops[1:]` are epilogue OpSpecs that run
+  *after* the head per-output-tile or per-K-tile (scope-determined)
+
+For decode opt2 we need MATH ops to run **before** the head GEMM (the
+online-softmax `m, l, O` update + emit `P` for the GEMM input). The
+current shape has no slot for that.
+
+Two failed paths considered earlier:
+- Embedding softmax_merge as an *epilogue* of #1 (Q·Kᵀ) — wrong: it
+  needs to read `Sj` (= #1's output) and writes a fresh `P` consumed by
+  the *next* GEMM (#2's P·V); circular.
+- A single mixed-engine recipe `flash_pv_merge` absorbing MATH + GEMM +
+  MATH — violates the engine boundary, complicates PE_SCHEDULER.
+
+The clean shape is: **drop the head/epilogue distinction from the
+command format**. Make `ops` a flat ordered tuple; let each op's
+`scope` + position determine its tile-loop placement. The "prologue"
+concept stays in the *user-facing* `tl.composite(...)` API as
+ergonomics — but the **HW command never sees it**.
+
+### Why a recipe (not a generic op list)
+
+Decode opt2's `#2` MATH chain is 8 steps with a fixed structure (online
+softmax merge). Letting the kernel author wire those 8 steps would be
+both verbose and a hazard surface (correctness depends on step order).
+A *recipe* — a single named op (`softmax_merge`) that TLContext expands
+to the 8 steps with addresses pre-filled — keeps the kernel concise and
+the structure under one source of truth. The recipe table lives in
+TLContext (the compiler analog); PE_SCHEDULER never reads it. See
+boundary in D5.
+
+## Decision
+
+### D1. `CompositeCmd` is a flat ordered op list
+
+```python
+@dataclass(frozen=True)
+class CompositeCmd:
+    completion: CompletionHandle
+    ops: tuple[OpSpec, ...]                       # ordered MATH/GEMM ops
+    rw_handles: tuple[TensorHandle, ...] = ()     # cross-composite hazard
+    data_op: bool = True
+
+    @property
+    def logical_bytes(self) -> int: ...           # ADR-0064 D2
+```
+
+The previous `(op, a, b, out_addr, out_nbytes, math_op)` fields are
+absorbed into `ops`. The first OpSpec in `ops` (if `kind == "gemm"`) is
+the head GEMM; preceding OpSpecs are pre-loop ops, following OpSpecs
+are post-loop / per-tile epilogue ops depending on `scope`.
+
+### D2. `OpSpec` carries named operands
+
+```python
+@dataclass(frozen=True)
+class OpSpec:
+    kind: str                                     # "gemm" | "rmax" | ...
+    scope: Scope                                  # KERNEL | K_TILE | OUTPUT_TILE
+    operands: dict[str, TensorHandle]             # named — see D5
+    extra: dict[str, Any]                         # scalars, axes, m/k/n, etc.
+    out: TensorHandle | None = None               # explicit write-back handle
+
+    @property
+    def logical_bytes(self) -> int: ...
+```
+
+Named operands let PE_SCHEDULER route inputs to engine ports (e.g., GEMM
+takes `operands["a"]`, `operands["b"]`) without positional convention.
+
+### D3. Position + scope determines tile-loop placement
+
+**Semantics split.** *Position* determines the **phase** (where in the
+tile-loop lifecycle the op fires); *scope* determines the **repetition**
+within that phase. `Scope.KERNEL` means "**once per `CompositeCmd`
+invocation**", not "once per kernel launch" — its single firing happens
+at the op's position in the sequence. The same KERNEL value carries
+different *phases* depending on whether the OpSpec sits before, at, or
+after the GEMM op.
+
+PE_SCHEDULER scans `cmd.ops` for the GEMM op (≤1 per CompositeCmd, see
+D6). Calling its index `g`:
+
+| OpSpec position | scope | Phase | Placement in plan |
+|---|---|---|---|
+| `0 .. g-1` | KERNEL | pre-tile-loop | run once before tile loop |
+| `g` | KERNEL | head GEMM | drives tile loop with `extra["m"], extra["k"], extra["n"]` |
+| `g+1 .. ` | K_TILE | in-accumulation | per K-tile epilogue (inside K accumulation loop) |
+| `g+1 .. ` | OUTPUT_TILE | per-output-tile | per (m, n) tile epilogue (after K accumulation) |
+| `g+1 .. ` | KERNEL | post-tile-loop | run once after tile loop |
+
+A composite with no GEMM op (e.g., MATH-only composite) treats all ops
+as KERNEL-scope sequential — phase collapses to "single-shot in order".
+
+### D4. PE_SCHEDULER auto-inserts DMAs from the operand `pinned` flag
+
+PE_SCHEDULER scans the GEMM's `operands`. For each:
+- **not `pinned`** (data lives in HBM, must be streamed) → emit DMA_READ
+  Stage before the FETCH/GEMM Stages
+- **`pinned`** (already staged in TCM via a prior `tl.load`, or a recipe's
+  TCM scratch / primary-out `P`) → no DMA Stage, consumed in place
+
+The kernel never emits explicit DMA cmds for composite operands;
+PE_SCHEDULER infers them from handles. `tl.ref` operands are not pinned
+(DMA_READ); `tl.load` results and recipe scratch / primary-out are pinned
+(in place).
+
+> **As-built note (P3).** An earlier draft of D4 derived the DMA decision
+> from a `space` tag (`hbm`→DMA, `tcm`→in place) and would have flipped the
+> current `space` values (today `tl.load`=`hbm`, `tl.ref`=`tcm`). The
+> implementation **keeps the existing `pinned` flag** as the DMA signal
+> instead — `pinned` already encodes "already in TCM, skip DMA", flipping
+> `space` would change existing bench Stage sequences for no functional
+> gain. Unifying `space` and retiring `pinned` is a deferred low-priority
+> cleanup.
+
+**Prologue recipe ops are TCM-only (Phase 1 limit, enforced at TLContext
+emit).** Every operand of a *prologue recipe* (non-GEMM) OpSpec **must** be
+TCM-resident; passing an HBM handle as a recipe operand raises a validation
+error at emit time. This bounds DMA staging to the head op's operands in
+Phase 1. The **head op itself** (gemm *or* math) keeps the existing
+DMA-staged-from-HBM behavior — D4's TCM-only rule does **not** apply to it.
+This invariant is also restated in D6 #7.
+
+### D5. RECIPE_DESCRIPTORS — TLContext-internal, NOT in `pe_commands.py`
+
+Recipes live in a new TLContext-adjacent module
+(`src/kernbench/triton_emu/tl_recipes.py`). PE_SCHEDULER does **not**
+import it. The first recipe:
+
+```python
+@dataclass(frozen=True)
+class RecipeDescriptor:
+    # Operand contract: name → "R" (read-only) | "RW" (read-write)
+    operands: dict[str, str]
+    # Type / shape rule for the implicit primary output (auto-bound to
+    # head's first matrix operand). None means recipe emits no value.
+    primary_out: PrimaryOutSpec | None
+    # Single-shot prologue (full reduction) vs tile-aligned execution.
+    tile_alignment: Literal["single_shot", "tile_aligned"]
+    # Internal scratch budget (bytes), function of operand dims.
+    internal_scratch_bytes_fn: Callable[..., int]
+    # Engine micro-op sequence: TLContext expands this into flat OpSpecs.
+    engine_seq: tuple[EngineOp, ...]
+
+
+RECIPE_DESCRIPTORS["softmax_merge"] = RecipeDescriptor(
+    operands={"s": "R", "m": "RW", "l": "RW", "O": "RW"},
+    primary_out=PrimaryOutSpec(
+        from_shape="s", from_dtype="s", transform="identity",
+    ),
+    tile_alignment="single_shot",
+    internal_scratch_bytes_fn=lambda G, TILE, d, bpe: bpe * (
+        G + G + G + G * TILE + G  # m_loc + m_new + corr + P + l_loc
+    ),
+    engine_seq=(
+        EngineOp("MATH", "rmax",       src="s",         dst="m_loc",  reduce_axis=-1),
+        EngineOp("MATH", "max_elem",   src_a="m",       src_b="m_loc", dst="m_new"),
+        EngineOp("MATH", "exp_diff",   src_a="m",       src_b="m_new", dst="corr"),
+        EngineOp("MATH", "exp_diff",   src_a="s",       src_b="m_new", dst="P",     bcast_axis=0),
+        EngineOp("MATH", "rsum",       src="P",         dst="l_loc",   reduce_axis=-1),
+        EngineOp("MATH", "fma",        src_a="l",       src_b="corr",  src_c="l_loc", dst="l"),
+        EngineOp("MATH", "mul_bcast",  src_a="O",       src_b="corr",  dst="O",     bcast_axis=1),
+        EngineOp("MATH", "copy",       src="m_new",     dst="m"),
+    ),
+)
+```
+
+TLContext, on `tl.composite(prologue=[{"op": "softmax_merge", ...}],
+op="gemm", ...)`:
+1. Validates operand R/RW against `RECIPE_DESCRIPTORS["softmax_merge"]`.
+2. Allocates scratch (m_loc, m_new, corr, P, l_loc) via the scratch_scope
+   helper (ADR-0063).
+3. Derives the primary output P's shape from `s.shape` (identity).
+4. Expands `engine_seq` into 8 flat MATH OpSpecs with all addresses /
+   sizes filled. Each OpSpec gets `scope=KERNEL`.
+5. **Auto-binds (with conflict check, D6.6).** If the head GEMM does
+   *not* explicitly provide `operands["a"]`, auto-bind
+   `operands["a"] = P_handle`. If the kernel *did* provide `a`
+   explicitly, emit a validation error — ambiguous primary-output
+   replacement.
+6. Emits `CompositeCmd(ops=(8 MATH + 1 GEMM + epilogue), rw_handles=
+   (m, l, O))`.
+
+**RECIPE_DESCRIPTORS appears nowhere in the HW path**. PE_SCHEDULER
+sees only the flat ops list.
+
+### D6. Invariants
+
+1. **GEMM count.** `0 ≤ count(op.kind == "gemm" in cmd.ops) ≤ 1`. A
+   composite with 0 GEMMs is MATH-only.
+2. **softmax_merge has no V operand.** `RECIPE_DESCRIPTORS["softmax_merge"
+   ].operands` does not include any HBM-resident handle → no DMA inserted
+   during the prologue → V DMA can only happen during the GEMM head →
+   natural K-before-V DMA priority across #1 / #2 of decode opt2.
+3. **Strict FIFO across composites.** PE_SCHEDULER tracks in-flight
+   `rw_handles`. A new composite whose `rw_handles` (or any of its op
+   inputs) intersect with an in-flight composite's `rw_handles` waits
+   for all earlier composites to complete — no out-of-order reorder.
+4. **No backwards-incompatibility for legacy callers.** The user-facing
+   API `tl.composite(op="gemm", a, b, epilogue=[...])` is preserved.
+   TLContext internally lowers it to a flat-ops `CompositeCmd`.
+5. **PE_SCHEDULER recipe-free.** PE_SCHEDULER does not import
+   RECIPE_DESCRIPTORS, does not know that "softmax_merge" exists, and
+   has no recipe-specific code branches.
+6. **No implicit operand replacement (auto-bind conflict).** If a
+   prologue recipe declares a `primary_out` AND the head op's auto-bind
+   target operand (e.g., GEMM `a`) is *also* provided by the kernel
+   explicitly, TLContext raises a validation error at emit time. The
+   kernel must either let the recipe bind (omit the operand) **or**
+   specify it explicitly (omit the prologue, or use a recipe without
+   `primary_out`). This prevents the ambiguity of "which value wins".
+7. **Prologue MATH operands TCM-only.** Restatement of D4: every operand
+   of a *prologue recipe* (non-GEMM) OpSpec must be TCM-resident. The head
+   op (gemm or math) is exempt — it may stream from HBM. Phase 1 enforced
+   at TLContext emit.
+
+### D7. Boundary summary (compiler vs scheduler vs engine)
+
+```
+HOST                                       DEVICE
+TLContext (compiler analog)          PE_SCHEDULER (scheduler/dispatcher)
+  - validate against RECIPE_DESC.       - scan flat ops list
+  - derive primary_out shape            - identify GEMM (≤1)
+  - allocate scratch                    - position-based stage placement
+  - expand recipe → flat OpSpecs        - auto-insert DMA stages from space
+  - compute rw_handles                  - strict-FIFO RW hazard tracking
+  - emit CompositeCmd                   - emit Stage list
+  imports: RECIPE_DESCRIPTORS            imports: nothing recipe-related
+
+                ENGINES (PE_MATH / PE_GEMM / PE_DMA)
+                  - read Stage.params (op_kind, addresses, n_elements)
+                  - existing latency models unchanged
+                  imports: nothing recipe-related
+```
+
+### D8. Composite output handle + output-space DMA
+
+`tl.composite(...)` returns the **output `TensorHandle`** (not a bare
+`CompletionHandle`), so a composite's result can feed a downstream op the
+same way `tl.dot`'s result does. The handle's `pending` field references
+the composite's completion; downstream consumers auto-await it
+(`_await_pending`, ADR-0062 lazy pattern) and `tl.wait(handle)` works.
+
+**`out` is always a `TensorHandle`** (never a raw address) — its location is
+the kernel's choice, never guessed:
+
+- **HBM output** → `out=tl.ref(addr, shape)`. `tl.ref` returns a handle with
+  `space="hbm"`, so the composite writes it back via DMA_WRITE. The
+  STORE + DMA_WRITE stay **inside** the composite pipeline — they are part of
+  the fused op, *not* a separate `tl.store`.
+- **In-place TCM** → `out=<existing TCM handle>` (e.g. the recipe accumulator
+  `O`). `space="tcm"` → STORE only, result stays in TCM (chainable).
+- **Omitted** → TLContext **auto-allocates a TCM scratch** (like `tl.dot`)
+  and returns it.
+
+For ergonomics, `out_ptr: int` is accepted as **shorthand for an HBM
+output** — equivalent to `out=tl.ref(out_ptr, <output shape>)` (the compiler
+wraps it into an `space="hbm"` handle). When no output is specified the
+compiler owns the scratch address, never the kernel. (`tl.ref` returns
+`space="hbm"` — it references HBM data. The operand-**input** DMA decision
+stays `pinned`-based per D4, so this label does not affect input streaming.)
+
+The output handle's **`space` drives the tile loop's write-back** (the
+output analog of the operand `pinned` rule in D4):
+
+| `out.space` | tile-loop write-back |
+| --- | --- |
+| `"tcm"` | STORE only — result stays in TCM (chainable; **no DMA_WRITE**) |
+| `"hbm"` | STORE + DMA_WRITE — final HBM output |
+
+This makes a TCM-resident composite output (e.g. a recipe accumulator) legal
+**and** keeps a TCM-scratch address out of the DMA engine's PA decoder
+(which would otherwise reject the high-bit scratch address). Existing benches
+wrap their HBM output via `tl.ref` → write-back unchanged (byte-equal).
+
+## Alternatives
+
+### A1. Three composites per tile (no prologue concept)
+
+Split #2 further into `softmax_merge` + `gemm(P·V)` + `add(O)` as three
+separate CompositeCmds. Pros: no flat-ops restructuring (`add` reuses
+existing epilogue path). Cons: 3 dispatches per tile instead of 2 —
+diverges from ADR-0060 §8 item 4 sizing ("split into **two**"); loses
+~32 ns × N_tiles of fixed-cost savings. **Rejected.**
+
+### A2. Mixed-engine recipe (`flash_pv_merge`)
+
+One recipe absorbing MATH + GEMM + MATH (ADR-0060 §5.6 original
+phrasing). Pros: tightest dispatch (1 composite). Cons: PE_SCHEDULER
+must know engine-boundary recipes, violates the "engines see only
+Stage.op_kind" boundary, recipe table becomes ISA-like and pollutes
+HW interface. **Rejected.**
+
+### A3. Generic op-list DAG (`tl.ex_composite([...])`)
+
+User-defined sequence of any ops with named slots and dataflow analysis.
+Pros: maximally flexible. Cons: a single user (softmax_merge) does not
+justify a generic DAG language; PE_SCHEDULER would need dataflow
+analysis; reintroduces ADR-0060 §8 item 4's "absorb the inner loop"
+trap. **Rejected (Simplicity First).**
+
+### A4. RW-aware scheduler reorder (not strict FIFO)
+
+PE_SCHEDULER could allow a later CompositeCmd to overtake an in-flight
+one when their `rw_handles` don't intersect (e.g., next-tile #1 could
+overlap with this-tile #2 since #1 doesn't touch m/l/O). Pros: more
+overlap. Cons: more scheduler state, gains are bounded (next #1 waits
+for K DMA anyway). **Deferred** — strict FIFO ships first; RW-aware
+reorder may be a future follow-up ADR.
+
+### A5. Embedding softmax_merge as an epilogue of #1 (Q·Kᵀ)
+
+Wrong: softmax_merge's output `P` is the input of #2's P·V GEMM, which
+runs *after* #1's epilogue. Embedding it would force #2 to read its own
+input from #1's epilogue scratch — circular tile-loop dependency.
+**Rejected (incorrect).**
+
+## Consequences
+
+### Positive
+- ADR-0060 §5.6 / §8 item 4 carve-out implemented exactly: opt2 = 2
+  composites per tile, K-before-V DMA priority natural.
+- Composite contract becomes uniform: flat ops list with scope-based
+  placement. No structural distinction between "head" and "epilogue"
+  in the HW cmd.
+- DMAs auto-inferred from operand `space`; kernel surface simpler.
+- HW interface (`CompositeCmd` shape) stays small even as recipes grow
+  — recipe expansion happens host-side.
+- New fused patterns (linear+gelu+norm, rmsnorm+linear, …) can be added
+  by inserting one `RECIPE_DESCRIPTORS` entry; no PE_SCHEDULER changes.
+
+### Negative
+- CompositeCmd's struct shape changes — a refactor for all current
+  callers. Semantics preserved (D6.4); existing goldens unchanged once
+  ADR-0064 Revision 2 lands.
+- `OpSpec.operands` changes from positional `tuple[Any, ...]` to
+  `dict[str, TensorHandle]` — touches the existing epilogue lowering.
+- New TLContext code: recipe expansion + scratch-slot allocation +
+  primary_out auto-bind. ~80–120 LOC.
+- PE_SCHEDULER's `_generate_plan` gets a position-based scan + DMA
+  auto-insertion + strict-FIFO RW tracker. ~60–100 LOC.
+- **Strict-FIFO RW hazard tracking is safe but conservative.** When
+  two composites share any `rw_handle`, the new one waits for *all*
+  prior overlapping composites to fully complete — even when their
+  compute could in principle overlap (e.g., next-tile #1 = Q·Kᵀ does
+  not touch m/l/O and could overlap with this-tile's #2, but FIFO
+  still serializes any path that touches the same handles down the
+  line). This **may under-expose composite-level overlap** relative to
+  a smarter scheduler. RW-aware reorder (A4) is the deferred
+  improvement.
+
+## Open review items
+
+1. **Recipe shape derivation rule beyond identity.** Future recipes may
+   need non-identity transforms (e.g., transposed primary_out). The
+   `PrimaryOutSpec.transform` field is forward-compat; add new
+   transforms as needed.
+2. **Recipe scratch allocator** integration with ADR-0063's
+   `scratch_scope`. Initial wiring: TLContext allocates inside the
+   current `scratch_scope` context if active, else in a kernel-scope
+   persistent slot.
+3. **`tl_recipes.py` location.** New module under `triton_emu/`. Keeps
+   recipe knowledge co-located with the compiler analog (TLContext).
+4. **GEMM `out=TensorHandle`.** New form alongside legacy `out_ptr:
+   int`. Both accepted; handle form preferred for new code (enables
+   strict-FIFO RW tracking by handle identity).
+
+## Test Requirements
+
+1. **CompositeCmd flat-ops refactor — meaning preservation.** Every
+   existing bench's op_log is byte-equal pre/post refactor (legacy API
+   `tl.composite(op="gemm", a, b, epilogue=[...])` lowers to the same
+   Stage sequence).
+2. **softmax_merge recipe lowering.** TLContext call `tl.composite(
+   prologue=[{"op": "softmax_merge", "s": Sj, "m": m, "l": l, "O": O}],
+   op="gemm", b=V_ref, out=O, epilogue=[{"op": "add", "other": O}])`
+   produces a `CompositeCmd` with exactly 10 ops in the expected order
+   (8 MATH + 1 GEMM + 1 MATH(add)) and `rw_handles == (m, l, O)`.
+3. **K-before-V DMA priority invariant.** op_log of decode opt2 shows
+   no V-related DMA during #2's MATH prologue (#2's V DMA begins only
+   when the GEMM head starts).
+4. **Strict-FIFO RW serialization.** Two consecutive composites sharing
+   any `rw_handle` complete in dispatch order; their stages do not
+   interleave.
+5. **DMA auto-insertion from `pinned`.** A GEMM composite with a
+   not-`pinned` operand (e.g. `b=tl.ref(...)`) emits a DMA_READ Stage; a
+   `pinned` operand (e.g. a recipe's TCM scratch / primary-out, or a
+   `tl.load` result) does not. (As-built D4 note: the DMA decision uses
+   the existing `pinned` flag, not a `space` tag.)
+6. **opt2 numeric parity with opt3.** In data mode, opt2's final
+   `(m, l, O)` matches opt3's within fp tolerance.
+7. **opt2 dispatch ratio (after ADR-0064 Rev2) — robust.** opt3 vs
+   opt2 per-tile PE_CPU dispatch cycles satisfies `opt3 > 2 × opt2`
+   (qualitative gate; calibration-independent). The default-calibration
+   model expectation is ≈ 4.0× (FIXED=40 cycles, R=0.0625 cycles/byte)
+   — see DDD-0065 §11 for the model-derived numbers; only the loose
+   `> 2×` bound is the test gate so it survives calibration changes.
+   Ratio is FIXED-dominated, reflecting command-count reduction as the
+   primary signal.
+8. **GEMM-count invariant.** A composite with two GEMM OpSpecs raises a
+   validation error at TLContext emit.
+9. **Within composite size cap (ADR-0064 D7).** Decode opt2's `#2`
+   composite (10 ops, ~322 logical bytes — per-op summation of operand
+   handles) fits comfortably under the default
+   `MAX_COMPOSITE_LOGICAL_BYTES=1024` — no segmentation. The recipe
+   lowering test asserts `cmd.logical_bytes < 1024`.
+10. **Auto-bind conflict.** `tl.composite(op="gemm", a=A_handle,
+    prologue=[{"op": "softmax_merge", ...}], ...)` where softmax_merge
+    declares `primary_out` raises a validation error at emit time
+    (D6.6).
+11. **MATH operand TCM invariant.** `tl.composite(op="math",
+    a=tl.ref(addr, shape), ...)` (passing an HBM-resident handle as a
+    MATH operand) raises a validation error at emit time (D6.7).
+
+## Dependencies
+
+- **ADR-0060** §5.6, §8 item 4 — the carve-out this ADR implements.
+- **ADR-0064 Revision 2** — dispatch cost model; lands first, makes
+  opt2's fewer-issues win measurable.
+- **ADR-0063** `scratch_scope` — recipe intermediates allocated inside;
+  `m, l, O` allocated outside as persistent arena (per DDD-0060 §6.2).
+- **ADR-0046** `tl_context` contract — extended by `prologue=[...]`
+  kwarg, `out=TensorHandle`, RECIPE_DESCRIPTORS lookup.
+- **ADR-0042** tile plan generators — extended (not redesigned) to
+  process the flat ops list and auto-insert DMAs.
+- **ADR-0014** PE pipeline — boundary preserved: scheduler stays
+  recipe-free, engines stay op_kind-opaque.
+
+## Migration
+
+> **Implementation ordering note.** The ADR-0064-first ordering below was
+> the original plan. The actual implementation **reversed steps 1 and 2** —
+> ADR-0065 Phase 1 (flat-ops `CompositeCmd`) landed *before* ADR-0064 Rev2.
+> Reason: `logical_bytes` (ADR-0064 D2) is naturally defined on the flat-ops
+> shape, so doing the refactor first avoids a throwaway transitional
+> `logical_bytes`. Each step is independently safe: P1 is a meaning-
+> preserving refactor (op_log byte-equal) under the pre-existing cost path;
+> ADR-0064 Rev2 then swaps in the structural formula on the already-flat
+> shape. The end state is identical either way.
+
+Land order (original plan; see note above for the as-built reversal of 1↔2):
+1. **ADR-0064 Revision 2** (separate PR): structural dispatch cost +
+   `logical_bytes` property + topology config override + one-time
+   goldens regeneration.
+2. **ADR-0065 Phase 1** (this ADR, PR a): `CompositeCmd` flat-ops
+   restructure + legacy-API lowering. Refactor only; goldens unchanged.
+3. **ADR-0065 Phase 2** (this ADR, PR b): `tl_recipes.py` +
+   `softmax_merge` recipe + PE_SCHEDULER position-scan + DMA
+   auto-insertion + strict-FIFO RW tracker.
+4. **ADR-0065 Phase 3** (this ADR, PR c): `_gqa_decode_long.py` opt2
+   variant + numeric parity test + dispatch-ratio measurement against
+   opt3 baseline.
+
+Detailed implementation plan: see **DDD-0065**.

docs/adr/INDEX.md

@@ -0,0 +1,175 @@
+# ADR Index
+
+Auto-generated by `tools/generate_adr_index.py`. Total ADRs: **47**.
+
+Classification mirrors the `/report` skill's section assignment. When adding a new ADR, also add an entry to the `CLASSIFICATION` table in `tools/generate_adr_index.py`.
+
+## Design Principles
+
+- [ADR-0013](./ADR-0013-ver-verification-strategy.md) — Verification Strategy and Phase 1 Test Plan
+- [ADR-0033](./ADR-0033-lat-latency-model-assumptions.md) — Latency Model: Assumptions and Known Simplifications
+
+## High-level Architecture
+
+- [ADR-0003](./ADR-0003-dev-target-system-hierarchy.md) — Target System Hierarchy & Modeling Scope  _(System hierarchy (Tray / SIP / CUBE / PE))_
+- [ADR-0007](./ADR-0007-api-runtime-api-boundaries.md) — Runtime API and Simulation Engine Boundaries  _(Runtime API ↔ sim_engine boundaries)_
+- [ADR-0016](./ADR-0016-dev-iochiplet-noc-and-memory-path.md) — IOChiplet NOC and Memory Data Path  _(IOChiplet NOC and memory data path)_
+- [ADR-0017](./ADR-0017-dev-cube-noc-and-hbm-connectivity.md) — Cube NOC and HBM Connectivity  _(Cube NOC and HBM connectivity)_
+
+## Detailed Architecture
+
+One subsection per component file under `src/kernbench/components/builtin/`.
+
+### forwarding
+
+- [ADR-0037](./ADR-0037-dev-forwarding-component.md) — Forwarding Component (forwarding_v1)
+
+### hbm_ctrl
+
+- [ADR-0034](./ADR-0034-dev-hbm-controller-internal-design.md) — HBM Controller Internal Design
+
+### io_cpu
+
+- [ADR-0036](./ADR-0036-dev-io-cpu-component-model.md) — IO_CPU Component Model
+
+### m_cpu
+
+- [ADR-0035](./ADR-0035-dev-m-cpu-and-m-cpu-dma-component-model.md) — M_CPU and M_CPU.DMA Component Model
+
+### pcie_ep
+
+- [ADR-0038](./ADR-0038-dev-pcie-ep-component-model.md) — PCIE_EP Component Model
+
+### pe_cpu
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE Pipeline Execution Model
+
+### pe_dma
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE Pipeline Execution Model
+- [ADR-0023](./ADR-0023-dev-ipcq-pe-collective.md) — PE-level IPCQ — Inter-PE Collective Communication
+
+### pe_fetch_store
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE Pipeline Execution Model
+
+### pe_gemm
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE Pipeline Execution Model
+
+### pe_ipcq
+
+- [ADR-0023](./ADR-0023-dev-ipcq-pe-collective.md) — PE-level IPCQ — Inter-PE Collective Communication
+
+### pe_math
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE Pipeline Execution Model
+
+### pe_mmu
+
+- [ADR-0039](./ADR-0039-dev-pe-mmu-component-model.md) — PE_MMU Component Model — Component + Utility Dual Role
+
+### pe_scheduler
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE Pipeline Execution Model
+
+### pe_tcm
+
+- [ADR-0040](./ADR-0040-dev-pe-tcm-component-model.md) — PE_TCM Component Model — Dual-Channel BW Serialization
+
+### sram
+
+- [ADR-0041](./ADR-0041-dev-cube-sram-component-model.md) — Cube SRAM Component Model — terminal scratchpad on cube NoC
+
+### tiling
+
+- [ADR-0042](./ADR-0042-prog-tile-plan-generators.md) — Tile Plan Generators — GEMM/Math Pipeline Plan Builders
+
+## Implementation Decisions
+
+### Address Scheme
+
+- [ADR-0001](./ADR-0001-mem-physaddr-layout.md) — 51-bit Physical Address Layout & Decoding Contract
+- [ADR-0011](./ADR-0011-mem-memory-addressing-simplification.md) — Memory Addressing — PA / VA / LA Address Models
+
+### Routing & Helper API
+
+- [ADR-0002](./ADR-0002-lat-routing-distance.md) — Routing Distance, Ordering & Bypass Rules
+- [ADR-0051](./ADR-0051-lat-routing-helper-api.md) — Routing Helper API — `AddressResolver` + `PathRouter`
+
+### Memory Semantics & Local-HBM Bandwidth
+
+- [ADR-0004](./ADR-0004-mem-memory-semantics-local-hbm.md) — Memory Semantics & Local-HBM Bandwidth Guarantee
+
+### Topology Compilation, Diagrams & Builder Algorithms
+
+- [ADR-0005](./ADR-0005-dev-diagram-views-distance-layout.md) — Diagram Views & Distance-Aware Layout Rules
+- [ADR-0006](./ADR-0006-dev-topology-compilation-distance-diagram.md) — Topology Compilation, Distance Extraction, and Automatic Diagram Generation
+- [ADR-0053](./ADR-0053-dev-topology-builder-algorithms.md) — Topology Builder + Visualizer Algorithms
+
+### Tensor Deployment and Allocation
+
+- [ADR-0008](./ADR-0008-api-tensor-deploy-and-allocation.md) — Tensor Deployment and Allocation (Host Allocator, PA-first)
+
+### Kernel Execution and Host-Device Messaging
+
+- [ADR-0009](./ADR-0009-api-kernel-execution-messaging.md) — Kernel Execution Messaging and Completion Semantics
+- [ADR-0012](./ADR-0012-api-host-io-message-schema.md) — Host ↔ IO_CPU Message Schema (PA-first, PE-tagged)
+
+### CLI Surface and Semantics
+
+- [ADR-0010](./ADR-0010-api-cli-surface-and-semantics.md) — Command Line Interface and Execution Semantics
+
+### Component Port/Wire Fabric Model
+
+- [ADR-0015](./ADR-0015-dev-component-port-wire-model.md) — Component Port/Wire Model and Fabric Routing
+
+### Two-Pass Data Execution
+
+- [ADR-0020](./ADR-0020-prog-data-execution-two-pass.md) — 2-Pass Data Execution Model (Timing / Data Separation)
+
+### 2D Grid Program Identity
+
+- [ADR-0022](./ADR-0022-prog-program-id-2d-grid.md) — 2D Grid program_id Semantics
+
+### Parallelism (Launcher, DP, TP, AHBM backend, CCL algorithm)
+
+- [ADR-0024](./ADR-0024-par-sip-tp-launcher.md) — SIP-level Launcher — rank = SIP
+- [ADR-0026](./ADR-0026-par-dppolicy-intra-device.md) — DPPolicy = Intra-Device Only — remove sip/num_sips fields
+- [ADR-0027](./ADR-0027-par-megatron-tp.md) — Megatron-style Tensor Parallelism API
+- [ADR-0047](./ADR-0047-par-ahbm-ccl-backend.md) — AHBM CCL Backend — `torch.distributed`-compat shim
+- [ADR-0050](./ADR-0050-par-ccl-algorithm-module-contract.md) — CCL Algorithm Module Contract — `ccl/algorithms/*.py`
+
+### IPCQ Direction Addressing
+
+- [ADR-0025](./ADR-0025-algo-ipcq-direction-addressing.md) — IPCQ Direction Addressing — address-based matching
+
+### Intercube All-Reduce
+
+- [ADR-0032](./ADR-0032-algo-intercube-allreduce.md) — Intercube All-Reduce — pe0 cube-mesh reduce + multi-SIP exchange
+
+### Evaluation Harnesses
+
+- [ADR-0043](./ADR-0043-eval-allreduce-harness.md) — Allreduce Evaluation Harness — `tests/sccl/`
+- [ADR-0044](./ADR-0044-eval-gemm-harness.md) — GEMM Evaluation Harness — `scripts/gemm_sweep.py` + `tests/gemm/`
+- [ADR-0054](./ADR-0054-eval-milestone-benches.md) — Milestone Eval Benches — self-contained sweep + figure benches
+
+### Bench Module Contract
+
+- [ADR-0045](./ADR-0045-prog-bench-module-contract.md) — Bench Module Contract — registration, dispatch, and authoring
+
+### Kernel-side tl.* API (TLContext)
+
+- [ADR-0046](./ADR-0046-prog-tl-context-contract.md) — TLContext — Kernel-side `tl.*` API Contract
+
+### Memory Allocator Algorithms
+
+- [ADR-0048](./ADR-0048-mem-allocator-algorithms.md) — Memory Allocator Algorithms — VirtualAllocator + PEMemAllocator
+
+### Probe Subcommand
+
+- [ADR-0049](./ADR-0049-ver-probe-subcommand.md) — `kernbench probe` Subcommand — Traffic-Pattern Verification Harness
+
+### Sim-engine Op Log and Memory Store Schemas
+
+- [ADR-0052](./ADR-0052-dev-oplog-memory-store-schemas.md) — OpLog + MemoryStore Schemas — sim_engine internals

docs/ddd/DDD-0065-flat-ops-composite-softmax-merge-detailed-design.md

@@ -0,0 +1,699 @@
+# Detailed Design Document — Flat-ops `CompositeCmd` + `softmax_merge` recipe (ADR-0065)
+
+**Status:** Accepted (companion to ADR-0065, Accepted) — reconciled to
+the as-built implementation.
+**Scope:** the `CompositeCmd` flat-ops restructure, the `softmax_merge`
+recipe, the decode opt2 variant in `_gqa_attention_decode_opt2.py`, and the
+data-mode numeric path that makes opt2's recipe compute.
+**Audience:** reviewer / implementer. Read **ADR-0065** and **ADR-0064
+Revision 2** first.
+
+---
+
+## 0. As-built reconciliation (read first)
+
+This DDD was authored in the design round; the implementation diverged in
+several places. The sections below have been corrected, but the headline
+deltas are collected here. (ADR-0065 itself was kept in sync; see commits
+`47e2c78`,`184f654`,`55f025c`,`17fb940`,`35453cc`,`4a55ae5`,`e8d6c28`,
+`4089e18`,`3d4a3d4`,`5bbe293`,`e9d62b9`,`f0c5b9c`.)
+
+| Topic | Original DDD | As-built |
+| --- | --- | --- |
+| **Phase order** | P0 (ADR-0064 Rev2) lands first | **P1 (flat-ops) first**, then P0 — `logical_bytes` is naturally defined on the flat shape, so the refactor goes first |
+| **Operand-input DMA** | from `handle.space` (`space=="hbm"`→DMA_READ) | **from the `pinned` flag** (already-in-TCM→skip). The `space` flip was avoided (it would break byte-equality); `tl.ref`→hbm, but the *decision* stays pinned-based |
+| **Output DMA / handle** | (not addressed; `out=O`) | **ADR-0065 D8** — `tl.composite` returns the output `TensorHandle`; `out`=handle (HBM via `tl.ref`) \| `out_ptr` (HBM shorthand) \| omitted→auto-TCM; **DMA_WRITE gated on `out.space=="hbm"`** (TCM output stays on-chip, chainable) |
+| **Size cap (D7)** | deterministic segmentation | **hard cap → `ValueError`** at emit (no segmentation) |
+| **Prologue execution** | feeder yields `prologue_stages` first | feeder feeds each prologue stage as a **standalone 1-stage tile** (folding broke PE_MATH→PE_DMA routing) |
+| **D6.7 (MATH TCM-only)** | every non-GEMM op | **prologue recipe ops only**; the head op (gemm *or* math) keeps DMA-staged-from-HBM |
+| **Numeric parity** | "just works" via the existing data executor | composites did **not** compute in data mode at all — built in **N1–N4** (below): composite-as-gemm record, recipe MATH compute, GEMM accumulate, opt2 ordering |
+| **Dispatch ratio** | model ≈ 4.0× | **measured 5.22×** per-tile (gate `>2×`); R-sweep 4.20×/5.22×/5.45× |
+| **opt2↔opt3 parity** | exact in data mode | exact recipe math verified by N2/N3 (isolated, non-trivial); e2e opt2↔opt3 holds in the **near-uniform regime** (kernbench K reshape-as-transpose caveat, shared by both kernels) |
+
+**Data-mode numerics (N1–N4), absent from the original DDD — see §15.**
+
+---
+
+## 1. Goal and success criteria
+
+**Goal:** make decode opt2 *runnable*, *numerically equivalent to opt3*,
+and *measurably cheaper in PE_CPU dispatch* under the ADR-0064 Rev2 cost
+model. Existing benches keep their (post-ADR-0064 Rev2) goldens
+unchanged.
+
+**Success criteria:**
+1. **Refactor safety.** `CompositeCmd` flat-ops restructure preserves the
+   op_log byte-for-byte for every existing bench
+   (`tests/<existing>` all green, no goldens diff).
+2. **Recipe correctness.** opt2 final `(m, l, O)` matches opt3 within fp
+   tolerance in data mode (`enable_data=True`).
+3. **Boundary preservation.** PE_SCHEDULER imports nothing recipe-related;
+   PE_MATH/PE_GEMM/PE_DMA see only `Stage.params["op_kind"]` (no new
+   engine code).
+4. **Dispatch ratio (robust).** Per-tile PE_CPU dispatch cycles satisfy
+   `opt3 > 2 × opt2` (qualitative invariant, calibration-independent).
+   Default ADR-0064 Rev2 calibration (FIXED=40 cycles, R=0.0625
+   cycles/byte) gives model-expected ≈ 4.0× — informative only, not the
+   test gate.
+5. **K-before-V invariant.** op_log of decode opt2 shows zero V-related
+   DMA during #2's MATH prologue.
+
+**Non-goals:** prefill (causal `if` in kernel control flow, separate
+work); RW-aware scheduler reorder (strict FIFO ships); user-defined
+recipes; new engine micro-ops in PE_MATH (it stays op_kind-opaque).
+
+---
+
+## 2. Where the design sits in the current code
+
+```
+runtime_api (host)
+  RuntimeContext.zeros/empty/launch          context.py
+        │  KernelLaunchMsg
+        ▼
+sim_engine
+  GraphEngine(enable_data=…)                 engine.py
+  MemoryStore                                memory_store.py
+        │
+        ▼
+components / PE pipeline (per PE)
+  PE_CPU → PE_SCHEDULER → PE_{DMA,GEMM,MATH,IPCQ,TCM}
+  greenlet kernel ↔ SimPy                    kernel_runner.py, pe_cpu.py
+  CompositeCmd → tile plan → engines         pe_scheduler.py (ADR-0014 D6)
+        │
+        ▼
+tl programming model                         tl_context.py
+  composite / dot / load / store / scratch_scope
+  ── ADR-0065 NEW ──────────────────────────
+  tl_recipes.py                              RECIPE_DESCRIPTORS, EngineOp,
+                                             RecipeDescriptor, PrimaryOutSpec
+```
+
+The new module `tl_recipes.py` sits beside `tl_context.py`. PE_SCHEDULER
+(`pe_scheduler.py`) and `tiling.py` extend to consume the flat ops list
+without importing `tl_recipes`.
+
+---
+
+## 3. File plan
+
+### 3.1 Modified production files
+
+| File | Change | ADR ref |
+|---|---|---|
+| `src/kernbench/common/pe_commands.py` | `CompositeCmd` flat-ops dataclass; `OpSpec.operands: dict`; new `Scope` value semantics (KERNEL position-based); `logical_bytes` property on every `PeCommand`. | 0065 D1/D2, 0064 D2 |
+| `src/kernbench/triton_emu/tl_context.py` | `composite(prologue=[...], out=TensorHandle, …)`; recipe lowering pass (RECIPE_DESCRIPTORS lookup, scratch alloc, OpSpec expansion, primary_out auto-bind); legacy API preserved (`op="gemm", a, b, epilogue=[...]` still works, internally lowers to flat ops). | 0065 D1/D5 |
+| `src/kernbench/triton_emu/tl_recipes.py` (NEW) | `RecipeDescriptor`, `PrimaryOutSpec`, `EngineOp` dataclasses; `RECIPE_DESCRIPTORS["softmax_merge"]` entry; expansion helper. | 0065 D5 |
+| `src/kernbench/components/builtin/pe_scheduler.py` | `_generate_plan` rewritten for flat ops: scan for GEMM (≤1), position-based KERNEL placement, OUTPUT_TILE/K_TILE epilogue, strict-FIFO RW tracker keyed by handle id. | 0065 D3/D6 |
+| `src/kernbench/components/builtin/tiling.py` | `generate_plan_from_ops(ops, tile_*, bpe, pe_prefix)` wraps `generate_gemm_plan`/`generate_math_plan`. Input DMA_READ from the **`pinned`** flag; **output DMA_WRITE gated on `out.space=="hbm"`** (D8). `_math_stage` carries operand+output addrs on recipe prologue stages (N2). | 0065 D3/D4/D8 |
+| `src/kernbench/common/pe_cost_model.py` (NEW) | `PeCostModel` (FIXED/R/cap) + `from_node_attrs`; replaces the removed Rev1 `cpu_issue_cost.py`. | 0064 D1/D4 |
+| `src/kernbench/components/builtin/pe_cpu.py` | reads `pe_cost_model:` + `clock_freq_ghz` from node attrs, threads to TLContext (legacy + greenlet KernelRunner). | 0064 D1/D4 |
+| `src/kernbench/sim_engine/op_log.py` | `_extract_op_info(CompositeCmd)` scans for the GEMM op + `accumulate` flag (N1/N3); `record_end` promotes an address-carrying MATH stage to op_kind="math" (N2). | 0065 N1-N3 |
+| `src/kernbench/sim_engine/data_executor.py` | `_compute_math` +5 recipe ops; `_execute_gemm` best-effort skip + accumulate (N1-N3). | 0065 N1-N3 |
+| `src/kernbench/benches/_gqa_attention_decode_opt2.py` (NEW) | opt2 variant — tile 0 establishes (m,l,O) with primitives; the merge tile uses `#1 scores = tl.composite(op="gemm", a=Q, b=K_T1)` + `#2 tl.composite(prologue=[softmax_merge], op="gemm", b=V_ref, out=O, epilogue=[add])`. | 0065 §Migration P3 |
+
+### 3.2 New tests
+
+| File | Covers |
+|---|---|
+| `tests/test_composite_flat_ops_refactor.py` | P1 flat-ops lowering shape (legacy + math + epilogue) |
+| `tests/test_pe_cost_model.py` | ADR-0064 Rev2 — `logical_bytes`/formula, override, `PeCpuOverheadCmd` bypass, D7 cap→error, live wiring/path-parity |
+| `tests/test_tl_recipe_softmax_merge.py` | recipe lowering — 10 ops in order + `rw_handles == (m, l, O)` + size cap |
+| `tests/test_tl_composite_validation.py` | auto-bind conflict (D6.6); prologue MATH operand TCM-only (D6.7) |
+| `tests/test_pe_scheduler_flat_ops_plan.py` | position-scan; pinned-based DMA; prologue stages; MATH-only |
+| `tests/test_pe_scheduler_strict_fifo.py` | `_RwHazardTracker` admit/retire (RAW/WAW/disjoint/legacy) |
+| `tests/test_composite_output_handle.py` | D8 — composite returns handle (auto-TCM); chaining; output-space DMA; `tl.ref`→hbm |
+| `tests/test_recipe_data_mode.py` | **N2/N3** — recipe MATH compute m/l/O-rescale + full accumulate (O = O·corr + P@V) vs numpy |
+| `tests/test_e2e_pipeline.py` | **N1** — composite GEMM computes a@b in data mode |
+| `tests/attention/test_gqa_decode_opt2.py` | dispatch ratio (>2×) + R-sweep; K-before-V; **N4** opt2↔opt3 parity |
+
+---
+
+## 4. Data model — `CompositeCmd` flat-ops
+
+### 4.1 New shape
+
+```python
+@dataclass(frozen=True)
+class CompositeCmd:
+    completion: CompletionHandle
+    ops: tuple[OpSpec, ...]                       # ordered MATH/GEMM ops
+    rw_handles: tuple[TensorHandle, ...] = ()     # cross-composite hazard
+    data_op: bool = True
+
+    @property
+    def logical_bytes(self) -> int:
+        return (
+            4                                     # framing
+            + 1 + sum(op.logical_bytes for op in self.ops)
+            + 1 + 8 * len(self.rw_handles)
+        )
+```
+
+The previous `op`, `a`, `b`, `out_addr`, `out_nbytes`, `math_op` fields
+are **removed**. All info lives in `ops[i]`.
+
+### 4.2 `OpSpec` extended
+
+```python
+@dataclass(frozen=True)
+class OpSpec:
+    kind: str                                     # "gemm" | "rmax" | "exp_diff" | ...
+    scope: Scope                                  # KERNEL | K_TILE | OUTPUT_TILE
+    operands: dict[str, TensorHandle]
+    extra: dict[str, Any] = field(default_factory=dict)
+    out: TensorHandle | None = None
+
+    @property
+    def logical_bytes(self) -> int:
+        return (
+            1 + 1                                        # opcode + scope enum
+            + 1 + 8 * len(self.operands)                 # named handles
+            + (8 if self.out is not None else 0)
+            + 1 + 4 * len(self.extra)                    # scalars
+        )
+```
+
+Backward compat: existing `operands: tuple[Any, ...]` use sites
+(epilogue `bias`, `scale`, `add`, etc.) migrate to `operands={"other":
+h}` / `operands={"src": h}` per a small lookup in
+`pe_commands.EPILOGUE_OPS`.
+
+### 4.3 `Scope` semantics — phase vs repetition (ADR-0065 D3)
+
+`Scope` keeps its three values (`KERNEL`, `K_TILE`, `OUTPUT_TILE`).
+What changes is the **semantics split**:
+- **Position** determines the *phase* (where in the tile-loop lifecycle
+  the op fires — pre-loop / in-loop / post-loop).
+- **Scope** determines the *repetition* within the phase
+  (KERNEL = once per `CompositeCmd` invocation; K_TILE = per K-tile
+  iteration; OUTPUT_TILE = per output tile).
+
+`KERNEL` does *not* mean "once per kernel launch" or "once anywhere".
+It means "once per `CompositeCmd` invocation, at this op's position".
+A composite with multiple KERNEL-scope ops (e.g., softmax_merge's 8
+prologue ops) executes each exactly once, in sequence order, before
+the GEMM tile loop begins.
+
+Concretely:
+- KERNEL ops *before* GEMM → pre-tile-loop (single shot, in order)
+- KERNEL ops *after* GEMM → post-tile-loop (single shot, in order)
+- K_TILE / OUTPUT_TILE ops *after* GEMM → inside tile loop as today
+
+---
+
+## 5. Recipe: `softmax_merge`
+
+### 5.1 `tl_recipes.py` definitions
+
+```python
+@dataclass(frozen=True)
+class PrimaryOutSpec:
+    from_shape: str          # operand name to copy shape from
+    from_dtype: str          # operand name to copy dtype from
+    transform: str           # "identity" | "trans" (forward-compat)
+
+
+@dataclass(frozen=True)
+class EngineOp:
+    engine: Literal["MATH", "GEMM", "DMA"]
+    op_kind: str             # PE_MATH op_kind label (opaque to engine)
+    # Slot names: any of "src", "src_a", "src_b", "src_c", "dst" — these
+    # are RecipeDescriptor.operands names or internal scratch names.
+    src: str | None = None
+    src_a: str | None = None
+    src_b: str | None = None
+    src_c: str | None = None
+    dst: str | None = None
+    reduce_axis: int | None = None
+    bcast_axis: int | None = None
+
+
+@dataclass(frozen=True)
+class RecipeDescriptor:
+    operands: dict[str, str]                            # name → "R" | "RW"
+    primary_out: PrimaryOutSpec | None                  # implicit output spec
+    tile_alignment: Literal["single_shot", "tile_aligned"]
+    internal_scratch_bytes_fn: Callable[..., int]
+    engine_seq: tuple[EngineOp, ...]
+
+
+RECIPE_DESCRIPTORS: dict[str, RecipeDescriptor] = { ... }
+```
+
+### 5.2 Engine sequence (8 MATH ops)
+
+| # | engine | op_kind | reads (slots) | writes (slots) | extras |
+|---|---|---|---|---|---|
+| 1 | MATH | `rmax`       | s              | m_loc          | `reduce_axis=-1` |
+| 2 | MATH | `max_elem`   | m, m_loc       | m_new          | — |
+| 3 | MATH | `exp_diff`   | m, m_new       | corr           | — |
+| 4 | MATH | `exp_diff`   | s, m_new       | P (primary_out)| `bcast_axis=0` |
+| 5 | MATH | `rsum`       | P              | l_loc          | `reduce_axis=-1` |
+| 6 | MATH | `fma`        | l, corr, l_loc | l              | — |
+| 7 | MATH | `mul_bcast`  | O, corr        | O              | `bcast_axis=1` |
+| 8 | MATH | `copy`       | m_new          | m              | — |
+
+Intermediates `m_loc, m_new, corr, P, l_loc` are scratch slots allocated
+by TLContext before lowering. `P` is the recipe's primary output; auto-
+bound to the head GEMM's `operands["a"]`.
+
+### 5.3 PE_MATH op_kind additions
+
+PE_MATH is op_kind-opaque ([pe_math.py:90-99](src/kernbench/components/builtin/pe_math.py#L90-L99)).
+Latency = `ceil(num_elements / vector_width) / clock`. Adding new
+op_kind labels (rmax, max_elem, exp_diff, rsum, fma, mul_bcast, copy)
+requires **no PE_MATH code changes** — they pass through as string
+labels in `Stage.params["op_kind"]`. Latency stays consistent for the
+same `num_elements`.
+
+(Reductions like `rmax` / `rsum` have log-depth final stages in real HW
+not modelled by `ceil(num_elements / vector_width)`; this is an existing
+approximation, not introduced by ADR-0065.)
+
+---
+
+## 6. TLContext lowering pipeline
+
+Input (user-facing API, preserved):
+```python
+tl.composite(
+    op="gemm", b=tl.ref(v_tile(j), (TILE, d)), out=O,
+    prologue=[{"op": "softmax_merge", "s": Sj, "m": m, "l": l, "O": O}],
+    epilogue=[{"op": "add", "other": O}],
+)
+```
+
+Lowering steps in `TLContext.composite()`:
+
+1. **Validate prologue items.** For each, lookup
+   `RECIPE_DESCRIPTORS[item["op"]]`. Check all operands present, type
+   tags (R / RW) match handle states (e.g., RW must not be `tl.ref`).
+2. **Allocate internal scratch.** Call
+   `recipe.internal_scratch_bytes_fn(G, TILE, d, bpe)`; allocate via
+   `scratch_scope` helper (ADR-0063). Get slot addresses for `m_loc,
+   m_new, corr, P, l_loc`.
+3. **Derive primary_out handle.** From `recipe.primary_out`: `P` handle =
+   `TensorHandle(shape=Sj.shape, dtype=Sj.dtype, addr=<P slot>,
+   space="tcm")`.
+4. **Expand `engine_seq` → flat OpSpecs.** For each `EngineOp`:
+   - Resolve slot names (operands or scratch) to `TensorHandle`s.
+   - Build `OpSpec(kind=op_kind, scope=Scope.KERNEL, operands={...},
+     extra={...}, out=...)`.
+   - `out` is the handle for `dst` slot; `operands` is `{src*: handle}`.
+5. **Auto-bind head GEMM (with conflict check, ADR-0065 D6.6).** GEMM
+   OpSpec built from `op="gemm"`, `operands={"a": P_handle, "b":
+   V_handle}`, `out=<output handle>`, `extra={"m": G, "k": TILE, "n":
+   d}`. **Validation:** if the kernel passed `a=<some handle>`
+   explicitly AND there is a prologue recipe with `primary_out` that
+   would auto-bind into `a`, raise a validation error (auto-bind
+   conflict). The kernel must omit `a` (let recipe bind) or omit the
+   prologue / use a `primary_out=None` recipe (no auto-bind).
+6. **Build epilogue OpSpecs.** Each `epilogue=[{...}]` entry maps to
+   `OpSpec(kind=..., scope=Scope.OUTPUT_TILE, operands={...}, ...)`.
+   The `add` epilogue specifically gets `operands={"other": O_handle}`
+   so the Stage later carries `other_handle` (GAP 9 fix).
+7. **Compute `rw_handles`.** Collect all handles tagged RW in prologue
+   recipes + the head GEMM's `out` (if accumulating) + any RW epilogue
+   targets. For softmax_merge: `rw_handles = (m, l, O)`.
+8. **Operand validation (ADR-0065 D6.7, as-built scope).** Checked on the
+   **prologue recipe ops only**: every operand handle must be
+   `space == "tcm"` (HBM recipe operand → `ValueError`). The head op (gemm
+   *or* math) is **exempt** — it may stream from HBM. softmax_merge's 8 MATH
+   ops only reference TCM-resident handles → the invariant naturally holds.
+   Also: GEMM-count ≤ 1 (D6.1).
+9. **Check composite size cap (ADR-0064 D7, as-built).** If
+   `cmd.logical_bytes > max_composite_logical_bytes` (default 1024), raise
+   `ValueError` at emit — **hard cap, no auto-segmentation**. softmax_merge's
+   10-op composite (~322 bytes) is well under the cap.
+10. **Resolve the output handle (ADR-0065 D8) + emit.** `out`=handle (HBM
+    `tl.ref` \| in-place TCM) \| `out_ptr` (HBM shorthand) \| omitted→auto-TCM
+    scratch. Emit `CompositeCmd(ops=(8 MATH + 1 GEMM + 1 MATH(add)),
+    rw_handles=(m,l,O))`, attach a `CompositeFuture` (carrying the
+    completion) to the output handle, and **return that handle** so the
+    result chains like `tl.dot`'s (downstream auto-awaits via
+    `_await_pending`).
+
+**Legacy lowering.** Pre-existing
+`tl.composite(op="gemm", a=A, b=B, epilogue=[{"op":"bias","bias":h}])`
+goes through the same emit path: no prologue, head OpSpec from `op="gemm"`,
+epilogue OpSpecs built from `epilogue=[...]`. Output `CompositeCmd` has
+flat ops list; semantics preserved.
+
+---
+
+## 7. PE_SCHEDULER tile plan generation (flat-ops)
+
+```python
+def _generate_plan(cmd: CompositeCmd) -> PipelinePlan:
+    pp = self._pe_prefix
+    bpe = 2
+
+    # 1. Identify GEMM op (≤1).
+    gemm_idx = next(
+        (i for i, o in enumerate(cmd.ops) if o.kind == "gemm"),
+        None,
+    )
+
+    # 2. MATH-only composite (no GEMM): single tile of all ops in order.
+    if gemm_idx is None:
+        return _math_only_plan(cmd.ops, pp)
+
+    gemm = cmd.ops[gemm_idx]
+    M, K, N = gemm.extra["m"], gemm.extra["k"], gemm.extra["n"]
+
+    # 3. Partition ops by position.
+    pre_loop_ops  = cmd.ops[:gemm_idx]              # KERNEL scope
+    after_ops     = cmd.ops[gemm_idx+1:]            # mixed scopes
+
+    # 4. Pre-loop MATH Stages (single-shot, before tile loop).
+    pre_stages = [_math_stage_from_op(op, pp) for op in pre_loop_ops]
+
+    # 5. Tile loop (existing generate_gemm_plan pattern, but:
+    #    - DMA_READ inserted only when operand.space == "hbm"
+    #    - K_TILE epilogue from after_ops with scope == K_TILE
+    #    - OUTPUT_TILE epilogue from after_ops with scope == OUTPUT_TILE
+    #    - DMA_WRITE inserted only when out.space == "hbm")
+    tile_plans = _generate_tile_loop(
+        gemm, M, K, N, after_ops, pp, bpe,
+    )
+
+    # 6. Post-loop KERNEL Stages (after tile loop completes).
+    post_loop_ops = [op for op in after_ops if op.scope == Scope.KERNEL]
+    post_stages   = [_math_stage_from_op(op, pp) for op in post_loop_ops]
+
+    return PipelinePlan(
+        prologue_stages=tuple(pre_stages),
+        tiles=tile_plans,
+        epilogue_stages=tuple(post_stages),
+        m_tiles=..., k_tiles=..., n_tiles=...,
+    )
+```
+
+`PipelinePlan` gets two new fields:
+
+```python
+@dataclass(frozen=True)
+class PipelinePlan:
+    tiles: list[TilePlan]
+    m_tiles: int
+    k_tiles: int
+    n_tiles: int
+    prologue_stages: tuple[Stage, ...] = ()
+    epilogue_stages: tuple[Stage, ...] = ()
+```
+
+Feeder change (`_feed_loop`): yield `prologue_stages` Stages first, then
+tile loop as today, then `epilogue_stages`. For the single-Stage case
+(no prologue/epilogue), behaviour identical to today.
+
+---
+
+## 8. Strict-FIFO RW hazard tracker
+
+```python
+class _RwHazardTracker:
+    def __init__(self):
+        self._inflight: list[tuple[CompletionHandle, set[str]]] = []
+        # (completion, set of handle ids written or RW)
+
+    def admit(self, env, cmd: CompositeCmd) -> Generator:
+        """Block until no in-flight composite has overlapping RW."""
+        rw_ids = {h.id for h in cmd.rw_handles}
+        # Strict FIFO: wait for ALL prior composites whose rw_set
+        # intersects, in order.
+        while True:
+            blocking = [c for c, s in self._inflight if s & rw_ids]
+            if not blocking:
+                break
+            # Wait on the *first* blocking composite's completion.
+            yield blocking[0].done_event
+        self._inflight.append((cmd.completion, rw_ids))
+
+    def retire(self, completion: CompletionHandle) -> None:
+        self._inflight = [
+            (c, s) for c, s in self._inflight if c != completion
+        ]
+```
+
+`PE_SCHEDULER` invokes `admit()` before queueing a CompositeCmd's tile
+plan and `retire()` upon `done_event.succeed()`. Strict FIFO: even when
+overlap is partial, the new composite waits for all prior overlapping
+composites to *fully* drain — simple, easy to reason about, easy to
+test. (RW-aware reorder deferred per ADR-0065 A4.)
+
+**Conservatism trade-off (ADR-0065 Consequences/Negative).** This rule
+is safe but may under-expose overlap. Concretely for decode opt2:
+next-tile's #1 (Q·Kᵀ) does not touch `(m, l, O)` and *could* in
+principle run while this-tile's #2 is still computing on those handles.
+Under strict FIFO, next-tile's #1 still gets dispatched (no rw overlap
+with prior in-flight) — so #1 cross-tile pipelining is preserved.
+However, if a future bench composed multiple #2-like recipes that share
+handles, FIFO would serialize them even when their compute could
+overlap. Acceptable for Phase 1; revisit when a real workload shows
+the gap.
+
+---
+
+## 9. Phased implementation plan (test-first per CLAUDE.md)
+
+Each phase is an independent Phase-1 → Phase-2 cycle per CLAUDE.md Part 1.
+
+| Phase | Deliverable | Gate |
+|---|---|---|
+| **P0 (= ADR-0064 Rev2)** | `logical_bytes` property + dispatch formula + topology override + size cap; one-time goldens regeneration | typical 1-OpSpec composite ≈ 43 ns; dispatch ratio test infrastructure in place |
+| **P1** | `CompositeCmd` flat-ops dataclass refactor + `OpSpec.operands: dict` + legacy `tl.composite` lowering preserved | every existing bench's op_log byte-equal pre/post |
+| **P2** | `tl_recipes.py` + `RECIPE_DESCRIPTORS["softmax_merge"]` + TLContext lowering for `prologue=[...]` | recipe lowering test (10 ops, rw_handles, addresses) passes |
+| **P3** | PE_SCHEDULER `_generate_plan` flat-ops + `PipelinePlan.prologue_stages` + DMA auto-insertion + `_feed_loop` extension | all existing bench Stage sequences unchanged; new MATH chain visible in op_log for opt2 mock |
+| **P4** | Strict-FIFO RW hazard tracker + `_RwHazardTracker` integration | two-composite RW conflict test serializes correctly |
+| **P5** | `_gqa_decode_long.py` opt2 variant; data-mode numeric parity check; K-before-V invariant check | opt2 matches opt3 within fp tolerance; no V DMA during prologue |
+| **P6** | Dispatch-ratio measurement: opt3 vs opt2 per-tile PE_CPU cycles + R sensitivity sweep | `opt3 > 2 × opt2` at default calibration (gate); model-expected ≈ 4.0× (informative); `opt2 < opt3` at all `R ∈ {0.25, 0.0625, 0.03125}` (ADR-0064 Test #9) |
+
+**As-built ordering:** the implementation reversed P0↔P1 — **P1 (flat-ops
+refactor) landed first**, then P0 (ADR-0064 Rev2 cost model), because
+`logical_bytes` is naturally defined on the flat shape (doing P1 first avoids
+a throwaway transitional definition). P1 is a pure, byte-equal refactor; P0
+then swaps in the structural formula and regenerates goldens (in practice no
+golden churn occurred). P5 was split: **P5(B)** shipped the opt2 kernel +
+op_log dispatch measurement first (numeric parity deferred), then **D8 +
+N1–N4** (§15) built the data-mode numeric path. P6 closed the measurement.
+
+---
+
+## 10. Verification plan (concrete)
+
+Mirrors ADR-0065 §Test Requirements; grounded in SPEC R2/R5, ADR-0023/
+0025, ADR-0046, ADR-0054.
+
+**Refactor safety (P1):**
+- Parametrize over `tests/<existing bench>.py` baseline op_logs.
+- Pre/post refactor diff must be empty for *every* bench.
+- Failure → revert P1 atomically; do not land partial.
+
+**Recipe lowering (P2):**
+- Given `Sj=(8,64), m=(8,), l=(8,), O=(8,128)`, lower
+  `tl.composite(prologue=[{"op":"softmax_merge", "s":Sj,...}], op="gemm",
+   b=V_ref, out=O, epilogue=[{"op":"add", "other":O}])`.
+- Assert `len(cmd.ops) == 10`, `cmd.ops[0..7].kind == [rmax, max_elem,
+  exp_diff, exp_diff, rsum, fma, mul_bcast, copy]`,
+  `cmd.ops[8].kind == "gemm"`, `cmd.ops[9].kind == "add"`.
+- Assert `cmd.rw_handles == (m, l, O)`.
+- Assert all OpSpec addresses are concrete (no None).
+
+**Scheduler plan (P3):**
+- For the above CompositeCmd, assert PipelinePlan has 8 prologue_stages
+  (MATH), tile loop with FETCH/GEMM/MATH(add), no epilogue_stages.
+- Assert no DMA_READ for prologue MATH stages (operands TCM-resident).
+- Assert one DMA_READ for V (`b.space == "hbm"`).
+- Assert no DMA_WRITE (`out.space == "tcm"`).
+
+**RW hazard (P4):**
+- Submit composite A (rw_handles=(O,)) then composite B (operands
+  reading O). Assert B's Stages start only after A's done_event.
+
+**Numeric parity (P5):**
+- In data mode, run opt2 and opt3 for `(G=8, T_q=1, S_kv=64, C=1, P=1)`.
+- Assert `O_opt2 ≈ O_opt3` within `fp16` tolerance.
+
+**K-before-V invariant (P5):**
+- op_log inspection: for each tile's #2 CompositeCmd, no
+  `memory/dma_read` events between #2 start and the GEMM compute begin
+  that are not the V tile DMA itself.
+
+**Dispatch ratio (P6):**
+- For `S_kv=64, n_tiles=16`, measure total PE_CPU dispatch cycles for
+  opt3 vs opt2 paths.
+- Assert `opt3 > 2 × opt2` (robust gate — calibration-independent).
+- Record observed ratio; the default-calibration model expects ≈ 4.0×
+  (DDD §11). The recorded number is informative for performance
+  tracking, not a test gate.
+- **Sensitivity sweep.** Repeat with `R ∈ {0.25, 0.0625, 0.03125}`
+  cycles/byte; assert opt2 < opt3 in all three. (ADR-0064 Test #9.)
+
+**Invariant guards (P2):**
+- Auto-bind conflict: `tl.composite(op="gemm", a=A_handle,
+  prologue=[{"op": "softmax_merge", "s": Sj, "m": m, "l": l, "O": O}],
+  ...)` → `pytest.raises(ValidationError, match="auto-bind conflict")`.
+- MATH operand TCM-only: `tl.composite(op="math", a=tl.ref(addr,
+  shape), ...)` → `pytest.raises(ValidationError, match="MATH operand
+  must be TCM")`.
+- Per-op logical_bytes summation: a synthetic composite with the same
+  handle `O` in 3 OpSpecs + in `rw_handles` reports
+  `logical_bytes` higher than a deduplicated equivalent by exactly
+  `3 × 8` (operand references) + `8` (rw_handles entry) = 32 bytes.
+
+**Determinism (all phases):**
+- Identical inputs → identical op_log + latency (SPEC §0.1).
+
+---
+
+## 11. Performance model (expected)
+
+Per-tile PE_CPU dispatch cycles, ADR-0064 Rev2 default calibration
+(FIXED=40 cycles, R=0.0625 cycles/byte, 1 cycle = 1 ns at 1 GHz):
+
+```
+opt3: 10 cmds × 40 + total_bytes(≈232) × 0.0625 = 400  + 14.5 ≈ 414.5 ns
+opt2:  2 cmds × 40 + total_bytes(≈380) × 0.0625 =  80  + 23.75 ≈ 103.75 ns
+ratio = 414.5 / 103.75 ≈ 4.0×
+```
+
+The FIXED term dominates (~96% of the differential), reflecting
+ADR-0064's framing: the primary signal is command-count reduction.
+The byte term is a secondary refinement (visible: opt2 carries more
+total bytes than opt3, so its byte cost is higher; this is more than
+offset by 8× fewer FIXED payments).
+
+Per-tile engine timing dominates total latency; dispatch is one
+component. The ratio above is the *dispatch-only* relative cost. Total
+latency improvement depends on whether dispatch was on the critical
+path — measured per workload in P6.
+
+**R sensitivity** (ADR-0064 Test #9 — sanity at neighbouring R):
+
+| R (cycles/byte) | opt3 ns | opt2 ns | ratio |
+|---|---|---|---|
+| 0.25 (4 B/cy)     | 458    | 175    | 2.6× |
+| 0.0625 (16 B/cy) ✓| 414.5  | 103.75 | 4.0× |
+| 0.03125 (32 B/cy) | 407.25 | 91.875 | 4.4× |
+
+`opt2 < opt3` at all three; ratio is monotonic in `1/R` (more
+FIXED-dominated as R decreases).
+
+**As-built (measured, `test_gqa_decode_opt2.py`).** The per-tile emitters
+give opt3=960 ns, opt2=184 ns → **5.22×** at the default calibration (higher
+than the model's ≈4.0× because the opt3 emitter includes the full online
+merge); R-sweep ratios **4.20× / 5.22× / 5.45×** for R∈{0.25,0.0625,0.03125},
+monotonically increasing as R decreases. The gate is the calibration-robust
+`opt3 > 2 × opt2`.
+
+---
+
+## 12. Open items (status)
+
+Most ADR-0065 §Open review items map to specific implementation choices
+already documented in §§5–8. Live remaining items:
+
+| Item | Resolution venue |
+|---|---|
+| Non-identity `PrimaryOutSpec.transform` (e.g., `"trans"`) | Future recipe addition; current code uses string dispatch in step 3 of §6 |
+| `_gqa_decode_long.py` opt2 — TILE size sweep | Open (existing §B-3, §9 of ADR-0060) |
+| Larger composite descriptors approaching ADR-0064 D7 cap (1024 bytes) | Hard cap → `ValueError` at emit (no segmentation); kernel must split the recipe |
+| `scratch_scope` budget vs softmax_merge intermediates (G·TILE for P) | ADR-0063 sizing check at S_kv = 256K |
+| Scheduler plan-gen cost for large composites (ADR-0064 Open Item 2) | If P6 measurements show scheduler-bound behaviour near the cap, expose per-op-count `overhead_ns` on PE_SCHEDULER |
+
+---
+
+## 13. Risks
+
+| Risk | Likelihood | Mitigation |
+|---|---|---|
+| `CompositeCmd` refactor (P1) changes op_log for some bench | med | Parametrized regression over all benches; revert atomically |
+| `OpSpec.operands` migration (positional → dict) breaks epilogue site | med | Per-EPILOGUE_OPS migration table; existing epilogue tests cover |
+| `_RwHazardTracker` false-positive serializes too aggressively | low | Strict FIFO is the safe direction; RW-aware reorder is a future ADR |
+| ADR-0064 Rev2 and ADR-0065 P1 land in wrong order → stale goldens | med | PR ordering enforced; CI gate requires both to be green |
+| `softmax_merge` scratch budget exceeds `scratch_scope` at long context | med | ADR-0063 sizing check; if needed, tile_alignment promotion (future) |
+| PE_MATH op_kind label collision with new strings | low | Labels are opaque strings — no enum conflict; just add to allowed set if validated |
+
+---
+
+## 14. Glossary & references
+
+- **flat ops list** — `CompositeCmd.ops: tuple[OpSpec, ...]` with no
+  head/epilogue field; placement determined by op position + scope.
+- **prologue (user API)** — `tl.composite(prologue=[...])` kwarg for
+  recipe ops that lower into pre-GEMM KERNEL-scope OpSpecs. *HW command
+  does not carry this concept.*
+- **recipe** — a named macro op (e.g., `softmax_merge`) that TLContext
+  expands to multiple flat OpSpecs at emit time. Registry =
+  `RECIPE_DESCRIPTORS`.
+- **primary_out** — recipe's implicit output (e.g., `P` for
+  softmax_merge). Auto-bound to head GEMM's first matrix operand.
+- **handle-identity strict FIFO** — cross-composite serialization rule:
+  a composite waits for all prior composites whose `rw_handles`
+  intersect with its own to complete.
+- **logical_bytes** — HW-logical byte size of a PE command, derived
+  from a structural rule (ADR-0064 D2), not from Python's
+  `sys.getsizeof`.
+
+Source anchors: `tl_context.py` (tl API), `pe_scheduler.py:104-143`
+(composite tile plan), `pe_dma.py:45,89` (read channel), `pe_math.py`
+(op_kind opaque latency), `pe_cpu.py` (dispatch hook),
+`pe_commands.py` (CompositeCmd, OpSpec, EPILOGUE_OPS).
+
+ADRs: **0065** (this design), **0064 Revision 2** (dispatch cost),
+**0060** (GQA fused attention, §5.6 / §8 item 4 carve-out), **0063**
+(scratch_scope), **0046** (tl_context contract), **0042** (tile plan
+generators), **0014** (PE pipeline execution model).
+
+---
+
+## 15. Data-mode numerics (D8 + N1–N4) — as-built
+
+The original DDD assumed opt2's numeric parity "just works" via the existing
+DataExecutor. It does not: a composite in the builtin tiled path is recorded
+**only as latency TileTokens** (`op_kind="unknown"`, no addresses) — its
+matmul/MATH is never computed; the output address is never written. Making
+opt2's recipe compute required building the composite/recipe data-mode path.
+
+**D8 — output handle + output-space DMA (commit `e8d6c28`).** Prerequisite:
+a composite's output must be a real, addressable, chainable handle, and a
+TCM-resident output must NOT be DMA_WRITTEN (its high-bit scratch address
+crashes the DMA PA decoder). `tl.composite` returns the output `TensorHandle`
+(auto-TCM if `out` omitted); the tile loop's DMA_WRITE is gated on
+`out.space=="hbm"`. This is what lets opt2 *run* in data mode at all.
+
+**N1 — composite GEMM record (commit `4089e18`).** `op_log._extract_op_info
+(CompositeCmd)` scans `ops` for the GEMM (not `ops[0]` — a recipe's GEMM sits
+after the prologue) and emits one `composite_gemm` record (a/b/out
+addrs+spaces). `PE_SCHEDULER` records the composite's numeric op (no-op
+without an op_logger). `DataExecutor._execute_gemm` is best-effort: skips on
+(KeyError, ValueError) reading inputs (torch.empty / sharded operands).
+
+**N2 — recipe MATH ops (commit `3d4a3d4`).** `_compute_math` gains
+`rmax`/`rsum` (keepdims) + `max_elem`/`exp_diff`/`mul_bcast` (binary; numpy
+broadcasting covers `bcast_axis`). `tiling._math_stage` carries operand+output
+addrs/shapes/spaces+axis on the **recipe prologue** stages only;
+`op_log.record_end` promotes a MATH stage to `op_kind="math"` *only* when it
+carries `input_addrs` (legacy epilogue stages stay opaque → byte-equal). A
+latent P2 bug was fixed: reductions keep the reduced axis as size 1 (keepdims)
+so they broadcast against `m=(G,1)`. Verified: the 8 MATH ops produce m, l
+(full online-softmax update) and O (rescaled by corr) vs a numpy reference.
+
+**N3 — GEMM-epilogue accumulate (commit `5bbe293`).** The composite GEMM
+record is marked `accumulate=True` when an `add` epilogue targets the GEMM's
+own output; `_execute_gemm` then does `O = O_existing + a@b`. The composite's
+numeric op is recorded at **completion** (`_retire_on_done`), not dispatch, so
+its `t_start` sorts AFTER its prologue MATH ops → the GEMM reads the computed
+P and the rescaled O. Verified: full recipe → `O = O_old·corr + P@V` vs numpy.
+
+**N4 — opt2 parity + ordering (commit `e9d62b9`).** `composite()` awaits its
+prologue recipe operands (the cross-composite RAW: #2's `s` is #1's score
+tile) so #2 waits for #1. With N1–N3, opt2 computes a correct attention output
+e2e in data mode and matches opt3 within fp tolerance. **Scope:** the e2e
+opt2↔opt3 parity holds in the near-uniform-attention regime — kernbench's K
+reshape-as-transpose (opt3 docstring: correct for zero/symmetric inputs)
+makes opt2's per-sub-tile K interpretation differ from opt3's single-tile one
+for high-contrast inputs (a pre-existing kernbench limitation shared by both
+kernels, not the recipe). The **exact** recipe math is verified non-trivially
+by N2/N3 (`test_recipe_data_mode.py`).

docs/report/1H-codesign-paper/.gitignore

@@ -0,0 +1,2 @@
+# Downloaded Tectonic toolchain binary (large, machine-local; re-fetched by /paper)
+build/.tools/

docs/report/1H-codesign-paper/figures/gqa_latency.json

@@ -0,0 +1,97 @@
+{
+  "version": 1,
+  "panels": [
+    "single_user_prefill_gqa",
+    "multi_user_prefill_gqa",
+    "single_user_decode_gqa",
+    "multi_user_decode_gqa"
+  ],
+  "rows": [
+    {
+      "panel": "single_user_prefill_gqa",
+      "kind": "prefill",
+      "C": 1,
+      "S_kv": 16,
+      "latency_ns": 445.1180000000004,
+      "op_log_summary": {
+        "gemm_count": 2,
+        "ipcq_copy_count": 0,
+        "dma_read_count": 3,
+        "dma_write_count": 1
+      },
+      "engine_occupancy_ns": {
+        "pe_gemm": 2.048000000000002,
+        "pe_math": 5.0,
+        "pe_dma": 72.0,
+        "pe_fetch_store": 0,
+        "pe_ipcq": 0,
+        "pe_cpu": 0
+      }
+    },
+    {
+      "panel": "multi_user_prefill_gqa",
+      "kind": "prefill",
+      "C": 4,
+      "S_kv": 16,
+      "latency_ns": 4630.408000000019,
+      "op_log_summary": {
+        "gemm_count": 32,
+        "ipcq_copy_count": 24,
+        "dma_read_count": 12,
+        "dma_write_count": 4
+      },
+      "engine_occupancy_ns": {
+        "pe_gemm": 8.192000000000917,
+        "pe_math": 236.0,
+        "pe_dma": 1446.945000000003,
+        "pe_fetch_store": 0,
+        "pe_ipcq": 0,
+        "pe_cpu": 0
+      }
+    },
+    {
+      "panel": "single_user_decode_gqa",
+      "kind": "decode",
+      "C": 1,
+      "P": 8,
+      "S_kv": 64,
+      "latency_ns": 3631.730500000015,
+      "op_log_summary": {
+        "gemm_count": 16,
+        "ipcq_copy_count": 21,
+        "dma_read_count": 24,
+        "dma_write_count": 1
+      },
+      "engine_occupancy_ns": {
+        "pe_gemm": 16.383999999998196,
+        "pe_math": 160.0,
+        "pe_dma": 1450.1150000000016,
+        "pe_fetch_store": 0,
+        "pe_ipcq": 0,
+        "pe_cpu": 0
+      }
+    },
+    {
+      "panel": "multi_user_decode_gqa",
+      "kind": "decode",
+      "C": 4,
+      "P": 8,
+      "S_kv": 128,
+      "latency_ns": 6692.578999999866,
+      "op_log_summary": {
+        "gemm_count": 64,
+        "ipcq_copy_count": 93,
+        "dma_read_count": 96,
+        "dma_write_count": 1
+      },
+      "engine_occupancy_ns": {
+        "pe_gemm": 32.76799999998184,
+        "pe_math": 688.0000000000009,
+        "pe_dma": 15919.751500000013,
+        "pe_fetch_store": 0,
+        "pe_ipcq": 0,
+        "pe_cpu": 0
+      }
+    }
+  ]
+}

docs/report/1H-codesign-paper/main.tex

@@ -0,0 +1,46 @@
+\documentclass[10pt,twocolumn]{article}
+
+\usepackage[margin=0.75in]{geometry}
+\usepackage{graphicx}
+\usepackage{booktabs}
+\usepackage{amsmath}
+\usepackage{amssymb}
+\usepackage{siunitx}
+\DeclareSIUnit\flop{FLOP}
+\DeclareSIUnit\cycle{cycle}
+\DeclareSIUnit\byte{B}
+\usepackage{xcolor}
+\usepackage{hyperref}
+\hypersetup{colorlinks=true,linkcolor=blue!50!black,citecolor=blue!50!black,urlcolor=blue!50!black}
+\usepackage{caption}
+\captionsetup{font=small,labelfont=bf}
+\usepackage{microtype}
+\usepackage{tikz}
+\usetikzlibrary{arrows.meta,positioning,calc,fit}
+
+\graphicspath{{figures/}}
+
+\title{\textbf{Hardware--Software Co-Design for Grouped-Query Attention on AHBM}}
+
+\author{Mukesh Garg \and Jiyoon Lee \and Yangwook Kang}
+
+\date{
+\small
+AGI Computing Lab, System Technology Group\\
+2026 H1 Report
+}
+
+\begin{document}
+\maketitle
+
+\input{sections/00-exec-summary}
+\input{sections/01-introduction}
+\input{sections/02-platform}
+\input{sections/03-gemm}
+\input{sections/04-allreduce}
+\input{sections/05-gqa}
+\input{sections/06-discussion}
+\input{sections/07-conclusion}
+\input{sections/08-future-work}
+
+\end{document}

docs/report/1H-codesign-paper/sections/00-exec-summary.tex

@@ -0,0 +1,36 @@
+\section*{Executive Summary}
+\addcontentsline{toc}{section}{Executive Summary}
+
+Grouped-Query Attention (GQA) has largely replaced GPT-3-style
+multi-head attention in modern decoder-only LLMs (e.g., Llama 3 and
+Mistral) due to its reduced memory and bandwidth requirements. Mapping
+GQA efficiently onto AHBM introduces three architectural requirements:
+optimized placement of KV caches and weights to minimize inter-PE
+communication, low-overhead support for unavoidable cross-PE traffic,
+and efficient pipelining of memory accesses and computation within each
+PE.
+
+To address these requirements, this report introduces three
+hardware--software co-design mechanisms: GQA-aware placement
+of KV caches and weights across TCM, SRAM, and HBM; PE\_IPCQ,
+an efficient on-device collective communication primitive; and a
+composite-command GEMM pipeline that tightly pipelines memory
+and compute operations within each PE under PE\_SCHEDULER control.
+Kernbench-based evaluation show up to \SI{38}{\percent} lower collective
+communication overhead, \SIrange{25}{35}{\percent} lower all-reduce
+latency than a mesh-based interconnect, and up to \SI{78}{\percent} of
+peak MAC efficiency for compute-intensive GEMM kernels. These results
+demonstrate that the fused GQA kernel can efficiently exploit AHBM's HBM bandwidth.
+
+While GQA serves as the motivating workload in this study, the
+resulting architectural mechanisms are broadly applicable across the
+AHBM software stack. PE\_IPCQ provides a scalable communication
+substrate for collective operations and communication-intensive workloads, while
+composite-command execution enables efficient fusion of memory
+movement, GEMM kernels, normalization, and other element-wise
+operations. Together with the hierarchical data-placement framework,
+these capabilities form a reusable foundation for future AI kernels
+and communication libraries on AHBM. Looking ahead, these same mechanisms provide
+the basis for optimizing FFN-intensive workloads such as Mixture-of-Experts (MoE)
+layers and for developing integrated optimization strategies spanning both attention
+and MoE components within next-generation AI models.

docs/report/1H-codesign-paper/sections/01-introduction.tex

@@ -0,0 +1,83 @@
+\section{Introduction}
+\label{sec:intro}
+
+AHBM integrates compute units directly into the HBM stack. Each
+processing element (PE) is paired with a dedicated slice of HBM and uses
+local TCM and SRAM to stage data between memory and the MAC array. On
+this memory-centric architecture, kernel performance depends not only on
+compute throughput but also on how effectively data is placed, moved,
+and shared across the memory hierarchy and between PEs. Optimizing AI
+kernels on AHBM therefore requires hardware--software co-design, in
+which kernel algorithms and architectural mechanisms are developed
+together.
+
+To enable detailed performance analysis and rapid design exploration, we
+developed \textbf{KernBench}, a source-level discrete-event simulation
+platform for AHBM. KernBench implements the AHBM execution model,
+including memory-system latencies, the PE execution model, inter-PE
+communication, and host-side orchestration, while executing both kernel
+and host software directly from source code. This provides fine-grained
+visibility into execution behavior. It enables systematic evaluation of hardware--software co-design choices
+independently of higher-level software stacks such as compilers and
+runtimes. All
+results presented in this report are obtained using KernBench, which
+serves as the common evaluation platform for all mechanisms and kernels
+discussed in this study.
+
+This report focuses on Grouped-Query Attention (GQA), one of the most
+performance- and bandwidth-critical components of LLM inference. Modern
+decoder-only models such as Llama~3 and Mistral have largely transitioned
+from GPT-3-style multi-head attention (MHA) to GQA, in which multiple
+query heads share a single KV head to reduce KV cache capacity and
+memory-bandwidth requirements. While GQA improves system efficiency at
+the model level, mapping it efficiently onto AHBM introduces three
+architectural requirements: optimized placement of KV caches and weights
+to minimize inter-PE communication, low-overhead support for unavoidable
+cross-PE traffic, and efficient pipelining of memory accesses and
+computation within each PE.
+
+To address these requirements, this report introduces three
+hardware--software co-design mechanisms. First, GQA-aware data placement
+distributes KV caches and weights across the TCM/SRAM/HBM hierarchy to
+reduce communication overhead and improve data locality. Second,
+PE\_IPCQ provides an efficient on-device collective communication
+primitive for reductions and other communication-intensive operations.
+Third, a composite-command GEMM pipeline tightly pipelines memory
+movement and computation within each PE under PE\_SCHEDULER control,
+reducing command overhead while keeping the MAC array efficiently
+utilized.
+
+The compute enabler (the composite-command GEMM pipeline,
+\S\ref{sec:gemm}) and the communication enabler (PE\_IPCQ,
+\S\ref{sec:allreduce}) are first developed and evaluated independently.
+The fused GQA kernel (\S\ref{sec:gqa}) then combines them with
+GQA-aware data placement to demonstrate an end-to-end attention
+implementation on AHBM. The correspondence is direct: attention's
+$QK^{\top}$ and $PV$ products are precisely the GEMMs that benefit from
+the composite-command pipeline, while its KV reductions are precisely
+the collective operations that benefit from PE\_IPCQ. Together, these
+mechanisms enable the fused GQA kernel to efficiently exploit AHBM's HBM
+bandwidth.
+
+Although GQA serves as the motivating workload for this study, the
+resulting mechanisms are intended as reusable building blocks for a much
+broader class of AI kernels. PE\_IPCQ can support collective
+communication across distributed and communication-intensive workloads,
+while composite-command execution can be applied to GEMM-based kernels,
+feed-forward networks (FFNs), normalization, and other fused operator
+pipelines. Together with the hierarchical data-placement framework,
+these mechanisms form a foundation for future AI kernels and
+communication libraries on AHBM. In the second half of 2026, this
+foundation will be extended to FFN- and MoE-dominated workloads and to
+end-to-end optimization of complete LLM execution.
+
+The remainder of this report is organized as follows.
+Section~\ref{sec:platform} describes the KernBench platform and the AHBM
+configuration used throughout this study. Sections~\ref{sec:gemm},
+\ref{sec:allreduce}, and \ref{sec:gqa} present the composite-command
+GEMM pipeline, PE\_IPCQ collective communication, and the fused GQA
+kernel, respectively. Section~\ref{sec:discussion} discusses the broader
+architectural implications of these results. Finally,
+Sections~\ref{sec:conclusion} and \ref{sec:future} summarize the key
+findings and outline future work on FFN, MoE, and full-model
+optimization.

docs/report/1H-codesign-paper/sections/02-platform.tex

@@ -0,0 +1,319 @@
+\section{The KernBench Platform}
+\label{sec:platform}
+
+All results in this report are produced on \emph{KernBench}, a
+system-level discrete-event simulator for LLM kernels running on AHBM.
+This section explains why the platform exists, the device and
+execution model it presents to a kernel writer, how its latency model
+turns that execution into a number, and the concrete hardware
+configuration used for every experiment that follows.
+
+\subsection{Why KernBench: source-level kernels without a software stack}
+\label{sec:why}
+
+In a production end-to-end (E2E) stack, kernel performance is entangled
+with every layer above the hardware: the compiler's tiling and
+scheduling choices, the framework's operator dispatch, the collective
+library, and the runtime. Good E2E numbers require \emph{all} of those
+layers to be co-optimized, which makes it hard to answer a narrower but
+more fundamental question: \emph{given the hardware, how fast can a
+well-written kernel be, and which hardware features actually make it
+faster?}
+
+KernBench is built to answer exactly that question. Kernels are written
+and executed at the \emph{source level}---as algorithmic descriptions
+in a small tile-oriented kernel API---with no dependency on a compiler
+or any other software-stack layer. The simulator takes the kernel and a
+hardware topology and reports the latency the modeled hardware would
+deliver. This isolation is deliberate: it lets us study algorithm-level
+optimizations (how to tile a GEMM, how to schedule a collective, how to
+fuse an attention kernel) and the hardware features that support them,
+without the confound of compiler maturity or framework overhead. The
+cost is that KernBench numbers are \emph{not} E2E latencies; they are
+the achievable-kernel latencies an ideal software stack would expose.
+
+\subsection{Device and execution model}
+\label{sec:exec}
+
+KernBench models AHBM as a hierarchy of SIPs, CUBEs, and processing
+elements (PEs). At the system level, multiple SIPs are connected
+through inter-package links, while each SIP contains a collection of
+CUBEs joined by an on-package interconnect
+(Fig.~\ref{fig:sip-arch}).
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{sip_architecture.pdf}
+\caption{Modeled hardware graph at the \emph{SIP} level (one
+example configuration; specific parameters in \S\ref{sec:hw} /
+Table~\ref{tab:hw}). The SIP holds a $4{\times}4$ mesh of CUBEs and
+an IO chiplet, with each line a directed link labelled by its
+physical distance and bandwidth.}
+\label{fig:sip-arch}
+\end{figure}
+
+Each CUBE contains eight PEs, shared SRAM, HBM controllers, an
+\textsf{M\_CPU} control processor, and an intra-CUBE router mesh
+(Fig.~\ref{fig:cube-arch}). Together these components form the
+execution substrate for all kernels evaluated in this report.
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{cube_architecture.pdf}
+\caption{Modeled hardware graph at the \emph{CUBE} level---a
+zoom-in on one CUBE node from Fig.~\ref{fig:sip-arch}. The CUBE
+holds 8 PEs (each with its own HBM channels), the \textsf{M\_CPU}
+control processor, a shared SRAM, an intra-CUBE NoC router mesh, and
+UCIe links to neighbouring CUBEs.}
+\label{fig:cube-arch}
+\end{figure}
+
+Within a PE, commands are dispatched by \textsf{PE\_CPU} to
+\textsf{PE\_SCHED}, which routes work to specialized execution
+engines---DMA, FETCH/STORE, GEMM, vector-math, and IPCQ
+(Fig.~\ref{fig:pe-arch}). Commands come in two flavours. \emph{Atomic}
+commands target a single engine---a plain DMA read, a GEMM tile, a
+vector-math op, an IPCQ send/receive---and are the natural unit for
+short or special-purpose work; the \textsf{PE\_CPU} itself also runs
+control-plane work directly. \emph{Composite} commands, by contrast,
+carry an ordered pipeline of operations across multiple engines
+(\textsf{DMA\_READ} $\rightarrow$ \textsf{FETCH} $\rightarrow$
+\textsf{GEMM}/\textsf{MATH} $\rightarrow$ \textsf{STORE} $\rightarrow$
+\textsf{DMA\_WRITE}) that the \textsf{PE\_SCHED} tiles and streams
+without per-tile redispatch. The composite form is the substrate for
+the GEMM optimization (\S\ref{sec:gemm}) and, combined with on-PE
+collectives, for fused attention (\S\ref{sec:gqa}).
+
+\begin{figure*}[t]
+\centering
+\includegraphics[width=\linewidth]{pe_architecture.png}
+\caption{Modeled PE architecture used in this report---one example
+configuration whose specific parameters are listed in
+\S\ref{sec:hw} (Table~\ref{tab:hw}); KernBench is not tied to this
+particular decomposition and supports arbitrary PE block layouts
+provided each component is given a port and bandwidth model.
+\textsf{PE\_CPU} dispatches commands to the \textsf{PE\_SCHED}, which
+routes tile-token streams through the \textsf{PE\_DMA},
+\textsf{PE\_FETCH\_STORE}, and \textsf{GEMM}/\textsf{MATH} engines
+along on-chip links; the \textsf{PE\_IPCQ} provides the control plane
+for on-device collective communication.}
+\label{fig:pe-arch}
+\end{figure*}
+
+KernBench is layered along the flow of a request:
+
+\begin{itemize}
+  \item The \textbf{runtime API} is host-facing and
+  topology-agnostic---it deploys tensors and launches kernels but knows
+  nothing about routing or interconnect.
+  \item The \textbf{simulation engine} schedules discrete events and
+  routes every request through the modeled graph.
+  \item The \textbf{components} are device-side nodes that model
+  hardware behaviour: the per-PE blocks (scheduler, DMA, GEMM and
+  vector-math engines, TCM, IPCQ), the NoC routers, the HBM
+  controllers, and the inter-chiplet links.
+\end{itemize}
+
+Data and timing are handled in two passes, so that a kernel's numeric
+results and its latency are computed consistently but independently.
+
+\subsection{Latency model: graph traversal and contention}
+\label{sec:latency}
+
+The modeled hardware hierarchy described above is represented
+internally as a directed graph. Nodes correspond to hardware
+components---PEs, routers, memory controllers, SRAM blocks, IO
+chiplets---while edges represent communication links with associated
+bandwidth (\si{\giga\byte\per\second}) and propagation
+(\si{\nano\second}) attributes. Every operation in KernBench---DMA
+transfers, remote memory accesses, collective communication, command
+dispatch---is modelled as a traversal through this graph. End-to-end
+latency is decomposed into four contributions accumulated along the
+traversal path (Fig.~\ref{fig:latency-model}): \textbf{per-node
+overhead} at each component, \textbf{per-edge transmission} on each
+wire, \textbf{drain} (per-flit service occupancy) at the destination,
+and \textbf{queuing delay} at the shared FIFOs that the wire and the
+destination share between concurrent transactions.
+
+\begin{figure*}[t]
+\centering
+\includegraphics[width=\textwidth]{latency_model.png}
+\caption{Conceptual schematic of the latency model. Two source nodes
+(Requester A/B) inject flits through a chain of routers into a
+destination node; on the shared edge between routers, flits from the
+two transactions are interleaved flit-by-flit by the wire's FIFO
+arrival order. End-to-end latency is the sum of four contributions:
+\textbf{per-node overhead} (the switch's fixed processing cost,
+shown in light yellow---the same colour as the destination's
+processing-logic block), \textbf{per-edge transmission}
+($\textit{flit\_size}/\textit{BW}$ on each wire),
+\textbf{drain} (per-flit service occupancy at the destination's
+channel), and \textbf{queuing delay} (waiting in a FIFO when a
+shared resource is busy). The places where queuing actually
+accumulates are highlighted in green: the router output queue and
+the destination's input queue. This is the model, not a
+measurement; specific bandwidths and overheads are listed in
+\S\ref{sec:hw} (Table~\ref{tab:hw}).}
+\label{fig:latency-model}
+\end{figure*}
+
+\paragraph{The hardware as a graph.} The topology is compiled once at
+configuration time into this graph and is never mutated during a run.
+There are no hidden shortcuts, implicit bypasses, or magic paths: if a
+request reaches its destination, the path it took is explicit in the
+graph, and the latency it incurred is the sum of the per-node and
+per-edge costs paid along that path. The same graph representation
+applies recursively at every hierarchy level---system, SIP
+(Fig.~\ref{fig:sip-arch}), CUBE (Fig.~\ref{fig:cube-arch}), and PE
+(Fig.~\ref{fig:pe-arch}).
+
+\paragraph{From graph to discrete-event simulation.} The graph is
+driven by a discrete-event engine. Two kinds of events advance
+simulation time: \emph{node events} (component switching overhead,
+service completions such as an HBM channel commit or a GEMM tile
+finish) and \emph{edge events} (the flit-by-flit serialization of a
+payload across a bandwidth-limited link). The engine maintains a
+priority queue of pending events ordered by their scheduled time and
+fires them one at a time, with ties broken under a deterministic
+policy so that the same kernel on the same topology always yields the
+same trace. Per-request correlation IDs are stamped at injection and
+carried through every hop, so the path from injection to completion
+is fully traceable. Every nanosecond in a reported latency
+corresponds to exactly one of these events on exactly one node or
+edge---there is no slack in the budget.
+
+\paragraph{Latency contributions.} Three kinds of latency accumulate
+along a traversal: (i) \emph{per-node fixed overhead}---each component
+carries a small switching cost (router decode, controller pickup,
+scheduler handoff); (ii) \emph{per-edge transfer time}---each link's
+payload is decomposed into fixed-size flits (default
+\SI{256}{\byte}), and each flit arrives at
+$\text{prop}+\text{flit\_bytes}/\text{bw}$ after the previous one,
+giving wormhole semantics across multi-hop paths; and (iii)
+\emph{per-service occupancy}---memory controllers, GEMM stages, and
+collective engines hold the request for their service time before
+releasing it downstream. Each of these is attached to a specific node
+or edge in the graph; together they make up the entire latency budget.
+
+\paragraph{Congestion: where bottlenecks emerge.} The simulator's
+sharpness comes from how it models contention for those nodes and
+edges. \emph{Every directed edge has a FIFO}: an arriving flit takes
+its bandwidth-limited transfer time on top of whatever earlier flits
+are still being served, so a busy link queues later traffic behind
+earlier traffic rather than transferring everything at peak BW.
+\emph{HBM is modelled with per-pseudo-channel parallelism}: a
+stateless array of channel-availability timestamps with address-based
+channel selection captures the bank-level concurrency that real HBM
+exposes, so 64 channels per CUBE deliver real parallelism on uniform
+addresses but a hot channel surfaces as the bottleneck on skewed ones.
+\emph{Every component has a serial worker}: a router carrying two
+heavy streams interleaves them at flit granularity in arrival order
+rather than fanning out for free, so two concurrent collectives sharing
+a link share its bandwidth, not double it. Without these mechanisms
+the simulator would simply re-confirm the peak-BW roofline; with them,
+it reveals where the real bottlenecks form and which hardware levers
+actually relieve them---which is exactly the question the codesign
+work in this report turns on.
+
+\paragraph{Control-plane (issue) cost model.} The cost of \emph{issuing}
+a command is modelled structurally rather than with a per-operation
+calibration table. The PE control processor charges, per command,
+\[
+  d_{\text{cmd}} = \textsf{FIXED} + b_{\text{logical}} \cdot R,
+\]
+where $b_{\text{logical}}$ is the command's hardware-logical byte size,
+\textsf{FIXED} captures the fixed per-command cost (queue-tail update,
+completion registration) and $R$ captures the per-byte cost of
+serializing the command descriptor into the scheduler queue. The
+default anchoring (\textsf{FIXED} $= 40$ cycles, $R = 0.0625$
+cycles/byte, i.e.\ \SI{16}{\byte\per\cycle}, at \SI{1}{\giga\hertz})
+places a typical composite at roughly \SI{43}{\nano\second}, and a
+hard cap on a composite's descriptor size prevents the model from
+rewarding arbitrarily large fused commands beyond what real descriptor
+queues accept. In the configurations measured here, command issue is
+not the bottleneck---data movement is---so this term stays small
+relative to DMA and collective time.
+
+\paragraph{Accuracy.} The model is precise about the effects that
+dominate kernel latency on this class of hardware: per-edge bandwidth
+occupancy and flit-level serialization, HBM pseudo-channel parallelism,
+and per-component switching overhead. Two independent cross-checks
+drawn from the experiments in this report confirm that this precision
+translates into physically reasonable kernel latencies. First, in the
+GEMM study (\S\ref{sec:gemm}), simulator-measured MAC efficiency
+tracks an analytic ideal-pipeline model within roughly
+\SIrange{10}{20}{\percent} across a wide range of tile counts; the
+residual gap is attributable to pipeline-fill and DMA effects the
+analytic model omits. Second, in the all-reduce study
+(\S\ref{sec:allreduce}, Fig.~\ref{fig:allreduce-cmp}), simulator
+latency for a 2D-torus over six devices follows the expected
+startup-plus-per-packet shape across the entire payload sweep---tight
+at small payloads where startup dominates, and within a single-digit
+multiplicative factor at the largest payloads, where the residual gap
+is explained by per-router switching the analytic shape elides. A
+single-device point from an external full-system simulator (FSIM) at
+the largest payload sits an order of magnitude above the KernBench
+multi-device torus, illustrating the well-known gap between an
+achievable-kernel number and a full end-to-end-stack number rather
+than a model error. The known simplifications---FIFO router
+arbitration (instead of round-robin), HBM scheduler without
+write-buffer reordering, no bank conflict, no refresh or thermal
+effects, and no upstream backpressure---are the price of a
+deterministic, inspectable model. These simplifications bound the
+absolute accuracy, but the agreement with both analytic models and the
+external full-system simulator data above indicates that KernBench is
+sufficiently accurate for evaluating the \emph{relative}
+hardware--software design trade-offs (tiling A vs.\ B, topology X
+vs.\ Y, with vs.\ without composite command, mesh vs.\ torus) that
+are the primary objective of this work.
+
+\subsection{Modeled hardware configuration}
+\label{sec:hw}
+
+Table~\ref{tab:hw} summarizes the hardware configuration used for
+every experiment in this report. It is read directly from the
+simulator's topology description; per-experiment workload parameters
+(matrix shapes, collective sizes, sequence lengths) are stated in
+their respective sections rather than here.
+
+\begin{table}[t]
+\centering
+\caption{Modeled hardware configuration (shared by all experiments).}
+\label{tab:hw}
+\small
+\begin{tabular}{@{}ll@{}}
+\toprule
+\textbf{Parameter} & \textbf{Value} \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Hierarchy}} \\
+SIPs & 2 (1D ring) \\
+CUBEs per SIP & 16 ($4\times4$ mesh) \\
+PEs per CUBE & 8 (4 corners $\times$ 2) \\
+PEs total & 256 \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Processing element (PE)}} \\
+GEMM engine peak & \SI{8}{\tera\flop\per\second} (f16) \\
+TCM (on-PE) & \SI{16}{\mega\byte}, \SI{512}{\giga\byte\per\second} R/W \\
+\quad kernel scratch & \SI{1}{\mega\byte} \\
+DMA engines & 1 read + 1 write \\
+CPU / scheduler overhead & \SI{2}{\nano\second} / \SI{1}{\nano\second} \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Memory (per CUBE)}} \\
+HBM capacity & \SI{48}{\giga\byte} (8 slices) \\
+HBM aggregate BW & \SI{1024}{\giga\byte\per\second} \\
+HBM pseudo-channels & 64 (8 per PE), \SI{32}{\giga\byte\per\second} each \\
+SRAM (shared) & \SI{32}{\mega\byte}, \SI{128}{\giga\byte\per\second} link \\
+HBM burst & \SI{256}{\byte} \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Interconnect}} \\
+Intra-CUBE NoC link & \SI{256}{\giga\byte\per\second}, \SI{2}{\nano\second}/router \\
+Inter-CUBE (UCIe PHY) & \SI{512}{\giga\byte\per\second}, \SI{8}{\nano\second}, XY routing \\
+Inter-SIP (PCIe) & \SI{768}{\giga\byte\per\second} per endpoint \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Command-issue cost model (defaults)}} \\
+FIXED per command & 40 cycles \\
+per-byte rate $R$ & 0.0625 cycles/byte (\SI{16}{\byte\per\cycle}) \\
+composite size cap & \SI{1024}{\byte} \\
+\bottomrule
+\end{tabular}
+\end{table}

docs/report/1H-codesign-paper/sections/03-gemm.tex

@@ -0,0 +1,110 @@
+\section{GEMM Acceleration via the Composite Command}
+\label{sec:gemm}
+
+\subsection{Why it is needed}
+
+GEMM is the compute core of every transformer block---the QKV
+projections, the attention score and context products, and the
+feed-forward matrices are all matrix multiplications. On a tiled
+accelerator a single logical GEMM expands into many hardware tiles, and
+each tile must be read from HBM, fetched into the register file, computed,
+stored, and written back. If every one of those tile-stage steps were an
+independently issued command, two costs would dominate. First, the host
+and the PE control processor would pay a per-command issue overhead
+$O(\text{tiles}\times\text{stages})$ times, which for a deep-$K$ reduction
+is hundreds to thousands of issues. Second, with stages dispatched
+one-at-a-time the scheduler cannot overlap the DMA of the next tile with
+the compute of the current one---the pipeline never fills, and the MAC
+array sits idle waiting for data. The hardware question is therefore: what
+issue mechanism lets a single GEMM saturate the MAC array without drowning
+in command overhead?
+
+\subsection{Design}
+
+The answer is the \emph{composite command}. A single command carries the
+ordered tile pipeline
+\[
+\textsf{DMA\_READ}\rightarrow\textsf{FETCH}\rightarrow\textsf{GEMM}
+\rightarrow\textsf{STORE}\rightarrow\textsf{DMA\_WRITE},
+\]
+and the PE scheduler splits the payload into hardware tiles
+(here $32\times64\times32$), emitting one tile token per tile. Subsequent
+stages are reached by \emph{token self-routing} between the on-PE engines,
+so a tile flows DMA\,$\rightarrow$\,fetch\,$\rightarrow$\,GEMM\,$%
+\rightarrow$\,store without returning to the scheduler between stages.
+Because the whole pipeline is described by one command, the issue cost is
+paid once per GEMM rather than once per tile-stage, and the scheduler is
+free to keep every stage busy on different tiles simultaneously---tile
+$i$'s GEMM overlaps tile $i{+}1$'s DMA read. A multi-operation composite
+additionally lets an epilogue (for example a vector-math step) ride the
+same tile loop, firing per $K$-tile, per output tile, or once per kernel
+according to its declared scope; this is what later allows an attention
+kernel to fuse its softmax work into the GEMM pipeline
+(\S\ref{sec:gqa}).
+
+\subsection{Results}
+
+We sweep eight GEMM shapes spanning square, tall, wide, and deep-$K$
+geometries, under three operand-staging variants
+(\textsf{ref\_ref}, both operands streamed from HBM; \textsf{load\_ref},
+one operand resident in TCM; \textsf{load\_load}, both resident).
+Figure~\ref{fig:gemm-util} reports MAC utilization and efficiency, and
+Figure~\ref{fig:gemm-stages} breaks the kernel into per-stage engine
+busy time.
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gemm_mac_utilization_theoretical_vs_measured.png}
+\caption{GEMM MAC utilization and efficiency, theoretical vs.\ measured.
+Tile-fill sets the ceiling: under-tile shapes (marked $\ast$) such as
+$M{=}K{=}N{=}32$, $M{=}8$, and $K{=}8$ cannot fill the MAC tile and cap at
+\SI{50}{\percent}, \SI{25}{\percent}, \SI{12.5}{\percent}. For
+tile-filling shapes, efficiency climbs with tile count---from
+\textasciitilde\SI{23}{\percent} at one tile to \textasciitilde%
+\SI{78}{\percent} measured at 48 tiles ($K{=}3072$)---and the measured
+bars track the analytic prediction within
+\SIrange{10}{20}{\percent}.}
+\label{fig:gemm-util}
+\end{figure}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gemm_stage_breakdown.png}
+\caption{Per-stage engine wall-clock (DMA in, Fetch, GEMM, DMA out) under
+\textsf{load\_ref} staging. For the deep-$K$ shape ($K{=}3072$, 48 tiles)
+the Fetch and GEMM stages are large and comparable
+(\textasciitilde\SI{770}{} and \SI{785}{\nano\second}) while DMA-out is
+negligible---a compute-rich, well-pipelined regime. For the low-reuse
+shape ($M{=}128,K{=}8,N{=}128$) DMA-out grows to
+\textasciitilde\SI{350}{\nano\second} and compute is small---a
+data-movement-bound regime.}
+\label{fig:gemm-stages}
+\end{figure}
+
+Two regularities stand out. (i) Utilization is governed by how completely
+the problem fills the MAC tile: the three under-tile shapes are hard-capped
+well below \SI{100}{\percent}, independent of how the kernel is issued.
+(ii) Among tile-filling shapes, efficiency is governed by tile count---more
+tiles amortize the one-time pipeline fill and issue cost, so the deep-$K$
+shape reaches \textasciitilde\SI{78}{\percent} of peak while a single-tile
+shape reaches only \textasciitilde\SI{23}{\percent}. The stage breakdown
+explains why: with 48 tiles the GEMM and Fetch stages overlap and stay
+busy, whereas the low-reuse shape spends most of its time moving data in
+and out.
+
+\subsection{Analysis and meaning}
+
+The composite command does not manufacture bandwidth or MAC throughput; it
+removes the two software-shaped obstacles between a GEMM and its hardware
+roofline. By paying issue cost once per GEMM and chaining tile-stages
+through token self-routing, it lets the scheduler keep the pipeline full,
+so compute-rich shapes actually reach the efficiency their arithmetic
+intensity allows ($\sim$\SI{78}{\percent} measured at 48 tiles), and
+data-bound shapes actually reach their DMA bound instead of stalling on
+command overhead. The close agreement between measured and theoretical
+efficiency (Figure~\ref{fig:gemm-util}) is also the report's primary
+validation that KernBench's latency model is faithful in the regime that
+matters. The hardware implication is concrete: a single-command,
+self-routing tile pipeline is the issue mechanism that makes the MAC array
+usable, and it is a prerequisite---not a luxury---for the fused attention
+kernel of \S\ref{sec:gqa}.

docs/report/1H-codesign-paper/sections/04-allreduce.tex

@@ -0,0 +1,113 @@
+\section{PE\_IPCQ and Collective Communication}
+\label{sec:allreduce}
+
+\subsection{Why it is needed}
+
+Distributing a transformer across devices turns every tensor-parallel
+layer into a collective: partial results computed on different PEs, CUBEs,
+and SIPs must be summed and redistributed with an all-reduce. If that
+collective is handled by the host or by a generic DMA path, three problems
+appear. The reduction traffic competes with the kernel's own compute DMA
+on the same links, causing head-of-line blocking; there is no efficient
+peer-to-peer ring primitive, so data takes extra hops; and the ordering is
+hard to make deterministic. The hardware question is how to perform
+collectives \emph{on the device}, overlapped with compute and reproducible
+run-to-run.
+
+\subsection{Design}
+
+KernBench models a dedicated per-PE collective engine, \textbf{PE\_IPCQ}
+(inter-PE communication queue). It is a control-plane block: it owns the
+ring-buffer address arithmetic, head/tail pointers, peer-pointer caches,
+backpressure, and the four-direction (N/S/E/W) neighbor map, with eight
+ring buffers per PE (four directions $\times$ \{tx, rx\}). Crucially,
+PE\_IPCQ does \emph{not} move data itself---it delegates the actual
+transfer to PE\_DMA, keeping a clean control/data split. To stop
+collective traffic from blocking compute, PE\_DMA is extended into a
+two-channel virtual-channel model: \texttt{vc\_compute} carries tile
+load/store for GEMM and vector math, \texttt{vc\_comm} carries IPCQ sends,
+each with an independent state machine. The same physical link is shared
+but progresses in chunks (\SI{256}{\byte}), so a large GEMM DMA does not
+lock the link end-to-end against a pending reduction. On top of this
+substrate the collective runs a hierarchical local-reduce / global
+all-reduce-broadcast schedule across whatever inter-device topology the
+configuration specifies.
+
+\subsection{Results}
+
+Following the milestone-evaluation convention, the collective sweep builds
+its own six-device (six-SIP, $2\times3$) configurations---distinct from
+the two-SIP default of Table~\ref{tab:hw}---and measures all-reduce
+latency as a function of payload size for three inter-device topologies:
+a 1D ring, a 2D mesh (no wrap), and a 2D torus. Table~\ref{tab:allreduce}
+and Figure~\ref{fig:allreduce-cmp} report the result.
+
+\begin{table}[t]
+\centering
+\caption{All-reduce latency (ns) across six devices, by topology and
+per-PE payload. Lower is better; the torus wins at every size.}
+\label{tab:allreduce}
+\small
+\begin{tabular}{@{}rrrr@{}}
+\toprule
+\textbf{Bytes/PE} & \textbf{2D mesh} & \textbf{Ring 1D} & \textbf{2D torus} \\
+\midrule
+256     & 2667  & 2365  & 1701 \\
+4{,}096  & 4450  & 4082  & 3038 \\
+16{,}384 & 8900  & 8217  & 6403 \\
+65{,}536 & 26705 & 24766 & 19865 \\
+98{,}304 & 38574 & 35798 & 28840 \\
+\bottomrule
+\end{tabular}
+\end{table}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{allreduce_comparison.png}
+\caption{All-reduce latency vs.\ per-PE payload for the three topologies,
+against the analytic torus model and an external full-system simulator
+(FSIM) reference. The measured torus tracks the analytic curve within a
+small constant factor; the FSIM single-device point
+(\SI{366}{\micro\second}) sits an order of magnitude above the
+KernBench algorithmic latency, illustrating the difference between an
+achievable-kernel number and a full end-to-end-stack number.}
+\label{fig:allreduce-cmp}
+\end{figure}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{allreduce_buffer_kind.png}
+\caption{Effect of IPCQ staging-buffer placement (2D torus). At
+\SI{64}{\kibi\byte}/PE, TCM staging (\SI{19865}{\nano\second}) beats HBM
+(\SI{23081}{\nano\second}) by \textasciitilde\SI{14}{\percent} and SRAM
+(\SI{32201}{\nano\second}) by \textasciitilde\SI{38}{\percent}; at small
+payloads the three are indistinguishable.}
+\label{fig:allreduce-buf}
+\end{figure}
+
+Three findings follow. First, topology matters and the effect grows with
+size: at \SI{96}{\kilo\byte}/PE the 2D torus is \SI{25}{\percent} faster
+than the mesh and \SI{19}{\percent} faster than the ring, because its
+wrap-around links shorten the worst-case reduction path. Second, the
+measured curves grow smoothly with payload and stay within a small
+constant factor of the analytic torus model, with the gap reflecting real
+link serialization that the analytic model idealizes away. Third, where
+the IPCQ staging buffer lives is a first-order knob: keeping it in on-PE
+TCM is materially faster than HBM or shared SRAM at large payloads
+(Figure~\ref{fig:allreduce-buf}).
+
+\subsection{Analysis and meaning}
+
+PE\_IPCQ turns the collective from an off-device, contention-prone
+operation into an on-device primitive whose latency tracks the
+interconnect's physical limits. The control/data split keeps the engine
+small while reusing PE\_DMA for movement, and the virtual-channel split is
+what lets a reduction proceed without stalling the compute DMA that feeds
+the very GEMMs producing the partials---the same property the fused
+attention kernel relies on when it interleaves KV reduction with score
+computation (\S\ref{sec:gqa}). For the hardware roadmap the results argue
+two things: provisioning wrap-around (torus) inter-device links is worth
+roughly a \SI{20}{}--\SI{25}{\percent} collective-latency reduction at
+scale, and giving the IPCQ a fast on-PE staging buffer (TCM) rather than
+forcing it through HBM or SRAM is worth another \SI{14}{}--%
+\SI{38}{\percent} at large payloads.

docs/report/1H-codesign-paper/sections/05-gqa.tex

@@ -0,0 +1,123 @@
+\section{Fused Grouped-Query Attention}
+\label{sec:gqa}
+
+\subsection{Why it is needed}
+
+Attention is the 1H focus, and it is where the two preceding optimizations
+have to come together. Grouped-Query Attention (GQA) shrinks the KV cache
+by sharing each KV head across a group of query heads (here $h_q=8$ query
+heads to $h_{kv}=1$ KV head, a group factor $G=8$), which makes decoding
+feasible at long context but also makes it acutely memory-bound: a decode
+step processes a single query position ($T_q=1$) against the entire KV
+history, so its arithmetic intensity is low and its time is dominated by
+streaming the KV cache out of HBM. FlashAttention-style tiling with an
+online-softmax merge avoids ever materializing the full score matrix, but
+realizing it as a fast \emph{fused} kernel needs both building blocks from
+this report: efficient GEMM issue (\S\ref{sec:gemm}) for the
+$Q\!\cdot\!K^{\top}$ and $P\!\cdot\!V$ products, and an efficient on-device
+reduction (\S\ref{sec:allreduce}) for the multi-user and
+sequence-parallel KV reductions. This section is the capstone: the fused
+kernel that uses the composite command and PE\_IPCQ at the same time.
+Multi-head attention (MHA) was studied in prior work and serves here as
+the established baseline rather than being re-derived.
+
+\subsection{Design}
+
+The fused GQA kernel issues its matrix products as scheduler-managed
+composite commands and keeps the online-softmax merge and the cross-device
+KV reduction inside the kernel, on PE\_IPCQ. Two kernel families cover the
+two phases. The \emph{prefill} kernel is head-parallel and rotates the KV
+shards around an inter-CUBE ring (``Ring KV''). The \emph{decode} kernel
+is head-replicated with a statically sharded KV cache and reduces partial
+attention outputs through an M-fold intra-CUBE chain and, for multiple
+users, a two-level reduce-to-root. Two further primitives make long
+context practical: a \emph{lazy load} that issues the KV \textsf{DMA\_READ}
+and returns immediately, auto-waiting only at first use so that KV load
+overlaps score computation; and per-tile \emph{scratch recycling} that
+keeps the running softmax accumulators ($m,\ell,O$) in a persistent arena
+while freeing per-tile temporaries, so the kernel fits the
+\SI{1}{\mebi\byte} scratch budget across many tiles. A further refinement
+that restructures the decode step into two stateful composites (a named
+\textsf{softmax\_merge} recipe) is designed but not yet wired into the
+measured path; results below reflect the implemented kernel only.
+
+\subsection{Results}
+
+We measure four headline panels that vary the user count $C$ and the phase:
+single- and multi-user prefill ($T_q=4$, $S_{kv}=16$), and single- and
+multi-user decode ($P=8$ PEs, $S_{kv}=64$ and $128$), all at $d_{\text{head}}=64$
+and $G=8$. For each panel we harvest end-to-end latency
+(max event end minus min event start, the same window convention as the
+GEMM study) together with the per-engine busy time and the operation mix.
+Figure~\ref{fig:gqa-lat} and Figure~\ref{fig:gqa-break} report the result;
+the underlying numbers are in Table~\ref{tab:gqa}.
+
+\begin{table}[t]
+\centering
+\caption{Fused GQA per-panel latency and operation mix. Compute (GEMM,
+MATH) is a tiny fraction of DMA occupancy; IPCQ copies grow with users and
+PEs.}
+\label{tab:gqa}
+\small
+\begin{tabular}{@{}lrrrr@{}}
+\toprule
+\textbf{Panel} & \textbf{Lat.\ (ns)} & \textbf{GEMM} & \textbf{IPCQ} & \textbf{DMA rd} \\
+\midrule
+prefill C=1        & 445   & 2  & 0  & 3  \\
+prefill C=4 (Ring) & 4630  & 32 & 24 & 12 \\
+decode C=1, P=8    & 3632  & 16 & 21 & 24 \\
+decode C=4, P=8    & 6693  & 64 & 93 & 96 \\
+\bottomrule
+\end{tabular}
+\end{table}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_latency_by_panel.png}
+\caption{Fused GQA end-to-end latency. Latency grows from
+\SI{445}{\nano\second} (single-user prefill) to \SI{6693}{\nano\second}
+(four-user decode) as the KV history and the number of participating
+devices grow.}
+\label{fig:gqa-lat}
+\end{figure}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_op_engine_breakdown.png}
+\caption{Where the work goes. Left: operation counts---GEMM and IPCQ-copy
+volume both scale with users and PEs. Right: summed engine occupancy on a
+log scale---the DMA engine dominates by two to three orders of magnitude
+over the GEMM and MATH engines in every panel.}
+\label{fig:gqa-break}
+\end{figure}
+
+The dominant observation is in Figure~\ref{fig:gqa-break}: the compute
+engines are almost idle. The GEMM engine accumulates only
+\SIrange{2}{33}{\nano\second} of busy time across the panels and the
+vector-math engine \SIrange{5}{688}{\nano\second}, while the DMA engine
+accumulates \SIrange{72}{15920}{\nano\second}. Fused GQA, as modeled here,
+is overwhelmingly data-movement bound. The operation mix shows why the
+collective machinery matters: IPCQ-copy count rises from zero (single-user
+prefill) to 93 (four-user decode) as the kernel reduces partial outputs
+across more PEs and CUBEs, and DMA-read count rises in step as more KV
+shards are streamed. The PE control-processor dispatch cost registered as
+zero in this configuration---command issue is simply not on the critical
+path when data movement is this dominant.
+
+\subsection{Analysis and meaning}
+
+These panels are the clearest statement of the codesign thesis in the
+report. Because the composite command keeps GEMM issue cheap and the MAC
+array barely occupied, the fused attention kernel's latency is set almost
+entirely by data movement: streaming the KV cache and reducing partials
+across devices. That is precisely the cost that the communication-side
+work targets---PE\_IPCQ for the on-device reduction, the lazy load for
+load/compute overlap, fast TCM staging and torus links for the reduction
+itself. In other words, the two enablers are not independent features that
+happen to appear in the same kernel; the GEMM optimization is what
+\emph{exposes} the data-movement bottleneck (by removing the compute and
+issue overhead that would otherwise hide it), and the communication
+optimization is what \emph{attacks} it. For an attention-dominated decoder
+the meaningful hardware investments are therefore the ones that move data
+faster and reduce it on-device---not additional MAC throughput, which this
+workload cannot use.

docs/report/1H-codesign-paper/sections/06-discussion.tex

@@ -0,0 +1,51 @@
+\section{Discussion: which hardware changes are meaningful}
+\label{sec:discussion}
+
+Read together, the three studies point to a consistent ranking of where
+hardware investment pays off for attention-centric decoding.
+
+\paragraph{The composite command is foundational, but indirectly.} Its
+direct effect---driving compute-rich GEMMs to
+\textasciitilde\SI{78}{\percent} of peak---matters most for the
+compute-bound parts of a model (the large feed-forward and projection
+matrices). For attention itself, its more important effect is
+\emph{diagnostic}: by making issue and compute nearly free, it removes the
+overhead that would otherwise mask the true bottleneck, and the fused GQA
+results then show unambiguously that the kernel is data-movement bound.
+Without a cheap, self-routing issue mechanism we would not be able to tell
+whether attention is slow because of compute or because of data movement;
+with it, the answer is clear.
+
+\paragraph{The communication path is where attention latency actually
+lives.} Every all-reduce and fused-GQA measurement says the same thing:
+the limiting resource is moving and reducing data, not multiplying it. That
+makes the PE\_IPCQ design and the choices around it the highest-value
+hardware levers for this workload:
+\begin{itemize}
+  \item \textbf{On-device collectives with a compute/communication
+  virtual-channel split.} Performing the reduction on the device, with
+  \texttt{vc\_comm} separated from \texttt{vc\_compute} so the reduction
+  does not stall the compute DMA, is what lets attention overlap KV
+  movement with score computation at all.
+  \item \textbf{Fast on-PE staging memory.} Keeping the IPCQ buffer in TCM
+  rather than HBM or SRAM is worth \SI{14}{}--\SI{38}{\percent} of
+  collective latency at large payloads---a pure placement decision with a
+  first-order effect.
+  \item \textbf{Wrap-around (torus) inter-device links.} A torus fabric
+  buys \SI{20}{}--\SI{25}{\percent} over a mesh or ring at scale by
+  shortening the worst-case reduction path.
+\end{itemize}
+
+\paragraph{Raw MAC throughput is not the constraint for attention.} The
+GQA panels leave the GEMM and vector-math engines two to three orders of
+magnitude below the DMA engine in busy time. Adding MAC area would not move
+decode latency; the workload cannot use it. This is the single most
+actionable finding for an attention-dominated roadmap.
+
+\paragraph{Caveats.} These conclusions are achievable-kernel results from a
+deterministic model, not E2E measurements; absolute numbers carry the
+model's idealizations (\S\ref{sec:latency}), though the relative rankings
+that drive the recommendations are robust to them. The headline GQA panels
+are also at modest scale (up to four users, single SIP), and one designed
+refinement---the two-composite \textsf{softmax\_merge} decode---is not yet
+in the measured path. Larger-scale and multi-SIP headline runs are 2H work.

docs/report/1H-codesign-paper/sections/07-conclusion.tex

@@ -0,0 +1,27 @@
+\section{Conclusion}
+\label{sec:conclusion}
+
+This 1H work set out to make attention-centric LLM kernels fast through
+hardware--software codesign, and to do so on a platform that isolates
+algorithm-level behavior from the rest of the software stack. The result is
+a coherent picture rather than three separate optimizations. A composite
+command that issues a tiled GEMM as one self-routing pipeline makes the MAC
+array usable---reaching \textasciitilde\SI{78}{\percent} of peak on
+compute-rich shapes with measured efficiency tracking theory---and, just as
+importantly, makes compute cheap enough that the real bottleneck becomes
+visible. A per-PE on-device collective engine, PE\_IPCQ, turns all-reduce
+into a primitive whose latency follows the interconnect's physical limits,
+with topology and staging-memory choices each worth tens of percent. Fused
+Grouped-Query Attention then combines the two and shows the payoff and the
+lesson at once: the kernel is data-movement bound, so the optimizations
+that move and reduce data---not those that add arithmetic---are what
+determine its speed.
+
+The practical conclusion for the hardware roadmap is therefore specific.
+The changes worth keeping are the single-command self-routing GEMM
+pipeline, the on-device PE\_IPCQ collective with its compute/communication
+virtual-channel split, fast on-PE staging memory, and wrap-around
+inter-device links. Additional MAC throughput is not, for this workload, a
+meaningful investment. KernBench made these conclusions measurable by
+holding everything except the algorithm and the hardware fixed; the next
+half extends the same method beyond attention to the rest of the decoder.

docs/report/1H-codesign-paper/sections/08-future-work.tex

@@ -0,0 +1,40 @@
+\section{Future Work --- 2H}
+\label{sec:future}
+
+The 1H work covered attention end to end. The natural next step is to
+complete the decoder and then to ask how compute and data should be
+distributed for realistic, agentic workloads.
+
+\paragraph{From attention to the full decoder: FFN and MoE.} A decoder
+block is attention followed by a feed-forward network (FFN), and in modern
+models that FFN is increasingly a mixture-of-experts (MoE) layer. The FFN
+is the compute-rich counterpart to attention---it is where the composite
+command's MAC-efficiency gains (\S\ref{sec:gemm}) should matter most---so
+adding it gives a balanced view of a full block instead of its
+memory-bound half alone. MoE adds a new dimension: a routing step selects a
+few experts per token, turning the dense FFN GEMM into a sparse, data-
+dependent dispatch. The open questions are how to issue expert GEMMs as
+composites under data-dependent token counts, and how to move tokens to
+experts efficiently---an all-to-all-shaped communication pattern distinct
+from the all-reduce studied here, and a natural extension of the PE\_IPCQ
+work.
+
+\paragraph{Compute and data distribution for full LLM decoding.} With both
+attention and FFN/MoE in hand, the question becomes where each layer's
+compute and state should live. Attention is KV-bound and favors keeping the
+KV cache close to the PEs that consume it; FFN/MoE is compute-bound and
+favors spreading GEMM work across PEs; MoE routing makes the optimal
+placement token- and time-dependent. A 2H goal is to use KernBench to
+explore these placement and parallelization trade-offs---tensor vs.\ expert
+vs.\ sequence parallelism---under a single, software-stack-independent
+model, so the interconnect and memory implications of each choice are
+measured rather than assumed.
+
+\paragraph{Agentic workloads.} Agentic inference interleaves many
+short, bursty decode requests with tool use and long shared contexts,
+which stresses the system differently from a single long generation:
+context reuse across requests, dynamic batching, and uneven expert load all
+change how compute and data should be dispersed. Characterizing how total
+compute and data movement distribute across the SIP/CUBE/PE hierarchy under
+such workloads---and which of the 1H hardware levers still dominate when the
+workload is this irregular---is the broader 2H agenda.

docs/report/1H-codesign-paper/toc.md

@@ -0,0 +1,73 @@
+# KernBench — 2026 1H HW-SW Codesign Report — Table of Contents
+
+> Agreed outline (the `/paper` contract). `/paper build` follows this.
+> Period: 2026 1H. Language: English. References: external literature only
+> (no ADR/SPEC named anywhere in the paper). PDF via Tectonic.
+
+## Section order → files
+
+1. **Executive Summary** — `sections/00-exec-summary.tex`
+   The attention-optimization goal, the two enabling optimizations (GEMM
+   composite command + PE_IPCQ communication), the fused GQA capstone, the
+   headline results, and the bottom-line recommendation on which HW changes
+   are worth keeping. Front-loads conclusions; may run slightly longer than
+   a terse abstract.
+
+2. **Introduction** — `sections/01-introduction.tex`
+   1H focus is **attention-kernel optimization**: FlashAttention-style
+   tiling and Grouped Query Attention (GQA), building on prior work on
+   Multi-Head Attention (MHA, studied earlier — referenced as the
+   established baseline, not re-derived). Optimizing a fused attention
+   kernel requires two enabling optimizations, each studied in its own
+   right and then **combined inside the fused kernel**:
+   (a) **GEMM optimization** (the composite command), and
+   (b) **communication optimization** (PE_IPCQ).
+   Narrative arc: *two enablers → the fused GQA capstone that uses both.*
+   Motivate why HW-SW codesign (not software alone) is required.
+
+3. **The KernBench Platform** — `sections/02-platform.tex`
+   - 3.1 *Why KernBench* — source-level kernel execution, no compiler /
+     SW-stack dependency; isolate algorithm-level optimization.
+   - 3.2 *Execution model* — discrete-event graph; runtime API → sim_engine
+     → components; PE pipeline + composite commands; 2-pass data/timing.
+   - 3.3 *Latency model & accuracy* — traversal-based golden invariants;
+     structural CPU-dispatch cost model; known simplifications + calibration.
+   - 3.4 *Modeled Hardware Configuration* — shared platform-wide config from
+     `topology.yaml` (hierarchy counts, clocks, capacities/BW, link params,
+     cost-model constants). Per-experiment workload params live in §4–§6.
+
+4. **GEMM Acceleration via Composite Command** — `sections/03-gemm.tex`
+   necessity · design · results · analysis. (+ GEMM workload params)
+
+5. **All-Reduce Acceleration via PE_IPCQ** — `sections/04-allreduce.tex`
+   necessity · design · results · analysis. (+ All-Reduce topologies/sizes)
+
+6. **Fused Grouped Query Attention** (composite + PE_IPCQ) — `sections/05-gqa.tex`
+   necessity · design · results · analysis. (+ GQA seq/head/user configs)
+
+7. **Discussion** — `sections/06-discussion.tex`
+   Which HW changes are meaningful, and under what regimes (cross-cutting).
+
+8. **Conclusion** — `sections/07-conclusion.tex`
+   The codesign thesis, stated plainly, supported by the measured results.
+
+9. **Future Work — 2H** — `sections/08-future-work.tex`
+   Add the FFN/MoE layer toward full LLM decoding; how compute & data should
+   be distributed for agentic / MoE workloads.
+
+10. **References** *(optional)* — external literature only (FlashAttention,
+    Megatron-LM, GPT-3, Llama 3). No ADR/SPEC entries.
+
+## Per-section structure (§4/§5/§6)
+
+Each of the three optimization sections follows the same four beats:
+**(a) necessity → (b) design → (c) experimental results → (d) analysis & meaning.**
+
+## Build notes
+
+- GEMM (§4) and All-Reduce (§5): reuse committed figures/CSVs under
+  `src/kernbench/benches/1H_milestone_output/{gemm,ccl}/`.
+- GQA (§6): committed `sweep.json` holds op-counts only — latency + figures
+  are generated by an isolated harness under `scripts/paper/` at build time.
+  Measure only the implemented path; mark proposed designs (e.g. the
+  flat-ops `softmax_merge` decode opt2) as designed-not-measured.

docs/report/architecture-2026-1H.md

@@ -0,0 +1,836 @@
+# KernBench — Architecture Design Document
+*2026 1H*
+
+KernBench is a system-level, discrete-event simulator for AI-accelerator
+chiplet systems. It models the data-movement and control paths across
+the full hardware hierarchy and reports end-to-end execution latency
+for kernels dispatched to the device's compute units.
+
+This document is a public summary of the architecture as designed and
+implemented in the first half of 2026. It assumes no prior knowledge of
+the simulator's internal documents; terms specific to the system are
+defined on first use.
+
+---
+
+## Design Principles
+
+KernBench is grounded in two foundational commitments: every measured
+latency must trace to explicit, modeled events on the simulator's graph,
+and every behavioral claim must be verifiable through tests that target
+spec-level invariants rather than incidental implementation details.
+
+<!-- src: ADR-0013 Context, Decision -->
+The verification posture is verification-driven. Tests are written to
+validate the architectural contracts that the simulator exposes —
+correct routing, deterministic results, monotonic latency under
+increasing hop counts — rather than to mirror the call graph of the
+implementation. Two phases coexist: a fast timing phase that exercises
+the simulator's discrete-event engine and produces a log of operations
+with timestamps, and an optional data-replay phase that uses that log
+to compute real numerical results. Tests can target either phase.
+
+<!-- src: ADR-0033 Context, Decision -->
+The latency model is intentionally abstract rather than
+cycle-accurate. Each modeled node contributes a configurable per-node
+overhead, each link contributes wire delay plus byte-over-bandwidth
+serialization, and each terminal service contributes its own service
+time. The simulator does not attempt to reproduce cache coherence
+protocols, microarchitectural pipelines, or full PCIe/UCIe protocol
+correctness; those are explicitly outside the scope. The aim is a
+simulator that compares system-level configurations meaningfully and
+deterministically, not one that ships microarchitectural truths.
+
+<!-- src: ADR-0033 Decision, Consequences -->
+Determinism is a hard requirement. Given identical inputs — topology,
+routing policy, and request stream — the simulator must produce
+identical outputs, hop traces included. This rules out reliance on
+unordered set iteration on the critical path and forces every latency
+contribution to come from an explicitly scheduled event on a modeled
+component or link. There are no implicit waits, no hardcoded magic
+delays, and no shortcuts that bypass the modeled graph.
+
+---
+
+## High-level Architecture
+
+<!-- src: ADR-0003 Context, Decision -->
+The simulated system is a four-level hierarchy. A **Tray** holds one or
+more **SIPs** (system-in-package), each containing a 2D mesh of
+**CUBEs** plus one or more **IO chiplets** that connect the SIP to the
+host. Each CUBE contains a regular grid of **PEs** (processing
+elements) plus its own attached resources — high-bandwidth memory
+(HBM), a per-cube SRAM scratchpad, and a management CPU (M_CPU). The PE
+itself is a composite of nine sub-components rather than a monolithic
+core. This hierarchy is fixed; the parameters along each axis (counts,
+mesh dimensions, link widths) are configurable through the topology
+spec.
+
+<!-- src: ADR-0007 Context, Decision -->
+A clean separation runs along the request flow. A **runtime API** at
+the top is the host-facing surface; it exposes tensor and kernel
+operations, owns host-side allocation metadata, and is topology-
+agnostic — it does not route or fan out. Below it the **simulation
+engine** decomposes runtime operations into discrete graph requests
+(memory writes, memory reads, kernel launches, MMU map installs) and
+schedules events deterministically. At the bottom, **components** model
+device behavior on a graph of nodes connected by links; they
+implement the actual latency contributions and pass requests along.
+No component reaches up into the runtime API, and no runtime call
+shortcuts the engine.
+
+<!-- DIAGRAM: Four-level system hierarchy — Tray containing SIPs, each SIP showing its 2D cube mesh and IO chiplet; one cube blown up to show the router mesh, attached PEs, M_CPU, SRAM, and HBM partition. -->
+
+### Tray
+
+<!-- src: ADR-0003 Decision -->
+The Tray is the outermost boundary. It owns the host CPU on one side
+and one or more SIPs on the other, connected through a fabric switch.
+For collective communication that must traverse multiple SIPs, the
+fabric switch acts as the common rendezvous: device-side outbound
+traffic from one SIP routes through the switch and back into the
+target SIP's IO chiplet.
+
+### SIP
+
+<!-- src: ADR-0003 Decision, ADR-0017 Context -->
+A SIP packages a 2D mesh of CUBEs and one or more IO chiplets. The
+default topology used by the simulator is a 4×4 cube mesh; the
+mesh dimensions are configurable. Each cube on the boundary of the
+mesh connects to its neighbors over UCIe (die-to-die) links arranged
+on the four cardinal sides — north, south, east, and west. The IO
+chiplets sit on one side of the SIP and provide the bridge to the host
+across PCIe.
+
+<!-- src: ADR-0016 Context, Decision -->
+The IO chiplet itself contains its own internal network. A
+host-facing PCIe endpoint passes traffic to a small NOC ("network on
+chip"); from there it can branch to a control-plane CPU that processes
+kernel-launch messages, or it can take the direct memory data path to
+the cube's HBM controller. The decision to provide a direct memory
+path that bypasses the control CPU was a deliberate concession to
+keep host-issued memory writes from paying control-plane overhead on
+the data path.
+
+### CUBE
+
+<!-- src: ADR-0017 Decision -->
+Each CUBE owns a 2D mesh of NOC routers and a set of attached
+resources: PEs, the cube-local SRAM scratchpad, the management CPU
+(M_CPU), and the HBM partition (split across multiple PE-private
+slices for bandwidth). The router mesh uses deterministic XY routing.
+Attached components do not connect to each other directly — they all
+sit on the router mesh, and every cube-internal transfer pays the
+mesh distance from source to destination.
+
+<!-- src: ADR-0017 Decision -->
+The HBM partition is per-PE: each PE owns one HBM slice, and the
+controller exposes per-PE channels so that the same PE always
+addresses the same set of HBM channels. This makes the local-HBM
+bandwidth from a PE to its own slice predictable, while accesses to
+another PE's slice — or a different cube's slice — pay the mesh
+distance and any UCIe crossings.
+
+### PE
+
+<!-- src: ADR-0014 Context, Decision -->
+A PE is not a monolithic core. Internally it is a set of nine
+sub-components, each modeling one stage of a request's flow: a small
+control CPU, a tile-pipeline scheduler, a DMA engine, a fetch-store
+engine that moves data between the on-PE scratchpad and the register
+file, a GEMM compute engine, a math compute engine, the tightly-
+coupled memory (TCM, the on-PE scratchpad), an MMU for virtual-to-
+physical address translation, and an inter-PE collective queue
+(IPCQ). The scheduler decomposes higher-level operations into per-tile
+stage sequences, and tile tokens self-route from one sub-component
+to the next.
+
+<!-- DIAGRAM: PE internal layout — the nine sub-components and the edges that connect them; tile token flowing through DMA_READ → FETCH → GEMM → STORE → DMA_WRITE. -->
+
+---
+
+## Detailed Architecture
+
+This section describes each modeled device-side component in turn.
+Components are listed in the alphabetical order used by the
+simulator's source tree.
+
+### forwarding
+
+<!-- src: ADR-0037 Context, Decision -->
+The forwarding component is the generic routing relay used wherever a
+node only needs to apply a small processing overhead and pass the
+request to the next hop. NOC routers, conn nodes, and ucie phys all
+reduce to this. Its first act on receiving a request is to apply the
+per-node overhead configured for it in the topology spec; after the
+overhead it simply hands the request to the next hop along the path.
+
+<!-- src: ADR-0037 Decision, Consequences -->
+The decision to share one implementation across these roles was made
+to keep the simulator's component set small without sacrificing
+modeling fidelity. Each instance still carries its own overhead and
+its own link bandwidth contributions, so different roles still produce
+different timing. What is shared is the dispatcher loop, not the
+parameter values.
+
+### hbm_ctrl
+
+<!-- src: ADR-0034 Context, Decision -->
+The HBM controller is the terminal node for all memory traffic that
+reaches HBM. Internally it owns a number of pseudo channels, partitioned
+per-PE so that each PE addresses a deterministic subset. On a request
+arrival the controller first selects the right pseudo channel from the
+target address, then enters a chunk-loop that drains the requested
+size in fixed-size flits over the channel's bandwidth.
+
+<!-- src: ADR-0034 Decision, Consequences -->
+The chunk-loop pattern replaces an earlier all-at-once drain. The
+benefit is that the controller no longer presents a flit-aware fabric
+with a single bulk transfer; instead it emits flits at a paced rate
+matching the channel bandwidth, which makes cross-flow contention
+visible. The bandwidth budget is calibrated against the configured
+HBM total bandwidth divided across the channel count.
+
+### io_cpu
+
+<!-- src: ADR-0036 Context, Decision -->
+The IO_CPU is the control-plane processor sitting inside the IO chiplet.
+It receives kernel-launch messages from the host, decodes them, and
+dispatches per-cube launches to the cube's management CPU. Pure memory
+operations bypass it entirely, taking the direct data path established
+inside the IO chiplet.
+
+<!-- src: ADR-0036 Decision -->
+On receiving a kernel-launch message, the IO_CPU consults the message's
+shard list — which already names the target SIP, cube, and PE for each
+piece of the tensor argument — and forwards a per-cube launch to each
+cube the kernel needs to reach. This makes the IO_CPU a deterministic
+fan-out point: it does not decode physical addresses to route, it just
+follows the explicit per-shard targets it was handed.
+
+### m_cpu
+
+<!-- src: ADR-0035 Context, Decision -->
+The M_CPU is the cube's management processor. It owns two distinct
+roles: as a control-plane fan-out point for kernel launches arriving
+from the IO chiplet, and as a DMA endpoint for host-initiated memory
+writes that need to land in this cube's HBM. The control role
+forwards launches to the right PE control CPUs; the DMA role places
+the actual bytes into HBM through the router mesh.
+
+<!-- src: ADR-0035 Decision -->
+The component model deliberately distinguishes the two roles because
+their routing differs: the control fan-out path uses command-kind
+links that do not appear on data-path routes, while the DMA path uses
+the same router mesh as PE-initiated DMA, with PE-internal nodes
+excluded. The routing layer knows about both modes and selects the
+appropriate adjacency at request time.
+
+### pcie_ep
+
+<!-- src: ADR-0038 Context, Decision -->
+The PCIE endpoint is the protocol boundary at the host-device edge.
+Its first act on each incoming request is to apply a configured
+protocol-processing overhead; after that it simply forwards. There is
+no internal queuing model, no retry, and no TLP-level fidelity — those
+are deliberately outside scope. The endpoint is bidirectional: host →
+device traffic (memory writes, kernel launches) flows one way, and
+device-side outbound traffic (cross-SIP collective sends) flows the
+other.
+
+<!-- src: ADR-0038 Decision, Alternatives Considered, Consequences -->
+A more detailed PCIe model was considered and rejected. The simulator
+is targeting system-level latency comparisons; making the endpoint
+heavier with credit-management and retry logic would not improve the
+metrics being studied. The decision keeps the endpoint as the
+documented protocol-boundary node, named consistently so routing
+helpers can locate it by SIP and IO instance.
+
+### pe_cpu
+
+<!-- src: ADR-0014 Decision -->
+The PE control CPU is the entry point for kernel work arriving from
+the cube's management CPU. It receives kernel-launch messages, resolves
+the kernel function by name, and hands execution to the scheduler with
+the resolved tensor arguments. From the scheduler's point of view, the
+PE_CPU is the upstream source of high-level commands; from the rest
+of the system's point of view, the PE_CPU is where a kernel's
+execution begins on a given PE.
+
+### pe_dma
+
+<!-- src: ADR-0014 Decision, ADR-0023 Decision -->
+The DMA engine on each PE has two distinct modes. In the standard PE
+pipeline it consumes tile tokens issued by the scheduler, acquires a
+read or write channel (modeled as a one-in-flight resource per
+direction), and runs the bytes to or from HBM through the mesh. In
+its collective mode it forwards send tokens for the cube's IPCQ into
+the fabric, snapshotting the source data at send time so later
+mutations cannot race the receiver's read. Both modes share the same
+channel resources but differ in their downstream handling — one
+returns when the round-trip completes, the other dispatches
+fire-and-forget.
+
+### pe_fetch_store
+
+<!-- src: ADR-0014 Decision -->
+The fetch-store engine is the bridge between the on-PE scratchpad
+(TCM) and the register file. It does not run DMA; it only moves bytes
+internally. On receiving a tile-stage token it sends a short request
+to the TCM, waits for the bandwidth-serialized delay, and continues
+the pipeline. The split between this engine and the TCM lets the
+scratchpad model its own read/write bandwidth independently.
+
+### pe_gemm
+
+<!-- src: ADR-0014 Decision -->
+The GEMM engine is the matrix-multiply compute unit. Tile tokens
+arriving at this stage carry the per-tile dimensions, and the engine
+contributes a service time accounting for one fused multiply-add over
+the tile's macs. Composite operations (where the same tensor pair is
+streamed across many tiles) reuse the engine through the scheduler;
+the engine itself is stateless between tiles.
+
+### pe_ipcq
+
+<!-- src: ADR-0023 Context, Decision -->
+The IPCQ — inter-process communication queue — is each PE's
+collective-communication endpoint. It owns ring buffers that hold
+inbound messages from neighbor PEs and bookkeeping for send credits.
+Direction names ("N", "S", "E", "W" for cube-internal neighbors and
+"global_*" for cross-SIP neighbors) are resolved to physical peer
+endpoints by a neighbor table installed at process-group creation
+time. The component itself does not move bytes — it issues DMA tokens
+through the local PE_DMA, which performs the actual cross-PE
+transfer.
+
+<!-- src: ADR-0023 Decision, Consequences -->
+A key invariant is that the inbound terminal — where data lands at
+the receiver — pays the link bandwidth drain plus any cube-internal
+mesh hop to the slot's backing memory. This prevents IPCQ from
+silently outpacing raw DMA at large transfer sizes. Outbound sends
+are fire-and-forget; credit return is the only backpressure signal.
+
+### pe_math
+
+<!-- src: ADR-0014 Decision -->
+The math engine handles element-wise and reduction operations. It
+consumes tile tokens carrying an operation kind (`exp`, `sum`, `max`,
+`where`, etc.) and contributes a service time proportional to the
+number of elements processed. Like the GEMM engine it is stateless;
+chained epilogues (a sequence of math operations after a GEMM tile)
+are scheduled as separate stages.
+
+### pe_mmu
+
+<!-- src: ADR-0039 Context, Decision -->
+The MMU has two roles, exposed through one component. As a node on
+the cube NOC it receives MMU-map and MMU-unmap messages and updates
+its internal page table, so that the runtime API can install
+virtual-to-physical mappings with measured fabric latency. As a
+utility object held inside the PE it offers synchronous translate
+calls to the PE's DMA and GEMM engines without taking simulator time
+itself; the calling engine pays any configured TLB overhead in its
+own process.
+
+<!-- src: ADR-0039 Decision, Alternatives Considered -->
+The page table supports multiple disjoint regions inside a single
+page, with later-write-wins semantics on overlap. This is a deliberate
+simulator stopgap to support parallelization policies that shard data
+at sub-page granularity without silent mis-routing through a real
+hardware MMU's one-PA-per-entry assumption. A real MMU does not work
+this way; the model documents this as a simplification.
+
+### pe_scheduler
+
+<!-- src: ADR-0014 Decision -->
+The scheduler is the sole dispatcher inside a PE. Simple commands are
+routed directly to the right engine. Composite commands generate a
+tile plan, and the resulting tile tokens are fed into the pipeline.
+Self-routing keeps the scheduler off the per-stage hot path: each
+engine, on finishing a stage, advances the token to the next stage's
+component itself, so the scheduler only does initial dispatch and
+completion tracking.
+
+### pe_tcm
+
+<!-- src: ADR-0040 Context, Decision -->
+The TCM is the per-PE tightly-coupled scratchpad memory. It models
+time only, not data — the actual payload lives in the simulator's
+memory store. Read and write are independent channels: each is
+modeled as a one-in-flight resource, so same-direction requests
+serialize but a read and a write can overlap. The bandwidth of each
+direction is configured separately and applied as bytes-over-bandwidth
+on each request.
+
+<!-- src: ADR-0040 Decision, Alternatives Considered -->
+The decision to keep read and write on separate channels was made
+because the PE pipeline's normal case overlaps fetch (read) and store
+(write). Collapsing them into a single shared channel would have
+artificially serialized that overlap and produced an incorrect
+bandwidth ceiling.
+
+### sram
+
+<!-- src: ADR-0041 Context, Decision -->
+The cube SRAM is a per-cube scratchpad attached to one of the cube's
+routers. As a node it applies a configured access overhead, pays the
+link-bandwidth drain stamped on the incoming request, and sends a
+response on the reverse path. It is a terminal — it does not forward.
+
+<!-- src: ADR-0041 Decision, Consequences -->
+A second role is as one of three backing-memory tiers (TCM, SRAM, HBM)
+that an inter-PE collective slot can live in. When the slot lives in
+SRAM, the PE_DMA pays the slot read or write latency directly using
+the configured SRAM bandwidth and overhead; the SRAM component does
+not need to know about collective semantics. This separation keeps
+the SRAM component agnostic to the collective subsystem.
+
+### tiling
+
+<!-- src: ADR-0042 Context, Decision -->
+The tile-plan generator is not a runtime component — it is a pure
+module of functions that take a problem shape (matrix dimensions, tile
+sizes) and produce an ordered list of tile-stage sequences. The
+scheduler consumes this list. Each tile's stage sequence depends on
+how its operands are staged: operands streamed from HBM produce
+DMA_READ stages, operands already resident in TCM (because they were
+loaded eagerly upfront) skip them.
+
+<!-- src: ADR-0042 Decision, Consequences -->
+The plan generator is intentionally pure — given the same input it
+returns the same plan, with no simulator events created. This lets
+the rest of the system reason about tile sequences as data, and it
+makes the plan testable in isolation without simulator state. New
+plan variants (for example, K-major or DTensor-aware plans) can be
+added as new functions following the same shape.
+
+---
+
+## Implementation Decisions
+
+This section collects cross-cutting decisions — algorithms, policies,
+schemes, and contracts — that span multiple components rather than
+living inside one.
+
+### Address Scheme
+
+<!-- src: ADR-0001 Context, Decision -->
+Every physical address in the simulator decodes into a structured
+location. A fixed-width physical address carries the SIP id, the
+cube id within the SIP, a type discriminator (HBM vs PE-resource vs
+others), and a type-specific offset. HBM addresses additionally encode
+the per-PE slice offset so the controller can determine which PE
+owns the target slice without external lookup. The layout is
+deliberately reserved rather than packed-to-fit, so new sub-units can
+be added at the type-discriminator level without rewriting existing
+addresses.
+
+<!-- src: ADR-0011 Context, Decision -->
+On top of physical addressing, the simulator supports three address
+models that the runtime API selects between. Direct physical
+addressing is retained as a fallback. Virtual addressing — the
+current default — gives each tensor a contiguous virtual range at
+deployment, with the per-PE MMU translating per access; an
+alternative logical-address scheme remains a future option. The
+virtual-address path is what every modern test path takes; the PA
+fallback is used by the MMU itself when no mapping exists for an
+address (a deliberate signal, not an error).
+
+<!-- src: ADR-0011 Decision, Consequences -->
+Tensor placement is represented as a list of physical-address shards,
+each tagged with target SIP, cube, and PE, plus a single tensor-wide
+virtual base. This means a kernel sees one virtual base for the whole
+tensor while the host driver and the engine still know exactly where
+each shard lives. Replicated tensors get per-cube local PA mappings;
+sharded tensors broadcast their mapping across cubes within a SIP.
+
+### Routing, Distance & Helper API
+
+<!-- src: ADR-0002 Context, Decision -->
+Routing is policy-driven, deterministic, and topology-aware. Given a
+source, a destination, and an intent — for example, PE-initiated
+DMA versus host-initiated memory write versus a generic
+component-to-component query — the routing layer picks the right
+path. The intent matters because different traffic types must avoid
+different categories of edges: PE-initiated DMA should not traverse
+command-only links; M_CPU DMA should not pass through PE-internal
+pipeline edges; cube-local transfers should not use the
+zero-distance UCIe bus that would otherwise look attractive to a
+shortest-path search.
+
+<!-- src: ADR-0051 Decision -->
+The routing layer therefore maintains four separate adjacency graphs
+at construction, each excluding a different category of edges, and
+picks the appropriate one per intent. On top of the graphs sits a
+helper API that hides the topology's naming convention: callers ask
+for the PCIe endpoint of a given SIP, the M_CPU of a given cube, or
+the HBM destination for a given physical address, and receive the
+corresponding node id. No component constructs node-id strings
+directly; if the naming convention ever changes, the change is local
+to the helper layer.
+
+<!-- src: ADR-0051 Decision, Consequences -->
+Path-finding itself uses Dijkstra with explicit per-edge weights
+(routing weight is allowed to differ from physical distance — for
+example, UCIe is configured to be routing-preferable). Tie-breaks
+follow insertion order, which keeps results deterministic. Paths
+between unreachable nodes raise rather than returning empty, surfacing
+topology errors immediately.
+
+### Memory Semantics and Local-HBM Bandwidth
+
+<!-- src: ADR-0004 Context, Decision -->
+A PE accessing its own HBM slice through its own cube's NOC must see
+the full local HBM bandwidth — that is the model's intent. Memory
+traffic accumulates latency from per-component overhead and
+bytes-over-link-bandwidth serialization along the path, but the
+controller does not throttle below the slice's allotted bandwidth.
+Cross-PE-slice accesses inside the same cube, cross-cube accesses
+through UCIe, and cross-SIP accesses through PCIe each pay
+progressively more overhead as the path grows.
+
+### Topology Compilation, Diagrams & Builder Algorithms
+
+<!-- src: ADR-0006 Context, Decision -->
+Topology is configurable, not hardcoded. The simulator reads a YAML
+spec, compiles it into a flat graph of nodes and edges plus four
+view projections at different abstraction levels — system, SIP, cube,
+PE — and uses the compiled graph as the single source for both
+execution and visualization. Distance metadata used by routing is
+extracted at compile time so that diagrams and routing decisions
+agree by construction.
+
+<!-- src: ADR-0005 Context, Decision -->
+Diagrams are derived artifacts of the compiled topology. The visualizer
+produces one SVG per view at the appropriate abstraction level; nothing
+in the diagrams is hand-drawn or hand-positioned. Distance-aware
+layout rules place nodes in the diagrams using the same coordinates
+that routing uses to compute distance, so a diagram that "looks
+wrong" is a signal that the topology itself has a problem, not the
+visualizer.
+
+<!-- src: ADR-0053 Decision -->
+Inside a cube the router mesh is generated automatically. PE corner
+positions are fixed by convention; the relay-column algorithm
+inserts additional grid columns whenever the gap between adjacent PE
+columns would exceed a tunable maximum. HBM occupies a central
+exclusion zone — router slots inside the zone are deliberately empty,
+since HBM controllers attach as separate named nodes. M_CPU and SRAM
+attach to the nearest router by Euclidean distance from their
+configured placement coordinates, and UCIe physical lanes distribute
+along the boundary rows and columns. The whole mesh is cached
+beside the topology spec and invalidated only when one of a small set
+of layout-relevant fields changes.
+
+<!-- DIAGRAM: One cube's router mesh — rows × cols of routers with HBM exclusion zone in the middle, PEs/M_CPU/SRAM attaching to nearest routers, UCIe PHYs along the perimeter. -->
+
+### Tensor Deployment and Allocation
+
+<!-- src: ADR-0008 Context, Decision -->
+Tensor deployment in the runtime API produces a list of physical-address
+shards plus a single tensor-wide virtual base. The host allocator
+walks the data-parallelism policy, computes per-shard placement, and
+emits the per-shard physical addresses through the per-PE allocators.
+No separate "allocate then later attach to a device" RPC exists —
+allocation and deployment are a single operation that produces a
+deployed tensor handle.
+
+### Memory Allocator Algorithms
+
+<!-- src: ADR-0048 Context, Decision -->
+Each per-PE allocator owns two channels — HBM slice and TCM — each
+backed by an offset-keyed free-list. Allocation is first-fit; freeing
+coalesces with adjacent free blocks. A device-wide virtual allocator
+sits above the per-PE allocators, aligns requests up to the configured
+page size, and coalesces on free in the same way. The trade-off is
+explicit: first-fit is simpler and cheaper than best-fit or buddy
+allocation, and the simulator's workload is stack-like enough
+(deploy / kernel / free in matched order) that fragmentation is not
+a practical concern.
+
+<!-- src: ADR-0048 Decision, Consequences -->
+Allocation failure raises rather than silently returning a partial
+result. A partial tensor reaching the engine would route over wrong
+PAs and silently corrupt simulator output, so an out-of-memory signal
+is preferred. The free path trusts its caller to pass back exactly
+what was allocated; the small risk of caller error in exchange for
+fast common-case freeing is documented as a deliberate trade.
+
+### Kernel Execution and Host-Device Messaging
+
+<!-- src: ADR-0009 Context, Decision -->
+Kernel execution decomposes into a small set of messages that travel
+the device graph. The host issues a single kernel-launch message; the
+IO_CPU fans it out per-cube; the cube M_CPU fans it out per-PE; the
+PE CPU resolves the kernel and runs it through the scheduler.
+Completion flows back the same way, gated by per-shard completion
+tracking. Memory operations follow the same pattern: a memory write
+or read travels as one message that the engine routes to the right
+HBM controller, with a response taking the reverse path.
+
+<!-- src: ADR-0012 Context, Decision -->
+The schema between the host and the device-side IO CPU is PA-first
+and shard-tagged. Every byte of host-issued payload arrives with an
+explicit target SIP, cube, PE, and physical address. The IO_CPU does
+not decode addresses to derive placement — placement is named
+explicitly by the shard list. This makes the host-device interface
+deterministic and keeps the routing helper free of host-derived
+intent.
+
+### CLI Surface and Semantics
+
+<!-- src: ADR-0010 Context, Decision -->
+The command-line interface exposes four subcommands. A bench runner
+loads a topology, resolves a registered benchmark by name or index,
+and runs it on a selected device. A bench-listing command enumerates
+the registered benchmarks. A probe utility runs a fixed catalog of
+traffic patterns through the engine for latency and bandwidth
+verification. A web viewer renders the topology in a browser. A
+benchmark instance is always single-device by convention; multi-SIP
+collective work happens inside the benchmark through the launcher
+abstraction, not by multiplexing the CLI.
+
+### Component Port and Wire Fabric Model
+
+<!-- src: ADR-0015 Context, Decision -->
+Every modeled component exposes input and output ports, and every
+edge in the topology connects an output port on one component to an
+input port on another. Bandwidth and propagation delay are properties
+of the wire between ports, not of the component endpoints. A
+component's responsibility is to apply its configured per-node
+overhead and either forward to the next hop or terminate; the wire
+charges the byte-over-bandwidth serialization separately.
+
+<!-- src: ADR-0015 Decision, Consequences -->
+This separation lets components be swapped behind their port
+interface without changing the rest of the model, and it keeps
+bandwidth contention at the wire level where multiple components may
+contend for the same edge. Future component models can refine
+internal behavior without disturbing the fabric.
+
+### Two-Pass Data Execution
+
+<!-- src: ADR-0020 Context, Decision -->
+The simulator runs in two passes. The first pass — fast and always
+on — runs the discrete-event engine and records every data operation
+in an operation log with timestamps, component identifiers, and per-
+operation parameters. The second pass — optional, opt-in — replays
+the log against an in-memory tensor store to produce actual numerical
+results. Tests that only need timing skip the second pass; tests that
+need to verify correctness opt in.
+
+<!-- src: ADR-0020 Decision, Consequences -->
+The split lets the timing engine remain unconcerned with data
+semantics: kernels move handles around, not bytes. The replay phase
+recovers data semantics from the recorded operations, in their
+original time order with a small set of secondary-sort rules. The
+op-log records carry enough metadata — input snapshots for compute
+operations, source snapshots for cross-component copies — that the
+replay phase cannot mis-order with respect to in-flight mutations.
+
+### Sim-engine Op Log and Memory Store Schemas
+
+<!-- src: ADR-0052 Context, Decision -->
+The operation log holds typed records with seven fields each: start
+and end timestamps, the component that issued the operation, an
+operation kind ("memory", "gemm", "math"), an operation name, a
+parameter dictionary, and a (currently unused) dependency list.
+Records are kept in stable timestamp order. The parameter dictionary
+varies by operation: a DMA read carries source address and byte count;
+a GEMM carries operand shapes, dtypes, and address spaces; a math
+operation carries input addresses and snapshots.
+
+<!-- src: ADR-0052 Decision, Consequences -->
+The companion memory store is a two-level dictionary keyed by
+address space ("hbm", "tcm", "sram", others) and integer address.
+Reads and writes are reference-based — no copy by default — so
+callers wanting to detach a snapshot must copy explicitly. This is
+deliberate: the engine-internal snapshot paths copy at well-defined
+points (math input capture, HBM source capture for DMA writes,
+inbound collective copies) and downstream replay code therefore
+sees stable data even when slot or scratch addresses are reused by
+later operations.
+
+### 2D Grid Program Identity
+
+<!-- src: ADR-0022 Context, Decision -->
+Inside a kernel the program identity is two-dimensional. The
+first axis corresponds to the PE index within a cube; the second
+corresponds to the cube index within a SIP. Together they let a
+kernel address its position both within its cube and within the
+larger system without needing to know the full topology. Total
+program counts along each axis are exposed symmetrically.
+
+### Parallelism — SIP Launcher, DPPolicy, Megatron-TP, AHBM Backend, and CCL Algorithm Module
+
+<!-- src: ADR-0024 Context, Decision -->
+The launcher model treats each SIP as one rank. Inside a process the
+launcher spawns one greenlet per SIP rank; the rank is bound to its
+greenlet so that any code running in that worker sees the right
+distributed-style rank. This is a deliberately PyTorch-compatible
+shape: a benchmark looks like a small DDP training script — initialize
+a process group, spawn workers, each worker runs the same body.
+
+<!-- src: ADR-0026 Context, Decision -->
+Data-parallelism policy lives in a single object that names the
+sharding strategy along the cube axis (replicate, row-wise,
+column-wise) and along the PE axis (same set of values), and optionally
+overrides the number of cubes or PEs participating. The policy is
+intra-device — it does not cross SIP boundaries. SIP-level parallelism
+is the launcher's responsibility, and the two axes compose
+orthogonally.
+
+<!-- src: ADR-0027 Context, Decision -->
+A Megatron-style tensor-parallel API sits on top of the launcher and
+the DP policy. Layer-level building blocks — column-parallel linear,
+row-parallel linear, all-reduce — name their sharding intent in terms
+the launcher and the placement policy can compose. This is the layer
+that bench code typically writes against.
+
+<!-- src: ADR-0047 Context, Decision -->
+For collective operations the runtime exposes a PyTorch-compatible
+distributed backend named "ahbm". On process-group initialization the
+backend loads the configured collective-algorithm module, resolves
+the world size (priority: explicit ccl.yaml override → defaults
+section → topology SIP count), imports the algorithm module
+dynamically, derives the SIP topology kind, and pushes the inter-PE
+neighbor table to every participating PE. From that point on, an
+all-reduce call dispatches the algorithm's kernel function across
+all ranks.
+
+<!-- src: ADR-0050 Context, Decision -->
+A collective-algorithm module is a Python module with a small, fixed
+contract. It exposes topology-kind integer constants, a name-to-kind
+mapping for the YAML configuration, a kernel-arguments builder, and
+a kernel function — the kernel function being aliased to the name
+`kernel` so the backend can find it generically. The kernel itself
+takes the tensor pointer, the per-cube element count, cube mesh
+width and height, the world size, the current rank, and the SIP
+topology dimensions; the backend appends those last four arguments
+automatically. New collectives slot in by adding a new module that
+follows this shape.
+
+<!-- src: ADR-0027 Decision, Consequences -->
+The combination is deliberate: bench authors get to write code that
+looks like a regular distributed training script, while the launcher,
+backend, and placement policies behind it remain free to redirect
+work to the right SIP, cube, and PE without exposing topology to the
+kernel.
+
+### IPCQ Direction Addressing
+
+<!-- src: ADR-0025 Context, Decision -->
+Inside a collective algorithm, peer PEs are named by direction —
+"N", "S", "E", "W" for cube-internal neighbors, and "global_*" for
+cross-SIP neighbors. Direction addressing is the addressing scheme:
+the algorithm names a direction, the IPCQ neighbor table installed
+at process-group time resolves the direction to the peer endpoint's
+physical-address coordinates, and the PE_DMA performs the actual
+transfer. The algorithm itself does not see PA arithmetic — direction
+is the user-facing handle.
+
+### Intercube All-Reduce
+
+<!-- src: ADR-0032 Context, Decision -->
+The default all-reduce algorithm uses a center-rooted bidirectional
+phase inside each SIP's cube mesh followed by an inter-SIP exchange
+on the mesh's root cube, and then a bidirectional broadcast back
+out. Center-rooting halves the in-cube hop count compared with a
+corner-rooted walk. The inter-SIP exchange itself follows the
+configured SIP topology — ring, torus, or non-wrapping mesh —
+selected at runtime through the SIP-topology kind integer the
+backend passes to the kernel.
+
+### Evaluation Harnesses
+
+<!-- src: ADR-0043 Context, Decision -->
+The all-reduce evaluation harness drives correctness and the
+latency/buffer-kind sweeps through the public distributed path —
+initialize process group, spawn workers, call all-reduce — rather
+than the lower-level engine interface. A shared helper module factors
+out the setup; sweep tests cover the buffer-kind tiers (TCM, SRAM,
+HBM) and the inter-SIP topology variants. The plots produced by the
+harness are part of its output contract; the harness regenerates them
+on demand.
+
+<!-- src: ADR-0044 Context, Decision -->
+The GEMM evaluation harness is split into two layers. A heavy
+shape-and-variant sweep lives as a manual script — it runs the same
+composite-GEMM benchmark across many shapes and operand-staging
+variants, harvests the resulting op-log, and writes a JSON summary.
+A faster figure-generation layer lives in the test suite and consumes
+that JSON to render plots. The split keeps the heavy data
+generation explicit and out of the regular test path.
+
+### Bench Module Contract
+
+<!-- src: ADR-0045 Context, Decision -->
+Adding a new benchmark requires only dropping a file into the
+benchmarks directory. The file registers one or more benchmark
+functions through a small decorator that takes a kebab-case name and
+a human-readable description. The decorator is the registration
+mechanism — there is no separate manifest. Each benchmark function
+takes one argument, conventionally named `torch`, which is the
+runtime context exposing tensor allocation, kernel launch,
+distributed APIs, and process-spawning. The function name is `run` by
+convention.
+
+<!-- src: ADR-0045 Decision, Consequences -->
+A benchmark must submit at least one operation, or the runner
+returns an error. A benchmark instance is single-device by default;
+when a benchmark is collective, it uses the distributed-process-spawn
+pattern internally — one worker greenlet per rank, with each worker
+binding to its rank. Multi-device benchmark patterns outside that
+shape are not supported.
+
+### Kernel-side `tl.*` API
+
+<!-- src: ADR-0046 Context, Decision -->
+Inside a kernel function, the `tl` argument exposes the kernel-side
+API in a shape that mirrors the conventions of established
+GPU-kernel languages. Categories: reference handles that name HBM
+data without issuing DMA; data movement (load, store) that does
+issue DMA; GEMM and math compute (dot, composite, the unary and
+binary math operations, reductions); index and scalar helpers
+(program identity, range-builders); metadata-only operations like
+transpose; and the collective primitives (send, receive,
+non-blocking receive). Tensor handles support arithmetic operators
+via a thread-local active context so kernel code reads naturally.
+
+<!-- src: ADR-0046 Decision, Consequences -->
+The API supports two execution modes. A command-list mode records
+operations into a list without consuming simulator time — useful for
+inspection and lightweight tests. A greenlet-driven mode runs the
+kernel as a child greenlet that switches back to the simulator on
+each `tl.*` call; the simulator drives the event scheduler and hands
+real data back to the kernel as DMA reads complete. The two modes
+share the same surface; the kernel does not know which one it is
+running under.
+
+### Probe Subcommand
+
+<!-- src: ADR-0049 Context, Decision -->
+The probe utility runs three families of traffic patterns through
+the engine — host-to-device writes at increasing hop counts,
+device-to-host reads at increasing hop counts, and PE-initiated DMA
+across the cube mesh — and reports actual latency, the analytical
+formula breakdown, effective bandwidth, bottleneck bandwidth, and
+utilization. A fixed reference size is used for the summary table;
+a separate utilization-versus-size sweep covers a logarithmic range
+of transfer sizes. Each case runs in its own engine instance so
+cases do not perturb each other.
+
+<!-- src: ADR-0049 Decision, Consequences -->
+The probe also checks a small set of invariants automatically:
+monotonic latency increase with hop count, device-to-host latency
+at least as large as host-to-device for the same hop count, and a
+faster best-case path than worst-case for cross-cube PE DMA. Failures
+print prominently. The output is meant for human reading; automated
+parsing should not depend on column widths or whitespace.
+
+---
+
+This document summarizes 46 architecture decisions captured during
+the first half of 2026. It is regenerated mechanically from the
+decision corpus; sources are recorded in HTML comments throughout.

scripts/gemm_sweep.py

@@ -1,237 +1,20 @@
 """Sweep GEMM shapes through kernbench and dump PE_accelerator engine times.
 
-For each shape:
-  - run benches.matmul_composite via the same run_bench path the CLI uses
-  - read result.engine.op_log
-  - filter to per-PE engines: pe_dma, pe_fetch_store, pe_gemm, pe_math
-  - record sum-of-durations (engine occupancy) AND wall-clock active interval
+Thin wrapper: the sweep logic now lives in
+``kernbench.benches.milestone_1h_gemm`` (the single home, ADR-0054, also the
+``milestone-1h-gemm`` bench). This script remains the manual entry point for
+regenerating ``docs/diagrams/gemm_sweep.json`` on demand and honors the same
+``SWEEP_SHAPES`` / ``SWEEP_TOPOLOGY`` env overrides.
 
-Output: docs/diagrams/gemm_sweep.json
+    python scripts/gemm_sweep.py
 """
 from __future__ import annotations
 
-import json
-import os
-import sys
-import time
-from pathlib import Path
-
-# Default sweep covering under-tile, single-tile, multi-tile, and asymmetric regimes.
-# Each entry is either a single integer (square M=K=N=S) or "MxKxN".
-# Override via env: SWEEP_SHAPES="16,32,16x2048x16,..."
-DEFAULT_SHAPES = [
-    "32x32x32",       # 1 tile, K=32 < TILE_K=64 → under-tile in K
-    "32x64x32",       # 1 tile, exact single-tile fit
-    "32x128x32",      # 2 tiles, aligned
-    "32x128x128",     # 8 tiles, aligned
-    "32x3072x32",     # 48 tiles, all K-axis (tall-skinny)
-    "8x128x128",      # 8 tiles, but M=8 < TILE_M=32 → MAC array under-fed
-    "128x8x128",      # 16 tiles, but K=8 < TILE_K=64 → MAC array under-fed
-    "512",            # 2048 tiles, fully aligned — "well-pipelined" reference
-]
-
-# Operand-staging variants exercised per shape.
-VARIANTS = ["ref_ref", "load_ref", "load_load"]
-
-# Engines whose timings we collect (component_id suffix match).
-ENGINES = ["pe_dma", "pe_fetch_store", "pe_gemm", "pe_math"]
-
-# Per-stage breakdown labels (StageType enum names from pe_types.py).
-STAGES = ["DMA_READ", "DMA_WRITE", "FETCH", "STORE", "GEMM", "MATH"]
-
-# Scheduler tile sizes (mirror of PeSchedulerComponent.TILE_M/K/N).
-TILE_M, TILE_K, TILE_N = 32, 64, 32
-
-OUT_PATH = Path(__file__).parent.parent / "docs" / "diagrams" / "gemm_sweep.json"
-
-
-def _engine_wall_ns(records, suffix: str) -> float:
-    """Wall-clock interval the engine was active (union of overlapping ops)."""
-    intervals = [(r.t_start, r.t_end) for r in records
-                 if r.component_id.endswith("." + suffix)]
-    if not intervals:
-        return 0.0
-    intervals.sort()
-    merged_end = intervals[0][1]
-    merged_start = intervals[0][0]
-    total = 0.0
-    for s, e in intervals[1:]:
-        if s <= merged_end:
-            merged_end = max(merged_end, e)
-        else:
-            total += merged_end - merged_start
-            merged_start, merged_end = s, e
-    total += merged_end - merged_start
-    return total
-
-
-def _engine_occupancy_ns(records, suffix: str) -> float:
-    return sum(r.t_end - r.t_start for r in records
-               if r.component_id.endswith("." + suffix))
-
-
-def _engine_count(records, suffix: str) -> int:
-    return sum(1 for r in records if r.component_id.endswith("." + suffix))
-
-
-def _stage_occupancy_ns(records, stage_type: str) -> float:
-    """Sum t_end - t_start over op_log records whose params.stage_type matches.
-
-    Requires op_log records produced post the TileToken stage_type capture
-    (sim_engine/op_log.py).
-    """
-    return sum(
-        r.t_end - r.t_start
-        for r in records
-        if r.params.get("stage_type") == stage_type
-    )
-
-
-def _stage_wall_ns(records, stage_type: str) -> float:
-    """Interval-union wall-clock for records whose stage_type matches."""
-    intervals = sorted(
-        (r.t_start, r.t_end) for r in records
-        if r.params.get("stage_type") == stage_type
-    )
-    if not intervals:
-        return 0.0
-    total = 0.0
-    cs, ce = intervals[0]
-    for s, e in intervals[1:]:
-        if s <= ce:
-            ce = max(ce, e)
-        else:
-            total += ce - cs
-            cs, ce = s, e
-    total += ce - cs
-    return total
-
-
-def _stage_count(records, stage_type: str) -> int:
-    return sum(1 for r in records if r.params.get("stage_type") == stage_type)
-
-
-def _run_one(M: int, K: int, N: int, topology: str, variant: str = "ref_ref") -> dict:
-    os.environ["MATMUL_M"] = str(M)
-    os.environ["MATMUL_K"] = str(K)
-    os.environ["MATMUL_N"] = str(N)
-    os.environ["MATMUL_VARIANT"] = variant
-
-    # Late imports so env vars are read by matmul_composite at module load.
-    # Force re-import to pick up new env values.
-    for mod_name in [m for m in list(sys.modules) if m.startswith("kernbench.benches.matmul_composite")]:
-        del sys.modules[mod_name]
-
-    from kernbench.benches.registry import resolve as resolve_bench
-    from kernbench.runtime_api.bench_runner import run_bench
-    from kernbench.runtime_api.types import resolve_device
-    from kernbench.sim_engine.engine import GraphEngine
-    from kernbench.topology.builder import resolve_topology
-
-    topo = resolve_topology(topology)
-    bench = resolve_bench("matmul-composite").run
-    device = resolve_device(None)
-
-    t0 = time.time()
-    result = run_bench(
-        topology=topo, bench_fn=bench, device=device,
-        engine_factory=lambda t, d: GraphEngine(
-            getattr(t, "topology_obj", t), enable_data=True,
-        ),
-    )
-    wall = time.time() - t0
-
-    op_log = result.engine.op_log
-    if not result.completion.ok:
-        raise RuntimeError(f"bench failed at M={M},K={K},N={N}: {result.completion}")
-
-    # Bytes touched at f16 (2 B): full A + full B + full out (each operand
-    # streamed once through HBM by the composite plan).
-    bytes_total = (M * K + K * N + M * N) * 2
-    row = {
-        "M": M, "K": K, "N": N,
-        "variant": variant,
-        "flops": 2 * M * K * N,
-        "bytes_hbm": bytes_total,
-        "arith_intensity": (2 * M * K * N) / bytes_total,  # flops/byte
-        "tile_count_expected": _ceil(M, TILE_M) * _ceil(N, TILE_N) * _ceil(K, TILE_K),
-        "sim_wall_clock_s": round(wall, 3),
-        "engines": {},
-    }
-    for eng in ENGINES:
-        row["engines"][eng] = {
-            "occupancy_ns": _engine_occupancy_ns(op_log, eng),
-            "wall_ns":      _engine_wall_ns(op_log, eng),
-            "record_count": _engine_count(op_log, eng),
-        }
-    row["stages"] = {}
-    for stage in STAGES:
-        row["stages"][stage] = {
-            "occupancy_ns": _stage_occupancy_ns(op_log, stage),
-            "wall_ns":      _stage_wall_ns(op_log, stage),
-            "record_count": _stage_count(op_log, stage),
-        }
-    # Kernel-window wall-clock = max t_end - min t_start over PE engine records.
-    pe_records = [r for r in op_log
-                  if any(r.component_id.endswith("." + e) for e in ENGINES)]
-    if pe_records:
-        row["pe_window_ns"] = max(r.t_end for r in pe_records) \
-                              - min(r.t_start for r in pe_records)
-    else:
-        row["pe_window_ns"] = 0.0
-    stage_records = [r for r in op_log
-                     if r.params.get("stage_type") in STAGES]
-    if stage_records:
-        row["composite_window_ns"] = max(r.t_end for r in stage_records) \
-                                     - min(r.t_start for r in stage_records)
-    else:
-        row["composite_window_ns"] = 0.0
-    return row
-
-
-def _ceil(a: int, b: int) -> int:
-    return (a + b - 1) // b
+from kernbench.benches.milestone_1h_gemm import run_sweep
 
 
 def main() -> int:
-    shapes_env = os.environ.get("SWEEP_SHAPES")
-    raw = (shapes_env.split(",") if shapes_env else DEFAULT_SHAPES)
-    shapes: list[tuple[int, int, int]] = []
-    for s in raw:
-        s = s.strip()
-        if not s:
-            continue
-        if "x" in s.lower():
-            parts = s.lower().split("x")
-            shapes.append((int(parts[0]), int(parts[1]), int(parts[2])))
-        else:
-            v = int(s)
-            shapes.append((v, v, v))
-    topology = os.environ.get("SWEEP_TOPOLOGY", "topology.yaml")
-
-    rows = []
-    for M, K, N in shapes:
-        for variant in VARIANTS:
-            print(f"[sweep] M={M} K={K} N={N} variant={variant} ...", flush=True)
-            row = _run_one(M, K, N, topology, variant=variant)
-            rows.append(row)
-            eng_dma = row["engines"]["pe_dma"]
-            eng_gem = row["engines"]["pe_gemm"]
-            print(f"   tiles={row['tile_count_expected']:>6}  "
-                  f"pe_window={row['pe_window_ns']:8.1f}ns  "
-                  f"dma_occ={eng_dma['occupancy_ns']:9.1f}  "
-                  f"gemm_occ={eng_gem['occupancy_ns']:8.1f}  "
-                  f"(sim {row['sim_wall_clock_s']:.1f}s)")
-
-    OUT_PATH.parent.mkdir(parents=True, exist_ok=True)
-    OUT_PATH.write_text(json.dumps({
-        "tile_sizes": {"M": TILE_M, "K": TILE_K, "N": TILE_N},
-        "engines": ENGINES,
-        "stages": STAGES,
-        "variants": VARIANTS,
-        "rows": rows,
-    }, indent=2))
-    print(f"\n[sweep] wrote {OUT_PATH}")
+    run_sweep()
     return 0
 
 

scripts/paper/paper_gqa_latency.py

@@ -0,0 +1,101 @@
+"""Report harness: GQA end-to-end latency + op/engine breakdown.
+
+Isolated under ``scripts/paper/`` for the 1H codesign report only — it is
+NOT a registered bench and does not touch other people's benches. It
+reuses the existing GQA headline panels (the real GQA kernels wired in
+``milestone_gqa_headline``) but, unlike that milestone (which records only
+op-counts), it also harvests per-panel end-to-end latency and per-engine
+occupancy from ``result.engine.op_log``.
+
+Latency definition (same window convention as ``milestone_1h_gemm``):
+    end_to_end_ns = max(r.t_end) - min(r.t_start)  over all op_log records.
+
+Output: docs/report/1H-codesign-paper/figures/gqa_latency.json
+
+Run:
+    python scripts/paper/paper_gqa_latency.py
+"""
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+from kernbench.benches.milestone_gqa_headline import (
+    _PANEL_DISPATCH,
+    _PANELS,
+    _make_bench_fn,
+    _summarize_op_log,
+)
+
+_REPORT_DIR = Path(__file__).resolve().parents[2] / "docs" / "report" / "1H-codesign-paper"
+_FIG_DIR = _REPORT_DIR / "figures"
+_OUT_JSON = _FIG_DIR / "gqa_latency.json"
+
+# PE engine component suffixes whose occupancy we break out.
+_ENGINES = ("pe_gemm", "pe_math", "pe_dma", "pe_fetch_store", "pe_ipcq", "pe_cpu")
+
+
+def _occupancy_ns(op_log, suffix: str) -> float:
+    return sum(
+        r.t_end - r.t_start
+        for r in op_log
+        if r.component_id.endswith("." + suffix)
+    )
+
+
+def _end_to_end_ns(op_log) -> float:
+    if not op_log:
+        return 0.0
+    return max(r.t_end for r in op_log) - min(r.t_start for r in op_log)
+
+
+def _run_panel(panel: str, topology: str) -> dict:
+    from kernbench.runtime_api.bench_runner import run_bench
+    from kernbench.runtime_api.types import resolve_device
+    from kernbench.sim_engine.engine import GraphEngine
+    from kernbench.topology.builder import resolve_topology
+
+    topo = resolve_topology(topology)
+    result = run_bench(
+        topology=topo,
+        bench_fn=_make_bench_fn(panel),
+        device=resolve_device(None),
+        engine_factory=lambda t, d: GraphEngine(
+            getattr(t, "topology_obj", t), enable_data=True,
+        ),
+    )
+    if not result.completion.ok:
+        raise RuntimeError(f"panel {panel!r} failed: {result.completion}")
+
+    op_log = result.engine.op_log
+    kind, params = _PANEL_DISPATCH[panel]
+    return {
+        "panel": panel,
+        "kind": kind,
+        **params,
+        "latency_ns": _end_to_end_ns(op_log),
+        "op_log_summary": _summarize_op_log(op_log),
+        "engine_occupancy_ns": {
+            eng: _occupancy_ns(op_log, eng) for eng in _ENGINES
+        },
+    }
+
+
+def main() -> None:
+    topology = "topology.yaml"
+    rows = [_run_panel(panel, topology) for panel in _PANELS]
+    _FIG_DIR.mkdir(parents=True, exist_ok=True)
+    out = {"version": 1, "panels": list(_PANELS), "rows": rows}
+    _OUT_JSON.write_text(json.dumps(out, indent=2))
+    print(f"wrote {_OUT_JSON}")
+    for r in rows:
+        s = r["op_log_summary"]
+        print(
+            f"  {r['panel']:24s} latency={r['latency_ns']:10.1f} ns  "
+            f"gemm={s['gemm_count']:3d} ipcq={s['ipcq_copy_count']:3d} "
+            f"dma_rd={s['dma_read_count']:3d} dma_wr={s['dma_write_count']:2d}"
+        )
+
+
+if __name__ == "__main__":
+    main()

scripts/paper/paper_latency_model_diagram.py

@@ -0,0 +1,535 @@
+#!/usr/bin/env python3
+"""Conceptual latency-model diagram for the KernBench paper (v5).
+
+Generic naming (Requester Node A/B, Router, Destination Node), with
+Router internals visible (two input ports -> switch -> output queue ->
+output port) so the queuing delay can be located *at the out port*
+rather than at the box edge.  The Destination Node shows queue -> drain
+slot -> processing logic.
+
+Highlights:
+  * Same-size Routers, each with explicit switching logic + output queue.
+  * Edge labels removed -- the wires speak for themselves.
+  * Annotation lines are strictly vertical; arrow heads use a larger
+    mutation_scale so no thin line protrudes past the tip; `shrinkA`
+    pulls the tail away from the text so they no longer overlap.
+  * "flit-level interleaving on wires" (not "shared FIFO").
+  * Destination: queue -> drain -> processing logic (no "served").
+
+Output:
+    docs/report/1H-codesign-paper/figures/latency_model.png
+
+Per /paper isolation: this is a report-only harness under scripts/paper/.
+"""
+
+import math
+from pathlib import Path
+
+import matplotlib.patches as patches
+import matplotlib.pyplot as plt
+
+OUT = Path(
+    "/Users/ywkang/kernbench/docs/report/1H-codesign-paper/figures/latency_model.png"
+)
+OUT.parent.mkdir(parents=True, exist_ok=True)
+
+
+# --- Colours --------------------------------------------------------------
+C_A      = "#E07B5E"
+C_B      = "#5E9BD1"
+C_NODE   = "#FFFFFF"
+C_BORD   = "#222222"
+C_WIRE   = "#222222"
+C_ANN    = "#333333"
+C_DIM    = "#777777"
+C_SWITCH = "#FAFAFA"
+C_OQUEUE = "#EFEFEF"
+C_PROC   = "#F4ECDC"
+C_QUEUE  = "#3CB371"   # medium-sea-green: distinctive vs A/B flit colours
+
+
+# --- Canvas ---------------------------------------------------------------
+fig = plt.figure(figsize=(17.0, 5.0))
+ax = fig.add_subplot(111)
+ax.set_xlim(0, 33)
+ax.set_ylim(2.2, 11.2)
+ax.axis("off")
+
+
+# --- Top header: end-to-end latency formula -----------------------------
+ax.text(
+    16.5, 10.6,
+    r"End-to-end latency  =  $\Sigma$ per-node overhead  +  "
+    r"$\Sigma$ per-edge transmission  +  drain  +  "
+    r"queuing delay",
+    ha="center", fontsize=12, color=C_ANN, weight="bold",
+)
+ax.plot([1.5, 31.5], [10.05, 10.05], color=C_DIM, lw=0.6)
+
+
+# --- Helpers --------------------------------------------------------------
+def box(cx, cy, w, h, label, fs=11, fweight="normal"):
+    ax.add_patch(patches.FancyBboxPatch(
+        (cx - w / 2, cy - h / 2), w, h,
+        boxstyle="round,pad=0.05",
+        facecolor=C_NODE, edgecolor=C_BORD, linewidth=1.6,
+    ))
+    if label:
+        ax.text(cx, cy, label, ha="center", va="center",
+                fontsize=fs, weight=fweight)
+
+
+def draw_flit_aligned(cx, cy, angle_rad, w, h, color, label):
+    """Draw a flit polygon rotated to align with the wire angle."""
+    cos_a, sin_a = math.cos(angle_rad), math.sin(angle_rad)
+    corners = []
+    for dx, dy in [(-w / 2, -h / 2), (w / 2, -h / 2),
+                   (w / 2, h / 2), (-w / 2, h / 2)]:
+        rx = dx * cos_a - dy * sin_a
+        ry = dx * sin_a + dy * cos_a
+        corners.append((cx + rx, cy + ry))
+    ax.add_patch(patches.Polygon(
+        corners, facecolor=color, edgecolor="black", linewidth=0.5,
+    ))
+    ax.text(cx, cy, label, ha="center", va="center",
+            fontsize=9, color="white", weight="bold")
+
+
+def draw_wire_with_flits(x0, y0, x1, y1, n_flits, label_char,
+                         color=None, flit_w=0.66, flit_h=0.95, gap=0.10):
+    """Wire (line) + endpoint arrowhead + flits centred on the line.
+
+    Wires are drawn in the neutral C_WIRE colour: the wire itself is
+    the place where the *transmission* delay accumulates (flit_size /
+    BW), not where flits queue. Queuing happens before the wire, at
+    the FIFO at the egress side -- coloured separately.
+    """
+    head_back = 0.30
+    dx, dy = x1 - x0, y1 - y0
+    wire_len = math.hypot(dx, dy)
+    ux, uy = dx / wire_len, dy / wire_len
+    x_line_end = x1 - ux * head_back
+    y_line_end = y1 - uy * head_back
+    ax.plot([x0, x_line_end], [y0, y_line_end],
+            color=C_WIRE, lw=1.6, zorder=1)
+    head_len, head_half = 0.30, 0.16
+    bx = x1 - ux * head_len
+    by = y1 - uy * head_len
+    px, py = -uy, ux
+    tri = [
+        (x1, y1),
+        (bx + px * head_half, by + py * head_half),
+        (bx - px * head_half, by - py * head_half),
+    ]
+    ax.add_patch(patches.Polygon(tri, facecolor=C_WIRE,
+                                  edgecolor=C_WIRE, linewidth=0.0))
+    # Flit train
+    angle_rad = math.atan2(dy, dx)
+    train_len = n_flits * flit_w + (n_flits - 1) * gap
+    centre_p = 0.5
+    start_p  = centre_p - (train_len / 2) / wire_len
+    step_p   = (flit_w + gap) / wire_len
+    if isinstance(label_char, str):
+        labels = [label_char] * n_flits
+        cols   = [color] * n_flits
+    else:
+        labels = [c[0] for c in label_char]
+        cols   = [c[1] for c in label_char]
+    for i in range(n_flits):
+        p = start_p + (i + 0.5) * step_p
+        cx = x0 + p * dx
+        cy = y0 + p * dy
+        draw_flit_aligned(cx, cy, angle_rad,
+                          flit_w, flit_h, cols[i], labels[i])
+
+
+def draw_router(cx, cy, w, h, two_inputs=True):
+    """Same-size Router with explicit internal switching logic and an
+    output queue.  Returns the wire-attachment (x,y) for each input
+    port and the single output port:
+      ((in1_x, in1_y), (in2_x, in2_y) or None, (out_x, out_y))
+    """
+    # Outer box
+    box(cx, cy, w, h, "")
+    ax.text(cx, cy + h / 2 - 0.32, "Router",
+            ha="center", fontsize=10, weight="bold")
+
+    # Input ports (small circles on the left edge)
+    in_x = cx - w / 2
+    in_x_internal = in_x + 0.25
+    if two_inputs:
+        in_y_top = cy + 0.55
+        in_y_bot = cy - 0.55
+        for iy in (in_y_top, in_y_bot):
+            ax.add_patch(patches.Circle(
+                (in_x_internal, iy), 0.12,
+                facecolor="white", edgecolor=C_BORD, linewidth=0.9,
+            ))
+    else:
+        in_y_top = None
+        in_y_bot = cy
+        ax.add_patch(patches.Circle(
+            (in_x_internal, in_y_bot), 0.12,
+            facecolor="white", edgecolor=C_BORD, linewidth=0.9,
+        ))
+
+    # Switch (small box, centre-left-ish). sw_cy = cy so that the
+    # output queue and (single-input) input port are at the same y as
+    # the router centre -- this keeps all router-to-router edges
+    # strictly horizontal.
+    sw_w, sw_h = 0.85, 1.10
+    sw_cx = cx - 0.55
+    sw_cy = cy
+    # The switch is the router's processing logic -- colour it the
+    # same as the Destination Node's processing-logic block so the
+    # two read as the same "processing" concept.
+    ax.add_patch(patches.Rectangle(
+        (sw_cx - sw_w / 2, sw_cy - sw_h / 2), sw_w, sw_h,
+        facecolor=C_PROC, edgecolor=C_BORD, linewidth=0.8,
+    ))
+    ax.text(sw_cx, sw_cy, "switch", ha="center", va="center",
+            fontsize=7.5, style="italic")
+
+    # Short feeder lines (no arrowhead) from input ports to the
+    # switch.  We deliberately drop arrowheads here: the head + port
+    # circle were too small at this scale and read as an overlap.
+    sw_left_x = sw_cx - sw_w / 2
+    if two_inputs:
+        for iy in (in_y_top, in_y_bot):
+            ax.plot(
+                [in_x_internal + 0.12, sw_left_x],
+                [iy, sw_cy + 0.25 * ((iy - cy) / 0.55)],
+                color=C_DIM, lw=0.7, zorder=1,
+            )
+    else:
+        ax.plot(
+            [in_x_internal + 0.12, sw_left_x],
+            [in_y_bot, sw_cy],
+            color=C_DIM, lw=0.7, zorder=1,
+        )
+
+    # Output queue (small queue holding a couple of flits). This *is*
+    # a real queueing location, so its fill takes the C_QUEUE family.
+    oq_w, oq_h = 0.95, 0.55
+    oq_cx = cx + 0.55
+    oq_cy = sw_cy
+    ax.add_patch(patches.Rectangle(
+        (oq_cx - oq_w / 2, oq_cy - oq_h / 2), oq_w, oq_h,
+        facecolor="#D9F0E1", edgecolor=C_QUEUE, linewidth=0.9,
+    ))
+    ax.text(oq_cx, oq_cy + oq_h / 2 + 0.18, "FIFO",
+            ha="center", fontsize=7, color=C_DIM, style="italic")
+    # Two small flits inside (A and B) hinting at the in-flight contents
+    mini_w, mini_h = 0.22, 0.34
+    for i, (col, lab) in enumerate([(C_A, "A"), (C_B, "B")]):
+        mx = oq_cx - 0.30 + i * (mini_w + 0.06)
+        ax.add_patch(patches.Rectangle(
+            (mx, oq_cy - mini_h / 2), mini_w, mini_h,
+            facecolor=col, edgecolor="black", linewidth=0.3,
+        ))
+        ax.text(mx + mini_w / 2, oq_cy, lab,
+                ha="center", va="center",
+                fontsize=5.5, color="white", weight="bold")
+
+    # Short feeder line (no arrowhead) from switch into out queue
+    ax.plot(
+        [sw_cx + sw_w / 2, oq_cx - oq_w / 2],
+        [sw_cy, oq_cy],
+        color=C_DIM, lw=0.7, zorder=1,
+    )
+    # Short feeder line from out queue to out port
+    out_x_internal = cx + w / 2 - 0.25
+    ax.plot(
+        [oq_cx + oq_w / 2, out_x_internal - 0.12],
+        [oq_cy, oq_cy],
+        color=C_DIM, lw=0.7, zorder=1,
+    )
+    # Output port circle on the right edge (a port marker, not a
+    # queue -- the queue is on the wire that follows, not the port).
+    ax.add_patch(patches.Circle(
+        (out_x_internal, oq_cy), 0.12,
+        facecolor="white", edgecolor=C_BORD, linewidth=0.9,
+    ))
+
+    return (
+        (in_x_internal, in_y_top) if two_inputs else None,
+        (in_x_internal, in_y_bot),
+        (out_x_internal, oq_cy),
+        oq_cx,   # also return the out-queue centre so callouts can target it
+        oq_cy,
+    )
+
+
+# --- Box layout ---------------------------------------------------------
+# Requester centres set to match Router 1's input-port y so that
+# Edge 1A and Edge 1B are strictly horizontal. Box heights are kept
+# small enough that a visible gap separates the two Requester boxes.
+R1_TMP_CY = 7.0
+ReqA = (2.8, R1_TMP_CY + 0.55)   # = 7.55, matches in_y_top of R1
+ReqB = (2.8, R1_TMP_CY - 0.55)   # = 6.45, matches in_y_bot of R1
+box(*ReqA, 3.0, 0.85, "Requester\nNode A", fs=10)
+box(*ReqB, 3.0, 0.85, "Requester\nNode B", fs=10)
+
+R_W, R_H = 3.0, 2.8
+R1 = (10.5, 7.0)
+R2 = (20.0, 7.0)
+
+r1_in_top, r1_in_bot, r1_out, r1_oq_cx, r1_oq_cy = draw_router(
+    *R1, R_W, R_H, two_inputs=True,
+)
+_, r2_in_bot, r2_out, r2_oq_cx, r2_oq_cy = draw_router(
+    *R2, R_W, R_H, two_inputs=False,
+)
+
+# Destination Node (same height as Router; wide enough so queue +
+# drain + processing-logic all fit on a single horizontal row).
+Dst = (28.4, 7.0)
+Dst_W, Dst_H = 6.0, R_H   # match the Router height
+box(*Dst, Dst_W, Dst_H, "")
+ax.text(Dst[0], Dst[1] + Dst_H / 2 - 0.25, "Destination Node",
+        ha="center", fontsize=10, weight="bold")
+
+
+# --- Edges: requester -> router 1 (Edge 1A & 1B, with flits) ------------
+draw_wire_with_flits(
+    ReqA[0] + 1.6, ReqA[1],
+    r1_in_top[0] - 0.02, r1_in_top[1],
+    n_flits=4, color=C_A, label_char="A",
+)
+draw_wire_with_flits(
+    ReqB[0] + 1.6, ReqB[1],
+    r1_in_bot[0] - 0.02, r1_in_bot[1],
+    n_flits=4, color=C_B, label_char="B",
+)
+
+# Edge 2: router1 out -> router2 in (horizontal, interleaved)
+labels_e2 = [("A", C_A) if i % 2 == 0 else ("B", C_B) for i in range(8)]
+E2_y = r1_out[1]
+draw_wire_with_flits(
+    r1_out[0] + 0.02, r1_out[1],
+    r2_in_bot[0] - 0.02, r2_in_bot[1],
+    n_flits=8, label_char=labels_e2,
+)
+
+# Edge 3: router2 out -> destination (horizontal, interleaved)
+labels_e3 = [("A", C_A) if i % 2 == 0 else ("B", C_B) for i in range(4)]
+draw_wire_with_flits(
+    r2_out[0] + 0.02, r2_out[1],
+    Dst[0] - Dst_W / 2, r2_out[1],
+    n_flits=4, label_char=labels_e3,
+)
+
+
+# --- Destination internals: queue -> drain -> processing logic ----------
+# Light-green halo behind the queue area marks it as a queueing point.
+C_QUEUE_BG = "#D9F0E1"
+
+dy_main = Dst[1] - 0.10
+Dst_left   = Dst[0] - Dst_W / 2
+
+# Queue (4 flits, FIFO; rightmost is the next to drain, so the
+# order is set so the serve sequence after the in-flight A alternates
+# A (in drain) -> B -> A -> B -> A -- giving a clean BABA queue.
+qW, qH, qGap = 0.42, 0.62, 0.07
+q_labels = ["A", "B", "A", "B"]
+q_cols   = [C_A,  C_B,  C_A,  C_B]
+q_total  = len(q_labels) * qW + (len(q_labels) - 1) * qGap
+q_x_start = Dst_left + 0.35
+
+# Green halo behind the queue boxes (the queue itself is a queueing
+# location -- mark it with the C_QUEUE colour family)
+ax.add_patch(patches.FancyBboxPatch(
+    (q_x_start - 0.10, dy_main - qH / 2 - 0.08),
+    q_total + 0.20, qH + 0.16,
+    boxstyle="round,pad=0.02",
+    facecolor=C_QUEUE_BG, edgecolor=C_QUEUE,
+    linewidth=0.9, zorder=1.5,
+))
+
+for i, (lab, col) in enumerate(zip(q_labels, q_cols)):
+    qx = q_x_start + i * (qW + qGap)
+    ax.add_patch(patches.Rectangle(
+        (qx, dy_main - qH / 2), qW, qH,
+        facecolor=col, edgecolor="black", linewidth=0.4,
+        zorder=2,
+    ))
+    ax.text(qx + qW / 2, dy_main, lab,
+            ha="center", va="center",
+            fontsize=8, color="white", weight="bold", zorder=3)
+# Common y for the "queue" / "drain" labels. Matches the FIFO label
+# spacing inside the Router (0.18 above the box top) and uses the same
+# font size for visual consistency.
+DST_LABEL_Y = dy_main + qH / 2 + 0.18
+ax.text(q_x_start + q_total / 2, DST_LABEL_Y,
+        "queue",
+        ha="center", fontsize=7, style="italic", color=C_DIM)
+
+# Drain slot (height matched to queue boxes; font matched to FIFO/queue
+# labels for visual consistency)
+dr_W, dr_H = 0.85, qH
+dr_x = q_x_start + q_total + 0.35
+dr_cx = dr_x + dr_W / 2
+dr_cy = dy_main
+ax.add_patch(patches.FancyBboxPatch(
+    (dr_x, dr_cy - dr_H / 2), dr_W, dr_H,
+    boxstyle="round,pad=0.03",
+    facecolor=C_A, edgecolor="black", linewidth=1.0,
+))
+ax.text(dr_cx, dr_cy, "A", ha="center", va="center",
+        fontsize=9, color="white", weight="bold")
+ax.text(dr_cx, DST_LABEL_Y,
+        "drain",
+        ha="center", fontsize=7, style="italic", color=C_DIM)
+
+# Drain-time double-arrow underneath the drain slot. Move the "drain
+# time" text a touch further down so the descender does not collide
+# with the arrowhead.
+dt_y = dr_cy - dr_H / 2 - 0.28
+ax.annotate(
+    "", xy=(dr_x + dr_W, dt_y), xytext=(dr_x, dt_y),
+    arrowprops=dict(arrowstyle="<->", lw=0.9, color=C_ANN),
+)
+ax.text(dr_cx, dt_y - 0.42, "drain time",
+        ha="center", fontsize=8.5, color=C_ANN, style="italic")
+
+# Processing-logic block (after drain) — wider so the two-line text
+# does not collide with the box edges
+pl_W, pl_H = 1.75, 1.00
+pl_x = dr_x + dr_W + 0.40
+pl_cx = pl_x + pl_W / 2
+pl_cy = dr_cy
+ax.add_patch(patches.FancyBboxPatch(
+    (pl_x, pl_cy - pl_H / 2), pl_W, pl_H,
+    boxstyle="round,pad=0.03",
+    facecolor=C_PROC, edgecolor=C_BORD, linewidth=1.0,
+))
+ax.text(pl_cx, pl_cy, "processing\nlogic",
+        ha="center", va="center", fontsize=9)
+
+# Small gray arrows along the queue -> drain -> processing chain
+ax.annotate(
+    "", xy=(dr_x - 0.04, dr_cy),
+    xytext=(q_x_start + q_total + 0.13, dr_cy),
+    arrowprops=dict(arrowstyle="-|>", lw=0.7,
+                    color=C_DIM, mutation_scale=7),
+)
+ax.annotate(
+    "", xy=(pl_x - 0.04, dr_cy),
+    xytext=(dr_x + dr_W + 0.04, dr_cy),
+    arrowprops=dict(arrowstyle="-|>", lw=0.7,
+                    color=C_DIM, mutation_scale=7),
+)
+
+
+# --- Annotations (strictly vertical, no diagonals, no overlap) ---------
+
+def vcallout(text, xy, x_text_top_y, fs=10, clearance=0.55,
+             marker_color=None):
+    """Vertical dotted callout. Optional `marker_color` paints a small
+    coloured circle just left of the text -- used to tie the label to
+    a colour code in the diagram.
+    """
+    text_y = x_text_top_y
+    if text_y < xy[1]:
+        line_top_y = text_y + clearance
+    else:
+        line_top_y = text_y - clearance
+    ax.plot([xy[0], xy[0]], [xy[1], line_top_y],
+            color=C_ANN, lw=0.9, linestyle=":", zorder=2)
+    if marker_color is not None:
+        # Text-width estimate at fs=10 (~0.18 data units per char) so
+        # the marker is placed clearly to the left of the text.
+        text_w = len(text) * 0.18
+        marker_x = xy[0] - text_w / 2 - 0.30
+        ax.add_patch(patches.Circle(
+            (marker_x, text_y), 0.14,
+            facecolor=marker_color, edgecolor=C_BORD, linewidth=0.6,
+            zorder=3,
+        ))
+    ax.text(xy[0], text_y, text,
+            ha="center", va="center",
+            fontsize=fs, color=C_ANN, zorder=3)
+
+
+# Transmission delay -> flit on Edge 2 (text BELOW the wire).
+# Extra clearance ~ 2 x text height so the line stops well clear of
+# the label.
+e2_total = 8 * 0.66 + 7 * 0.10
+e2_centre = (r1_out[0] + r2_in_bot[0]) / 2
+e2_start = e2_centre - e2_total / 2
+trans_idx = 5
+trans_cx  = e2_start + (trans_idx + 0.5) * (0.66 + 0.10)
+vcallout(
+    "transmission delay  =  flit_size / BW",
+    xy=(trans_cx, E2_y - 0.55),
+    x_text_top_y=E2_y - 2.6,
+    clearance=0.45,
+)
+
+# Flit-level interleaving -> ABOVE the wire, anchored on a flit near
+# the centre/right of the wire. Text y matches the queuing-delay
+# callout so the two labels sit on the same horizontal row, separated
+# horizontally so they don't collide.
+int_idx = 4
+int_cx  = e2_start + (int_idx + 0.5) * (0.66 + 0.10)
+vcallout(
+    "flit-level interleaving on wires",
+    xy=(int_cx, E2_y + 0.55),
+    x_text_top_y=R1[1] + R_H / 2 + 0.85,
+)
+
+# Queuing delay -> at Router 1 out-queue. Text positioned just above
+# the Router-1 box.
+vcallout(
+    "queuing delay",
+    xy=(r1_oq_cx, r1_oq_cy + 0.30),
+    x_text_top_y=R1[1] + R_H / 2 + 0.85,
+    marker_color=C_QUEUE,
+)
+
+# Drain -> below the drain slot (text is below).
+# Extra clearance ~ 2.5 x text height so the line stops further from
+# the label.
+vcallout(
+    "drain  =  per-flit service occupancy",
+    xy=(dr_cx, dt_y - 0.66),   # just *inside* the Destination Node
+                                  # bottom edge -- the dotted line then
+                                  # penetrates the box slightly rather
+                                  # than hanging below it
+    x_text_top_y=E2_y - 2.6,
+    clearance=0.60,
+)
+
+# Per-node overhead -> points at Router 1's switch block (where the
+# component's fixed processing cost lives). Text on the same row as
+# transmission delay and drain so the three "below the wire" callouts
+# sit on one horizontal baseline.
+R1_SW_CX = R1[0] - 0.55         # sw_cx for Router 1
+R1_SW_BOTTOM_Y = R1[1] - 0.55   # bottom of sw box
+vcallout(
+    "per-node overhead",
+    xy=(R1_SW_CX, R1_SW_BOTTOM_Y),
+    x_text_top_y=E2_y - 2.6,
+    clearance=0.60,
+    marker_color=C_PROC,
+)
+
+
+# --- Legend (close to the diagram) --------------------------------------
+LX, LY = 1.0, 2.6
+ax.add_patch(patches.Rectangle((LX, LY), 0.7, 0.45,
+                                facecolor=C_A, edgecolor="black",
+                                linewidth=0.4))
+ax.text(LX + 0.95, LY + 0.22, "Transaction A flit",
+        va="center", fontsize=10)
+
+ax.add_patch(patches.Rectangle((LX + 4.8, LY), 0.7, 0.45,
+                                facecolor=C_B, edgecolor="black",
+                                linewidth=0.4))
+ax.text(LX + 5.75, LY + 0.22, "Transaction B flit",
+        va="center", fontsize=10)
+
+
+# --- Save ----------------------------------------------------------------
+fig.savefig(OUT, dpi=140, bbox_inches="tight", facecolor="white")
+print(f"Wrote {OUT}")

scripts/paper/paper_plot_gqa.py

@@ -0,0 +1,106 @@
+"""Report harness: render GQA figures from ``gqa_latency.json``.
+
+Isolated under ``scripts/paper/`` for the 1H codesign report only.
+Reads the JSON emitted by ``paper_gqa_latency.py`` and writes two PNGs into
+``docs/report/1H-codesign-paper/figures/``:
+
+  gqa_latency_by_panel.png      end-to-end latency per panel
+  gqa_op_engine_breakdown.png   op-counts + engine occupancy per panel
+
+Run (after paper_gqa_latency.py):
+    python scripts/paper/paper_plot_gqa.py
+"""
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+import matplotlib
+
+matplotlib.use("Agg")
+import matplotlib.pyplot as plt  # noqa: E402
+
+_FIG_DIR = Path(__file__).resolve().parents[2] / "docs" / "report" / "1H-codesign-paper" / "figures"
+_IN_JSON = _FIG_DIR / "gqa_latency.json"
+
+_LABELS = {
+    "single_kv_group_prefill_gqa_c8_p8":
+        "prefill\nC=8, P=8\n(single KV group)",
+}
+
+
+def _load() -> list[dict]:
+    return json.loads(_IN_JSON.read_text())["rows"]
+
+
+def _plot_latency(rows: list[dict]) -> Path:
+    labels = [_LABELS[r["panel"]] for r in rows]
+    lat = [r["latency_ns"] for r in rows]
+    fig, ax = plt.subplots(figsize=(7.0, 4.0))
+    bars = ax.bar(labels, lat, color="#3b6ea5", width=0.6)
+    ax.set_ylabel("end-to-end latency (ns)")
+    ax.set_title("Fused GQA — end-to-end latency per panel")
+    ax.bar_label(bars, fmt="%.0f", padding=3, fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    ax.set_ylim(0, max(lat) * 1.15)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_latency_by_panel.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def _plot_breakdown(rows: list[dict]) -> Path:
+    labels = [_LABELS[r["panel"]] for r in rows]
+    x = range(len(rows))
+    fig, (ax1, ax2) = plt.subplots(1, 2, figsize=(11.0, 4.2))
+
+    # Left: op-count breakdown (grouped bars).
+    keys = ["gemm_count", "ipcq_copy_count", "dma_read_count", "dma_write_count"]
+    disp = ["GEMM", "IPCQ copy", "DMA read", "DMA write"]
+    colors = ["#3b6ea5", "#c0504d", "#9bbb59", "#8064a2"]
+    w = 0.2
+    for i, (k, d, c) in enumerate(zip(keys, disp, colors)):
+        vals = [r["op_log_summary"][k] for r in rows]
+        ax1.bar([xi + (i - 1.5) * w for xi in x], vals, width=w, label=d, color=c)
+    ax1.set_xticks(list(x))
+    ax1.set_xticklabels(labels, fontsize=8)
+    ax1.set_ylabel("op count")
+    ax1.set_title("Per-panel op-count breakdown")
+    ax1.legend(fontsize=8)
+    ax1.grid(axis="y", ls=":", alpha=0.5)
+
+    # Right: engine occupancy (log scale — compute vs data movement).
+    engs = ["pe_gemm", "pe_math", "pe_dma"]
+    edisp = ["GEMM engine", "MATH engine", "DMA engine"]
+    ecolors = ["#3b6ea5", "#e8a33d", "#c0504d"]
+    w2 = 0.25
+    for i, (e, d, c) in enumerate(zip(engs, edisp, ecolors)):
+        vals = [max(r["engine_occupancy_ns"][e], 0.1) for r in rows]
+        ax2.bar([xi + (i - 1) * w2 for xi in x], vals, width=w2, label=d, color=c)
+    ax2.set_yscale("log")
+    ax2.set_xticks(list(x))
+    ax2.set_xticklabels(labels, fontsize=8)
+    ax2.set_ylabel("summed engine occupancy (ns, log)")
+    ax2.set_title("Compute vs data-movement occupancy")
+    ax2.legend(fontsize=8)
+    ax2.grid(axis="y", ls=":", alpha=0.5)
+
+    fig.suptitle("Fused GQA — where the work goes", fontsize=12)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_op_engine_breakdown.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def main() -> None:
+    rows = _load()
+    p1 = _plot_latency(rows)
+    p2 = _plot_breakdown(rows)
+    print(f"wrote {p1}")
+    print(f"wrote {p2}")
+
+
+if __name__ == "__main__":
+    main()

scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py

@@ -0,0 +1,153 @@
+"""Comparative figures for milestone-gqa-decode-long-ctx-4cases.
+
+Reads sweep.json (emitted by ``kernbench run --bench
+milestone-gqa-decode-long-ctx-4cases``) and writes three PNGs into
+``docs/report/1H-codesign-paper/figures/``:
+
+  gqa_decode_long_ctx_4cases_latency.png   end-to-end latency per case
+  gqa_decode_long_ctx_4cases_traffic.png   ipcq/dma op-count breakdown
+  gqa_decode_long_ctx_4cases_memory.png    KV bytes per cube per case
+
+Run (after the bench):
+    GQA_DECODE_LONG_CTX_4CASES_RUN=1 python -m kernbench.cli.main run \\
+        --bench milestone-gqa-decode-long-ctx-4cases --topology topology.yaml
+    python scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py
+"""
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+import matplotlib
+
+matplotlib.use("Agg")
+import matplotlib.pyplot as plt  # noqa: E402
+
+_REPO_ROOT = Path(__file__).resolve().parents[2]
+_FIG_DIR = _REPO_ROOT / "docs" / "report" / "1H-codesign-paper" / "figures"
+_SWEEP_JSON = (
+    _REPO_ROOT / "src" / "kernbench" / "benches"
+    / "1H_milestone_output" / "gqa_decode_long_ctx_4cases" / "sweep.json"
+)
+
+# Panel name → (short label, case ordinal for left-to-right plot order).
+_CASE_INFO = {
+    "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp": (
+        "Case 1\nCube-SP × PE-TP", 1),
+    "single_kv_group_decode_long_ctx_gqa_cube_repl_pe_tp": (
+        "Case 2\nCube-Repl × PE-TP", 2),
+    "single_kv_group_decode_long_ctx_gqa_cube_repl_pe_sp": (
+        "Case 3\nCube-Repl × PE-SP", 3),
+    "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_sp": (
+        "Case 4 ★\nCube-SP × PE-SP", 4),
+}
+
+
+def _load() -> list[dict]:
+    return json.loads(_SWEEP_JSON.read_text())["rows"]
+
+
+def _sorted_by_case(rows: list[dict]) -> list[dict]:
+    return sorted(rows, key=lambda r: _CASE_INFO[r["panel"]][1])
+
+
+def _plot_latency(rows: list[dict]) -> Path:
+    rows = _sorted_by_case(rows)
+    labels = [_CASE_INFO[r["panel"]][0] for r in rows]
+    lat_us = [r["latency_ns"] / 1e3 for r in rows]
+    colors = ["#888", "#888", "#888", "#3b6ea5"]  # Case 4 highlighted
+    fig, ax = plt.subplots(figsize=(8.0, 4.5))
+    bars = ax.bar(labels, lat_us, color=colors, width=0.6)
+    ax.set_ylabel("end-to-end latency (µs)")
+    ax.set_title(
+        "Long-context decode 4-cases — end-to-end latency per case\n"
+        "LLaMA-3.1-70B single-KV-head group (8 cubes × 8 PEs)"
+    )
+    ax.bar_label(bars, fmt="%.1f", padding=3, fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    ax.set_ylim(0, max(lat_us) * 1.15)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_decode_long_ctx_4cases_latency.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def _plot_traffic(rows: list[dict]) -> Path:
+    rows = _sorted_by_case(rows)
+    labels = [_CASE_INFO[r["panel"]][0] for r in rows]
+    x = list(range(len(rows)))
+    keys = ["ipcq_copy_count", "dma_read_count", "dma_write_count"]
+    disp = ["IPCQ copy", "DMA read", "DMA write"]
+    colors = ["#c0504d", "#9bbb59", "#8064a2"]
+    w = 0.25
+    fig, ax = plt.subplots(figsize=(9.0, 4.5))
+    for i, (k, d, c) in enumerate(zip(keys, disp, colors)):
+        vals = [r["op_log_summary"][k] for r in rows]
+        ax.bar([xi + (i - 1) * w for xi in x], vals, width=w, label=d, color=c)
+    ax.set_xticks(list(x))
+    ax.set_xticklabels(labels, fontsize=9)
+    ax.set_ylabel("op count")
+    ax.set_title("Long-context decode 4-cases — op-count breakdown per case")
+    ax.legend(fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_decode_long_ctx_4cases_traffic.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def _kv_bytes_per_cube(panel: str, *, S_kv: int, h_kv: int,
+                       d_head: int, C: int) -> int:
+    """KV bytes a single cube's HBM holds (K + V together, f16)."""
+    # Cube-replicate ⇒ each cube holds full S_kv.
+    # Cube-SP        ⇒ each cube holds S_kv / C.
+    # PE replicate vs row_wise share the cube's HBM and don't change
+    # per-cube bytes.
+    S_per_cube = S_kv if "cube_repl" in panel else S_kv // C
+    return 2 * S_per_cube * h_kv * d_head * 2  # K + V, f16 (2 B/elem)
+
+
+def _plot_memory(rows: list[dict]) -> Path:
+    rows = _sorted_by_case(rows)
+    labels = [_CASE_INFO[r["panel"]][0] for r in rows]
+    mib_per_cube = [
+        _kv_bytes_per_cube(
+            r["panel"], S_kv=r["S_kv"], h_kv=r["h_kv"],
+            d_head=r["d_head"], C=r["C"],
+        ) / (1024 * 1024)
+        for r in rows
+    ]
+    # SP = blue (efficient); Repl = red (wasteful).
+    colors = ["#3b6ea5", "#c0504d", "#c0504d", "#3b6ea5"]
+    fig, ax = plt.subplots(figsize=(8.0, 4.5))
+    bars = ax.bar(labels, mib_per_cube, color=colors, width=0.6)
+    ax.set_ylabel("KV bytes per cube (MiB, K + V, f16)")
+    ax.set_title(
+        "Long-context decode 4-cases — KV memory per cube\n"
+        "(one KV-head group; per-layer, per-token state)"
+    )
+    ax.bar_label(bars, fmt="%.1f", padding=3, fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    ax.set_ylim(0, max(mib_per_cube) * 1.15)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_decode_long_ctx_4cases_memory.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def main() -> None:
+    rows = _load()
+    _FIG_DIR.mkdir(parents=True, exist_ok=True)
+    p1 = _plot_latency(rows)
+    p2 = _plot_traffic(rows)
+    p3 = _plot_memory(rows)
+    print(f"wrote {p1}")
+    print(f"wrote {p2}")
+    print(f"wrote {p3}")
+
+
+if __name__ == "__main__":
+    main()

src/kernbench/benches/1H_milestone_output/ccl/AllReduce_LRAB_2Dtorus_6SiP_2x3_with_TCM_SRAM_HBM.csv

@@ -0,0 +1,13 @@
+buffer_kind,sip_topology,n_sips,n_elem,bytes_per_pe,latency_ns
+hbm,torus_2d,6,128,256,2120.040000000012
+hbm,torus_2d,6,1024,2048,2717.2783333333473
+hbm,torus_2d,6,8192,16384,7315.184999999989
+hbm,torus_2d,6,32768,65536,23081.26500000037
+sram,torus_2d,6,128,256,2060.040000000012
+sram,torus_2d,6,1024,2048,2909.2783333333473
+sram,torus_2d,6,8192,16384,9523.184999999869
+sram,torus_2d,6,32768,65536,32201.265000000385
+tcm,torus_2d,6,128,256,1964.040000000012
+tcm,torus_2d,6,1024,2048,2477.2783333333473
+tcm,torus_2d,6,8192,16384,6403.185000000109
+tcm,torus_2d,6,32768,65536,19865.265000000378

src/kernbench/benches/1H_milestone_output/ccl/summary.csv

@@ -0,0 +1,37 @@
+algorithm,sip_topology,n_sips,n_elem,bytes_per_pe,bytes_per_sip,latency_ns
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8,16,256,2666.552500000015
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,2747.7400000000152
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,2855.990000000018
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,3072.490000000019
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,3337.1133333333582
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,3708.0333333333692
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,4449.873333333393
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,5933.020000000124
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,8900.379999999863
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,14835.099999999224
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,26704.540000000765
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,38573.97999999701
+lrab_hierarchical_allreduce,ring_1d,6,8,16,256,2365.255833333347
+lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,2436.9433333333473
+lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,2532.526666666683
+lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,2723.693333333349
+lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,3048.635000000021
+lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,3393.4016666666957
+lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,4082.401666666714
+lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,5458.80166666677
+lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,8216.934999999943
+lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,13733.201666665835
+lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,24765.73500000064
+lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,35798.268333331536
+lrab_hierarchical_allreduce,torus_2d,6,8,16,256,1700.6025000000095
+lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,1753.2900000000102
+lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,1823.540000000012
+lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,1964.040000000012
+lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,2196.8183333333463
+lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,2477.2783333333473
+lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,3038.1983333333583
+lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,4159.5050000000665
+lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,6403.185000000109
+lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,10890.5449999995
+lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,19865.265000000378
+lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,28839.98500000059

src/kernbench/benches/1H_milestone_output/gqa_headline/sweep.json

@@ -0,0 +1,31 @@
+{
+  "version": 1,
+  "panels": [
+    "single_kv_group_prefill_gqa_c8_p8"
+  ],
+  "config": {
+    "T_q_prefill": 4,
+    "T_q_decode": 1,
+    "S_kv_prefill": 16,
+    "h_q_decode": 8,
+    "h_kv_decode": 1,
+    "d_head": 64
+  },
+  "rows": [
+    {
+      "panel": "single_kv_group_prefill_gqa_c8_p8",
+      "kind": "prefill",
+      "C": 8,
+      "P": 8,
+      "T_q": 1024,
+      "S_kv": 1024,
+      "d_head": 128,
+      "op_log_summary": {
+        "gemm_count": 1024,
+        "ipcq_copy_count": 896,
+        "dma_read_count": 192,
+        "dma_write_count": 64
+      }
+    }
+  ]
+}

src/kernbench/benches/_gqa_attention_decode_long_ctx_cube_repl_pe_sp.py

@@ -0,0 +1,143 @@
+"""GQA decode kernel — Case 3 (Cube-Repl × PE-SP).
+
+Per GQA_full_deck.pptx slide 11:
+  - K, V replicated across all 8 cubes (the 8× memory waste).
+  - PEs SP on S_kv inside each cube: each PE attends to its
+    ``S_local = S_kv / P`` slice.
+  - Intra-CUBE 8-way reduce on the partial ``(m, ℓ, O)`` triple
+    (row chain + col bridge over the 2×4 PE grid, same structural
+    pattern as Case 4's intra-CUBE phase).
+  - NO inter-CUBE communication — every cube ends with the full
+    answer (8× redundant compute is the inherent cost of Case 3).
+  - Designated writer: only PE 0 of CUBE 0 stores ``O`` to avoid
+    8 redundant DMA writes.
+
+Tensor layout:
+  Q : (T_q, h_q · d_head) replicated on every rank.
+  K : (S_kv, h_kv · d_head) with cube=replicate, pe=row_wise —
+      each PE owns its (S_local, h_kv·d_head) shard within its cube.
+  V : same as K.
+  O : (T_q, h_q · d_head) — only PE 0 of CUBE 0 stores.
+
+Topology / SFR:
+  - ``configure_sfr_intercube_multisip`` provides the intra_* /
+    E/W/N/S namespaces; only the intra_* lanes are exercised here.
+"""
+from __future__ import annotations
+
+
+TILE_S_KV = 1024  # ADR-0063 §A.2 S_kv-axis tile sweep (per-tile width).
+
+
+def _merge_running(m_local, l_local, O_local, m_other, l_other, O_other, *, tl):
+    """Online-softmax merge of two partial ``(m, ℓ, O)`` triples."""
+    m_new = tl.maximum(m_local, m_other)
+    scale_old = tl.exp(m_local - m_new)
+    scale_new = tl.exp(m_other - m_new)
+    l_new = l_local * scale_old + l_other * scale_new
+    O_new = O_local * scale_old + O_other * scale_new
+    return m_new, l_new, O_new
+
+
+def gqa_attention_decode_long_ctx_cube_repl_pe_sp_kernel(
+    q_ptr: int,
+    k_ptr: int,
+    v_ptr: int,
+    o_ptr: int,
+    T_q: int,
+    S_kv: int,
+    h_q: int,
+    h_kv: int,
+    d_head: int,
+    C: int,
+    P: int,
+    *,
+    tl,
+) -> None:
+    """Case-3 decode: PE-SP attention; replicated KV per cube; intra-CUBE AR only."""
+    G = h_q // h_kv
+    S_local = S_kv // P  # each PE owns S_kv/P (cube=replicate, pe=row_wise)
+    pe_id = tl.program_id(axis=0)
+    cube_id = tl.program_id(axis=1)
+
+    # ── Local attention (S_kv-axis tile sweep, ADR-0063 §A.2) ──
+    Q = tl.load(q_ptr, shape=(G * T_q, d_head), dtype="f16")
+    n_tiles = (S_local + TILE_S_KV - 1) // TILE_S_KV
+    KV_ROW_BYTES = d_head * 2  # f16
+
+    tile_s0 = min(TILE_S_KV, S_local)
+    K_T = tl.load(k_ptr, shape=(d_head, tile_s0), dtype="f16")
+    V = tl.load(v_ptr, shape=(tile_s0, d_head), dtype="f16")
+    scores = tl.dot(Q, K_T)
+    m_local = tl.max(scores, axis=-1)
+    centered = scores - m_local
+    exp_scores = tl.exp(centered)
+    l_local = tl.sum(exp_scores, axis=-1)
+    O_local = tl.dot(exp_scores, V)
+
+    for tile_idx in range(1, n_tiles):
+        tile_start = tile_idx * TILE_S_KV
+        tile_s = min(TILE_S_KV, S_local - tile_start)
+        with tl.scratch_scope():
+            K_T_t = tl.load(k_ptr + tile_start * KV_ROW_BYTES,
+                            shape=(d_head, tile_s), dtype="f16")
+            V_t = tl.load(v_ptr + tile_start * KV_ROW_BYTES,
+                          shape=(tile_s, d_head), dtype="f16")
+            scores_t = tl.dot(Q, K_T_t)
+            m_tile = tl.max(scores_t, axis=-1)
+            centered_t = scores_t - m_tile
+            exp_scores_t = tl.exp(centered_t)
+            l_tile = tl.sum(exp_scores_t, axis=-1)
+            O_tile = tl.dot(exp_scores_t, V_t)
+            m_new, l_new, O_new = _merge_running(
+                m_local, l_local, O_local, m_tile, l_tile, O_tile, tl=tl,
+            )
+            tl.copy_to(m_local, m_new)
+            tl.copy_to(l_local, l_new)
+            tl.copy_to(O_local, O_new)
+
+    # ── Intra-CUBE 8-way reduce-to-PE0 (row chain + col bridge) ──
+    PE_GRID_COLS = 4
+    pe_col = pe_id % PE_GRID_COLS
+    pe_row = pe_id // PE_GRID_COLS
+    pe_cols_used = min(PE_GRID_COLS, P)
+    pe_rows_used = (P + PE_GRID_COLS - 1) // PE_GRID_COLS
+
+    if pe_cols_used > 1:
+        if pe_col < pe_cols_used - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="intra_E", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="intra_E", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="intra_E", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        if pe_col > 0:
+            tl.send(dir="intra_W", src=m_local)
+            tl.send(dir="intra_W", src=l_local)
+            tl.send(dir="intra_W", src=O_local)
+
+    if pe_col == 0 and pe_rows_used > 1:
+        if pe_row < pe_rows_used - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="intra_S", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="intra_S", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="intra_S", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        if pe_row > 0:
+            tl.send(dir="intra_N", src=m_local)
+            tl.send(dir="intra_N", src=l_local)
+            tl.send(dir="intra_N", src=O_local)
+
+    # ── Final normalise + store (designated writer: cube 0, PE 0) ──
+    if pe_id == 0 and cube_id == 0:
+        O_final = O_local / l_local
+        tl.store(o_ptr, O_final)

src/kernbench/benches/_gqa_attention_decode_long_ctx_cube_repl_pe_tp.py

@@ -0,0 +1,101 @@
+"""GQA decode kernel — Case 2 (Cube-Repl × PE-TP).
+
+Per GQA_full_deck.pptx slide 11:
+  - K, V replicated across all 8 cubes × 8 PEs (the 8 KB/tok/PE
+    memory waste — this is the inherent cost Case 2 demonstrates).
+  - PEs nominally split on the batch dim (PE-TP). For B=1 (single-
+    user decode, the slide-11 default), only PE 0 of CUBE 0 has
+    work; the other 63 ranks idle.
+  - NO inter-rank communication (each rank has full KV — slide 11
+    lists comm cost as "none").
+
+This kernel is the simplest of the 4 cases by design: one active
+rank does the full attention locally; everyone else early-returns.
+The DPPolicy at the call site models the cluster-wide 8× memory
+waste even though only one rank reads from HBM.
+
+Tensor layout (B=1):
+  Q : (T_q, h_q · d_head) replicated on every rank; loaded as
+      (G · T_q, d_head) on the active rank.
+  K : (S_kv, h_kv · d_head) replicated on every rank.
+  V : (S_kv, h_kv · d_head) replicated on every rank.
+  O : (T_q, h_q · d_head) — only PE 0 of CUBE 0 stores.
+
+Topology / SFR:
+  - configure_sfr_intercube_multisip is fine but not strictly required
+    (no inter-rank sends/recvs happen).
+"""
+from __future__ import annotations
+
+
+TILE_S_KV = 1024  # match decode_long — per-tile S_kv width (ADR-0063 §A.2).
+
+
+def gqa_attention_decode_long_ctx_cube_repl_pe_tp_kernel(
+    q_ptr: int,
+    k_ptr: int,
+    v_ptr: int,
+    o_ptr: int,
+    T_q: int,
+    S_kv: int,
+    h_q: int,
+    h_kv: int,
+    d_head: int,
+    C: int,
+    P: int,
+    *,
+    tl,
+) -> None:
+    """Case-2 decode: single-rank attention; full KV per rank; no comm."""
+    pe_id = tl.program_id(axis=0)
+    cube_id = tl.program_id(axis=1)
+    # B=1 single-user decode + PE-TP: only one rank has work.
+    # Slide-11 acknowledges this PE-TP waste at B=1.
+    if pe_id != 0 or cube_id != 0:
+        return
+
+    G = h_q // h_kv
+    n_tiles = (S_kv + TILE_S_KV - 1) // TILE_S_KV
+    KV_ROW_BYTES = d_head * 2  # f16
+
+    # ── Load Q (full; replicated; M-folded for GQA reuse) ──
+    Q = tl.load(q_ptr, shape=(G * T_q, d_head), dtype="f16")
+
+    # ── Tile 0: bootstrap persistent (m, ℓ, O) ──
+    tile_s0 = min(TILE_S_KV, S_kv)
+    K_T = tl.load(k_ptr, shape=(d_head, tile_s0), dtype="f16")
+    V = tl.load(v_ptr, shape=(tile_s0, d_head), dtype="f16")
+    scores = tl.dot(Q, K_T)
+    m_local = tl.max(scores, axis=-1)
+    centered = scores - m_local
+    exp_scores = tl.exp(centered)
+    l_local = tl.sum(exp_scores, axis=-1)
+    O_local = tl.dot(exp_scores, V)
+
+    # ── Tiles 1..N: fold via online-softmax merge in scratch_scope ──
+    for tile_idx in range(1, n_tiles):
+        tile_start = tile_idx * TILE_S_KV
+        tile_s = min(TILE_S_KV, S_kv - tile_start)
+        with tl.scratch_scope():
+            K_T_t = tl.load(k_ptr + tile_start * KV_ROW_BYTES,
+                            shape=(d_head, tile_s), dtype="f16")
+            V_t = tl.load(v_ptr + tile_start * KV_ROW_BYTES,
+                          shape=(tile_s, d_head), dtype="f16")
+            scores_t = tl.dot(Q, K_T_t)
+            m_tile = tl.max(scores_t, axis=-1)
+            centered_t = scores_t - m_tile
+            exp_scores_t = tl.exp(centered_t)
+            l_tile = tl.sum(exp_scores_t, axis=-1)
+            O_tile = tl.dot(exp_scores_t, V_t)
+            m_new = tl.maximum(m_local, m_tile)
+            scale_old = tl.exp(m_local - m_new)
+            scale_new = tl.exp(m_tile - m_new)
+            l_new = l_local * scale_old + l_tile * scale_new
+            O_new = O_local * scale_old + O_tile * scale_new
+            tl.copy_to(m_local, m_new)
+            tl.copy_to(l_local, l_new)
+            tl.copy_to(O_local, O_new)
+
+    # ── Final normalise + store ──
+    O_final = O_local / l_local
+    tl.store(o_ptr, O_final)

src/kernbench/benches/_gqa_attention_decode_long_ctx_cube_sp_pe_sp.py

@@ -0,0 +1,295 @@
+"""GQA decode kernel — Case 4 (Cube-SP × PE-SP). ★ slide-11 optimal.
+
+Per GQA_full_deck.pptx slide 11:
+  - K, V split 64-way (cube=row_wise, pe=row_wise); each rank owns
+    ``S_local = S_kv / (C·P)``.
+  - Local attention via per-rank S_kv-axis tile sweep (ADR-0063 §A.2)
+    keeps scratch bounded by ``TILE_S_KV``.
+  - Two-level reduce on the partial ``(m, ℓ, O)``:
+      • intra-CUBE 8-way (row chain along intra_W + col bridge along
+        intra_N) over the 2×4 PE grid;
+      • inter-CUBE 2-phase lrab over the 4×2 CUBE sub-mesh (ADR-0060
+        §4.2 lrab-adapted center-root reduce): Phase 1 row reduce
+        converges at root_col=2; Phase 2 col reduce on root_col
+        converges at root_row=1; root cube id = 6.
+  - PE 0 of CUBE 6 stores ``O``.
+
+Deviation from slide 13: slide prescribes AllReduce (every rank has
+the answer); the kernel does reduce-to-root (only cube 6 has it)
+per ADR-0060 §4. Treated as the kernbench Case-4 baseline.
+
+Tensor layout:
+  Q : (T_q, h_q · d_head) replicated; loaded as (G·T_q, d_head).
+  K : (S_kv, h_kv · d_head) sharded row_wise by (cube, pe); each rank
+      loads its (d_head, S_local) slice via byte-conserving reshape
+      of (S_local, h_kv·d_head) — correct for zero / symmetric
+      inputs (ADR-0060 §3 reshape-not-transpose caveat).
+  V : (S_kv, h_kv · d_head) sharded row_wise by (cube, pe); each rank
+      loads its (S_local, d_head) slice.
+  O : (T_q, h_q · d_head) — only PE 0 of CUBE 6 stores.
+
+Topology / SFR:
+  - Requires ``configure_sfr_intercube_multisip`` (provides disjoint
+    ``intra_*`` and ``E/W/N/S`` namespaces).
+  - Intra-CUBE PEs arranged as a 2×4 grid (no wrap).
+  - Inter-CUBE: 4×2 CUBE sub-mesh starting at origin (0, 0).
+"""
+from __future__ import annotations
+
+
+TILE_S_KV = 1024  # ADR-0063 §A.2 S_kv-axis tile sweep (per-tile width).
+
+# lrab geometry for the C=8 single-KV-head group (4×2 cube sub-mesh).
+_SUB_W = 4
+_SUB_H = 2
+_ROOT_COL = _SUB_W // 2   # 2
+_ROOT_ROW = _SUB_H // 2   # 1
+_ROOT_CUBE = _ROOT_ROW * _SUB_W + _ROOT_COL  # 6
+
+
+def _merge_running(m_local, l_local, O_local, m_other, l_other, O_other, *, tl):
+    """Online-softmax merge of two partial ``(m, ℓ, O)`` triples."""
+    m_new = tl.maximum(m_local, m_other)
+    scale_old = tl.exp(m_local - m_new)
+    scale_new = tl.exp(m_other - m_new)
+    l_new = l_local * scale_old + l_other * scale_new
+    O_new = O_local * scale_old + O_other * scale_new
+    return m_new, l_new, O_new
+
+
+def gqa_attention_decode_long_ctx_cube_sp_pe_sp_kernel(
+    q_ptr: int,
+    k_ptr: int,
+    v_ptr: int,
+    o_ptr: int,
+    T_q: int,
+    S_kv: int,
+    h_q: int,
+    h_kv: int,
+    d_head: int,
+    C: int,
+    P: int,
+    *,
+    tl,
+) -> None:
+    """Case-4 decode: Cube-SP × PE-SP, lrab center-root at cube 6."""
+    G = h_q // h_kv
+    n_ranks = C * P
+    S_local = S_kv // n_ranks
+    pe_id = tl.program_id(axis=0)
+    cube_id = tl.program_id(axis=1)
+
+    # ── Local attention (S_kv-axis tile sweep, ADR-0063 §A.2) ──
+    Q = tl.load(q_ptr, shape=(G * T_q, d_head), dtype="f16")
+    n_tiles = (S_local + TILE_S_KV - 1) // TILE_S_KV
+    KV_ROW_BYTES = d_head * 2  # f16
+
+    # Tile 0: establishes persistent (m_local, l_local, O_local). Cannot
+    # fold into Tiles 1..N loop (kernbench-only): persistent tensors
+    # must live outside tl.scratch_scope or scope teardown discards
+    # them; there's no scratch-backed (-inf, 0, 0) initializer.
+    tile_s0 = min(TILE_S_KV, S_local)
+    K_T = tl.load(k_ptr, shape=(d_head, tile_s0), dtype="f16")
+    V = tl.load(v_ptr, shape=(tile_s0, d_head), dtype="f16")
+    scores = tl.dot(Q, K_T)
+    m_local = tl.max(scores, axis=-1)
+    centered = scores - m_local
+    exp_scores = tl.exp(centered)
+    l_local = tl.sum(exp_scores, axis=-1)
+    O_local = tl.dot(exp_scores, V)
+
+    for tile_idx in range(1, n_tiles):
+        tile_start = tile_idx * TILE_S_KV
+        tile_s = min(TILE_S_KV, S_local - tile_start)
+        with tl.scratch_scope():
+            K_T_t = tl.load(k_ptr + tile_start * KV_ROW_BYTES,
+                            shape=(d_head, tile_s), dtype="f16")
+            V_t = tl.load(v_ptr + tile_start * KV_ROW_BYTES,
+                          shape=(tile_s, d_head), dtype="f16")
+            scores_t = tl.dot(Q, K_T_t)
+            m_tile = tl.max(scores_t, axis=-1)
+            centered_t = scores_t - m_tile
+            exp_scores_t = tl.exp(centered_t)
+            l_tile = tl.sum(exp_scores_t, axis=-1)
+            O_tile = tl.dot(exp_scores_t, V_t)
+            m_new, l_new, O_new = _merge_running(
+                m_local, l_local, O_local, m_tile, l_tile, O_tile, tl=tl,
+            )
+            tl.copy_to(m_local, m_new)
+            tl.copy_to(l_local, l_new)
+            tl.copy_to(O_local, O_new)
+
+    # ── Intra-CUBE reduce: row chain (intra_W) + col bridge (intra_N) ──
+    PE_GRID_COLS = 4
+    pe_col = pe_id % PE_GRID_COLS
+    pe_row = pe_id // PE_GRID_COLS
+    pe_cols_used = min(PE_GRID_COLS, P)
+    pe_rows_used = (P + PE_GRID_COLS - 1) // PE_GRID_COLS
+
+    if pe_cols_used > 1:
+        if pe_col < pe_cols_used - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="intra_E", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="intra_E", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="intra_E", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        if pe_col > 0:
+            tl.send(dir="intra_W", src=m_local)
+            tl.send(dir="intra_W", src=l_local)
+            tl.send(dir="intra_W", src=O_local)
+
+    if pe_col == 0 and pe_rows_used > 1:
+        if pe_row < pe_rows_used - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="intra_S", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="intra_S", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="intra_S", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        if pe_row > 0:
+            tl.send(dir="intra_N", src=m_local)
+            tl.send(dir="intra_N", src=l_local)
+            tl.send(dir="intra_N", src=O_local)
+
+    # ── Inter-CUBE lrab-adapted center-root reduce (ADR-0060 §4.2) ──
+    # Only PE 0 of each CUBE participates. Adapts Phases 1-2 of
+    # lrab_hierarchical_allreduce.py: bidirectional row reduce converges
+    # at root_col; bidirectional col reduce on root_col converges at
+    # root_row. Plain ``+`` replaced by log-sum-exp ``_merge_running``.
+    if pe_id == 0:
+        row = cube_id // _SUB_W
+        col = cube_id % _SUB_W
+
+        # Phase 1: row reduce — converge at col == _ROOT_COL.
+        if col == 0:
+            tl.send(dir="E", src=m_local)
+            tl.send(dir="E", src=l_local)
+            tl.send(dir="E", src=O_local)
+        elif 0 < col < _ROOT_COL:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="W", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="W", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="W", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+            tl.send(dir="E", src=m_local)
+            tl.send(dir="E", src=l_local)
+            tl.send(dir="E", src=O_local)
+        elif col == _ROOT_COL:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="W", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="W", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="W", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="E", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="E", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="E", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        elif _ROOT_COL < col < _SUB_W - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="E", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="E", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="E", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+            tl.send(dir="W", src=m_local)
+            tl.send(dir="W", src=l_local)
+            tl.send(dir="W", src=O_local)
+        elif col == _SUB_W - 1:
+            tl.send(dir="W", src=m_local)
+            tl.send(dir="W", src=l_local)
+            tl.send(dir="W", src=O_local)
+
+        # Phase 2: col reduce on col == _ROOT_COL — converge at row == _ROOT_ROW.
+        if col == _ROOT_COL:
+            if row == 0:
+                tl.send(dir="S", src=m_local)
+                tl.send(dir="S", src=l_local)
+                tl.send(dir="S", src=O_local)
+            elif 0 < row < _ROOT_ROW:
+                with tl.scratch_scope():
+                    m_other = tl.recv(dir="N", shape=m_local.shape, dtype="f16")
+                    l_other = tl.recv(dir="N", shape=l_local.shape, dtype="f16")
+                    O_other = tl.recv(dir="N", shape=O_local.shape, dtype="f16")
+                    m_new, l_new, O_new = _merge_running(
+                        m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                    )
+                    tl.copy_to(m_local, m_new)
+                    tl.copy_to(l_local, l_new)
+                    tl.copy_to(O_local, O_new)
+                tl.send(dir="S", src=m_local)
+                tl.send(dir="S", src=l_local)
+                tl.send(dir="S", src=O_local)
+            elif row == _ROOT_ROW:
+                with tl.scratch_scope():
+                    m_other = tl.recv(dir="N", shape=m_local.shape, dtype="f16")
+                    l_other = tl.recv(dir="N", shape=l_local.shape, dtype="f16")
+                    O_other = tl.recv(dir="N", shape=O_local.shape, dtype="f16")
+                    m_new, l_new, O_new = _merge_running(
+                        m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                    )
+                    tl.copy_to(m_local, m_new)
+                    tl.copy_to(l_local, l_new)
+                    tl.copy_to(O_local, O_new)
+                if _SUB_H - 1 > _ROOT_ROW:
+                    with tl.scratch_scope():
+                        m_other = tl.recv(dir="S", shape=m_local.shape, dtype="f16")
+                        l_other = tl.recv(dir="S", shape=l_local.shape, dtype="f16")
+                        O_other = tl.recv(dir="S", shape=O_local.shape, dtype="f16")
+                        m_new, l_new, O_new = _merge_running(
+                            m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                        )
+                        tl.copy_to(m_local, m_new)
+                        tl.copy_to(l_local, l_new)
+                        tl.copy_to(O_local, O_new)
+            elif _ROOT_ROW < row < _SUB_H - 1:
+                with tl.scratch_scope():
+                    m_other = tl.recv(dir="S", shape=m_local.shape, dtype="f16")
+                    l_other = tl.recv(dir="S", shape=l_local.shape, dtype="f16")
+                    O_other = tl.recv(dir="S", shape=O_local.shape, dtype="f16")
+                    m_new, l_new, O_new = _merge_running(
+                        m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                    )
+                    tl.copy_to(m_local, m_new)
+                    tl.copy_to(l_local, l_new)
+                    tl.copy_to(O_local, O_new)
+                tl.send(dir="N", src=m_local)
+                tl.send(dir="N", src=l_local)
+                tl.send(dir="N", src=O_local)
+            elif row == _SUB_H - 1 and _SUB_H - 1 > _ROOT_ROW:
+                tl.send(dir="N", src=m_local)
+                tl.send(dir="N", src=l_local)
+                tl.send(dir="N", src=O_local)
+
+    # ── Final normalise + store (root only: PE 0 of CUBE 6) ──
+    if pe_id == 0 and cube_id == _ROOT_CUBE:
+        O_final = O_local / l_local
+        tl.store(o_ptr, O_final)

src/kernbench/benches/_gqa_attention_decode_long_ctx_cube_sp_pe_tp.py

@@ -0,0 +1,246 @@
+"""GQA decode kernel — Case 1 (Cube-SP × PE-TP).
+
+Per GQA_full_deck.pptx slide 11:
+  - K, V split S_kv-wise across the 8 cubes (cube=row_wise);
+    replicated within each cube (pe=replicate). Each active rank
+    owns S_local = S_kv / C.
+  - PEs nominally split on the batch dim (PE-TP). At B=1 (the
+    slide-11 single-user default) only PE 0 of each cube has work;
+    PEs 1-7 of every cube idle (PE-TP-at-B=1 waste).
+  - NO intra-CUBE communication (only PE 0 holds a valid partial).
+  - Inter-CUBE 8-way reduce on (m, ℓ, O) via the lrab-adapted
+    center-root pattern (ADR-0060 §4.2): Phase 1 row reduce
+    converges at root_col; Phase 2 col reduce on root_col converges
+    at root_row. For sub_w=4, sub_h=2: root_col=2, root_row=1,
+    root_cube=6 (lrab geometric center).
+
+Tensor layout:
+  Q : (T_q, h_q · d_head) replicated on every rank.
+  K : (S_kv, h_kv · d_head) with cube=row_wise, pe=replicate —
+      each cube's PE 0 owns the full (S_local, h_kv·d_head) shard.
+  V : same as K.
+  O : (T_q, h_q · d_head) — only PE 0 of CUBE 6 stores.
+
+Topology / SFR:
+  - Requires ``configure_sfr_intercube_multisip`` for the E/W/N/S
+    inter-CUBE lanes used by the lrab reduce.
+"""
+from __future__ import annotations
+
+
+TILE_S_KV = 1024  # ADR-0063 §A.2 S_kv-axis tile sweep (per-tile width).
+
+
+def _merge_running(m_local, l_local, O_local, m_other, l_other, O_other, *, tl):
+    """Online-softmax merge of two partial ``(m, ℓ, O)`` triples."""
+    m_new = tl.maximum(m_local, m_other)
+    scale_old = tl.exp(m_local - m_new)
+    scale_new = tl.exp(m_other - m_new)
+    l_new = l_local * scale_old + l_other * scale_new
+    O_new = O_local * scale_old + O_other * scale_new
+    return m_new, l_new, O_new
+
+
+def gqa_attention_decode_long_ctx_cube_sp_pe_tp_kernel(
+    q_ptr: int,
+    k_ptr: int,
+    v_ptr: int,
+    o_ptr: int,
+    T_q: int,
+    S_kv: int,
+    h_q: int,
+    h_kv: int,
+    d_head: int,
+    C: int,
+    P: int,
+    *,
+    tl,
+) -> None:
+    """Case-1 decode: PE 0 per cube does attention; inter-CUBE lrab AR only."""
+    # B=1 single-user decode + PE-TP: only PE 0 per cube has work.
+    pe_id = tl.program_id(axis=0)
+    cube_id = tl.program_id(axis=1)
+    if pe_id != 0:
+        return
+
+    # Cube-SP geometry: 4×2 sub-mesh, lrab center-root cube = 6.
+    sub_w = 4
+    sub_h = C // sub_w
+    if sub_w < 2 or sub_h < 2 or sub_w * sub_h != C:
+        raise ValueError(
+            f"Case 1 requires C decomposable as sub_w=4, sub_h>=2; got C={C}"
+        )
+    root_col = sub_w // 2
+    root_row = sub_h // 2
+    root_cube = root_row * sub_w + root_col
+
+    G = h_q // h_kv
+    S_local = S_kv // C  # cube=row_wise, pe=replicate ⇒ S_local per cube
+    KV_ROW_BYTES = d_head * 2  # f16
+
+    # ── Local attention (S_kv-axis tile sweep, ADR-0063 §A.2) ──
+    Q = tl.load(q_ptr, shape=(G * T_q, d_head), dtype="f16")
+    n_tiles = (S_local + TILE_S_KV - 1) // TILE_S_KV
+
+    tile_s0 = min(TILE_S_KV, S_local)
+    K_T = tl.load(k_ptr, shape=(d_head, tile_s0), dtype="f16")
+    V = tl.load(v_ptr, shape=(tile_s0, d_head), dtype="f16")
+    scores = tl.dot(Q, K_T)
+    m_local = tl.max(scores, axis=-1)
+    centered = scores - m_local
+    exp_scores = tl.exp(centered)
+    l_local = tl.sum(exp_scores, axis=-1)
+    O_local = tl.dot(exp_scores, V)
+
+    for tile_idx in range(1, n_tiles):
+        tile_start = tile_idx * TILE_S_KV
+        tile_s = min(TILE_S_KV, S_local - tile_start)
+        with tl.scratch_scope():
+            K_T_t = tl.load(k_ptr + tile_start * KV_ROW_BYTES,
+                            shape=(d_head, tile_s), dtype="f16")
+            V_t = tl.load(v_ptr + tile_start * KV_ROW_BYTES,
+                          shape=(tile_s, d_head), dtype="f16")
+            scores_t = tl.dot(Q, K_T_t)
+            m_tile = tl.max(scores_t, axis=-1)
+            centered_t = scores_t - m_tile
+            exp_scores_t = tl.exp(centered_t)
+            l_tile = tl.sum(exp_scores_t, axis=-1)
+            O_tile = tl.dot(exp_scores_t, V_t)
+            m_new, l_new, O_new = _merge_running(
+                m_local, l_local, O_local, m_tile, l_tile, O_tile, tl=tl,
+            )
+            tl.copy_to(m_local, m_new)
+            tl.copy_to(l_local, l_new)
+            tl.copy_to(O_local, O_new)
+
+    # ── Inter-CUBE lrab-adapted center-root reduce (ADR-0060 §4.2) ──
+    row = cube_id // sub_w
+    col = cube_id % sub_w
+
+    # Phase 1: row reduce — converge at col == root_col.
+    if col == 0 and root_col > 0:
+        tl.send(dir="E", src=m_local)
+        tl.send(dir="E", src=l_local)
+        tl.send(dir="E", src=O_local)
+    elif 0 < col < root_col:
+        with tl.scratch_scope():
+            m_other = tl.recv(dir="W", shape=m_local.shape, dtype="f16")
+            l_other = tl.recv(dir="W", shape=l_local.shape, dtype="f16")
+            O_other = tl.recv(dir="W", shape=O_local.shape, dtype="f16")
+            m_new, l_new, O_new = _merge_running(
+                m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+            )
+            tl.copy_to(m_local, m_new)
+            tl.copy_to(l_local, l_new)
+            tl.copy_to(O_local, O_new)
+        tl.send(dir="E", src=m_local)
+        tl.send(dir="E", src=l_local)
+        tl.send(dir="E", src=O_local)
+    elif col == root_col:
+        if root_col > 0:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="W", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="W", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="W", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        if sub_w - 1 > root_col:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="E", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="E", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="E", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+    elif root_col < col < sub_w - 1:
+        with tl.scratch_scope():
+            m_other = tl.recv(dir="E", shape=m_local.shape, dtype="f16")
+            l_other = tl.recv(dir="E", shape=l_local.shape, dtype="f16")
+            O_other = tl.recv(dir="E", shape=O_local.shape, dtype="f16")
+            m_new, l_new, O_new = _merge_running(
+                m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+            )
+            tl.copy_to(m_local, m_new)
+            tl.copy_to(l_local, l_new)
+            tl.copy_to(O_local, O_new)
+        tl.send(dir="W", src=m_local)
+        tl.send(dir="W", src=l_local)
+        tl.send(dir="W", src=O_local)
+    elif col == sub_w - 1 and sub_w - 1 > root_col:
+        tl.send(dir="W", src=m_local)
+        tl.send(dir="W", src=l_local)
+        tl.send(dir="W", src=O_local)
+
+    # Phase 2: col reduce on col == root_col — converge at row == root_row.
+    if col == root_col:
+        if row == 0 and root_row > 0:
+            tl.send(dir="S", src=m_local)
+            tl.send(dir="S", src=l_local)
+            tl.send(dir="S", src=O_local)
+        elif 0 < row < root_row:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="N", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="N", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="N", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+            tl.send(dir="S", src=m_local)
+            tl.send(dir="S", src=l_local)
+            tl.send(dir="S", src=O_local)
+        elif row == root_row:
+            if root_row > 0:
+                with tl.scratch_scope():
+                    m_other = tl.recv(dir="N", shape=m_local.shape, dtype="f16")
+                    l_other = tl.recv(dir="N", shape=l_local.shape, dtype="f16")
+                    O_other = tl.recv(dir="N", shape=O_local.shape, dtype="f16")
+                    m_new, l_new, O_new = _merge_running(
+                        m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                    )
+                    tl.copy_to(m_local, m_new)
+                    tl.copy_to(l_local, l_new)
+                    tl.copy_to(O_local, O_new)
+            if sub_h - 1 > root_row:
+                with tl.scratch_scope():
+                    m_other = tl.recv(dir="S", shape=m_local.shape, dtype="f16")
+                    l_other = tl.recv(dir="S", shape=l_local.shape, dtype="f16")
+                    O_other = tl.recv(dir="S", shape=O_local.shape, dtype="f16")
+                    m_new, l_new, O_new = _merge_running(
+                        m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                    )
+                    tl.copy_to(m_local, m_new)
+                    tl.copy_to(l_local, l_new)
+                    tl.copy_to(O_local, O_new)
+        elif root_row < row < sub_h - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="S", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="S", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="S", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+            tl.send(dir="N", src=m_local)
+            tl.send(dir="N", src=l_local)
+            tl.send(dir="N", src=O_local)
+        elif row == sub_h - 1 and sub_h - 1 > root_row:
+            tl.send(dir="N", src=m_local)
+            tl.send(dir="N", src=l_local)
+            tl.send(dir="N", src=O_local)
+
+    # ── Final normalise + store (root only: PE 0 of cube 6) ──
+    if cube_id == root_cube:
+        O_final = O_local / l_local
+        tl.store(o_ptr, O_final)

src/kernbench/benches/_gqa_attention_decode_opt2.py

@@ -0,0 +1,85 @@
+"""GQA decode opt2 variant — two-composite inner attention (ADR-0065).
+
+opt2 of decode (ADR-0060 §8 item 4 / ADR-0065): the per-tile attention is
+expressed as **two composites** instead of opt3's long primitive sequence —
+
+  #1  Q·Kᵀ      → a GEMM composite writing the score tile ``scores``.
+  #2  softmax   → a ``softmax_merge`` recipe composite (online-softmax merge
+      + P·V        of ``(m, l, O)``) whose head GEMM is P·V, with an ``add``
+                   epilogue folding the P·V contribution into ``O``.
+
+Fewer PE_CPU-issued commands ⇒ lower dispatch cost under the ADR-0064 Rev2
+structural model (the CPU-offload win). Tile 0 establishes the running
+state with primitives (kernbench has no scratch-backed ``-inf`` initializer,
+same limitation as the opt3 kernel); the next tile exercises the recipe.
+
+**Scope (P5 B):** runnable in op_log mode (latency only). Full data-mode
+numeric parity (computing the recipe's 8 MATH ops in the DataExecutor) is a
+separate follow-up — see DDD-0065 / the P5-numerics note.
+"""
+from __future__ import annotations
+
+
+def gqa_attention_decode_opt2_kernel(
+    q_ptr: int,
+    k_ptr: int,
+    v_ptr: int,
+    o_ptr: int,
+    T_q: int,
+    S_kv: int,
+    h_q: int,
+    h_kv: int,
+    d_head: int,
+    C: int,
+    P: int,
+    *,
+    tl,
+) -> None:
+    """Single-rank (C=P=1) GQA decode using the opt2 two-composite form.
+
+    Layout mirrors ``gqa_attention_decode_long_ctx_cube_sp_pe_sp_kernel``
+    (M-fold Q, K loaded as ``(d_head, S_local)``). The KV slice is split
+    into two sub-tiles so the second one drives the ``softmax_merge``
+    recipe composite.
+    """
+    G = h_q // h_kv
+    S_local = S_kv // (C * P)
+    KV_ROW_BYTES = d_head * 2  # f16
+
+    Q = tl.load(q_ptr, shape=(G * T_q, d_head), dtype="f16")
+
+    # Split the local KV slice in two: tile 0 establishes the running state,
+    # tile 1 merges via the opt2 two-composite path.
+    half = max(1, S_local // 2)
+    rest = S_local - half
+
+    # ── Tile 0: establish running (m_local, l_local, O_local) ──
+    K_T0 = tl.load(k_ptr, shape=(d_head, half), dtype="f16")
+    V0 = tl.load(v_ptr, shape=(half, d_head), dtype="f16")
+    scores0 = tl.dot(Q, K_T0)
+    m_local = tl.max(scores0, axis=-1)
+    exp0 = tl.exp(scores0 - m_local)
+    l_local = tl.sum(exp0, axis=-1)
+    O_local = tl.dot(exp0, V0)
+
+    # ── Tile 1: opt2 two-composite merge (skipped when S_local == 1) ──
+    if rest > 0:
+        K_T1 = tl.load(k_ptr + half * KV_ROW_BYTES,
+                       shape=(d_head, rest), dtype="f16")
+        V1 = tl.ref(v_ptr + half * KV_ROW_BYTES, shape=(rest, d_head),
+                    dtype="f16")
+        # #1: Q·Kᵀ composite → score tile. No `out` → auto-allocated TCM
+        # scratch returned as a handle (ADR-0065 D8), consumed by #2.
+        scores1 = tl.composite(op="gemm", a=Q, b=K_T1)
+        # #2: softmax_merge recipe (online merge of (m,l,O)) + P·V + add.
+        tl.composite(
+            prologue=[{"op": "softmax_merge", "s": scores1,
+                       "m": m_local, "l": l_local, "O": O_local}],
+            op="gemm", b=V1, out=O_local,
+            epilogue=[{"op": "add", "other": O_local}],
+        )
+
+    # ── Final normalise + store (root only) ──
+    if tl.program_id(axis=0) == 0 and tl.program_id(axis=1) == 0:
+        O_final = O_local / l_local
+        tl.store(o_ptr, O_final)

src/kernbench/benches/_gqa_attention_decode_short.py

@@ -0,0 +1,165 @@
+"""GQA fused-attention decode kernel — short context (ADR-0060 §B.split.2).
+
+Short context (``S_kv < 256K``): each CUBE owns ``kv_per_cube`` whole
+KV heads, with no S_kv sharding across CUBEs and no inter-CUBE reduce.
+PE-SP within each CUBE: the ``P`` PEs split into ``kv_per_cube`` groups
+of ``P/kv_per_cube`` PEs each; each group does PE-SP across the group
+for one owned head, then the group's root PE stores its head's output.
+
+The local attention uses an S_kv-axis tile sweep (ADR-0063 §A.2) so
+per-rank scratch is bounded by ``TILE_S_KV``.
+
+Group layout on the 2×4 PE grid:
+  kv_per_cube=1, group=8 PEs (full 2×4): row chain + col bridge.
+  kv_per_cube=2, group=4 PEs (one row):  row chain only.
+  kv_per_cube=4, group=2 PEs (adj cols): 1-step chain.
+  kv_per_cube=8, group=1 PE:             no chain — direct write.
+
+Layout caveats:
+  - K, V: ``(h_kv·S_kv, d_head)`` head-stacked, deployed with
+    ``dp = (cube=row_wise, pe=row_wise)`` so each PE's chunk is
+    contiguously ``(S_local, d_head)`` at its own shard. K loaded as
+    ``(d_head, S_local)`` via byte-conserving reshape (ADR-0060 §3).
+  - Q: replicated ``(T_q, h_q·d_head)``, reshaped byte-conservingly to
+    ``(h_q·T_q, d_head)``. Kernel computes attention for ALL Q rows
+    against the group's owned K head; only the group's owned head rows
+    are semantically meaningful (correct for zero / symmetric inputs).
+  - O: replicated; each group root writes its head's
+    ``(h_q·T_q, d_head)`` result at disjoint PE-local addresses.
+  - GEMMs use ``tl.dot`` (no composite epilogue / ``softmax_scale``).
+"""
+from __future__ import annotations
+
+
+TILE_S_KV = 1024  # ADR-0063 §A.2 S_kv-axis tile sweep (per-tile width).
+
+
+def _merge_running(m_local, l_local, O_local, m_other, l_other, O_other, *, tl):
+    """Online-softmax merge of two partial ``(m, ℓ, O)`` triples."""
+    m_new = tl.maximum(m_local, m_other)
+    scale_old = tl.exp(m_local - m_new)
+    scale_new = tl.exp(m_other - m_new)
+    l_new = l_local * scale_old + l_other * scale_new
+    O_new = O_local * scale_old + O_other * scale_new
+    return m_new, l_new, O_new
+
+
+def gqa_attention_decode_short_kernel(
+    q_ptr: int,
+    k_ptr: int,
+    v_ptr: int,
+    o_ptr: int,
+    T_q: int,
+    S_kv: int,
+    h_q: int,
+    h_kv: int,
+    d_head: int,
+    C: int,
+    P: int,
+    kv_per_cube: int,
+    *,
+    tl,
+) -> None:
+    """Short-context GQA decode with PE-parallel heads + intra-group PE-SP."""
+    group_size = P // kv_per_cube
+    pe_id = tl.program_id(axis=0)
+    pe_in_group = pe_id % group_size
+    S_local = S_kv // group_size
+
+    # ── Local attention (S_kv-axis tile sweep, ADR-0063 §A.2) ──
+    Q = tl.load(q_ptr, shape=(h_q * T_q, d_head), dtype="f16")
+    n_tiles = (S_local + TILE_S_KV - 1) // TILE_S_KV
+    KV_ROW_BYTES = d_head * 2  # f16
+
+    # Tile 0: establishes persistent (m_local, l_local, O_local).
+    #
+    # Cannot be folded into the Tiles 1..N loop (kernbench-only limitation):
+    #   - persistent (m, ℓ, O) must live OUTSIDE ``tl.scratch_scope``,
+    #     otherwise scope teardown discards them before the next tile's
+    #     merge can read them;
+    #   - kernbench has no scratch-backed initializer — ``tl.zeros`` /
+    #     ``tl.full`` return addr=0 handles with no backing storage, so
+    #     they cannot be overwritten via ``tl.copy_to`` to seed (-inf, 0, 0).
+    # So Tile 0 computes the initial running state directly; Tiles 1..N
+    # fold into it. Triton port: limitation does not apply (SSA tensors
+    # stay live across iterations) — a single unified loop suffices.
+    tile_s0 = min(TILE_S_KV, S_local)
+    K_T = tl.load(k_ptr, shape=(d_head, tile_s0), dtype="f16")
+    V = tl.load(v_ptr, shape=(tile_s0, d_head), dtype="f16")
+    scores = tl.dot(Q, K_T)
+    m_local = tl.max(scores, axis=-1)
+    centered = scores - m_local
+    exp_scores = tl.exp(centered)
+    l_local = tl.sum(exp_scores, axis=-1)
+    O_local = tl.dot(exp_scores, V)
+
+    # Tiles 1..n_tiles-1: fold into running state via online-softmax merge.
+    # Triton port: drop the ``with tl.scratch_scope():`` line and replace
+    # each ``copy_to`` with a Python rebind.
+    for tile_idx in range(1, n_tiles):
+        tile_start = tile_idx * TILE_S_KV
+        tile_s = min(TILE_S_KV, S_local - tile_start)
+        with tl.scratch_scope():
+            K_T_t = tl.load(k_ptr + tile_start * KV_ROW_BYTES,
+                            shape=(d_head, tile_s), dtype="f16")
+            V_t = tl.load(v_ptr + tile_start * KV_ROW_BYTES,
+                          shape=(tile_s, d_head), dtype="f16")
+            scores_t = tl.dot(Q, K_T_t)
+            m_tile = tl.max(scores_t, axis=-1)
+            centered_t = scores_t - m_tile
+            exp_scores_t = tl.exp(centered_t)
+            l_tile = tl.sum(exp_scores_t, axis=-1)
+            O_tile = tl.dot(exp_scores_t, V_t)
+            m_new, l_new, O_new = _merge_running(
+                m_local, l_local, O_local, m_tile, l_tile, O_tile, tl=tl,
+            )
+            tl.copy_to(m_local, m_new)
+            tl.copy_to(l_local, l_new)
+            tl.copy_to(O_local, O_new)
+
+    # ── Communication: within-group chain reduce-to-root (Level-2 only) ──
+    group_cols = min(4, group_size)
+    group_rows = (group_size + group_cols - 1) // group_cols
+    pe_col_in_group = pe_in_group % group_cols
+    pe_row_in_group = pe_in_group // group_cols
+
+    # Row chain (within group's row, along intra_W, leftward).
+    if group_cols > 1:
+        if pe_col_in_group < group_cols - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="intra_E", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="intra_E", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="intra_E", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        if pe_col_in_group > 0:
+            tl.send(dir="intra_W", src=m_local)
+            tl.send(dir="intra_W", src=l_local)
+            tl.send(dir="intra_W", src=O_local)
+
+    # Col bridge (within group, along intra_N, row-1 → row-0).
+    if pe_col_in_group == 0 and group_rows > 1:
+        if pe_row_in_group < group_rows - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="intra_S", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="intra_S", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="intra_S", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        if pe_row_in_group > 0:
+            tl.send(dir="intra_N", src=m_local)
+            tl.send(dir="intra_N", src=l_local)
+            tl.send(dir="intra_N", src=O_local)
+
+    # ── Final normalise + store (group root only) ──
+    if pe_in_group == 0:
+        O_final = O_local / l_local
+        tl.store(o_ptr, O_final)

src/kernbench/benches/_gqa_attention_prefill_long.py

@@ -0,0 +1,190 @@
+"""GQA fused-attention prefill kernel — long context (ADR-0060 §5.5).
+
+Head-parallel: each CUBE owns one Q head and one KV slice. Over C ring
+steps the KV blocks rotate around the C CUBEs (W-send / E-recv) so every
+CUBE sees every block; the online-softmax merge folds each step into the
+running ``(m, ℓ, O)``. No inter-CUBE reduce — each CUBE writes its own
+head's output.
+
+The ring is **tile-granular** (ADR-0060 §5.5 + ADR-0063 §A.2): each
+ring step transmits ``n_tiles`` tiles of size ``TILE_S_KV`` rather than
+one full ``S_local`` slice. The kernel loop is nested over
+``(ring_step, tile_idx)``, with the bootstrap at ``(k=0, t=0)``
+establishing the persistent ``(m, ℓ, O)`` and every subsequent
+iteration folding its tile in inside a ``tl.scratch_scope``. Per-rank
+persistent scratch is ``(m, ℓ, O)`` only (~1 KB); per-tile in-scope
+scratch is bounded by ``TILE_S_KV`` regardless of ``S_local``.
+
+Topology / SFR:
+  - Requires ``configure_sfr_intercube_ring`` — either ``ring_size=C``
+    (1D row, ``C ≤ mesh_w``) or ``submesh_shape=(rows, cols)`` (snake
+    Hamiltonian cycle through a rectangular sub-mesh, for ``C > mesh_w``).
+  - ``P == 1`` (default): only PE 0 of each CUBE participates;
+    intra-CUBE PE parallelism is disabled.
+  - ``P > 1``: all P PEs of each CUBE participate via query-axis
+    split (ADR-0060 §5.5 last bullet + §B-item-3) — each PE owns
+    ``T_q/P`` disjoint query rows. Output rows are disjoint per PE,
+    so no intra-CUBE reduce is needed. Each PE drives its own
+    same-lane ring via independent IPCQ channels (P parallel rings).
+    Requires ``T_q % P == 0``.
+
+Layout caveats:
+  - GEMMs use ``tl.dot`` (no composite epilogue / ``softmax_scale``).
+  - K loaded as ``(d_head, tile_s)`` via byte-conserving reshape of
+    the deployed ``(tile_s, d_head)`` shard (ADR-0060 §3
+    reshape-not-transpose caveat).
+  - No causal masking / step-skip; blocking ``tl.recv``.
+  - ``TILE_S_KV`` is assumed to divide ``S_local`` evenly; last-tile
+    padding is not modelled.
+"""
+from __future__ import annotations
+
+
+TILE_S_KV = 1024  # ADR-0063 §A.2 S_kv-axis tile sweep (per-tile width).
+
+
+def _merge_running(m, l, O, m_step, l_step, O_step, *, tl):
+    """Online-softmax merge of two partial ``(m, ℓ, O)`` triples."""
+    m_new = tl.maximum(m, m_step)
+    scale_old = tl.exp(m - m_new)
+    scale_step = tl.exp(m_step - m_new)
+    l_new = l * scale_old + l_step * scale_step
+    O_new = O * scale_old + O_step * scale_step
+    return m_new, l_new, O_new
+
+
+def gqa_attention_prefill_long_kernel(
+    q_ptr: int,
+    k_ptr: int,
+    v_ptr: int,
+    o_ptr: int,
+    T_q: int,
+    S_kv: int,
+    d_head: int,
+    C: int,
+    P: int = 1,
+    *,
+    tl,
+) -> None:
+    """Head-parallel prefill attention with tile-granular Ring KV (ADR-0060 §5.5).
+
+    Tensor layout (per-rank shapes):
+      Q : (T_q_local, d_head)     one head per CUBE; replicated within
+          a CUBE for P=1, or sharded ``pe="row_wise"`` for P>1 so each
+          PE owns ``T_q_local = T_q // P`` query rows.
+      K : (S_kv, d_head)          sharded cube_row_wise → each CUBE owns
+          (S_local, d_head). Tiles loaded as (d_head, tile_s) via
+          byte-conserving reshape.
+      V : (S_kv, d_head)          sharded cube_row_wise → each CUBE owns
+          (S_local, d_head). Tiles loaded as (tile_s, d_head).
+      O : (T_q * C, d_head)       sharded cube_row_wise → each CUBE
+          writes its own ``T_q`` rows; for P>1 each PE writes a
+          disjoint ``(T_q_local, d_head)`` slice (no intra-CUBE
+          reduce — disjoint output rows). NO inter-CUBE reduce.
+
+    Algorithm: nested loop over (ring_step k, tile_idx t). At k=0 each
+    rank loads its own block's tiles from HBM; at k > 0 it receives
+    tiles from its E neighbour. Tiles are forwarded W to the next
+    ring step. Each tile's partial is folded into the running
+    (m, ℓ, O) via online-softmax merge.
+    """
+    pe_id = tl.program_id(axis=0)
+    if P == 1:
+        # Head-parallel only: PE 0 of each CUBE participates.
+        if pe_id != 0:
+            return
+        T_q_local = T_q
+    else:
+        # Intra-CUBE PE-SP (ADR-0060 §5.5 last bullet + §B-item-3):
+        # query-axis split across P PEs of each CUBE. Output rows are
+        # disjoint per PE ⇒ no intra-CUBE reduce. Each PE drives its
+        # own same-lane ring via independent IPCQ channels.
+        if pe_id >= P:
+            return
+        if T_q % P != 0:
+            raise ValueError(
+                f"T_q={T_q} must be divisible by P={P} when P > 1; "
+                "use P=1 for the PE-0-only fallback path."
+            )
+        T_q_local = T_q // P
+
+    S_local = S_kv // C
+    n_tiles = (S_local + TILE_S_KV - 1) // TILE_S_KV
+    KV_ROW_BYTES = d_head * 2  # f16
+    Q = tl.load(q_ptr, shape=(T_q_local, d_head), dtype="f16")
+
+    # ── Bootstrap: (t=0, k=0) — load my own tile 0, establish (m, ℓ, O) ──
+    #
+    # Cannot be folded into the (t, k) loop (kernbench-only limitation):
+    #   - persistent (m, ℓ, O) must live OUTSIDE ``tl.scratch_scope``,
+    #     otherwise scope teardown discards them before the next tile's
+    #     merge can read them;
+    #   - kernbench has no scratch-backed initializer — ``tl.zeros`` /
+    #     ``tl.full`` return addr=0 handles with no backing storage, so
+    #     they cannot be overwritten via ``tl.copy_to`` to seed (-inf, 0, 0);
+    #   - the ring-step ordering also requires k=0 to send its loaded
+    #     tile W before any k>0 iteration can recv from E, so the very
+    #     first (t=0, k=0) step has to run outside the scoped loop body
+    #     where the send is conditional on ``C > 1``.
+    # Triton port: limitation does not apply (SSA tensors stay live
+    # across iterations); the bootstrap collapses into the loop.
+    tile_s = min(TILE_S_KV, S_local)
+    K_t = tl.load(k_ptr, shape=(d_head, tile_s), dtype="f16")
+    V_t = tl.load(v_ptr, shape=(tile_s, d_head), dtype="f16")
+    if C > 1:
+        tl.send(dir="W", src=K_t)
+        tl.send(dir="W", src=V_t)
+    scores = tl.dot(Q, K_t)
+    m = tl.max(scores, axis=-1)
+    exp_scores = tl.exp(scores - m)
+    l = tl.sum(exp_scores, axis=-1)
+    O = tl.dot(exp_scores, V_t)
+
+    # ── Nested loop: outer tile, inner ring step ──
+    # Each tile propagates all the way through the ring before the next
+    # tile starts; IPCQ in-flight depth stays at 1 per direction.
+    #
+    # Per outer ``t``:
+    #   k=0      : load my own tile ``t`` from HBM
+    #   k=1..C-1 : receive tile ``t`` of a peer's block from E
+    #              (sent by my E neighbour during their previous k step)
+    #   k<C-1    : forward the tile to W
+    #
+    # The ``(t=0, k=0)`` iteration is the bootstrap above; the loop skips
+    # it via ``k_start``. Every other iteration uses ``scratch_scope`` +
+    # ``copy_to`` to recycle per-tile intermediates.
+    #
+    # Triton port: drop ``with tl.scratch_scope():`` and replace each
+    # ``copy_to`` with a Python rebind (e.g. ``m = m_new``).
+    for t in range(n_tiles):
+        tile_start = t * TILE_S_KV
+        tile_s = min(TILE_S_KV, S_local - tile_start)
+        k_start = 1 if t == 0 else 0
+        for k in range(k_start, C):
+            with tl.scratch_scope():
+                if k == 0:
+                    K_t = tl.load(k_ptr + tile_start * KV_ROW_BYTES,
+                                  shape=(d_head, tile_s), dtype="f16")
+                    V_t = tl.load(v_ptr + tile_start * KV_ROW_BYTES,
+                                  shape=(tile_s, d_head), dtype="f16")
+                else:
+                    K_t = tl.recv(dir="E", shape=(d_head, tile_s), dtype="f16")
+                    V_t = tl.recv(dir="E", shape=(tile_s, d_head), dtype="f16")
+                if k < C - 1:
+                    tl.send(dir="W", src=K_t)
+                    tl.send(dir="W", src=V_t)
+                scores = tl.dot(Q, K_t)
+                m_step = tl.max(scores, axis=-1)
+                exp_scores = tl.exp(scores - m_step)
+                l_step = tl.sum(exp_scores, axis=-1)
+                O_step = tl.dot(exp_scores, V_t)
+                m_new, l_new, O_new = _merge_running(
+                    m, l, O, m_step, l_step, O_step, tl=tl,
+                )
+                tl.copy_to(m, m_new)
+                tl.copy_to(l, l_new)
+                tl.copy_to(O, O_new)
+
+    # ── Final normalise + store (each CUBE writes its own head) ──
+    O_final = O / l
+    tl.store(o_ptr, O_final)

src/kernbench/benches/_gqa_attention_prefill_short.py

@@ -0,0 +1,147 @@
+"""GQA fused-attention prefill kernel — short context (ADR-0060 §B.split.2).
+
+Prefill analogue of ``_gqa_decode_short.py`` — same CUBE/PE layout
+(``kv_per_cube`` heads per CUBE, group-PE-SP, within-group chain reduce,
+no inter-CUBE reduce). The only structural difference from short decode:
+``T_q`` may be > 1 (prefill processes multiple query tokens) and Q is
+shaped ``(T_q, h_kv·d_head)`` — one Q head per KV head, no GQA M-fold.
+
+The local attention uses an S_kv-axis tile sweep (ADR-0063 §A.2) so
+per-rank scratch is bounded by ``TILE_S_KV``.
+
+No Ring KV here — each owned head is fully resident at its CUBE.
+"""
+from __future__ import annotations
+
+
+TILE_S_KV = 1024  # ADR-0063 §A.2 S_kv-axis tile sweep (per-tile width).
+
+
+def _merge_running(m_local, l_local, O_local, m_other, l_other, O_other, *, tl):
+    """Online-softmax merge of two partial ``(m, ℓ, O)`` triples."""
+    m_new = tl.maximum(m_local, m_other)
+    scale_old = tl.exp(m_local - m_new)
+    scale_new = tl.exp(m_other - m_new)
+    l_new = l_local * scale_old + l_other * scale_new
+    O_new = O_local * scale_old + O_other * scale_new
+    return m_new, l_new, O_new
+
+
+def gqa_attention_prefill_short_kernel(
+    q_ptr: int,
+    k_ptr: int,
+    v_ptr: int,
+    o_ptr: int,
+    T_q: int,
+    S_kv: int,
+    h_kv: int,
+    d_head: int,
+    C: int,
+    P: int,
+    kv_per_cube: int,
+    *,
+    tl,
+) -> None:
+    """Short-context prefill with PE-parallel heads + intra-group PE-SP."""
+    group_size = P // kv_per_cube
+    pe_id = tl.program_id(axis=0)
+    pe_in_group = pe_id % group_size
+    S_local = S_kv // group_size
+
+    # ── Local attention (S_kv-axis tile sweep, ADR-0063 §A.2) ──
+    Q = tl.load(q_ptr, shape=(h_kv * T_q, d_head), dtype="f16")
+    n_tiles = (S_local + TILE_S_KV - 1) // TILE_S_KV
+    KV_ROW_BYTES = d_head * 2  # f16
+
+    # Tile 0: establishes persistent (m_local, l_local, O_local).
+    #
+    # Cannot be folded into the Tiles 1..N loop (kernbench-only limitation):
+    #   - persistent (m, ℓ, O) must live OUTSIDE ``tl.scratch_scope``,
+    #     otherwise scope teardown discards them before the next tile's
+    #     merge can read them;
+    #   - kernbench has no scratch-backed initializer — ``tl.zeros`` /
+    #     ``tl.full`` return addr=0 handles with no backing storage, so
+    #     they cannot be overwritten via ``tl.copy_to`` to seed (-inf, 0, 0).
+    # So Tile 0 computes the initial running state directly; Tiles 1..N
+    # fold into it. Triton port: limitation does not apply (SSA tensors
+    # stay live across iterations) — a single unified loop suffices.
+    tile_s0 = min(TILE_S_KV, S_local)
+    K_T = tl.load(k_ptr, shape=(d_head, tile_s0), dtype="f16")
+    V = tl.load(v_ptr, shape=(tile_s0, d_head), dtype="f16")
+    scores = tl.dot(Q, K_T)
+    m_local = tl.max(scores, axis=-1)
+    centered = scores - m_local
+    exp_scores = tl.exp(centered)
+    l_local = tl.sum(exp_scores, axis=-1)
+    O_local = tl.dot(exp_scores, V)
+
+    # Tiles 1..n_tiles-1: fold into running state via online-softmax merge.
+    # Triton port: drop the ``with tl.scratch_scope():`` line and replace
+    # each ``copy_to`` with a Python rebind.
+    for tile_idx in range(1, n_tiles):
+        tile_start = tile_idx * TILE_S_KV
+        tile_s = min(TILE_S_KV, S_local - tile_start)
+        with tl.scratch_scope():
+            K_T_t = tl.load(k_ptr + tile_start * KV_ROW_BYTES,
+                            shape=(d_head, tile_s), dtype="f16")
+            V_t = tl.load(v_ptr + tile_start * KV_ROW_BYTES,
+                          shape=(tile_s, d_head), dtype="f16")
+            scores_t = tl.dot(Q, K_T_t)
+            m_tile = tl.max(scores_t, axis=-1)
+            centered_t = scores_t - m_tile
+            exp_scores_t = tl.exp(centered_t)
+            l_tile = tl.sum(exp_scores_t, axis=-1)
+            O_tile = tl.dot(exp_scores_t, V_t)
+            m_new, l_new, O_new = _merge_running(
+                m_local, l_local, O_local, m_tile, l_tile, O_tile, tl=tl,
+            )
+            tl.copy_to(m_local, m_new)
+            tl.copy_to(l_local, l_new)
+            tl.copy_to(O_local, O_new)
+
+    # ── Communication: within-group chain reduce-to-root (Level-2 only) ──
+    group_cols = min(4, group_size)
+    group_rows = (group_size + group_cols - 1) // group_cols
+    pe_col_in_group = pe_in_group % group_cols
+    pe_row_in_group = pe_in_group // group_cols
+
+    # Row chain (within group's row, along intra_W, leftward).
+    if group_cols > 1:
+        if pe_col_in_group < group_cols - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="intra_E", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="intra_E", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="intra_E", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        if pe_col_in_group > 0:
+            tl.send(dir="intra_W", src=m_local)
+            tl.send(dir="intra_W", src=l_local)
+            tl.send(dir="intra_W", src=O_local)
+
+    # Col bridge (within group, along intra_N, row-1 → row-0).
+    if pe_col_in_group == 0 and group_rows > 1:
+        if pe_row_in_group < group_rows - 1:
+            with tl.scratch_scope():
+                m_other = tl.recv(dir="intra_S", shape=m_local.shape, dtype="f16")
+                l_other = tl.recv(dir="intra_S", shape=l_local.shape, dtype="f16")
+                O_other = tl.recv(dir="intra_S", shape=O_local.shape, dtype="f16")
+                m_new, l_new, O_new = _merge_running(
+                    m_local, l_local, O_local, m_other, l_other, O_other, tl=tl,
+                )
+                tl.copy_to(m_local, m_new)
+                tl.copy_to(l_local, l_new)
+                tl.copy_to(O_local, O_new)
+        if pe_row_in_group > 0:
+            tl.send(dir="intra_N", src=m_local)
+            tl.send(dir="intra_N", src=l_local)
+            tl.send(dir="intra_N", src=O_local)
+
+    # ── Final normalise + store (group root only) ──
+    if pe_in_group == 0:
+        O_final = O_local / l_local
+        tl.store(o_ptr, O_final)