analytical-viz: capacity planning — symbol legend + plain-English formulas

Section 1 rewritten so a layman can read it without a glossary:

1. **Symbol legend table** listing every letter that appears in the
   three-axis formulas (N, b, N·b, HBM_per_PE, S_kv, kv_bpt,
   users_per_replica, TPOT SLO, step_latency, B_at_SLO, n_users,
   N_replicas, ⌈ ⌉). Each row: Symbol / Plain English / Value now
   for the current sidebar model + chip. For example:
     - N = 'Total number of model parameters (attention weights +
       FFN weights, across ALL layers)' → 6.98 billion params for
       Llama 3 8B
     - kv_bpt = 'KV cache bytes per token per user, summed across all
       layers (= 2·H_kv·d_head·b·layers)' → 128 KB/token

2. **Three-axis table** widened from (Formula / Meaning / Grows with)
   to (In symbols / In plain English / What it says / Value now).
   Axis A now shows the actual substituted arithmetic:
   ⌈ 13.96 GB / 6.0 GB ⌉ = ⌈ 2.33 ⌉ = 3 PEs.
   B and C flag 'see calculator' because they depend on user inputs.

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,434 @@
+# 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 에 에러 없음.
+
+### D8. Single-op-cmd fast-path (cmd 타입별 FIXED)
+
+모든 **single-op** cmd — DMA descriptor (`DmaReadCmd`, `DmaWriteCmd`)
+*그리고* single-op compute dispatch (`GemmCmd`, `MathCmd`, `CopyCmd`)
+— 는 일반 control-path 비용과 분리된, 가벼운 FIXED 를 지불한다:
+**`fixed_per_single_op_cmd_cycles`** (기본 **8 cycles**) 가 이들 전부에서
+`fixed_per_cmd_cycles` 를 대체한다. **`CompositeCmd` 만이 일반
+control-path 비용** `fixed_per_cmd_cycles` (= 40 cycles) 를 유지하는
+유일한 cmd 다.
+
+구체적으로 D1 의 공식이 다음과 같이 확장된다:
+```
+dispatch_cycles(cmd) = FIXED(type(cmd)) + cmd.logical_bytes × R
+FIXED(CompositeCmd) = fixed_per_cmd_cycles            (= 40)
+FIXED(나머지 전부, 즉 모든 single-op cmd:
+  DmaReadCmd / DmaWriteCmd / GemmCmd / MathCmd / CopyCmd)
+                    = fixed_per_single_op_cmd_cycles  (= 8)
+```
+
+**왜 single-op cmd 에 가벼운 FIXED 를 따로 두나.** D1 은 모든 발행 cmd 에
+동일한 FIXED 를 매겼다 — 가장 단순한 모델이라서. 그러나 실제로 single-op
+cmd — DMA descriptor 든, 엔진에 단일 GEMM / elementwise / copy 를
+dispatch 하는 것이든 — 는 40-cycle FIXED 가 모델링하던 일반 control-path
+비용보다 훨씬 가볍다. scheduler 측 plan 을 만들 필요 없이, descriptor 나
+instruction 하나를 엔진 queue 에 push 하는 것뿐이다:
+
+- NVIDIA Hopper TMA: 단일 PTX 명령 (`cp.async.bulk`) 으로 bulk async DMA
+  발사 — SM 측에선 ~1 ISA cycle, TMA 엔진이 fan-out / chunking 자체 처리.
+- NVIDIA Ampere `cp.async`: warp-level async copy ~1–2 cycles.
+- 단일 MMA / tensor-core 발행 (`wgmma`, `mma.sync`): math pipe 에 명령
+  하나, control-path round trip 아님.
+- AMD AQL packet 을 HSA queue 에 write: ~5–15 cycles.
+- 일반 descriptor-ring-push 디자인 (Synopsys/Xilinx-style DMA IP): MMIO
+  write 들 합쳐 ~5–20 cycles.
+- Tenstorrent tile descriptor 발행, Habana Gaudi TPC descriptor RAM
+  + start register: descriptor 당 한 자릿수 cycles.
+
+40-cycle FIXED 는 `CompositeCmd` 에 한해 정당하다. composite 는 단일
+dispatch 가 아니기 때문이다: scheduler 가 tile-feeder plan 을 만들고,
+내부 `DMA_READ → FETCH → GEMM → STORE → DMA_WRITE` stage 들에 걸쳐
+per-tile read/write hazard 를 추적하고, completion handle 을 배선해야
+한다. single-op `GemmCmd` 는 그런 기구 없이 엔진 발행 하나일 뿐이다 —
+이를 composite 와 같은 40 cycles 로 매기면 단일 명령을 스케줄된 plan
+전체와 혼동하는 것이다.
+
+**왜 이게 커널 평가에 중요한가.** 일정한 40-cycle FIXED 는 user 가 직접
+orchestrate 하는 single-op-primitive 커널의 dispatch overhead 를 부풀린다.
+`K = N_K · TILE_K` 의 chunked-prefetching async GEMM 은 K-tile 마다
+`DmaReadCmd` (B-chunk load) **와** `GemmCmd` (per-chunk `tl.dot`) **와**
+`MathCmd` (누적 `tl.add`) 를 발행하므로, per-cmd FIXED 가 dispatch
+예산을 지배한다. single-op fast-path 는 D1 이 잡으려던 *구조적 cmd 수
+신호* (`N` 개의 single-op cmd 를 발행하는 recipe 가 `N` tile 을 내부로
+묶는 composite 하나보다 더 비싸야 한다는 점) 를 유지하면서, 40-cycle
+일반 control-path 비용을 덧씌우는 모델링 과잉 청구는 분리한다.
+
+**왜 굳이 8 cycles 인가.** 위 조사 범위 (TMA ~1, descriptor-ring ~15) 의
+중간값. 의도적인 *모델링 디폴트* — 측정된 HW 수치가 아님. 이 디폴트의 역할
+은 정성적 동작에서 single-op dispatch path 를 composite control path 와
+구별하는 것. Topology override 가능 (D4).
+
+**Composite 수치에 미치는 영향.** Composite tile 내부 stage (HW tile 당
+`DMA_READ`, `FETCH`, `GEMM`, `STORE`, `DMA_WRITE`) 는 scheduler 의
+tile-feeder loop 가 발행하므로 host-side TLContext 의 `_charge_dispatch`
+를 안 거친다. 단일 `CompositeCmd` 자체는 여전히 40-cycle control-path
+FIXED 를 지불한다. D8 은 user-side single-op primitive (`tl.load` /
+`tl.store` / `tl.dot` / `tl.add` / …) 에 매기는 FIXED 만 바꾼다.
+Composite 측정치는 안정 유지.
+
+## 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/ADR-0070-algo-gqa-short-context-attention.md

@@ -0,0 +1,243 @@
+# ADR-0070: GQA Short-Context Attention — Unified A1/A2/A4/B Mapping for Prefill and Decode
+
+## Status
+
+Proposed — short-context (single tile to a handful of KV tiles) GQA
+attention benchmark covering **both prefill and decode** in **four
+mapping modes** (A1/A2/A4/B) and **three composite tiers** (without
+composite, GEMM-only composite, composite + softmax_merge fused).
+Compute-only attention — weight DMA, ITL batching, and end-to-end
+LLM driver scope are out of scope (separate ADRs).
+
+## Context
+
+A GQA layer at the LLaMA-3.1-70B headline shape (`h_q=64`, `h_kv=8`,
+`d_head=128`, GQA group `G = h_q/h_kv = 8`) on a 4×4 cube SIP can be
+mapped to AHBM in several ways that trade KV-cache memory, HBM read
+volume per cube, and intra-cube IPCQ traffic against PE parallelism.
+
+> **Bench-vs-headline shape.** The kernels and tests use a
+> `d_head = 64` proxy (`tests/attention/test_gqa_short_context.py`
+> `D_HEAD = 64`) so the per-tile working set fits comfortably in
+> scratch across all four modes. The mapping decisions are
+> `d_head`-independent — only the per-PE GEMM tile bytes scale.
+
+The four candidate mappings (per *GQA-mapping-on-AHBM-short-context* note §1):
+
+```
+Mode  kv_per_cube  C        group_size  Reduce / broadcast topology
+----  -----------  -------  ----------  ----------------------------------
+A1            1    h_kv      P (=8)     row chain + col bridge (2×4 mesh)
+A2            2    h_kv/2    P/2 (=4)   row chain only
+A4            4    h_kv/4    P/4 (=2)   single intra_W / intra_E hop
+B             8    1         1          no broadcast / no reduce (single PE)
+```
+
+Within each cube the 8 PEs split into `kv_per_cube` groups of
+`group_size = P/kv_per_cube` PEs; each group owns one KV head.
+
+The kernel must be able to compare these mappings at multiple context
+lengths (8K..64K KV tokens) and across the three composite tiers so
+the team can size the per-cube HBM / IPCQ / GEMM tradeoff empirically.
+
+## Decision
+
+### 1. Mapping (unified A1/A2/A4/B)
+
+Two unified kernels select mode at launch via `kv_per_cube ∈ {1,2,4,8}`:
+
+- **Prefill** (`gqa_attention_prefill_short_kernel`):
+  - Q-tile split — `T_q` rows floor-balanced across the group's
+    `group_size` PEs.
+  - FA2 head fusion — `G` Q heads fused into the M dimension of one
+    batched GEMM per PE per tile.
+  - **IPCQ KV broadcast** — group root (`pe_in_group == 0`) loads K/V
+    from HBM and IPCQ-broadcasts each tile to the rest of its group;
+    q-tiles are independent so no intra-group reduce.
+
+- **Decode** (`gqa_attention_decode_short_kernel`):
+  - Sequence-shard — each PE owns `S_local = S_kv/group_size` tokens
+    of the group's KV head.
+  - FA2 head fusion — same as prefill.
+  - **IPCQ chain reduce** — per-PE partial `(m, ℓ, O)` chain-reduce
+    up to group root (PE 0), which normalizes and stores.
+
+Geometry (2×4 mesh):
+
+```
+group_size = 8  : 2 row × 4 col   (A1)
+group_size = 4  : 1 row × 4 col   (A2)
+group_size = 2  : 1 row × 2 col   (A4)
+group_size = 1  : single PE       (B)
+```
+
+### 2. Shard addressing (ADR-0011 D-VA1 contract)
+
+Deploy is per-(sip, cube, pe) but `tl.load` receives a single global
+VA per tensor. The kernel computes its own shard base offset from
+`tl.program_id(axis=0)` (PE id) and `tl.program_id(axis=1)` (cube id):
+
+```python
+# Decode shape (pe=row_wise): K is split across PEs by sequence shard.
+cube_K_base       = cube_id        * kv_per_cube * K_HEAD_BYTES
+head_K_base       = group_id_in_cube * K_HEAD_BYTES
+pe_K_seq_offset   = pe_in_group     * n_tiles_per_pe * K_TILE_BYTES
+k_shard_base      = k_ptr + cube_K_base + head_K_base + pe_K_seq_offset
+```
+
+For **prefill** the K/V dp is `pe=replicate` (group root reads and
+broadcasts) so the `pe_in_group` term drops out:
+`k_head_shard_base = k_ptr + cube_K_base + head_K_base`.
+
+Skipping `cube_id` collapses all cubes onto cube 0's HBM region
+(observed pre-fix as an 11.5× per-cube DMA imbalance).
+
+### 3. Three composite tiers
+
+The same mapping is exercised against three GEMM/MATH dispatch styles
+so the contribution of the composite API vs the recipe-driven fusion
+can be isolated:
+
+```
+(1) without composite                  primitives only (tl.dot, tl.exp, …)
+(2) with composite (GEMM-only)         tl.composite(op="gemm")
+(3) with composite + softmax_merge     tl.composite(prologue=[softmax_merge],
+                                                    op="gemm", out=O,
+                                                    epilogue=[add])
+```
+
+File layout (`src/kernbench/benches/gqa_helpers/short_ctx/`):
+
+```
+_gqa_attention_{prefill,decode}_short.py                   (1)
+_gqa_attention_{prefill,decode}_short_composite.py         (2)
+_gqa_attention_{prefill,decode}_short_composite_fused.py   (3)
+```
+
+Tier (2) is **GEMM-only by definition** — no recipe-driven fusion,
+so both prefill's and decode's P·V stay `tl.dot`. Recipe-driven
+fusion is exactly what tier (3) adds.
+
+Tier (3) on **multi-cube modes (A1/A2/A4) of prefill** relies on the
+D4 supplement: any composite operand that is an IPCQ recv'd slot
+(non-root PE's K_T in Q·Kᵀ, V in P·V) is pinned and read in place
+instead of being DMA-streamed from HBM. Without the supplement every
+recv slot fed to a composite — both Q·Kᵀ's `b=K_T` and P·V's `b=V`
+— PageFaults on PA decode (PE scratch addresses overflow the 51-bit
+PA range).
+
+### 4. Caller contract — `_validate_config`
+
+Each kernel module exposes a single `_validate_config(...)` helper
+called by the bench wrapper before launch. The kernel itself is lean
+(no inline `if` guards): the contract block is enforced caller-side,
+sim cost zero, but catches every silent-shard-corruption foot-gun
+(`kv_per_cube=3 → group_size=2` via integer division, non-integer
+GQA group, mismatched cube count, partial-tile `S_kv`, …) before any
+address arithmetic runs.
+
+### 5. Tensor layouts (host-side, mode-invariant byte totals)
+
+- **Q**: `(kv_per_cube·T_q, h_q·d_head/kv_per_cube)`
+  `dp=(cube=column_wise, pe=replicate)` over `C = h_kv/kv_per_cube`.
+- **K**: `(h_kv·S_kv·d_head/TILE_S_KV, TILE_S_KV)` tile-major.
+  Decode: `dp=(cube=row_wise, pe=row_wise)`.
+  Prefill: `dp=(cube=row_wise, pe=replicate)` (broadcast model).
+- **V**: `(h_kv·S_kv, d_head)` native. Same dp as K.
+- **O**: same dp as Q.
+
+Caller pre-scales Q by `1/√d_head`.
+
+## Verification
+
+### Smoke + regression (`tests/attention/test_gqa_short_context.py`)
+
+- Prefill / decode smoke (all 4 modes, single-tile S_kv).
+- Multi-tile coverage (all 4 modes, S_kv chosen so each PE owns
+  ≥2 tiles).
+- Op-count invariants — prefill `dma_writes = group_size·kv_per_cube·C`,
+  `dma_reads = P·C + 2·kv_per_cube·C·n_tiles`; decode
+  `dma_writes = kv_per_cube·C`.
+- IPCQ topology — short-context kernel never emits inter-CUBE E/W
+  IPCQ.
+- **ADR-0011 D-VA1 regression** — `per_cube_disjoint_src_addrs`,
+  `per_cube_dma_balanced` (max/min DMA busy ratio < 1.1×) per mode
+  per phase.
+
+### Composite smoke
+
+- `test_{prefill,decode}_composite_smoke[A1/A2/A4/B]` — tier (2),
+  4 modes × 2 phases.
+- `test_{prefill,decode}_composite_fused_smoke[A1/A2/A4/B]` — tier
+  (3), 4 modes × 2 phases.
+- `test_decode_composite_fused_multitile[A1/A2/A4/B]` —
+  `n_tiles_per_pe == 2` per mode so the recipe-fused tile loop
+  actually executes.
+
+### Mode × context sweep (`tests/attention/test_gqa_short_context_sweep_*.py`)
+
+Six sweep files (3 tiers × {prefill, decode}) measure each
+`(mode, S_kv)` cell and dump the same metric set to CSV:
+
+```
+wall_us, n_pe, gemm_util, math_count, math_pipeline_us, hbm_bw_util,
+hbm_read_mb, hbm_write_kb, ipcq_kb, kv_cache_per_cube_mb
+```
+
+S_kv ∈ {8K, 16K, 32K, 64K}; outputs under `docs/sweeps/`.
+
+Composite/fused sweeps count both `gemm_f16` (non-pipeline) **and**
+`TileToken/GEMM` (pipeline composite path); fused additionally counts
+`TileToken/MATH`. Baseline sweeps count `gemm_f16` only.
+
+## Known limitations
+
+1. **Causal mask** — kernel is non-causal. Adding causal masking is
+   orthogonal to mapping and is scheduled as a separate ADR.
+2. **f16 online-softmax accumulator** — `(m, ℓ, O)` are f16 throughout.
+   Long-context numerical drift will need an f32 accumulator before
+   correctness-grade use.
+3. **End-to-end numeric validation** — current tests are op-count /
+   topology / per-cube-DMA invariants. No full-kernel
+   `np.allclose` against a torch/numpy reference yet.
+4. **Skinny-M GEMM underfill** — decode (`T_q=1`) gives `M = G = 8`,
+   well below `PE_SCHEDULER` supertile `TILE_M = 32`. The two GEMM
+   paths handle this differently:
+   - **Primitive `tl.dot`** dispatches the GEMM at the actual
+     `m = 8` (no padding); GEMM time reflects the real 8-row work.
+   - **`tl.composite(op="gemm")`** tiles at `TILE_M = 32`, padding
+     `M` 4× with zeros; GEMM time reflects the padded 32-row work.
+   Consequence: when sweep CSVs show tier (2)/(3) `gemm_util`
+   higher than tier (1), most of that gap is supertile padding
+   overhead, **not** extra useful work or fusion savings — read it as
+   "composite tile shape doesn't match decode-skinny shape" rather
+   than "composite/fusion is more compute-intensive." This is the
+   **intended decode shape** for this benchmark; lifting it requires
+   batched-M (multi-request inference, `B_sys > 1`), tracked as a
+   separate batched variant — see "Future work".
+5. **Multi-cube prefill composite without the ADR-0065 D4 supplement**
+   PageFaults. The supplement — IPCQ recv'd slots are pinned and read
+   in place as composite operands — ships alongside this ADR.
+
+## Future work
+
+- Batched-M variant (`B_sys ∈ {1,2,4,8,16}`) so composite/fused
+  pipeline overlap shows up in wall_clock rather than only in
+  per-engine utilization.
+- Long-context (S_kv ≥ 128K) sweep extension for the headline LLaMA
+  decode target.
+- Causal-mask + f32 accumulator promotion.
+- 3-variant comparison plots from the sweep CSVs (wall, gemm util,
+  hbm bw util, ipcq, kv cache) — generated under `docs/diagrams/`.
+
+## References
+
+- ADR-0011 — PhysAddr + VA contract (D-VA1: kernel-computed shard
+  offset).
+- ADR-0060 — GQA fused attention on AHBM (precursor of the unified
+  mapping).
+- ADR-0064 — CPU issue-cost model (composite supertile motivation).
+- ADR-0065 — Flat ops, composite, softmax_merge recipe (with the D4
+  supplement applied here).
+- ADR-0070 supersedes ADR-0060 §B.split.2 short-prefill clause.
+- *GQA-mapping-on-AHBM-short-context* note (research design rationale).

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-0014-dev-pe-pipeline-execution-model.md

@@ -184,8 +184,15 @@ The empty-`ops` form is the legacy single-op path.
 - `DMA_READ`: `simpy.Resource(capacity=1)`.
 - `DMA_WRITE`: `simpy.Resource(capacity=1)`.
 - Both channels run concurrently (READ ∥ WRITE allowed).
-- Within a channel, requests serialize (READ ∥ READ disallowed; same
+- Within a channel, requests serialize **at the issue path** (READ ∥ READ
-  for WRITE).
+  issued out-of-order is disallowed; same for WRITE). The channel is
+  held only until the request is enqueued onto the next hop (router),
+  then released — it does **not** block until the HBM round-trip
+  completes. Multiple in-flight requests are therefore allowed, and
+  HBM-level serialization is the responsibility of the HBM controller's
+  per-PC `available_at` timestamps (ADR-0033 D1). This is what allows
+  back-to-back small-tile DMAs to amortize the per-request head latency
+  through the fabric.
 - `vc_comm` is an orthogonal channel for IPCQ traffic defined in
   ADR-0023 D8 — out of scope for this ADR.
 

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,471 @@
+# 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.
+
+### D8. Single-op-cmd fast-path (per-cmd-type FIXED)
+
+Every **single-op** command — the DMA descriptors (`DmaReadCmd`,
+`DmaWriteCmd`) *and* the single-op compute dispatches (`GemmCmd`,
+`MathCmd`, `CopyCmd`) — carries a separate, lighter FIXED than the
+general control-path cost: **`fixed_per_single_op_cmd_cycles`** (default
+**8 cycles**) replaces `fixed_per_cmd_cycles` for all of them.
+**`CompositeCmd` is the only command that stays on the general
+control-path cost** `fixed_per_cmd_cycles` (= 40 cycles).
+
+Concretely, the D1 formula becomes
+```
+dispatch_cycles(cmd) = FIXED(type(cmd)) + cmd.logical_bytes × R
+FIXED(CompositeCmd) = fixed_per_cmd_cycles         (= 40)
+FIXED(everything else, i.e. every single-op cmd:
+  DmaReadCmd / DmaWriteCmd / GemmCmd / MathCmd / CopyCmd)
+                    = fixed_per_single_op_cmd_cycles  (= 8)
+```
+
+**Why a separate, lighter FIXED for single-op cmds.** D1 treated every
+emitted command uniformly because that is the simplest defensible
+model. In practice an single-op command — whether a DMA descriptor or a
+single GEMM / elementwise / copy dispatch to an engine — is *much*
+cheaper than the generic control-path cost the 40-cycle FIXED was
+modeling. It is a single descriptor or instruction pushed to an engine
+queue, with no scheduler-side plan to build:
+
+- NVIDIA Hopper TMA: a single PTX instruction (`cp.async.bulk`) initiates
+  a bulk async DMA — ~1 ISA cycle on the SM side, with the TMA engine
+  doing fan-out and chunking itself.
+- NVIDIA Ampere `cp.async`: ~1–2 cycles per warp-level async copy.
+- A single MMA / tensor-core issue (`wgmma`, `mma.sync`): one
+  instruction to the math pipe, not a control-path round trip.
+- AMD AQL packet write to a HSA queue: ~5–15 cycles.
+- Generic descriptor-ring-push designs (Synopsys/Xilinx-style DMA IP):
+  ~5–20 cycles for the MMIO writes.
+- Tenstorrent tile descriptor emission, Habana Gaudi TPC descriptor RAM
+  + start register: single-digit cycles per descriptor.
+
+The 40-cycle FIXED is justified for `CompositeCmd` specifically, because
+a composite is *not* a single dispatch: the scheduler must build a
+tile-feeder plan, track per-tile read/write hazards across its internal
+`DMA_READ → FETCH → GEMM → STORE → DMA_WRITE` stages, and wire a
+completion handle. An single-op `GemmCmd` is one engine issue with none of
+that machinery; charging it the same 40 cycles as a composite conflates
+a single instruction with a whole scheduled plan.
+
+**Why this matters for kernel evaluation.** A uniform 40-cycle FIXED
+inflates the dispatch overhead of user-orchestrated single-op-primitive
+kernels. A chunked-prefetching async GEMM at `K = N_K · TILE_K` emits,
+per K-tile, a `DmaReadCmd` (B-chunk load) **and** a `GemmCmd` (the
+per-chunk `tl.dot`) **and** a `MathCmd` (the running `tl.add`
+accumulate) — so the per-command FIXED dominates its dispatch budget.
+The single-op fast-path keeps the *structural* command-count signal that
+D1 was designed to capture (a recipe that emits `N` single-op commands
+*does* pay more than a recipe that emits one composite covering `N`
+tiles internally) without conflating it with a modeling overcharge that
+would be specific to a 40-cycle generic control path.
+
+**Why 8 cycles specifically.** 8 sits at the mid-range of the survey
+above (TMA ~1 to descriptor-ring ~15). It is intentionally a *modeling
+default*, not a measured HW number — the role of this default is to
+distinguish the single-op dispatch path from the composite control path in
+qualitative behaviour. The value is overridable per topology (D4).
+
+**Effect on composite numbers.** Composite tile-internal stages
+(`DMA_READ`, `FETCH`, `GEMM`, `STORE`, `DMA_WRITE` per HW tile) are
+emitted by the scheduler's tile-feeder loop, not by the host-side
+TLContext, so they do **not** go through `_charge_dispatch`. The single
+`CompositeCmd` itself still pays the 40-cycle control-path FIXED. D8
+only changes the FIXED charged to user-side single-op primitives
+(`tl.load` / `tl.store` / `tl.dot` / `tl.add` / …). Composite
+measurements remain stable.
+
+## 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/diagrams/allreduce_latency_plots/AllReduce_LRAB_2Dtorus_6SiP_2x3_with_TCM_SRAM_HBM.csv

@@ -1,13 +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,128,256,2345.040000000012
-hbm,torus_2d,6,1024,2048,2717.2783333333473
+hbm,torus_2d,6,1024,2048,2942.2783333333473
-hbm,torus_2d,6,8192,16384,7315.184999999989
+hbm,torus_2d,6,8192,16384,7540.184999999989
-hbm,torus_2d,6,32768,65536,23081.26500000037
+hbm,torus_2d,6,32768,65536,23306.26500000037
-sram,torus_2d,6,128,256,2060.040000000012
+sram,torus_2d,6,128,256,2285.040000000012
-sram,torus_2d,6,1024,2048,2909.2783333333473
+sram,torus_2d,6,1024,2048,3134.2783333333527
-sram,torus_2d,6,8192,16384,9523.184999999869
+sram,torus_2d,6,8192,16384,9748.184999999869
-sram,torus_2d,6,32768,65536,32201.265000000385
+sram,torus_2d,6,32768,65536,32426.265000000385
-tcm,torus_2d,6,128,256,1964.040000000012
+tcm,torus_2d,6,128,256,2189.040000000012
-tcm,torus_2d,6,1024,2048,2477.2783333333473
+tcm,torus_2d,6,1024,2048,2702.2783333333473
-tcm,torus_2d,6,8192,16384,6403.185000000109
+tcm,torus_2d,6,8192,16384,6628.18500000005
-tcm,torus_2d,6,32768,65536,19865.265000000378
+tcm,torus_2d,6,32768,65536,20090.265000000378

docs/diagrams/allreduce_latency_plots/summary.csv

@@ -1,37 +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,8,16,256,2918.5525000000157
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,2747.7400000000152
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,2999.740000000016
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,2855.990000000018
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,3107.990000000019
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,3072.490000000019
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,3324.4900000000207
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,3337.1133333333582
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,3589.1133333333582
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,3708.0333333333692
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,3960.0333333333692
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,4449.873333333393
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,4701.873333333393
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,5933.020000000124
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,6185.020000000124
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,8900.379999999863
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,9152.379999999861
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,14835.099999999224
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,15087.099999999224
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,26704.540000000765
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,26956.540000000765
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,38573.97999999701
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,38825.97999999701
-lrab_hierarchical_allreduce,ring_1d,6,8,16,256,2365.255833333347
+lrab_hierarchical_allreduce,ring_1d,6,8,16,256,2628.2558333333477
-lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,2436.9433333333473
+lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,2699.943333333348
-lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,2532.526666666683
+lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,2795.526666666683
-lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,2723.693333333349
+lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,2986.693333333351
-lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,3048.635000000021
+lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,3311.635000000021
-lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,3393.4016666666957
+lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,3656.4016666666957
-lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,4082.401666666714
+lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,4345.401666666714
-lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,5458.80166666677
+lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,5721.801666666768
-lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,8216.934999999943
+lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,8479.934999999887
-lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,13733.201666665835
+lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,13996.201666665835
-lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,24765.73500000064
+lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,25028.73500000064
-lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,35798.268333331536
+lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,36061.26833333154
-lrab_hierarchical_allreduce,torus_2d,6,8,16,256,1700.6025000000095
+lrab_hierarchical_allreduce,torus_2d,6,8,16,256,1925.6025000000104
-lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,1753.2900000000102
+lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,1978.290000000011
-lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,1823.540000000012
+lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,2048.540000000012
-lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,1964.040000000012
+lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,2189.040000000012
-lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,2196.8183333333463
+lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,2421.8183333333463
-lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,2477.2783333333473
+lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,2702.2783333333473
-lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,3038.1983333333583
+lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,3263.1983333333583
-lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,4159.5050000000665
+lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,4384.5050000000665
-lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,6403.185000000109
+lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,6628.18500000005
-lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,10890.5449999995
+lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,11115.5449999995
-lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,19865.265000000378
+lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,20090.265000000378
-lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,28839.98500000059
+lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,29064.985000000597

docs/diagrams/pe2pe_latency_plots/summary.csv

@@ -1,81 +1,81 @@
 hop,label,size_bytes,path,total_ns
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),128,ipcq,24.88749999999891
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),128,ipcq,42.88749999999891
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),128,raw,33.57999999999811
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),128,raw,51.57999999999811
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),256,ipcq,28.13749999999891
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),256,ipcq,46.13749999999891
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),256,raw,36.07999999999811
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),256,raw,54.07999999999811
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),384,ipcq,29.88749999999891
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),384,ipcq,47.88749999999891
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),384,raw,37.07999999999811
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),384,raw,55.07999999999811
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),512,ipcq,31.63749999999891
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),512,ipcq,49.63749999999891
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),512,raw,38.07999999999811
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),512,raw,56.07999999999811
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),768,ipcq,35.13749999999891
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),768,ipcq,53.13749999999891
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),768,raw,40.07999999999811
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),768,raw,58.07999999999811
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),1024,ipcq,38.63749999999891
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),1024,ipcq,56.63749999999891
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),1024,raw,42.07999999999811
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),1024,raw,60.07999999999811
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),2048,ipcq,52.63749999999891
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),2048,ipcq,70.63749999999891
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),2048,raw,50.07999999999811
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),2048,raw,68.07999999999811
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),4096,ipcq,80.63750000000073
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),4096,ipcq,98.63750000000073
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),4096,raw,66.08000000000175
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),4096,raw,84.08000000000175
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),8192,ipcq,136.63750000000073
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),8192,ipcq,154.63750000000073
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),8192,raw,98.08000000000175
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),8192,raw,116.08000000000175
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),10240,ipcq,164.63750000000073
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),10240,ipcq,182.63750000000073
-latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),10240,raw,114.08000000000175
+latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),10240,raw,132.08000000000175
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),128,ipcq,38.49749999999585
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),128,ipcq,56.49749999999585
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),128,raw,47.18999999999505
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),128,raw,65.18999999999505
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),256,ipcq,43.24749999999585
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),256,ipcq,61.24749999999585
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),256,raw,51.18999999999505
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),256,raw,69.18999999999505
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),384,ipcq,44.99749999999585
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),384,ipcq,62.99749999999585
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),384,raw,52.18999999999505
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),384,raw,70.18999999999505
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),512,ipcq,46.74749999999585
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),512,ipcq,64.74749999999585
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),512,raw,53.18999999999505
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),512,raw,71.18999999999505
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),768,ipcq,50.24749999999585
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),768,ipcq,68.24749999999585
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),768,raw,55.18999999999505
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),768,raw,73.18999999999505
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),1024,ipcq,53.74749999999585
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),1024,ipcq,71.74749999999585
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),1024,raw,57.18999999999505
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),1024,raw,75.18999999999505
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),2048,ipcq,67.74749999999585
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),2048,ipcq,85.74749999999585
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),2048,raw,65.18999999999505
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),2048,raw,83.18999999999505
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),4096,ipcq,95.74750000000131
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),4096,ipcq,113.74750000000131
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),4096,raw,81.19000000000233
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),4096,raw,99.19000000000233
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),8192,ipcq,151.7475000000013
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),8192,ipcq,169.7475000000013
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),8192,raw,113.19000000000233
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),8192,raw,131.19000000000233
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),10240,ipcq,179.7475000000013
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),10240,ipcq,197.7475000000013
-latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),10240,raw,129.19000000000233
+latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),10240,raw,147.19000000000233
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),128,ipcq,81.15999999999804
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),128,ipcq,99.15999999999804
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),128,raw,89.28999999999724
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),128,raw,107.28999999999724
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),256,ipcq,88.65999999999804
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),256,ipcq,106.65999999999804
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),256,raw,95.53999999999724
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),256,raw,113.53999999999724
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),384,ipcq,90.90999999999804
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),384,ipcq,108.90999999999804
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),384,raw,96.53999999999724
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),384,raw,114.53999999999724
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),512,ipcq,93.15999999999804
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),512,ipcq,111.15999999999804
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),512,raw,97.53999999999724
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),512,raw,115.53999999999724
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),768,ipcq,97.65999999999804
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),768,ipcq,115.65999999999804
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),768,raw,99.53999999999724
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),768,raw,117.53999999999724
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),1024,ipcq,103.15999999999804
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),1024,ipcq,121.15999999999804
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),1024,raw,102.53999999999724
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),1024,raw,120.53999999999724
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),2048,ipcq,125.15999999999804
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),2048,ipcq,143.15999999999804
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),2048,raw,114.53999999999724
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),2048,raw,132.53999999999724
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),4096,ipcq,169.15999999999985
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),4096,ipcq,187.15999999999985
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),4096,raw,138.54000000000087
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),4096,raw,156.54000000000087
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),8192,ipcq,257.15999999999985
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),8192,ipcq,275.15999999999985
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),8192,raw,186.54000000000087
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),8192,raw,204.54000000000087
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),10240,ipcq,301.15999999999985
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),10240,ipcq,319.15999999999985
-latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),10240,raw,210.54000000000087
+latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),10240,raw,228.54000000000087
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),128,ipcq,103.15999999999804
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),128,ipcq,121.15999999999804
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),128,raw,111.28999999999724
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),128,raw,129.28999999999724
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),256,ipcq,112.65999999999804
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),256,ipcq,130.65999999999804
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),256,raw,119.53999999999724
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),256,raw,137.53999999999724
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),384,ipcq,114.90999999999804
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),384,ipcq,132.90999999999804
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),384,raw,120.53999999999724
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),384,raw,138.53999999999724
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),512,ipcq,117.15999999999804
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),512,ipcq,135.15999999999804
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),512,raw,121.53999999999724
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),512,raw,139.53999999999724
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),768,ipcq,121.65999999999804
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),768,ipcq,139.65999999999804
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),768,raw,123.53999999999724
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),768,raw,141.53999999999724
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),1024,ipcq,127.15999999999804
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),1024,ipcq,145.15999999999804
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),1024,raw,126.53999999999724
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),1024,raw,144.53999999999724
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),2048,ipcq,149.15999999999804
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),2048,ipcq,167.15999999999804
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),2048,raw,138.53999999999724
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),2048,raw,156.53999999999724
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),4096,ipcq,193.15999999999985
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),4096,ipcq,211.15999999999985
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),4096,raw,162.54000000000087
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),4096,raw,180.54000000000087
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),8192,ipcq,281.15999999999985
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),8192,ipcq,299.15999999999985
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),8192,raw,210.54000000000087
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),8192,raw,228.54000000000087
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),10240,ipcq,325.15999999999985
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),10240,ipcq,343.15999999999985
-latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),10240,raw,234.54000000000087
+latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),10240,raw,252.54000000000087

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/Agentic_Runtime_Architecture.md

@@ -0,0 +1,848 @@
+# AHBM Agentic Runtime Architecture
+
+## Scope
+
+This note organizes the current design discussion for executing agentic workloads on a memory-centric AHBM-style architecture.
+
+The scope of this version is limited to:
+
+1. Motivation
+2. AHBM hardware assumptions
+3. Attention execution
+4. Attention execution policies
+5. Fan-in and join
+
+MoE execution and whole-model layer-aware scheduling are intentionally deferred to a later section.
+
+---
+
+# 1. Motivation
+
+Agentic workload patterns: loop, fan-out / fan-in. Loop 은 일반적인 LLM execution 과 다르지 않다. Sub-agent 가 만들어질 때 발생하는 Fan out / Fan in 이 주로 고려되어야 할 대상이다.
+
+## 1.1 Agentic fan-out
+
+An agentic workload often starts from one shared conversation or task context and then forks into several specialized sub-agents.
+
+```text
+Shared prefix
+   ├─ Agent A: Analyze performance
+   ├─ Agent B: Check correctness
+   └─ Agent C: Find alternatives
+```
+
+Each sub-agent sees:
+
+```text
+Shared prefix
++
+Private role or instruction
++
+Private generated continuation
+```
+
+The shared prefix may be long, while each private branch may initially contain only a small number of new tokens.
+
+This creates two important execution properties:
+
+1. All branches reuse the same prefix KV cache.
+2. New query tokens from multiple branches can potentially be batched.
+
+## 1.2 Why the workload is different from ordinary batching
+
+Ordinary batching groups unrelated requests that happen to arrive at similar times.
+
+Agentic fan-out is different because the branches are structurally related:
+
+```text
+Agent A context = Shared prefix + A suffix
+Agent B context = Shared prefix + B suffix
+Agent C context = Shared prefix + C suffix
+```
+
+The requests therefore have:
+
+- identical prefix KV,
+- different private suffix KV,
+- potentially synchronized execution points,
+- a later fan-in stage that combines their results.
+
+This structure creates opportunities that are not available in unrelated-request batching.
+
+## 1.3 Main optimization opportunity
+
+For attention, the key operation is:
+
+```text
+Q × Kᵀ
+```
+
+Multiple sub-agents can have different query rows while reading the same shared prefix KV.
+
+Instead of executing:
+
+```text
+Q_A × K_shared
+Q_B × K_shared
+Q_C × K_shared
+```
+
+independently, the runtime can combine the query rows:
+
+```text
+Q_combined =
+[
+  Q_A
+  Q_B
+  Q_C
+]
+```
+
+and execute:
+
+```text
+Q_combined × K_shared
+```
+
+The arithmetic is still row-independent, but the larger GEMM can improve utilization and amortize scheduling overhead.
+
+## 1.4 Main design questions
+
+The architecture must answer:
+
+- How should shared KV be distributed across CUBEs and PEs?
+- How should query rows from multiple sub-agents be grouped?
+- How should online softmax state be reduced across sequence shards?
+- Should Q be replicated or partitioned?
+- How should large fan-out results be joined back into the main agent?
+- Which data should be reused, recomputed, summarized, or materialized on demand?
+
+---
+
+# 2. AHBM Hardware Assumptions
+
+## 2.2 Sequence parallelism
+
+The sequence dimension of one KV head is distributed across the 32 PEs.
+
+```text
+KV head sequence
+   ├─ PE0  owns sequence shard 0
+   ├─ PE1  owns sequence shard 1
+   ├─ ...
+   └─ PE31 owns sequence shard 31
+```
+
+Each PE stores a different part of the sequence for the same KV head. A query row must attend to all 32 sequence shards, so every PE computes a partial attention result for its local KV shard.
+
+## 2.3 Q placement
+
+The baseline follows an AHBM-style replicated-Q execution model.
+
+Logically:
+
+```text
+The same Q rows are visible to all 32 PEs.
+```
+
+This does not require 32 independent physical copies. A possible implementation is:
+
+```text
+Chip-level Q source
+        ↓
+CUBE multicast
+        ↓
+Shared Q buffer within each CUBE
+        ↓
+8 PEs consume the same Q tile
+```
+
+Thus Q is logically replicated across PEs while physical traffic is reduced through multicast and shared buffering.
+
+## 2.4 KV placement
+
+KV remains stationary near the owning PE.
+
+```text
+PE0  → KV shard 0
+PE1  → KV shard 1
+...
+PE31 → KV shard 31
+```
+
+The baseline avoids:
+
+- KV remapping,
+- KV replication,
+- remote KV reads,
+- page-table reconstruction for every query group.
+
+## 2.5 GEMM engine assumptions
+
+Representative PE GEMM tile shapes include:
+
+```text
+16 × 16 × 16
+```
+
+and:
+
+```text
+8 × 64 × 8
+```
+
+For:
+
+```text
+8 sub-agents
+20 query tokens per sub-agent
+```
+
+the combined number of Q rows is:
+
+```text
+M = 8 × 20 = 160
+```
+
+If each PE owns 256 KV columns, a representative local GEMM is:
+
+```text
+Q_local: 160 × d
+K_local: d × 256
+```
+
+Both `M = 160` and `N = 256` align well with the assumed tile shapes.
+
+---
+
+# 3. Attention Execution
+
+## 3.1 Fan-out context structure
+
+Assume the parent agent has already prefetched a shared prefix.
+
+```text
+Shared prefix KV pages
+```
+
+The runtime forks the logical context:
+
+```text
+Branch A page table:
+[Shared prefix pages] + [A private pages]
+
+Branch B page table:
+[Shared prefix pages] + [B private pages]
+
+Branch C page table:
+[Shared prefix pages] + [C private pages]
+```
+
+The shared pages are referenced by multiple branches without being copied. Only private suffix pages are branch-specific.
+
+## 3.2 Sub-agent query batching
+
+Assume:
+
+```text
+8 sub-agents
+20 new tokens per sub-agent
+```
+
+The runtime forms:
+
+```text
+Q_A: 20 × d
+Q_B: 20 × d
+...
+Q_H: 20 × d
+```
+
+and concatenates them along the row dimension:
+
+```text
+Q_combined: 160 × d
+```
+
+The combined operation is:
+
+```text
+(160 × d) × (d × S)
+```
+
+where `S` is the total sequence length represented by all KV shards.
+
+Each query row remains logically independent. Batching changes the execution shape, not the attention semantics.
+
+## 3.3 Per-PE local attention
+
+Each PE owns only a local sequence shard.
+
+For PE `p`:
+
+```text
+Scores_p = Q_combined × K_pᵀ
+```
+
+The PE then computes local online-softmax state:
+
+```text
+m_p[row]
+l_p[row]
+o_p[row, :]
+```
+
+For 160 rows, each PE conceptually produces:
+
+```text
+m_p[160]
+l_p[160]
+o_p[160, d_v]
+```
+
+These may be processed in smaller row tiles for pipelining.
+
+## 3.4 Online softmax merge
+
+Each row has an independent online-softmax state.
+
+The reduction is always:
+
+```text
+same row index across different sequence shards
+```
+
+It is never:
+
+```text
+different query rows reduced together
+```
+
+Therefore 160 query rows do not imply 160 serialized communication rounds. The implementation exchanges vector or tiled payloads such as:
+
+```text
+m[160]
+l[160]
+o[160, d_v]
+```
+
+or:
+
+```text
+m[16]
+l[16]
+o[16, d_v]
+```
+
+The row states can be communicated and merged in parallel.
+
+## 3.5 Hierarchical reduction
+
+Because one KV head spans four CUBEs and 32 PEs, reduction is hierarchical.
+
+```text
+32 PE partial states
+        ↓
+8-PE reduction inside each CUBE
+        ↓
+4 CUBE-level states
+        ↓
+4-CUBE reduction
+        ↓
+Final attention outputs
+```
+
+The same online-softmax merge primitive is used at both levels.
+
+## 3.6 Prefill versus decode
+
+### Decode
+
+Decode typically has a very small number of Q rows.
+
+```text
+small Q
+   ↓ multicast
+32 PEs read local KV shards
+   ↓
+hierarchical reduction
+```
+
+The dominant concerns are usually KV bandwidth and reduction overhead.
+
+### Prefill
+
+Fan-out can make the query dimension much larger.
+
+```text
+Per-agent Q rows: 20
+Number of agents: 8
+Combined Q rows: 160
+```
+
+The larger `M` dimension can produce a much better GEMM shape. Agentic batching is therefore especially attractive for prefill or multi-token private suffix processing.
+
+## 3.7 Compute-cost clarification
+
+Replicating Q across 32 PEs does not multiply total attention FLOPs by 32. Each PE computes a different KV-column region.
+
+```text
+32 PEs × 160 rows × 256 local columns
+=
+160 rows × 8192 total columns
+```
+
+If Q is temporally tiled into four groups of 40 rows:
+
+```text
+4 × 32 PEs × 40 rows × 256 columns
+=
+160 rows × 8192 columns
+```
+
+The total arithmetic is identical. The policy changes Q traffic, reduction traffic, scheduling granularity, utilization, and buffering—not the mathematical amount of attention work.
+
+---
+
+# 4. Attention Execution Policies
+
+## 4.1 Policy A: Replicated Q, stationary KV
+
+Policy A keeps the existing KV placement unchanged.
+
+```text
+Q_combined
+    ↓ multicast
+PE0  computes against KV shard 0
+PE1  computes against KV shard 1
+...
+PE31 computes against KV shard 31
+```
+
+Each PE executes a local GEMM such as:
+
+```text
+(160 × d) × (d × 256)
+```
+
+### Advantages
+
+- No KV movement
+- No KV replication
+- No remote KV access
+- No page-table regrouping
+- Natural compatibility with sequence-parallel attention
+- Large local GEMM shapes
+- Simple hierarchical softmax reduction
+
+### Costs
+
+- Q must be distributed to all CUBEs
+- Every row requires a 32-way logical reduction
+- Large Q batches increase multicast and softmax-state traffic
+
+Policy A is the natural baseline for the assumed AHBM mapping.
+
+## 4.2 Policy B: Partitioned Q groups
+
+A possible alternative is to divide query rows among PE groups.
+
+```text
+Q group 0 → PE group 0
+Q group 1 → PE group 1
+...
+```
+
+However, every query row must still attend to the complete KV sequence. Because the 32 PEs already represent 32 sequence shards, spatially partitioning Q means each Q group must somehow access all KV shards.
+
+This requires one of the following:
+
+1. Regroup KV shards for each Q group.
+2. Read remote KV through symmetric memory.
+3. Replicate KV across Q-processing groups.
+4. Time-multiplex the same PEs over Q groups.
+
+The first three add memory-system complexity. The fourth is mainly temporal tiling and does not provide true spatial Q partitioning.
+
+## 4.3 Why Policy B is not automatically better
+
+The comparison is:
+
+```text
+Policy A cost
+=
+Q multicast
++
+hierarchical reduction
+```
+
+versus:
+
+```text
+Policy B cost
+=
+KV regrouping, replication, or remote access
++
+additional scheduling complexity
+```
+
+Policy B becomes attractive only if:
+
+```text
+Q distribution cost + reduction cost
+>
+remote or regrouped KV cost
+```
+
+Relevant factors include Q size, multicast bandwidth, reduction bandwidth, symmetric-memory bandwidth, remote-KV latency, KV replication capacity, page-table overhead, and GEMM utilization.
+
+## 4.4 Recommended baseline
+
+For:
+
+```text
+1 KV head = 4 CUBEs = 32 PEs
+```
+
+use:
+
+```text
+Policy A:
+replicated Q + stationary sequence-sharded KV
+```
+
+Policy B should be treated as an adaptive or future policy for cases where Q batches become extremely large, multicast becomes a bottleneck, reduction traffic dominates, or remote KV access becomes inexpensive.
+
+## 4.5 Runtime decision model
+
+A future runtime can estimate:
+
+```text
+T_A =
+T_Q_multicast
++
+T_local_GEMM
++
+T_hierarchical_reduce
+```
+
+and:
+
+```text
+T_B =
+T_Q_partition
++
+T_remote_or_replicated_KV
++
+T_local_GEMM
++
+T_group_reduce
++
+T_remap
+```
+
+The runtime selects the lower-cost policy for the current sequence length, agent count, Q size, KV-head mapping, bandwidth state, and interconnect congestion.
+
+---
+
+# 5. Fan-In and Join
+
+## 5.1 Fan-in problem
+
+After fan-out, each sub-agent produces a private result.
+
+```text
+Agent A → result A
+Agent B → result B
+Agent C → result C
+```
+
+The main agent must synthesize these results into a final continuation.
+
+## 5.2 Why branch KV cannot be concatenated
+
+Each branch has a different causal token history.
+
+```text
+A history:
+Shared prefix + A instruction + A reasoning
+
+B history:
+Shared prefix + B instruction + B reasoning
+
+C history:
+Shared prefix + C instruction + C reasoning
+```
+
+The K/V tensors of a token depend on token content, position, preceding causal context, and every transformer layer.
+
+Therefore:
+
+```text
+A private KV
++
+B private KV
++
+C private KV
+```
+
+does not form the KV cache of any valid single token sequence.
+
+## 5.3 Baseline text join
+
+The standard framework behavior is:
+
+```text
+Sub-agent result text
+        ↓
+Framework gathers and formats results
+        ↓
+Main-agent join prompt
+        ↓
+Continuation prefill
+        ↓
+New main-agent KV suffix
+```
+
+The framework normally aggregates the inputs, while the main LLM performs final synthesis.
+
+## 5.4 Main-agent KV after join
+
+The main-agent page table becomes:
+
+```text
+[Shared prefix KV pages]
++
+[Join-input KV pages]
++
+[Main continuation KV pages]
+```
+
+The private A/B/C pages remain separate and are not attached directly.
+
+## 5.5 Cost of full-text gather
+
+Assume:
+
+```text
+8 sub-agents
+1000 output tokens per agent
+```
+
+A full-text join creates:
+
+```text
+8000 join-input tokens
+```
+
+These tokens must pass through all transformer layers during continuation prefill. This increases prefill work, KV allocation, future decode KV reads, and context-window occupancy.
+
+## 5.6 Schema-constrained join
+
+A practical optimization is to constrain each sub-agent to a compact result schema.
+
+```json
+{
+  "claim": "memory_bandwidth_bottleneck",
+  "confidence": 0.91,
+  "evidence_ids": [17, 24]
+}
+```
+
+The framework can serialize it compactly:
+
+```text
+Finding: memory bandwidth bottleneck
+Confidence: 0.91
+Evidence: E17, E24
+```
+
+Example:
+
+```text
+Before:
+8 × 1000 = 8000 tokens
+
+After:
+8 × 50 = 400 tokens
+```
+
+The main LLM still performs final reasoning, but the continuation prefill is much shorter.
+
+## 5.7 Deduplication and aggregation
+
+Different sub-agents may produce overlapping findings.
+
+```json
+[
+  {"claim": "memory_bw", "confidence": 0.91},
+  {"claim": "memory_bw", "confidence": 0.87},
+  {"claim": "compute", "confidence": 0.42}
+]
+```
+
+The framework can combine duplicates:
+
+```text
+Primary finding:
+- Memory-bandwidth bottleneck
+- Supported by 2 agents
+- Maximum confidence: 0.91
+
+Alternative:
+- Compute bottleneck
+- Confidence: 0.42
+```
+
+This is preprocessing, not final reasoning.
+
+## 5.8 Pointer-based join
+
+Detailed evidence can remain outside the initial main-agent context.
+
+```text
+Agent A:
+- Finding: memory-bandwidth bottleneck
+- Evidence handle: E17
+
+Agent B:
+- Finding: reduction error
+- Evidence handle: E24
+```
+
+The main agent initially receives only summaries and handles. The framework retrieves and appends evidence only when requested.
+
+```text
+Effective input
+=
+summary tokens
++
+tokens for evidence actually used
+```
+
+The transformer does not directly dereference the handle; the framework resolves it through retrieval or a tool call.
+
+## 5.9 Hierarchical reduction
+
+For large fan-out width, local reducer agents can summarize groups of branches.
+
+```text
+16 sub-agents
+      ↓
+4 local reducers
+      ↓
+4 summaries
+      ↓
+Main agent
+```
+
+Benefits:
+
+- Reducers operate in parallel.
+- The final join prompt is shorter.
+- Main-agent KV growth is smaller.
+- Fan-in traffic is organized hierarchically.
+
+Costs:
+
+- Additional reducer inference
+- Possible loss of detail
+- Need for fallback evidence retrieval
+
+## 5.10 Latent-state join
+
+A more aggressive approach is to replace long text outputs with learned latent representations.
+
+```text
+Sub-agent hidden states
+        ↓
+Compression or projection
+        ↓
+Small latent-token set
+        ↓
+Main model
+```
+
+Example:
+
+```text
+1000 text tokens
+    ↓
+16 latent tokens
+```
+
+This could reduce join prefill and KV growth, but it requires training the main model to consume latent tokens, aligning branch representations, defining causal and positional semantics, and validating accuracy. It is a model-system co-design direction rather than a drop-in runtime optimization.
+
+## 5.11 KV and compute impact
+
+The shared parent prefix KV is already present. The primary optimization target is the newly created join suffix.
+
+```text
+New main KV size
+∝
+number of join-input tokens
+```
+
+Reducing join input reduces continuation-prefill work, new KV allocation, Q rows processed during join, later decode-time KV reads, and context-window consumption.
+
+A practical flow is:
+
+```text
+Sub-agent full outputs
+        ↓
+Schema-constrained results
+        ↓
+Deduplication and ranking
+        ↓
+Summaries + evidence handles
+        ↓
+Selective evidence materialization
+        ↓
+Main-agent continuation prefill
+```
+
+## 5.12 Recommended practical join design
+
+For an unchanged LLaMA-style model on AHBM:
+
+1. Reuse shared prefix KV through page-table references.
+2. Keep branch-private KV isolated.
+3. Require schema-constrained sub-agent outputs.
+4. Deduplicate repeated claims in the framework.
+5. Pass compact summaries, confidence values, and evidence handles.
+6. Retrieve detailed evidence only when requested.
+7. Introduce hierarchical reducer agents for wide fan-out.
+8. Keep the main LLM responsible for final synthesis.
+
+---
+
+# Current Baseline Summary
+
+```text
+Shared prefix prefill
+        ↓
+Shared KV page reuse
+        ↓
+Agentic fan-out
+        ↓
+Combine private Q rows
+        ↓
+Replicated-Q attention over stationary sequence-sharded KV
+        ↓
+Hierarchical online-softmax reduction
+        ↓
+Independent private branch continuations
+        ↓
+Schema/pointer-based fan-in
+        ↓
+Main-agent continuation prefill
+```
+
+Main design principles:
+
+- Reuse shared prefix KV without copying it.
+- Batch sub-agent Q rows when they read the same KV.
+- Keep KV stationary and multicast Q.
+- Reduce softmax state hierarchically by matching row index.
+- Do not directly merge branch-private KV.
+- Reduce fan-in cost by shortening and selectively materializing join inputs.

docs/report/1H-codesign-paper/figures/cube_architecture.svg

@@ -0,0 +1,312 @@
+<svg xmlns="http://www.w3.org/2000/svg" width="860" height="798" viewBox="55 2 860 798">
+  <title>cube</title>
+  <rect width="970" height="900" fill="#ffffff"/>
+  <text x="485" y="22" text-anchor="middle" font-family="monospace" font-size="18" font-weight="bold" fill="#1f2937">CUBE TOPOLOGY — 17.0×14.0mm | 6×6 Router Mesh | n_to_one mode | 64 pseudo-ch</text>
+  <text x="485" y="40" text-anchor="middle" font-family="monospace" font-size="15" fill="#ffffff">Per-PE: 8 ch × 32.0 GB/s = 256.0 GB/s | Cube total: 64 × 32.0 = 2048.0 GB/s</text>
+  <rect x="60" y="60" width="850.0" height="700.0" rx="6" fill="none" stroke="#94a3b8" stroke-width="2" stroke-dasharray="8,4"/>
+  <rect x="260" y="285" width="450" height="250" rx="6" fill="#ecfdf5" stroke="#047857" stroke-width="2" opacity="0.6"/>
+  <text x="485" y="395" text-anchor="middle" font-family="monospace" font-size="16" font-weight="bold" fill="#047857">HBM_CTRL | 64 pseudo channels</text>
+  <text x="485" y="412" text-anchor="middle" font-family="monospace" font-size="14" fill="#059669">Total BW: 2048 GB/s</text>
+  <rect x="270.0" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
+  <rect x="283.4" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
+  <rect x="296.9" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
+  <rect x="310.3" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
+  <rect x="323.8" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
+  <rect x="337.2" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
+  <rect x="350.6" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
+  <rect x="364.1" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
+  <rect x="377.5" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
+  <rect x="390.9" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
+  <rect x="404.4" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
+  <rect x="417.8" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
+  <rect x="431.2" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
+  <rect x="444.7" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
+  <rect x="458.1" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
+  <rect x="471.6" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
+  <rect x="485.0" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
+  <rect x="498.4" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
+  <rect x="511.9" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
+  <rect x="525.3" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
+  <rect x="538.8" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
+  <rect x="552.2" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
+  <rect x="565.6" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
+  <rect x="579.1" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
+  <rect x="592.5" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
+  <rect x="605.9" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
+  <rect x="619.4" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
+  <rect x="632.8" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
+  <rect x="646.2" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
+  <rect x="659.7" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
+  <rect x="673.1" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
+  <rect x="686.6" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
+  <text x="324" y="286" text-anchor="middle" font-family="monospace" font-size="12" fill="#3b82f6">PE0×8ch</text>
+  <text x="431" y="286" text-anchor="middle" font-family="monospace" font-size="12" fill="#60a5fa">PE1×8ch</text>
+  <text x="539" y="286" text-anchor="middle" font-family="monospace" font-size="12" fill="#475569">PE2×8ch</text>
+  <text x="646" y="286" text-anchor="middle" font-family="monospace" font-size="12" fill="#94a3b8">PE3×8ch</text>
+  <rect x="270.0" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
+  <rect x="283.4" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
+  <rect x="296.9" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
+  <rect x="310.3" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
+  <rect x="323.8" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
+  <rect x="337.2" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
+  <rect x="350.6" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
+  <rect x="364.1" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
+  <rect x="377.5" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
+  <rect x="390.9" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
+  <rect x="404.4" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
+  <rect x="417.8" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
+  <rect x="431.2" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
+  <rect x="444.7" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
+  <rect x="458.1" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
+  <rect x="471.6" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
+  <rect x="485.0" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
+  <rect x="498.4" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
+  <rect x="511.9" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
+  <rect x="525.3" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
+  <rect x="538.8" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
+  <rect x="552.2" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
+  <rect x="565.6" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
+  <rect x="579.1" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
+  <rect x="592.5" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
+  <rect x="605.9" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
+  <rect x="619.4" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
+  <rect x="632.8" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
+  <rect x="646.2" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
+  <rect x="659.7" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
+  <rect x="673.1" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
+  <rect x="686.6" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
+  <text x="324" y="539" text-anchor="middle" font-family="monospace" font-size="12" fill="#f59e0b">PE4×8ch</text>
+  <text x="431" y="539" text-anchor="middle" font-family="monospace" font-size="12" fill="#fbbf24">PE5×8ch</text>
+  <text x="539" y="539" text-anchor="middle" font-family="monospace" font-size="12" fill="#ef4444">PE6×8ch</text>
+  <text x="646" y="539" text-anchor="middle" font-family="monospace" font-size="12" fill="#f87171">PE7×8ch</text>
+  <line x1="135" y1="135" x2="285" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="135" x2="135" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="135" x2="435" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="135" x2="285" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="435" y1="135" x2="585" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="435" y1="135" x2="435" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="585" y1="135" x2="685" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="585" y1="135" x2="585" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="135" x2="835" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="135" x2="685" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="835" y1="135" x2="835" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="260" x2="285" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="260" x2="135" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="260" x2="435" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="260" x2="285" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="435" y1="260" x2="585" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="435" y1="260" x2="435" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="585" y1="260" x2="685" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="585" y1="260" x2="585" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="260" x2="835" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="260" x2="685" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="835" y1="260" x2="835" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="335" x2="285" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="335" x2="135" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="335" x2="685" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="335" x2="285" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="335" x2="835" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="335" x2="685" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="835" y1="335" x2="835" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="485" x2="285" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="485" x2="135" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="485" x2="685" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="485" x2="285" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="485" x2="835" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="485" x2="685" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="835" y1="485" x2="835" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="560" x2="285" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="560" x2="135" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="560" x2="435" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="560" x2="285" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="435" y1="560" x2="585" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="435" y1="560" x2="435" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="585" y1="560" x2="685" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="585" y1="560" x2="585" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="560" x2="835" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="560" x2="685" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="835" y1="560" x2="835" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="135" y1="685" x2="285" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="285" y1="685" x2="435" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="435" y1="685" x2="585" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="585" y1="685" x2="685" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <line x1="685" y1="685" x2="835" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
+  <circle cx="135" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="135" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c0</text>
+  <rect x="119" y="81" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
+  <text x="135" y="92" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE0</text>
+  <line x1="135" y1="127" x2="149" y2="97" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
+  <circle cx="285" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="285" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c1</text>
+  <rect x="269" y="81" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
+  <text x="285" y="92" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE1</text>
+  <line x1="285" y1="127" x2="299" y2="97" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
+  <circle cx="435" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="435" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c2</text>
+  <circle cx="585" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="585" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c3</text>
+  <circle cx="685" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="685" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c4</text>
+  <circle cx="835" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="835" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c5</text>
+  <circle cx="135" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="135" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c0</text>
+  <circle cx="285" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="285" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c1</text>
+  <circle cx="435" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="435" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c2</text>
+  <rect x="419" y="206" width="32" height="16" rx="3" fill="#ffffff" stroke="#f59e0b" stroke-width="1"/>
+  <text x="435" y="217" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#f59e0b">M_CPU</text>
+  <line x1="435" y1="252" x2="449" y2="222" stroke="#f59e0b" stroke-width="1" opacity="0.6"/>
+  <circle cx="585" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="585" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c3</text>
+  <circle cx="685" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="685" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c4</text>
+  <rect x="669" y="206" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
+  <text x="685" y="217" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE2</text>
+  <line x1="685" y1="252" x2="699" y2="222" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
+  <circle cx="835" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="835" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c5</text>
+  <rect x="819" y="206" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
+  <text x="835" y="217" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE3</text>
+  <line x1="835" y1="252" x2="849" y2="222" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
+  <circle cx="135" cy="335" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="135" y="338" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r2c0</text>
+  <circle cx="285" cy="335" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="285" y="338" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r2c1</text>
+  <circle cx="685" cy="335" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="685" y="338" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r2c4</text>
+  <circle cx="835" cy="335" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="835" y="338" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r2c5</text>
+  <circle cx="135" cy="485" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="135" y="488" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r3c0</text>
+  <rect x="119" y="523" width="32" height="16" rx="3" fill="#ffffff" stroke="#d97706" stroke-width="1"/>
+  <text x="135" y="534" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#d97706">SRAM</text>
+  <line x1="135" y1="493" x2="149" y2="523" stroke="#d97706" stroke-width="1" opacity="0.6"/>
+  <circle cx="285" cy="485" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="285" y="488" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r3c1</text>
+  <circle cx="685" cy="485" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="685" y="488" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r3c4</text>
+  <circle cx="835" cy="485" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="835" y="488" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r3c5</text>
+  <circle cx="135" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="135" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c0</text>
+  <rect x="119" y="598" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
+  <text x="135" y="609" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE4</text>
+  <line x1="135" y1="568" x2="149" y2="598" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
+  <circle cx="285" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="285" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c1</text>
+  <rect x="269" y="598" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
+  <text x="285" y="609" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE5</text>
+  <line x1="285" y1="568" x2="299" y2="598" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
+  <circle cx="435" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="435" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c2</text>
+  <circle cx="585" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="585" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c3</text>
+  <circle cx="685" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="685" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c4</text>
+  <circle cx="835" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="835" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c5</text>
+  <circle cx="135" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="135" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c0</text>
+  <circle cx="285" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="285" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c1</text>
+  <circle cx="435" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="435" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c2</text>
+  <circle cx="585" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="585" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c3</text>
+  <circle cx="685" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="685" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c4</text>
+  <rect x="669" y="723" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
+  <text x="685" y="734" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE6</text>
+  <line x1="685" y1="693" x2="699" y2="723" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
+  <circle cx="835" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
+  <text x="835" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c5</text>
+  <rect x="819" y="723" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
+  <text x="835" y="734" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE7</text>
+  <line x1="835" y1="693" x2="849" y2="723" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
+  <polyline points="135,143 208,216 251,216 324,289" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
+  <text x="239" y="216" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
+  <polyline points="285,143 358,216 358,216 431,289" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
+  <text x="368" y="216" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
+  <polyline points="685,268 674,278 549,278 539,289" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
+  <text x="622" y="278" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
+  <polyline points="835,268 824,278 657,278 646,289" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
+  <text x="751" y="278" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
+  <polyline points="135,552 146,542 313,542 324,531" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
+  <text x="239" y="542" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
+  <polyline points="285,552 296,542 421,542 431,531" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
+  <text x="368" y="542" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
+  <polyline points="685,677 612,604 612,604 539,531" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
+  <text x="622" y="604" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
+  <polyline points="835,677 762,604 719,604 646,531" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
+  <text x="751" y="604" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
+  <rect x="65" y="360" width="50" height="100" rx="3" fill="#ffffff" stroke="#475569" stroke-width="1.5" opacity="0.9"/>
+  <text x="90" y="357" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#475569">UCIe-W</text>
+  <rect x="67" y="362" width="46" height="23" rx="2" fill="#cbd5e1" opacity="0.7"/>
+  <text x="90" y="376" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c0</text>
+  <polyline points="127,135 120,142 120,366 113,374" fill="none" stroke="#cbd5e1" stroke-width="1" opacity="0.5"/>
+  <rect x="67" y="386" width="46" height="23" rx="2" fill="#94a3b8" opacity="0.7"/>
+  <text x="90" y="400" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c1</text>
+  <polyline points="127,260 120,267 120,390 113,398" fill="none" stroke="#94a3b8" stroke-width="1" opacity="0.5"/>
+  <rect x="67" y="410" width="46" height="23" rx="2" fill="#64748b" opacity="0.7"/>
+  <text x="90" y="424" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c2</text>
+  <polyline points="127,560 120,553 120,428 113,422" fill="none" stroke="#64748b" stroke-width="1" opacity="0.5"/>
+  <rect x="67" y="434" width="46" height="23" rx="2" fill="#374151" opacity="0.7"/>
+  <text x="90" y="448" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c3</text>
+  <polyline points="127,685 120,678 120,452 113,446" fill="none" stroke="#374151" stroke-width="1" opacity="0.5"/>
+  <rect x="435" y="65" width="100" height="50" rx="3" fill="#ffffff" stroke="#475569" stroke-width="1.5" opacity="0.9"/>
+  <text x="485" y="62" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#475569">UCIe-N</text>
+  <rect x="437" y="67" width="23" height="46" rx="2" fill="#cbd5e1" opacity="0.7"/>
+  <text x="448" y="93" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c0</text>
+  <polyline points="135,127 142,120 442,120 448,113" fill="none" stroke="#cbd5e1" stroke-width="1" opacity="0.5"/>
+  <rect x="461" y="67" width="23" height="46" rx="2" fill="#94a3b8" opacity="0.7"/>
+  <text x="472" y="93" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c1</text>
+  <polyline points="285,127 292,120 466,120 472,113" fill="none" stroke="#94a3b8" stroke-width="1" opacity="0.5"/>
+  <rect x="485" y="67" width="23" height="46" rx="2" fill="#64748b" opacity="0.7"/>
+  <text x="496" y="93" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c2</text>
+  <polyline points="685,127 678,120 504,120 496,113" fill="none" stroke="#64748b" stroke-width="1" opacity="0.5"/>
+  <rect x="509" y="67" width="23" height="46" rx="2" fill="#374151" opacity="0.7"/>
+  <text x="520" y="93" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c3</text>
+  <polyline points="835,127 828,120 528,120 520,113" fill="none" stroke="#374151" stroke-width="1" opacity="0.5"/>
+  <rect x="855" y="360" width="50" height="100" rx="3" fill="#ffffff" stroke="#475569" stroke-width="1.5" opacity="0.9"/>
+  <text x="880" y="357" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#475569">UCIe-E</text>
+  <rect x="857" y="362" width="46" height="23" rx="2" fill="#cbd5e1" opacity="0.7"/>
+  <text x="880" y="376" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c0</text>
+  <polyline points="843,135 850,142 850,367 857,374" fill="none" stroke="#cbd5e1" stroke-width="1" opacity="0.5"/>
+  <rect x="857" y="386" width="46" height="23" rx="2" fill="#94a3b8" opacity="0.7"/>
+  <text x="880" y="400" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c1</text>
+  <polyline points="843,260 850,267 850,391 857,398" fill="none" stroke="#94a3b8" stroke-width="1" opacity="0.5"/>
+  <rect x="857" y="410" width="46" height="23" rx="2" fill="#64748b" opacity="0.7"/>
+  <text x="880" y="424" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c2</text>
+  <polyline points="843,560 850,553 850,428 857,422" fill="none" stroke="#64748b" stroke-width="1" opacity="0.5"/>
+  <rect x="857" y="434" width="46" height="23" rx="2" fill="#374151" opacity="0.7"/>
+  <text x="880" y="448" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c3</text>
+  <polyline points="843,685 850,678 850,452 857,446" fill="none" stroke="#374151" stroke-width="1" opacity="0.5"/>
+  <rect x="435" y="705" width="100" height="50" rx="3" fill="#ffffff" stroke="#475569" stroke-width="1.5" opacity="0.9"/>
+  <text x="485" y="702" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#475569">UCIe-S</text>
+  <rect x="437" y="707" width="23" height="46" rx="2" fill="#cbd5e1" opacity="0.7"/>
+  <text x="448" y="733" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c0</text>
+  <polyline points="135,693 142,700 442,700 448,707" fill="none" stroke="#cbd5e1" stroke-width="1" opacity="0.5"/>
+  <rect x="461" y="707" width="23" height="46" rx="2" fill="#94a3b8" opacity="0.7"/>
+  <text x="472" y="733" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c1</text>
+  <polyline points="285,693 292,700 466,700 472,707" fill="none" stroke="#94a3b8" stroke-width="1" opacity="0.5"/>
+  <rect x="485" y="707" width="23" height="46" rx="2" fill="#64748b" opacity="0.7"/>
+  <text x="496" y="733" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c2</text>
+  <polyline points="685,693 678,700 504,700 496,707" fill="none" stroke="#64748b" stroke-width="1" opacity="0.5"/>
+  <rect x="509" y="707" width="23" height="46" rx="2" fill="#374151" opacity="0.7"/>
+  <text x="520" y="733" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c3</text>
+  <polyline points="835,693 828,700 528,700 520,707" fill="none" stroke="#374151" stroke-width="1" opacity="0.5"/>
+  <rect x="60" y="775" width="10" height="10" rx="2" fill="#3b82f6" stroke="#94a3b8" stroke-width="0.5"/>
+  <text x="74" y="784" font-family="monospace" font-size="13" fill="#1f2937">PE Router</text>
+  <rect x="147" y="775" width="10" height="10" rx="2" fill="#f59e0b" stroke="#94a3b8" stroke-width="0.5"/>
+  <text x="161" y="784" font-family="monospace" font-size="13" fill="#1f2937">M_CPU / SRAM</text>
+  <rect x="255" y="775" width="10" height="10" rx="2" fill="#475569" stroke="#94a3b8" stroke-width="0.5"/>
+  <text x="269" y="784" font-family="monospace" font-size="13" fill="#1f2937">UCIe</text>
+  <rect x="307" y="775" width="10" height="10" rx="2" fill="#1f2937" stroke="#94a3b8" stroke-width="0.5"/>
+  <text x="321" y="784" font-family="monospace" font-size="13" fill="#1f2937">Relay</text>
+  <rect x="366" y="775" width="10" height="10" rx="2" fill="#10b981" stroke="#94a3b8" stroke-width="0.5"/>
+  <text x="380" y="784" font-family="monospace" font-size="13" fill="#1f2937">HBM Link</text>
+  <rect x="446" y="775" width="10" height="10" rx="2" fill="#1f2937" stroke="#94a3b8" stroke-width="0.5"/>
+  <text x="460" y="784" font-family="monospace" font-size="13" fill="#1f2937">Mesh Link</text>
+</svg>

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
+      }
+    }
+  ]
+}