paper(gqa): single-rank decode-streaming study — composite's masked bandwidth win

Adds the memory-bound mirror of the compute-bound prefill figure. New single-rank decode sweep (T_q=1, M=8) reuses the prefill_compute_bound kernels, sweeping per-rank S_kv with NO cross-CUBE reduce — isolating the local attention that the 64-way Case-6 decode masks under its reduce tail.

Finding: even memory-bound decode benefits from the composite command. Its scheduler-streamed concurrent per-tile DMAs reach ~233 GB/s (91% of the 256 GB/s per-rank roofline) while the primitive's blocking tl.dot serializes one tile DMA and plateaus at ~166 GB/s — a 25-28% latency win that widens with context. This refines the paper's 'decode is latency-neutral' claim: neutral at 64-way production scale (reduce-dominated), but the composite extracts bandwidth at the local level. The bandwidth roofline here mirrors the MAC roofline in prefill.

Adds: gqa_decode_streaming.py sweep (wired into the milestone umbrella as GQA_1H_SWEEPS=decode_streaming), paper_plot_gqa_decode_streaming.py, the figure, and a new fig:gqa-decode-stream + paragraph in 05-gqa.tex; rebuilt main.pdf. Prefill (Experiment B) re-verified byte-identical (794.1/668.1/646.9us).

Co-Authored-By: Claude Opus 4.8 (1M context) <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.
+```

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` 완전 제거
 
@@ -113,9 +115,9 @@ 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
-                pe=ls.local_pe,                   # local within cube (explicit name)
+                sip=target_sip,                              # from caller (current_device)
+                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-0062-prog-tl-async-load.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Proposed
+Accepted
 
 > **ADR-0060**(AHBM GQA Fused Attention)의 보조 ADR. Decode와 long-context
 > 어텐션은 **KV-load-bound**이며, load/compute 오버랩이 지배적 레버다. 본

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-proposed/ADR-0060-algo-gqa-fused-attention-ahbm-ko.md

@@ -1,747 +0,0 @@
-# ADR-0060: AHBM GQA Fused Attention 커널 (Llama3-70B)
-
-## Status
-
-Proposed
-
-**Context model:** Llama3-70B.
-**Decision drivers:** agentic workload → 낮은 batch, 긴 context;
-KV-load-bound decode; long-context prefill용 sequence-parallel (Ring KV).
-
-**Supersedes / extends:** 기존 mesh-native 어텐션 커널
-`_attention_mesh_kv`(prefill)와 `_attention_mesh_mlo`(decode), 그리고
-`milestone-gqa-llama70b` eval bench. *§A. 기존 kernbench 작업과의 관계* 참조 —
-그 코드가 본 ADR이 진짜 GQA·causal·long-context 커널로 업그레이드하는
-baseline이다.
-
-**Supporting ADRs** (efficiency / scale enabler — *GQA blocker 아님*; §8 정정
-참조): **ADR-0063** `tl.scratch_scope`(per-tile scratch 재활용 — 현실적 context
-길이에 필요), **ADR-0062** lazy `tl.load`(첫 사용 시점 auto-wait를 갖는
-non-blocking load → load/compute 오버랩), **ADR-0061** `tl.broadcast`(선택적
-mask/범용 편의). 두 GEMM(Q·Kᵀ, P·V)은 scheduler가 관리하는 `tl.composite`
-커맨드로 발행된다(기존 `CompositeCmd`; 새 커맨드 종류 없음); 진짜 GQA 자체는
-커널 재구조화만 필요(§5.2).
-
-**Algorithm lineage.** 이 커널은 **FlashAttention**(tiling + online/streaming
-softmax, P·V fused — full score matrix 미생성)이다. §4의 KV-parallel
-split-and-combine는 **FlashDecoding**(split-KV + log-sum-exp merge). §5.5의 Ring
-경로는 **Ring Attention**(KV 블록을 mesh 주위로 회전, 같은 online softmax로
-fold). 새 수학은 도입하지 않으며, 본 ADR은 이 알려진 알고리즘을 kernbench의
-**greenlet `tl` 프로그래밍 모델**(ADR-0020, ADR-0046)과 **IPCQ** PE↔PE
-collective(ADR-0023/0025)에 매핑한다.
-
----
-
-## TL;DR — 최종 GQA 커널 (composite hybrid) pseudocode
-
-결정(§1)을 한 곳에: **GEMM → `tl.composite`; softmax 머지 + tree reduction →
-커널; `tl.load`은 lazy.** KV head별로, `G`를 matmul M 차원에 fold(진짜 GQA,
-broadcast 없음). 이것이 reference 형태이며, 산문 섹션이 각 부분을 부연한다.
-
-```python
-def gqa_attention(q_ptr, k_ptr, v_ptr, o_ptr,
-                  counter, start_pe, N, q_block, scale, *, tl):
-    # ---- geometry (kernel arithmetic, §2) ----
-    pe_id   = tl.program_id(axis=rank_axis)
-    G       = H_q // H_kv                       # query heads per KV head (=8)
-    causal  = q_block.is_prefill
-
-    for kv in range(H_kv):                       # one KV head per iteration (§5.2)
-        # Q group: G rows folded into M → [G·T_q, d]. Lazy load (ADR-0062):
-        # issue now, auto-wait at first use inside the first composite.
-        q_g     = tl.load(q_base(kv), (G * q_block.T_q, d))      # [G·T_q, d]
-        my_len  = valid_len(counter, start_pe, pe_id, N) if not causal else S_kv
-        n_tiles = ceil(my_len / TILE)
-
-        # persistent arena (outside scratch_scope, ADR-0063): -inf, 0, zeros
-        m, l, O = init_running(G * q_block.T_q, d)
-
-        for j in range(n_tiles):
-            if causal and tile_all_future(j, q_block):
-                continue                                         # causal skip (kernel if)
-            with tl.scratch_scope():                             # per-tile temporaries (ADR-0063)
-                # --- GEMM #1 on the scheduler: Q·Kⱼᵀ; K streamed by composite ---
-                Sj = tl.composite("gemm", a=q_g,
-                                  b=tl.ref(k_tile(kv, j), (d, TILE))) * scale   # Kᵀ pre-stored [d,TILE]
-                if causal and tile_partial(j, q_block):
-                    Sj = Sj + causal_mask(j, q_block)            # additive boundary mask
-                # --- online softmax merge in the kernel (MATH ops) ---
-                m_new = tl.maximum(m, tl.max(Sj, axis=-1))
-                P     = tl.exp(Sj - m_new)
-                corr  = tl.exp(m - m_new)
-                l     = l * corr + tl.sum(P, axis=-1)
-                # --- GEMM #2 on the scheduler: P·Vⱼ; V streamed by composite ---
-                Oj    = tl.composite("gemm", a=P,
-                                     b=tl.ref(v_tile(kv, j), (TILE, d)))
-                O     = O * corr + Oj                            # running merge
-                m     = m_new
-
-        # ---- cross-PE combine (§4): log-sum-exp tree to root, or store ----
-        if N == 1:
-            tl.store(o_base(kv), O / l)
-        else:
-            tree_reduce_and_store(m, l, O, pe_id, N, o_base(kv))  # tl.send/tl.recv
-```
-
-> **이 형태인 이유:** 두 `tl.composite` 호출이 tiling + K/V DMA 스트리밍 +
-> cross-tile 파이프라이닝을 PE_SCHEDULER에 offload한다(CPU는 coarse descriptor를
-> 내고 앞서 나가 → 엔진이 saturate 유지, §1); softmax 머지와 IPCQ tree reduction은
-> 커널에 남는다(기존 `CompositeCmd`가 cross-tile 상태를 못 들기 때문). `K`는
-> reshape-not-transpose caveat를 피하려 transpose된 `[d, TILE]`로 미리 저장된다
-> (§3, §B). 전용 "flash-composite" 커맨드 종류는 도입하지 **않는다**(§8 항목 4).
-
----
-
-## A. 기존 kernbench 작업과의 관계 (먼저 읽을 것)
-
-kernbench는 **오늘날 이미 IPCQ 상에서 online-softmax `(m, ℓ, O)` 머지로
-FlashAttention을 돌린다.** 두 커널이 존재한다:
-
-| File | Role | Mechanism |
-|---|---|---|
-| `src/kernbench/benches/_attention_mesh_kv.py` | prefill (Ring K/V) | per-rank partial attention, bidirectional K/V fan-out, online-softmax fold |
-| `src/kernbench/benches/_attention_mesh_mlo.py` | decode (split-KV) | per-rank one-shot partial attention, bidirectional `(m,ℓ,O)` fan-out, log-sum-exp merge |
-
-둘 다 `milestone-gqa-llama70b`(4 패널:
-`{single,multi}_user × {prefill,decode}`,
-`src/kernbench/benches/milestone_gqa_llama70b.py`)이 구동하며
-`tests/attention/test_milestone_gqa_llama70b.py`에서 테스트된다.
-
-**이들은 greenlet `tl` API로 작성되었다:** `tl.load`, `tl.dot`,
-`tl.softmax`/`tl.max`/`tl.sum`/`tl.exp`, `tl.send`/`tl.recv`, 그리고
-`TensorHandle`에 대한 Python `-`/`*`/`/`(각각 `MathCmd` emit) — GEMM은 composite가
-아니라 **blocking `tl.dot`**로. 실행 중 `(m, ℓ, O)`는 루프를 관통하는 Python
-`TensorHandle`일 뿐이다. 본 ADR은 실행 상태와 softmax 머지는 커널에 유지하되
-**두 GEMM은 scheduler가 관리하는 `tl.composite` 경로로 옮긴다**(§1 참조) — 이것이
-**설계에 중요**하다.
-
-**baseline의 의도된 세 한계** — 효율적 GQA 커널이 정확히 들어내야 하는 것:
-
-1. **GQA 재사용 없음.** `h_q == h_kv == 1`
-   (`test_milestone_gqa_llama70b.py:137-142`). 테스트는 이를 *broadcast view*에
-   대한 MemoryStore byte-conservation 실패로 돌리지만, 그 실패는 baseline의
-   **head-packing 핵**의 속성이다(`_view(K, (h_q·d, S_kv))`는 모든 head를 하나의
-   matmul 차원에 뭉치고 `h_q == h_kv`일 때만 byte를 보존). 올바른 수정은
-   broadcast op이 아니라 **커널 재구조화**다: **한 번에 한 KV head**를 처리하고
-   `G` group 행을 matmul **M** 차원에 fold(§5.2). 이는 byte-보존 reshape만 쓰므로
-   진짜 GQA(`h_q = G·h_kv`)가 **신규 primitive 없이** 돈다 — §8 참조.
-2. **O(N) reduction.** baseline은 all-to-all bidirectional fan-out을 하여 *모든*
-   rank가 full 답을 갖는다(`n_ranks − 1` 단계). 어텐션은 query owner에서만 `O`가
-   필요 → root로의 **tree reduction**은 `⌈log₂ N⌉` 단계(§4).
-3. **검증 스케일만.** `S = 16`인 이유는 1 MiB scratch bump allocator가 per-tile
-   임시값을 누수하고(`test_milestone_gqa_llama70b.py:123-148`) causal masking /
-   tiling이 없기 때문 → **ADR-0063**(재활용) + §5(tiling, causal skip) + composite
-   K/V 스트리밍(§3) + **ADR-0062**(lazy load 오버랩)으로 해결.
-
-**문서 부채(범위 밖이나 기록):** baseline은 ADR-0055/0056/0057/0058/0059를
-인용하나 **파일로 존재하지 않는다** — ghost 참조다. 본 ADR은 이를 소급
-작성하지 않으며; 권고는 Detailed Design Document의 *Open Decisions* 참조.
-
----
-
-## 0. 참조 차원 (Llama3-70B)
-
-| Symbol | Meaning | Value |
-|---|---|---|
-| `H_q` | query heads | 64 |
-| `H_kv` | KV heads | 8 |
-| `G` | GQA group size = `H_q / H_kv` | 8 |
-| `d` | head dim | 128 |
-| `L` | layers | 80 |
-| `D` | model dim | 8192 |
-
-하드웨어 요약:
-- **AHBM (chip)** = **CUBE**(메모리 cube, 각각 **PE**를 담은 logic die) 집합 +
-  **IO die**(ADR-0003).
-- **IPCQ**: PE↔PE 큐, PE당 4 mesh-방향 queue-pair
-  (`N/S/E/W`, ADR-0023 D3; inter-SIP용 `global_*`, ADR-0032). 커널
-  API: `tl.send(dir, src)` / `tl.recv(dir, shape, dtype)`
-  (`tl_context.py:402-499`).
-- **Composite command** (`CompositeCmd`, `pe_commands.py:144-162`): 단일
-  GEMM(또는 MATH) *head* + element-wise *epilogue* 단계
-  (`bias/relu/scale/add/...`), PE_SCHEDULER에 **non-blocking**으로 발행되며,
-  scheduler가 tile plan을 생성하고 타일당 DMA→GEMM→write를 스트리밍한다
-  (ADR-0014 D6; `pe_scheduler.py:104-143`). 이는 일반적 multi-op DAG가 **아니다**:
-  두 GEMM을 chain할 수 없고, 인스턴스 간 register 상태를 못 들며, IPCQ를
-  pop/wait 못 한다. 따라서 본 ADR은 두 어텐션 GEMM(Q·Kᵀ, P·V)을 **각각** 자체
-  composite로 발행하고 cross-GEMM softmax 머지 + IPCQ reduction은 커널에 유지한다
-  — 새 "flash-composite" 커맨드 종류는 **필요 없다**(§1, §8 참조).
-- **Allocation policy (SP):** multi-user 패널에서 CUBE당 query head 하나; KV
-  group의 `G` query head는 한 CUBE의 PE 안에 매핑.
-
-**직교하는 두 매핑 레이어** (round-robin placement):
-- **Layer 1 (KV-parallel, intra-request):** KV token `i`는 PE
-  `(start_pe + i) mod N`에 안착 — round-robin이라 한 긴 request가 `N` PE에 걸쳐
-  분할, ≤1 token으로 균형.
-- **Layer 2 (inter-request):** `start_pe = request_id mod N`이 request별 "PE-0
-  역할"을 회전.
-- 결합: `pe = (request_id + global_token_idx) mod N`.
-
-`N` = 하나의 (query head, request)에 대한 KV-parallel reduction group의 PE 수.
-Reduction은 **intra-CUBE** 유지(query head는 CUBE를 가로지르지 않음 — 명시적
-non-goal).
-
----
-
-## 0.5 커널 경계, 전제조건, I/O 계약
-
-### 0.5.1 디코더 레이어 내 위치
-
-```
-1. RMSNorm
-2. QKV projection (GEMM)        ─┐  qkv_rope kernel (SEPARATE, upstream)
-3. RoPE on Q and K              ─┤
-4. write new K,V → KV cache     ─┘
-5. ===== THIS KERNEL: FlashAttention =====   (post-RoPE Q, K-cache, V-cache → O)
-6. Output projection (GEMM)        out_proj kernel (SEPARATE, downstream)
-7. residual add → FFN ...
-```
-
-**전제조건 (upstream `qkv_rope`, 본 커널 아님):**
-- **P1.** Q는 이미 RoPE-회전됨. 여기서 회전 없음.
-- **P2.** K-cache는 **post-RoPE** K 저장(어텐션 시점에 재회전 안 함 — post-RoPE
-  캐싱의 이유).
-- **P3.** decode의 경우, 새 step의 K 행은 RoPE-회전되어 그 소유 PE의 K-cache
-  슬롯에 **본 커널 launch 전에 `qkv_rope`가** 추가한다. ⇒ 본 커널은 KV cache에
-  대해 **pure read**.
-- **P4.** V는 회전 안 함; V-cache는 raw projected V 보유.
-- **P5.** upstream RoPE 위치는 token의 **절대 global 위치**
-  `global_idx = local_slot·N + ((pe_id − start_pe) mod N)`(§2.1). Round-robin
-  placement ≠ RoPE 위치.
-
-### 0.5.2 Shape 기호
-`T_q` = 이번 launch의 query 길이(decode: 1; prefill/chunk: chunk 폭).
-`S` = 전체 context 길이. `S_pe` = 이 PE 소유 key ≈ `⌈S/N⌉`.
-저장 `bf16`(numpy proxy `f16`, `memory_store.py:16`); `m,ℓ,O` 누산기는 정밀도가
-중요한 곳에서 `f32`.
-
-### 0.5.3 INPUTS (커널 launch당)
-
-| Input | Shape (per KV head) | Location | Notes |
-|---|---|---|---|
-| `Q` | `[G, T_q, d]` | per-PE HBM (loaded to TCM) | post-RoPE (P1). group의 `G` query 행 batched. |
-| `K_cache` | `[S_pe, d]` | per-PE HBM, base `K_base[pe]`, contiguous | post-RoPE (P2). Read-only. 로컬은 dense, global index는 strided (§2.1). |
-| `V_cache` | `[S_pe, d]` | per-PE HBM, base `V_base[pe]` | raw V (P4). Read-only. |
-| `global_token_counter` | scalar | launch arg | 커널이 `S_pe`, slot↔global, causal bound 도출. |
-| `start_pe` (= `request_id mod N`) | scalar | launch arg | Layer-2 회전. |
-| `pe_id`, `N` | scalars | launch / `tl.program_id` | reduction-group geometry. |
-| `q_block_meta` `{q_start, T_q}` | launch arg | prefill/SP causal masking & skip. |
-| `O_base` | address | launch arg | 최종 O가 쓰이는 곳 (root PE만). |
-| `softmax_scale` | scalar | launch arg | `1/√d`. |
-
-**Ring Attention (§5.5)**에는 ring step마다 추가: IPCQ로 들어오는 `K_block,
-V_block`을 ping-pong 버퍼에(post-RoPE), 그리고 causal step-skip용
-`step_kv_global_range`.
-
-### 0.5.4 OUTPUTS
-
-| Output | Shape | Location | Notes |
-|---|---|---|---|
-| `O` (final) | `[G, T_q, d]` | `O_base`, **root PE만** | 정규화 `O_acc / ℓ_acc`, 저장 dtype으로 cast. |
-
-**중간값 (non-root PE, IPCQ 상 — 커널-가시 출력 아님):**
-`(m_i, ℓ_i, O_i)` (`m,ℓ`: `[G, T_q]`; `O_i`: `[G, T_q, d]` 비정규화)이
-`tl.send`로 `tree_parent`에 push(§4). `O_i`가 무거운 부분.
-
-**No-SP (`N=1`):** IPCQ partial 없음; 단일 PE의 실행 중 `(m,ℓ,O)`를 제자리에서
-정규화하여 `O_base`에 기록.
-
-**명시적 비-출력:** KV-cache write(upstream, P3), output projection(downstream),
-score `S` 및 prob `P`(미생성).
-
----
-
-## 1. 결정 (메커니즘)
-
-**커널을 *하이브리드*로 구현한다: 두 GEMM(Q·Kᵀ, P·V)을 scheduler가 관리하는
-`tl.composite(op="gemm")` 커맨드로 발행하고, online-softmax 머지와 cross-PE
-reduction은 커널 수준 `tl` op으로 유지한다.** `tl.load`은 **lazy**다(non-blocking;
-wait는 load된 데이터의 첫 사용 시점에 자동 삽입 — ADR-0062). 따라서 명시적 HBM
-load가 뒤따르는 compute와 오버랩된다. kernbench의 실행 + latency 모델에 근거한
-근거:
-
-- **GEMM tiling이 PE_SCHEDULER에 offload된다.** `CompositeCmd`는 non-blocking
-  (`kernel_runner.py:182-191`, `pe_scheduler.py:104-121`): 커널이 **하나의 coarse
-  descriptor**(M = `G·T_q`, PE당 전체 tile sweep)를 push하면 scheduler가 tile
-  plan을 생성하고 타일당 DMA→GEMM→write를 스트리밍한다(ADR-0014 D6). K/V는
-  scheduler가 HBM에서 스트리밍하는 `tl.ref` operand이므로, 타일당 **K/V prefetch는
-  scheduler의 일**이다 — 명시적 prefetch op 없음. CPU(greenlet)는 현재 composite가
-  도는 동안 **다음** composite를 낼 수 있게 풀려나, scheduler가 타일을 가로질러
-  GEMM 엔진을 saturate 유지한다.
-- **이는 하드웨어를 반영**하며 CPU issue-rate를 execution-rate로부터 decouple한다.
-  대조적으로 blocking per-op `tl.dot` 경로는 매 GEMM마다 CPU를 멈추고, 끼어드는
-  softmax MATH op 동안 GEMM 엔진 **버블**을 남긴다; CPU가 fine-grained per-tile
-  발행을 따라잡을 수 있을 때만 현실적이다.
-- **실행 중 `(m, ℓ, O)` flash 상태는 Python `TensorHandle`로** 루프를 관통한다
-  (baseline이 이미 그렇게 함); softmax 머지(max/exp/sum/rescale)는 두 GEMM
-  composite **사이**의 커널 수준 `tl` MATH다. 기존 `CompositeCmd`는 두 GEMM을
-  chain하거나 cross-tile register 상태를 못 들므로(§0), 머지는 필연적으로 커널에
-  산다 — 이것이 하이브리드 분할이지, 우회한 한계가 아니다.
-- **Cross-PE 결합**은 P·V 후 `(m, ℓ, O)`에 대한 log-sum-exp **tree**로, 커널 수준
-  `tl.send`/`tl.recv`(§4)를 통해 — 불변.
-
-그래서 per-tile 내부 파이프라인은:
-
-```
-q_g = tl.load(Q group)                         # lazy; auto-wait at first use
-per tile j:
-  Sⱼ = tl.composite("gemm", a=q_g, b=tl.ref(Kⱼ))  → Sⱼ   # scheduler streams Kⱼ DMA + GEMM
-  Sⱼ += maskⱼ                                          # kernel MATH, boundary tile only
-  online-softmax: mⱼ, m_new, P, corr, ℓ               # kernel MATH
-  Oⱼ = tl.composite("gemm", a=P, b=tl.ref(Vⱼ))   → Oⱼ   # scheduler streams Vⱼ DMA + GEMM
-  O = O*corr + Oⱼ;  m = m_new                         # kernel MATH (running merge)
-```
-
-각 타일의 MATH 임시값을 `tl.scratch_scope`(ADR-0063)로 감싸 scratch를 O(1)로
-유지하고, 다음 타일의 composite를 현재 타일 결과를 wait하기 전에 발행
-(non-blocking 핸들)하여 scheduler가 타일을 가로질러 파이프라인되게 한다.
-
-제어 흐름(tile skip, mask 생성, reduction 스케줄링, 주소 산술)은 **커널**에 산다
-(greenlet 본문의 평범한 Python `if`/산술). 이는 kernbench의 greenlet 모델이 이미
-허용하는 바 그대로다(`kernel_runner.py`, ADR-0020 D3).
-
-> **이것이 무엇을 대체하나.** 이전 iteration은 "composite는 latency 이점 없음"이라는
-> 논거로 순수 greenlet primitive 경로(전부 `tl.dot`, composite 없음)를 제안했다.
-> 그것은 **오직** 시뮬레이터가 현재 per-op CPU 발행 비용을 **0**으로 청구하기
-> 때문에만(`dispatch_cycles=0`, `pe_cpu.py`) 성립한다 — descriptor offload가 숨기려
-> 존재하는 바로 그 CPU issue-rate / DMA-program 비용을 모델이 지워버린 것이다.
-> 하이브리드가 효율 커널의 충실한 표현이다. 이점의 **측정 가능한** 크기(CPU가 많은
-> 타일에 대해 엔진을 saturate할 수 있는가?)는 op 종류별 발행 비용 모델링에
-> 좌우되며, future work로 추적된다(cost model; §9). `dispatch_cycles=0`에서도
-> non-blocking composite 경로는 blocking `tl.dot` 경로가 남기는 GEMM 엔진 버블을
-> 채운다.
->
-> **이것이 효율적 선택인 이유.** GEMM tiling + DMA 스트리밍 + cross-tile
-> 파이프라이닝은 검증된 `CompositeCmd` scheduler 경로에 offload되고; softmax 머지와
-> 검증된 IPCQ collective는 커널에 남는다. 진짜 새 기계장치는 작고 범용적인 두
-> primitive뿐이다(ADR-0062 lazy `tl.load`, ADR-0063 `tl.scratch_scope`); reduction은
-> `tl.send`/`tl.recv`를 재사용한다. 전용 "flash-composite" 커맨드(softmax 머지 +
-> carried register 상태 + IPCQ-push epilogue를 내재화하는 한 종류)는 만들지
-> **않는다** — 크고 특수목적이며, 하이브리드 대비 유일한 delta(완전 softmax
-> offload)가 현재 모델링 충실도에서 정당화되지 않는다; §8 참조.
-
----
-
-## 2. 메모리 레이아웃 & 드라이버 책임
-
-### 2.1 KV cache 할당
-- per-PE KV 버퍼는 per-PE 최대 context `⌈max_context / N⌉ × d × dtype`로 K와 V를
-  각각 sizing. kernbench에서는 sequence 차원에 대해 `pe`(또는 `cube`) 축이
-  `row_wise`인 `DPPolicy`로 배치된다(`policy/placement/dp.py`; baseline은 K/V에
-  `DPPolicy(pe="row_wise")`, Q에 `replicate`).
-- PE 안에서 할당된 token은 **연속 append**된다(slot 0,1,2,…). global index는
-  strided(`i, i+N, …`)이나 per-PE 버퍼는 dense ⇒ DMA가 연속 유지.
-- Slot → global: `global_idx = local_slot·N + ((pe_id − start_pe) mod N)`.
-
-### 2.2 드라이버 launch당 의무 (최소)
-드라이버는 launch마다 base + counter + rotation을 공급하고; 커널이 나머지를
-도출한다:
-
-| Launch arg | Purpose |
-|---|---|
-| `K_base[pe]`, `V_base[pe]` | per-PE KV 버퍼 base (tensor VA에서) |
-| `O_base` | 어텐션 출력 목적지 |
-| `global_token_counter` | 현재 sequence 위치 |
-| `start_pe = request_id mod N` | Layer-2 회전 |
-| `pe_id` (`tl.program_id`), `N` | geometry |
-| `q_block_meta` | prefill/SP query block 시작 + 길이 |
-
-커널 도출(평범한 산술):
-- **이번 step에 내가 쓸 차례:** `(start_pe + counter) mod N == pe_id`
-- **내 valid 길이:** `base = counter // N; rem = counter % N;
-  my_len = base + (1 if ((pe_id − start_pe) mod N) < rem else 0)`
-- **read 범위:** tile `0 .. ⌈my_len / TILE⌉`.
-- **causal bound / per-tile skip:** query block과 각 tile의 global 위치에서.
-
-드라이버 = base + counter + rotation. 단일 공식
-`(request_id + token_idx) mod N`이 placement policy 전부다.
-
----
-
-## 3. Per-tile op 시퀀스 (greenlet `tl`)
-
-한 iteration = 한 PE의 한 KV tile. 두 GEMM은 `tl.composite(op="gemm")`
-(scheduler 관리 tiling + K/V DMA 스트리밍); softmax 머지는 그 사이의 커널 `tl`
-MATH다. 실제 `tl` 이름(`tl_context.py`), lazy `tl.load`(ADR-0062),
-`tl.scratch_scope`(ADR-0063) 사용:
-
-```python
-# running state (persistent arena — allocated once, outside the scope)
-# m: [G, T_q]  l: [G, T_q]  O: [G, T_q, d]
-q_g = tl.load(Q_group_ptr, (G*T_q, d))               # lazy; auto-wait at first use (ADR-0062)
-
-with tl.scratch_scope():                              # per-tile MATH temporaries recycled
-    Sj    = tl.composite("gemm", a=q_g,                # [G·T_q, TILE]; scheduler streams Kⱼ DMA
-                         b=tl.ref(K_base + j*TILE*d, (TILE, d))) * softmax_scale
-    if mask_j is not None:
-        Sj = Sj + mask_j                              # additive causal mask (boundary tile)
-    m_j   = tl.max(Sj, axis=-1)
-    m_new = tl.maximum(m, m_j)
-    P     = tl.exp(Sj - m_new)                        # no full-matrix softmax; streaming
-    corr  = tl.exp(m - m_new)                         # rescale factor for old accumulators
-    l     = l * corr + tl.sum(P, axis=-1)
-    Oj    = tl.composite("gemm", a=P,                  # [G·T_q, d]; scheduler streams Vⱼ DMA
-                         b=tl.ref(V_base + j*TILE*d, (TILE, d)))
-    O     = O * corr + Oj                             # running merge (kernel MATH)
-    m     = m_new
-```
-
-Notes:
-- `q_g`는 `[G·T_q, d]`로 reshape된 GQA-batched query(`G` group 행을 matmul M
-  차원에 fold; byte-보존). 한 K/V tile이 모든 `G·T_q` 행을 서비스 — GQA 재사용
-  레버 — broadcast 없이.
-- **K/V는 `tl.ref` operand**로 composite scheduler가 HBM에서 타일당 스트리밍한다
-  (`pe_scheduler.py:104-143`): 그것이 *바로* prefetch/파이프라인이므로 명시적
-  prefetch op이 없다. 타일 `j+1`의 composite를 타일 `j`를 wait하기 전에 발행
-  (non-blocking 핸들)하면 scheduler가 타일을 가로질러 파이프라인되고 GEMM 엔진이
-  saturate 유지된다.
-- `tl.trans`는 kernbench에서 **메타데이터 전용**이고(`tl_context.py:390`)
-  `MemoryStore.read`는 transpose가 아니라 *reshape*한다(`memory_store.py:73`).
-  zero/structural 실행엔 무해하나; 비자명 수치 데이터에선 reshape-not-transpose가
-  되어 transpose된 K로의 Q·Kᵀ는 주의가 필요하다(§11) — K를 transpose해 `[d, TILE]`로
-  미리 저장하거나, 실제 `tl.transpose`를 추가(후보 primitive, 시뮬레이터의
-  성능-모델링 목적상 아마 불필요).
-- Masking: 커널이 query/KV global offset에서 boundary-tile mask를 만들어 더한다
-  (`Sj + mask_j`); full-past tile은 `None` 전달; full-future tile은 **skip**(`if`이
-  enqueue 안 함).
-- **새** "flash-composite" 커맨드 종류(softmax 머지 + carried `(m,ℓ,O)`를
-  내재화하는 것)는 쓰지 **않는다**; 기존 `CompositeCmd`가 각 GEMM을 담당하고 머지는
-  커널에 남는다(§1, §8 항목 4).
-
----
-
-## 4. Reduction (KV-parallel / SP combine) — root로의 tree
-
-tile sweep 후 각 PE는 `(m_i, ℓ_i, O_i)`(비정규화)를 가진다. 결합/교환 가능한
-log-sum-exp 머지로 결합(baseline의 fold와 동일 수학, `_attention_mesh_mlo.py:117-122`):
-
-```python
-def merge(m_a, l_a, O_a, m_b, l_b, O_b):
-    m = tl.maximum(m_a, m_b)
-    sa, sb = tl.exp(m_a - m), tl.exp(m_b - m)
-    return m, l_a*sa + l_b*sb, O_a*sa + O_b*sb
-# final: O = O_root / l_root      # normalise once at the tree root
-```
-
-- **타이밍:** 교환은 **P·V 후**(각 PE가 최종 `O_i`를 가짐).
-- **토폴로지:** mesh tree, depth `⌈log₂ N⌉`, 물리 이웃 상. `N=8`의 경우: level-0
-  `(0↔1,2↔3,4↔5,6↔7)`, level-1 `(1↔3,5↔7)`, level-2 `(3↔7)`, root = `7`. 각 tree
-  pair는 물리 `N/S/E/W` 이웃이어야 `tl.send(dir)`/`tl.recv(dir)`이 실제 link를
-  쓴다(튜닝 항목 §9; 이웃을 wiring하는 SFR install은
-  `configure_sfr_intracube_pe_ring`, baseline이 사용).
-- **왜 baseline fan-out이 아니라 tree인가:** 어텐션은 한 곳(`O_base`를 쓰는 PE)에서
-  `O`가 필요하다. tree는 `⌈log₂ N⌉` 단계(`N=8`에 3) vs baseline all-to-all의 `N−1`
-  (7). per-PE sweep이 짧은 decode에서 reduction이 지배적 비용이므로 실제 이득이다.
-  (baseline fan-out은 답을 모든 rank에 복제 — 여기선 불필요.)
-- **Payload:** `O_i`(`d`-vector)가 무겁고; `m,ℓ`은 `(g, T_q)`당 scalar. 별도
-  `tl.send`로 전송(baseline은 triplet을 세 send로 — 같은 패턴).
-
-커널 구조(greenlet):
-
-```python
-# leaf / internal node: fold children that send to me, then forward up
-for child_dir in tree_children_dirs(pe_id, N):
-    m_c = tl.recv(child_dir, m.shape); l_c = tl.recv(child_dir, l.shape)
-    O_c = tl.recv(child_dir, O.shape)
-    m, l, O = merge(m, l, O, m_c, l_c, O_c)
-if not is_root(pe_id):
-    tl.send(parent_dir(pe_id), m); tl.send(parent_dir(pe_id), l); tl.send(parent_dir(pe_id), O)
-else:
-    tl.store(O_base, O / l)
-```
-
-`tl.recv`는 데이터 도착까지 blocking(`tl_context.py:446-499`); 머지 순서가 정적
-tree로 고정되므로 data-dependent `pop`이 불필요 — **하드웨어/composite
-`pop`-as-dependency 변경이 필요 없는** 이유(§6).
-
----
-
-## 5. Case별 커널
-
-네 case 모두 §3(tile sweep) + §4(reduction)를 공유하며; `N`, query-block 폭,
-어느 `tile_*` predicate가 발동하는지에서만 다르다.
-
-### 5.1 공통 skeleton
-
-```python
-def attn_kernel(q_ptr, k_ptr, v_ptr, o_ptr, counter, start_pe, N,
-                q_block, softmax_scale, *, tl):
-    pe_id   = tl.program_id(axis=rank_axis)
-    my_len  = valid_len(counter, start_pe, pe_id, N)
-    n_tiles = ceil(my_len / TILE)
-    q_g     = load_Q_group(q_ptr)            # [G·T_q, d]; lazy tl.load, G folded into M (no broadcast)
-    m, l, O = init_running()                 # persistent arena: -inf, 0, zeros
-    for j in range(n_tiles):                 # K/V streamed per tile by the composite scheduler (§3)
-        if tile_all_future(j, q_block):      # causal skip (kernel if)
-            continue
-        run_tile(j, mask_or_null(j, q_block))   # §3: 2 composites + softmax MATH, in tl.scratch_scope
-    if N == 1:
-        tl.store(o_ptr, O / l)               # no reduction
-    else:
-        tree_reduce_and_store(pe_id, N, o_ptr)  # §4
-```
-
-### 5.2 DECODE, no SP (`N=1`, 한 PE가 head의 KV 소유)
-- `T_q = 1`, 모든 과거 KV에 attend ⇒ future tile 없음, 마지막 ragged tile만 mask.
-- **GQA 재사용이 전부**(decode는 KV-load-bound): KV head의 `G=8` query 행을 matmul
-  **M** 차원에 fold(`q_g`를 `[G, T_q, d] → [G·T_q, d]`로 reshape, byte-보존).
-  그러면 `Q·Kᵀ`는 `composite([G·T_q, d], Kᵀ[d, TILE]) → [G·T_q, TILE]`이고
-  `P·V`는 `composite([G·T_q, TILE], V[TILE, d]) → [G·T_q, d]`. KV tile(`[TILE, d]`)이
-  공유 `K`/`V` operand — **한 번 스트리밍되어 모든 `G·T_q` 행이 자동 재사용** —
-  그것들이 GEMM의 M 행이기 때문. K/V broadcast 불필요; composite의 tile plan 내
-  `m = G·T_q`가 timing이 모든 `G` 행의 작업을 올바르게 세게 한다(leading batch
-  축은 세어지지 *않음* — §8 참조).
-- `S_pe`가 scratch에 맞으면(작은/중간 context) 이것은 **one-shot** partial
-  attention으로 축약(Q·Kᵀ용 composite 하나, softmax 하나, P·V용 composite 하나) —
-  바로 baseline `_attention_mesh_mlo`의 `_partial_attention`, 단지 GQA-batched에
-  composite 경로. Tiling(§3)은 `S_pe`가 scratch scope의 tile 예산을 초과할 때만
-  발동.
-
-### 5.3 DECODE, with SP / KV-parallel (`N=8`)
-- 한 request의 KV가 8 PE에 round-robin; 각각 ≈`my_len` token 소유. 각 PE가 자기
-  로컬 tile에 대해 `G=8` query 행을 GQA-batch.
-- `T_q=1` ⇒ 짧은 per-PE sweep ⇒ **reduction 지배** ⇒ §4 tree(3단계)가 구조적으로
-  중요한 부분. Reduction latency는 batch될 때 다른 동시 decode token으로, 혹은 긴
-  single-stream context의 긴 per-PE tile sweep으로 숨겨진다.
-
-### 5.4 PREFILL, no SP
-- 전체 prompt 상주; query는 `T_q` token 블록(chunk).
-- causality가 실재, query block `[qs, qe)` vs KV tile `[ks, ke)`:
-  - `ke ≤ qs` → `tile_all_past` → `mask=None`, full compute.
-  - `ks ≥ qe` → `tile_all_future` → **skip**(커널 `if`).
-  - overlap → `tile_partial` → 커널이 삼각 additive mask 생성, tile op 시퀀스가
-    더함(`Sj + mask_j`).
-- GQA는 group의 query head를 decode와 같은 축으로 batch.
-
-### 5.5 PREFILL, with SP (Ring KV)
-- KV가 `N` PE(ring)에 sharded; 각 step이 peer의 KV 블록을 IPCQ로 전달. 실행 중
-  `(m,ℓ,O)`는 **ring step을 가로질러** carry(Python 핸들, `_attention_mesh_kv`가
-  오늘 하는 그대로).
-- IPCQ가 다음 step의 KV 수신을 현재 step의 compute와 `tl.recv_async`/`tl.wait`로
-  오버랩(`tl_context.py:543-560` — 이미 존재).
-- **Causal ring skip:** 들어오는 step의 KV 블록이 전부 로컬 query block *이후*면 그
-  compute를 skip(recv-consume 안 함 / fold 안 함) — causal 어텐션에서 약 절반 step
-  제거.
-- 수신 버퍼는 ping-pong(persistent arena, 재활용 안 함).
-- ring 후, `(m,ℓ,O)`는 query-소유 PE별 최종; KV-parallel 분할이 남으면 §4가
-  tree-merge한 뒤 정규화 + store.
-
-baseline `_attention_mesh_kv`가 이미 ring fold를 구현한다; 본 ADR은 GQA 재사용,
-causal step-skip, `recv_async` 오버랩을 추가한다.
-
----
-
-## 6. 왜 하드웨어 / composite 변경이 불필요한가
-
-- HW/composite `pop`-as-dependency가 도움 될 유일한 곳은 poll을 숨길 다른 작업이
-  없는 외로운 reduction — 즉 batch=1 **그리고** 짧은 context. 타깃은 **agentic =
-  낮은 batch, 긴 context** ⇒ 각 PE가 많은 KV tile을 가짐; §4 tree의 `tl.recv`
-  blocking은 동시 token / 긴 context의 sweep 작업으로 가려진다.
-- §4 reduction은 **정적** tree를 쓰므로 recv 순서가 컴파일 타임에 고정 —
-  data-dependent pop 없음. `tl.recv`(blocking)로 충분.
-- **결정: greenlet `tl.send`/`tl.recv` collective로 출시.** 짧은-context,
-  single-stream, latency-critical 타깃이 나타날 때만 HW `pop` 재검토.
-
----
-
-## 7. 제어 vs 실행 분할 (load-bearing 원칙)
-
-| Concern | Owner |
-|---|---|
-| tile skip (future), mask 생성, causal bound | **커널** (greenlet 본문의 Python `if` + 산술) |
-| 주소 / offset / valid-length 산술 | **커널** (counter에서) |
-| reduction 스케줄링, IPCQ send/recv 순서 | **커널** (정적 tree) |
-| Q·Kᵀ, P·V (per-tile K/V DMA 스트리밍 + tiling 포함) | **`tl.composite`** → PE_SCHEDULER |
-| Q load, mask add, softmax math, 실행 중 `(m,ℓ,O)` 머지 | **`tl` op** PE 엔진 상 (커널 발행) |
-
-커널이 결정하고; GEMM은 composite로 scheduler에 offload되며; 나머지 `tl` op은
-이미 결정된 작업을 실행한다. 이는 kernbench가 이미 지원하는 greenlet + composite
-모델 그대로 — 새 제어 추상화 없음.
-
----
-
-## 8. 필요한 kernbench 변경
-
-**설계 iteration의 정정:** 진짜 GQA(`h_q > h_kv`)는 **신규 primitive 불필요** —
-§5.2의 커널 재구조화(KV head별, `G`를 M에 fold, byte-보존 reshape)만 필요. 보조
-ADR들은 *efficiency / scale* enabler이지 GQA blocker가 아니다.
-
-**커널 내 알고리즘 작업 (신규 primitive 없음; 기존 `tl` API):**
-
-- **GQA Q축 batching**(재사용 레버) — KV head별로 `G·T_q`를 matmul M 차원에 fold
-  (§5.2); `_view` 식 byte-보존 reshape; GEMM은 M = `G·T_q`인
-  `tl.composite(op="gemm")`. 오늘 timing/data 모드 모두 동작.
-- **composite를 통한 GEMM**(§1/§3) — Q·Kᵀ와 P·V를 각각 non-blocking
-  `tl.composite(op="gemm")`로 발행; PE_SCHEDULER가 tiling하고 `tl.ref` K/V
-  operand의 DMA를 스트리밍(기존 `CompositeCmd`; 새 커맨드 종류 없음).
-- root로의 tree reduction(§4)이 baseline all-to-all fan-out 대체 — `tl.send`/
-  `tl.recv` 상의 순수 커널 제어 흐름.
-- causal tile skip + additive boundary mask(§3/§5.4) — 커널 `if` + `+`로 더한
-  mask 텐서.
-- round-robin KV placement / valid-length 산술(§2) — launch-arg 산술 +
-  `DPPolicy(pe="row_wise")`.
-
-**efficiency / scale을 위한 신규 primitive (각각 보조 ADR 있음):**
-
-1. **per-tile scratch 재활용** — **ADR-0063**(`tl.scratch_scope`). *스케일에 필요*:
-   `S=16` 상한(1 MiB bump allocator)을 제거해 현실적 context 길이가 돌게 함. 셋 중
-   최고 가치.
-2. **Lazy `tl.load`** — **ADR-0062**(non-blocking load + 첫 사용 시점 auto-wait; API
-   표면 불변). *효율*: 명시적 load(Q group, 비-composite 커널)를 뒤따르는 compute와
-   오버랩. 타일당 **K/V** prefetch는 composite scheduler가 처리(§1)하므로, 이것은
-   나머지 명시적 load를 담당. 전역 시맨틱 변경 → 기존 골든 재생성(ADR-0062 D3).
-3. **GQA head / mask broadcast** — **ADR-0061**(`tl.broadcast`). *선택적 편의*, GQA
-   blocker 아님(위 정정 참조). `G·T_q` 행에 걸친 additive-mask 구성과 범용 커널에
-   유용; `np.matmul`이 data 모드에서 이미 broadcast하므로 정확성엔 불필요. 최저
-   우선순위.
-
-**명시적 REJECTED (효율적 대안 선택):**
-
-4. ~~내부 루프 전체를 내재화하는 전용 "flash-composite" 커맨드 종류 —
-   DMA→MM→VEC→DMA→MM→VEC + carried `(m,ℓ,O)` register 상태 + 꼬리 IPCQ push.~~ 두
-   GEMM은 기존 `CompositeCmd`를 **쓴다**(§1/§3) — 그것이 scheduler 관리 tiling, K/V
-   DMA 스트리밍, cross-tile 파이프라이닝을 준다. 기각되는 것은 softmax 머지 +
-   cross-tile register 수명 + IPCQ-push epilogue까지 흡수하는 **새** 커맨드 종류다:
-   크고 특수목적이며, 하이브리드 대비 유일한 delta(완전 softmax offload)가 현재
-   모델링 충실도(per-op CPU 발행 비용 = 0; §1 참조)에서 정당화되지 않는다. cost
-   model(§9)이 완전 offload를 측정 가능하게 가치 있게 만들면 재검토.
-5. ~~하드웨어 `pop`-as-dependency.~~ 범위 밖(§6).
-6. ~~본 커널 내 RoPE / QKV projection / KV-cache write.~~ Upstream `qkv_rope`
-   (P1–P5). RoPE를 fold하면 매 decode step마다 과거 tile을 재회전해야 하고 Ring
-   Attention의 post-RoPE pass-through를 깬다.
-
----
-
-## 9. Open 튜닝 항목 (kernbench에서 측정, blocking 아님)
-
-1. **Composite tile-pipeline depth** — KV-load-bound에 지배적; 커널이 wait 전에
-   non-blocking composite를 얼마나 앞서 발행하는지, 그리고 scheduler의 타일당
-   스트리밍 depth.
-2. **reduction tree의 PE↔mesh-이웃 매핑** — 각 depth-`⌈log₂ N⌉` pair가 물리
-   `N/S/E/W` 이웃인지 보장; 나쁜 매핑은 hop 추가. SFR install 대조 검증.
-3. **TILE 크기** — scratch 상주(`S/P tile + O_acc + G-way GQA`)와 DMA 효율 균형;
-   ADR-0063 및 scheduler의 `TILE_M/K/N`(`pe_scheduler.py`)과 상호작용.
-4. §5.5의 **Ring 버퍼 ping-pong vs `recv_async` depth**.
-5. **decode의 one-shot vs tiled 교차점**(§5.2) — tiling이 단일 composite를 이기는
-   `S_pe` 임계값.
-6. **per-op CPU 발행 비용 (cost model)** — 현재 `dispatch_cycles=0`(`pe_cpu.py`)이라
-   composite-vs-primitive 발행 오버헤드가 보이지 않음. op 종류별 차등 발행 비용
-   (`tl.composite` descriptor push ≫ primitive op)이 하이브리드의 CPU-saturation
-   이점을 **측정 가능**하게 한다(§1). **ADR-0064**에 명세; 별도 future work로 추적.
-
----
-
-## 10. Coverage 요약
-
-| Case | KV placement | Inner loop | Reduction | Masking |
-|---|---|---|---|---|
-| Decode, no SP | 1 PE, all KV | one-shot or tiled, GQA-batched | none | last tile only |
-| Decode, SP | round-robin `N` PEs | tiled, GQA-batched | §4 tree (`⌈log₂ N⌉`) | last tile only |
-| Prefill, no SP | resident | tiled per q-block | none | skip future / triangular boundary |
-| Prefill, SP (Ring) | ring-sharded | per ring step (recv_async overlap) | running state + §4 tree | causal step-skip + boundary |
-
-**Case별 I/O** (전체 계약 §0.5):
-
-| Case | Inputs | Output | IPCQ partials |
-|---|---|---|---|
-| Decode, no SP | `Q[G,1,d]`, full `K/V[S,d]` on 1 PE | `O[G,1,d]` at `O_base` | none |
-| Decode, SP | `Q[G,1,d]`, per-PE `K/V[S_pe,d]` | `O[G,1,d]` (root) | `(m,ℓ,O_i)` per PE → tree |
-| Prefill, no SP | `Q[G,T_q,d]`, `K/V[≤end,d]` | `O[G,T_q,d]` at `O_base` | none |
-| Prefill, SP (Ring) | `Q[G,T_q,d]`, ring-delivered `K/V` blocks | `O[G,T_q,d]` (root) | running `(m,ℓ,O)` + tree |
-
-모든 case에서: KV-cache write와 RoPE는 **upstream**에서; output projection은
-**downstream**에서; score `S`와 prob `P`는 미생성.
-
----
-
-## 11. 검증 계획 (Phase 1 테스트 개요)
-
-SPEC/ADR coverage: R5 (PE↔PE IPCQ, PE↔HBM), R2 (traversal에 의한 latency),
-ADR-0023/0025 (IPCQ), ADR-0046 (`tl` 계약), ADR-0054 (eval bench).
-
-시뮬레이터의 계약은 bit-exact 수치가 아니라 **traversal에 의한 latency + 결정성 +
-구조적 정확성**(SPEC §0, §0.1)이다 — Phase 2 데이터는 주로 data path를
-exercise하려 존재하고, `tl.trans`는 reshape-not-transpose이며, `bf16`은 `f16`으로
-모델된다. 따라서 검증은 **구조/타이밍 우선**, 수치 parity는 제한된 2차 체크.
-
-1. **data 모드 실행(`enable_data=True`):** GQA 커널(`h_q = G·h_kv`)이 baseline
-   head-packing이 부딪히는 byte-conservation 에러 없이 — 네 case 모두 — 완료. (이는
-   §5.2 재구조화가 필요하지, 신규 primitive가 *아님*.)
-   **수치 parity(2차):** reshape-as-transpose가 정확한 symmetric/identity 입력에
-   대해, 커널 `O`가 numpy FlashAttention reference와 fp tolerance 내 일치. 완전
-   asymmetric parity는 실제 `tl.transpose`에 좌우됨(범위 밖; 플래그됨).
-2. **GQA 재사용:** `h_q = G·h_kv`에서 K/V `dma_read_count`가 `G`와 무관(타일당 한
-   load, group에 걸쳐 재사용), 반면 GEMM 작업은 `G`로 스케일. 레버가 실제 발동함을
-   assert.
-3. **tree reduction 단계 수:** decode-SP가 `⌈log₂ N⌉` reduction 라운드(`N−1`이
-   아님)를 발행; op_log `ipcq_send`/`recv` 수가 tree와 일치.
-4. **Causal skip:** prefill이 모든 `tile_all_future` tile을 skip — GEMM 수가 full
-   grid가 아니라 하삼각 tile 수와 일치.
-5. **Long context (ADR-0063):** scope 없이 1 MiB를 넘기는 `S` sweep이 완료하고
-   reference와 일치.
-6. **Load/compute 오버랩:** tiled sweep의 end-to-end latency가 직렬
-   `Σ(load+compute)`보다 낮음 — composite scheduler의 타일당 K/V 스트리밍(§1/§3)과
-   Q load를 오버랩하는 lazy `tl.load`(ADR-0062)에서. (오버랩은 실제 모델 동시성,
-   빼기가 아님.)
-7. **Composite GEMM offload (구조적):** 각 tile의 Q·Kᵀ와 P·V가 blocking `tl.dot`이
-   아니라 `CompositeCmd`(non-blocking)를 PE_SCHEDULER에 emit; op_log가 composite
-   tile plan을 보이고 커널이 다음 tile의 composite를 wait 전에 발행(cross-tile
-   파이프라이닝).
-8. **결정성:** 동일 입력 → 동일 op_log + latency (SPEC §0.1).
-
----
-
-## B. 하이브리드 전환에서 나온 Open 설계 항목 (추후 검토)
-
-이들은 결정이 순수 greenlet primitive 경로에서 **composite hybrid + lazy
-`tl.load`**(이번 개정)로 옮겨갈 때 생겼다. 설계를 막는 것은 없으며; 각각 구현 중
-검증 패스가 필요하다. (묻지 않고) 작업 합의에 따라 여기 기록한다 — 권고는 내가
-예측한 기본값이며; 검토 시 수정.
-
-1. **DDD-0060이 아직 미동기화.** Detailed Design Document는 여전히 옛
-   `tl.load_async` 더블버퍼 경로와 primitive `tl.dot` 내부 루프를 기술한다(그
-   §4.3/§5/§10). 하이브리드(composite GEMM, lazy load, K 사전 transpose)로
-   갱신해야 한다. DDD가 파생 how-to이고 큰 재작성이라 *검토용으로 남김*; ADR이 이제
-   권위 있는 기록이다. **권고:** 구현 시작 전 후속으로 DDD 동기화.
-
-2. **composite GEMM의 K operand 방향.** Q·Kᵀ는 `b = [d, TILE]`가 필요하나 KV
-   cache는 K를 `[S_pe, d]`로 저장한다. `tl.trans`는 메타데이터 전용이고
-   `MemoryStore.read`는 transpose가 아니라 reshape(`memory_store.py:73`) — 비자명
-   데이터엔 런타임 transpose가 틀리다. **권고:** cache에 K를 **사전 transpose**
-   `[d, S_pe]`로 저장(pseudocode와 §3이 이를 가정)하여 `tl.ref(k_tile, (d, TILE))`이
-   연속 slice가 되게. upstream `qkv_rope` write 레이아웃이 이를 지원하는지 검증,
-   아니면 실제 `tl.transpose` 추가(더 무거움; 연기).
-
-3. **composite 출력 버퍼 vs `tl.scratch_scope`.** 각 Q·Kᵀ composite는 `Sj`를
-   `out_addr`에 쓰고; 커널이 softmax MATH용으로 그것을 읽는다. 그 출력 버퍼와
-   in-flight composite의 타깃은 per-tile `scratch_scope`(ADR-0063)가 소비 전에
-   재활용하지 **않는** 곳에 살아야 한다 — in-flight lazy load와 같은 규율(ADR-0062
-   D-Negative). **권고:** *현재* tile의 composite 출력은 scoped arena에(같은
-   iteration에 소비); persistent `(m,ℓ,O)`는 바깥에. 다음 tile의 composite를 일찍
-   발행할 때(cross-tile 파이프라이닝) use-after-recycle이 없음을 검증.
-
-4. **composite 스트리밍 하의 GQA `dma_read_count` 레버.** 레버(§11.2: K/V
-   `dma_read_count`가 `G`와 무관)는 composite가 모든 `G·T_q` M-행에 재사용되는
-   **하나의** K/V tile DMA를 emit한다고 가정한다. scheduler의 `generate_gemm_plan`은
-   `TILE_M/K/N`(`pe_scheduler.py:35-37`, 32/64/32)으로 tiling한다 — `G·T_q`에 대한
-   M-tiling이 M-tile마다 공유 K/V tile DMA를 재발행하지 **않음**을 확인(즉 operand
-   DMA가 M-tile에 걸쳐 공유되거나, 아니면 레버가 약해짐). **권고:** levers 테스트에서
-   assert; 위반 시 GQA 이득은 DMA가 아니라 compute에만 — 여전히 정확하나 헤드라인이
-   바뀜.
-
-5. **커널 TILE vs scheduler `TILE_M/K/N`.** 커널은 논리적 KV `TILE`을 다루고;
-   scheduler는 고정 `TILE_M/K/N`으로 내부 재tiling한다. 두 tiling 레이어가
-   상호작용(scratch 상주, 파이프라인 depth). **권고:** 커널 TILE을 K/V 스트리밍
-   granularity로 보고 scheduler가 GEMM을 sub-tile하게; DDD에 관계를 문서화하고 둘 다
-   sweep(§9 항목 1, 3).
-
-6. **cost model은 별도 ADR.** 하이브리드의 CPU-saturation 이점은
-   `dispatch_cycles=0`인 동안 보이지 않는다. op 종류별 issue-cost 모델은
-   **ADR-0064**에 명세; 본 ADR의 §1/§9는 *측정 가능한*(구조적뿐 아닌) 이점을 위해
-   그것에 의존. **권고:** eval에서 하이브리드 latency 이점을 주장하기 전에 ADR-0064
-   모델을 도입.
-
-7. **Ring 경로(§5.5) GEMM.** §5.5는 여전히 primitive op + `recv_async`로 ring
-   fold를 기술한다. 일관성을 위해 ring의 per-step Q·Kᵀ / P·V도 composite여야 하고;
-   IPCQ `recv_async` 오버랩은 직교하며 유지. **권고:** 구현 중 ring step에 같은
-   하이브리드 형태 적용; 낮은 리스크, §3 미러.

docs/adr-proposed/ADR-0060-algo-gqa-fused-attention-ahbm.md

@@ -1,817 +0,0 @@
-# ADR-0060: AHBM GQA Fused Attention Kernel (Llama3-70B)
-
-## Status
-
-Proposed
-
-**Context model:** Llama3-70B.
-**Decision drivers:** agentic workload → low batch, long context;
-KV-load-bound decode; sequence-parallel (Ring KV) for long-context prefill.
-
-**Supersedes / extends:** the existing mesh-native attention kernels
-`_attention_mesh_kv` (prefill) and `_attention_mesh_mlo` (decode) and the
-`milestone-gqa-llama70b` eval bench. See *§A. Relationship to existing
-kernbench work* — that code is the baseline this ADR upgrades to a real
-GQA, causal, long-context kernel.
-
-**Supporting ADRs** (efficiency / scale enablers — *not* GQA blockers;
-see §8 correction): **ADR-0063** `tl.scratch_scope` (per-tile scratch
-recycling — required for realistic context length), **ADR-0062** lazy
-`tl.load` (non-blocking load with auto-wait on first use → load/compute
-overlap), **ADR-0061** `tl.broadcast` (optional mask/general
-convenience). The two GEMMs (Q·Kᵀ, P·V) are issued as scheduler-managed
-`tl.composite` commands (the existing `CompositeCmd`; no new command
-kind); real GQA itself needs only kernel restructuring (§5.2).
-
-**Algorithm lineage.** This kernel is **FlashAttention** (tiling +
-online/streaming softmax with fused P·V — no full score matrix
-materialised). The KV-parallel split-and-combine in §4 is
-**FlashDecoding** (split-KV with log-sum-exp merge). The Ring path in
-§5.5 is **Ring Attention** (KV blocks rotated around the mesh, folded by
-the same online softmax). No new math is introduced; this ADR maps those
-known algorithms onto the kernbench **greenlet `tl` programming model**
-(ADR-0020, ADR-0046) and the **IPCQ** PE↔PE collective (ADR-0023/0025).
-
----
-
-## TL;DR — final GQA kernel (composite hybrid) pseudocode
-
-The decision (§1) in one place: **GEMMs → `tl.composite`; softmax merge +
-tree reduction → kernel; `tl.load` is lazy.** Per-KV-head, `G` folded into
-the matmul M dim (real GQA, no broadcast). This is the reference shape; the
-prose sections elaborate each piece.
-
-```python
-def gqa_attention(q_ptr, k_ptr, v_ptr, o_ptr,
-                  counter, start_pe, N, q_block, scale, *, tl):
-    # ---- geometry (kernel arithmetic, §2) ----
-    pe_id   = tl.program_id(axis=rank_axis)
-    G       = H_q // H_kv                       # query heads per KV head (=8)
-    causal  = q_block.is_prefill
-
-    for kv in range(H_kv):                       # one KV head per iteration (§5.2)
-        # Q group: G rows folded into M → [G·T_q, d]. Lazy load (ADR-0062):
-        # issue now, auto-wait at first use inside the first composite.
-        q_g     = tl.load(q_base(kv), (G * q_block.T_q, d))      # [G·T_q, d]
-        my_len  = valid_len(counter, start_pe, pe_id, N) if not causal else S_kv
-        n_tiles = ceil(my_len / TILE)
-
-        # persistent arena (outside scratch_scope, ADR-0063): -inf, 0, zeros
-        m, l, O = init_running(G * q_block.T_q, d)
-
-        for j in range(n_tiles):
-            if causal and tile_all_future(j, q_block):
-                continue                                         # causal skip (kernel if)
-            with tl.scratch_scope():                             # per-tile temporaries (ADR-0063)
-                # --- GEMM #1 on the scheduler: Q·Kⱼᵀ; K streamed by composite ---
-                Sj = tl.composite("gemm", a=q_g,
-                                  b=tl.ref(k_tile(kv, j), (d, TILE))) * scale   # Kᵀ pre-stored [d,TILE]
-                if causal and tile_partial(j, q_block):
-                    Sj = Sj + causal_mask(j, q_block)            # additive boundary mask
-                # --- online softmax merge in the kernel (MATH ops) ---
-                m_new = tl.maximum(m, tl.max(Sj, axis=-1))
-                P     = tl.exp(Sj - m_new)
-                corr  = tl.exp(m - m_new)
-                l     = l * corr + tl.sum(P, axis=-1)
-                # --- GEMM #2 on the scheduler: P·Vⱼ; V streamed by composite ---
-                Oj    = tl.composite("gemm", a=P,
-                                     b=tl.ref(v_tile(kv, j), (TILE, d)))
-                O     = O * corr + Oj                            # running merge
-                m     = m_new
-
-        # ---- cross-PE combine (§4): log-sum-exp tree to root, or store ----
-        if N == 1:
-            tl.store(o_base(kv), O / l)
-        else:
-            tree_reduce_and_store(m, l, O, pe_id, N, o_base(kv))  # tl.send/tl.recv
-```
-
-> **Why this shape:** the two `tl.composite` calls offload tiling + K/V DMA
-> streaming + cross-tile pipelining to PE_SCHEDULER (CPU issues coarse
-> descriptors and runs ahead → engines stay saturated, §1); the softmax
-> merge and the IPCQ tree reduction stay in the kernel because the existing
-> `CompositeCmd` cannot carry cross-tile state. `K` is pre-stored
-> transposed `[d, TILE]` to sidestep the reshape-not-transpose caveat
-> (§3, §B). A bespoke "flash-composite" command kind is **not** introduced
-> (§8 item 4).
-
----
-
-## A. Relationship to existing kernbench work (read first)
-
-kernbench **already runs FlashAttention with an online-softmax `(m, ℓ, O)`
-merge over IPCQ today.** Two kernels exist:
-
-| File | Role | Mechanism |
-|---|---|---|
-| `src/kernbench/benches/_attention_mesh_kv.py` | prefill (Ring K/V) | per-rank partial attention, bidirectional K/V fan-out, online-softmax fold |
-| `src/kernbench/benches/_attention_mesh_mlo.py` | decode (split-KV) | per-rank one-shot partial attention, bidirectional `(m,ℓ,O)` fan-out, log-sum-exp merge |
-
-Both are driven by `milestone-gqa-llama70b` (4 panels:
-`{single,multi}_user × {prefill,decode}`,
-`src/kernbench/benches/milestone_gqa_llama70b.py`) and tested in
-`tests/attention/test_milestone_gqa_llama70b.py`.
-
-**They are written in the greenlet `tl` API:** `tl.load`, `tl.dot`,
-`tl.softmax`/`tl.max`/`tl.sum`/`tl.exp`, `tl.send`/`tl.recv`, and Python
-`-`/`*`/`/` on `TensorHandle` (each emits a `MathCmd`) — the GEMMs as
-**blocking `tl.dot`**, not composites. The running `(m, ℓ, O)` is just
-Python `TensorHandle`s threaded through the loop. This ADR keeps the
-running state and the softmax merge in the kernel but **moves the two
-GEMMs onto the scheduler-managed `tl.composite` path** (see §1) — this
-**matters for the design**.
-
-**Three deliberate limitations of the baseline** — exactly what an
-*efficient GQA* kernel must lift:
-
-1. **No GQA reuse.** `h_q == h_kv == 1`
-   (`test_milestone_gqa_llama70b.py:137-142`). The test attributes this to
-   a MemoryStore byte-conservation failure on a *broadcast view*, but that
-   failure is a property of the baseline's **head-packing hack**
-   (`_view(K, (h_q·d, S_kv))`, which conflates all heads into one matmul
-   dim and only conserves bytes when `h_q == h_kv`). The correct fix is
-   **kernel restructuring**, not a broadcast op: process **one KV head at
-   a time** and fold the `G` group rows into the matmul **M** dimension
-   (§5.2). That uses only byte-conserving reshapes, so real GQA
-   (`h_q = G·h_kv`) runs with **no new primitive** — see §8.
-2. **O(N) reduction.** The baseline does an all-to-all bidirectional
-   fan-out so *every* rank ends with the full answer (`n_ranks − 1`
-   steps). Attention only needs `O` at the query owner → a **tree
-   reduction** to root is `⌈log₂ N⌉` steps (§4).
-3. **Validation scale only.** `S = 16` because the 1 MiB scratch bump
-   allocator leaks per-tile temporaries
-   (`test_milestone_gqa_llama70b.py:123-148`) and there is no causal
-   masking / tiling → fixed by **ADR-0063** (recycling) + §5 (tiling,
-   causal skip) + composite K/V streaming (§3) + **ADR-0062** (lazy load
-   overlap).
-
-**Documentation debt (out of scope but recorded):** the baseline cites
-ADR-0055/0056/0057/0058/0059, **none of which exist as files** — they are
-ghost references. This ADR does not retro-write them; see the Detailed
-Design Document's *Open Decisions* for the recommendation.
-
----
-
-## 0. Reference dimensions (Llama3-70B)
-
-| Symbol | Meaning | Value |
-|---|---|---|
-| `H_q` | query heads | 64 |
-| `H_kv` | KV heads | 8 |
-| `G` | GQA group size = `H_q / H_kv` | 8 |
-| `d` | head dim | 128 |
-| `L` | layers | 80 |
-| `D` | model dim | 8192 |
-
-Hardware recap:
-- **AHBM (chip)** = set of **CUBE**s (memory cubes, each with a logic die
-  containing **PE**s) + **IO die** (ADR-0003).
-- **IPCQ**: PE↔PE queues, 4 mesh-direction queue-pairs per PE
-  (`N/S/E/W`, ADR-0023 D3; `global_*` for inter-SIP, ADR-0032). Kernel
-  API: `tl.send(dir, src)` / `tl.recv(dir, shape, dtype)`
-  (`tl_context.py:402-499`).
-- **Composite command** (`CompositeCmd`, `pe_commands.py:144-162`): a
-  single GEMM (or MATH) *head* plus element-wise *epilogue* stages
-  (`bias/relu/scale/add/...`), issued **non-blocking** to PE_SCHEDULER,
-  which generates a tile plan and streams DMA→GEMM→write per tile
-  (ADR-0014 D6; `pe_scheduler.py:104-143`). It is **not** a general
-  multi-op DAG: it cannot chain two GEMMs, cannot carry register state
-  across instances, and cannot pop/wait on IPCQ. This ADR therefore issues
-  **each** of the two attention GEMMs (Q·Kᵀ, P·V) as its own composite and
-  keeps the cross-GEMM softmax merge + the IPCQ reduction in the kernel —
-  it does **not** need a new "flash-composite" command kind (see §1, §8).
-- **Allocation policy (SP):** one query head per CUBE in the multi-user
-  panels; the `G` query heads of a KV group map within a CUBE's PEs.
-
-**Two orthogonal mapping layers** (round-robin placement):
-- **Layer 1 (KV-parallel, intra-request):** KV token `i` lands on PE
-  `(start_pe + i) mod N` — round-robin so one long request is split
-  across `N` PEs, balanced to ≤1 token.
-- **Layer 2 (inter-request):** `start_pe = request_id mod N` rotates the
-  "PE-0 role" per request.
-- Combined: `pe = (request_id + global_token_idx) mod N`.
-
-`N` = number of PEs in the KV-parallel reduction group for one
-(query head, request). Reduction stays **intra-CUBE** (a query head never
-spans CUBEs — explicit non-goal).
-
----
-
-## 0.5 Kernel boundary, preconditions, I/O contract
-
-### 0.5.1 Position in the decoder layer
-
-```
-1. RMSNorm
-2. QKV projection (GEMM)        ─┐  qkv_rope kernel (SEPARATE, upstream)
-3. RoPE on Q and K              ─┤
-4. write new K,V → KV cache     ─┘
-5. ===== THIS KERNEL: FlashAttention =====   (post-RoPE Q, K-cache, V-cache → O)
-6. Output projection (GEMM)        out_proj kernel (SEPARATE, downstream)
-7. residual add → FFN ...
-```
-
-**Preconditions (upstream `qkv_rope`, NOT this kernel):**
-- **P1.** Q is already RoPE-rotated. No rotation here.
-- **P2.** K-cache stores **post-RoPE** K (never re-rotated at attention
-  time — the reason for post-RoPE caching).
-- **P3.** For decode, the new step's K row is RoPE-rotated and appended
-  to its owning PE's K-cache slot **by `qkv_rope` before this kernel
-  launches.** ⇒ this kernel is **pure read** on the KV cache.
-- **P4.** V is not rotated; V-cache holds raw projected V.
-- **P5.** RoPE position upstream is the token's **absolute global
-  position** `global_idx = local_slot·N + ((pe_id − start_pe) mod N)`
-  (§2.1). Round-robin placement ≠ RoPE position.
-
-### 0.5.2 Shape symbols
-`T_q` = query length this launch (decode: 1; prefill/chunk: chunk width).
-`S` = total context length. `S_pe` = keys owned by this PE ≈ `⌈S/N⌉`.
-Storage `bf16` (numpy proxy `f16`, `memory_store.py:16`); `m,ℓ,O`
-accumulators `f32` where precision matters.
-
-### 0.5.3 INPUTS (per kernel launch)
-
-| Input | Shape (per KV head) | Location | Notes |
-|---|---|---|---|
-| `Q` | `[G, T_q, d]` | per-PE HBM (loaded to TCM) | post-RoPE (P1). `G` query rows of the group batched. |
-| `K_cache` | `[S_pe, d]` | per-PE HBM, base `K_base[pe]`, contiguous | post-RoPE (P2). Read-only. Dense locally, strided in global index (§2.1). |
-| `V_cache` | `[S_pe, d]` | per-PE HBM, base `V_base[pe]` | raw V (P4). Read-only. |
-| `global_token_counter` | scalar | launch arg | kernel derives `S_pe`, slot↔global, causal bounds. |
-| `start_pe` (= `request_id mod N`) | scalar | launch arg | Layer-2 rotation. |
-| `pe_id`, `N` | scalars | launch / `tl.program_id` | reduction-group geometry. |
-| `q_block_meta` `{q_start, T_q}` | launch arg | prefill/SP causal masking & skip. |
-| `O_base` | address | launch arg | where final O is written (root PE only). |
-| `softmax_scale` | scalar | launch arg | `1/√d`. |
-
-For **Ring Attention (§5.5)** add per ring step: incoming `K_block,
-V_block` via IPCQ into ping-pong buffers (post-RoPE), plus
-`step_kv_global_range` for the causal step-skip.
-
-### 0.5.4 OUTPUTS
-
-| Output | Shape | Location | Notes |
-|---|---|---|---|
-| `O` (final) | `[G, T_q, d]` | `O_base`, **root PE only** | normalised `O_acc / ℓ_acc`, cast to storage dtype. |
-
-**Intermediate (non-root PE, on IPCQ — not a kernel-visible output):**
-`(m_i, ℓ_i, O_i)` (`m,ℓ`: `[G, T_q]`; `O_i`: `[G, T_q, d]` unnormalised)
-pushed to `tree_parent` via `tl.send` (§4). `O_i` is the heavy part.
-
-**No-SP (`N=1`):** no IPCQ partials; the single PE's running `(m,ℓ,O)` is
-normalised in place and written to `O_base`.
-
-**Explicitly NOT outputs:** KV-cache writes (upstream, P3), output
-projection (downstream), score `S` and probs `P` (never materialised).
-
----
-
-## 1. Decision (mechanism)
-
-**Implement the kernel as a *hybrid*: issue the two GEMMs (Q·Kᵀ and P·V)
-as scheduler-managed `tl.composite(op="gemm")` commands, and keep the
-online-softmax merge and the cross-PE reduction as kernel-level `tl`
-ops.** `tl.load` is **lazy** (non-blocking; the wait is auto-inserted at
-first use of the loaded data — ADR-0062), so explicit HBM loads overlap
-the compute that follows. Rationale grounded in kernbench's execution +
-latency model:
-
-- **GEMM tiling is offloaded to PE_SCHEDULER.** A `CompositeCmd` is
-  non-blocking (`kernel_runner.py:182-191`, `pe_scheduler.py:104-121`):
-  the kernel pushes **one coarse descriptor** (M = `G·T_q`, the whole
-  per-PE tile sweep) and the scheduler generates the tile plan and streams
-  DMA→GEMM→write per tile (ADR-0014 D6). K/V are `tl.ref` operands the
-  scheduler streams from HBM, so per-tile **K/V prefetch is the
-  scheduler's job** — no explicit prefetch op. The CPU (greenlet) is freed
-  to issue the **next** composite while the current one runs, so the
-  scheduler keeps the GEMM engine saturated across tiles.
-- **This reflects the hardware** and decouples CPU issue-rate from
-  execution-rate. The blocking per-op `tl.dot` path, by contrast, stalls
-  the CPU on every GEMM and leaves GEMM-engine **bubbles** during the
-  interleaved softmax MATH ops; it is realistic only if the CPU can keep
-  up with fine-grained per-tile issue.
-- **The running `(m, ℓ, O)` flash state stays Python `TensorHandle`s**
-  threaded through the loop (the baseline already does this); the softmax
-  merge (max/exp/sum/rescale) is kernel-level `tl` MATH **between** the two
-  GEMM composites. The existing `CompositeCmd` cannot chain two GEMMs or
-  carry cross-tile register state (§0), so the merge necessarily lives in
-  the kernel — this is the hybrid split, not a limitation worked around.
-- **Cross-PE combination** is a log-sum-exp **tree** over `(m, ℓ, O)`
-  after P·V, via kernel-level `tl.send`/`tl.recv` (§4) — unchanged.
-
-So the per-tile inner pipeline is:
-
-```
-q_g = tl.load(Q group)                         # lazy; auto-wait at first use
-per tile j:
-  Sⱼ = tl.composite("gemm", a=q_g, b=tl.ref(Kⱼ))  → Sⱼ   # scheduler streams Kⱼ DMA + GEMM
-  Sⱼ += maskⱼ                                          # kernel MATH, boundary tile only
-  online-softmax: mⱼ, m_new, P, corr, ℓ               # kernel MATH
-  Oⱼ = tl.composite("gemm", a=P, b=tl.ref(Vⱼ))   → Oⱼ   # scheduler streams Vⱼ DMA + GEMM
-  O = O*corr + Oⱼ;  m = m_new                         # kernel MATH (running merge)
-```
-
-with each tile's MATH temporaries wrapped in `tl.scratch_scope`
-(ADR-0063) so scratch stays O(1), and the next tile's composites issued
-before the current tile's results are waited on (non-blocking handles) so
-the scheduler pipelines across tiles.
-
-Control flow (tile skip, mask generation, reduction scheduling, address
-arithmetic) lives in the **kernel** (plain Python `if`/arithmetic in the
-greenlet body). This is exactly what kernbench's greenlet model already
-permits (`kernel_runner.py`, ADR-0020 D3).
-
-> **What this supersedes.** An earlier iteration proposed a pure greenlet
-> primitive path (all `tl.dot`, no composite) on the argument that
-> "composite yields no latency benefit." That holds **only because** the
-> simulator currently charges **zero** per-op CPU issue cost
-> (`dispatch_cycles=0`, `pe_cpu.py`) — it models away exactly the CPU
-> issue-rate / DMA-program cost that descriptor offload exists to hide.
-> The hybrid is the faithful representation of an efficient kernel. The
-> **measurable** size of the win (can the CPU saturate the engines for
-> many tiles?) is gated on modelling an op-type-differentiated issue cost,
-> tracked as future work (cost model; §9). Even at `dispatch_cycles=0` the
-> non-blocking composite path fills the GEMM-engine bubbles the blocking
-> `tl.dot` path leaves.
->
-> **Why this is the efficient choice.** GEMM tiling + DMA streaming +
-> cross-tile pipelining are offloaded to the proven `CompositeCmd`
-> scheduler path; the softmax merge and the proven IPCQ collective stay in
-> the kernel. The only genuinely new machinery is two small, general
-> primitives (ADR-0062 lazy `tl.load`, ADR-0063 `tl.scratch_scope`); the
-> reduction reuses `tl.send`/`tl.recv`. A bespoke "flash-composite"
-> command (one kind internalising the softmax merge + carried register
-> state + an IPCQ-push epilogue) is **not** built — large, special-purpose,
-> and its only delta over this hybrid (full softmax offload) is not
-> justified at the current modelling fidelity; see §8.
-
----
-
-## 2. Memory layout & driver responsibilities
-
-### 2.1 KV cache allocation
-- Per-PE KV buffers sized to per-PE max context `⌈max_context / N⌉ × d ×
-  dtype`, for K and V separately. In kernbench these are deployed with a
-  `DPPolicy` whose `pe` (or `cube`) axis is `row_wise` over the sequence
-  dimension (`policy/placement/dp.py`; the baseline uses
-  `DPPolicy(pe="row_wise")` for K/V and `replicate` for Q).
-- Within a PE, assigned tokens are **appended contiguously** (slot
-  0,1,2,…). Global indices are strided (`i, i+N, …`) but the per-PE
-  buffer is dense ⇒ DMA stays contiguous.
-- Slot → global: `global_idx = local_slot·N + ((pe_id − start_pe) mod N)`.
-
-### 2.2 Driver per-launch duties (minimal)
-The driver supplies, per launch, the bases + counter + rotation; the
-kernel derives the rest:
-
-| Launch arg | Purpose |
-|---|---|
-| `K_base[pe]`, `V_base[pe]` | per-PE KV buffer bases (from tensor VA) |
-| `O_base` | attention output destination |
-| `global_token_counter` | current sequence position |
-| `start_pe = request_id mod N` | Layer-2 rotation |
-| `pe_id` (`tl.program_id`), `N` | geometry |
-| `q_block_meta` | prefill/SP query block start + length |
-
-Kernel-derived (plain arithmetic):
-- **my turn to write this step:** `(start_pe + counter) mod N == pe_id`
-- **my valid length:** `base = counter // N; rem = counter % N;
-  my_len = base + (1 if ((pe_id − start_pe) mod N) < rem else 0)`
-- **read range:** tiles `0 .. ⌈my_len / TILE⌉`.
-- **causal bounds / per-tile skip:** from global positions of the query
-  block vs each tile.
-
-Driver = bases + counter + rotation. The single formula
-`(request_id + token_idx) mod N` is the entire placement policy.
-
----
-
-## 3. Per-tile op sequence (greenlet `tl`)
-
-One iteration = one KV tile on one PE. The two GEMMs are
-`tl.composite(op="gemm")` (scheduler-managed tiling + K/V DMA streaming);
-the softmax merge is kernel `tl` MATH between them. Real `tl` names
-(`tl_context.py`), with lazy `tl.load` (ADR-0062), `tl.scratch_scope`
-(ADR-0063):
-
-```python
-# running state (persistent arena — allocated once, outside the scope)
-# m: [G, T_q]  l: [G, T_q]  O: [G, T_q, d]
-q_g = tl.load(Q_group_ptr, (G*T_q, d))               # lazy; auto-wait at first use (ADR-0062)
-
-with tl.scratch_scope():                              # per-tile MATH temporaries recycled
-    Sj    = tl.composite("gemm", a=q_g,                # [G·T_q, TILE]; scheduler streams Kⱼ DMA
-                         b=tl.ref(K_base + j*TILE*d, (TILE, d))) * softmax_scale
-    if mask_j is not None:
-        Sj = Sj + mask_j                              # additive causal mask (boundary tile)
-    m_j   = tl.max(Sj, axis=-1)
-    m_new = tl.maximum(m, m_j)
-    P     = tl.exp(Sj - m_new)                        # no full-matrix softmax; streaming
-    corr  = tl.exp(m - m_new)                         # rescale factor for old accumulators
-    l     = l * corr + tl.sum(P, axis=-1)
-    Oj    = tl.composite("gemm", a=P,                  # [G·T_q, d]; scheduler streams Vⱼ DMA
-                         b=tl.ref(V_base + j*TILE*d, (TILE, d)))
-    O     = O * corr + Oj                             # running merge (kernel MATH)
-    m     = m_new
-```
-
-Notes:
-- `q_g` is the GQA-batched query reshaped to `[G·T_q, d]` (the `G` group
-  rows folded into the matmul M dim; byte-conserving). One K/V tile serves
-  all `G·T_q` rows — the GQA reuse lever — with no broadcast.
-- **K/V are `tl.ref` operands** the composite scheduler streams from HBM
-  per tile (`pe_scheduler.py:104-143`): that *is* the prefetch/pipeline,
-  so there is no explicit prefetch op. Issuing tile `j+1`'s composites
-  before waiting on tile `j` (non-blocking handles) keeps the scheduler
-  pipelined across tiles and the GEMM engine saturated.
-- `tl.trans` is **metadata-only** in kernbench (`tl_context.py:390`) and
-  `MemoryStore.read` *reshapes* rather than transposes
-  (`memory_store.py:73`). For zero/structural runs this is harmless; for
-  non-trivial numeric data it yields a reshape-not-transpose, so Q·Kᵀ via
-  a transposed K needs care (§11) — store K pre-transposed `[d, TILE]`, or
-  add a real `tl.transpose` (a candidate further primitive, likely
-  unnecessary given the simulator's performance-modeling purpose).
-- Masking: the kernel builds the boundary-tile mask from query/KV global
-  offsets and adds it (`Sj + mask_j`); full-past tiles pass `None`;
-  full-future tiles are **skipped** (the `if` never enqueues them).
-- A **new** "flash-composite" command kind (one that internalises the
-  softmax merge + carried `(m,ℓ,O)`) is **not** used; the existing
-  `CompositeCmd` covers each GEMM and the merge stays in the kernel
-  (§1, §8 item 4).
-
----
-
-## 4. Reduction (KV-parallel / SP combine) — tree to root
-
-After its tile sweep each PE holds `(m_i, ℓ_i, O_i)` (unnormalised).
-Combine via the associative/commutative log-sum-exp merge (identical math
-to the baseline's fold, `_attention_mesh_mlo.py:117-122`):
-
-```python
-def merge(m_a, l_a, O_a, m_b, l_b, O_b):
-    m = tl.maximum(m_a, m_b)
-    sa, sb = tl.exp(m_a - m), tl.exp(m_b - m)
-    return m, l_a*sa + l_b*sb, O_a*sa + O_b*sb
-# final: O = O_root / l_root      # normalise once at the tree root
-```
-
-- **Timing:** exchange is **after P·V** (each PE has its final `O_i`).
-- **Topology:** mesh tree, depth `⌈log₂ N⌉`, over physical neighbours.
-  For `N=8`: level-0 `(0↔1,2↔3,4↔5,6↔7)`, level-1 `(1↔3,5↔7)`, level-2
-  `(3↔7)`, root = `7`. Each tree pair must be a physical `N/S/E/W`
-  neighbour so `tl.send(dir)`/`tl.recv(dir)` use a real link
-  (tuning item §9; the SFR install that wires neighbours is
-  `configure_sfr_intracube_pe_ring`, used by the baseline).
-- **Why a tree, not the baseline fan-out:** attention needs `O` at one
-  place (the PE that writes `O_base`). A tree is `⌈log₂ N⌉` steps
-  (3 for `N=8`) vs the baseline all-to-all's `N−1` (7). For decode, where
-  the per-PE sweep is short, the reduction is the dominant cost, so this
-  is a real win. (The baseline's fan-out replicates the answer to all
-  ranks — unnecessary here.)
-- **Payload:** `O_i` (`d`-vector) is heavy; `m,ℓ` are scalars per `(g,
-  T_q)`. Sent as separate `tl.send`s (the baseline sends the triplet as
-  three sends — same pattern).
-
-Kernel structure (greenlet):
-
-```python
-# leaf / internal node: fold children that send to me, then forward up
-for child_dir in tree_children_dirs(pe_id, N):
-    m_c = tl.recv(child_dir, m.shape); l_c = tl.recv(child_dir, l.shape)
-    O_c = tl.recv(child_dir, O.shape)
-    m, l, O = merge(m, l, O, m_c, l_c, O_c)
-if not is_root(pe_id):
-    tl.send(parent_dir(pe_id), m); tl.send(parent_dir(pe_id), l); tl.send(parent_dir(pe_id), O)
-else:
-    tl.store(O_base, O / l)
-```
-
-`tl.recv` blocks until data arrives (`tl_context.py:446-499`); the merge
-order is fixed by the static tree, so no data-dependent `pop` is needed —
-which is why **no hardware/composite `pop`-as-dependency change is
-required** (§6).
-
----
-
-## 5. Per-case kernels
-
-All four cases share §3 (tile sweep) + §4 (reduction); they differ only
-in `N`, query-block width, and which `tile_*` predicates fire.
-
-### 5.1 Common skeleton
-
-```python
-def attn_kernel(q_ptr, k_ptr, v_ptr, o_ptr, counter, start_pe, N,
-                q_block, softmax_scale, *, tl):
-    pe_id   = tl.program_id(axis=rank_axis)
-    my_len  = valid_len(counter, start_pe, pe_id, N)
-    n_tiles = ceil(my_len / TILE)
-    q_g     = load_Q_group(q_ptr)            # [G·T_q, d]; lazy tl.load, G folded into M (no broadcast)
-    m, l, O = init_running()                 # persistent arena: -inf, 0, zeros
-    for j in range(n_tiles):                 # K/V streamed per tile by the composite scheduler (§3)
-        if tile_all_future(j, q_block):      # causal skip (kernel if)
-            continue
-        run_tile(j, mask_or_null(j, q_block))   # §3: 2 composites + softmax MATH, in tl.scratch_scope
-    if N == 1:
-        tl.store(o_ptr, O / l)               # no reduction
-    else:
-        tree_reduce_and_store(pe_id, N, o_ptr)  # §4
-```
-
-### 5.2 DECODE, no SP (`N=1`, one PE owns the head's KV)
-- `T_q = 1`, attends to all past KV ⇒ no future tiles, mask only on the
-  final ragged tile.
-- **GQA reuse is the whole game** (decode is KV-load-bound): the `G=8`
-  query rows of the KV head are folded into the matmul **M** dimension
-  (`q_g` reshaped `[G, T_q, d] → [G·T_q, d]`, a byte-conserving reshape).
-  Then `Q·Kᵀ` is `composite([G·T_q, d], Kᵀ[d, TILE]) → [G·T_q, TILE]` and
-  `P·V` is `composite([G·T_q, TILE], V[TILE, d]) → [G·T_q, d]`. The KV tile
-  (`[TILE, d]`) is the shared `K`/`V` operand — **streamed once, reused by
-  all `G·T_q` rows automatically** because they are the M rows of the
-  GEMM. No broadcast of K/V is needed; `m = G·T_q` in the composite's tile
-  plan also makes the timing count all `G` rows' work correctly (a leading
-  batch axis would *not* be counted — see §8).
-- If `S_pe` fits in scratch (small/medium context) this degenerates to a
-  **one-shot** partial attention (one composite for Q·Kᵀ, one softmax, one
-  composite for P·V) — exactly the baseline `_attention_mesh_mlo`
-  `_partial_attention`, just GQA-batched and on the composite path. Tiling
-  (§3) only kicks in when `S_pe` exceeds the scratch scope's tile budget.
-
-### 5.3 DECODE, with SP / KV-parallel (`N=8`)
-- One request's KV is round-robin across 8 PEs; each owns ≈`my_len`
-  tokens. Each PE GQA-batches its `G=8` query rows over its local tiles.
-- `T_q=1` ⇒ short per-PE sweep ⇒ **reduction dominates** ⇒ the §4 tree
-  (3 steps) is the structurally important part. Reduction latency is
-  hidden by other concurrent decode tokens when batched, or by the long
-  per-PE tile sweep for long single-stream context.
-
-### 5.4 PREFILL, no SP
-- Whole prompt resident; query is a block of `T_q` tokens (chunked).
-- Causality is real, per query block `[qs, qe)` vs KV tile `[ks, ke)`:
-  - `ke ≤ qs` → `tile_all_past` → `mask=None`, full compute.
-  - `ks ≥ qe` → `tile_all_future` → **skip** (kernel `if`).
-  - overlap → `tile_partial` → kernel builds a triangular additive mask,
-    the tile op sequence adds it (`Sj + mask_j`).
-- GQA batches the group's query heads along the same axis as decode.
-
-### 5.5 PREFILL, with SP (Ring KV)
-- KV sharded across `N` PEs (the ring); each step delivers a peer's KV
-  block over IPCQ. Running `(m,ℓ,O)` is carried **across ring steps**
-  (Python handles, exactly as `_attention_mesh_kv` does today).
-- IPCQ overlaps next step's KV receive with current step's compute via
-  `tl.recv_async`/`tl.wait` (`tl_context.py:543-560` — already exists).
-- **Causal ring skip:** if the incoming step's KV block is entirely
-  *after* the local query block, skip its compute (don't recv-consume /
-  don't fold) — eliminates ≈half the steps for causal attention.
-- Receive buffers ping-pong (persistent arena, not recycled).
-- After the ring, `(m,ℓ,O)` is final per query-owning PE; if KV-parallel
-  splits remain, §4 tree-merges them, then normalise + store.
-
-The baseline `_attention_mesh_kv` already implements the ring fold; this
-ADR adds GQA reuse, causal step-skip, and `recv_async` overlap.
-
----
-
-## 6. Why no hardware / composite change is needed
-
-- The only place a HW/composite `pop`-as-dependency would help is a lone
-  reduction with no other work to hide a poll — i.e. batch=1 **and** short
-  context. The target is **agentic = low batch, long context** ⇒ each PE
-  has many KV tiles; the §4 tree's `tl.recv` blocks are covered by the
-  sweep work of concurrent tokens / long context.
-- The §4 reduction uses a **static** tree, so the recv order is fixed at
-  compile time — no data-dependent pop. `tl.recv` (blocking) suffices.
-- **Decision: ship with the greenlet `tl.send`/`tl.recv` collective.**
-  Revisit a HW `pop` only if a short-context, single-stream,
-  latency-critical target emerges.
-
----
-
-## 7. Control vs execution split (the load-bearing principle)
-
-| Concern | Owner |
-|---|---|
-| tile skip (future), mask generation, causal bounds | **kernel** (Python `if` + arithmetic in greenlet body) |
-| address / offset / valid-length arithmetic | **kernel** (from counter) |
-| reduction scheduling, IPCQ send/recv ordering | **kernel** (static tree) |
-| Q·Kᵀ, P·V (incl. per-tile K/V DMA streaming + tiling) | **`tl.composite`** → PE_SCHEDULER |
-| Q load, mask add, softmax math, running `(m,ℓ,O)` merge | **`tl` ops** on PE engines (kernel-issued) |
-
-The kernel decides; the GEMMs are offloaded to the scheduler as
-composites; the remaining `tl` ops execute already-decided work. This is
-exactly the greenlet + composite model kernbench already supports — no new
-control abstraction.
-
----
-
-## 8. Required kernbench changes
-
-**Correction from design iteration:** real GQA (`h_q > h_kv`) needs **no
-new primitive** — only the kernel restructuring in §5.2 (per KV head,
-`G` folded into M, byte-conserving reshapes). The supporting ADRs are
-*efficiency / scale* enablers, not GQA blockers.
-
-**Algorithm work in the kernel (no new primitive; existing `tl` API):**
-
-- **GQA Q-axis batching** (the reuse lever) — fold `G·T_q` into the matmul
-  M dim per KV head (§5.2); `_view`-style byte-conserving reshape; the
-  GEMM is a `tl.composite(op="gemm")` with M = `G·T_q`. Runs today in both
-  timing and data mode.
-- **GEMMs via composite** (§1/§3) — Q·Kᵀ and P·V each issued as a
-  non-blocking `tl.composite(op="gemm")`; PE_SCHEDULER tiles them and
-  streams the `tl.ref` K/V operands' DMA (existing `CompositeCmd`; no new
-  command kind).
-- Tree reduction to root (§4) replacing the baseline all-to-all fan-out —
-  pure kernel control flow over `tl.send`/`tl.recv`.
-- Causal tile skip + additive boundary mask (§3/§5.4) — kernel `if` +
-  a mask tensor added with `+`.
-- Round-robin KV placement / valid-length arithmetic (§2) — launch-arg
-  arithmetic + `DPPolicy(pe="row_wise")`.
-
-**New primitives for efficiency / scale (each has a supporting ADR):**
-
-1. **Per-tile scratch recycling** — **ADR-0063** (`tl.scratch_scope`).
-   *Required for scale*: removes the `S=16` ceiling (1 MiB bump
-   allocator) so realistic context lengths run. Highest-value of the
-   three.
-2. **Lazy `tl.load`** — **ADR-0062** (non-blocking load + auto-wait on
-   first use; API surface unchanged). *Efficiency*: overlaps explicit
-   loads (the Q group, non-composite kernels) with following compute. The
-   per-tile **K/V** prefetch is handled by the composite scheduler (§1),
-   so this covers the remaining explicit loads. Global semantics change →
-   existing goldens regenerate (ADR-0062 D3).
-3. **GQA head / mask broadcast** — **ADR-0061** (`tl.broadcast`).
-   *Optional convenience*, not a GQA blocker (see correction above).
-   Useful for additive-mask construction across the `G·T_q` rows and for
-   general kernels; `np.matmul` already broadcasts in data mode, so it is
-   not needed for correctness. Lowest priority.
-
-**Explicitly REJECTED (efficient alternative chosen):**
-
-4. ~~A bespoke "flash-composite" command kind that internalises the whole
-   inner loop — DMA→MM→VEC→DMA→MM→VEC with carried `(m,ℓ,O)` register
-   state and a tail IPCQ push.~~ The two GEMMs **do** use the existing
-   `CompositeCmd` (§1/§3) — that gives scheduler-managed tiling, K/V DMA
-   streaming, and cross-tile pipelining. What is rejected is a **new**
-   command kind that also absorbs the softmax merge + cross-tile register
-   lifetime + an IPCQ-push epilogue: it is large and special-purpose, and
-   its only delta over the hybrid (full softmax offload) is not justified
-   at the current modelling fidelity (per-op CPU issue cost = 0; see §1).
-   Revisit if the cost model (§9) makes full offload measurably worthwhile.
-5. ~~Hardware `pop`-as-dependency.~~ Out of scope (§6).
-6. ~~RoPE / QKV projection / KV-cache write inside this kernel.~~ Upstream
-   `qkv_rope` (P1–P5). Folding RoPE in would force re-rotating past tiles
-   every decode step and break Ring Attention's post-RoPE pass-through.
-
----
-
-## 9. Open tuning items (measured in kernbench, not blocking)
-
-1. **Composite tile-pipeline depth** — dominant for KV-load-bound; how far
-   ahead the kernel issues non-blocking composites before waiting, and the
-   scheduler's per-tile streaming depth.
-2. **PE↔mesh-neighbour mapping for the reduction tree** — ensure each
-   depth-`⌈log₂ N⌉` pair is a physical `N/S/E/W` neighbour; bad mapping
-   adds hops. Verify against the SFR install.
-3. **TILE size** — balance scratch residency (`S/P tiles + O_acc + G-way
-   GQA`) against DMA efficiency; interacts with ADR-0063 and the
-   scheduler's `TILE_M/K/N` (`pe_scheduler.py`).
-4. **Ring buffer ping-pong vs `recv_async` depth** in §5.5.
-5. **One-shot vs tiled crossover for decode** (§5.2) — the `S_pe`
-   threshold where tiling beats a single composite.
-6. **Per-op CPU issue cost (cost model)** — currently `dispatch_cycles=0`
-   (`pe_cpu.py`), so composite-vs-primitive issue overhead is invisible.
-   An op-type-differentiated issue cost (a `tl.composite` descriptor push
-   ≫ a primitive op) is what makes the hybrid's CPU-saturation win
-   **measurable** (§1). Specified in **ADR-0064**; tracked as separate
-   future work.
-
----
-
-## 10. Coverage summary
-
-| Case | KV placement | Inner loop | Reduction | Masking |
-|---|---|---|---|---|
-| Decode, no SP | 1 PE, all KV | one-shot or tiled, GQA-batched | none | last tile only |
-| Decode, SP | round-robin `N` PEs | tiled, GQA-batched | §4 tree (`⌈log₂ N⌉`) | last tile only |
-| Prefill, no SP | resident | tiled per q-block | none | skip future / triangular boundary |
-| Prefill, SP (Ring) | ring-sharded | per ring step (recv_async overlap) | running state + §4 tree | causal step-skip + boundary |
-
-**I/O per case** (full contract §0.5):
-
-| Case | Inputs | Output | IPCQ partials |
-|---|---|---|---|
-| Decode, no SP | `Q[G,1,d]`, full `K/V[S,d]` on 1 PE | `O[G,1,d]` at `O_base` | none |
-| Decode, SP | `Q[G,1,d]`, per-PE `K/V[S_pe,d]` | `O[G,1,d]` (root) | `(m,ℓ,O_i)` per PE → tree |
-| Prefill, no SP | `Q[G,T_q,d]`, `K/V[≤end,d]` | `O[G,T_q,d]` at `O_base` | none |
-| Prefill, SP (Ring) | `Q[G,T_q,d]`, ring-delivered `K/V` blocks | `O[G,T_q,d]` (root) | running `(m,ℓ,O)` + tree |
-
-In all cases: KV-cache writes and RoPE happen **upstream**; output
-projection **downstream**; score `S` and probs `P` are never materialised.
-
----
-
-## 11. Verification plan (Phase 1 test outline)
-
-SPEC/ADR coverage: R5 (PE↔PE IPCQ, PE↔HBM), R2 (latency by traversal),
-ADR-0023/0025 (IPCQ), ADR-0046 (`tl` contract), ADR-0054 (eval bench).
-
-The simulator's contract is **latency by traversal + determinism +
-structural correctness** (SPEC §0, §0.1), not bit-exact numerics — Phase 2
-data exists mainly to exercise the data path, `tl.trans` is
-reshape-not-transpose, and `bf16` is modelled as `f16`. Verification is
-therefore **structural/timing-first**, with numeric parity as a bounded
-secondary check.
-
-1. **Runs in data mode (`enable_data=True`):** the GQA kernel
-   (`h_q = G·h_kv`) completes without the byte-conservation error that the
-   baseline head-packing hits — for all four cases. (This needs the §5.2
-   restructuring, *not* a new primitive.)
-   **Numeric parity (secondary):** for symmetric/identity inputs where
-   reshape-as-transpose is exact, kernel `O` matches a numpy
-   FlashAttention reference within fp tolerance. Full asymmetric parity is
-   gated on a real `tl.transpose` (out of scope; flagged).
-2. **GQA reuse:** with `h_q = G·h_kv`, the K/V `dma_read_count` is
-   independent of `G` (one load per tile, reused across the group), while
-   GEMM work scales with `G`. Asserts the lever actually fires.
-3. **Tree reduction step count:** decode-SP issues `⌈log₂ N⌉` reduction
-   rounds (not `N−1`); op_log `ipcq_send`/`recv` counts match the tree.
-4. **Causal skip:** prefill skips all `tile_all_future` tiles — GEMM
-   count equals the lower-triangular tile count, not the full grid.
-5. **Long context (ADR-0063):** a sweep at `S` that overflows 1 MiB
-   without scopes completes and matches the reference.
-6. **Load/compute overlap:** end-to-end latency of the tiled sweep is
-   below the serial `Σ(load+compute)` — from the composite scheduler
-   streaming K/V per tile (§1/§3) and lazy `tl.load` (ADR-0062) overlapping
-   the Q load. (Overlap is real modelled concurrency, not a subtraction.)
-7. **Composite GEMM offload (structural):** each tile's Q·Kᵀ and P·V emit a
-   `CompositeCmd` (non-blocking) to PE_SCHEDULER, not a blocking `tl.dot`;
-   op_log shows the composite tile plan and the kernel issues the next
-   tile's composites before waiting (cross-tile pipelining).
-8. **Determinism:** identical inputs → identical op_log + latency
-   (SPEC §0.1).
-
----
-
-## B. Open design items from the hybrid pivot (review later)
-
-These arose when the decision moved from a pure greenlet primitive path to
-the **composite hybrid + lazy `tl.load`** (this revision). None blocks the
-design; each needs a verification pass during implementation. Recorded here
-(rather than asked) per the working agreement — the recommendation is my
-predicted default; revise on review.
-
-1. **DDD-0060 is not yet synced.** The Detailed Design Document still
-   describes the old `tl.load_async` double-buffer path and primitive
-   `tl.dot` inner loop (its §4.3/§5/§10). It must be updated to the hybrid
-   (composite GEMMs, lazy load, K pre-transposed). *Left for review*
-   because the DDD is a derived how-to and a large rewrite; the ADR is now
-   the authoritative record. **Recommend:** sync DDD as a follow-up before
-   implementation starts.
-
-2. **K operand orientation for the composite GEMM.** Q·Kᵀ needs `b =
-   [d, TILE]`, but the KV cache stores K as `[S_pe, d]`. `tl.trans` is
-   metadata-only and `MemoryStore.read` reshapes, not transposes
-   (`memory_store.py:73`) — so a runtime transpose is wrong for non-trivial
-   data. **Recommend:** store K **pre-transposed** `[d, S_pe]` in the cache
-   (the pseudocode and §3 assume this), making `tl.ref(k_tile, (d, TILE))`
-   a contiguous slice. Verify the upstream `qkv_rope` write layout supports
-   this, or add a real `tl.transpose` (heavier; deferred).
-
-3. **Composite output buffer vs `tl.scratch_scope`.** Each Q·Kᵀ composite
-   writes `Sj` to an `out_addr`; the kernel then reads it for the softmax
-   MATH. That output buffer, and the in-flight composites' targets, must
-   live where the per-tile `scratch_scope` (ADR-0063) will **not** recycle
-   them before they are consumed — same discipline as in-flight lazy loads
-   (ADR-0062 D-Negative). **Recommend:** composite outputs for the *current*
-   tile live in the scoped arena (consumed same iteration); the persistent
-   `(m,ℓ,O)` stays outside. Verify no use-after-recycle when the next
-   tile's composites are issued early (cross-tile pipelining).
-
-4. **GQA `dma_read_count` lever under composite streaming.** The lever
-   (§11.2: K/V `dma_read_count` independent of `G`) assumes the composite
-   emits **one** K/V tile DMA reused across all `G·T_q` M-rows. The
-   scheduler's `generate_gemm_plan` tiles by `TILE_M/K/N`
-   (`pe_scheduler.py:35-37`, 32/64/32) — confirm the M-tiling over `G·T_q`
-   does **not** re-issue the shared K/V tile DMA per M-tile (i.e. operand
-   DMA is shared across M-tiles, or the lever weakens). **Recommend:**
-   assert it in the levers test; if violated, the GQA win is in compute
-   only, not DMA — still correct, but the headline changes.
-
-5. **Kernel TILE vs scheduler `TILE_M/K/N`.** The kernel reasons about a
-   logical KV `TILE`; the scheduler re-tiles internally at fixed
-   `TILE_M/K/N`. Two tiling layers interact (scratch residency, pipeline
-   depth). **Recommend:** treat the kernel TILE as the K/V streaming
-   granularity and let the scheduler sub-tile the GEMM; document the
-   relationship in the DDD and sweep both (§9 items 1, 3).
-
-6. **Cost model is a separate ADR.** The hybrid's CPU-saturation benefit is
-   invisible while `dispatch_cycles=0`. The per-op-type issue-cost model is
-   specified in **ADR-0064**; this ADR's §1/§9 depend on it for the
-   *measurable* (not just structural) win. **Recommend:** land ADR-0064's
-   model before claiming hybrid latency wins in the eval.
-
-7. **Ring path (§5.5) GEMMs.** §5.5 still describes the ring fold with
-   primitive ops + `recv_async`. For consistency the ring's per-step Q·Kᵀ /
-   P·V should also be composites; the IPCQ `recv_async` overlap is
-   orthogonal and stays. **Recommend:** apply the same hybrid shape to the
-   ring step during implementation; low risk, mirrors §3.

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

@@ -32,9 +32,8 @@ the group is the decode-time efficiency win (ADR-0060 §5.2).
 
 ### Why it does not work today
 
-The existing mesh kernels reshape tensors with a metadata-only helper
-`_view` (`src/kernbench/benches/_attention_mesh_kv.py:25-36`,
-`_attention_mesh_mlo.py:25-36`):
+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):
@@ -62,8 +61,7 @@ A reshape conserves bytes; a **broadcast does not** (`[S_kv, 1, d]` →
 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 current test suite documents this as a deliberate limitation
-(`tests/attention/test_milestone_gqa_llama70b.py:137-142`):
+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
@@ -196,7 +194,7 @@ entire point of GQA (it `G×`'s KV-cache HBM footprint and bandwidth).
 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-llama70b` panel
+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

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

@@ -1,166 +0,0 @@
-# ADR-0064: op 종류별 CPU 발행 비용 모델 (command construct + dispatch)
-
-## Status
-
-Proposed
-
-> **ADR-0060**(AHBM GQA Fused Attention)의 보조 ADR. 거기서의 하이브리드
-> 결정(GEMM은 `tl.composite`, softmax 머지는 커널)이 이기는 이유는
-> **타일링을 PE_SCHEDULER에 offload하여 CPU가 coarse descriptor를 내리고
-> 앞서 나가 엔진을 saturate**하기 때문이다. 그 이점이 현재 시뮬레이터에서는
-> per-op CPU 발행 비용이 0이라 **보이지 않는다**. 본 ADR은 발행 비용을
-> 실재화하고 **op 종류별로 차등**화하여, composite-vs-primitive 트레이드오프
-> (그리고 CPU saturation)가 측정 가능하도록 한다.
-
-## Context
-
-### 현재 상태
-
-- 모든 `tl.*` op은 커맨드를 emit하기 전에 `_emit_dispatch_overhead()`를
-  호출하며(`tl_context.py:123-125, 190, 227, 235, 612, …`), 이는
-  `dispatch_cycles > 0`일 때**만** `PeCpuOverheadCmd(cycles=dispatch_cycles)`를
-  emit한다. 즉 발행 비용은 `tl.load`, `tl.dot`, MATH op, `tl.composite`에
-  동일하게 적용되는 **단일 균일 노브**다.
-- 그 노브는 두 실행 경로 모두에서 **0**으로 하드코딩되어 있다
-  (`pe_cpu.py:101` greenlet runner, `:195` legacy replay). ⇒ 커맨드 발행
-  — *descriptor를 construct하고 scheduler 큐에 push* — 이 현재 PE_CPU에서
-  **0 ns**다.
-- `PeCpuOverheadCmd`는 PE_CPU(`kernel_runner.py:131-132`)와
-  PE_SCHEDULER(`pe_scheduler.py:97-100`)에서 `yield env.timeout(cmd.cycles)`로
-  소비된다. 수동 `tl.cycles(n)`도 존재한다(`tl_context.py:695`).
-
-### 본 ADR을 규정하는 두 가지 발견 (ADR-0060 검토에서)
-
-- **Q1 — composite 발행 비용.** `tl.composite`를 construct + push 하는 비용은
-  CPU 시간 기준 **~40 ns** 정도(descriptor 빌드 + 큐 push)로 예상된다. 현재는
-  **0**이다. 훅은 존재하며, 값(그리고 op 종류별 차등)만 빠져 있다.
-- **Q2 — scheduler dispatch vs DMA latency.** PE_SCHEDULER의 composite
-  dispatch는 **non-blocking**이다: `_dispatch_composite`는 tile plan을 생성하고
-  feeder에 enqueue한 뒤 즉시 반환한다(`pe_scheduler.py:104-121`). 실제 **DMA
-  latency는 타일이 흐를 때 PE_DMA에서 부과**되며(`drain_ns =
-  compute_drain_ns(path, nbytes)`, `pe_dma.py:89`), scheduler dispatch에
-  덩어리로 들어가지 **않는다** ⇒ **중복 계상 없음**, DMA는 modeled component에
-  머문다(SPEC §0.1). scheduler의 plan 생성 자체는 현재 sim 시간 **0**이다.
-
-### 균일-그리고-0이 하이브리드에 틀린 이유
-
-하이브리드의 논거 전체는 **하나의** composite descriptor가 `N_tiles`만큼의
-GEMM 타일링을 offload하므로, CPU가 `O(N_tiles × ops/tile)`개의 fine 커맨드
-대신 `O(1)`개의 coarse 커맨드를 낸다는 것이다. 발행 비용 = 0(그리고 균일)이면
-모델은 다음을 보여줄 수 없다:
-- CPU가 충분히 빨리 못 밀어넣을 때 primitive 경로가 엔진을 **saturate 못 할 수
-  있음**(핵심 ADR-0060 §1 주장), 그리고
-- composite가 단일 primitive보다 construct 비용이 **더 크지만** 그것이 대체하는
-  다수의 primitive보다는 훨씬 작음.
-
-단일 균일 `dispatch_cycles`로는 이를 표현할 수 없다: 실제 construct 비용에서
-`tl.load` ≪ `tl.composite`이기 때문이다.
-
-## Decision
-
-### D1. PE_CPU의 op 종류별 발행 비용 테이블
-
-단일 `dispatch_cycles` 스칼라를 **커맨드 종류별 비용 테이블**로 교체하고,
-발행 시점(커맨드가 scheduler로 dispatch되기 전)에 PE_CPU에서 부과한다.
-greenlet이 이 시간을 지불하며, 이것이 바로 CPU가 얼마나 빨리 work를 밀어넣을
-수 있는지를 gate한다 — saturation 레버.
-
-시작 추정치(추후 calibrate — 검토 항목 참조):
-
-| 발행 op | CPU 발행 비용 (construct + push) |
-|---|---|
-| `tl.composite` (GEMM/MATH descriptor) | ~40 ns (Q1 추정) |
-| `tl.load` / `tl.store` (DMA descriptor) | 작음 (수 ns) |
-| `tl.dot` / MATH / IPCQ send/recv | 작음 (수 ns) |
-
-요점은 **비율**이다: composite construct ≫ primitive 발행, 그러나 그것이
-대체하는 primitive 발행들의 합 ≪. 정확한 ns는 설정 가능하다.
-
-### D2. 실행 latency는 엔진에 유지 (Q2 — 변경 없음)
-
-DMA/GEMM/MATH 실행 latency는 오늘처럼 PE_DMA / PE_GEMM / PE_MATH에서 부과된다.
-발행 비용(D1)은 **CPU 측에만** 추가되며, `drain_ns`를 건드리지 않으므로 중복
-계상이 없다. 이는 SPEC §0.1(latency는 modeled component에서)을 보존한다.
-
-### D3. scheduler plan 생성 비용 — 0에서 시작, 재검토
-
-PE_SCHEDULER의 tile-plan 생성은 오늘 sim 시간 0이다. 일단 유지한다(지배적
-레버는 CPU 발행 비용 D1); scheduler 측 비용이 유의미하다고 판명되면 기존
-`overhead_ns` node attr로 노출한다. 여기서 결정하지 않고 검토 항목으로 둔다.
-
-### D4. 설정 가능한 값; 골든 재생성
-
-비용 테이블은 설정 가능하다(per-topology / node attrs, 기본값은 D1 추정치).
-발행 비용을 0이 아니게 하면 **모든** bench latency가 바뀌므로, 골든 latency를
-한 번 **재생성**한다 — ADR-0062의 전역 lazy-load 회귀와 동일한 자세의 모델
-충실도 개선이다.
-
-## Alternatives
-
-### A1. 단일 균일 `dispatch_cycles > 0` 유지
-
-기각: op construct 비용은 자릿수 단위로 다르다(`tl.load` vs `tl.composite`).
-균일값은 primitive를 과대 청구하거나 composite를 과소 청구하며, 어느 쪽이든
-하이브리드가 의존하는 composite-vs-primitive 트레이드오프를 왜곡한다.
-
-### A2. 발행 비용을 PE_CPU 대신 PE_SCHEDULER에 부과
-
-기각: saturation 질문은 *"CPU가 엔진을 바쁘게 유지할 만큼 descriptor를 빨리
-밀어넣을 수 있는가?"*이며 — 이는 **PE_CPU**의 issue-bandwidth 속성이다.
-scheduler에 부과하면 CPU back-pressure를 모델하지 못한다.
-
-### A3. DMA program/setup 시간을 별도의 고정 per-descriptor 비용으로 모델
-
-실제 HW는 전송 시간과 별개로 DMA descriptor program 비용을 지불한다. 연기:
-초기에는 descriptor-program 비용을 **발행 op**의 D1 비용에 합친다(DMA를
-유발하는 `tl.load` / composite). calibration이 유의미함을 보이면 PE_DMA 고정
-setup으로 분리한다. 검토 항목.
-
-## Consequences
-
-### Positive
-- 하이브리드의 CPU-offload / saturation 이점(ADR-0060 §1)이 구조적일 뿐
-  아니라 **측정 가능**해진다.
-- composite-vs-primitive 및 타일링 granularity 트레이드오프가 latency에
-  보여 eval 서사를 가능케 한다.
-- 하드웨어에 더 충실(issue bandwidth는 실제 병목).
-
-### Negative
-- **모든** bench 골든이 이동 → 한 번의 재생성(D4); CI 골든 fixture 갱신.
-- op 종류별 값은 calibration 필요; ~40 ns composite 수치는 추정이고
-  primitive는 미지정 — 결과는 숫자만큼만 정확하다(calibrate 전까지 절대
-  latency를 과대 주장 금지).
-- 발행 경로에 비용 테이블 lookup 추가(런타임 영향 미미).
-
-## Open review items (자율 결정; 검토 시 수정)
-
-1. **op 종류별 값의 calibration 출처.** composite ≈ 40 ns는 작업 추정치;
-   primitive 발행 비용은 placeholder. *권고:* 문서화된 가정(instruction-issue +
-   큐 push)에서 값을 택하고, 실제 reference가 생기기 전까지 절대 latency를
-   잠정으로 취급; **비율**을 방어 가능하게 유지.
-2. **scheduler plan-gen 비용 (D3).** *권고:* 초기에는 0 유지; scheduler-bound
-   동작을 보이는 workload가 나오면 `overhead_ns`로 노출.
-3. **DMA program 시간 (A3).** *권고:* 먼저 발행 op의 비용에 합치고; 필요 시
-   PE_DMA setup으로 분리.
-4. **테이블 위치.** *권고:* 흩어진 node attr이 아니라, 커맨드 종류로 키잉되고
-   per-topology 오버라이드 가능한 작은 중앙 cost 모듈 — 값을 한 곳에서 검토
-   가능하도록.
-5. **회귀 롤아웃.** *권고:* lazy-load(ADR-0062)와 본 ADR의 cost model을 한
-   번의 골든 재생성 패스로 함께 도입하여 두 번의 churn을 피함.
-6. **legacy replay 경로와의 상호작용**(`pe_cpu.py:_execute_legacy`) — greenlet과
-   replay 경로가 동일한 cost 테이블을 읽어 결과가 일치하도록 보장. 검증 요.
-
-## Test Requirements
-
-1. **composite는 한 번, primitive는 op마다 청구.** `N_tiles`에 걸쳐 하나의
-   `tl.composite`를 내는 커널은 PE_CPU에서 composite 발행 비용을 한 번 청구하고;
-   동등한 `N_tiles × ops` primitive 커널은 op당 비용을 `N_tiles × ops`번 청구.
-   PE_CPU busy time이 그에 따라 다름을 assert.
-2. **DMA latency 불변 (Q2).** 고정된 전송에 대해 PE_DMA `drain_ns`는 ADR 이전
-   값과 동일 — 발행 비용은 DMA에 합쳐지지 않고 PE_CPU에 가산(중복 계상 없음).
-3. **saturation 관측 가능.** per-op 발행 비용이 0이 아니면, many-tile primitive
-   sweep는 GEMM 엔진 idle(CPU-bound 발행)을 보이고 composite sweep는 바쁘게
-   유지 — ADR-0060 §1 레버.
-4. **결정성:** 동일 입력 → 동일 op_log + latency (SPEC §0.1).
-5. **경로 일치:** greenlet과 legacy-replay 경로가 동일 커널에 대해 동일한 발행
-   비용 회계를 산출.

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

@@ -1,179 +0,0 @@
-# ADR-0064: Per-op-type CPU issue cost model (command construct + dispatch)
-
-## Status
-
-Proposed
-
-> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention). 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. This ADR makes the issue cost real and **op-type-differentiated**,
-> so composite-vs-primitive trade-offs (and CPU saturation) are measurable.
-
-## Context
-
-### What exists today
-
-- Every `tl.*` op calls `_emit_dispatch_overhead()` before emitting its
-  command (`tl_context.py:123-125, 190, 227, 235, 612, …`), which emits
-  `PeCpuOverheadCmd(cycles=dispatch_cycles)` **only if** `dispatch_cycles
-  > 0`. So the issue cost is a **single uniform knob** applied identically
-  to `tl.load`, `tl.dot`, a MATH op, and `tl.composite`.
-- That knob is hardcoded to **0** in both live execution paths
-  (`pe_cpu.py:101` greenlet runner, `:195` legacy 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`) and on PE_SCHEDULER
-  (`pe_scheduler.py:97-100`); a manual `tl.cycles(n)` also exists
-  (`tl_context.py:695`).
-
-### Two findings that frame this ADR (from ADR-0060 review)
-
-- **Q1 — composite issue cost.** Constructing + pushing a `tl.composite`
-  is expected to cost on the order of **~40 ns** of CPU time (descriptor
-  build + queue push). Today it is **0**. The hook exists; only the value
-  (and its per-op-type differentiation) is missing.
-- **Q2 — scheduler dispatch vs DMA latency.** `PE_SCHEDULER`'s composite
-  dispatch is **non-blocking**: `_dispatch_composite` generates the tile
-  plan and enqueues to the feeder, returning immediately
-  (`pe_scheduler.py:104-121`). The actual **DMA latency is charged on
-  PE_DMA** as each tile flows through (`drain_ns = compute_drain_ns(path,
-  nbytes)`, `pe_dma.py:89`), **not** lumped into the scheduler dispatch ⇒
-  **no double-counting**, and DMA stays on its modelled component
-  (SPEC §0.1). The scheduler's own plan-generation currently costs **0**
-  sim time.
-
-### Why uniform-and-zero is wrong for the hybrid
-
-The hybrid's whole 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 (and
-uniform), the model cannot show:
-- that the primitive path may **fail to saturate** the engines when the
-  CPU cannot push fast enough (the core ADR-0060 §1 claim), nor
-- that a composite **costs more to construct** than a single primitive but
-  far less than the many primitives it replaces.
-
-A single uniform `dispatch_cycles` cannot express this: `tl.load` ≪
-`tl.composite` in real construct cost.
-
-## Decision
-
-### D1. Per-op-type issue-cost table on PE_CPU
-
-Replace the single `dispatch_cycles` scalar with a **cost table keyed by
-command type**, charged on PE_CPU at issue (before the command is
-dispatched to the scheduler). The greenlet pays this time, which is
-exactly what gates how fast the CPU can push work — the saturation lever.
-
-Starting estimates (calibrate later — see review items):
-
-| Issued op | CPU issue cost (construct + push) |
-|---|---|
-| `tl.composite` (GEMM/MATH descriptor) | ~40 ns (Q1 estimate) |
-| `tl.load` / `tl.store` (DMA descriptor) | small (a few ns) |
-| `tl.dot` / MATH / IPCQ send/recv | small (a few ns) |
-
-The point is the **ratio**: a composite construct ≫ a primitive issue, but
-≪ the sum of the primitive issues it replaces. Exact ns are configurable.
-
-### D2. Execution latency stays on the engines (Q2 — no change)
-
-DMA/GEMM/MATH execution latency remains charged on PE_DMA / PE_GEMM /
-PE_MATH as today. The issue cost (D1) is **CPU-side only** and additive; it
-does **not** touch `drain_ns`, so there is no double-count. This preserves
-SPEC §0.1 (latency on modelled components).
-
-### D3. Scheduler plan-generation cost — start at 0, revisit
-
-`PE_SCHEDULER`'s tile-plan generation costs 0 sim time today. Keep that for
-now (the dominant lever is CPU issue cost, D1); expose it later via the
-existing `overhead_ns` node attr if a scheduler-side cost proves material.
-Marked as a review item, not decided here.
-
-### D4. Configurable values; goldens regenerate
-
-The cost table is configurable (per-topology / node attrs, default to the
-D1 estimates). Turning issue cost non-zero changes **every** bench's
-latency, so golden latencies are **regenerated** once — a model-fidelity
-improvement, same posture as ADR-0062's global lazy-load regression.
-
-## Alternatives
-
-### A1. Keep a single uniform `dispatch_cycles > 0`
-
-Rejected: op construct costs differ by an order of magnitude (`tl.load` vs
-`tl.composite`); a uniform value either over-charges primitives or
-under-charges composites, and in both cases misrepresents the
-composite-vs-primitive trade-off the hybrid depends on.
-
-### A2. Charge the issue cost 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 it on the scheduler would not model CPU back-pressure.
-
-### A3. Model DMA program/setup time as a separate fixed per-descriptor cost
-
-Real HW pays a DMA descriptor program cost distinct from transfer time.
-Deferred: initially fold the descriptor-program cost into the **issuing
-op's** D1 cost (a `tl.load` / composite that triggers DMA). Split it out to
-a PE_DMA fixed setup only if calibration shows it matters. Review item.
-
-## Consequences
-
-### Positive
-- The hybrid's CPU-offload / saturation win (ADR-0060 §1) becomes
-  **measurable**, not just structural.
-- Composite-vs-primitive and tiling-granularity trade-offs are visible in
-  latency, enabling the eval narrative.
-- More faithful to hardware (issue bandwidth is a real bottleneck).
-
-### Negative
-- **All** bench goldens shift → one-time regeneration (D4); CI golden
-  fixtures update.
-- Per-op-type values need calibration; the ~40 ns composite figure is an
-  estimate, primitives are unspecified — results are only as good as the
-  numbers (do not over-claim absolute latencies until calibrated).
-- Adds a cost-table lookup on the issue path (negligible runtime).
-
-## Open review items (decided autonomously; revise on review)
-
-1. **Calibration source for the per-op-type values.** Composite ≈ 40 ns is
-   a working estimate; primitive issue costs are placeholders. *Recommend:*
-   pick values from a documented assumption (instruction-issue + queue
-   push) and treat absolute latency as provisional until a real reference
-   exists; keep the **ratios** defensible.
-2. **Scheduler plan-gen cost (D3).** *Recommend:* keep 0 initially; expose
-   via `overhead_ns` if a workload shows scheduler-bound behaviour.
-3. **DMA program time (A3).** *Recommend:* fold into the issuing op's cost
-   first; split to PE_DMA setup only if needed.
-4. **Where the table lives.** *Recommend:* a small central cost module
-   keyed by command type, overridable per topology — not scattered node
-   attrs — so the values are reviewable in one place.
-5. **Regression rollout.** *Recommend:* land lazy-load (ADR-0062) and this
-   ADR's cost model in one golden-regeneration pass to avoid two churns.
-6. **Interaction with the legacy replay path** (`pe_cpu.py:_execute_legacy`)
-   — ensure both the greenlet and replay paths read the same cost table so
-   results match. Verify.
-
-## Test Requirements
-
-1. **Composite charges once; primitives charge per-op.** A kernel issuing
-   one `tl.composite` over `N_tiles` charges one composite issue cost on
-   PE_CPU; the equivalent `N_tiles × ops` primitive kernel charges the
-   per-op cost `N_tiles × ops` times. Assert the PE_CPU busy time differs
-   accordingly.
-2. **DMA latency unchanged (Q2).** For a fixed transfer, PE_DMA `drain_ns`
-   is identical to the pre-ADR value — the issue cost is additive on
-   PE_CPU, not folded into DMA (no double-count).
-3. **Saturation is observable.** With non-zero per-op issue cost, a
-   many-tile primitive sweep shows GEMM-engine idle (CPU-bound issue)
-   whereas the composite sweep keeps it busy — the ADR-0060 §1 lever.
-4. **Determinism:** identical inputs → identical op_log + latency
-   (SPEC §0.1).
-5. **Path parity:** greenlet and legacy-replay paths produce identical
-   issue-cost accounting for the same kernel.

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

@@ -1,39 +1,42 @@
 # Detailed Design Document — AHBM GQA Fused Attention (ADR-0060)
 
 **Status:** Draft (companion to ADR-0060, Proposed)
-**Scope:** implementation-ready design for an efficient GQA FlashAttention
-kernel on AHBM in the kernbench simulator.
-**Audience:** reviewer resuming with *"GQA 검토 시작하자"*. Read ADR-0060
-first (decision record), then this (the how), then the *Open Decisions*
-section (§10) for the choices made autonomously that need your sign-off.
+**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.
 
-> **How this document was produced.** It is the result of several
-> design↔evaluation iterations against the real kernbench source
-> (`src/kernbench/triton_emu/tl_context.py`, `sim_engine/*`,
-> `components/builtin/pe_*`, the `_attention_mesh_*` kernels, and the GQA
-> tests). Every claim about what exists is grounded in a file path; every
-> claim about what is new is tied to a supporting ADR (0061/0062/0063).
+> **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:** an *efficient* GQA fused attention kernel that runs on AHBM in
-kernbench, across the four operating points decode/prefill ×
-single-/multi-user, at realistic (not just validation) scale, in both
+**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: the three levers actually fire and are observable in op_log:
-   GQA K/V-load amortisation, KV prefetch overlap, `⌈log₂ N⌉` reduction.
+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-CUBE query-head split;
-cycle-accurate microarchitecture (SPEC §5).
+`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).
 
 ---
 
@@ -41,31 +44,33 @@ cycle-accurate microarchitecture (SPEC §5).
 
 ```
 runtime_api (host)
-  RuntimeContext.zeros/empty/launch        context.py  (tensor deploy, DPPolicy, kernel launch)
+  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   ← ADR-0061 adds _execute_broadcast
-  MemoryStore (nbytes check)               memory_store.py    ← the GQA blocker lives here
+  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,FETCH_STORE,TCM,IPCQ}
-  greenlet kernel ↔ SimPy                  kernel_runner.py, pe_cpu.py
+  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/store/dot/softmax/max/sum/exp/maximum/send/recv/trans/…
-  ← ADR-0061 adds broadcast/repeat
-  ← ADR-0062 adds load_async (+ wait arm)
-  ← ADR-0063 adds scratch_scope
+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)
 ```
 
-The attention kernels are **bench-side Python** (`src/kernbench/benches/`)
-invoked via `ctx.launch(name, kernel_fn, *args)`; they run as greenlet
-kernels that emit `tl.*` ops. No component (device-side) code changes are
-needed for the algorithm — only the three `tl`/sim_engine primitives.
+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`.
 
 ---
 
@@ -75,424 +80,241 @@ needed for the algorithm — only the three `tl`/sim_engine primitives.
 
 | File | Change | ADR |
 |---|---|---|
-| `src/kernbench/triton_emu/tl_context.py` | add `broadcast`/`repeat`; `load_async` + `wait` arm; `scratch_scope` ctx-mgr | 0061/0062/0063 |
-| `src/kernbench/common/pe_commands.py` | add `BroadcastCmd`; `DmaReadCmd.blocking` flag (or `DmaReadAsyncCmd`) | 0061/0062 |
-| `src/kernbench/sim_engine/data_executor.py` | add `_execute_broadcast`; handle async dma_read resolution | 0061/0062 |
-| `src/kernbench/components/builtin/pe_dma.py` | non-blocking read path returns a future-resolving event | 0062 |
-| `src/kernbench/sim_engine/event_log.py` / op_log | recognise `broadcast` op_name (math) | 0061 |
+| `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) |
 
-### 3.2 New kernel + bench files
+No new GEMM command kind: the two GEMMs use the **existing** `CompositeCmd`.
+
+### 3.2 Kernel + bench files
 
 | File | Role |
 |---|---|
-| `src/kernbench/benches/_gqa_attention.py` | the unified GQA kernel (4 cases via flags), with a shared `_flash_tile_sweep` helper and `_tree_reduce` helper |
-| `src/kernbench/benches/_gqa_helpers.py` | `valid_len`, `tile_*` predicates, mask builders, tree topology (children/parent dirs), GQA-batched Q load |
-| extend `src/kernbench/benches/milestone_gqa_llama70b.py` | add headline-scale panels with real `G`, prefetch, tiling |
+| `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_correctness.py` | numpy reference parity, all 4 cases, real GQA |
-| `tests/attention/test_gqa_levers.py` | GQA load amortisation, tree step count, causal skip, prefetch overlap |
-| `tests/test_tl_broadcast.py` | ADR-0061 unit |
-| `tests/test_tl_load_async.py` | ADR-0062 unit |
+| `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 (DPPolicy)
+### 4.1 Tensor placement — CUBE Group, 2-level SP, contiguous KV
 
-Using `DPPolicy(cube, pe, num_cubes, num_pes)` (`policy/placement/dp.py`):
+A `CUBE Group` is `C` CUBEs in one SIP owning one KV head (ADR-0060 §0).
+Placement via `DPPolicy(..., cube_start=...)`:
 
-| Tensor | single-user (PE ring in 1 cube) | multi-user (cube ring) |
+| Tensor | Decode+SP | Prefill+SP (`C=G`) |
 |---|---|---|
-| `Q` (group, replicated) | `cube=replicate, pe=replicate, num_pes=N` | `cube=replicate, pe=replicate, num_cubes=N` |
-| `K`, `V` (sequence-sharded) | `cube=replicate, pe=row_wise, num_pes=N` | `cube=row_wise, pe=replicate, num_cubes=N` |
-| `O` (group, replicated) | as `Q` | as `Q` |
+| `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 |
 
-This matches the baseline `_PANEL_DISPATCH`
-(`milestone_gqa_llama70b.py:64-81`). `row_wise` over the sequence (M) axis
-realises the round-robin shard (§2.1 of ADR-0060); the contiguous per-PE
-buffer is the row-wise slice.
+**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.
 
-`rank_axis` selects which `tl.program_id` carries the ring rank: `0` for
-single-user (PE id in cube), `1` for multi-user (cube level; only PE 0 of
-each cube participates) — kept from the baseline so the SFR installs
-(`configure_sfr_intracube_pe_ring` / `configure_sfr_intercube_multisip`)
-apply unchanged.
+**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.
 
-### 4.2 Launch signature
+`--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
-ctx.launch(
-    f"{panel}_gqa", gqa_attention_kernel,
-    q, k, v, o,                       # Tensors (TensorArg)
-    S_q, S_kv_per_rank, H_q, H_kv, D_HEAD, n_ranks,   # scalars
-    rank_axis, mode, counter, start_pe,               # scalars (mode: 0=decode,1=prefill)
-    softmax_scale,                                     # float scalar
-    _auto_dim_remap=False,            # keep d_head from colliding (baseline note)
-)
+# 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)
 ```
 
-`H_q > H_kv` is the change from the baseline's `H_q=H_kv=1`; it requires
-ADR-0061 to survive Phase 2.
+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-batched Q load — fold `G` into the matmul M dim (no broadcast)
+### 4.3 GQA reuse — fold `G` into the matmul M dim (decode)
 
-For a KV head `kv`, its group is query heads `[kv·G, (kv+1)·G)`. Load
-those `G` rows and **reshape the group into the matmul M dimension** —
-this is the key to GQA reuse and it needs only a byte-conserving reshape,
-**not** a broadcast:
-
-```python
-def load_q_group(q_ptr, kv, G, T_q, d, tl):
-    # Q laid out [H_q, T_q, d]; the group is contiguous in the head axis.
-    qg = tl.load(q_ptr + kv*G*T_q*d, shape=(G, T_q, d), dtype="f16")
-    return _view(qg, (G*T_q, d))            # byte-conserving reshape → M = G*T_q
-```
-
-Then per KV tile (`Kj`, `Vj` are `[TILE, d]`, one KV head, loaded once):
-
-```python
-Sj  = tl.dot(q_g, tl.trans(Kj)) * scale     # [G*T_q, d]·[d, TILE] -> [G*T_q, TILE]
-Oj  = tl.dot(P,   Vj)                        # [G*T_q, TILE]·[TILE, d] -> [G*T_q, d]
-```
-
-K/V are the **shared operands** (the matmul's `b`), loaded once and reused
-across all `G·T_q` M-rows automatically. The emitted `GemmCmd` carries
-`m = G·T_q`, so Phase-1 timing counts all `G` rows' work correctly.
-
-> **Why not a leading batch axis** (`[G,T_q,d]·[G,d,TILE]`)? `tl.dot`
-> derives `m,k,n` from the **last two dims only** (`tl_context.py:220`)
-> and `GemmCmd` carries no batch dim — so a leading `G` axis is computed
-> correctly by `np.matmul` in Phase 2 (it broadcasts) but **under-counted
-> by `G×` in Phase-1 timing**. Folding into M fixes both. This supersedes
-> the earlier batched-vs-loop question.
-
-> **Transpose caveat.** `tl.trans` is metadata-only and `MemoryStore.read`
-> *reshapes* (not transposes) (`memory_store.py:73`,
-> `data_executor.py:153-168` does `np.matmul` on the reshaped operands).
-> Fine for timing/structure and for the baseline's zero data; for
-> non-trivial numeric parity, store K pre-transposed as `[d, S_pe]` or add
-> a real `tl.transpose` (§10.10).
+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 (concrete, real `tl`)
+## 5. Kernel design
 
-### 5.1 Shared tile sweep helper
+**The full pseudocode for both kernels (and the 3 decode CPU-pipelining
+variants) lives in ADR-0060 TL;DR + §5.** Implementation notes only here.
 
-```python
-def flash_tile_sweep(q_g, k_base, v_base, n_tiles, TILE, d, scale,
-                     q_block, causal, tl, PREFETCH=2):
-    """One PE's FlashAttention sweep over its local KV tiles.
-    Returns running (m, l, O) in the persistent arena."""
-    m, l, O = _init_running(q_g.shape, d, tl)           # persistent arena
-    f_k, f_v = {}, {}
-    for j in range(min(PREFETCH, n_tiles)):             # prime (ADR-0062)
-        if causal and tile_all_future(j, q_block): continue
-        f_k[j] = tl.load_async(k_base + j*TILE*d, (TILE, d))
-        f_v[j] = tl.load_async(v_base + j*TILE*d, (TILE, d))
-    for j in range(n_tiles):
-        if causal and tile_all_future(j, q_block): continue
-        with tl.scratch_scope():                        # recycle temporaries (ADR-0063)
-            Kj = tl.wait(f_k[j])
-            Sj = tl.dot(q_g, tl.trans(Kj)) * scale
-            mask = make_mask_tile(j, q_block) if (causal and tile_partial(j, q_block)) else None
-            if mask is not None: Sj = Sj + mask
-            m_j   = tl.max(Sj, axis=-1)
-            m_new = tl.maximum(m, m_j)
-            P     = tl.exp(Sj - m_new)
-            corr  = tl.exp(m - m_new)
-            l     = l * corr + tl.sum(P, axis=-1)
-            Vj    = tl.wait(f_v[j])
-            O     = O * corr + tl.dot(P, Vj)
-            m     = m_new
-            nxt = j + PREFETCH                           # refill pipeline
-            if nxt < n_tiles and not (causal and tile_all_future(nxt, q_block)):
-                f_k[nxt] = tl.load_async(k_base + nxt*TILE*d, (TILE, d))
-                f_v[nxt] = tl.load_async(v_base + nxt*TILE*d, (TILE, d))
-    return m, l, O
-```
+### 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.
 
-> Note: `m, l, O` are reassigned inside the `with` but their **storage**
-> must be the persistent arena, not the recycled scope. Concretely the
-> merge ops write to pre-allocated persistent handles (double-buffered);
-> ADR-0063 §D3 describes the persistent-vs-scoped arena split. This is the
-> one subtlety the implementer must get right — see §10.5.
+### 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.2 Tree reduction helper
-
-```python
-def tree_reduce_and_store(m, l, O, pe_id, N, o_ptr, tl):
-    for cdir in tree_children_dirs(pe_id, N):           # static tree (§4 ADR-0060)
-        m_c = tl.recv(cdir, m.shape); l_c = tl.recv(cdir, l.shape)
-        O_c = tl.recv(cdir, O.shape)
-        mn  = tl.maximum(m, m_c); sa = tl.exp(m - mn); sb = tl.exp(m_c - mn)
-        l   = l*sa + l_c*sb; O = O*sa + O_c*sb; m = mn
-    if is_root(pe_id, N):
-        tl.store(o_ptr, O / l)
-    else:
-        pdir = parent_dir(pe_id, N)
-        tl.send(pdir, m); tl.send(pdir, l); tl.send(pdir, O)
-```
-
-`tree_children_dirs` / `parent_dir` map the binary tree (level pairs in
-ADR-0060 §4) onto physical `N/S/E/W` neighbours — this mapping is the
-tuning item ADR-0060 §9.2 and an Open Decision (§10.2).
-
-### 5.3 Top-level kernel (all four cases)
-
-```python
-def gqa_attention_kernel(q_ptr,k_ptr,v_ptr,o_ptr,
-                         S_q,S_kv_per_rank,H_q,H_kv,d,n_ranks,
-                         rank_axis,mode,counter,start_pe,scale,*,tl):
-    if rank_axis != 0 and tl.program_id(axis=0) != 0:   # multi-user: only PE0 per cube
-        return
-    pe_id = tl.program_id(axis=rank_axis)
-    G = H_q // H_kv
-    causal = (mode == 1)                                # prefill
-    q_block = (counter, S_q) if causal else (counter, 1)
-    # decode small-S fast path = one tile; otherwise tiled (TILE from config)
-    for kv in range(H_kv):                              # one KV head per iter (or PE-mapped)
-        q_g  = load_q_group(q_ptr, kv, G, S_q, d, tl)
-        my_len  = valid_len(counter, start_pe, pe_id, n_ranks) if mode==0 else S_kv_per_rank
-        n_tiles = ceil(my_len / TILE)
-        m,l,O = flash_tile_sweep(q_g, k_base(kv), v_base(kv), n_tiles, TILE, d,
-                                 scale, q_block, causal, tl)
-        if n_ranks == 1:
-            tl.store(o_ptr_for(kv), O / l)
-        elif mode == 1 and is_ring(...):                # prefill ring already folded across steps
-            ring_fold_then_store(...)                   # §5.4
-        else:
-            tree_reduce_and_store(m,l,O, pe_id, n_ranks, o_ptr_for(kv), tl)
-```
-
-This collapses the baseline's two separate kernels (`_attention_mesh_kv`,
-`_attention_mesh_mlo`) into one parameterised kernel; the `mode`/`n_ranks`
-flags select the case. (Keeping them separate is also fine — Open
-Decision §10.6.)
-
-### 5.4 Ring (prefill SP) specifics
-
-The Ring path keeps the baseline's structure (`_attention_mesh_kv.py`)
-with three upgrades:
-- `recv_async` for the next step's K/V (overlap) — `tl_context.py:543`.
-- causal step-skip: `if step_all_future(step, q_block): continue` before
-  folding (skips ≈half the steps).
-- GQA: `q_g` is the `G`-row group; `tl.broadcast` for the KV head.
-
-The fold across ring steps is the same online-softmax merge as §5.2's
-inner update — running `(m,l,O)` carried across steps as Python handles
-(exactly as the baseline does today).
+### 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. The three new primitives — integration points
+## 6. Supporting primitives — integration points
 
-### 6.1 ADR-0061 `tl.broadcast` (OPTIONAL — not on the GQA critical path)
-GQA reuse is achieved by folding `G` into M (§4.3), so broadcast is **not**
-required to unblock `h_q > h_kv`. It remains a clean primitive for additive
-causal-mask construction across the `G·T_q` rows and for general kernels.
-- `tl_context.py`: new `broadcast(x, new_shape)` → emits `BroadcastCmd`,
-  output handle from `_make_compute_out` (so `nbytes` is correct).
-- `data_executor.py`: `_execute_broadcast` = `np.broadcast_to(src,
-  shape).copy()` → `store.write`.
-- Latency: `pe_math._compute_ns(prod(new_shape))`; logs `math/broadcast`,
-  **no** `dma_read`.
-- Lowest priority; deliver after the efficiency/scale features land.
+### 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-0062 `tl.load_async`
-- `tl_context.py`: `load_async(ptr, shape, dtype) → LoadFuture`; extend
-  `wait` to resolve `LoadFuture`.
-- `pe_dma.py`: non-blocking read path; transfer occupies the read channel
-  in parallel; completion is a scheduled event.
-- op_log unchanged (`memory/dma_read`) ⇒ `dma_read_count` metric 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-0063 `tl.scratch_scope`
-- `tl_context.py`: context manager saving/restoring `_scratch_cursor`.
-- Persistent arena: `m,l,O` (+ double-buffer) and live `LoadFuture`
-  buffers allocated outside the scope.
-- Removes the `S=16` cap (`test_milestone_gqa_llama70b.py:123-148`).
+### 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. Ordered so each builds
-on a green predecessor; every phase keeps the existing baseline green.
+Each phase is an independent Phase-1→Phase-2 cycle; each keeps the baseline
+green.
 
 | Phase | Deliverable | Gate |
 |---|---|---|
-| **P1** | Real GQA by restructuring (per KV head, `G`→M fold); one-shot, no tiling. **No new primitive.** | data-mode run completes for `h_q=G·h_kv`; K/V `dma_read_count` independent of `G`; `gemm` m=`G·T_q` |
-| **P2** | Tree reduction replacing fan-out (decode SP) | `⌈log₂ N⌉` reduction rounds; structure preserved |
-| **P3** | ADR-0063 `tl.scratch_scope` + tiled sweep | long-`S` decode completes (today caps at S=16) |
-| **P4** | ADR-0062 `tl.load_async` + prefetch in sweep | latency < serial Σ(load+compute) |
-| **P5** | Causal masking + tile/step skip (prefill, ring) | GEMM count = triangular tile count |
-| **P6** | Headline-scale milestone panels (real `G`, tiling, prefetch) + figures | sweep.json headline rows; figures render |
-| **P7** *(opt)* | ADR-0061 `tl.broadcast` (mask convenience); optional `tl.transpose` for numeric parity | broadcast unit test; asymmetric numeric parity |
+| **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 alone delivers "GQA actually runs" — and it needs no new feature**,
-only the §4.3/§5 restructuring (the iteration-2 finding). P3 is the most
-important *new* feature (removes the S=16 scale ceiling). P4 is the
-efficiency lever. P2/P5 are algorithmic. P7 is optional polish.
+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)
 
-Grounded in SPEC R2/R5, ADR-0023/0025 (IPCQ), ADR-0046 (`tl`),
-ADR-0054 (eval bench). Mirrors ADR-0060 §11.
+Mirrors ADR-0060 §11; grounded in SPEC R2/R5, ADR-0023/0025, ADR-0046, ADR-0054.
 
-**Runs + numeric (Phase 2, `enable_data=True`):**
-- `test_gqa_correctness.py`: for `(G∈{1,8}, T_q∈{1,16}, S, N∈{1,8})` the
-  GQA kernel **completes** (no byte-conservation error) — the `G=8` cases
-  that the baseline head-packing cannot express. This is gated on the
-  §4.3 restructuring, **not** on a new primitive.
+**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 where reshape-as-transpose is exact; full
-  asymmetric parity is gated on an optional `tl.transpose` (§10.10).
+  symmetric/identity inputs (reshape-as-transpose exact); full asymmetric
+  parity gated on a real `tl.transpose` (deferred).
 
-**Levers (op_log assertions):**
-- GQA amortisation: K/V `dma_read_count` constant as `G` grows 1→8;
-  `gemm` work grows with `G`.
-- Tree: decode-SP `ipcq_send`/`recv` counts = tree edges, reduction
-  rounds = `⌈log₂ N⌉` (not `N−1`).
-- Causal skip: prefill `gemm_count` = lower-triangular tile count.
-- Prefetch overlap: tiled-sweep latency strictly < `Σ(load+compute)`.
-- Scratch: peak `_scratch_cursor` bounded by one tile, independent of
-  `n_tiles`.
+**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 inputs → identical op_log + latency (SPEC §0.1).
-- Every routed request latency > 0 (SPEC §0.1).
+**Invariants:** determinism (identical op_log + latency); every routed
+request latency > 0 (SPEC §0.1).
 
 ---
 
-## 9. Performance model (what we expect to see)
-
-Rough per-PE decode cost (KV-load-bound), with prefetch overlap:
+## 9. Performance model (expected)
 
+Per-rank decode (KV-load-bound), composite-streamed + overlapped:
 ```
-t_decode_pe ≈ max( DMA(K+V over S_pe tiles),  compute(QKᵀ + PV over S_pe, ×G amortised) )
-            + t_tree_reduce(⌈log₂ N⌉ hops, payload O[G,d])
+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 and their expected signature in op_log / latency:
-- **GQA reuse** → KV bytes moved are `H_kv·S·d` not `H_q·S·d` (8× less),
-  with compute still `H_q`-scaled. Decode being KV-load-bound, this is the
-  dominant decode win.
-- **Prefetch** → DMA hidden behind compute ⇒ `t ≈ max(...)` not sum.
-- **Tree** → reduction `⌈log₂ N⌉` (3 for N=8) vs baseline `N−1` (7).
-- **Causal skip** (prefill) → ≈½ the tile/step work.
-
-These are the quantities the milestone bench's op_log summary + a new
-latency column should report (Open Decision §10.3 on adding latency to the
-sweep JSON).
+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 Decisions / Assumptions made autonomously (REVIEW THESE)
+## 10. Open items (status)
 
-Choices I made to keep progressing without your input. Each is reversible;
-none is yet implemented in production code (docs only). Marked
-**[recommend]** = my default if you don't object.
+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:
 
-### 10.1 Greenlet `tl` model over composites **[recommend: greenlet]**
-The biggest design call. ADR-0060 §1/§8 argues the greenlet path is
-equally efficient in kernbench's per-op latency model and far smaller to
-build than a flash-composite. **Risk:** if you specifically want the
-*composite* abstraction exercised for the eval narrative (e.g. to study
-scheduler-managed tiling), we'd instead invest in a flash-composite
-command kind. I judged that out of proportion to the goal. *Confirm the
-greenlet direction.*
+| 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) |
 
-### 10.2 Tree reduction topology / neighbour mapping **[recommend: binary tree on N/S/E/W]**
-ADR-0060 §4 fixes a binary tree (N=8: root=7). The mapping of tree edges
-to physical mesh directions depends on the SFR install's neighbour
-wiring. I assumed `configure_sfr_intracube_pe_ring` gives a ring on which
-the tree pairs are 1-hop neighbours. **Needs verification** against the
-actual installed neighbour table; if pairs aren't 1-hop, either relabel
-PEs or accept multi-hop sends. *Confirm or hand me the intended mapping.*
-
-### 10.3 Keep the baseline all-to-all as an option? **[recommend: replace with tree]**
-The baseline fan-out gives the answer on *every* rank; the tree gives it
-on root only. Attention needs root-only, so tree wins. If any downstream
-consumer needs O replicated (none found), we'd keep fan-out. *Assuming
-root-only is fine.*
-
-### 10.4 GQA matmul shape — RESOLVED in iteration 2 **[fold G into M]**
-Earlier this asked "batched leading axis vs per-`g` loop." Resolved: do
-**neither** — fold `G·T_q` into the matmul M dimension (§4.3). A leading
-`G` axis would be computed by `np.matmul` in data mode but **under-counted
-`G×` in Phase-1 timing** (`GemmCmd` carries no batch dim,
-`tl_context.py:220`). Folding into M gives correct timing *and* correct
-data with a single byte-conserving reshape and one shared K/V operand. No
-new primitive, no per-`g` loop.
-
-### 10.5 Persistent-vs-scoped arena split **[recommend: two explicit arenas]**
-ADR-0063 requires `m,l,O` to live outside the recycled scope. The cleanest
-implementation reserves a small fixed persistent region at kernel start
-and double-buffers the merge. I assumed this is acceptable extra
-bookkeeping inside the shared helper. *Alternative:* make `scratch_scope`
-take an explicit "keep-alive" handle list. I picked the two-arena model
-for clarity.
-
-### 10.6 One unified kernel vs two (decode/prefill) **[recommend: unify]**
-§5.3 unifies into one parameterised kernel; the baseline has two. Unifying
-reduces duplication but makes the kernel branch-heavy. *Alternative:* keep
-`_gqa_decode` and `_gqa_prefill` separate, sharing helpers. Low stakes;
-either is fine.
-
-### 10.7 Ghost ADRs 0055–0059 **[recommend: backfill as Draft, separate task]**
-The baseline cites ADR-0055/0056/0057/0058/0059 which **do not exist**
-(`grep` finds refs in code/tests, no files; no git history of them). This
-is pre-existing documentation debt, not created by this work. I did **not**
-backfill them (out of scope, and reverse-engineering intent is risky).
-*Recommendation:* a separate `Draft` ADR-backfill task to document the
-existing mesh kernels + milestone bench + SFR install, so ADR-0060 has a
-real predecessor to "supersede". *Confirm you want this as follow-up.*
-
-### 10.8 Numbering **[recommend: as chosen]**
-GQA ADR = **0060**; supporting = **0061/0062/0063**. I avoided 0055–0059
-(claimed by ghost refs) to prevent collisions. The file was renamed from
-`ADR-XXXX-…` to `ADR-0060-algo-…` (title was erroneously "ADR-001").
-
-### 10.9 bf16 modelled as f16 **[assumption, pre-existing]**
-`memory_store.py:16` maps `bf16→float16` (numpy has no bf16). Numeric
-parity tests must use tolerances that accommodate this, and we should not
-claim bf16 *accuracy* results — only timing/structure. Pre-existing
-simulator limitation; flagged so results aren't over-interpreted.
-
-### 10.10 `tl.trans` is reshape-not-transpose **[decision: structural-first verify]**
-Discovered in iteration 2: `tl.trans` only swaps shape metadata
-(`tl_context.py:390`) and `MemoryStore.read` *reshapes* the bytes
-(`memory_store.py:73`); `data_executor._execute_gemm` then `np.matmul`s
-the reshaped operands (`data_executor.py:153-168`). So for **non-trivial
-numeric data**, `Q·Kᵀ` via `tl.trans(K)` is mathematically a reshape, not
-a transpose — wrong numbers (the baseline never notices: it runs on
-all-zeros and only asserts op_log structure). **Decision:** treat the
-simulator as a *performance* model (SPEC §0) — verify structure + timing +
-determinism first; treat numeric parity as a bounded secondary check
-(symmetric/identity inputs, or store K pre-transposed). A real
-data-materialising `tl.transpose` is a *candidate* further primitive if
-true numeric validation is ever required — **recommend deferring it**; it
-is not needed for the performance-modeling goal. *Confirm you're content
-with structural-first verification.*
-
-### 10.11 GQA reuse = fold `G` into M, NOT a broadcast op **[finding, iteration 2]**
-The headline mechanism finding. The deferred-GQA limitation
-(`test_milestone_gqa_llama70b.py:137-142`) was blamed on a missing
-broadcast / the MemoryStore nbytes check. Iteration 2 shows the real
-cause is the baseline's head-packing reshape, and the fix is kernel
-restructuring (§4.3), which needs **no new primitive**. This demoted
-ADR-0061 from "the blocker" to "optional convenience" and reordered the
-phase plan (§7). Recorded explicitly because it changes the project's
-critical path: *GQA itself is not gated on any new feature; scale
-(ADR-0063) and efficiency (ADR-0062) are the features that matter.*
+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.
 
 ---
 
@@ -500,33 +322,34 @@ critical path: *GQA itself is not gated on any new feature; scale
 
 | Risk | Likelihood | Mitigation |
 |---|---|---|
-| Numeric parity blocked by trans-as-reshape | med | structural-first verify (§10.10); pre-transpose K or defer `tl.transpose` |
-| Tree pairs not 1-hop on installed mesh | med | verify SFR neighbour table (§10.2); relabel PEs |
-| scratch_scope use-after-scope bug | med | two-arena discipline in shared helper; debug poison |
-| prefetch buffer recycled while in flight | med | live `LoadFuture` buffers in persistent arena (§10.5) |
-| GQA broadcast `.copy()` memory blow-up | low | bounded by TILE in tiled design |
-| Effort underestimated (6 phases) | med | P0–P1 alone deliver the core unblock; later phases independent |
+| 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** — grouped-query attention; `G = H_q/H_kv` query heads share
-  one KV head. Llama3-70B: `G=8`.
-- **FlashAttention / FlashDecoding / Ring Attention** — see ADR-0060
-  lineage note.
-- **IPCQ** — inter-PE queues; `tl.send`/`tl.recv` over `N/S/E/W`
-  (ADR-0023/0025/0032).
-- **Greenlet `tl` model** — kernel interleaves with SimPy; ops emit
-  commands scheduled on PE engines (ADR-0020, ADR-0046).
-- **Persistent vs scoped arena** — ADR-0063 scratch split.
+- **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.
 
-Key source anchors: `tl_context.py` (tl API), `memory_store.py:67-73`
-(nbytes check), `data_executor.py` (Phase 2), `pe_dma.py`/`pe_gemm.py`/
-`pe_math.py` (latency), `_attention_mesh_kv.py` / `_attention_mesh_mlo.py`
-(baseline kernels), `milestone_gqa_llama70b.py` (eval bench),
-`tests/attention/test_milestone_gqa_llama70b.py` (current limits).
+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), **0061** (broadcast), **0062** (async load),
-**0063** (scratch scope); related accepted **0014/0020/0023/0025/0042/
-0045/0046/0052/0054**.
+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
-  for WRITE).
+- Within a channel, requests serialize **at the issue path** (READ ∥ READ
+  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
 
@@ -119,9 +121,9 @@ 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
-                pe=ls.local_pe,                   # local within cube (explicit name)
+                sip=target_sip,                              # from caller (current_device)
+                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-0062-prog-tl-async-load.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Proposed
+Accepted
 
 > Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention). Decode and
 > long-context attention are **KV-load-bound**; load/compute overlap is a

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

@@ -2,7 +2,7 @@
 
 ## Status
 
-Proposed
+Accepted
 
 > Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention). Long-context
 > attention sweeps many K/V tiles; each tile's intermediates
@@ -102,6 +102,68 @@ kernel keeps two arenas:
 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

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/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,1024,2048,2717.2783333333473
-hbm,torus_2d,6,8192,16384,7315.184999999989
-hbm,torus_2d,6,32768,65536,23081.26500000037
-sram,torus_2d,6,128,256,2060.040000000012
-sram,torus_2d,6,1024,2048,2909.2783333333473
-sram,torus_2d,6,8192,16384,9523.184999999869
-sram,torus_2d,6,32768,65536,32201.265000000385
-tcm,torus_2d,6,128,256,1964.040000000012
-tcm,torus_2d,6,1024,2048,2477.2783333333473
-tcm,torus_2d,6,8192,16384,6403.185000000109
-tcm,torus_2d,6,32768,65536,19865.265000000378
+hbm,torus_2d,6,128,256,2345.040000000012
+hbm,torus_2d,6,1024,2048,2942.2783333333473
+hbm,torus_2d,6,8192,16384,7540.184999999989
+hbm,torus_2d,6,32768,65536,23306.26500000037
+sram,torus_2d,6,128,256,2285.040000000012
+sram,torus_2d,6,1024,2048,3134.2783333333527
+sram,torus_2d,6,8192,16384,9748.184999999869
+sram,torus_2d,6,32768,65536,32426.265000000385
+tcm,torus_2d,6,128,256,2189.040000000012
+tcm,torus_2d,6,1024,2048,2702.2783333333473
+tcm,torus_2d,6,8192,16384,6628.18500000005
+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,32,64,1024,2747.7400000000152
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,2855.990000000018
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,3072.490000000019
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,3337.1133333333582
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,3708.0333333333692
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,4449.873333333393
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,5933.020000000124
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,8900.379999999863
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,14835.099999999224
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,26704.540000000765
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,38573.97999999701
-lrab_hierarchical_allreduce,ring_1d,6,8,16,256,2365.255833333347
-lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,2436.9433333333473
-lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,2532.526666666683
-lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,2723.693333333349
-lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,3048.635000000021
-lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,3393.4016666666957
-lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,4082.401666666714
-lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,5458.80166666677
-lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,8216.934999999943
-lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,13733.201666665835
-lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,24765.73500000064
-lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,35798.268333331536
-lrab_hierarchical_allreduce,torus_2d,6,8,16,256,1700.6025000000095
-lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,1753.2900000000102
-lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,1823.540000000012
-lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,1964.040000000012
-lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,2196.8183333333463
-lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,2477.2783333333473
-lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,3038.1983333333583
-lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,4159.5050000000665
-lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,6403.185000000109
-lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,10890.5449999995
-lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,19865.265000000378
-lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,28839.98500000059
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8,16,256,2918.5525000000157
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,2999.740000000016
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,3107.990000000019
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,3324.4900000000207
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,3589.1133333333582
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,3960.0333333333692
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,4701.873333333393
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,6185.020000000124
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,9152.379999999861
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,15087.099999999224
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,26956.540000000765
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,38825.97999999701
+lrab_hierarchical_allreduce,ring_1d,6,8,16,256,2628.2558333333477
+lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,2699.943333333348
+lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,2795.526666666683
+lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,2986.693333333351
+lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,3311.635000000021
+lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,3656.4016666666957
+lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,4345.401666666714
+lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,5721.801666666768
+lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,8479.934999999887
+lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,13996.201666665835
+lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,25028.73500000064
+lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,36061.26833333154
+lrab_hierarchical_allreduce,torus_2d,6,8,16,256,1925.6025000000104
+lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,1978.290000000011
+lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,2048.540000000012
+lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,2189.040000000012
+lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,2421.8183333333463
+lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,2702.2783333333473
+lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,3263.1983333333583
+lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,4384.5050000000665
+lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,6628.18500000005
+lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,11115.5449999995
+lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,20090.265000000378
+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,raw,33.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,raw,36.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,raw,37.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,raw,38.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,raw,40.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,raw,42.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,raw,50.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,raw,66.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,raw,98.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,raw,114.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,raw,47.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,raw,51.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,raw,52.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,raw,53.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,raw,55.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,raw,57.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,raw,65.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,raw,81.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,raw,113.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,raw,129.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,raw,89.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,raw,95.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,raw,96.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,raw,97.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,raw,99.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,raw,102.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,raw,114.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,raw,138.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,raw,186.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,raw,210.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,raw,111.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,raw,119.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,raw,120.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,raw,121.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,raw,123.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,raw,126.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,raw,138.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,raw,162.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,raw,210.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,raw,234.54000000000087
+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,51.57999999999811
+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,54.07999999999811
+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,55.07999999999811
+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,56.07999999999811
+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,58.07999999999811
+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,60.07999999999811
+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,68.07999999999811
+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,84.08000000000175
+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,116.08000000000175
+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,132.08000000000175
+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,65.18999999999505
+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,69.18999999999505
+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,70.18999999999505
+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,71.18999999999505
+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,73.18999999999505
+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,75.18999999999505
+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,83.18999999999505
+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,99.19000000000233
+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,131.19000000000233
+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,147.19000000000233
+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,107.28999999999724
+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,113.53999999999724
+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,114.53999999999724
+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,115.53999999999724
+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,117.53999999999724
+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,120.53999999999724
+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,132.53999999999724
+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,156.54000000000087
+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,204.54000000000087
+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,228.54000000000087
+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,129.28999999999724
+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,137.53999999999724
+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,138.53999999999724
+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,139.53999999999724
+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,141.53999999999724
+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,144.53999999999724
+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,156.53999999999724
+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,180.54000000000087
+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,228.54000000000087
+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,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/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
+      }
+    }
+  ]
+}

docs/report/1H-codesign-paper/figures/probe_pe_dma_summary.txt

@@ -0,0 +1,57 @@
+
+=== H2D Write Latency (IO->HBM, data=32768B) ===
+  Case           Target           Hops     Actual    Ovhd  Drain  Wire  Ovhd%  Drain%     Eff.BW      BN.BW    Util%
+  -------------------------------------------------------------------------------------------------------------------
+  h2d-1hop       cube0.pe0            1     289.53    50.0  256.0  0.18  17.3%  88.4%     113.17      128.0    88.4%
+  h2d-2hop       cube4.pe0            2     326.04    78.0  256.0  0.20  23.9%  78.5%     100.50      128.0    78.5%
+  h2d-3hop       cube8.pe0            3     362.56   106.0  256.0  0.20  29.2%  70.6%      90.38      128.0    70.6%
+  h2d-4hop       cube12.pe0            4     399.06   134.0  256.0  0.21  33.6%  64.1%      82.11      128.0    64.1%
+  [v] Monotonic increase: PASS
+
+  BW Saturation (Util% by data size):
+  Case                           4KB    16KB    64KB   256KB     1MB
+  ------------------------------------------------------------------
+  h2d-1hop                     38.9%   71.8%   91.1%   97.6%   99.4%
+  h2d-2hop                     29.0%   62.1%   86.8%   96.3%   99.1%
+  h2d-3hop                     23.2%   54.7%   82.8%   95.1%   98.7%
+  h2d-4hop                     19.3%   48.8%   79.2%   93.8%   98.4%
+
+=== D2H Read Latency (HBM->IO, data=32768B) ===
+  Case           Source           Hops     Actual    Ovhd  Drain  Wire  Ovhd%  Drain%     Eff.BW      BN.BW    Util%
+  -------------------------------------------------------------------------------------------------------------------
+  d2h-1hop       cube0.pe0            1     571.20    23.0  256.0  0.04   4.0%  44.8%      57.37      128.0    44.8%
+  d2h-2hop       cube4.pe0            2     635.72    51.0  256.0  0.05   8.0%  40.3%      51.54      128.0    40.3%
+  d2h-3hop       cube8.pe0            3     700.24    79.0  256.0  0.06  11.3%  36.6%      46.80      128.0    36.6%
+  d2h-4hop       cube12.pe0            4     764.76   107.0  256.0  0.07  14.0%  33.5%      42.85      128.0    33.5%
+  [v] Monotonic increase: PASS
+
+  BW Saturation (Util% by data size):
+  Case                           4KB    16KB    64KB   256KB     1MB
+  ------------------------------------------------------------------
+  d2h-1hop                      9.2%   28.9%   61.9%   86.7%   96.3%
+  d2h-2hop                      7.8%   25.2%   57.4%   84.4%   95.6%
+  d2h-3hop                      6.7%   22.4%   53.5%   82.2%   94.9%
+  d2h-4hop                      5.9%   20.1%   50.2%   80.1%   94.2%
+  [v] D2H >= H2D (reverse data path): PASS
+
+=== PE DMA Latency (pe_dma -> router -> HBM, data=32768B) ===
+  Case                       Target                         Actual    Ovhd  Drain  Wire  Ovhd%  Drain%     Eff.BW      BN.BW    Util%
+  ----------------------------------------------------------------------------------------------------------------------------
+  pe-local-hbm               c0.pe0->c0.slice0              141.00     2.0  128.0  0.00   1.4%  90.8%     232.40      256.0    90.8%
+  pe-same-half-hbm           c0.pe0->c0.slice1              147.87     4.0  128.0  0.03   2.7%  86.6%     221.60      256.0    86.6%
+  pe-cross-half-hbm          c0.pe0->c0.slice4              161.17    10.0  128.0  0.09   6.2%  79.4%     203.31      256.0    79.4%
+  pe-cross-cube-hbm-best     c0.pe0->c1.slice0              330.52    30.0  256.0  0.01   9.1%  77.5%      99.14      128.0    77.5%
+  pe-cross-cube-hbm-worst    c0.pe0->c15.slice0             677.12   180.0  256.0  0.06  26.6%  37.8%      48.39      128.0    37.8%
+  * Local BN: 256.0 GB/s, Remote BN: 256.0 GB/s
+  [v] Cross-cube best < worst: PASS  (330.52ns < 677.12ns)
+
+  BW Saturation (Util% by data size):
+  Case                           4KB    16KB    64KB   256KB     1MB
+  ------------------------------------------------------------------
+  pe-local-hbm                 88.9%   97.0%   99.2%   99.8%  100.0%
+  pe-same-half-hbm             79.9%   94.1%   98.5%   99.6%   99.9%
+  pe-cross-half-hbm            61.3%   86.4%   96.2%   99.0%   99.8%
+  pe-cross-cube-hbm-best       51.6%   81.0%   94.5%   98.6%   99.6%
+  pe-cross-cube-hbm-worst      15.1%   41.6%   74.0%   91.9%   97.8%
+
+============================================================

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

@@ -0,0 +1,94 @@
+<svg xmlns="http://www.w3.org/2000/svg" width="508" height="495" viewBox="70 40 508 495">
+  <title>sip</title>
+  <rect width="648" height="648" fill="#ffffff"/>
+  <line x1="108.0" y1="144.0" x2="252.0" y2="144.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="180.0" y="140.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="108.0" y1="144.0" x2="108.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="108.0" y="200.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="252.0" y1="144.0" x2="396.0" y2="144.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="324.0" y="140.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="252.0" y1="144.0" x2="252.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="252.0" y="200.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="396.0" y1="144.0" x2="540.0" y2="144.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="468.0" y="140.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="396.0" y1="144.0" x2="396.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="396.0" y="200.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="540.0" y1="144.0" x2="540.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="540.0" y="200.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="108.0" y1="264.0" x2="252.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="180.0" y="260.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="108.0" y1="264.0" x2="108.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="108.0" y="320.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="252.0" y1="264.0" x2="396.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="324.0" y="260.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="252.0" y1="264.0" x2="252.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="252.0" y="320.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="396.0" y1="264.0" x2="540.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="468.0" y="260.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="396.0" y1="264.0" x2="396.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="396.0" y="320.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="540.0" y1="264.0" x2="540.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="540.0" y="320.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="108.0" y1="384.0" x2="252.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="180.0" y="380.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="108.0" y1="384.0" x2="108.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="108.0" y="440.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="252.0" y1="384.0" x2="396.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="324.0" y="380.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="252.0" y1="384.0" x2="252.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="252.0" y="440.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="396.0" y1="384.0" x2="540.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="468.0" y="380.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="396.0" y1="384.0" x2="396.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="396.0" y="440.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="540.0" y1="384.0" x2="540.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="540.0" y="440.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="108.0" y1="504.0" x2="252.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="180.0" y="500.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="252.0" y1="504.0" x2="396.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="324.0" y="500.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <line x1="396.0" y1="504.0" x2="540.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="468.0" y="500.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
+  <polyline points="324.0,56.0 108.0,56.0 108.0,144.0" fill="none" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="216.0" y="96.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">2.5mm 512GB/s</text>
+  <polyline points="324.0,56.0 252.0,56.0 252.0,144.0" fill="none" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="288.0" y="96.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">2.5mm 512GB/s</text>
+  <polyline points="324.0,56.0 396.0,56.0 396.0,144.0" fill="none" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="360.0" y="96.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">2.5mm 512GB/s</text>
+  <polyline points="324.0,56.0 540.0,56.0 540.0,144.0" fill="none" stroke="#000000" stroke-width="1" opacity="0.8"/>
+  <text x="432.0" y="96.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">2.5mm 512GB/s</text>
+  <rect x="84.0" y="128.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="108.0" y="148.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (0,0)</text>
+  <rect x="228.0" y="128.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="252.0" y="148.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (1,0)</text>
+  <rect x="372.0" y="128.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="396.0" y="148.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (2,0)</text>
+  <rect x="516.0" y="128.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="540.0" y="148.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (3,0)</text>
+  <rect x="84.0" y="248.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="108.0" y="268.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (0,1)</text>
+  <rect x="228.0" y="248.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="252.0" y="268.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (1,1)</text>
+  <rect x="372.0" y="248.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="396.0" y="268.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (2,1)</text>
+  <rect x="516.0" y="248.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="540.0" y="268.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (3,1)</text>
+  <rect x="84.0" y="368.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="108.0" y="388.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (0,2)</text>
+  <rect x="228.0" y="368.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="252.0" y="388.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (1,2)</text>
+  <rect x="372.0" y="368.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="396.0" y="388.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (2,2)</text>
+  <rect x="516.0" y="368.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="540.0" y="388.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (3,2)</text>
+  <rect x="84.0" y="488.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="108.0" y="508.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (0,3)</text>
+  <rect x="228.0" y="488.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="252.0" y="508.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (1,3)</text>
+  <rect x="372.0" y="488.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="396.0" y="508.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (2,3)</text>
+  <rect x="516.0" y="488.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
+  <text x="540.0" y="508.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (3,3)</text>
+  <rect x="308.0" y="50.0" width="32.0" height="12.0" rx="4" fill="#0ea5e9" stroke="#000000" stroke-width="1"/>
+  <text x="324.0" y="60.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#ffffff">IO io0</text>
+</svg>

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

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

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

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

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

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

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

@@ -0,0 +1,73 @@
+\section{서론}
+\label{sec:intro}
+
+AHBM은 연산 유닛을 HBM 스택 내부에 직접 통합한 메모리 중심 가속기
+아키텍처이다. 각 처리 요소(PE)는 전용 HBM 슬라이스와 짝지어지며,
+TCM과 SRAM을 거쳐 메모리와 MAC 어레이 사이의 데이터를 단계적으로
+이동시킨다. 이러한 메모리 중심 구조에서 커널 성능은 연산 처리량만이
+아니라, 메모리 계층과 PE 사이에 데이터를 얼마나 효율적으로 배치하고,
+이동시키며, 공유하는가에 의해 결정된다. 따라서 AHBM에서의 AI 커널
+최적화는 커널 알고리즘과 아키텍처 메커니즘을 함께 발전시키는
+하드웨어--소프트웨어 코디자인을 필요로 한다.
+
+상세한 성능 분석과 신속한 설계 탐색을 위해, 우리는
+\textbf{KernBench}---AHBM을 위한 소스 수준 이산 사건 시뮬레이션
+플랫폼---을 개발하였다. KernBench는 메모리 시스템 지연, PE 실행
+모델, PE 간 통신, 호스트 측 오케스트레이션을 포함한 AHBM 실행 모델을
+구현하면서, 커널과 호스트 소프트웨어를 소스 코드로부터 직접 실행한다.
+이를 통해 실행 거동을 세밀하게 관찰할 수 있을 뿐만 아니라, 컴파일러나
+런타임과 같은 상위 소프트웨어 스택과 독립적으로 하드웨어--소프트웨어
+코디자인 선택을 체계적으로 평가할 수 있다. 본 보고서에 제시된 모든
+결과는 KernBench를 사용하여 얻은 것이며, KernBench는 본 연구에서
+다루는 모든 메커니즘과 커널의 공통 평가 플랫폼 역할을 한다.
+
+본 보고서는 LLM 추론에서 가장 성능 및 대역폭 임계 구성요소 중 하나인
+Grouped-Query Attention(GQA)에 초점을 맞춘다. Llama~3, Mistral과
+같은 최신 디코더 전용 모델은 GPT-3 형태의 다중-헤드 어텐션(MHA)으로
+부터 GQA---여러 질의 헤드가 하나의 KV 헤드를 공유함으로써 KV 캐시
+용량과 메모리 대역폭 요구량을 줄이는 방식---로 대거 전환되었다. GQA는
+모델 수준에서 시스템 효율을 개선하지만, 이를 AHBM에 효율적으로
+매핑하기 위해서는 세 가지 아키텍처 요구사항이 발생한다: PE 간 통신을
+최소화하기 위한 KV 캐시와 가중치의 최적 배치, 불가피한 PE 간 트래픽에
+대한 저오버헤드 지원, 그리고 각 PE 내에서 메모리 접근과 연산의
+효율적인 파이프라이닝이다.
+
+이러한 요구사항을 충족하기 위해, 본 보고서는 세 가지 하드웨어--
+소프트웨어 코디자인 메커니즘을 제안한다. 첫째, GQA-aware 데이터
+배치는 TCM/SRAM/HBM 계층에 KV 캐시와 가중치를 분산 배치하여 통신
+오버헤드를 줄이고 데이터 지역성을 향상시킨다. 둘째, PE\_IPCQ는 리덕션
+등 통신 집약적 연산을 위한 효율적인 온디바이스 집합 통신 프리미티브를
+제공한다. 셋째, composite-command GEMM 파이프라인은 PE\_SCHEDULER의
+제어 하에 각 PE 내부에서 메모리 이동과 연산을 긴밀하게 파이프라이닝
+하여, 명령 오버헤드를 줄이면서도 MAC 어레이의 가동률을 효율적으로
+유지한다.
+
+연산 측 인에이블러(composite-command GEMM 파이프라인,
+\S\ref{sec:gemm})와 통신 측 인에이블러(PE\_IPCQ,
+\S\ref{sec:allreduce})는 우선 독립적으로 개발 및 평가된다. 이후 융합
+GQA 커널(\S\ref{sec:gqa})은 이들을 GQA-aware 데이터 배치와 결합하여
+AHBM 상에서 종단간 어텐션 구현을 시연한다. 대응 관계는 직접적이다:
+어텐션의 $QK^{\top}$와 $PV$ 곱은 정확히 composite-command 파이프라인이
+이득을 주는 GEMM이며, KV 리덕션은 정확히 PE\_IPCQ가 이득을 주는 집합
+연산이다. 이러한 메커니즘들이 결합되어, 융합 GQA 커널은 AHBM의 HBM
+대역폭을 효율적으로 활용할 수 있게 된다.
+
+본 연구에서 GQA가 주된 동기 부여 워크로드 역할을 하지만, 도출된
+메커니즘들은 훨씬 광범위한 AI 커널 분류군에서 재사용 가능한 구성
+요소로 의도되었다. PE\_IPCQ는 분산 및 통신 집약적 워크로드 전반에
+걸쳐 집합 통신을 지원할 수 있으며, composite-command 실행은 GEMM
+기반 커널, 피드포워드 네트워크(FFN), 정규화, 그 외 융합 연산자
+파이프라인에 적용 가능하다. 계층적 데이터 배치 프레임워크와 함께 이
+메커니즘들은 AHBM에서의 향후 AI 커널과 통신 라이브러리의 기반을
+형성한다. 2026년 하반기에는 이 기반을 FFN 및 MoE 주도 워크로드와
+완전한 LLM 실행의 종단간 최적화로 확장할 예정이다.
+
+본 보고서의 나머지 부분은 다음과 같이 구성된다.
+\S\ref{sec:platform}에서는 KernBench 플랫폼과 본 연구 전반에 걸쳐
+사용된 AHBM 구성을 기술한다. \S\ref{sec:gemm},
+\S\ref{sec:allreduce}, \S\ref{sec:gqa}에서는 각각 composite-command
+GEMM 파이프라인, PE\_IPCQ 집합 통신, 융합 GQA 커널을 다룬다.
+\S\ref{sec:discussion}에서는 이러한 결과의 광범위한 아키텍처적 함의를
+논의한다. 마지막으로 \S\ref{sec:conclusion}과 \S\ref{sec:future}에서는
+주요 결과를 요약하고 FFN, MoE, 전체 모델 최적화에 관한 향후 과제를
+제시한다.

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

@@ -0,0 +1,465 @@
+\begin{figure*}[t]
+\centering
+\begin{subfigure}[b]{0.495\textwidth}
+  \centering
+  \includegraphics[width=\linewidth,height=0.7\linewidth,keepaspectratio]{sip_architecture.pdf}
+  \caption{SIP Architecture}
+  \label{fig:sip-arch}
+\end{subfigure}\hfill
+\begin{subfigure}[b]{0.42\textwidth}
+  \centering
+  \includegraphics[width=\linewidth,height=0.84\linewidth]{cube_architecture.pdf}
+  \caption{CUBE Architecture}
+  \label{fig:cube-arch}
+\end{subfigure}
+
+\vspace{0.6em}
+
+\begin{subfigure}[t]{\textwidth}
+  \centering
+  \includegraphics[width=\linewidth]{pe_architecture.png}
+  \caption{PE level (zoom-in of one PE in (b)): \textsf{PE\_CPU}
+  dispatches commands to \textsf{PE\_SCHED}, which routes tile-token
+  streams through \textsf{PE\_DMA}, \textsf{PE\_FETCH\_STORE}, and
+  \textsf{GEMM}/\textsf{MATH} engines; \textsf{PE\_IPCQ} is the
+  on-device collective control plane.}
+  \label{fig:pe-arch}
+\end{subfigure}
+\caption{Modeled hardware graph at the SIP, CUBE, and PE levels.
+This figure is an \emph{illustrative topology} chosen for
+readability; the \emph{experimental configuration} used throughout
+this report (port counts, slice counts, and link parameters) is
+specified in \S\ref{sec:hw} / Table~\ref{tab:hw}, and numbers in the
+two may differ. KernBench is not tied to this particular
+arrangement: each box is a modeled component node, each line a
+directed link with bandwidth and propagation attributes, and any
+topology that respects those attributes is supported.}
+\label{fig:hw-arch}
+\end{figure*}
+
+\begin{figure*}[t]
+\centering
+\includegraphics[width=\textwidth]{latency_model.png}
+\caption{Conceptual schematic of the latency model. Two source nodes
+(Requester A/B) inject flits through a chain of routers into a
+destination node; on the shared edge between routers, flits from the
+two transactions are interleaved flit-by-flit by the wire's FIFO
+arrival order. End-to-end latency is the sum of four contributions:
+\textbf{per-node overhead} (the switch's fixed processing cost,
+shown in light yellow---the same colour as the destination's
+processing-logic block), \textbf{per-edge transmission}
+($\textit{flit\_size}/\textit{BW}$ on each wire),
+\textbf{drain} (per-flit service occupancy at the destination's
+channel), and \textbf{queuing delay} (waiting in a FIFO when a
+shared resource is busy). The places where queuing actually
+accumulates are highlighted in green: the router output queue and
+the destination's input queue. This is the model, not a
+measurement; specific bandwidths and overheads are listed in
+\S\ref{sec:hw} (Table~\ref{tab:hw}).}
+\label{fig:latency-model}
+\end{figure*}
+
+\section{The KernBench Platform}
+\label{sec:platform}
+
+All results in this report are produced on \emph{KernBench}, a
+system-level discrete-event simulator for LLM kernels running on AHBM.
+This section explains why the platform exists, the device and
+execution model it presents to a kernel writer, how its latency model
+turns that execution into a number, and the concrete hardware
+configuration used for every experiment that follows.
+
+\subsection{Why KernBench}
+\label{sec:why}
+
+In a production end-to-end (E2E) stack, kernel performance is entangled
+with every layer above the hardware: the compiler's tiling and
+scheduling choices, the framework's operator dispatch, the collective
+library, and the runtime. Good E2E numbers require \emph{all} of those
+layers to be co-optimized, which makes it hard to answer a narrower but
+more fundamental question: \emph{given the hardware, how fast can a
+well-written kernel be, and which hardware features actually make it
+faster?}
+
+KernBench is built to answer exactly that question. Kernels are written
+and executed at the \emph{source level}---as algorithmic descriptions
+in a small tile-oriented kernel API---with no dependency on a compiler
+or any other software-stack layer. The simulator takes the kernel and a
+hardware topology and reports the latency the modeled hardware would
+deliver. This isolation is deliberate: it lets us study algorithm-level
+optimizations (how to tile a GEMM, how to schedule a collective, how to
+fuse an attention kernel) and the hardware features that support them,
+without the confound of compiler maturity or framework overhead. The
+cost is that KernBench numbers are \emph{not} E2E latencies; they are
+the achievable-kernel latencies an ideal software stack would expose.
+
+\subsection{Device and execution model}
+\label{sec:exec}
+
+KernBench models AHBM as a hierarchy of SIPs, CUBEs, and processing
+elements (PEs). At the system level, multiple SIPs are connected
+through inter-package links, while each SIP contains a collection of
+CUBEs joined by an on-package interconnect
+(Fig.~\ref{fig:hw-arch}\subref{fig:sip-arch}).
+
+Each CUBE contains eight PEs, shared SRAM, HBM controllers, an
+\textsf{M\_CPU} control processor, and an intra-CUBE router mesh
+(Fig.~\ref{fig:hw-arch}\subref{fig:cube-arch}). Together these
+components form the execution substrate for all kernels evaluated in
+this report. This organization reflects the memory-centric nature of
+AHBM: each PE is paired with a dedicated slice of HBM bandwidth and
+local on-PE storage (TCM), so compute lives next to the data it
+consumes rather than fetching it through a far-away memory
+controller. The platform's performance question is therefore not
+``how many FLOPs can the chip do'' but ``how well can a kernel keep
+each compute engine fed from its locally-attached memory while
+moving the unavoidable traffic between PEs efficiently.''
+
+Within a PE, commands are dispatched by \textsf{PE\_CPU} to
+\textsf{PE\_SCHED}, which routes work to specialized execution
+engines---DMA, FETCH/STORE, GEMM, vector-math, and IPCQ
+(Fig.~\ref{fig:hw-arch}\subref{fig:pe-arch}). Commands come in two
+flavours. \emph{Atomic} commands target a single engine---a plain
+DMA read, a GEMM tile, a vector-math op, an IPCQ send/receive---and
+are the natural unit for short or special-purpose work; the
+\textsf{PE\_CPU} itself also runs control-plane work directly.
+\emph{Composite} commands, by contrast, carry an ordered pipeline of
+operations across multiple engines (\textsf{DMA\_READ} $\rightarrow$
+\textsf{FETCH} $\rightarrow$ \textsf{GEMM}/\textsf{MATH} $\rightarrow$
+\textsf{STORE} $\rightarrow$ \textsf{DMA\_WRITE}) that the
+\textsf{PE\_SCHED} tiles and streams without per-tile redispatch. The
+composite form is the substrate for the GEMM optimization
+(\S\ref{sec:gemm}) and, combined with on-PE collectives, for fused
+attention (\S\ref{sec:gqa}).
+
+KernBench is layered along the flow of a request:
+
+\begin{itemize}
+  \item The \textbf{runtime API} is host-facing and
+  topology-agnostic---it deploys tensors and launches kernels but knows
+  nothing about routing or interconnect.
+  \item The \textbf{simulation engine} schedules discrete events and
+  routes every request through the modeled graph.
+  \item The \textbf{components} are device-side nodes that model
+  hardware behaviour: the per-PE blocks (scheduler, DMA, GEMM and
+  vector-math engines, TCM, IPCQ), the NoC routers, the HBM
+  controllers, and the inter-chiplet links.
+\end{itemize}
+
+Data and timing are handled in two passes, so that a kernel's
+numeric results and its latency are computed consistently but
+independently. \textbf{Pass~1 (timing)} runs the kernel under
+the discrete-event engine: memory ops (\textsf{tl.load},
+\textsf{tl.store}) execute against a host-side \textsf{MemoryStore}
+and return real tensor data, while compute ops (GEMM, vector-math)
+only emit records into an op-log carrying their operands, shapes,
+dtypes, and scheduled time. \textbf{Pass~2 (data)} replays that
+op-log offline in numpy, producing the actual numeric outputs and
+comparing them against a reference computation under per-dtype
+tolerances (e.g.\ \texttt{rtol}/\texttt{atol} $=10^{-3}$ for f16).
+Pass~2 is optional---runs that need only latency skip it---but when
+enabled it guarantees that every reported timing number corresponds
+to a kernel whose numeric output has been verified end-to-end.
+
+\subsection{Latency model: graph traversal and contention}
+\label{sec:latency}
+
+The modeled hardware hierarchy described above is represented
+internally as a directed graph. Nodes correspond to hardware
+components---PEs, routers, memory controllers, SRAM blocks, IO
+chiplets---while edges represent communication links with associated
+bandwidth (\si{\giga\byte\per\second}) and propagation
+(\si{\nano\second}) attributes. Every operation in KernBench---DMA
+transfers, remote memory accesses, collective communication, command
+dispatch---is modelled as a traversal through this graph. End-to-end
+latency is decomposed into four contributions accumulated along the
+traversal path (Fig.~\ref{fig:latency-model}): \textbf{per-node
+overhead} at each component, \textbf{per-edge transmission} on each
+wire, \textbf{drain} (per-flit service occupancy) at the destination,
+and \textbf{queuing delay} at the shared FIFOs that the wire and the
+destination share between concurrent transactions.
+
+\paragraph{The hardware as a graph.} The topology is compiled once at
+configuration time into this graph and is never mutated during a run.
+There are no hidden shortcuts, implicit bypasses, or magic paths: if a
+request reaches its destination, the path it took is explicit in the
+graph, and the latency it incurred is the sum of the per-node and
+per-edge costs paid along that path. The same graph representation
+applies recursively at every hierarchy level---system, SIP, CUBE, and
+PE (Fig.~\ref{fig:hw-arch}).
+
+\paragraph{From graph to discrete-event simulation.} The graph is
+driven by a discrete-event engine. Two kinds of events advance
+simulation time: \emph{node events} (component switching overhead,
+service completions such as an HBM channel commit or a GEMM tile
+finish) and \emph{edge events} (the flit-by-flit serialization of a
+payload across a bandwidth-limited link). The engine maintains a
+priority queue of pending events ordered by their scheduled time and
+fires them one at a time, with ties broken under a deterministic
+policy so that the same kernel on the same topology always yields the
+same trace. Per-request correlation IDs are stamped at injection and
+carried through every hop, so the path from injection to completion
+is fully traceable. Every modeled latency contribution
+corresponds to exactly one of these events on exactly one node or
+edge---there is no slack in the budget.
+
+\paragraph{Latency contributions.} Three kinds of latency accumulate
+along a traversal: (i) \emph{per-node fixed overhead}---each component
+carries a small switching cost (router decode, controller pickup,
+scheduler handoff); (ii) \emph{per-edge transfer time}---each link's
+payload is decomposed into fixed-size flits (default
+\SI{256}{\byte}), and each flit arrives at
+$\text{prop}+\text{flit\_bytes}/\text{bw}$ after the previous one,
+giving wormhole semantics across multi-hop paths; and (iii)
+\emph{per-service occupancy}---memory controllers, GEMM stages, and
+collective engines hold the request for their service time before
+releasing it downstream. Each of these is attached to a specific node
+or edge in the graph; together they make up the entire latency budget.
+
+Beyond data movement and execution latency, KernBench
+also models the control-plane cost required to issue work
+to accelerator engines. Since every DMA, GEMM, vector-
+math, and IPCQ operation is initiated through the
+\textsf{PE\_CPU} $\rightarrow$ \textsf{PE\_SCHED} path,
+command dispatch latency contributes to the end-to-end
+execution time of all kernels.
+
+\paragraph{Command dispatch overhead model.} The cost incurred by
+the \textsf{PE\_CPU} when it dispatches a command to one of the
+accelerator engines is modelled structurally rather than with a
+per-operation calibration table. The \textsf{PE\_CPU} charges, per
+command,
+\[
+  d_{\text{cmd}} = \textsf{FIXED} + b_{\text{logical}} \cdot R,
+\]
+This linear-in-size form reflects the serialization cost of writing a
+command descriptor into the scheduler queue: a fixed per-command
+bookkeeping cost plus a byte-wise descriptor-transfer cost.
+Concretely, $b_{\text{logical}}$ is the command's hardware-logical byte
+size, \textsf{FIXED} captures the fixed per-command cost (queue-tail
+update, completion registration) and $R$ captures the per-byte cost of
+serializing the command descriptor into the scheduler queue. This
+\textsf{FIXED} is distinct from Table~\ref{tab:hw}'s
+\SI{2}{\nano\second} / \SI{1}{\nano\second} \textsf{PE\_CPU} /
+\textsf{PE\_SCHED} fixed costs: the latter are per-component
+traversal overheads each command pays when it transits those nodes,
+whereas the \textsf{FIXED} term here is the command-descriptor
+issue cost. \textsf{FIXED} depends on the command class: a single-op
+command --- one engine operation, i.e.\ a single DMA descriptor, GEMM,
+or elementwise issue --- carries a lighter 8-cycle \textsf{FIXED},
+while the composite command, which the scheduler expands into a
+multi-stage tile-feeder plan, carries 40 cycles
+(\S\ref{sec:gemm-vs-async}). With the composite \textsf{FIXED} $= 40$
+cycles and $R = 0.0625$ cycles/byte (i.e.\ \SI{16}{\byte\per\cycle},
+at \SI{1}{\giga\hertz}) a typical composite lands at roughly
+\SI{43}{\nano\second}, and a hard cap on a composite's descriptor size
+prevents the model from rewarding arbitrarily large fused commands
+beyond what real descriptor queues accept. In the configurations measured here, command issue is
+not the bottleneck---data movement is---so this term stays small
+relative to DMA and collective time.
+
+\begin{table*}[t]
+\centering
+\caption{PE\,$\to$\,HBM DMA latency probe at varying hop distances
+(\SI{32}{\kibi\byte} transfer; output captured from \texttt{kernbench
+probe}). \emph{Util\%} is the achieved bandwidth as a fraction of the
+path's bottleneck-edge bandwidth.}
+\label{tab:probe-pe-dma}
+\small
+\begin{tabular}{@{}lrrr@{}}
+\toprule
+Case & Latency~(\si{\nano\second}) & Util\% (\,32\,KiB) & Util\% (\,1\,MiB) \\
+\midrule
+PE\,$\to$\,local HBM             & 141.0 & 90.8 & 100.0 \\
+PE\,$\to$\,same-half HBM         & 147.9 & 86.6 & ~99.9 \\
+PE\,$\to$\,cross-half HBM        & 161.2 & 79.4 & ~99.8 \\
+PE\,$\to$\,cross-CUBE (best)     & 330.5 & 77.5 & ~99.6 \\
+PE\,$\to$\,cross-CUBE (worst)    & 677.1 & 37.8 & ~97.8 \\
+\bottomrule
+\end{tabular}
+\end{table*}
+
+\subsubsection{Congestion and contention modeling}
+\label{sec:congestion}
+
+The simulator's
+sharpness comes from how it models contention for those nodes and
+edges. \emph{Every directed edge has a FIFO}: an arriving flit takes
+its bandwidth-limited transfer time on top of whatever earlier flits
+are still being served, so a busy link queues later traffic behind
+earlier traffic rather than transferring everything at peak BW.
+\emph{HBM is modelled with per-pseudo-channel parallelism}: a
+stateless array of channel-availability timestamps with address-based
+channel selection captures the bank-level concurrency that real HBM
+exposes, so 64 channels per CUBE deliver real parallelism on uniform
+addresses but a hot channel surfaces as the bottleneck on skewed ones.
+\emph{Every component has a serial worker}: a router carrying two
+heavy streams interleaves them at flit granularity in arrival order
+rather than fanning out for free, so two concurrent collectives sharing
+a link share its bandwidth, not double it. Without these mechanisms
+the simulator would simply re-confirm the peak-BW roofline; with them,
+it reveals where the real bottlenecks form and which hardware levers
+actually relieve them---which is exactly the question the codesign
+work in this report turns on.
+
+
+\subsection{Accuracy}
+\label{sec:accuracy}
+
+The model is precise about the effects that
+dominate kernel latency on this class of hardware: per-edge bandwidth
+occupancy and flit-level serialization, HBM pseudo-channel parallelism,
+and per-component switching overhead. Two independent cross-checks
+drawn from the experiments in this report confirm that this precision
+translates into physically reasonable kernel latencies.
+
+First, in the GEMM study (\S\ref{sec:gemm}), simulator-measured MAC
+efficiency tracks an analytic ideal-pipeline model within
+\SI{1.4}{ppt} across every swept shape---from a single-tile
+$M{=}K{=}N{=}32$ at $\sim\SI{7.7}{\percent}$ up to the deep-$K$
+$K{=}3072$ case at $\sim\SI{90}{\percent}$. The residual gap is
+fill/tail overhead the closed-form pipeline model omits.
+
+Second, in the all-reduce study (\S\ref{sec:allreduce},
+Fig.~\ref{fig:allreduce-cmp}), simulator latency for a 2D-torus over
+six devices follows the expected startup-plus-per-packet shape across
+the entire payload sweep---tight at small payloads where startup
+dominates, and within a single-digit multiplicative factor at the
+largest payloads, where the residual gap is explained by per-router
+switching the analytic shape elides.
+
+Third, the simulator's per-traversal behaviour can be probed
+directly. Running \texttt{kernbench probe} on the modelled topology
+issues a sequence of PE-to-HBM DMA reads at progressively greater
+hop distances and reports the per-component overhead, per-edge
+serialization, and per-PC drain that the model charges
+(Table~\ref{tab:probe-pe-dma}). Three properties stand out. First,
+the reported latency increases \emph{monotonically} with hop
+count---from \SI{141}{\nano\second} at the local HBM slice to
+\SI{677}{\nano\second} at the worst-case remote-CUBE slice---with the
+increment per added hop matching the per-router overhead and the
+UCIe-link cost in Table~\ref{tab:hw}. Second, the bandwidth-saturation
+curves track the wire's serialization model: at \SI{1}{\mebi\byte}
+transfers the local PE DMA reaches \SI{100}{\percent} of the per-edge
+bandwidth limit, while the longest cross-CUBE path saturates at
+\SI{97.8}{\percent}---exactly the gap that the per-edge propagation
+and per-router overhead model would predict. Third, the simulator
+passes the internal-consistency invariants the probe builds in
+(monotonic hop progression; D2H reads at least as long as the
+equivalent H2D writes because of the reverse-path acknowledgement;
+cross-CUBE best less than cross-CUBE worst). These are properties
+that hold \emph{only} when the underlying graph traversal, FIFO
+contention, and propagation models are internally
+self-consistent---a model error in any of them would surface as a
+non-monotonic or under-utilising curve here long before it polluted a
+kernel-level measurement.
+
+The known simplifications---FIFO router arbitration (instead of
+round-robin), HBM scheduler without write-buffer reordering, no bank
+conflict, no refresh or thermal effects, and no upstream
+backpressure---are the price of a deterministic, inspectable model.
+These simplifications bound the absolute accuracy, but the agreement
+with both analytic models and the probe's internal-consistency
+checks above indicates that KernBench is sufficiently accurate for
+evaluating the \emph{relative} hardware--software design trade-offs
+(tiling A vs.\ B, topology X vs.\ Y, with vs.\ without composite
+command, mesh vs.\ torus) that are the primary objective of this
+work.
+
+
+
+\subsection{Modeled hardware configuration}
+\label{sec:hw}
+
+Table~\ref{tab:hw} summarizes the hardware configuration
+used throughout this report. The intent of this
+configuration is not to model a specific product, but to
+represent a realistic memory-centric accelerator and to
+provide a consistent baseline for evaluating hardware--
+software co-design mechanisms.
+
+The compute capability within each CUBE is provisioned
+such that inference workloads can effectively saturate
+the available HBM bandwidth. The aggregate compute
+throughput is therefore balanced against memory-system
+bandwidth rather than being intentionally over- or
+under-provisioned. Concretely, 8 PEs $\times$
+\SI{8}{\tera\flop\per\second} give roughly
+\SI{64}{\tera\flop\per\second} of f16 compute per CUBE
+against \SI{2048}{\giga\byte\per\second} of HBM bandwidth,
+which places the balanced point at about
+$\sim$31~FLOP/byte; inference decode kernels typically
+operate well below that, so HBM bandwidth---not compute---is
+the natural ceiling on this configuration. For reference, the
+GQA decode kernels evaluated in \S\ref{sec:gqa} operate at
+arithmetic intensity well below this balance point, so their
+ceiling is set by HBM and inter-PE traffic rather than by GEMM
+throughput. HBM bandwidth
+is distributed evenly across the PEs within a CUBE, with
+each PE responsible for servicing approximately one-eighth
+of the CUBE's memory bandwidth through its dedicated HBM
+channels.
+
+The on-chip interconnect is configured using bandwidth
+and latency parameters representative of commercially
+available mesh-network IPs. The goal is not to study a
+particular NoC implementation, but rather to evaluate
+kernel behavior under a realistic baseline communication
+fabric.
+
+For die-to-die communication, UCIe-A bandwidth
+characteristics are used as the reference point for
+inter-CUBE links. Inter-SIP communication is modeled
+using PCIe-class links. Together, these assumptions
+provide a representative communication hierarchy for
+evaluating collective communication and distributed
+kernel execution.
+
+Unless otherwise stated, all experiments use this
+configuration. Workload-specific parameters such as
+matrix dimensions, sequence lengths, and collective
+payload sizes are introduced in their respective sections.
+
+\begin{table}[t]
+\centering
+\caption{Modeled hardware configuration}
+\label{tab:hw}
+\small
+\begin{tabular}{@{}ll@{}}
+\toprule
+\textbf{Parameter} & \textbf{Value} \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Hierarchy}} \\
+SIPs & 2 (1D ring) \\
+CUBEs per SIP & 16 ($4\times4$ mesh) \\
+PEs per CUBE & 8 (4 corners $\times$ 2) \\
+PEs total & 256 \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Processing element (PE)}} \\
+GEMM engine peak & \SI{8}{\tera\flop\per\second} (f16) \\
+TCM (on-PE) & \SI{16}{\mega\byte}, \SI{512}{\giga\byte\per\second} R/W \\
+\quad kernel scratch & \SI{1}{\mega\byte} \\
+DMA engines & 1 read + 1 write \\
+\textsf{PE\_CPU} fixed cost   & \SI{2}{\nano\second} \\
+\textsf{PE\_SCHED} fixed cost & \SI{1}{\nano\second} \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Memory (per CUBE)}} \\
+HBM capacity & \SI{48}{\giga\byte} (8 slices) \\
+HBM aggregate BW & \SI{2048}{\giga\byte\per\second} \\
+HBM pseudo-channels & 64 (8 per PE), \SI{32}{\giga\byte\per\second} each \\
+SRAM (shared) & \SI{32}{\mega\byte}, \SI{128}{\giga\byte\per\second} link \\
+HBM burst & \SI{256}{\byte} \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Interconnect}} \\
+Intra-CUBE NoC link & \SI{256}{\giga\byte\per\second}, \SI{2}{\nano\second}/router \\
+Inter-CUBE (UCIe PHY) & \SI{512}{\giga\byte\per\second}, \SI{8}{\nano\second}, XY routing \\
+Inter-SIP (PCIe) & \SI{768}{\giga\byte\per\second} per endpoint \\
+\midrule
+\multicolumn{2}{@{}l}{\emph{Command-issue cost model (defaults)}} \\
+FIXED per single-op command & 8 cycles \\
+FIXED per composite command & 40 cycles \\
+per-byte rate $R$ & 0.0625 cycles/byte (\SI{16}{\byte\per\cycle}) \\
+composite size cap & \SI{1024}{\byte} \\
+\bottomrule
+\end{tabular}
+\end{table}

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

@@ -0,0 +1,401 @@
+\section{GEMM Acceleration via the Composite Command}
+\label{sec:gemm}
+
+GEMM is the compute core of every transformer block---the QKV
+projections, the attention score and context products, and the
+feed-forward matrices are all matrix multiplications. On a tiled
+accelerator a single logical GEMM expands into many hardware tiles, and
+each tile must be read from HBM, fetched into the register file, computed,
+stored, and written back. If every one of those tile-stage steps were an
+independently issued command, two costs would dominate. First, the host
+and the PE control processor would pay a per-command issue overhead
+$O(\text{tiles}\times\text{stages})$ times, which for a deep-$K$ reduction
+is hundreds to thousands of issues. Second, with stages dispatched
+one-at-a-time the scheduler cannot overlap the DMA of the next tile with
+the compute of the current one---the pipeline never fills, and the MAC
+array sits idle waiting for data. The hardware question is therefore: what
+issue mechanism lets a single GEMM saturate the MAC array without drowning
+in command overhead?
+
+\subsection{Design}
+
+The answer is the \emph{composite command}. A single command carries
+the ordered five-stage tile pipeline (\textsf{DMA\_READ},
+\textsf{FETCH}, \textsf{GEMM}, \textsf{STORE}, \textsf{DMA\_WRITE}),
+and the PE scheduler splits the payload into hardware tiles
+(here $32\times64\times32$), emitting one tile token per tile.
+Subsequent stages are reached by \emph{token self-routing} between
+the on-PE engines, so a tile flows through the chain without
+returning to the scheduler between stages. Because the whole
+pipeline is described by one command, the issue cost is paid once
+per GEMM rather than once per tile-stage, and the scheduler is free
+to keep every stage busy on different tiles simultaneously---tile
+$i$'s GEMM overlaps tile $i{+}1$'s DMA read. A multi-operation composite
+additionally lets an epilogue (for example a vector-math step) ride the
+same tile loop, firing per $K$-tile, per output tile, or once per kernel
+according to its declared scope; this is what later allows an attention
+kernel to fuse its softmax work into the GEMM pipeline
+(\S\ref{sec:gqa}).
+
+\subsection{Results}
+
+We sweep eight GEMM shapes spanning square, tall, wide, and deep-$K$
+geometries under two operand-staging variants that bracket the
+realistic LLM cases. In \textsf{load\_ref}, the activation $A$ is
+pre-staged on chip and only the weight $W$ streams from HBM during
+the composite (the ``activation in TCM, weights from HBM'' case
+typical of decoding with a small batch). In \textsf{ref\_ref}, both
+operands stream from HBM during the composite (the ``cold both
+sides'' case, where the activation does not fit on-chip or is itself
+the output of a producer composite that wrote back to HBM). Both
+variants run the same composite command; only the up-front staging
+of $A$ differs.
+
+We evaluate the composite GEMM along two axes that together describe
+how well it lands on the hardware: \emph{HBM bandwidth utilization}
+(does the workload saturate the per-PE HBM bandwidth, i.e.\ is it
+memory-bound on this configuration?) and \emph{achieved GEMM
+throughput} (what fraction of the \SI{8}{\tera\flop\per\second}
+per-PE peak does the kernel actually deliver). Both metrics use the
+composite window as the denominator, so the up-front \textsf{tl.load}
+of $A$ in \textsf{load\_ref} is excluded — only HBM traffic and
+compute that happen \emph{inside} the composite count.
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gemm_hbm_bw_util.png}
+\caption{Per-PE HBM bandwidth utilization during the composite window,
+for the two staging variants. The ceiling is the per-PE HBM bandwidth
+(\SI{256}{\giga\byte\per\second}; dashed line at \SI{100}{\percent}).
+\textsf{ref\_ref} streams both $A$ and $W$ and so always sits at a
+higher BW-utilization than \textsf{load\_ref} on the same shape;
+both variants saturate ($>\SI{85}{\percent}$) once the kernel
+has enough total bytes to keep the link busy (deep-$K$ $K{=}3072$,
+or low-reuse output-dominated $M{=}128,K{=}8,N{=}128$). Small or
+single-tile shapes do not have enough total traffic to saturate the
+link and are bottlenecked by pipeline fill rather than HBM.}
+\label{fig:gemm-bw}
+\end{figure}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gemm_per_pe_tflops.png}
+\caption{Per-PE GEMM throughput delivered during the composite window
+(reference line at the \SI{8}{\tera\flop\per\second} per-PE engine
+peak). At the deep-$K$ corner $M{=}32,K{=}3072,N{=}32$ the
+activation-pre-staged kernel reaches
+\SI{7.18}{\tera\flop\per\second} ($\sim\SI{90}{\percent}$ of peak);
+the both-from-HBM variant drops to
+\SI{3.83}{\tera\flop\per\second} ($\sim\SI{48}{\percent}$) at the
+same shape because the second operand doubles HBM pressure and the
+kernel hits the BW ceiling shown in Fig.~\ref{fig:gemm-bw}.
+Under-tile shapes ($M{=}K{=}N{=}32$, $M{=}8$, $K{=}8$) are
+hard-capped by tile-fill at \SI{50}{\percent}, \SI{25}{\percent},
+\SI{12.5}{\percent} of peak regardless of staging.}
+\label{fig:gemm-tflops}
+\end{figure}
+
+The composite reaches the hardware roofline in two distinct regimes.
+\emph{Memory-bound}: $K{=}3072$ deep-$K$ saturates HBM (\SI{88}{\percent}
+load\_ref, \SI{94}{\percent} ref\_ref) and the low-reuse
+$M{=}128,K{=}8,N{=}128$ corner saturates even harder
+($>\SI{96}{\percent}$ in both), because back-to-back output
+writes dominate. \emph{Compute-rich}: at the same deep-$K$ shape,
+\textsf{load\_ref} delivers \SI{7.18}{\tera\flop\per\second}
+($\sim\SI{90}{\percent}$ of the per-PE peak) — the configuration's
+clean win. The \emph{distance between the two variants} at the same
+shape is the cost of \emph{not} pre-staging $A$: at $K{=}3072$ the
+throughput halves (\SI{7.18}{\tera\flop\per\second} $\to$
+\SI{3.83}{\tera\flop\per\second}) because the second operand doubles
+HBM pressure
+and the kernel hits the BW ceiling. This is the operational
+take-away — when the activation can live on-chip the composite
+delivers near-peak GEMM; when it cannot the BW ceiling dominates and
+throughput halves.
+
+Figure~\ref{fig:gemm-util} validates that simulator-measured
+efficiency tracks an analytic ideal-pipeline model (within
+\SI{1.4}{ppt} across every swept shape), and
+Figure~\ref{fig:gemm-stages} resolves the per-stage engine wall-clock
+that backs the above interpretation.
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gemm_mac_utilization_theoretical_vs_measured.png}
+\caption{GEMM MAC utilization and efficiency, theoretical vs.\ measured
+(\textsf{load\_ref} staging). Tile-fill sets the ceiling: under-tile
+shapes (marked $\ast$) such as $M{=}K{=}N{=}32$, $M{=}8$, and $K{=}8$
+cannot fill the MAC tile and cap at \SI{50}{\percent},
+\SI{25}{\percent}, \SI{12.5}{\percent}. For tile-filling shapes,
+efficiency climbs with tile count---from
+\textasciitilde\SI{15}{\percent} at one tile to \textasciitilde%
+\SI{90}{\percent} measured at 48 tiles ($K{=}3072$)---and the
+measured bars track the analytic ideal-pipeline prediction within
+\SI{1.4}{ppt} across every shape.}
+\label{fig:gemm-util}
+\end{figure}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gemm_stage_breakdown.png}
+\caption{Per-stage engine wall-clock (DMA in, Fetch, GEMM, DMA out)
+under \textsf{load\_ref} staging. For the deep-$K$ shape ($K{=}3072$,
+48 tiles) DMA in, Fetch, and GEMM are all close to
+\textasciitilde\SI{785}{\nano\second}, running concurrently in a
+tightly pipelined regime while DMA-out is negligible. For the
+low-reuse shape ($M{=}128,K{=}8,N{=}128$, 16 output tiles) DMA-out
+grows to \textasciitilde\SI{336}{\nano\second} while GEMM compute is
+\textasciitilde\SI{262}{\nano\second}---a data-movement-bound regime
+where the output write becomes the largest stage.}
+\label{fig:gemm-stages}
+\end{figure}
+
+\subsection{Analysis and meaning}
+
+The composite command does not manufacture bandwidth or MAC
+throughput; it removes the two software-shaped obstacles between a
+GEMM and its hardware roofline. By paying the issue cost once per
+GEMM rather than once per tile-stage, and by chaining tile-stages
+through token self-routing, it lets the scheduler keep every stage
+busy on a different tile, so compute-rich shapes actually reach the
+efficiency their arithmetic intensity allows
+($\sim$\SI{90}{\percent} of GEMM peak at the deep-$K$
+\textsf{load\_ref} corner) and data-bound shapes actually reach
+their HBM-BW ceiling instead of stalling on command overhead.
+
+The split between \textsf{load\_ref} and \textsf{ref\_ref} also makes
+the hardware-software boundary explicit: when the activation fits in
+on-chip storage the composite is BW-headroom-rich and saturates the
+GEMM engine; when it does not, both operands compete for the same
+HBM port and the kernel is BW-bound. This is the lever the GQA
+kernel of \S\ref{sec:gqa} reaches for next — keeping the right
+working set on-chip so the composite pipeline lands in the
+compute-rich regime rather than the BW-bound one.
+
+\subsection{Why composite, and not kernel-orchestrated async loading?}
+\label{sec:gemm-vs-async}
+
+A reader familiar with double-buffered GEMM kernels on conventional
+hardware may ask: why a hardware-side composite command at all? Why
+isn't the obvious kernel-level pattern --- async-load each operand,
+overlap with compute, accumulate --- sufficient?
+
+To answer this concretely we contrast composite against two
+kernel-orchestrated baselines that have access to the same single-op
+primitives the platform exposes (\textsf{tl.load} for async DMA into
+TCM, \textsf{tl.dot} for a single-op GEMM command on TCM-resident
+operands, \textsf{tl.store} for a DMA write-back). Both baselines
+pre-stage activation $A$ identically to \textsf{load\_ref}, so the
+window measured is the engine-pipeline window with $A$'s up-front DMA
+excluded for all three kernels.
+
+\paragraph{Async-full (naive).} A single \textsf{tl.load(A)} followed by
+a single \textsf{tl.load(B)} (async, queued behind $A$),
+\textsf{tl.dot(A, B)}, and \textsf{tl.store(out)}. The kernel issues
+four commands total. The decisive constraint is that
+\textsf{tl.dot}'s \textsf{\_await\_pending(b)} blocks the GEMM
+command until the \emph{entire} $B$ has landed in TCM --- there is no
+way at the runtime API surface to express ``start computing on
+tile 0 of $B$ while tile 1 is still in flight.'' Load-of-$B$ and
+GEMM therefore serialize.
+
+\paragraph{Async-tiled (chunked prefetch).} The kernel-level workaround is to
+split $B$ along $K$ into \textsf{TILE\_K}-sized chunks, issue async
+\textsf{tl.load}s for those chunks, issue one \textsf{tl.dot} per
+chunk (each blocking only on its own $b_i$), and accumulate via
+\textsf{out = out + tl.dot(...)}. This is the standard
+double-buffered-GEMM pattern transcribed to the single-op primitives.
+We measure two versions of this kernel that differ only in
+\emph{prefetch depth}:
+\textbf{depth-2 (TCM-bounded)} keeps at most two B-chunks in flight
+at any time by issuing the next \textsf{tl.load} just before each
+\textsf{tl.dot}; \textbf{depth-$\infty$ (queue-all)} issues all $N_K$
+B-chunk loads up front so the DMA engine has the deepest possible
+request queue. For a $K{=}3072$ shape both versions emit
+48 $A$-chunk loads + 48 $B$-chunk loads + 48 \textsf{tl.dot}s + 47
+elementwise adds + 1 store = 192 host-side commands; the only
+difference is the temporal interleaving of B-load and dot dispatches.
+
+\paragraph{Why two depths.} The depth distinction matters because the
+async-full kernel and the depth-$\infty$ async-tiled kernel both pin the
+\emph{entire} $B$ in TCM simultaneously --- $K \cdot N \cdot 2$
+bytes. The on-PE scratch is capped at \SI{1}{\mebi\byte}
+(\texttt{topology.yaml: pe\_tcm.kernel\_scratch\_mb=1}), so an LLM-scale
+attention $B = K_{\text{KV}} \times d_{\text{head}}$ at
+$K_{\text{KV}}=4096, d_{\text{head}}=128$ already needs
+\SI{1}{\mebi\byte}, and at $K_{\text{KV}}=8192$ it needs \SI{2}{\mebi\byte}
+--- past the cap. async-full and queue-all async-tiled are therefore
+not just slower than composite but \emph{architecturally infeasible}
+at LLM context length. The depth-2 async-tiled kernel is the only
+kernel-level
+variant whose peak TCM footprint stays
+$O(2 \cdot \textsf{TILE\_K} \cdot N)$ regardless of $K$, the same
+order as composite's per-tile streaming buffer. It is the apples-to-apples
+comparison.
+
+\paragraph{Per-PE throughput.} Figure~\ref{fig:gemm-async} reports
+per-PE TFLOP/s for all four kernels side-by-side. The single-op
+fast-path in the dispatch cost model (\S\ref{sec:congestion}) is
+enabled, so every single-op command the async kernels emit --- DMA
+descriptors and \textsf{tl.dot}/\textsf{tl.add} alike --- is charged the
+light 8-cycle \textsf{FIXED}, not the 40-cycle composite control-path
+cost; neither async-tiled variant carries an inflated per-command cost.
+
+\paragraph{The $K{=}3072$ corner, concretely.} We take this shape
+($M{=}32, K{=}3072, N{=}32$) as the running example throughout the
+mechanism discussion because it is the regime where the four
+kernels spread the widest --- the deepest $K$ in the sweep maps
+onto $K / \textsf{TILE\_K} = 48$ hardware tiles, so per-tile costs
+amplify into the largest measurable gap. The work content is
+identical for all four kernels: $\sim$6.3 M f16 MACs and
+$\sim$386 KiB of $B$ traffic from HBM, which together require
+$\sim$786 ns of GEMM-engine compute and $\sim$781 ns of DMA on a
+saturated per-PE link. What differs is the number of host commands
+the same work is decomposed into --- 2 for composite (one
+\textsf{tl.load(A)} plus one composite), 4 for async-full, and 192
+for either async-tiled variant (48 $A$-loads + 48 $B$-loads + 48
+\textsf{tl.dot}s + 47 elementwise adds + 1 store). The
+engine-pipeline-window throughput tracks that decomposition closely:
+composite reaches \SI{7.18}{\tera\flop\per\second} (post-overlap
+limit, only \SI{10}{\percent} below the \SI{8}{\tera\flop\per\second}
+per-PE GEMM peak), async-full \SI{3.91}{\tera\flop\per\second}
+(DMA and compute serialize on a single big dot), and both async-tiled
+variants $\sim$\SI{2.53}{\tera\flop\per\second} (192 commands' worth of
+structural dispatch cost --- even at the light per-command rate ---
+accumulates on the wall). The next paragraph attributes those gaps to
+specific simulator mechanisms.
+
+\begin{figure*}[t]
+\centering
+\includegraphics[width=\linewidth]{gemm_composite_vs_async_tflops.png}
+\caption{Per-PE achieved TFLOP/s for the same shape sweep run under
+four issuance patterns: composite (one command, scheduler streams
+per-tile internally), async-full (one \textsf{tl.dot} on
+fully-loaded $B$), async-tiled with depth-2 double-buffer
+(TCM-bounded; the only kernel-level variant that scales to LLM
+context length), and async-tiled with depth-$\infty$
+(all B-tiles queued up front; included as a sanity check that the
+prefetch depth is \emph{not} what separates the async-tiled kernel from
+composite). All curves exclude $A$'s up-front DMA from the
+measurement window. Composite wins at every full-tile shape; the
+depth-2 and depth-$\infty$ async-tiled kernels deliver
+\emph{indistinguishable} throughput (e.g. \SI{2.522}{} vs.\
+\SI{2.544}{\tera\flop\per\second} at $K{=}3072$), confirming that
+prefetch depth is not the lever --- the structural per-command
+dispatch cost is. The gap between composite and the async-tiled
+kernels grows with $K_{\text{useful}}$
+(\textbf{$\sim$$2.8\times$} at $K{=}3072$, where the async-tiled
+kernels emit 192 host commands while composite emits one).
+The under-tile corner $M{=}128,K{=}8,N{=}128$ inverts: composite's
+per-tile orchestration overhead exceeds the per-tile useful work,
+and all three async kernels beat it.}
+\label{fig:gemm-async}
+\end{figure*}
+
+\paragraph{Decomposing the gap.} Three structural mechanisms separate
+composite from the kernel-level baselines, and they layer.
+
+\emph{1. Inter-engine token routing happens below the host-side
+dispatch path.} The composite encodes the full
+\textsf{DMA\_READ}$\to$\textsf{FETCH}$\to$\textsf{GEMM}$\to$\textsf{STORE}$\to$\textsf{DMA\_WRITE}
+pipeline once. The scheduler's tile-feeder loop then emits one
+\emph{tile token} per HW tile inside that one composite, and each
+token self-routes between engines after each stage finishes. The
+per-tile token routing is a scheduler-internal event, not a fresh
+host command, so it does not pay the structural CPU dispatch cost.
+At $K{=}3072$ the composite emits 48 tile tokens that flow
+fully-pipelined through five stages each --- 240 inter-engine
+hand-offs total --- behind a single command from the host's point of
+view.
+
+\emph{2. \textsf{tl.dot} cannot replicate that per-tile pipeline at the
+kernel level.} A single-op GEMM command is handled on the GEMM engine as
+a single monolithic compute timeout for the supplied $M{\times}K{\times}N$;
+there is no internal token loop that would let a streaming DMA of
+$B[i{+}1]$ overlap with the GEMM of $B[i]$ inside one
+\textsf{tl.dot}. The user can only recover inter-tile overlap by
+emitting one \textsf{tl.dot} per chunk --- which is exactly what the
+async-tiled baseline does, at the price of $N$ host-side commands.
+
+\emph{3. The host-side dispatch cost the async-tiled baseline pays is
+structural, not modelling slack.} KernBench charges every host-emitted
+command a structural CPU dispatch cost $d_{\text{cmd}} = \textsf{FIXED} +
+b_{\text{logical}} \cdot R$ (\S\ref{sec:congestion}). The cost model is
+deliberately charitable to the async kernels here: only a
+\textsf{CompositeCmd} pays the 40-cycle control-path \textsf{FIXED} (it
+alone drives a scheduler-built tile-feeder plan), while \emph{every}
+single-op command --- the 96 DMA descriptors, the 48 \textsf{tl.dot}s,
+and the 47 elementwise adds alike --- pays only the light 8-cycle
+\textsf{FIXED}, calibrated to descriptor-ring-push / single-instruction
+issue patterns in modern accelerators (NVIDIA Hopper TMA $\sim$1 ISA
+cycle, a single \textsf{mma.sync} one instruction, AMD AQL packet writes
+$\sim$5--15 cycles). So the async-tiled baseline is \emph{not} penalized
+by an inflated per-command cost on \emph{any} of its operations. Yet it
+still emits 192 host commands against composite's one, so $\sim$192
+light dispatches accumulate to $\sim$\SI{1.5}{\micro\second} of
+structural PE\_CPU time that composite never pays --- composite hides
+its 48 tiles' 240 inter-engine hand-offs as scheduler-internal events
+(mechanism 1). Layered on top, the dependency chain on the running
+accumulator serializes the 47 adds on the math engine. The net result
+is that the async-tiled kernel runs $\sim$2.8$\times$ slower than
+composite at $K{=}3072$ despite getting the inter-chunk overlap right
+--- down from $\sim$6.3$\times$ under the earlier uniform-40 cost model,
+because D8 removed the per-command overcharge, but \emph{not} closed:
+the residual gap is the command-count structure (one composite vs.\ 192
+single-ops), not a modelling artifact.
+
+\emph{4. Prefetch depth is not the lever, command count is.} The
+depth-2 and depth-$\infty$ async-tiled kernels land within \SI{1}{\percent}
+of each other at every measured shape (Figure~\ref{fig:gemm-async}).
+This is the diagnostic against a natural objection: ``surely the
+async-tiled kernel was just under-prefetching; deepen the queue and the
+DMA$\to$GEMM overlap recovers.'' Deepening the prefetch queue
+changes \emph{when} DMA descriptors hit the engine but not their
+total number or the structural dispatch cost they each pay. The
+$\sim$2.8$\times$ gap to composite is not a prefetch-depth gap; it
+is the gap between
+``one composite command with internal per-tile token routing'' and
+``$N_K$ host-side dot/add commands, each charged separately.''
+The depth-2 kernel additionally constrains peak TCM occupancy to
+$O(2 \cdot \textsf{TILE\_K} \cdot N)$, matching composite's per-tile
+streaming buffer; the depth-$\infty$ kernel needs the full $B$ in
+TCM, which makes it infeasible at LLM context length even if the
+throughput were competitive.
+
+\paragraph{Where the composite advantage doesn't apply.} The shape
+$M{=}128,K{=}8,N{=}128$ inverts the picture: composite delivers
+\SI{0.66}{\tera\flop\per\second} and both async kernels reach
+\SI{1.21}{\tera\flop\per\second}. The reason is consistent with the
+analysis above and explicit in the simulator state. Composite emits
+16 tile tokens (one per output tile) for this shape, each carrying
+$K_{\text{useful}}{=}8$ MACs across a TILE\_K$=64$ pipeline ---
+$12.5\%$ of the hardware tile's MAC slots are useful, the rest is
+K-padding. The per-tile inter-engine hand-off cost stays the same
+regardless. When the per-tile useful work is small enough that
+hand-off overhead exceeds the GEMM work itself, a single-op
+\textsf{tl.dot} --- which submits one monolithic GEMM command with no
+per-tile orchestration --- wins. The take-away is the bound on
+composite's value: it amortizes useful per-tile compute, not
+padding-dominated under-tile shapes. Real kernels at this corner are
+better served by reshape-into-batched-GEMM transforms that move
+under-tile $K$ into a tile-filling dimension before reaching the GEMM
+engine, which is exactly what the GQA decode kernel of
+\S\ref{sec:gqa} does for its $K_{\text{useful}}=\text{head\_dim}=128$
+inner reduction.
+
+\paragraph{Summary of the comparison.} Composite is the only one of
+the four kernels to combine (a) macro-command dispatch at the host
+boundary (amortizing the structural CPU cost across all the work a
+single GEMM does), (b) scheduler-internal per-HW-tile streaming of
+DMA$\rightleftarrows$compute, and (c) TCM-bounded streaming buffer.
+Kernel-orchestrated async kernels can have any two of those, not all
+three: async-full pays one host dispatch (a) but forfeits per-tile
+overlap (b) and pins all of $B$ in TCM (c); depth-$\infty$ async-tiled
+achieves inter-chunk overlap but at $N_K$ host dispatches and
+full-$B$ TCM occupancy; depth-2 async-tiled fixes the TCM
+footprint (c) but still pays $N_K$ host dispatches. The two corners
+where async catches up
+(small-$K$ where composite has nothing useful to amortize;
+under-tile shapes where per-tile useful work is sub-token) are
+diagnostic of \emph{where composite is the wrong tool}, not of a
+slack the user kernel could close.

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

@@ -0,0 +1,288 @@
+\section{PE\_IPCQ and Collective Communication}
+\label{sec:allreduce}
+
+Distributing a transformer across devices turns every tensor-parallel
+layer into a collective: partial results computed on different PEs, CUBEs,
+and SIPs must be summed and redistributed with an all-reduce. Underneath
+the algorithm this is fundamentally a PE-to-PE problem---many short
+messages flowing between neighbors as the reduction proceeds. The natural
+software realization is a per-direction ring buffer whose head and tail
+pointers the producer and consumer update atomically and poll, but our
+H2 2025 report measured this scheme end-to-end and found that the
+atomic-pointer traffic together with the consumer's polling loop dominate
+the per-message cost: the queue itself becomes the bottleneck well before
+the link runs out of bandwidth, and a pure-SW collective spends most of
+its time on metadata rather than on actually moving partials. A second,
+orthogonal problem is link sharing---if the collective rides the kernel's
+generic DMA path it competes with the GEMM's compute traffic on the same
+wires, so a large tile transfer head-of-line-blocks a pending reduction.
+This section asks how a dedicated hardware primitive can lift queue
+management off the software path entirely and, in the same design, stop
+collective traffic from stalling behind compute traffic, so the all-reduce
+runs overlapped with compute and at the interconnect's physical limit.
+
+\subsection{Design}
+
+The proposed block is \textbf{PE\_IPCQ} (inter-PE communication queue),
+a small controller dropped into every PE next to PE\_DMA, PE\_GEMM,
+PE\_MATH, and PE\_TCM. It is a control-plane block---it holds no
+payload data---and consists of three pieces: a per-direction
+\emph{QPair register file} (\textasciitilde\SI{576}{\byte} of flip-flops
+covering up to eight directions), a combinational slot-address
+generator and backpressure comparator, and a credit injector/receiver
+wired to the NoC. Each QPair holds the local pointers
+(\texttt{my\_head}, \texttt{my\_tail}), shadowed views of the peer's
+(\texttt{peer\_head\_cache}, \texttt{peer\_tail\_cache}), the local and
+peer rx-buffer physical bases, ring depth and slot size (both
+power-of-two), and the peer's credit-target address. The ring data
+itself lives in a reserved slot region of TCM (or PE-local HBM, or
+cube-shared SRAM), addressed by the QPair registers; PE\_IPCQ never
+touches the bytes, only the pointers. The PE\_CPU sees the controller
+as an MMIO peripheral and the N/S/E/W direction labels as logical
+ports---one kernel image runs across a 1D ring, 2D mesh, or 2D torus
+because the topology only changes which peers the QPair registers
+point at.
+
+\emph{Initialization.} The host CCL backend brings the whole machine
+to a usable state before any kernel runs.
+\texttt{init\_process\_group(backend="ahbm")} loads \texttt{ccl.yaml},
+resolves the algorithm + topology + buffer\_kind + slot configuration,
+allocates an rx ring-buffer region on every participating PE so every
+rank now knows every other rank's \texttt{rx\_base\_pa}, and fans out
+an \texttt{IpcqInitMsg} that writes the QPair register file on each
+PE\_IPCQ over MMIO. The same fan-out wires the per-direction
+credit-return channel: each PE\_IPCQ records its peer's credit-target
+address and a back-pointer that the credit receiver will use to update
+the right QPair. After init, every register the runtime needs is
+preloaded; nothing in the kernel path allocates memory, walks a table,
+or talks to the host.
+
+\emph{Send.} When the kernel executes \texttt{tl.send(dir="E",
+src\_addr, nbytes)}, the PE\_CPU performs a single MMIO write into
+PE\_IPCQ describing the request (direction, source address/space,
+length, sender handle). The controller evaluates backpressure in one
+combinational compare---\texttt{(my\_head $-$ peer\_tail\_cache) $<$
+n\_slots}---and, on a hit, the slot-address generator returns
+\texttt{dst = peer\_rx\_base\_pa + (my\_head \% n\_slots) $\times$
+slot\_size} in one to two cycles. PE\_IPCQ then emits one
+\texttt{IpcqDmaToken} on its dedicated port to PE\_DMA's
+\texttt{vc\_comm} channel, carrying the data descriptor (\texttt{src,
+dst, nbytes}) plus a small piggyback header (\texttt{sender\_seq =
+my\_head, src\_coord, direction}). \texttt{my\_head} is incremented in
+the same cycle---a local flip-flop bump, not a cross-PE atomic---and
+\texttt{tl.send} returns fire-and-forget. If the backpressure compare
+fails, the controller stalls the CPU in either of two modes selected
+at init: \texttt{poll} (CPU re-reads a status CSR) or \texttt{sleep}
+(controller asserts a wake event when a credit arrives). Both modes
+are benchmarked in the results.
+
+\emph{Transport.} PE\_DMA is the only block that touches the fabric,
+and it has been extended for IPCQ in two ways. First, it now exposes
+two virtual channels---\texttt{vc\_compute} for GEMM/Math
+\texttt{TileToken}s and \texttt{vc\_comm} for \texttt{IpcqDmaToken}s
+---with independent state machines and a chunk-level
+(\SI{256}{\byte}) weighted round-robin arbiter on the shared physical
+link. A large GEMM tile DMA can no longer monopolize the link against
+a pending IPCQ send, enabling compute and communication to make
+progress independently as an architectural property of the design. Second, on the sender side PE\_DMA packs the piggyback
+header into the same flit train as the data, and on the receiver side
+it runs the I6 \emph{atomic terminal handler}: pay the per-flit
+bottleneck-BW drain, then, in one indivisible block, (i) write the
+payload into \texttt{MemoryStore} at \texttt{dst\_addr} (the
+receiver's ring slot) and (ii) forward an \texttt{IpcqMetaArrival}
+$\{\texttt{sender\_seq, dst\_addr}\}$ to the local PE\_IPCQ over a
+PE-internal wire. No yield, no PE\_CPU interaction, no cache-coherence
+round trip---the bytes land in TCM and the metadata reaches the
+controller in the same cycle.
+
+\emph{Receive and credit return.} PE\_IPCQ's Meta Extractor
+range-matches the incoming \texttt{dst\_addr} against each direction's
+$[\texttt{rx\_base}, \texttt{rx\_base} + \texttt{n\_slots} \times
+\texttt{slot\_size})$ window---unambiguous even when two directions
+share a peer, as in a 2-rank bidirectional ring---and updates
+\texttt{peer\_head\_cache[d] := max(prev, sender\_seq + 1)}, releasing
+any \texttt{tl.recv} blocked on direction $d$. When the kernel
+eventually consumes the slot, the controller increments
+\texttt{my\_tail} and emits a \SI{16}{\byte} credit packet on the
+dedicated fast-path channel preloaded at init; the latency is the
+real fabric path (per-node overhead + edge propagation + 16 B /
+bottleneck BW), so an in-cube credit returns faster than a cross-SIP
+credit and the model carries no magic constants. The sender's PE\_IPCQ
+absorbs the credit, advances \texttt{peer\_tail\_cache}, and
+de-asserts backpressure if it was stalled. The pointer-synchronization
+problem the software baseline solved with explicit atomic RMWs and
+polling loops has been split in two and dissolved into existing
+traffic: \emph{head} updates ride on the DMA payload itself, with the
+receiver's PE\_DMA performing the data + metadata write in a single
+atomic step, so the sender never blocks waiting for the receiver to
+see it; \emph{tail} updates ride on a dedicated 16 B credit on a
+side-channel, with no software in the loop. Send becomes a single
+MMIO write, receive becomes a flip-flop read, backpressure becomes
+one 64-bit subtract, and the per-message cost in IPCQ is dominated by
+the link traversal rather than by metadata bookkeeping---which is
+what the H2 2025 measurement showed the pure-software queue could
+not achieve. On top of this substrate the collective runs a
+hierarchical local-reduce / global all-reduce-broadcast schedule
+across whatever inter-device topology the configuration specifies,
+with the IPCQ ring buffer placed in on-PE TCM, PE-local HBM, or
+cube-shared SRAM---the third knob the results section sweeps.
+
+\subsection{Design alternatives}
+\label{sec:ipcq-alternatives}
+
+PE\_IPCQ is one point in a small space of hardware mechanisms for
+moving a short message from one PE to a neighbor and signalling its
+arrival. Three established alternatives anchor the space, each the
+HW realization of a familiar host-networking idea: a \emph{doorbell +
+polling} scheme (the classic MMIO doorbell---write the payload by DMA,
+write a doorbell, let the peer poll or take an interrupt); a
+\emph{hardware message queue} (HMQ, the NVLink-style descriptor engine
+that pushes a queue entry to the peer, with large payloads still
+riding a second DMA); and a \emph{completion-queue} design (RDMA-CQ,
+the InfiniBand/RoCE pattern where a DMA write auto-posts a completion
+entry the peer's CQ polls). PE\_IPCQ is the fourth: a hardware ring
+with credit return, splitting the control plane into PE\_IPCQ and the
+data plane into PE\_DMA, with head updates riding the payload and tail
+updates riding a 16\,B side-channel credit (\S\ref{sec:allreduce}).
+
+\begin{figure*}[t]
+\centering
+\includegraphics[width=0.78\linewidth]{ipcq_alternatives_architecture_flow.png}
+\caption{Per-send data and control flow for the four PE-to-PE
+signalling mechanisms (sender\,$\rightarrow$\,NoC\,$\rightarrow$\,receiver).
+Doorbell and RDMA-CQ each issue two fabric transactions (payload then
+doorbell / completion) and leave the peer polling or taking an
+interrupt; HMQ adds a dedicated descriptor engine but still moves large
+payloads on a second DMA; PE\_IPCQ folds head-pointer signalling into
+the payload flit train and returns the tail credit on a side channel,
+so a send is one MMIO write and a receive is a flip-flop read. This is
+a \emph{design schematic}, not a measured comparison.}
+\label{fig:ipcq-arch}
+\end{figure*}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{ipcq_alternatives_decision_matrix.png}
+\caption{Why the ring+credit design was chosen, across five criteria:
+single-send latency, whether the host CPU sits on the critical path,
+whether the receiver must poll or take a wake-up interrupt, whether the
+control and data datapaths are duplicated, and whether the mechanism is
+right-sized for single-owner PE-to-PE traffic (rather than a
+multi-tenant fabric). PE\_IPCQ is the only design that clears every
+criterion. The accompanying per-send step-count tally
+($\sim$28 control events for IPCQ versus $\sim$38 for HMQ, $\sim$53 for
+RDMA-CQ, and $\sim$56 for doorbell+polling) is an \emph{illustrative}
+order-of-magnitude comparator over hand-counted pipeline steps---not a
+simulator measurement. The measured, simulator-grounded results follow
+in the next subsection.}
+\label{fig:ipcq-decision}
+\end{figure}
+
+The qualitative comparison motivates the design but is not a
+quantitative claim: the cycle-step tallies above are hand-counted
+control events, deliberately separated from the measured latencies that
+follow. Everything in the results subsection runs on the PE\_IPCQ
+substrate and is simulator-grounded.
+
+\subsection{Results}
+
+All measurements in this section run on the PE\_IPCQ substrate
+described above; the topology sweep is intended to characterize how
+effectively the proposed mechanism exposes the underlying interconnect's
+properties, not to compare PE\_IPCQ against an alternative
+communication primitive. Following the milestone-evaluation convention,
+the collective sweep builds its own six-device (six-SIP, $2\times3$)
+configurations---distinct from the two-SIP default of
+Table~\ref{tab:hw}---and measures all-reduce latency as a function of
+payload size for three inter-device topologies: a 1D ring, a 2D mesh
+(no wrap), and a 2D torus (Figure~\ref{fig:allreduce-topo}).
+Table~\ref{tab:allreduce} and Figure~\ref{fig:allreduce-cmp} report the
+result.
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{allreduce_topology.png}
+\caption{The three six-device ($2\times3$) inter-device topologies the
+collective sweep runs over, and the hierarchical local-reduce /
+global all-reduce-broadcast schedule mapped onto each: a 1D ring, a 2D
+mesh (no wrap-around), and a 2D torus (wrap-around links on both axes).
+The torus's wrap links shorten the worst-case reduction path, which is
+what the latency sweep below rewards.}
+\label{fig:allreduce-topo}
+\end{figure}
+
+\begin{table}[t]
+\centering
+\caption{All-reduce latency (ns) across six devices, by topology and
+per-PE payload. Lower is better; the torus wins at every size.}
+\label{tab:allreduce}
+\small
+\begin{tabular}{@{}rrrr@{}}
+\toprule
+\textbf{Bytes/PE} & \textbf{2D mesh} & \textbf{Ring 1D} & \textbf{2D torus} \\
+\midrule
+256     & 4189  & 3883  & 2957 \\
+4{,}096  & 5566  & 5240  & 4031 \\
+16{,}384 & 10016 & 9376  & 7396 \\
+65{,}536 & 27821 & 25925 & 20858 \\
+98{,}304 & 39690 & 36957 & 29833 \\
+\bottomrule
+\end{tabular}
+\end{table}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{allreduce_comparison.png}
+\caption{All-reduce latency vs.\ per-PE payload for the three PE\_IPCQ
+topologies, against the analytic torus model. The isolated point in
+the top panel (\SI{366}{\micro\second}) is the only data point
+available from the H2 2025 software-queue measurement campaign---an
+intra-device all-reduce across 16 CUBEs on a single device. A true
+6-device inter-SIP measurement under the same software queue was not
+collected, so this point in fact \emph{understates} the SW-queue cost
+for the multidevice configuration KernBench measures here; even so the
+proposed PE\_IPCQ inter-device curves sit roughly an order of
+magnitude below it, which is the headline SW-vs-HW comparison this
+section makes. The measured torus tracks the analytic curve within a
+small constant factor, the gap reflecting real link serialization that
+the analytic model idealizes away.}
+\label{fig:allreduce-cmp}
+\end{figure}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{allreduce_buffer_kind.png}
+\caption{Effect of IPCQ staging-buffer placement (2D torus). At
+\SI{64}{\kibi\byte}/PE, TCM staging (\SI{20858}{\nano\second}) beats HBM
+(\SI{24074}{\nano\second}) by \textasciitilde\SI{13}{\percent} and SRAM
+(\SI{33194}{\nano\second}) by \textasciitilde\SI{37}{\percent}; at small
+payloads the three are indistinguishable.}
+\label{fig:allreduce-buf}
+\end{figure}
+
+Three findings follow. First, topology matters and the effect grows with
+size: at \SI{96}{\kilo\byte}/PE the 2D torus is \SI{25}{\percent} faster
+than the mesh and \SI{19}{\percent} faster than the ring, because its
+wrap-around links shorten the worst-case reduction path. Second, the
+measured curves grow smoothly with payload and stay within a small
+constant factor of the analytic torus model, with the gap reflecting real
+link serialization that the analytic model idealizes away. Third, where
+the IPCQ staging buffer lives is a first-order knob: keeping it in on-PE
+TCM is materially faster than HBM or shared SRAM at large payloads
+(Figure~\ref{fig:allreduce-buf}).
+
+\subsection{Analysis and meaning}
+
+PE\_IPCQ turns the collective from an off-device, contention-prone
+operation into an on-device primitive whose latency tracks the
+interconnect's physical limits. The control/data split keeps the engine
+small while reusing PE\_DMA for movement, and the virtual-channel split is
+what lets a reduction proceed without stalling the compute DMA that feeds
+the very GEMMs producing the partials---the same property the fused
+attention kernel relies on when it interleaves KV reduction with score
+computation (\S\ref{sec:gqa}). For the hardware roadmap the results argue
+two things: provisioning wrap-around (torus) inter-device links is worth
+roughly a \SI{20}{}--\SI{25}{\percent} collective-latency reduction at
+scale, and giving the IPCQ a fast on-PE staging buffer (TCM) rather than
+forcing it through HBM or SRAM is worth another \SI{14}{}--%
+\SI{38}{\percent} at large payloads.

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

@@ -0,0 +1,413 @@
+\section{Fused Grouped-Query Attention}
+\label{sec:gqa}
+
+Attention is the 1H focus, and it is where the two preceding optimizations
+have to come together. Grouped-Query Attention (GQA) shrinks the KV cache
+by sharing each KV head across a group of query heads (here $h_q=8$ query
+heads to $h_{kv}=1$ KV head, a group factor $G=8$), which makes decoding
+feasible at long context but also makes it acutely memory-bound: a decode
+step processes a single query position ($T_q=1$) against the entire KV
+history, so its arithmetic intensity is low and its time is dominated by
+streaming the KV cache out of HBM. FlashAttention-style tiling with an
+online-softmax merge avoids ever materializing the full score matrix, but
+realizing it as a fast \emph{fused} kernel needs both building blocks from
+this report: efficient GEMM issue (\S\ref{sec:gemm}) for the
+$Q\!\cdot\!K^{\top}$ and $P\!\cdot\!V$ products, and an efficient on-device
+reduction (\S\ref{sec:allreduce}) for the multi-user and
+sequence-parallel KV reductions. \emph{Fused} here is meant in the
+FlashAttention sense---$Q\!\cdot\!K^{\top}$, the online softmax, and
+$P\!\cdot\!V$ collapse into a single kernel that never materializes the
+score matrix---and, beyond that, the cross-device KV reduction is absorbed
+into the same kernel (on PE\_IPCQ) rather than issued as a separate
+all-reduce. This section is the capstone: the fused
+kernel that uses the composite command and PE\_IPCQ at the same time.
+Multi-head attention (MHA) was studied in prior work and serves here as
+the established baseline rather than being re-derived.
+
+\subsection{Data Placement Policy}
+\label{sec:gqa-placement}
+
+Long-context decode is bound by the KV cache, so the first-order design
+question is how to place that cache---and the running softmax state it
+feeds---across the two hardware axes the machine exposes: the CUBEs and,
+within each CUBE, the PEs (here $C{=}8$ CUBEs $\times$ $P{=}8$ PEs, for
+$C\!\cdot\!P{=}64$ attention engines over one KV-head group on the
+LLaMA-3.1-70B target). Each axis can \emph{replicate} the KV cache or
+\emph{shard} it, and a shard can run along the sequence dimension
+$S_{kv}$ or the head dimension $d_{\text{head}}$. The cross product is a
+small, enumerable taxonomy; six placements span its meaningful corners
+(Figure~\ref{fig:gqa-kv-sharding}).
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_long_ctx_6cases_kv_sharding_diagram.png}
+\caption{The six KV-placement strategies, drawn on the
+$S_{kv}\!\times\!d_{\text{head}}$ KV tensor (rows = sequence, columns =
+head dimension). Cube colour bands and dashed PE dividers show which
+axis each level shards. Cases~1--3 either replicate the cache or shard
+it on a single axis (8-way at most); Cases~4--6 reach a full 64-way
+split, three different ways: Case~4 splits $S_{kv}$ across CUBEs and
+$d_{\text{head}}$ across PEs, Case~5 the mirror, and Case~6~$\star$
+splits $S_{kv}$ on \emph{both} axes.}
+\label{fig:gqa-kv-sharding}
+\end{figure}
+
+Two quantities decide which placement is viable, and they pull against
+each other. The first is \textbf{per-PE KV memory}. With a per-PE HBM
+budget of \SI{6.0}{\giga\byte} and \SI{1.76}{\giga\byte} of attention
+weights resident, the KV headroom is \SI{4.24}{\giga\byte} per PE. At a
+production context of $S_{kv}{=}1\,\text{M}$ tokens the unsharded cache
+is \SI{40}{\giga\byte}/PE (Case~1), an 8-way shard is \SI{5}{\giga\byte}
+(Cases~2--3)---both \emph{over} the headroom---while only the 64-way
+placements bring it to \SI{640}{\mega\byte}/PE (Cases~4--6), comfortably
+inside budget. Memory alone therefore eliminates Cases~1--3 at long
+context. The second quantity is \textbf{communication per token}, and it
+is what separates the three survivors. Sharding $d_{\text{head}}$
+(Cases~4--5) makes each PE hold only a slice of every head, so the
+$Q\!\cdot\!K^{\top}$ score is \emph{partial} and must be all-reduced
+across the slice owners on every token---a reduction whose volume scales
+with $S_{kv}$ ($\sim$\SI{166}{\mega\byte}/token analytically, intra-CUBE
+on the NoC for Case~4, inter-CUBE on UCIe for Case~5). Case~6~$\star$
+instead shards $S_{kv}$ on both axes, so every PE computes a
+\emph{complete} score over its own token range and only the small running
+softmax state $(m,\ell,O)$ is merged across PEs
+($\sim$\SI{6.2}{\mega\byte}/token)---a $\sim$27$\times$ lighter collective
+than the $d_{\text{head}}$-split designs.
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_long_ctx_6cases_summary.png}
+\caption{Long-context placement analysis at $S_{kv}{=}1\,\text{M}$ tokens.
+\emph{Left:} the per-PE HBM budget---\SI{1.76}{\giga\byte} of attention
+weights leave \SI{4.24}{\giga\byte} of KV headroom (red line).
+\emph{Middle:} per-PE KV memory per case (log scale); only the 64-way
+placements (Cases~4--6, \SI{640}{\mega\byte}) clear the headroom, while
+the unsharded (\SI{40}{\giga\byte}) and 8-way (\SI{5}{\giga\byte}) cases
+overflow. \emph{Right:} analytical communication per token (log scale);
+the $d_{\text{head}}$-split Cases~4--5 pay a partial-score all-reduce
+($\sim$\SI{166}{\mega\byte}/token) that the both-axes-$S_{kv}$ split of
+Case~6~$\star$ avoids ($\sim$\SI{6.2}{\mega\byte}/token, merging only the
+softmax state). On these two axes Case~6 (marked $\star$ in the figure)
+is the single placement that lands both inside the memory budget and at
+low per-token communication; whether that combination is the right one to
+pick is a regime-specific question taken up next.}
+\label{fig:gqa-budget}
+\end{figure}
+
+These two costs---per-PE KV memory and per-token communication---are
+intrinsic properties of each placement, fixed by how it shards the cache
+and independent of the workload regime. They do not by themselves name a
+winner: a short prompt where the whole cache fits on one PE values low
+communication and tolerates replication, whereas a million-token decode
+is bound by the memory wall and will pay communication to escape it.
+Which placement is appropriate is therefore a per-regime question, which
+the short- and long-context subsections that follow answer by running the
+options on the simulator and reading off latency, traffic, and the
+redundant compute each one induces.
+
+\subsection{Inference with Short-Context Length}
+\label{sec:gqa-short}
+
+The fused GQA kernel issues its matrix products as scheduler-managed
+composite commands and keeps the online-softmax merge and the cross-device
+KV reduction inside the kernel, on PE\_IPCQ. Two kernel families cover the
+two phases. The \emph{prefill} kernel is head-parallel and rotates the KV
+shards around an inter-CUBE ring (``Ring KV''). The \emph{decode} kernel
+is head-replicated with a statically sharded KV cache and reduces partial
+attention outputs through an M-fold intra-CUBE chain and, for multiple
+users, a two-level reduce-to-root. Two further primitives make long
+context practical: a \emph{lazy load} that issues the KV \textsf{DMA\_READ}
+and returns immediately, auto-waiting only at first use so that KV load
+overlaps score computation; and per-tile \emph{scratch recycling} that
+keeps the running softmax accumulators ($m,\ell,O$) in a persistent arena
+while freeing per-tile temporaries, so the kernel fits the
+\SI{1}{\mebi\byte} scratch budget across many tiles. A further refinement
+restructures the decode step so the per-tile matrix products and the
+online-softmax merge are issued as \emph{composite} commands rather than
+hand-tiled primitives; \S\ref{sec:gqa-composite} measures it against the
+primitive baseline at long context.
+
+% TODO: CUBE <-> KV-head mapping diagram for the short-context regime
+% (h_kv=8 KV heads -> 8 CUBEs, 1:1; intra-CUBE PE usage).
+% Bench code: src/kernbench/benches/gqa_helpers/short_ctx/
+
+% TODO: prefill performance figure (latency, stage breakdown).
+% TODO: decode performance figure (latency, stage breakdown).
+%       Bench output for short_ctx to be generated.
+
+\subsection{Inference with Long-Context Length}
+\label{sec:gqa-long}
+
+% TODO: prefill long-context kernel implementation description
+% (Sequence-Parallel partition of S_kv, per-case mechanics).
+% Bench code: src/kernbench/benches/gqa_helpers/long_ctx/
+
+% TODO: prefill long-context performance figure.
+
+Long-context decode is the regime where the KV cache, not attention
+compute, sets serving cost, so the placement question of
+\S\ref{sec:gqa-placement} becomes decisive here. To pick the right
+placement for this regime we run each of the six options as a fused
+decode kernel on the simulator (one decode step on the LLaMA-3.1-70B
+single-KV-head-group target, $C{=}8$ CUBEs $\times$ $P{=}8$ PEs) at a
+tractable $S_{kv}{=}8192$, and read off end-to-end latency, on-device op
+traffic, and the redundant compute each one induces. The swept context is
+small enough that all six fit in memory at $S_{kv}{=}8192$; the
+placements the long-context memory budget rules out (Cases~1--3) are
+drawn in red, run here only to expose their issue and communication
+structure.
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_latency.png}
+\caption{Measured end-to-end decode latency per placement
+($S_{kv}{=}8192$; red = ruled out by the long-context memory budget,
+blue = the predicted Pareto choice). The fastest raw latency belongs to
+Case~3 (\SI{17.8}{\micro\second})---but Case~3 replicates the full KV
+cache into every CUBE, so it overflows the per-PE budget at production
+context and wastes 8$\times$ the compute
+(Figure~\ref{fig:gqa-6cases-par}). Among the placements that actually fit
+1\,M-token memory (Cases~4--6), Case~6~$\star$ is the fastest
+(\SI{30.6}{\micro\second}, versus \SI{31.4}{} and \SI{34.5}{\micro\second}
+for the $d_{\text{head}}$-split Cases~4 and~5)---making it the placement
+of choice for long-context decode.}
+\label{fig:gqa-6cases-lat}
+\end{figure}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_parallelism.png}
+\caption{Redundant compute per placement, measured as active-PE
+$\times$ $S_{\text{local}}$ (PE-tokens; lower means less wasted work).
+The minimum is \num{8192} PE-tokens---one pass over the sequence.
+Case~3 inflates this 8$\times$ to \num{65536} by replicating the KV
+cache across all eight CUBEs so every CUBE redundantly re-attends the
+whole sequence; the $d_{\text{head}}$-split Cases~4--5 likewise carry
+\num{65536} because each token is processed across eight head slices.
+Case~6~$\star$ achieves the full 64-way split at the minimal
+\num{8192} PE-tokens---fully parallel, no replication.}
+\label{fig:gqa-6cases-par}
+\end{figure}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_traffic.png}
+\caption{Measured on-device op traffic per placement. The unsharded
+Case~1 issues no IPCQ copies (each PE has the full cache, nothing to
+reduce); the single-axis Case~3 charges 168. Among the 64-way splits,
+the $d_{\text{head}}$-split Cases~4--5 charge the most---280 IPCQ copies
+and 8 DMA writes each, the partial-score all-reduce that head-slicing
+forces---while Case~6~$\star$ needs only 189 IPCQ copies and a single DMA
+write, because it merges just the running softmax state $(m,\ell,O)$
+rather than partial scores. This is the on-device collective traffic that
+PE\_IPCQ and the torus links of \S\ref{sec:allreduce} are provisioned to
+absorb at link speed.}
+\label{fig:gqa-6cases-traffic}
+\end{figure}
+
+The measurements select the right placement for this regime. The fastest
+raw latency (Case~3, \SI{17.8}{\micro\second}) comes from replicating the
+full KV cache into every CUBE---which is exactly the placement the
+long-context memory budget forbids, and which the parallelism panel shows
+wastes 8$\times$ the compute. Restricting attention to the placements that
+fit production-context memory (the 64-way splits, Cases~4--6), the choice
+is Case~6~$\star$: it is the fastest of the three
+(\SI{30.6}{\micro\second}), and the op-count panel shows why---its
+softmax-state-only reduction charges 189 IPCQ copies and one DMA write
+against the 280 copies and 8 DMA writes the $d_{\text{head}}$-split
+Cases~4--5 pay for their partial-score all-reduce. For long-context
+decode, then, the appropriate data placement is the both-axes sequence
+shard (Case~6): it is the cheapest-communicating member of the only
+memory-feasible family, and the cross-PE softmax reduction it does pay is
+precisely the traffic the communication-side codesign of this report is
+built to move quickly.
+
+\subsection{Use of Composite Commands}
+\label{sec:gqa-composite}
+
+The decode kernel of \S\ref{sec:gqa-long} issues its local attention as
+primitive operations: it walks each PE's $S_{\text{local}}$ token slice in
+\SI{1024}{}-token tiles, and for every tile issues a $Q\!\cdot\!K^{\top}$
+\textsf{dot}, the online-softmax primitives, and a $P\!\cdot\!V$
+\textsf{dot}, merging the running $(m,\ell,O)$ state by hand. The number
+of PE\_CPU commands this costs grows with the context. At a production
+context of $S_{kv}{=}1\,\text{M}$ tokens the Case-6 64-way split gives
+each PE $S_{\text{local}}{=}16384$ tokens, so its local attention is a
+$Q\!\cdot\!K^{\top}$ of $(8,128)\!\cdot\!(128,16384)$ and a
+$P\!\cdot\!V$ of $(8,16384)\!\cdot\!(16384,128)$---sixteen hand-issued
+tiles, each a fresh batch of CPU commands.
+
+The composite command lets the kernel hand that tiling to PE\_SCHEDULER.
+We compare three command forms of the \emph{same} Case-6 kernel---identical
+placement and identical $(m,\ell,O)$ reduce, differing only in how the
+local attention is issued:
+\begin{itemize}
+  \item \textbf{primitive}---the hand-tiled \textsf{dot}/softmax kernel
+    above (the baseline of \S\ref{sec:gqa-long}).
+  \item \textbf{composite}---each matrix product is one coarse
+    \textsf{composite} GEMM over the \emph{whole} $S_{\text{local}}$, with
+    $K$ and $V$ passed as HBM references so PE\_SCHEDULER streams and
+    tiles them on the fixed $32\!\times\!64\!\times\!32$ MAC tile; the
+    softmax stays primitive.
+  \item \textbf{composite\,+\,softmax\_merge}---additionally folds the
+    online-softmax merge and $P\!\cdot\!V$ into a single stateful
+    \textsf{softmax\_merge} recipe composite.
+\end{itemize}
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_decode_long_ctx_composite.png}
+\caption{Three command forms of the Case-6 decode kernel, swept over
+context length ($S_{\text{local}}{=}S_{kv}/64$ per PE). \emph{Right:}
+PE\_CPU commands issued. The hand-tiled primitive kernel rises
+$O(n_{\text{tiles}})$---from 96 commands at one tile to 426 at the
+1\,M-token, sixteen-tile production point---while both composite forms
+issue a context-\emph{independent} $O(1)$ count (94 and 98) that
+\emph{saturates}: one coarse descriptor offloads the entire per-tile
+fan-out. \emph{Left:} the consequence for wall-clock latency is none---all
+three land on the same curve (\SI{30.6}{}, \SI{231}{},
+\SI{461}{\micro\second} at 8\,K\,/\,64\,K\,/\,128\,K), because decode is
+bound by streaming the KV cache, not by issue. Command-count is measured
+at emit time (exact, to 1\,M); latency on the data-mode engine over the
+tractable range.}
+\label{fig:gqa-composite}
+\end{figure}
+
+Figure~\ref{fig:gqa-composite} reads off the two quantities that matter,
+and they point in opposite directions. The PE\_CPU command count (right)
+collapses from a context-growing $O(n_{\text{tiles}})$ to a flat $O(1)$:
+at 1\,M tokens the composite form issues \num{94} commands against the
+primitive kernel's \num{426}, a $4.5\times$ reduction
+($4.2\times$ in modeled dispatch cost), and---crucially---that number no
+longer grows with context. The wall-clock latency (left), by contrast,
+is unchanged across all three forms: decode is bound by streaming the KV
+cache out of HBM, so the command form does not move the critical path.
+
+That juxtaposition is the point. The composite command is not a latency
+optimization for this memory-bound decode \emph{at the 64-way production scale} (a
+single-rank caveat follows, Figure~\ref{fig:gqa-decode-stream}); it is a
+\emph{CPU-issue}
+optimization. Its value is removing the per-tile dispatch work that would
+otherwise grow without bound as context grows, freeing PE\_CPU to run
+ahead and keep the engines fed---which is exactly what lets the
+data-movement cost analyzed next show through as the true bottleneck
+rather than being masked by issue overhead. The \textsf{softmax\_merge}
+recipe folds the online merge into the same descriptor; on this
+memory-bound path its marginal cost over the plain GEMM composite is small
+(98 vs.\ 94 commands), and like the plain composite it keeps the issued
+count flat as context scales.
+
+\paragraph{Isolating the rank: the masked streaming win.} That neutrality
+is a property of the \emph{full 64-way} critical path, not of the local
+attention: at production scale the inter-CUBE $(m,\ell,O)$ reduce tail and
+shared-HBM contention set the wall clock, so a faster local attention does
+not surface. Stripping those away---a single rank, no cross-CUBE reduce,
+swept over the per-rank context $S_{kv}$ (so $S_{kv}{=}16$\,K here is the
+per-PE load of a 1M-token, 64-way-sharded decode)---exposes the local
+attention directly (Figure~\ref{fig:gqa-decode-stream}), and the composite
+\emph{does} win, by \SI{25}{}--\SI{28}{\percent}. The reason is the
+memory-bound mirror of prefill: its scheduler-streamed concurrent per-tile
+DMAs keep the HBM pipeline full and reach \SI{233}{\giga\byte\per\second}
+---\SI{91}{\percent} of the per-rank \SI{256}{\giga\byte\per\second}
+roofline---whereas the primitive kernel's blocking \textsf{tl.dot}
+serializes one tile DMA at a time and plateaus at
+\SI{166}{\giga\byte\per\second}. So even for memory-bound decode the
+composite is not \emph{only} a CPU-issue optimization---it also extracts
+bandwidth---but that latency benefit materializes only when the local
+attention is on the critical path, which at 64-way production scale it is
+not.
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_decode_streaming.png}
+\caption{Single-rank memory-bound decode ($T_q{=}1$, $M{=}8$), three
+command forms, swept over per-rank context. \emph{Left:} end-to-end
+latency---the composite forms run \SI{25}{}--\SI{28}{\percent} below the
+primitive, a gap that widens with context. \emph{Right:} achieved HBM
+bandwidth against the per-rank \SI{256}{\giga\byte\per\second} roofline.
+The primitive's blocking load$\rightarrow$dot serializes the KV stream and
+plateaus at \SI{166}{\giga\byte\per\second}; the composite forms pipeline
+concurrent per-tile DMAs through the scheduler and reach
+\SI{233}{\giga\byte\per\second}. This is the memory-bound mirror of the
+prefill result (Figure~\ref{fig:gqa-prefill-cb}): there the composite
+approaches the MAC roofline, here the bandwidth roofline. Capped at
+$16$\,K---the plain composite materializes the full $(M,S_{kv})$ scores in
+TCM, so beyond that only the \textsf{softmax\_merge} recipe, which tiles
+the softmax, stays within scratch.}
+\label{fig:gqa-decode-stream}
+\end{figure}
+
+\paragraph{The compute-bound mirror: prefill.} Decode's \emph{production}
+verdict---command form is latency-neutral at 64-way scale---is a property
+of its regime, not of the
+composite command. A decode step has $T_q{=}1$, so its score and context
+products are skinny ($M{=}G\,T_q{=}8$): the MAC array is barely fed and
+the kernel is bound by streaming the KV cache. Prefill is the opposite
+corner. It processes a block of query positions at once, so $M{=}G\,T_q$
+is large and tile-filling, the GEMMs carry real arithmetic intensity
+($\sim$$M$ flops/byte, well above the roofline ridge), and the kernel is
+\emph{compute-bound}. This is the regime the composite command was built
+for (\S\ref{sec:gemm}): it streams the per-HW-tile
+DMA$\rightleftarrows$compute pipeline so the MAC array stays fed, whereas
+the primitive kernel's blocking \textsf{tl.dot} serializes each tile's
+load and compute and starves the array between tiles. We run the same
+three command forms on a single-rank compute-bound prefill (FlashAttention
+$Q$-block $\times$ $S_{kv}$-tile, online softmax) and sweep the context
+length (Figure~\ref{fig:gqa-prefill-cb}).
+
+\begin{figure}[t]
+\centering
+\includegraphics[width=\linewidth]{gqa_prefill_compute_bound.png}
+\caption{Compute-bound prefill, three command forms, swept over context
+length ($M{=}8\,T_q$ tile-filling). \emph{Left:} end-to-end latency.
+\emph{Right:} MAC utilization (achieved $\div$ the
+\SI{8}{\tera\flop\per\second} per-PE peak). The hand-tiled primitive sits
+flat at $\sim$\SI{68}{\percent}---its serial load$\to$dot path leaves the
+MAC array idle between tiles regardless of context. The composite forms
+climb with context (\SI{67}{}$\to$\SI{80}{\percent} plain,
+\SI{67}{}$\to$\SI{83}{\percent} with the recipe) because a deeper $P\!\cdot
+\!V$ reduction gives more HW tiles to pipeline, and they convert that into
+wall-clock: at \num{1024} the recipe form is \SI{646.9}{} vs.\
+\SI{794.1}{\micro\second} (\SI{19}{\percent} faster). The margin
+\emph{grows} with context---the compute-bound mirror of the GEMM result of
+\S\ref{sec:gemm}.}
+\label{fig:gqa-prefill-cb}
+\end{figure}
+
+The two studies together state the composite command's value precisely. It
+has two distinct benefits, and which one matters is set by the workload's
+roofline position. The first is \emph{host-issue offload}: one macro
+command in place of $O(n_{\text{tiles}})$ fine ones, which removes
+PE\_CPU dispatch work and is regime-independent (it shows in the decode
+command count). The second is \emph{MAC-array feeding}: the
+scheduler-internal per-tile DMA$\rightleftarrows$compute pipeline, which
+only converts to latency when the workload is compute-bound enough to have
+a MAC array worth keeping busy (it shows in the prefill utilization). A
+memory-bound decode exercises only the first; a compute-bound prefill
+exercises both. The composite command is the single mechanism that
+delivers each where it applies.
+
+\subsection{Comprehensive Analysis}
+\label{sec:gqa-analysis}
+
+These panels are the clearest statement of the codesign thesis in the
+report. Because the composite command keeps GEMM issue cheap and the MAC
+array barely occupied, the fused attention kernel's latency is set almost
+entirely by data movement: streaming the KV cache and reducing partials
+across devices. That is precisely the cost that the communication-side
+work targets---PE\_IPCQ for the on-device reduction, the lazy load for
+load/compute overlap, fast TCM staging and torus links for the reduction
+itself. In other words, the two enablers are not independent features that
+happen to appear in the same kernel; the GEMM optimization is what
+\emph{exposes} the data-movement bottleneck (by removing the compute and
+issue overhead that would otherwise hide it), and the communication
+optimization is what \emph{attacks} it. For an attention-dominated decoder
+the meaningful hardware investments are therefore the ones that move data
+faster and reduce it on-device---not additional MAC throughput, which this
+workload cannot use.
+
+% TODO: cross-regime DP (data parallelism) applicability:
+%   - Does Case-4 long-context placement compose with batch-level DP
+%     without further changes?
+%   - Does the short-context placement compose the same way?
+%   - Implications for multi-user serving (single vs. mixed regimes).

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

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

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

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

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

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

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

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

scratch_dot_tiling_experiment.py

@@ -0,0 +1,204 @@
+#!/usr/bin/env python3
+"""SCRATCH EXPERIMENT (not production; do not commit).
+
+Question: does charging the *primitive* decode kernel for per-HW-tile
+(16x16x16) CPU dispatch flip the "composite gives no decode-latency
+benefit" conclusion?
+
+We monkeypatch TLContext.dot so that, in the primitive kernel, every
+tl.dot whose (M,K,N) exceeds the HW GEMM tile (mac_m/mac_k/mac_n) is
+split by the CPU into ceil(M/mac_m)*ceil(K/mac_k)*ceil(N/mac_n)
+HW-tile-sized GemmCmds.  Each tile GemmCmd is emitted through the normal
+_emit() path, so it (a) charges PE_CPU dispatch overhead via
+_charge_dispatch (PeCpuOverheadCmd), and (b) blocks like a normal
+single-op GemmCmd on PE_GEMM at the cycle-accurate ceil-product latency.
+
+We also inject mac_m/mac_k/mac_n into every pe_gemm topology node so BOTH
+the primitive-tiled and the composite variants run on the *same*
+cycle-accurate engine (fair comparison).  Composite is left untouched:
+the CPU emits ONE CompositeCmd, and PE_SCHEDULER tiles internally (no
+per-HW-tile CPU dispatch).
+
+Data correctness:
+  Inputs are ctx.zeros (q/k/v), so every matmul result is zeros and the
+  DataExecutor replay is trivial.  To keep replay numerically correct
+  regardless, exactly ONE emitted tile per dot carries the *real* full
+  operands+output handles (so the DataExecutor computes the true (M,N)
+  result via the recorded handle shapes), while its timing fields
+  (m,k,n) are the HW-tile size so the engine charges exactly one tile of
+  cycle time.  The remaining n_tiles-1 emitted tiles are timing-only
+  GemmCmds (16x16x16) writing to throwaway scratch.  Net: n_tiles tiles
+  of engine time + n_tiles dispatch charges, and a correct final output.
+
+Engine mode: enable_data=True (same as the production sweep's
+_engine_latency_ns), op_log end-to-end latency.
+"""
+from __future__ import annotations
+
+import json
+from math import ceil
+from pathlib import Path
+
+# --- HW GEMM tile under test -------------------------------------------------
+MAC_M = 16
+MAC_K = 16
+MAC_N = 16
+# Legacy alt to also try: (8, 16, 32)
+
+S_KV_LATENCY = (8192, 32_768, 65_536, 131_072)
+
+ROOT = Path(__file__).resolve().parent
+SWEEP_JSON = (
+    ROOT / "src" / "kernbench" / "benches" / "1H_milestone_output"
+    / "gqa" / "long_ctx" / "sweep_decode_composite.json"
+)
+
+
+# --- mac-dim topology override -----------------------------------------------
+def _topo_with_mac(mac_m: int, mac_k: int, mac_n: int):
+    """Compiled topology with mac dims injected into every pe_gemm node."""
+    from kernbench.topology.builder import resolve_topology
+
+    handle = resolve_topology("topology.yaml")
+    g = handle.topology_obj
+    n = 0
+    for node in g.nodes.values():
+        if node.kind == "pe_gemm":
+            node.attrs["mac_m"] = mac_m
+            node.attrs["mac_k"] = mac_k
+            node.attrs["mac_n"] = mac_n
+            n += 1
+    print(f"    injected mac=({mac_m},{mac_k},{mac_n}) into {n} pe_gemm nodes")
+    return handle
+
+
+# --- tiling monkeypatch for TLContext.dot ------------------------------------
+def _make_tiled_dot(orig_dot, mac_m: int, mac_k: int, mac_n: int):
+    from kernbench.common.pe_commands import GemmCmd
+
+    def tiled_dot(self, a, b):
+        if len(a.shape) < 2 or len(b.shape) < 2:
+            return orig_dot(self, a, b)
+        m, k = a.shape[-2], a.shape[-1]
+        k2, n = b.shape[-2], b.shape[-1]
+        if k != k2:
+            raise ValueError(f"dot shape mismatch: a.K={k} != b.K={k2}")
+
+        n_tiles = ceil(m / mac_m) * ceil(k / mac_k) * ceil(n / mac_n)
+
+        out_shape = (*a.shape[:-2], m, n)
+        out = self._make_compute_out(shape=out_shape, dtype=a.dtype)
+        self._await_pending(a, b)
+
+        if n_tiles <= 1:
+            self._emit(GemmCmd(a=a, b=b, out=out, m=m, k=k, n=n))
+            return out
+
+        # One real-data tile: full handles (so DataExecutor computes the
+        # true result), but timing fields = HW tile (one tile of cycles).
+        self._emit(GemmCmd(a=a, b=b, out=out, m=mac_m, k=mac_k, n=mac_n))
+        # Remaining timing-only tiles: throwaway scratch, 16x16x16.
+        scratch = self._make_compute_out(shape=(mac_m, mac_n), dtype=a.dtype)
+        for _ in range(n_tiles - 1):
+            self._emit(GemmCmd(a=a, b=b, out=scratch,
+                               m=mac_m, k=mac_k, n=mac_n))
+        return out
+
+    return tiled_dot
+
+
+# --- latency runner (replicates sweep's _engine_latency_ns) ------------------
+def _engine_latency_ns(variant: str, S_kv: int, topo) -> float:
+    from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_composite import (  # noqa: E501
+        _end_to_end_ns, _run_panel_fn,
+    )
+    from kernbench.runtime_api.bench_runner import run_bench
+    from kernbench.runtime_api.types import resolve_device
+    from kernbench.sim_engine.engine import GraphEngine
+
+    result = run_bench(
+        topology=topo, bench_fn=_run_panel_fn(variant, S_kv),
+        device=resolve_device(None),
+        engine_factory=lambda t, d: GraphEngine(
+            getattr(t, "topology_obj", t), enable_data=True,
+        ),
+    )
+    if not result.completion.ok:
+        raise RuntimeError(
+            f"{variant}@{S_kv} failed: {result.completion}"
+        )
+    return _end_to_end_ns(result.engine.op_log)
+
+
+def _emit_dispatch(variant: str, S_kv: int) -> int:
+    """PE_CPU command count at the center rank (cube 6, pe 0)."""
+    from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_composite import (  # noqa: E501
+        _emit_dispatch as prod_emit,
+    )
+    return prod_emit(variant, S_kv)[0]
+
+
+def main() -> None:
+    import kernbench.triton_emu.tl_context as tlc
+
+    # Baseline (A): primitive UNTILED latencies from the production sweep
+    # (mac=0 / TFLOPS model). Read straight off the committed sweep JSON.
+    sweep = json.loads(SWEEP_JSON.read_text())
+    base_A = {}
+    for r in sweep["rows"]:
+        if r["variant"] == "primitive" and r["latency_ns"] is not None:
+            base_A[r["S_kv"]] = r["latency_ns"]
+
+    print(f"== mac tile = ({MAC_M},{MAC_K},{MAC_N}) ==")
+
+    # --- command-count sanity (emit-time, mac-independent) ---------------
+    orig_dot = tlc.TLContext.dot
+    print("\n[dispatch counts @ S_kv=131072]")
+    n_prim_untiled = _emit_dispatch("primitive", 131072)
+    tlc.TLContext.dot = _make_tiled_dot(orig_dot, MAC_M, MAC_K, MAC_N)
+    try:
+        n_prim_tiled = _emit_dispatch("primitive", 131072)
+        n_comp = None
+    finally:
+        tlc.TLContext.dot = orig_dot
+    n_comp = _emit_dispatch("composite", 131072)
+    print(f"  primitive UNTILED PE_CPU cmds : {n_prim_untiled}")
+    print(f"  primitive TILED   PE_CPU cmds : {n_prim_tiled}  "
+          f"(x{n_prim_tiled / max(n_prim_untiled,1):.0f})")
+    print(f"  composite         PE_CPU cmds : {n_comp}")
+
+    # --- latency sweep ----------------------------------------------------
+    rows = []
+    topo = _topo_with_mac(MAC_M, MAC_K, MAC_N)
+
+    for S_kv in S_KV_LATENCY:
+        A = base_A.get(S_kv)
+
+        # (C) composite on the mac engine (untouched dot path)
+        C = _engine_latency_ns("composite", S_kv, topo)
+
+        # (B) primitive TILED on the mac engine
+        tlc.TLContext.dot = _make_tiled_dot(orig_dot, MAC_M, MAC_K, MAC_N)
+        try:
+            B = _engine_latency_ns("primitive", S_kv, topo)
+        finally:
+            tlc.TLContext.dot = orig_dot
+
+        gap_pct = (B - C) / C * 100.0 if C else float("nan")
+        rows.append((S_kv, A, B, C, gap_pct))
+        print(f"  S_kv={S_kv:>7}: A(untiled)={A!s:>12}  "
+              f"B(tiled)={B:12.2f}  C(comp)={C:12.2f}  (B-C)/C={gap_pct:+6.1f}%")
+
+    # --- final table ------------------------------------------------------
+    print("\n==================== RESULT TABLE ====================")
+    print(f"{'S_kv':>8} | {'A untiled(ns)':>14} | {'B tiled(ns)':>14} | "
+          f"{'C comp(ns)':>14} | {'(B-C)/C':>9}")
+    print("-" * 72)
+    for S_kv, A, B, C, gap in rows:
+        a_s = f"{A:.2f}" if A is not None else "n/a"
+        print(f"{S_kv:>8} | {a_s:>14} | {B:>14.2f} | {C:>14.2f} | "
+              f"{gap:>+8.1f}%")
+
+
+if __name__ == "__main__":
+    main()

scripts/paper/measure_gqa_decode_placement_comm.py

@@ -0,0 +1,362 @@
+"""Measured comm overlay for all 6 GQA decode KV placements.
+
+Runs the simulator for Cases 1-6 (the same 6 placements the chart in
+paper_plot_gqa_4cases_summary.py covers analytically) at S_kv = 64 K,
+sums actual IPCQ-copy bytes from the engine op_log, projects to
+per-token (x80 layers), scales the partial-score-AR component of
+Cases 4/5 from S_kv = 64 K -> S_kv = 1 M (linear in S_kv; other
+cases are S_kv-independent), adds the constant Wo + FFN AR
+(1.25 MB / token), and writes the result to JSON for
+paper_plot_gqa_4cases_summary.py to overlay on the analytical bars.
+
+Single layer of decode attention only — the projection × 80 takes
+the per-layer measurement to a per-token total.
+
+Usage:
+  python scripts/paper/measure_gqa_decode_placement_comm.py
+"""
+from __future__ import annotations
+
+import json
+import os
+import sys
+from pathlib import Path
+
+from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_repl_pe_sp import (
+    gqa_attention_decode_long_ctx_cube_repl_pe_sp_kernel as _case3_kernel,
+)
+from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_repl_pe_tp import (
+    gqa_attention_decode_long_ctx_cube_repl_pe_tp_kernel as _case1_kernel,
+)
+from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_sp_pe_sp import (
+    gqa_attention_decode_long_ctx_cube_sp_pe_sp_kernel as _case6_kernel,
+)
+from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_sp_pe_tp import (
+    gqa_attention_decode_long_ctx_cube_sp_pe_tp_kernel as _case2_kernel,
+)
+from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_sp_pe_tp_dhead import (
+    gqa_attention_decode_long_ctx_cube_sp_pe_tp_dhead_kernel as _case4_kernel,
+)
+from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_tp_dhead_pe_sp import (
+    gqa_attention_decode_long_ctx_cube_tp_dhead_pe_sp_kernel as _case5_kernel,
+)
+from kernbench.benches.gqa_helpers.shared._gqa_panel_helpers import _ccl_cfg
+from kernbench.ccl.sfr_config import configure_sfr_intercube_multisip
+from kernbench.policy.placement.dp import DPPolicy
+from kernbench.runtime_api.bench_runner import run_bench
+from kernbench.runtime_api.types import resolve_device
+from kernbench.sim_engine.engine import GraphEngine
+from kernbench.topology.builder import resolve_topology
+
+_C = 8
+_P = 8
+_N_LAYERS = 80
+_S_KV_MEAS = 64 * 1024       # per-run simulator S_kv (1/16th of headline)
+_S_KV_HEADLINE = 1 << 20     # 1 Mi tokens, the chart's headline S_kv
+
+# Per-cube S_kv share for the d_head-TP partial-score AR cost.
+#   Case 4 (Cube-SP × PE-TP_dhead): per-cube = S_kv / C
+#   Case 5 (Cube-TP_dhead × PE-SP): per-cube = S_kv (KV replicated across
+#       cubes for the cube-axis d_head-TP), so partial-score AR scales
+#       with the full S_kv.
+# Headline / measured per-cube ratios give the partial-score-AR scale-up
+# from S_kv = 64 K to S_kv = 1 M. (m,ℓ,O) AR is S_kv-independent.
+_PARTIAL_SCORE_SCALE = _S_KV_HEADLINE / _S_KV_MEAS   # = 16
+
+# Per-token Wo + FFN AR (constant across all cases, comes from the
+# attn-output and FFN-down all-reduces NOT measured by the attention-
+# only kernel run here).
+_WO_PER_LAYER_BYTES = 8 * 1024
+_FFN_PER_LAYER_BYTES = 8 * 1024
+_WO_FFN_PER_TOKEN_BYTES = (
+    (_WO_PER_LAYER_BYTES + _FFN_PER_LAYER_BYTES) * _N_LAYERS
+)   # 1.25 MB
+
+# Total PE count in one KV-head group — average per-PE comm = total / N.
+_NUM_PES = _C * _P
+
+# Total partial-score-AR slice produced by the attention compute when
+# d_head is sharded. Used to split measured IPCQ traffic into the
+# S_kv-scaling component (partial scores) vs the S_kv-independent
+# component ((m,ℓ,O) merge). Same as the analytical formula in
+# paper_plot_gqa_4cases_summary.py: h_q · S_q · per_cube_S_kv · 2 bytes.
+_H_Q = 8
+_S_Q = 1
+_BYTES_PER_ELEM = 2
+
+_PARAMS = dict(C=_C, P=_P, T_q=_S_Q, S_kv=_S_KV_MEAS,
+               d_head=128, h_q=_H_Q, h_kv=1)
+
+
+def _bench_fn_case1(ctx):
+    """Case 1: Cube-Repl x PE-repl (PE-TP doesn't shard KV)."""
+    p = _PARAMS
+    configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
+    dp = DPPolicy(cube="replicate", pe="replicate",
+                  num_cubes=p["C"], num_pes=p["P"])
+    q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp, name="q_c1")
+    k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp, name="k_c1")
+    v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp, name="v_c1")
+    o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp, name="o_c1")
+    ctx.launch("case1_repl_repl", _case1_kernel,
+               q, k, v, o,
+               p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
+               p["d_head"], p["C"], p["P"],
+               _auto_dim_remap=False)
+
+
+def _bench_fn_case2(ctx):
+    """Case 2: Cube-SP x PE-repl (PE-TP doesn't shard KV)."""
+    p = _PARAMS
+    configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
+    dp_full = DPPolicy(cube="replicate", pe="replicate",
+                       num_cubes=p["C"], num_pes=p["P"])
+    dp_kv = DPPolicy(cube="row_wise", pe="replicate",
+                     num_cubes=p["C"], num_pes=p["P"])
+    q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_full, name="q_c2")
+    k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="k_c2")
+    v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="v_c2")
+    o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_full, name="o_c2")
+    ctx.launch("case2_sp_repl", _case2_kernel,
+               q, k, v, o,
+               p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
+               p["d_head"], p["C"], p["P"],
+               _auto_dim_remap=False)
+
+
+def _bench_fn_case3(ctx):
+    """Case 3: Cube-Repl x PE-SP."""
+    p = _PARAMS
+    configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
+    dp_full = DPPolicy(cube="replicate", pe="replicate",
+                       num_cubes=p["C"], num_pes=p["P"])
+    dp_kv = DPPolicy(cube="replicate", pe="row_wise",
+                     num_cubes=p["C"], num_pes=p["P"])
+    q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_full, name="q_c3")
+    k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="k_c3")
+    v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="v_c3")
+    o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_full, name="o_c3")
+    ctx.launch("case3_repl_sp", _case3_kernel,
+               q, k, v, o,
+               p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
+               p["d_head"], p["C"], p["P"],
+               _auto_dim_remap=False)
+
+
+def _bench_fn_case4(ctx):
+    p = _PARAMS
+    configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
+    dp_full = DPPolicy(cube="replicate", pe="column_wise",
+                       num_cubes=p["C"], num_pes=p["P"])
+    dp_kv = DPPolicy(cube="row_wise", pe="column_wise",
+                     num_cubes=p["C"], num_pes=p["P"])
+    q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_full, name="q_c4")
+    k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="k_c4")
+    v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="v_c4")
+    o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_full, name="o_c4")
+    ctx.launch("case4_dhead_tp", _case4_kernel,
+               q, k, v, o,
+               p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
+               p["d_head"], p["C"], p["P"],
+               _auto_dim_remap=False)
+
+
+def _bench_fn_case5(ctx):
+    p = _PARAMS
+    configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
+    dp_q = DPPolicy(cube="column_wise", pe="replicate",
+                    num_cubes=p["C"], num_pes=p["P"])
+    dp_kv = DPPolicy(cube="column_wise", pe="row_wise",
+                     num_cubes=p["C"], num_pes=p["P"])
+    q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_q, name="q_c5")
+    k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="k_c5")
+    v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="v_c5")
+    o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_q, name="o_c5")
+    ctx.launch("case5_dhead_tp_inter", _case5_kernel,
+               q, k, v, o,
+               p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
+               p["d_head"], p["C"], p["P"],
+               _auto_dim_remap=False)
+
+
+def _bench_fn_case6(ctx):
+    p = _PARAMS
+    configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
+    dp_full = DPPolicy(cube="replicate", pe="replicate",
+                       num_cubes=p["C"], num_pes=p["P"])
+    dp_kv = DPPolicy(cube="row_wise", pe="row_wise",
+                     num_cubes=p["C"], num_pes=p["P"])
+    q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_full, name="q_c6")
+    k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="k_c6")
+    v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
+                  dtype="f16", dp=dp_kv, name="v_c6")
+    o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
+                  dtype="f16", dp=dp_full, name="o_c6")
+    ctx.launch("case6_sp_sp", _case6_kernel,
+               q, k, v, o,
+               p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
+               p["d_head"], p["C"], p["P"],
+               _auto_dim_remap=False)
+
+
+def _sum_ipcq_bytes(op_log) -> int:
+    """Sum nbytes across all ipcq_copy records."""
+    return sum(
+        r.params.get("nbytes", 0)
+        for r in op_log
+        if r.op_kind == "memory" and r.op_name == "ipcq_copy"
+    )
+
+
+def _partial_score_slices(case: int) -> int:
+    """Divisor that splits S_kv into partial-score tiles, per the
+    analytical model in paper_plot_gqa_4cases_summary.py:
+
+      partial_score_per_PE = h_q * S_q * (s_kv / slices) * bytes
+
+    Case 4 (Cube-SP x PE-TP-dhead): intra-cube AR over d_head shards
+        on PE axis -> partial tile per PE has per-cube S_kv = s_kv/C.
+    Case 5 (Cube-TP-dhead x PE-SP): inter-cube AR over d_head shards
+        on cube axis -> partial tile per PE has per-PE S_kv = s_kv/P.
+    Cases 1, 2, 3, 6: no partial-score AR (only (m,l,O) merge).
+    """
+    if case == 4:
+        return _C
+    if case == 5:
+        return _P
+    return 0
+
+
+def _split_attn_layer_bytes(case: int, total_ipcq_bytes: int,
+                            s_kv: int) -> tuple[int, int]:
+    """Split per-layer attention-time IPCQ bytes into:
+      (partial_score_component, mlo_component).
+
+    The partial-score component scales with s_kv (so it must be scaled
+    when projecting from the measure-time s_kv to the headline s_kv);
+    the (m,l,O) component is constant in s_kv.
+
+    Partial-score size is analytically known per-case (formula in
+    _partial_score_slices); the remainder is treated as (m,l,O) + any
+    other S_kv-independent overhead. Per-PE = total / NUM_PES.
+    """
+    per_pe_total = total_ipcq_bytes // _NUM_PES
+    slices = _partial_score_slices(case)
+    if slices == 0:
+        # No partial-score AR for this case.
+        return 0, per_pe_total
+    partial_score_per_pe = (
+        _H_Q * _S_Q * (s_kv // slices) * _BYTES_PER_ELEM
+    )
+    partial_score_per_pe = min(partial_score_per_pe, per_pe_total)
+    mlo_per_pe = per_pe_total - partial_score_per_pe
+    return partial_score_per_pe, mlo_per_pe
+
+
+_KERNELS = (
+    (1, "Case 1 (Cube-Repl x PE-repl)",      _bench_fn_case1),
+    (2, "Case 2 (Cube-SP   x PE-repl)",      _bench_fn_case2),
+    (3, "Case 3 (Cube-Repl x PE-SP)",        _bench_fn_case3),
+    (4, "Case 4 (Cube-SP   x PE-TP d_head)", _bench_fn_case4),
+    (5, "Case 5 (Cube-TP d_head x PE-SP)",   _bench_fn_case5),
+    (6, "Case 6 (Cube-SP   x PE-SP)  [*]",   _bench_fn_case6),
+)
+
+
+def main() -> int:
+    topology = os.environ.get("GQA_1H_TOPOLOGY", "topology.yaml")
+    topo = resolve_topology(topology)
+
+    out: dict = {
+        "S_kv_measured": _S_KV_MEAS,
+        "S_kv_headline": _S_KV_HEADLINE,
+        "n_layers": _N_LAYERS,
+        "num_pes": _NUM_PES,
+        "wo_ffn_per_token_bytes": _WO_FFN_PER_TOKEN_BYTES,
+        "cases": {},
+    }
+
+    print(f"Measuring at S_kv={_S_KV_MEAS:,} ; scaling partial-score AR "
+          f"to S_kv={_S_KV_HEADLINE:,} (×{int(_PARTIAL_SCORE_SCALE)})")
+    print()
+
+    for case_id, label, bench_fn in _KERNELS:
+        try:
+            res = run_bench(
+                topology=topo, bench_fn=bench_fn,
+                device=resolve_device(None),
+                engine_factory=lambda t, d: GraphEngine(
+                    getattr(t, "topology_obj", t), enable_data=True,
+                ),
+            )
+        except Exception as e:
+            print(f"  {label:<42} FAIL: {type(e).__name__}: {e}")
+            return 1
+        if not res.completion.ok:
+            print(f"  {label:<42} ENGINE FAIL: {res.completion}")
+            return 1
+
+        total_ipcq = _sum_ipcq_bytes(res.engine.op_log)
+        partial_pe, mlo_pe = _split_attn_layer_bytes(
+            case_id, total_ipcq, _S_KV_MEAS,
+        )
+        # Per-token attention-time comm at S_kv = 1 M:
+        #   (partial_score_per_layer × scale  +  mlo_per_layer) × 80 layers
+        scaled_partial_per_token = (
+            partial_pe * int(_PARTIAL_SCORE_SCALE) * _N_LAYERS
+        )
+        mlo_per_token = mlo_pe * _N_LAYERS
+        attn_per_token = scaled_partial_per_token + mlo_per_token
+        total_per_token = attn_per_token + _WO_FFN_PER_TOKEN_BYTES
+
+        out["cases"][str(case_id)] = {
+            "label": label,
+            "total_ipcq_bytes_one_layer": total_ipcq,
+            "per_pe_partial_score_bytes_one_layer": partial_pe,
+            "per_pe_mlo_bytes_one_layer": mlo_pe,
+            "per_pe_attn_bytes_per_token_at_1M": attn_per_token,
+            "per_pe_total_bytes_per_token_at_1M": total_per_token,
+        }
+
+        print(f"  {label:<42} "
+              f"ipcq_total={total_ipcq:>10,}  "
+              f"per_pe_attn(1L)={(partial_pe + mlo_pe):>9,}  "
+              f"per_pe_total/tok@1M={total_per_token / (1<<20):>7.2f} MB")
+
+    out_path = (
+        Path(__file__).resolve().parents[2]
+        / "src" / "kernbench" / "benches"
+        / "1H_milestone_output" / "gqa" / "long_ctx"
+        / "gqa_long_ctx_6cases_measured_comm.json"
+    )
+    out_path.parent.mkdir(parents=True, exist_ok=True)
+    out_path.write_text(json.dumps(out, indent=2))
+    print()
+    print(f"wrote {out_path}")
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

scripts/paper/paper_cube_architecture.py

@@ -0,0 +1,155 @@
+"""Re-emit cube_view.svg in an academic (white-background, large-font)
+palette and convert it to PDF for the 1H-codesign-paper Figure 2.
+
+Source of truth: docs/diagrams/cube_view.svg (generated by
+src/kernbench/topology/visualizer.py:_render_cube_view_svg, dark theme).
+
+This script does a targeted color/font/size remap on the dark palette so
+the resulting figure prints well on a white paper page. If the upstream
+palette or geometry in visualizer.py changes, the maps below must be
+reviewed.
+"""
+from __future__ import annotations
+
+import re
+import shutil
+import subprocess
+from pathlib import Path
+
+REPO = Path(__file__).resolve().parents[2]
+SRC_SVG = REPO / "docs" / "diagrams" / "cube_view.svg"
+OUT_DIR = REPO / "docs" / "report" / "1H-codesign-paper" / "figures"
+OUT_SVG = OUT_DIR / "cube_architecture.svg"
+OUT_PDF = OUT_DIR / "cube_architecture.pdf"
+
+# ── 1. Legend-icon safety: legend rects whose general color rule below
+#    would turn them white-on-white. Apply BEFORE the wholesale fill rules.
+LEGEND_FIXUP: list[tuple[str, str]] = [
+    # "Relay" legend icon (slate-700 router)
+    ('fill="#334155" stroke="#475569" stroke-width="0.5"',
+     'fill="#94a3b8" stroke="#475569" stroke-width="0.5"'),
+    # "Mesh Link" legend icon (slate-600)
+    ('fill="#475569" stroke="#475569" stroke-width="0.5"',
+     'fill="#94a3b8" stroke="#475569" stroke-width="0.5"'),
+]
+
+# ── 2. Color remap: dark theme -> academic (white) theme
+COLOR_MAP: list[tuple[str, str]] = [
+    # page background slate-900 -> white
+    ('fill="#0f172a"', 'fill="#ffffff"'),
+    # title text slate-400 -> slate-800
+    ('fill="#94a3b8"', 'fill="#1f2937"'),
+    # subtitle text slate-500 -> slate-600
+    ('fill="#64748b"', 'fill="#475569"'),
+    # router has-attach stroke slate-500 -> slate-600
+    ('stroke="#64748b"', 'stroke="#475569"'),
+    # router has-attach fill slate-600 -> white
+    ('fill="#475569"', 'fill="#ffffff"'),
+    # mesh lines + cube boundary stroke slate-600 -> slate-400
+    ('stroke="#475569"', 'stroke="#94a3b8"'),
+    # router no-attach fill slate-700 -> white
+    ('fill="#334155"', 'fill="#ffffff"'),
+    # component block dark fills -> white (PE / M_CPU / SRAM / UCIe)
+    ('fill="#2d1f3d"', 'fill="#ffffff"'),
+    ('fill="#451a03"', 'fill="#ffffff"'),
+    ('fill="#1c1917"', 'fill="#ffffff"'),
+    ('fill="#1e1b4b"', 'fill="#ffffff"'),
+    # HBM zone background emerald-950 -> emerald-50
+    ('fill="#052e16"', 'fill="#ecfdf5"'),
+    # router label text white -> slate-800
+    ('fill="white"', 'fill="#1f2937"'),
+    # boost low-alpha emerald annotations so PE/HBM BW labels read on light bg
+    ('fill="#10b98188"', 'fill="#047857"'),
+    ('fill="#05966988"', 'fill="#059669"'),
+]
+
+# ── 3. UCIe palette desaturation: bright violet/indigo/fuchsia is too
+#      attention-grabbing on a white page; remap to a slate gradient.
+#      NOTE: side-effect on PE2/PE3 HBM port-bar colors (which share
+#      #8b5cf6/#a78bfa with UCIe) — PE labels still convey identity.
+#      Applied AFTER the academic color map so the new slate values are
+#      not picked up by the rules above.
+UCIE_DESATURATE: list[tuple[str, str]] = [
+    # UCIe block stroke/text (violet-500) -> slate-600
+    ('"#8b5cf6"', '"#475569"'),
+    # UCIe cell 1 / PE3 port bar (violet-400) -> slate-400
+    ('"#a78bfa"', '"#94a3b8"'),
+    # UCIe cell 0 (indigo-400) -> slate-300
+    ('"#818cf8"', '"#cbd5e1"'),
+    # UCIe cell 2 (purple-400) -> slate-500
+    ('"#c084fc"', '"#64748b"'),
+    # UCIe cell 3 (fuchsia-400) -> gray-700
+    ('"#e879f9"', '"#374151"'),
+]
+
+# ── 4. Font-size bumps. The CUBE figure is rendered at half-text-width
+#      (~250pt) inside the side-by-side subfigure in 02-platform.tex, so
+#      native fonts get crushed ~3x by \linewidth scaling. We push the
+#      bumps to the legibility limit of the layout (router-label text
+#      stays inside a slightly enlarged circle; legend items may touch).
+FONT_MAP: dict[str, str] = {
+    "5":  "10",
+    "6":  "12",
+    "7":  "14",
+    "8":  "13",   # legend rect text — capped by upstream layout spacing
+                  # (advance = 7*len(label)+24 was sized for ~font 8);
+                  # font 13 keeps each item's text inside its slot.
+    "9":  "14",
+    "10": "15",
+    "11": "16",
+    "14": "18",   # title — kept moderate so it does not overflow canvas
+}
+
+# ── 5. Router circle radius bump (only circles use r="8"). Enlarged so
+#      the bumped router labels stay inside the circle.
+RADIUS_MAP: list[tuple[str, str]] = [
+    (' r="8"', ' r="17"'),
+]
+
+# ── 6. Tighten whitespace: move legend just below dashed box and crop
+#      the unused canvas margins so LaTeX's \linewidth scaling does not
+#      shrink the figure text any more than necessary.
+LAYOUT_FIXUP: list[tuple[str, str]] = [
+    # Move legend rects up (dashed box bottom is at y=760)
+    (' y="865"', ' y="775"'),
+    # Move legend text baselines up to match (offset = 9 px)
+    (' y="874"', ' y="784"'),
+    # Tight crop: 5 px margin around dashed box + cut top/bottom whitespace
+    ('<svg xmlns="http://www.w3.org/2000/svg" width="970" height="900" '
+     'viewBox="0 0 970 900">',
+     '<svg xmlns="http://www.w3.org/2000/svg" width="860" height="798" '
+     'viewBox="55 2 860 798">'),
+]
+
+
+def _bump_font(m: re.Match) -> str:
+    return f'font-size="{FONT_MAP.get(m.group(1), m.group(1))}"'
+
+
+def main() -> None:
+    if not SRC_SVG.exists():
+        raise SystemExit(f"source SVG missing: {SRC_SVG}")
+    rsvg = shutil.which("rsvg-convert")
+    if rsvg is None:
+        raise SystemExit("rsvg-convert not found (brew install librsvg)")
+
+    svg = SRC_SVG.read_text(encoding="utf-8")
+    for old, new in (LEGEND_FIXUP + COLOR_MAP + UCIE_DESATURATE
+                     + RADIUS_MAP + LAYOUT_FIXUP):
+        if old not in svg:
+            print(f"warn: pattern not present in source SVG: {old}")
+        svg = svg.replace(old, new)
+    svg = re.sub(r'font-size="(\d+)"', _bump_font, svg)
+
+    OUT_DIR.mkdir(parents=True, exist_ok=True)
+    OUT_SVG.write_text(svg, encoding="utf-8")
+    subprocess.run(
+        [rsvg, "-f", "pdf", "-o", str(OUT_PDF), str(OUT_SVG)],
+        check=True,
+    )
+    print(f"wrote {OUT_SVG.relative_to(REPO)}")
+    print(f"wrote {OUT_PDF.relative_to(REPO)}")
+
+
+if __name__ == "__main__":
+    main()

scripts/paper/paper_latency_model_diagram.py

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

scripts/paper/paper_plot_gemm_async_vs_composite.py

@@ -0,0 +1,261 @@
+"""Composite vs async-loading GEMM TFLOPS comparison.
+
+For each shape in the milestone sweep, run both benches:
+  - matmul-composite (load_ref variant — A pre-staged, B streamed by
+    scheduler inside one composite command)
+  - matmul-async (A and B both async-loaded via tl.load, then a single
+    tl.dot — no per-tile overlap of streaming-B with GEMM)
+
+Compute per-PE TFLOPS = 2*M*K*N / pe_window_ns for each kernel and emit a
+side-by-side bar chart PNG to
+  src/kernbench/benches/1H_milestone_output/gemm/gemm_composite_vs_async_tflops.png
+
+Run from repo root:
+  python scripts/paper/paper_plot_gemm_async_vs_composite.py
+"""
+from __future__ import annotations
+
+import json
+import os
+import sys
+from pathlib import Path
+
+REPO = Path(__file__).resolve().parents[2]
+OUT_DIR = REPO / "src" / "kernbench" / "benches" / "1H_milestone_output" / "gemm"
+OUT_PNG = OUT_DIR / "gemm_composite_vs_async_tflops.png"
+OUT_JSON = OUT_DIR / "gemm_composite_vs_async.json"
+
+TOPO_PATH = REPO / "topology.yaml"
+
+# Same 7 shapes as the milestone composite sweep (skip square=512: B alone
+# is 512 KiB; async needs all of A+B+out in TCM scratch and the 512×512
+# square pushes total scratch use past the 1 MiB cap).
+SHAPES: list[tuple[int, int, int]] = [
+    (32, 32, 32),     # 1 tile, K<TILE_K (under-tile in K)
+    (32, 64, 32),     # 1 tile
+    (32, 128, 32),    # 2 tiles
+    (32, 128, 128),   # 8 tiles
+    (32, 3072, 32),   # 48 tiles (deep-K)
+    (8, 128, 128),    # under-tile in M
+    (128, 8, 128),    # under-tile in K
+]
+
+ENGINES = ("pe_dma", "pe_fetch_store", "pe_gemm", "pe_math")
+STAGES = ("DMA_READ", "DMA_WRITE", "FETCH", "STORE", "GEMM", "MATH")
+
+
+def _pe_records(op_log):
+    return [r for r in op_log
+            if any(r.component_id.endswith("." + e) for e in ENGINES)]
+
+
+def _pe_window_ns(op_log) -> float:
+    pe = _pe_records(op_log)
+    if not pe:
+        return 0.0
+    return max(r.t_end for r in pe) - min(r.t_start for r in pe)
+
+
+def _composite_window_ns(op_log) -> float:
+    """For the composite kernel: window of records carrying a stage_type
+    set by the composite plan (DMA_READ, FETCH, GEMM, STORE, DMA_WRITE).
+    Excludes the initial up-front tl.load(A) because that record is an
+    atomic DmaReadCmd with no stage_type. Matches the existing
+    milestone_1h_gemm.py / gemm_per_pe_tflops.png methodology.
+    """
+    stage_records = [r for r in op_log
+                     if r.params.get("stage_type") in STAGES]
+    if not stage_records:
+        return 0.0
+    return max(r.t_end for r in stage_records) \
+           - min(r.t_start for r in stage_records)
+
+
+def _async_engine_window_ns(op_log) -> float:
+    """For the async kernel: engine pipeline window that excludes the
+    initial tl.load(A), paralleling composite_window's exclusion of the
+    up-front A pre-stage. The first DMA_READ record on pe_dma is the
+    tl.load(A); the window starts at the SECOND pe_dma record's t_start
+    (= tl.load(B)) and ends at the last engine record's t_end.
+    """
+    pe = _pe_records(op_log)
+    if not pe:
+        return 0.0
+    dma = sorted(
+        (r for r in op_log if r.component_id.endswith(".pe_dma")),
+        key=lambda r: r.t_start,
+    )
+    if len(dma) < 2:
+        return _pe_window_ns(op_log)
+    window_start = dma[1].t_start
+    window_end = max(r.t_end for r in pe)
+    return window_end - window_start
+
+
+def _run_one(bench_name: str, variant: str | None, M: int, K: int, N: int) -> dict:
+    os.environ["MATMUL_M"] = str(M)
+    os.environ["MATMUL_K"] = str(K)
+    os.environ["MATMUL_N"] = str(N)
+    if variant is not None:
+        os.environ["MATMUL_VARIANT"] = variant
+    elif "MATMUL_VARIANT" in os.environ:
+        del os.environ["MATMUL_VARIANT"]
+
+    from kernbench.benches.registry import resolve as resolve_bench
+    from kernbench.runtime_api.bench_runner import run_bench
+    from kernbench.runtime_api.types import resolve_device
+    from kernbench.sim_engine.engine import GraphEngine
+    from kernbench.topology.builder import resolve_topology
+
+    # Chunked async loads A in TILE_K-sized blocks at sequential offsets;
+    # those offsets do not match the row-major (M, K) layout of A, so the
+    # DataExecutor would fail on the resulting mismatched read. The
+    # simulator's *timing* path doesn't need correct data — only the
+    # number of bytes per DMA matters — so bypass DataExecutor for the
+    # chunked kernel. Composite/naive-async loads use the full (M, K)
+    # shape and remain data-correct.
+    if bench_name in ("matmul-async-chunked", "matmul-async-chunked-db"):
+        GraphEngine._flush_data_phase = lambda self: None
+
+    topo = resolve_topology(str(TOPO_PATH))
+    bench = resolve_bench(bench_name).run
+    device = resolve_device(None)
+    result = run_bench(
+        topology=topo, bench_fn=bench, device=device,
+        engine_factory=lambda t, d: GraphEngine(
+            getattr(t, "topology_obj", t), enable_data=True,
+        ),
+    )
+    if not result.completion.ok:
+        raise RuntimeError(f"{bench_name} failed at {M}x{K}x{N}: {result.completion}")
+    log = result.engine.op_log
+    pe_window = _pe_window_ns(log)
+    if bench_name == "matmul-composite":
+        engine_window = _composite_window_ns(log)
+    elif bench_name in ("matmul-async-chunked", "matmul-async-chunked-db"):
+        # First N_chunks pe_dma records are A pre-stage; engine window
+        # starts at the (N_chunks+1)-th pe_dma record (= first B-chunk
+        # load). Fall back to the naive analog for K <= TILE_K (kernel
+        # collapses to single load+dot+store).
+        TILE_K = 64
+        n_chunks = max(K // TILE_K, 1)
+        if n_chunks <= 1:
+            engine_window = _async_engine_window_ns(log)
+        else:
+            dma = sorted(
+                (r for r in log if r.component_id.endswith(".pe_dma")),
+                key=lambda r: r.t_start,
+            )
+            pe = _pe_records(log)
+            if len(dma) > n_chunks and pe:
+                engine_window = max(r.t_end for r in pe) - dma[n_chunks].t_start
+            else:
+                engine_window = _async_engine_window_ns(log)
+    else:
+        engine_window = _async_engine_window_ns(log)
+    flops = 2 * M * K * N
+    # flops / ns * 1e-3 = TFLOP/s  (since 1 flop/ns = 1 GFLOP/s)
+    return {
+        "M": M, "K": K, "N": N,
+        "bench": bench_name, "variant": variant,
+        "pe_window_ns": pe_window,
+        "engine_window_ns": engine_window,
+        "flops": flops,
+        "tflops": (flops / engine_window / 1000.0) if engine_window > 0 else 0.0,
+        "n_records": len(log),
+    }
+
+
+def collect() -> list[dict]:
+    rows: list[dict] = []
+    for M, K, N in SHAPES:
+        print(f"  shape M={M:4d} K={K:5d} N={N:4d} ...", flush=True)
+        comp = _run_one("matmul-composite", "load_ref", M, K, N)
+        asyn = _run_one("matmul-async", None, M, K, N)
+        chnk = _run_one("matmul-async-chunked", None, M, K, N)
+        chnkdb = _run_one("matmul-async-chunked-db", None, M, K, N)
+        print(f"    composite:    {comp['engine_window_ns']:8.1f} ns  "
+              f"{comp['tflops']:6.3f} TFLOPS")
+        print(f"    async naive:  {asyn['engine_window_ns']:8.1f} ns  "
+              f"{asyn['tflops']:6.3f} TFLOPS")
+        print(f"    chunked all:  {chnk['engine_window_ns']:8.1f} ns  "
+              f"{chnk['tflops']:6.3f} TFLOPS")
+        print(f"    chunked db=2: {chnkdb['engine_window_ns']:8.1f} ns  "
+              f"{chnkdb['tflops']:6.3f} TFLOPS")
+        rows.append(comp)
+        rows.append(asyn)
+        rows.append(chnk)
+        rows.append(chnkdb)
+    return rows
+
+
+def plot(rows: list[dict]) -> None:
+    import matplotlib.pyplot as plt
+    import numpy as np
+
+    shape_keys = []
+    for M, K, N in SHAPES:
+        shape_keys.append((M, K, N))
+
+    by_key: dict[tuple[int, int, int], dict[str, dict]] = {}
+    for r in rows:
+        key = (r["M"], r["K"], r["N"])
+        by_key.setdefault(key, {})[r["bench"]] = r
+
+    labels = [f"M={M}\nK={K}\nN={N}" for (M, K, N) in shape_keys]
+    comp_tflops = [by_key[k]["matmul-composite"]["tflops"] for k in shape_keys]
+    asyn_tflops = [by_key[k]["matmul-async"]["tflops"] for k in shape_keys]
+    chnk_tflops = [by_key[k]["matmul-async-chunked"]["tflops"] for k in shape_keys]
+    chnkdb_tflops = [by_key[k]["matmul-async-chunked-db"]["tflops"]
+                     for k in shape_keys]
+
+    x = np.arange(len(shape_keys))
+    width = 0.20
+
+    fig, ax = plt.subplots(figsize=(14, 5.5))
+    b_c = ax.bar(x - 1.5*width, comp_tflops, width,
+                 label="Composite (load_ref)", color="#10b981")
+    b_a = ax.bar(x - 0.5*width, asyn_tflops, width,
+                 label="Async naive (tl.load full + tl.dot)",
+                 color="#f59e0b")
+    b_db = ax.bar(x + 0.5*width, chnkdb_tflops, width,
+                  label="Async chunked-prefetch, depth=2 "
+                        "(TCM-bounded)", color="#a855f7")
+    b_k = ax.bar(x + 1.5*width, chnk_tflops, width,
+                 label="Async chunked-prefetch, depth=$\\infty$ "
+                       "(all B-tiles queued up front)", color="#3b82f6")
+
+    ax.axhline(8.0, linestyle="--", color="#94a3b8", linewidth=0.8,
+               label="Per-PE GEMM peak (8 TFLOP/s)")
+    ax.set_ylabel("Per-PE achieved TFLOP/s")
+    ax.set_title("Composite vs async-loading GEMM — per-PE throughput "
+                 "(engine pipeline window, A pre-stage excluded)")
+    ax.set_xticks(x)
+    ax.set_xticklabels(labels, fontsize=8)
+    ax.grid(axis="y", alpha=0.25)
+    ax.legend(loc="upper left", fontsize=9)
+
+    for bar in (*b_c, *b_a, *b_db, *b_k):
+        h = bar.get_height()
+        ax.annotate(f"{h:.2f}", xy=(bar.get_x() + bar.get_width()/2, h),
+                    xytext=(0, 2), textcoords="offset points",
+                    ha="center", va="bottom", fontsize=7, color="#475569")
+
+    fig.tight_layout()
+    OUT_DIR.mkdir(parents=True, exist_ok=True)
+    fig.savefig(OUT_PNG, dpi=150)
+    print(f"wrote {OUT_PNG.relative_to(REPO)}")
+    plt.close(fig)
+
+
+def main() -> int:
+    rows = collect()
+    OUT_DIR.mkdir(parents=True, exist_ok=True)
+    OUT_JSON.write_text(json.dumps(rows, indent=2))
+    print(f"wrote {OUT_JSON.relative_to(REPO)}")
+    plot(rows)
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

scripts/paper/paper_plot_gqa.py

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

scripts/paper/paper_plot_gqa_4cases_summary.py

@@ -0,0 +1,638 @@
+"""GQA 4-cases (+ 2 d_head-TP variants) — combined memory + comm summary.
+
+Single PNG with two panels side-by-side:
+  (left)  Per-PE KV memory at S_kv = 1 M tokens (across 80 layers).
+  (right) Per-PE communication per output token at S_kv = 1 M (decode).
+
+================================================================
+SYSTEM INPUTS (LLaMA-3.1-70B single-KV-head group)
+================================================================
+  N_layers  = 80          (transformer layers per token)
+  h_kv      = 1           (per KV group)
+  h_q       = 8           (query heads per KV group)
+  d_head    = 128
+  d_model   = 8192        (LLaMA-3.1-70B hidden dim — for Wo / FFN AR)
+  bytes     = 2           (FP16)
+  C         = 8 cubes per KV group     P = 8 PEs per cube
+  HBM_per_PE     = 6.0 GB
+  Attn weights   = 1.76 GB → KV headroom = 4.24 GB / PE
+  B = 1 user      S_q = 1 token (decode)
+
+================================================================
+(1) PER-PE KV MEMORY                                       (left panel)
+================================================================
+
+KV bytes per token across all 80 layers, single KV group:
+  KV/tok = 2 (K+V) · h_kv · d_head · bytes · N_layers
+         = 2 · 1   · 128    · 2     · 80
+         = 40 KB / token
+
+Per-PE share = KV/tok ÷ divisor, where divisor depends on the K/V
+tensor placement (NOT on the compute-side label):
+
+  Case 1   Cube-SP   × PE-replicate         divisor = C       =  8
+  Case 1'  Cube-SP   × PE-TP (on d_head)    divisor = C·P     = 64
+  Case 2   Cube-Repl × PE-replicate         divisor = 1       =  1
+  Case 3   Cube-Repl × PE-SP                divisor = P       =  8
+  Case 3'  Cube-TP   × PE-SP                divisor = C·P     = 64
+  Case 4   Cube-SP   × PE-SP    ★           divisor = C·P     = 64
+
+Per-PE bytes @ S_kv = 1 M = KV/tok · 1 M / divisor:
+  Case 1   →  5.0  GB ✗   (exceeds 4.24 GB headroom)
+  Case 1'  →  640  MB ✓
+  Case 2   → 40.0  GB ✗
+  Case 3   →  5.0  GB ✗
+  Case 3'  →  640  MB ✓
+  Case 4   →  640  MB ✓
+
+Max KV context per PE = 4.24 GB · 1024² / (KV/tok ÷ divisor):
+  Case 1, 3   :  889 K tokens
+  Case 1', 3', 4 :  7.11 M tokens ✓
+  Case 2      :  111 K tokens
+
+================================================================
+(2) PER-PE COMMUNICATION                                   (right panel)
+================================================================
+
+(A) WEIGHT AllReduces (constant across all cases)
+    Wo  AR  ≈ 8 KB / layer / PE          (Wq replicated → AR partial Y)
+    FFN AR  ≈ 8 KB / layer / PE
+    × 80 layers = 1.25 MB / token / PE   (the bottom blue stack on every bar)
+
+(B) ATTENTION-TIME collective (the differentiator)
+
+  Cases that compute partial attention locally (PE-SP / PE-repl):
+    Q is replicated and each rank attends to its own complete KV slice
+    (full d_head) → only the small (m, ℓ, O) triple needs AllReducing.
+    Payload (h_q · S_q · d_head · 2) ≈ 4 KB per AR step, hierarchical ≈ 32 KB / layer.
+    CONSTANT in S_kv.
+
+    Case 1   : inter-cube AR on (m,ℓ,O)          → ~32 KB / layer
+    Case 2   : nothing                            → 0
+    Case 3   : intra-cube AR on (m,ℓ,O)          → ~32 KB / layer
+    Case 4 ★ : intra + inter-cube AR on (m,ℓ,O)  → ~64 KB / layer
+
+  d_head-TP variants (1', 3'):
+    Each rank holds only d_head/divisor dims → Q·K^T produces only
+    partial sums → must AR PARTIAL SCORES before softmax. The score
+    tile per AR is (h_q · S_q · S_kv / slices_for_seq) · 2 bytes,
+    which SCALES with S_kv.
+
+    Case 1'  : intra-cube partial-score AR
+               payload = h_q · S_q · (S_kv/C) · 2 = 2 MB / layer at 1 M
+    Case 3'  : inter-cube partial-score AR (UCIe, slower than NoC)
+               payload = h_q · S_q · (S_kv/P) · 2 = 2 MB / layer at 1 M
+
+  Plus tiny (m,ℓ,O) AR for both 1' and 3' to combine partial attentions
+  across the remaining axis (~32 KB / layer).
+
+(C) Total per output token per PE (sum over 80 layers, S_kv = 1 M):
+  Case 1   ≈   4 MB     (1.25 MB Wo+FFN + 2.5 MB inter-cube AR)
+  Case 1'  ≈ 166 MB     (1.25 MB + 160 MB partial-score AR + ~5 MB other)
+  Case 2   ≈   1.2 MB   (Wo+FFN only)
+  Case 3   ≈   4 MB     (1.25 MB + 2.5 MB intra-cube AR)
+  Case 3'  ≈ 166 MB     (1.25 MB + 160 MB partial-score AR + ~5 MB other)
+  Case 4 ★ ≈   6 MB     (1.25 MB + 5 MB 2-phase AR)
+
+================================================================
+KEY TAKEAWAYS
+================================================================
+  - Memory winners (640 MB / PE @ 1M):  Cases 1', 3', 4 (any 64-way sharding).
+  - Comm winner among those three:  Case 4 (6 MB).  Cases 1' and 3' both
+    pay ~160 MB because the d_head-sharded variant ARs a 2 MB partial-
+    score tile every layer — the score tensor scales with S_kv while
+    the (m, ℓ, O) triple does not.
+  - Cases 1 and 3 (8-way sharding) don't fit 1M context (5 GB vs 4.24 GB
+    headroom).
+  - Case 2 has the cheapest comm but the most memory (40 GB / PE).
+  - Case 4 ★ is the Pareto-best:  fits 1M context AND lowest comm
+    among memory-feasible options.
+
+Output PNG:
+  src/kernbench/benches/1H_milestone_output/gqa/long_ctx/
+    gqa_4cases_summary.png
+"""
+from __future__ import annotations
+
+import json
+import textwrap
+from pathlib import Path
+
+import matplotlib.patches as mpatches
+import matplotlib.pyplot as plt
+
+# ── System constants (LLaMA-3.1-70B single-KV-head group, decode) ───
+_N_LAYERS          = 80
+_H_KV              = 1
+_H_Q               = 8
+_D_HEAD            = 128
+_D_MODEL           = 8192
+_BYTES_PER_ELEM    = 2          # FP16
+_C                 = 8          # cubes per KV group
+_P                 = 8          # PEs per cube
+_B                 = 1
+_S_Q               = 1          # decode
+_HBM_PER_PE_GB     = 6.0
+_WEIGHTS_PER_PE_GB = 1.76
+_HEADROOM_GB       = _HBM_PER_PE_GB - _WEIGHTS_PER_PE_GB    # 4.24 GB
+
+_HEADLINE_S_KV = 1 << 20        # 1 Mi tokens
+
+# Per-token KV bytes (single KV group, all 80 layers).
+_KV_PER_TOK_BYTES = (
+    2 * _H_KV * _D_HEAD * _BYTES_PER_ELEM * _N_LAYERS
+)   # 40 KB
+
+# Per-token Wo + FFN AR (constant across cases).
+_WO_PER_LAYER_BYTES  = 8 * 1024
+_FFN_PER_LAYER_BYTES = 8 * 1024
+
+# ── Per-PE attention weight breakdown (single KV-head group) ────────
+#
+# LLaMA-3.1-70B single-KV-head group dimensions:
+#   d_model           = 8192
+#   h_q  per group    = 8     (8 query heads attend to 1 KV head per group)
+#   h_kv per group    = 1
+#   d_head            = 128
+#   FP16  (2 bytes)
+#
+# Per layer, per PE:
+#   Wq  shape (d_model, h_q · d_head)  REPLICATED across all 64 PEs of group
+#         = 8192 · 8·128 · 2 bytes        = 16 MB / layer / PE
+#   Wk  shape (d_model, h_kv · d_head) REPLICATED
+#         = 8192 · 1·128 · 2 bytes        =  2 MB / layer / PE
+#   Wv  shape (d_model, h_kv · d_head) REPLICATED
+#         = 8192 · 1·128 · 2 bytes        =  2 MB / layer / PE
+#   Wo  shape (h_q · d_head, d_model)  ROW-SPLIT across C=8 cubes,
+#                                       replicated within cube
+#         = (1024/C) · 8192 · 2 bytes    =  2 MB / layer / PE
+#   ────────────────────────────────────────────────────────────────
+#   total attn weights / layer / PE       = 22 MB
+#   × 80 layers                           = 1.76 GB / PE   (decimal GB)
+#
+# KV headroom per PE  =  HBM_PER_PE - weights  =  6.0 - 1.76  =  4.24 GB.
+# (FFN weights are accounted for in a separate budget, not in this 4.24 GB.)
+
+# Slide-17 convention: per-layer values in binary MB (MiB), totals in
+# "GB" formed by ×80 layers ÷ 1000 — giving the canonical 1.76 GB total
+# and 4.24 GB headroom that the slide-17 chart reports.
+_WQ_MB_PER_LAYER = (_D_MODEL * _H_Q * _D_HEAD * _BYTES_PER_ELEM) / (1024 ** 2)   # 16
+_WK_MB_PER_LAYER = (_D_MODEL * _H_KV * _D_HEAD * _BYTES_PER_ELEM) / (1024 ** 2)  #  2
+_WV_MB_PER_LAYER = (_D_MODEL * _H_KV * _D_HEAD * _BYTES_PER_ELEM) / (1024 ** 2)  #  2
+_WO_MB_PER_LAYER = ((_H_Q * _D_HEAD) // _C * _D_MODEL
+                    * _BYTES_PER_ELEM) / (1024 ** 2)                              #  2
+
+_WQ_GB = _WQ_MB_PER_LAYER * _N_LAYERS / 1000    # 1.28 GB
+_WK_GB = _WK_MB_PER_LAYER * _N_LAYERS / 1000    # 0.16 GB
+_WV_GB = _WV_MB_PER_LAYER * _N_LAYERS / 1000    # 0.16 GB
+_WO_GB = _WO_MB_PER_LAYER * _N_LAYERS / 1000    # 0.16 GB
+_WEIGHTS_GB = _WQ_GB + _WK_GB + _WV_GB + _WO_GB  # 1.76 GB
+
+# (m, ℓ, O) AR cost per layer, derived from the kernel topology.
+#
+# Per-PE payload  T  for one (m, ℓ, O) merge step:
+#     O  — h_q · S_q · d_head · 2 bytes  (FP16)
+#     m  — h_q · S_q · 4 bytes           (FP32)
+#     ℓ  — h_q · S_q · 4 bytes           (FP32)
+#     T  = h_q · S_q · (d_head · 2 + 8)  ≈ 2.1 KB
+#
+# Reduce algorithm: the decode kernels use hierarchical reduce-only
+# (chain / tree) rather than ring all-reduce, because the merged result
+# only needs to land on the cube that runs the downstream Wo gemm —
+# not on every PE.  For a chain of N participants the total traffic is
+# (N-1)·T and the per-PE average is (N-1)/N · T.
+#
+# The previous version used a hard-coded 32 KB / layer placeholder which
+# overestimated the per-PE cost by ~6× (and ~3× even under a hypothetical
+# ring-AR assumption).  See `gqa_long_ctx_6cases_measured_comm.json` for
+# the measured numbers this matches against.
+_MLO_PAYLOAD_BYTES = _H_Q * _S_Q * (_D_HEAD * _BYTES_PER_ELEM + 8)
+
+
+def _reduce_chain_per_pe_bytes(n_participants: int) -> int:
+    """Per-PE average bytes for a single-stage chain/tree reduce of T."""
+    if n_participants <= 1:
+        return 0
+    return _MLO_PAYLOAD_BYTES * (n_participants - 1) // n_participants
+
+
+_MLO_INTRA_BYTES_PER_LAYER = _reduce_chain_per_pe_bytes(_P)  # PE-axis
+_MLO_INTER_BYTES_PER_LAYER = _reduce_chain_per_pe_bytes(_C)  # cube-axis
+
+# Cases — renumbered in MEMORY-DESCENDING ORDER (left to right):
+#   Case 1   (40 GB) : no sharding
+#   Cases 2, 3 ( 5 GB) : single-axis sharding (cube OR PE only)
+#   Cases 4, 5, 6 (640 MB) : two-axis sharding (cube AND PE)
+#                            — Case 6 ★ is the Pareto-best (lowest comm)
+_CASES = (1, 2, 3, 4, 5, 6)
+
+_DIVISOR = {
+    1: 1,             # Cube-Repl × PE-replicate                — no sharding
+    2: _C,            # Cube-SP   × PE-replicate                — cube-axis only
+    3: _P,            # Cube-Repl × PE-SP                       — PE-axis (S_kv) only
+    4: _C * _P,       # Cube-SP   × PE-TP  (d_head)             — 64-way (d_head intra)
+    5: _C * _P,       # Cube-TP   × PE-SP                       — 64-way (d_head inter)
+    6: _C * _P,       # Cube-SP   × PE-SP                       — 64-way (S_kv both axes) ★
+}
+
+_CASE_LABEL = {
+    1: "Case 1\nCube-Repl\nPE-repl",
+    2: "Case 2\nCube-SP\nPE-repl",
+    3: "Case 3\nCube-Repl\nPE-SP",
+    4: "Case 4\nCube-SP\nPE-TP",
+    5: "Case 5\nCube-TP\nPE-SP",
+    6: "Case 6 ★\nCube-SP\nPE-SP",
+}
+_CASE_COLOR = {
+    1: "#C0504D",   # red    — worst memory (no sharding)
+    2: "#E0834A",   # orange — single-axis sharded (cube)
+    3: "#EBA854",   # tan    — single-axis sharded (PE)
+    4: "#A6C2E0",   # light blue — d_head-TP 64-way (PE)
+    5: "#C7D8A0",   # light green — d_head-TP 64-way (cube)
+    6: "#8064A2",   # purple — Pareto winner ★ (S_kv 64-way)
+}
+_ATTN_DESC = {
+    1: "none",
+    2: "online-softmax (m,ℓ,O) — inter-cube",
+    3: "online-softmax (m,ℓ,O) — intra-cube",
+    4: "partial scores + (m,ℓ,O) merge  (d_head-TP)",
+    5: "partial scores + (m,ℓ,O) merge  (d_head-TP)",
+    6: "online-softmax (m,ℓ,O) — intra + inter",
+}
+
+_WO_COLOR   = "#9EC5E8"
+_FFN_COLOR  = "#4A78B8"
+_ATTN_COLOR = "#E07A3F"
+
+_OUT_DIR = (
+    Path(__file__).resolve().parents[2]
+    / "src" / "kernbench" / "benches"
+    / "1H_milestone_output" / "gqa" / "long_ctx"
+)
+_MEASURED_JSON = _OUT_DIR / "gqa_long_ctx_6cases_measured_comm.json"
+
+
+def _load_measured() -> dict[int, float] | None:
+    """Load measured per-PE comm bytes (already scaled to S_kv=1M).
+
+    Returns {case_id: per_token_total_bytes} or None if JSON missing.
+    Produced by scripts/paper/measure_gqa_decode_placement_comm.py.
+    """
+    if not _MEASURED_JSON.exists():
+        return None
+    data = json.loads(_MEASURED_JSON.read_text())
+    return {
+        int(cid): info["per_pe_total_bytes_per_token_at_1M"]
+        for cid, info in data["cases"].items()
+    }
+
+
+# ── Formulae ────────────────────────────────────────────────────────
+
+
+def kv_per_pe_bytes(case, s_kv: int) -> int:
+    """Per-PE KV bytes at the given S_kv."""
+    return _KV_PER_TOK_BYTES * s_kv // _DIVISOR[case]
+
+
+def max_s_kv(case) -> int:
+    """Max KV context per PE given the 4.24 GB headroom."""
+    headroom_bytes = int(_HEADROOM_GB * (1 << 30))
+    bytes_per_tok = _KV_PER_TOK_BYTES // _DIVISOR[case]
+    return headroom_bytes // bytes_per_tok
+
+
+def _partial_score_bytes_per_layer(s_kv: int, slices_for_seq: int) -> int:
+    """AR payload for partial scores in d_head-TP cases."""
+    return _H_Q * _S_Q * (s_kv // slices_for_seq) * _BYTES_PER_ELEM
+
+
+def attn_comm_per_layer_bytes(case: int, s_kv: int) -> int:
+    """Per-layer per-PE attention-time comm bytes (decode, B=1, S_q=1).
+
+    Case numbering follows the memory-descending order defined in
+    _CASES (1=no sharding, 6=Pareto winner).
+    """
+    if case == 1:                # Cube-Repl × PE-repl — no sharding
+        return 0
+    if case == 2:                # Cube-SP   × PE-repl — inter-cube AR on (m,ℓ,O)
+        return _MLO_INTER_BYTES_PER_LAYER
+    if case == 3:                # Cube-Repl × PE-SP   — intra-cube AR on (m,ℓ,O)
+        return _MLO_INTRA_BYTES_PER_LAYER
+    if case == 4:                # Cube-SP   × PE-TP(d_head) — partial-score AR
+        return (_partial_score_bytes_per_layer(s_kv, _C)
+                + _MLO_INTRA_BYTES_PER_LAYER
+                + _MLO_INTER_BYTES_PER_LAYER)
+    if case == 5:                # Cube-TP(d_head) × PE-SP   — partial-score AR
+        return (_partial_score_bytes_per_layer(s_kv, _P)
+                + _MLO_INTRA_BYTES_PER_LAYER
+                + _MLO_INTER_BYTES_PER_LAYER)
+    if case == 6:                # Cube-SP   × PE-SP   — 2-phase (m,ℓ,O) AR ★
+        return _MLO_INTRA_BYTES_PER_LAYER + _MLO_INTER_BYTES_PER_LAYER
+    raise ValueError(f"unknown case {case}")
+
+
+def per_token_bytes(case, s_kv: int) -> tuple[int, int, int]:
+    """(Wo AR, FFN AR, Attn) bytes per output token per PE — × 80 layers."""
+    wo = _WO_PER_LAYER_BYTES * _N_LAYERS
+    ffn = _FFN_PER_LAYER_BYTES * _N_LAYERS
+    attn = attn_comm_per_layer_bytes(case, s_kv) * _N_LAYERS
+    return wo, ffn, attn
+
+
+# ── Formatters ──────────────────────────────────────────────────────
+
+
+def _fmt_bytes(b: float) -> str:
+    if b >= (1 << 30):
+        return f"{b / (1 << 30):.2f} GB"
+    if b >= (1 << 20):
+        return f"{b / (1 << 20):.1f} MB"
+    if b >= (1 << 10):
+        return f"{b / (1 << 10):.0f} KB"
+    return f"{b:.0f} B"
+
+
+def _fmt_tokens(n: int) -> str:
+    if n >= 1_000_000:
+        return f"{n / 1_000_000:.2f} M"
+    if n >= 1_000:
+        return f"{n / 1_000:.0f} K"
+    return f"{n}"
+
+
+# ── Panels ──────────────────────────────────────────────────────────
+
+
+def _plot_budget(ax) -> None:
+    """HBM budget per PE — stacked weights + KV headroom + HBM ceiling.
+
+    Wq is the dominant weight slice (~1.28 GB). Wk, Wv, Wo are each
+    small (~0.16 GB) so their slice labels would overlap on the bar —
+    they're shown in the legend only, and only Wq + KV-headroom get
+    on-bar annotations.
+    """
+    components = [
+        ("Wq  (REPL)",       _WQ_GB, "#7B9CC4"),
+        ("Wk  (REPL)",       _WK_GB, "#A0BBD8"),
+        ("Wv  (REPL)",       _WV_GB, "#C5D6E8"),
+        ("Wo  (cube-split)", _WO_GB, "#E2EAF3"),
+        ("KV cache headroom",
+                             _HBM_PER_PE_GB - _WEIGHTS_GB, "#9BBB59"),
+    ]
+    bottom = 0.0
+    for label, val, color in components:
+        ax.bar(0, val, bottom=bottom, color=color, edgecolor="black",
+               width=0.7, label=f"{label}  ·  {val:.2f} GB")
+        # Only annotate slices thick enough to fit text without overlap.
+        if val >= 0.50:
+            ax.text(0, bottom + val / 2, f"{label}\n{val:.2f} GB",
+                    ha="center", va="center", fontsize=9, weight="bold")
+        bottom += val
+    ax.axhline(_HBM_PER_PE_GB, color="red", ls="--", lw=1.4,
+               label=f"HBM = {_HBM_PER_PE_GB} GB")
+    ax.set_xticks([0])
+    ax.set_xticklabels(["per-PE HBM"], fontsize=10)
+    ax.set_ylabel("GB per PE")
+    ax.set_ylim(0, _HBM_PER_PE_GB * 1.10)
+    ax.set_title(
+        f"Per-PE HBM budget\n"
+        f"weights {_WEIGHTS_GB:.2f} GB + KV = {_HBM_PER_PE_GB} GB",
+        fontsize=10,
+    )
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    ax.legend(loc="upper right", fontsize=7.5, framealpha=0.92)
+
+
+def _plot_memory(ax) -> None:
+    vals_gb = [kv_per_pe_bytes(c, _HEADLINE_S_KV) / (1 << 30) for c in _CASES]
+    labels = [_CASE_LABEL[c] for c in _CASES]
+    colors = [_CASE_COLOR[c] for c in _CASES]
+    x = list(range(len(_CASES)))
+    bars = ax.bar(x, vals_gb, color=colors, width=0.65)
+    ax.set_xticks(x)
+    ax.set_xticklabels(labels, fontsize=9)
+    ax.set_ylabel("KV bytes per PE (GB, log)")
+    ax.set_yscale("log")
+    ax.set_title(
+        f"Per-PE KV memory at S_kv = {_HEADLINE_S_KV:,} tokens "
+        f"(across {_N_LAYERS} layers, FP16)",
+        fontsize=11,
+    )
+    ax.grid(axis="y", ls=":", alpha=0.5, which="both")
+    ax.axhline(_HEADROOM_GB, color="red", ls="--", lw=1.4,
+               label=f"KV headroom = {_HEADROOM_GB} GB / PE")
+    for bar, v_gb in zip(bars, vals_gb):
+        v_bytes = v_gb * (1 << 30)
+        fits = v_gb <= _HEADROOM_GB
+        ax.text(bar.get_x() + bar.get_width() / 2,
+                v_gb * 1.10,
+                _fmt_bytes(v_bytes) + (" ✓" if fits else " ✗"),
+                ha="center", va="bottom", fontsize=9,
+                color="green" if fits else "red", weight="bold")
+    ax.legend(loc="upper right", fontsize=9)
+
+
+def _plot_comm(ax, *, mode: str = "analytical") -> None:
+    """Per-PE comm panel.
+
+    mode = "analytical": single solid bars from per_token_bytes formula.
+    mode = "paired"    : analytical (solid) + simulator-measured
+                         (hatched) side-by-side per case, when the
+                         measurement JSON is available.
+    """
+    labels = [_CASE_LABEL[c] for c in _CASES]
+
+    wo_mb_list: list[float] = []
+    ffn_mb_list: list[float] = []
+    attn_mb: list[float] = []
+    for c in _CASES:
+        wo, ffn, attn = per_token_bytes(c, _HEADLINE_S_KV)
+        wo_mb_list.append(wo / (1 << 20))
+        ffn_mb_list.append(ffn / (1 << 20))
+        attn_mb.append(attn / (1 << 20))
+
+    measured = _load_measured() if mode == "paired" else None
+    paired = measured is not None
+    source_tag = "analytical (solid) vs simulator-measured (hatched)" \
+                 if paired else "analytical"
+
+    n_cases = len(_CASES)
+    x = list(range(n_cases))
+    bar_w = 0.36 if paired else 0.65
+    x_ana  = [xi - bar_w / 2 for xi in x] if paired else x
+    x_meas = [xi + bar_w / 2 for xi in x] if paired else None
+
+    # Analytical bars (solid).
+    ax.bar(x_ana, wo_mb_list,  width=bar_w,
+           color=_WO_COLOR,  edgecolor="black")
+    ax.bar(x_ana, ffn_mb_list, width=bar_w, bottom=wo_mb_list,
+           color=_FFN_COLOR, edgecolor="black")
+    bottoms_attn = [w + f for w, f in zip(wo_mb_list, ffn_mb_list)]
+    ax.bar(x_ana, attn_mb, width=bar_w, bottom=bottoms_attn,
+           color=_ATTN_COLOR, edgecolor="black")
+
+    # Measured bars (hatched) — same Wo+FFN base, attn from op_log.
+    meas_attn_mb: list[float] = []
+    if paired:
+        for i, c in enumerate(_CASES):
+            meas_total = measured.get(c, 0) / (1 << 20)
+            meas_attn_mb.append(
+                max(meas_total - wo_mb_list[i] - ffn_mb_list[i], 0.0))
+        ax.bar(x_meas, wo_mb_list,  width=bar_w,
+               color=_WO_COLOR,  edgecolor="black",
+               hatch="///", alpha=0.85)
+        ax.bar(x_meas, ffn_mb_list, width=bar_w, bottom=wo_mb_list,
+               color=_FFN_COLOR, edgecolor="black",
+               hatch="///", alpha=0.85)
+        ax.bar(x_meas, meas_attn_mb, width=bar_w, bottom=bottoms_attn,
+               color=_ATTN_COLOR, edgecolor="black",
+               hatch="///", alpha=0.85)
+
+    ax.set_xticks(x)
+    ax.set_xticklabels(labels, fontsize=9)
+    ax.set_ylabel("Comm per token per PE (MB, log)")
+    ax.set_yscale("log")
+    title = (
+        f"Per-PE communication at S_kv = {_HEADLINE_S_KV:,} tokens "
+        f"(decode S_q=1, B=1; {_N_LAYERS} layers)  —  {source_tag}"
+    )
+    if paired:
+        title += "\n(simulator measured at S_kv = 8K; " \
+                 "partial-score AR scaled ×128 to S_kv = 1M)"
+    ax.set_title(title, fontsize=10)
+    ax.grid(axis="y", ls=":", alpha=0.5, which="both")
+
+    totals_ana = [wo_mb_list[i] + ffn_mb_list[i] + attn_mb[i]
+                  for i in range(n_cases)]
+    ymax = max(totals_ana)
+    if paired:
+        ymax = max(ymax, max(
+            (measured.get(c, 0) / (1 << 20)) for c in _CASES))
+    ax.set_ylim(top=ymax * 22)
+
+    for i, c in enumerate(_CASES):
+        ana_bytes = totals_ana[i] * (1 << 20)
+        if paired:
+            meas_bytes = measured.get(c, 0)
+            top_y = max(totals_ana[i], meas_bytes / (1 << 20))
+            label = (f"ana: {_fmt_bytes(ana_bytes)}\n"
+                     f"sim: {_fmt_bytes(meas_bytes)}")
+        else:
+            top_y = totals_ana[i]
+            label = _fmt_bytes(ana_bytes)
+        ax.text(x[i], top_y * 1.5, label,
+                ha="center", va="bottom",
+                fontsize=8 if paired else 9,
+                weight="bold", linespacing=1.05)
+
+        # Attention-time AR descriptor — placed at the attention-segment
+        # midpoint, centered between the analytical and measured bars so
+        # it visually spans both. Coloured the same as the attention bar
+        # (no background box) so it lives within the case's bar zone.
+        # Wrapped narrow so each line fits inside the bar-pair width.
+        wrapped = textwrap.fill(_ATTN_DESC[c], width=14)
+        if attn_mb[i] > 0:
+            mid = bottoms_attn[i] + attn_mb[i] / 2
+            ax.text(x[i], mid, wrapped,
+                    ha="center", va="center",
+                    fontsize=7 if paired else 8,
+                    color="black", weight="bold",
+                    linespacing=1.0)
+        else:
+            ax.text(x[i], totals_ana[i] * 0.4, wrapped,
+                    ha="center",
+                    fontsize=7 if paired else 8,
+                    color="grey", style="italic",
+                    linespacing=1.0)
+
+    legend_handles = [
+        mpatches.Patch(facecolor=_WO_COLOR, edgecolor="black",
+                       label=f"Wo  AR  (× {_N_LAYERS} layers)"),
+        mpatches.Patch(facecolor=_FFN_COLOR, edgecolor="black",
+                       label=f"FFN AR  (× {_N_LAYERS} layers)"),
+        mpatches.Patch(facecolor=_ATTN_COLOR, edgecolor="black",
+                       label="Attn-time collective"),
+    ]
+    if paired:
+        legend_handles.append(
+            mpatches.Patch(facecolor="white", edgecolor="black",
+                           hatch="///", label="simulator-measured"))
+    ax.legend(handles=legend_handles, loc="upper right", fontsize=8,
+              framealpha=0.92)
+
+
+def main() -> Path:
+    _OUT_DIR.mkdir(parents=True, exist_ok=True)
+
+    # (a) Per-PE HBM budget — standalone PNG.
+    fig_b, ax_b = plt.subplots(figsize=(4.0, 6.0))
+    _plot_budget(ax_b)
+    fig_b.tight_layout()
+    out_b = _OUT_DIR / "gqa_long_ctx_6cases_hbm_budget.png"
+    fig_b.savefig(out_b, dpi=150)
+    plt.close(fig_b)
+    print(f"wrote {out_b}")
+
+    # (b) Combined 3-panel summary — HBM budget + KV memory + comm.
+    fig = plt.figure(figsize=(22.0, 6.5))
+    gs = fig.add_gridspec(1, 3, width_ratios=[0.7, 1.6, 1.6], wspace=0.22)
+    ax_b2 = fig.add_subplot(gs[0, 0])
+    ax_m  = fig.add_subplot(gs[0, 1])
+    ax_c  = fig.add_subplot(gs[0, 2])
+    _plot_budget(ax_b2)
+    _plot_memory(ax_m)
+    _plot_comm(ax_c)
+    fig.tight_layout()
+    out = _OUT_DIR / "gqa_long_ctx_6cases_summary.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    print(f"wrote {out}")
+
+    # (c) 2-panel companion (analytical only).
+    fig2 = plt.figure(figsize=(18.0, 6.5))
+    gs2 = fig2.add_gridspec(1, 2, width_ratios=[1.0, 1.0], wspace=0.18)
+    ax_m2 = fig2.add_subplot(gs2[0, 0])
+    ax_c2 = fig2.add_subplot(gs2[0, 1])
+    _plot_memory(ax_m2)
+    _plot_comm(ax_c2, mode="analytical")
+    fig2.tight_layout()
+    out2 = _OUT_DIR / "gqa_long_ctx_6cases_memory_comm_analytical.png"
+    fig2.savefig(out2, dpi=150)
+    plt.close(fig2)
+    print(f"wrote {out2}")
+
+    # (d) 2-panel companion — analytical vs simulator-measured paired.
+    fig3 = plt.figure(figsize=(19.0, 6.5))
+    gs3 = fig3.add_gridspec(1, 2, width_ratios=[1.0, 1.0], wspace=0.18)
+    ax_m3 = fig3.add_subplot(gs3[0, 0])
+    ax_c3 = fig3.add_subplot(gs3[0, 1])
+    _plot_memory(ax_m3)
+    _plot_comm(ax_c3, mode="paired")
+    fig3.tight_layout()
+    out3 = _OUT_DIR / "gqa_long_ctx_6cases_memory_comm_paired.png"
+    fig3.savefig(out3, dpi=150)
+    plt.close(fig3)
+    print(f"wrote {out3}")
+
+    # Paper-ready table to stdout.
+    print()
+    print(f"  {'Case':<27} {'KV/tok·PE':>12} {'KV @ 1M':>11} "
+          f"{'Max S_kv':>11} {'Comm @ 1M':>12}")
+    print("  " + "-" * 80)
+    for c in _CASES:
+        kv_per_tok_pe = _KV_PER_TOK_BYTES / _DIVISOR[c]
+        mem_1m = kv_per_pe_bytes(c, _HEADLINE_S_KV)
+        max_s = max_s_kv(c)
+        comm_1m = sum(per_token_bytes(c, _HEADLINE_S_KV))
+        label = _CASE_LABEL[c].replace(chr(10), " ")
+        print(f"  {label:<27} "
+              f"{kv_per_tok_pe / 1024:>9.3f} KB "
+              f"{_fmt_bytes(mem_1m):>11} "
+              f"{_fmt_tokens(max_s):>11} "
+              f"{_fmt_bytes(comm_1m):>12}")
+    return out
+
+
+if __name__ == "__main__":
+    main()

scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py

@@ -0,0 +1,248 @@
+"""Comparative figures for milestone-gqa-decode-long-ctx-4cases.
+
+Reads sweep_decode.json (emitted by the milestone-1h-gqa bench) and
+writes four PNGs into the same bench-output dir
+(src/kernbench/benches/1H_milestone_output/gqa/long_ctx/):
+
+  gqa_decode_long_ctx_6cases_latency.png       end-to-end latency per case
+  gqa_decode_long_ctx_6cases_traffic.png       ipcq/dma op-count breakdown
+  gqa_decode_long_ctx_6cases_memory.png        per-PE KV bytes per case
+  gqa_decode_long_ctx_6cases_parallelism.png   per-PE S_local (compute work)
+
+Filename still says "4cases" for backwards compat, but the script now
+covers all SIX kv-sharding strategies from the analytical chart
+(`gqa_4cases_summary.png`) — the original 4 plus the two new
+d_head-TP variants:
+
+  Case 1   Cube-Repl × PE-repl      (PE-TP doesn't shard KV)
+  Case 2   Cube-SP   × PE-repl
+  Case 3   Cube-Repl × PE-SP
+  Case 4   Cube-SP   × PE-TP-d_head      ← NEW
+  Case 5   Cube-TP-d_head × PE-SP        ← NEW
+  Case 6 ★ Cube-SP   × PE-SP             (Pareto-best)
+
+Run (after the bench):
+    GQA_DECODE_LONG_CTX_4CASES_RUN=1 python -m kernbench.cli.main run \\
+        --bench milestone-gqa-decode-long-ctx-4cases --topology topology.yaml
+    python scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py
+"""
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+import matplotlib
+
+matplotlib.use("Agg")
+import matplotlib.pyplot as plt  # noqa: E402
+
+_REPO_ROOT = Path(__file__).resolve().parents[2]
+# Sweep JSON + PNGs live together under the bench output dir.
+_FIG_DIR = (
+    _REPO_ROOT / "src" / "kernbench" / "benches"
+    / "1H_milestone_output" / "gqa" / "long_ctx"
+)
+_SWEEP_JSON = _FIG_DIR / "sweep_decode.json"
+
+# Panel name → (short label, case ordinal, accent flag) using the
+# analytical chart's memory-descending ordering. PE-TP doesn't shard
+# KV memory, so the cube_repl_pe_tp panel maps to Case 1 (no
+# sharding, KV-wise) and cube_sp_pe_tp panel maps to Case 2.
+_NORMAL, _OVERFLOW, _PARETO = "normal", "overflow", "pareto"
+
+_CASE_INFO = {
+    # panel name                                                       label                            ord  flag
+    "single_kv_group_decode_long_ctx_gqa_cube_repl_pe_tp":            ("Case 1\nCube-Repl × PE-repl",   1, _OVERFLOW),
+    "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp":              ("Case 2\nCube-SP × PE-repl",     2, _OVERFLOW),
+    "single_kv_group_decode_long_ctx_gqa_cube_repl_pe_sp":            ("Case 3\nCube-Repl × PE-SP",     3, _OVERFLOW),
+    "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp_dhead":        ("Case 4\nCube-SP × PE-TP-d_head", 4, _NORMAL),
+    "single_kv_group_decode_long_ctx_gqa_cube_tp_dhead_pe_sp":        ("Case 5\nCube-TP-d_head × PE-SP", 5, _NORMAL),
+    "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_sp":              ("Case 6 ★\nCube-SP × PE-SP",     6, _PARETO),
+}
+
+# Bar fill colour per flag (used by every panel).
+_FLAG_COLOR = {
+    _NORMAL:   "#888888",     # neutral grey
+    _OVERFLOW: "#c0504d",     # red — fails the per-PE HBM budget
+    _PARETO:   "#3b6ea5",     # blue — Pareto-best
+}
+
+
+def _load() -> list[dict]:
+    return json.loads(_SWEEP_JSON.read_text())["rows"]
+
+
+def _sorted_by_case(rows: list[dict]) -> list[dict]:
+    return sorted(rows, key=lambda r: _CASE_INFO[r["panel"]][1])
+
+
+def _bar_colors(rows: list[dict]) -> list[str]:
+    return [_FLAG_COLOR[_CASE_INFO[r["panel"]][2]] for r in rows]
+
+
+def _plot_latency(rows: list[dict]) -> Path:
+    rows = _sorted_by_case(rows)
+    labels = [_CASE_INFO[r["panel"]][0] for r in rows]
+    lat_us = [r["latency_ns"] / 1e3 for r in rows]
+    fig, ax = plt.subplots(figsize=(12.0, 4.8))
+    bars = ax.bar(labels, lat_us, color=_bar_colors(rows), width=0.6)
+    ax.set_ylabel("end-to-end latency (µs)")
+    ax.set_title(
+        "Long-context decode 6-cases — end-to-end latency per case\n"
+        "LLaMA-3.1-70B single-KV-head group (8 cubes × 8 PEs)"
+    )
+    ax.bar_label(bars, fmt="%.1f", padding=3, fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    ax.set_ylim(0, max(lat_us) * 1.15)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_decode_long_ctx_6cases_latency.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def _plot_traffic(rows: list[dict]) -> Path:
+    rows = _sorted_by_case(rows)
+    labels = [_CASE_INFO[r["panel"]][0] for r in rows]
+    x = list(range(len(rows)))
+    keys = ["ipcq_copy_count", "dma_read_count", "dma_write_count"]
+    disp = ["IPCQ copy", "DMA read", "DMA write"]
+    colors = ["#c0504d", "#9bbb59", "#8064a2"]
+    w = 0.25
+    fig, ax = plt.subplots(figsize=(11.0, 4.5))
+    for i, (k, d, c) in enumerate(zip(keys, disp, colors)):
+        vals = [r["op_log_summary"][k] for r in rows]
+        ax.bar([xi + (i - 1) * w for xi in x], vals, width=w, label=d, color=c)
+    ax.set_xticks(list(x))
+    ax.set_xticklabels(labels, fontsize=9)
+    ax.set_ylabel("op count")
+    ax.set_title("Long-context decode 6-cases — op-count breakdown per case")
+    ax.legend(fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_decode_long_ctx_6cases_traffic.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def _s_local_per_pe(panel: str, *, S_kv: int, C: int, P: int) -> int:
+    """S_local (token count) each PE attends over locally.
+
+      cube_repl_pe_tp  (Case 1): S_kv        (no sharding, KV-wise)
+      cube_sp_pe_tp    (Case 2): S_kv / C    (cube splits S_kv, PEs replicate)
+      cube_repl_pe_sp  (Case 3): S_kv / P
+      cube_sp_pe_tp_dhead (Case 4): S_kv / C (cube splits S_kv, PE splits d_head)
+      cube_tp_dhead_pe_sp (Case 5): S_kv / P (cube splits d_head, PE splits S_kv)
+      cube_sp_pe_sp    (Case 6 ★): S_kv / (C·P)
+    """
+    cube_splits_s = "cube_sp" in panel
+    pe_splits_s   = "pe_sp" in panel
+    S_per_cube = S_kv // C if cube_splits_s else S_kv
+    return S_per_cube // P if pe_splits_s else S_per_cube
+
+
+def _d_head_per_pe(panel: str, *, d_head: int, C: int, P: int) -> int:
+    """d_head dims each PE owns (Cases 4 and 5 shard d_head)."""
+    if "cube_tp_dhead" in panel:    # Case 5: cube shards d_head
+        return d_head // C
+    if "pe_tp_dhead" in panel:      # Case 4: PE shards d_head
+        return d_head // P
+    return d_head                   # Cases 1, 2, 3, 6: full d_head per PE
+
+
+def _active_pe_count(panel: str, *, C: int, P: int) -> int:
+    """Number of PEs doing non-idle attention work.
+
+      cube_repl_pe_tp  (Case 1): 1     (PE-TP idle for B=1; only one PE works)
+      cube_sp_pe_tp    (Case 2): C     (PE 0 of each cube; 7 PEs idle per cube)
+      cube_repl_pe_sp  (Case 3): C·P   (all PEs busy, cube-side redundant)
+      cube_sp_pe_tp_dhead (Case 4): C·P (PE shards d_head — all 64 active)
+      cube_tp_dhead_pe_sp (Case 5): C·P (PE shards S_kv — all active)
+      cube_sp_pe_sp    (Case 6 ★): C·P (all 64 PEs doing unique work)
+    """
+    if "cube_repl" in panel and "pe_tp" in panel and "dhead" not in panel:
+        return 1
+    if "cube_sp" in panel and "pe_tp" in panel and "dhead" not in panel:
+        return C
+    return C * P
+
+
+def _kv_bytes_per_pe(panel: str, *, S_kv: int, h_kv: int,
+                     d_head: int, C: int, P: int) -> int:
+    """KV bytes a single PE references (K + V, f16, 2 B/elem)."""
+    s_local = _s_local_per_pe(panel, S_kv=S_kv, C=C, P=P)
+    d_local = _d_head_per_pe(panel, d_head=d_head, C=C, P=P)
+    return 2 * s_local * h_kv * d_local * 2
+
+
+def _plot_memory(rows: list[dict]) -> Path:
+    """Per-PE KV bytes — Case 6 ★ wins (64-way split)."""
+    rows = _sorted_by_case(rows)
+    labels = [_CASE_INFO[r["panel"]][0] for r in rows]
+    mib_per_pe = [
+        _kv_bytes_per_pe(
+            r["panel"], S_kv=r["S_kv"], h_kv=r["h_kv"],
+            d_head=r["d_head"], C=r["C"], P=r["P"],
+        ) / (1024 * 1024)
+        for r in rows
+    ]
+    fig, ax = plt.subplots(figsize=(12.0, 4.8))
+    bars = ax.bar(labels, mib_per_pe, color=_bar_colors(rows), width=0.6)
+    ax.set_ylabel("KV bytes per PE (MiB, K + V, f16)")
+    ax.set_title(
+        "Long-context decode 6-cases — KV memory per PE\n"
+        "(one KV-head group; per-layer, per-token state)"
+    )
+    ax.bar_label(bars, fmt="%.3f", padding=3, fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    ax.set_ylim(0, max(mib_per_pe) * 1.15)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_decode_long_ctx_6cases_memory.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def _plot_parallelism(rows: list[dict]) -> Path:
+    """Total active PE-token compute load — exposes redundant-work cases."""
+    rows = _sorted_by_case(rows)
+    labels = [_CASE_INFO[r["panel"]][0] for r in rows]
+    total_work = [
+        _active_pe_count(r["panel"], C=r["C"], P=r["P"])
+        * _s_local_per_pe(r["panel"], S_kv=r["S_kv"], C=r["C"], P=r["P"])
+        for r in rows
+    ]
+    fig, ax = plt.subplots(figsize=(12.0, 4.8))
+    bars = ax.bar(labels, total_work, color=_bar_colors(rows), width=0.6)
+    ax.set_ylabel("active-PE × S_local (PE-tokens; lower ⇒ less wasted work)")
+    ax.set_title(
+        "Long-context decode 6-cases — total compute load across active PEs\n"
+        "(Case 3 replicates KV across 8 cubes → 8× wasted PE-tokens; "
+        "Case 6 ★ is fully parallel without replication)"
+    )
+    ax.bar_label(bars, fmt="%d", padding=3, fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    ax.set_ylim(0, max(total_work) * 1.15)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_decode_long_ctx_6cases_parallelism.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def main() -> None:
+    rows = _load()
+    _FIG_DIR.mkdir(parents=True, exist_ok=True)
+    p1 = _plot_latency(rows)
+    p2 = _plot_traffic(rows)
+    p3 = _plot_memory(rows)
+    p4 = _plot_parallelism(rows)
+    print(f"wrote {p1}")
+    print(f"wrote {p2}")
+    print(f"wrote {p3}")
+    print(f"wrote {p4}")
+
+
+if __name__ == "__main__":
+    main()

scripts/paper/paper_plot_gqa_decode_long_ctx_composite.py

@@ -0,0 +1,146 @@
+"""Comparative figure for the Case-6 composite-command decode study.
+
+Reads sweep_decode_composite.json (emitted by milestone-1h-gqa, sweep
+``composite``) and writes one two-panel PNG into the bench-output dir:
+
+  gqa_decode_long_ctx_composite.png
+    Left  — end-to-end decode latency (µs) vs context length, per command
+            form (primitive / composite / composite_extended).
+    Right — PE_CPU command count vs context length: the hand-tiled
+            primitive kernel issues O(n_tiles) commands (rises with
+            context), while the coarse composite forms issue O(1) and
+            *saturate* — PE_SCHEDULER absorbs the per-tile fan-out.
+
+The x-axis is the global context length S_kv; each PE owns
+S_local = S_kv/(C·P=64) tokens, so at the 1M production point every PE
+runs Q·Kᵀ of (G·T_q, d_head)·(d_head, 16384) and P·V of
+(G·T_q, 16384)·(16384, d_head).
+
+Run (after the bench):
+    GQA_1H_RUN=1 GQA_1H_SWEEPS=composite python -m kernbench.cli.main run \\
+        --bench milestone-1h-gqa --topology topology.yaml
+    python scripts/paper/paper_plot_gqa_decode_long_ctx_composite.py
+"""
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+import matplotlib
+
+matplotlib.use("Agg")
+import matplotlib.pyplot as plt  # noqa: E402
+
+_REPO_ROOT = Path(__file__).resolve().parents[2]
+_FIG_DIR = (
+    _REPO_ROOT / "src" / "kernbench" / "benches"
+    / "1H_milestone_output" / "gqa" / "long_ctx"
+)
+_SWEEP_JSON = _FIG_DIR / "sweep_decode_composite.json"
+_PAPER_FIG_DIR = (
+    _REPO_ROOT / "docs" / "report" / "1H-codesign-paper" / "figures"
+)
+
+_N_RANKS = 64  # C·P for the Case-6 64-way split.
+
+# variant key → (display label, colour, marker)
+_VARIANT_STYLE = {
+    "primitive":          ("primitive (tl.dot, hand-tiled)", "#c0504d", "o"),
+    "composite":          ("composite GEMM",                 "#3b6ea5", "s"),
+    "composite_extended": ("composite + softmax_merge",      "#4f8a4f", "^"),
+}
+_ORDER = ("primitive", "composite", "composite_extended")
+
+
+def _load() -> dict:
+    return json.loads(_SWEEP_JSON.read_text())
+
+
+def _series(rows: list[dict], variant: str, key: str):
+    """Sorted (S_kv, value) series for a variant, skipping null values
+    (latency is only measured over the tractable S_kv subset)."""
+    pts = sorted(
+        ((r["S_kv"], r[key]) for r in rows
+         if r["variant"] == variant and r.get(key) is not None),
+        key=lambda t: t[0],
+    )
+    return [p[0] for p in pts], [p[1] for p in pts]
+
+
+def _xticklabels(s_kvs: list[int]) -> list[str]:
+    out = []
+    for s in s_kvs:
+        if s >= 1 << 20:
+            out.append(f"{s // (1 << 20)}M")
+        else:
+            out.append(f"{s // 1024}K")
+    return out
+
+
+def main() -> None:
+    sweep = _load()
+    rows = sweep["rows"]
+    s_kv_op = sweep["s_kv_opcount"]
+    s_kv_lat = sweep["s_kv_latency"]
+
+    fig, (ax_lat, ax_cmd) = plt.subplots(1, 2, figsize=(13.0, 4.8))
+
+    for v in _ORDER:
+        label, color, marker = _VARIANT_STYLE[v]
+        xs, lat = _series(rows, v, "latency_ns")
+        ax_lat.plot(xs, [y / 1e3 for y in lat], marker=marker,
+                    color=color, label=label, lw=2)
+        xs, cmds = _series(rows, v, "pe_cpu_cmd_count")
+        ax_cmd.plot(xs, cmds, marker=marker, color=color, label=label, lw=2)
+
+    ax_lat.set_xticks(s_kv_lat)
+    ax_lat.set_xticklabels(_xticklabels(s_kv_lat))
+    ax_cmd.set_xticks(s_kv_op)
+    ax_cmd.set_xticklabels(_xticklabels(s_kv_op), fontsize=8)
+    for ax in (ax_lat, ax_cmd):
+        ax.set_xscale("log", base=2)
+        ax.set_xlabel(
+            r"context length $S_{kv}$ "
+            r"($S_{\mathrm{local}}=S_{kv}/64$ per PE)"
+        )
+        ax.grid(True, ls=":", alpha=0.5)
+        ax.legend(fontsize=9)
+
+    ax_lat.set_ylabel("end-to-end decode latency (µs)")
+    ax_lat.set_title(
+        "Case-6 decode latency per command form\n"
+        "(memory-bound: command form does not move the critical path)"
+    )
+    ax_cmd.set_ylabel("PE_CPU commands issued")
+    ax_cmd.set_title(
+        "PE_CPU command count per command form\n"
+        "(primitive O(n$_\\mathrm{tiles}$) rises; composite O(1) saturates)"
+    )
+    ax_cmd.axvline(1 << 20, color="#888", ls="--", lw=1, alpha=0.7)
+    ax_cmd.annotate("1M production\ncontext", xy=(1 << 20, 0),
+                    xytext=(1 << 18, 0.6), fontsize=8,
+                    textcoords=("data", "axes fraction"), ha="right",
+                    color="#555")
+
+    fig.suptitle(
+        "Case-6 (Cube-SP × PE-SP) long-context decode — use of composite "
+        "commands\nLLaMA-3.1-70B single-KV-head group (8 cubes × 8 PEs), "
+        "$T_q{=}1$",
+        fontsize=11,
+    )
+    fig.tight_layout(rect=(0, 0, 1, 0.94))
+
+    out = _FIG_DIR / "gqa_decode_long_ctx_composite.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    print(f"wrote {out}")
+
+    # Mirror into the paper figures dir (derived artifact).
+    if _PAPER_FIG_DIR.is_dir():
+        dst = _PAPER_FIG_DIR / out.name
+        dst.write_bytes(out.read_bytes())
+        print(f"copied {dst}")
+
+
+if __name__ == "__main__":
+    main()

scripts/paper/paper_plot_gqa_decode_streaming.py

@@ -0,0 +1,124 @@
+"""Comparative figure for the memory-bound decode-streaming composite study.
+
+Reads sweep_decode_streaming.json (emitted by milestone-1h-gqa, sweep
+``decode_streaming``) and writes one two-panel PNG:
+
+  gqa_decode_streaming.png
+    Left  — end-to-end single-rank decode latency (µs) vs per-rank context.
+    Right — achieved HBM bandwidth (GB/s) vs context, against the
+            256 GB/s per-rank roofline.
+
+The memory-bound mirror of the compute-bound prefill figure. With T_q=1
+the GEMMs are skinny (M=8) and the kernel is bound by streaming the KV
+cache. Isolating a single rank (no inter-CUBE reduce) reveals what the
+64-way Case-6 decode masks: the composite command still wins, not by
+feeding the MAC array but by keeping the DMA pipeline full — its
+scheduler-streamed concurrent tile DMAs extract ~230 GB/s (near the
+256 GB/s roofline) while the primitive kernel's blocking tl.dot serializes
+one tile DMA at a time and plateaus at ~166 GB/s. That bandwidth gap is a
+~25-28 % latency win that grows nowhere near prefill's compute-bound
+margin but is decidedly not zero.
+
+Run (after the bench):
+    GQA_1H_RUN=1 GQA_1H_SWEEPS=decode_streaming python -m kernbench.cli.main \\
+        run --bench milestone-1h-gqa --topology topology.yaml
+    python scripts/paper/paper_plot_gqa_decode_streaming.py
+"""
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+import matplotlib
+
+matplotlib.use("Agg")
+import matplotlib.pyplot as plt  # noqa: E402
+
+_REPO_ROOT = Path(__file__).resolve().parents[2]
+_FIG_DIR = (
+    _REPO_ROOT / "src" / "kernbench" / "benches"
+    / "1H_milestone_output" / "gqa" / "long_ctx"
+)
+_SWEEP_JSON = _FIG_DIR / "sweep_decode_streaming.json"
+_PAPER_FIG_DIR = (
+    _REPO_ROOT / "docs" / "report" / "1H-codesign-paper" / "figures"
+)
+
+# Per-rank HBM roofline: 8 pseudo-channels × 32 GB/s (topology.yaml
+# hbm_ctrl.num_pcs / pc_bw_gbs; = pe_dma_to_noc_bw_gbs).
+_PEAK_HBM_GBS = 256.0
+
+_VARIANT_STYLE = {
+    "primitive":          ("primitive (tl.dot, hand-tiled)", "#c0504d", "o"),
+    "composite":          ("composite GEMM",                 "#3b6ea5", "s"),
+    "composite_extended": ("composite + softmax_merge",      "#4f8a4f", "^"),
+}
+_ORDER = ("primitive", "composite", "composite_extended")
+
+
+def _ctx_label(c: int) -> str:
+    return f"{c // 1024}K" if c >= 1024 else str(c)
+
+
+def _series(rows, variant, key):
+    pts = sorted(((r["s_kv"], r[key]) for r in rows
+                  if r["variant"] == variant), key=lambda t: t[0])
+    return [p[0] for p in pts], [p[1] for p in pts]
+
+
+def main() -> None:
+    sweep = json.loads(_SWEEP_JSON.read_text())
+    rows = sweep["rows"]
+    ctxs = sweep["s_kv_points"]
+
+    fig, (ax_lat, ax_bw) = plt.subplots(1, 2, figsize=(13.0, 4.8))
+
+    for v in _ORDER:
+        label, color, marker = _VARIANT_STYLE[v]
+        xs, lat = _series(rows, v, "latency_ns")
+        ax_lat.plot(xs, [y / 1e3 for y in lat], marker=marker,
+                    color=color, label=label, lw=2)
+        xs, bw = _series(rows, v, "achieved_bw_gbs")
+        ax_bw.plot(xs, bw, marker=marker, color=color, label=label, lw=2)
+
+    for ax in (ax_lat, ax_bw):
+        ax.set_xscale("log", base=2)
+        ax.set_xticks(ctxs)
+        ax.set_xticklabels([_ctx_label(c) for c in ctxs])
+        ax.set_xlabel(r"per-rank context length $S_{kv}$ ($T_q{=}1$)")
+        ax.grid(True, ls=":", alpha=0.5)
+        ax.legend(fontsize=9)
+
+    ax_lat.set_ylabel("end-to-end decode latency (µs)")
+    ax_lat.set_title("Single-rank memory-bound decode latency per command form")
+
+    ax_bw.set_ylabel("achieved HBM bandwidth (GB/s)")
+    ax_bw.set_title(
+        "HBM bandwidth — composite keeps the DMA pipe full; primitive plateaus"
+    )
+    ax_bw.axhline(_PEAK_HBM_GBS, color="#888", ls="--", lw=1, alpha=0.7)
+    ax_bw.text(ctxs[0], _PEAK_HBM_GBS - 8, "256 GB/s roofline",
+               fontsize=8, color="#555", va="top")
+    ax_bw.set_ylim(0, _PEAK_HBM_GBS * 1.08)
+
+    fig.suptitle(
+        "Memory-bound decode streaming — use of composite commands\n"
+        "single-rank, GQA single-KV-head group ($h_q{=}8$, $d_{\\text{head}}"
+        "{=}128$); $M{=}8$ skinny, KV-streaming-bound",
+        fontsize=11,
+    )
+    fig.tight_layout(rect=(0, 0, 1, 0.92))
+
+    out = _FIG_DIR / "gqa_decode_streaming.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    print(f"wrote {out}")
+
+    if _PAPER_FIG_DIR.is_dir():
+        dst = _PAPER_FIG_DIR / out.name
+        dst.write_bytes(out.read_bytes())
+        print(f"copied {dst}")
+
+
+if __name__ == "__main__":
+    main()

scripts/paper/paper_plot_gqa_kv_sharding_diagram.py

@@ -0,0 +1,324 @@
+"""6-case KV-sharding tensor diagram (the slide-13 PNG export).
+
+Flat 2-D rectangles, one per sharding case, with:
+    Y axis = S_kv (vertical)        — Cube-SP / PE-SP slice it
+    X axis = d_head (horizontal)    — Cube-TP-d_head / PE-TP-d_head slice it
+
+Drops the batch axis entirely (decode: B = 1, T_q = 1). Same case set
+and visual encoding as slide 13 of GQA_full_deck.pptx; matplotlib
+renders it cleanly so the PNG sits next to the other GQA summary
+artifacts in 1H_milestone_output/gqa/long_ctx/.
+"""
+from __future__ import annotations
+
+from pathlib import Path
+
+import matplotlib.patches as mpatches
+import matplotlib.pyplot as plt
+
+_C = 8
+_P = 8
+
+_GROUP_FILLS = [
+    "#A5D8FF", "#B2F2BB", "#FFD8A8", "#FFC9C9",
+    "#D0BFFF", "#99E9F2", "#FCC2D7", "#FFEC99",
+]
+
+_ACC = {
+    "red":    "#E03131",
+    "orange": "#FD7E14",
+    "blue":   "#1C7ED6",
+    "green":  "#37B24D",
+}
+
+# (label, accent, kv, comm, overflow, encoding-flags, axis-spec)
+#   y_split   = 8 horizontal Y bands  (Cube-SP on S_kv)
+#   x_split   = 8 vertical X bands    (Cube-TP-d_head)
+#   pe_y      = 7 fine horizontal dividers within each Y band
+#   pe_x      = 7 fine vertical dividers within each X band
+#   axes      = small annotation under the chip naming the axes
+#               that the cube/PE actually shard, so the reader can
+#               parse Case 5 (where cube colour fills run X instead
+#               of Y, breaking the visual symmetry of the rest).
+_CASES = [
+    dict(label="Case 1\nCube-Repl / PE-repl", accent=_ACC["red"],
+         kv="40 GB", comm="1.2 MB", overflow=True,
+         y_split=False, x_split=False, pe_y=False, pe_x=False,
+         axes="Cube: replicated     PE: replicated"),
+    dict(label="Case 2\nCube-SP / PE-repl",   accent=_ACC["orange"],
+         kv="5 GB",  comm="3.8 MB", overflow=True,
+         y_split=True,  x_split=False, pe_y=False, pe_x=False,
+         axes="Cube → Y (S_kv)     PE: replicated"),
+    dict(label="Case 3\nCube-Repl / PE-SP",   accent=_ACC["orange"],
+         kv="5 GB",  comm="3.8 MB", overflow=True,
+         y_split=False, x_split=False, pe_y=True,  pe_x=False,
+         axes="Cube: replicated     PE → Y (S_kv)"),
+    dict(label="Case 4\nCube-SP / PE-TP-d_head", accent=_ACC["blue"],
+         kv="640 MB", comm="166 MB", overflow=False,
+         y_split=True,  x_split=False, pe_y=False, pe_x=True,
+         axes="Cube → Y (S_kv)     PE → X (d_head)"),
+    dict(label="Case 5\nCube-TP-d_head / PE-SP", accent=_ACC["blue"],
+         kv="640 MB", comm="166 MB", overflow=False,
+         y_split=False, x_split=True,  pe_y=True,  pe_x=False,
+         axes="Cube → X (d_head)     PE → Y (S_kv)"),
+    dict(label="Case 6 ★\nCube-SP / PE-SP",   accent=_ACC["green"],
+         kv="640 MB", comm="6.2 MB", overflow=False,
+         y_split=True,  x_split=False, pe_y=True,  pe_x=False,
+         axes="Cube → Y (S_kv)     PE → Y (S_kv)"),
+]
+
+_OUT_DIR = (
+    Path(__file__).resolve().parents[2]
+    / "src" / "kernbench" / "benches"
+    / "1H_milestone_output" / "gqa" / "long_ctx"
+)
+
+
+def _draw_panel(ax, cfg):
+    """Draw one case's 2-D KV-tensor rectangle into a panel ax."""
+    ax.set_xlim(0, 1)
+    ax.set_ylim(1, 0)        # Y points down (S_kv ↓)
+    ax.set_aspect("auto")
+    ax.set_xticks([])
+    ax.set_yticks([])
+
+    cube_repl = not cfg["y_split"] and not cfg["x_split"]
+    pe_repl   = not cfg["pe_y"]   and not cfg["pe_x"]
+
+    # Cube-level colour fill.
+    if cfg["y_split"] and not cfg["x_split"]:
+        # 8 horizontal Y bands.
+        for c in range(_C):
+            ax.add_patch(mpatches.Rectangle(
+                (0, c / _C), 1, 1 / _C,
+                facecolor=_GROUP_FILLS[c], edgecolor="black", linewidth=0.6))
+            ax.text(0.04, c / _C + 0.5 / _C, f"C{c}",
+                    ha="left", va="center", fontsize=8,
+                    fontweight="bold", color="#333")
+    elif cfg["x_split"] and not cfg["y_split"]:
+        # 8 vertical X bands.
+        for c in range(_C):
+            ax.add_patch(mpatches.Rectangle(
+                (c / _C, 0), 1 / _C, 1,
+                facecolor=_GROUP_FILLS[c], edgecolor="black", linewidth=0.6))
+            ax.text(c / _C + 0.5 / _C, 0.04, f"C{c}",
+                    ha="center", va="top", fontsize=8,
+                    fontweight="bold", color="#333")
+    else:
+        ax.add_patch(mpatches.Rectangle(
+            (0, 0), 1, 1,
+            facecolor="#F5F5F5", edgecolor="black", linewidth=0.8))
+        ax.text(0.5, 0.5, "× 8 cubes\nfull KV",
+                ha="center", va="center",
+                fontsize=10, fontweight="bold",
+                fontstyle="italic", color="#666")
+
+    # PE-level fine dividers — distinguished from cube boundaries by
+    # using a dashed style + slightly stronger contrast. This is what
+    # makes Case 5's PE-SP (horizontal lines across vertical cube
+    # bands) read as "different axis from the cubes" at a glance.
+    if cfg["pe_y"]:
+        outer = _C if cfg["y_split"] else 1
+        band = 1 / outer
+        for o in range(outer):
+            for p in range(1, _P):
+                y = o * band + band * p / _P
+                ax.axhline(y, color="#222", linewidth=0.8,
+                           linestyle=(0, (3, 2)), alpha=0.75)
+    if cfg["pe_x"]:
+        outer = _C if cfg["x_split"] else 1
+        band = 1 / outer
+        for o in range(outer):
+            for p in range(1, _P):
+                x = o * band + band * p / _P
+                ax.axvline(x, color="#222", linewidth=0.8,
+                           linestyle=(0, (3, 2)), alpha=0.75)
+
+    # Heavy outline on top.
+    ax.add_patch(mpatches.Rectangle(
+        (0, 0), 1, 1, facecolor="none",
+        edgecolor="black", linewidth=1.2))
+
+    # Replication badges — small text-only badges in the corners of
+    # the rectangle, no ghost-card stacking (which mis-reads as a
+    # larger enclosing tensor).
+    badges: list[str] = []
+    if cube_repl:
+        badges.append("× 8 cube copies")
+    if pe_repl and (cfg["y_split"] or cfg["x_split"]):
+        # Cube is sharded but PEs in each cube replicate that shard.
+        badges.append("× 8 PEs / cube replicate")
+    elif pe_repl and cube_repl:
+        # Both replicated — PE replication adds to the cube one.
+        badges.append("× 8 PEs / cube replicate")
+    if badges:
+        ax.text(0.98, 0.02, "\n".join(badges),
+                ha="right", va="top", fontsize=7,
+                fontweight="bold", color="#444",
+                fontstyle="italic",
+                bbox=dict(facecolor="white", edgecolor="#888",
+                          boxstyle="round,pad=0.20", linewidth=0.5))
+
+
+def _make_table_png() -> Path:
+    """Slide-14 companion table: per-PE memory + comm for all 6 cases."""
+    headers = ["Case", "Sharding", "KV / PE", "Fit",
+               "Comm/tok\n(analytical)", "Notes"]
+    rows = [
+        ("Case 1",   "Cube-Repl  ·  PE-repl",      "40 GB",  "✗",
+         "1.2 MB",
+         "no sharding —\nfull KV on every PE"),
+        ("Case 2",   "Cube-SP    ·  PE-repl",      "5 GB",   "✗",
+         "3.8 MB",
+         "cube-axis\nsharded only"),
+        ("Case 3",   "Cube-Repl  ·  PE-SP",        "5 GB",   "✗",
+         "3.8 MB",
+         "PE-axis\nsharded only"),
+        ("Case 4",   "Cube-SP    ·  PE-TP-d_head", "640 MB", "✓",
+         "166 MB",
+         "d_head split intra-cube\npartial-score AR ∝ S_kv"),
+        ("Case 5",   "Cube-TP-d_head  ·  PE-SP",   "640 MB", "✓",
+         "166 MB",
+         "d_head split inter-cube\npartial-score AR on UCIe"),
+        ("Case 6 ★", "Cube-SP    ·  PE-SP",        "640 MB", "✓",
+         "6.2 MB",
+         "S_kv split both axes\n(m,ℓ,O) AR only"),
+    ]
+    accents = [_ACC["red"], _ACC["orange"], _ACC["orange"],
+               _ACC["blue"], _ACC["blue"], _ACC["green"]]
+
+    fig, ax = plt.subplots(figsize=(15.0, 5.0))
+    ax.set_axis_off()
+
+    cell_data = [headers] + [list(r) for r in rows]
+    tbl = ax.table(cellText=cell_data,
+                   colWidths=[0.07, 0.20, 0.09, 0.05, 0.14, 0.28],
+                   cellLoc="center", loc="center")
+    tbl.auto_set_font_size(False)
+    tbl.set_fontsize(10.5)
+    tbl.scale(1.0, 2.4)
+
+    n_cols = len(headers)
+    n_rows = len(rows) + 1   # +1 header
+    # Header styling.
+    for ci in range(n_cols):
+        cell = tbl[(0, ci)]
+        cell.set_facecolor("#1F4E79")
+        cell.set_text_props(color="white", weight="bold")
+        cell.set_edgecolor("#1F4E79")
+    # Body styling.
+    for ri, row in enumerate(rows, start=1):
+        is_pareto = row[0].endswith("★")
+        row_fill = "#E8F5E9" if is_pareto else (
+            "white" if ri % 2 == 1 else "#F5F5F7")
+        # Case-name cell uses accent.
+        case_cell = tbl[(ri, 0)]
+        case_cell.set_facecolor(accents[ri - 1])
+        case_cell.set_text_props(color="white", weight="bold")
+        # Remaining cells.
+        for ci in range(1, n_cols):
+            cell = tbl[(ri, ci)]
+            cell.set_facecolor(row_fill)
+            txt_kwargs = {"weight": "bold" if is_pareto else "normal",
+                          "color": "#333"}
+            if ci == 2:                              # KV / PE
+                txt_kwargs["color"] = (
+                    "#C62828" if row[3] == "✗" else "#2E7D32")
+                txt_kwargs["weight"] = "bold"
+            if ci == 3:                              # Fit
+                txt_kwargs["color"] = (
+                    "#C62828" if row[3] == "✗" else "#2E7D32")
+                txt_kwargs["weight"] = "bold"
+            cell.set_text_props(**txt_kwargs)
+        # Last-column (Notes) cells left-aligned for readability.
+        tbl[(ri, n_cols - 1)].get_text().set_ha("left")
+
+    # Force left-align on the Notes header too.
+    tbl[(0, n_cols - 1)].get_text().set_ha("left")
+
+    fig.suptitle(
+        "GQA decode KV-sharding — per-PE memory & communication\n"
+        "(LLaMA 70B GQA single KV-head group  ·  S_kv = 1 M, FP16, "
+        "80 layers)",
+        fontsize=11.5, y=0.94,
+    )
+    out = _OUT_DIR / "gqa_long_ctx_6cases_kv_sharding_table.png"
+    fig.savefig(out, dpi=150, bbox_inches="tight")
+    plt.close(fig)
+    print(f"wrote {out}")
+    return out
+
+
+def main() -> Path:
+    _OUT_DIR.mkdir(parents=True, exist_ok=True)
+    n = len(_CASES)
+    fig = plt.figure(figsize=(20.0, 7.0))
+    # Three rows per column: case chip · axis-spec annotation · rectangle.
+    gs = fig.add_gridspec(3, n,
+                          height_ratios=[0.55, 0.32, 8.5],
+                          hspace=0.05, wspace=0.20,
+                          left=0.04, right=0.99,
+                          top=0.93, bottom=0.06)
+
+    for i, cfg in enumerate(_CASES):
+        # Top: case chip header.
+        ax_chip = fig.add_subplot(gs[0, i])
+        ax_chip.set_xticks([])
+        ax_chip.set_yticks([])
+        for spine in ax_chip.spines.values():
+            spine.set_visible(False)
+        ax_chip.add_patch(mpatches.Rectangle(
+            (0, 0), 1, 1, transform=ax_chip.transAxes,
+            facecolor=cfg["accent"], edgecolor=cfg["accent"]))
+        ax_chip.text(0.5, 0.5, cfg["label"],
+                     ha="center", va="center",
+                     fontsize=10, fontweight="bold",
+                     color="white")
+
+        # Middle: axis-spec annotation — names which axis the cube
+        # shards on and which axis the PE shards on (essential for
+        # parsing Case 5 where the cube colour fills run X instead
+        # of Y, breaking the visual symmetry of the rest).
+        ax_axes = fig.add_subplot(gs[1, i])
+        ax_axes.set_xticks([])
+        ax_axes.set_yticks([])
+        for spine in ax_axes.spines.values():
+            spine.set_visible(False)
+        ax_axes.add_patch(mpatches.Rectangle(
+            (0, 0), 1, 1, transform=ax_axes.transAxes,
+            facecolor="#F5F5F7", edgecolor="#CCCCCC",
+            linewidth=0.6))
+        ax_axes.text(0.5, 0.5, cfg["axes"],
+                     ha="center", va="center",
+                     fontsize=8.5, fontweight="bold",
+                     color="#1F4E79")
+
+        # Bottom: the tensor rectangle.
+        ax = fig.add_subplot(gs[2, i])
+        _draw_panel(ax, cfg)
+        ax.set_xlabel("X : d_head = 128  →",
+                      fontsize=9, fontweight="bold",
+                      fontstyle="italic", color="#1F4E79")
+        ax.set_ylabel("Y : S_kv = 1 M  ↓",
+                      fontsize=9, fontweight="bold",
+                      fontstyle="italic", color="#1F4E79")
+
+    fig.suptitle(
+        "GQA decode KV-tensor sharding — 6 cases  ·  "
+        "LLaMA 70B GQA single KV-head group  ·  "
+        "C = 8 cubes × P = 8 PEs  ·  S_kv = 1 M, FP16, 80 layers",
+        fontsize=12, y=0.99,
+    )
+
+    out = _OUT_DIR / "gqa_long_ctx_6cases_kv_sharding_diagram.png"
+    fig.savefig(out, dpi=150, bbox_inches="tight")
+    plt.close(fig)
+    print(f"wrote {out}")
+
+    # Companion table PNG (slide-14 export).
+    _make_table_png()
+    return out
+
+
+if __name__ == "__main__":
+    main()