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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

docs/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()

scripts/paper/paper_plot_gqa_prefill_compute_bound.py

@@ -0,0 +1,111 @@
+"""Comparative figure for the compute-bound prefill composite study.
+
+Reads sweep_prefill_compute_bound.json (emitted by milestone-1h-gqa,
+sweep ``prefill_cb``) and writes one two-panel PNG:
+
+  gqa_prefill_compute_bound.png
+    Left  — end-to-end prefill latency (µs) vs context length.
+    Right — MAC utilization (achieved / 8 TFLOP·s⁻¹ per-PE peak) vs context.
+
+Unlike memory-bound decode (where command form is latency-neutral), in
+compute-bound prefill the composite command keeps the MAC array fed by
+streaming DMA↔compute per HW tile, so it wins on both latency and
+utilization — and the margin grows with context (deeper P·V reduction =
+more tiles to pipeline).
+
+Run (after the bench):
+    GQA_1H_RUN=1 GQA_1H_SWEEPS=prefill_cb python -m kernbench.cli.main run \\
+        --bench milestone-1h-gqa --topology topology.yaml
+    python scripts/paper/paper_plot_gqa_prefill_compute_bound.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_prefill_compute_bound.json"
+_PAPER_FIG_DIR = (
+    _REPO_ROOT / "docs" / "report" / "1H-codesign-paper" / "figures"
+)
+
+_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["ctx_len"], 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["ctx_points"]
+
+    fig, (ax_lat, ax_util) = 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, util = _series(rows, v, "mac_util")
+        ax_util.plot(xs, [u * 100 for u in util], marker=marker,
+                     color=color, label=label, lw=2)
+
+    for ax in (ax_lat, ax_util):
+        ax.set_xscale("log", base=2)
+        ax.set_xticks(ctxs)
+        ax.set_xticklabels([_ctx_label(c) for c in ctxs])
+        ax.set_xlabel(r"context length (= $T_q$ = $S_{kv}$)")
+        ax.grid(True, ls=":", alpha=0.5)
+        ax.legend(fontsize=9)
+
+    ax_lat.set_ylabel("end-to-end prefill latency (µs)")
+    ax_lat.set_title("Compute-bound prefill latency per command form")
+    ax_util.set_ylabel("MAC utilization (% of 8 TFLOP·s⁻¹ peak)")
+    ax_util.set_title(
+        "MAC utilization — composite keeps the array fed; primitive starves"
+    )
+    ax_util.axhline(100, color="#888", ls="--", lw=1, alpha=0.6)
+
+    fig.suptitle(
+        "Compute-bound prefill attention — use of composite commands\n"
+        "single-rank, GQA single-KV-head group ($h_q{=}8$, $d_{\\text{head}}"
+        "{=}128$); $M{=}8T_q$ tile-filling",
+        fontsize=11,
+    )
+    fig.tight_layout(rect=(0, 0, 1, 0.92))
+
+    out = _FIG_DIR / "gqa_prefill_compute_bound.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_prefill_long_ctx_4cases.py

@@ -0,0 +1,253 @@
+"""Comparative figures for milestone-gqa-prefill-long-ctx-4cases.
+
+Mirror of paper_plot_gqa_decode_long_ctx_4cases but for the prefill
+variant. Reads sweep.json (emitted by ``kernbench run --bench
+milestone-gqa-prefill-long-ctx-4cases``) and writes four PNGs into
+``docs/report/1H-codesign-paper/figures/``:
+
+  gqa_prefill_long_ctx_4cases_latency.png       end-to-end latency per case
+  gqa_prefill_long_ctx_4cases_traffic.png       ipcq/dma op-count breakdown
+  gqa_prefill_long_ctx_4cases_memory.png        per-PE KV bytes per case
+  gqa_prefill_long_ctx_4cases_parallelism.png   active-PE × S_local work load
+
+Run (after the bench):
+    GQA_PREFILL_LONG_CTX_4CASES_RUN=1 python -m kernbench.cli.main run \\
+        --bench milestone-gqa-prefill-long-ctx-4cases --topology topology.yaml
+    python scripts/paper/paper_plot_gqa_prefill_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_prefill.json"
+
+# Panel name → (short label, case ordinal for left-to-right plot order).
+_CASE_INFO = {
+    "single_kv_group_prefill_long_ctx_gqa_cube_sp_pe_tp": (
+        "Case 1\nCube-SP × PE-TP", 1),
+    "single_kv_group_prefill_long_ctx_gqa_cube_repl_pe_tp": (
+        "Case 2\nCube-Repl × PE-TP", 2),
+    "single_kv_group_prefill_long_ctx_gqa_cube_repl_pe_sp": (
+        "Case 3\nCube-Repl × PE-SP", 3),
+    "single_kv_group_prefill_long_ctx_gqa_cube_sp_pe_sp": (
+        "Case 4 ★\nCube-SP × PE-SP", 4),
+}
+
+
+def _load() -> list[dict]:
+    return json.loads(_SWEEP_JSON.read_text())["rows"]
+
+
+def _sorted_by_case(rows: list[dict]) -> list[dict]:
+    return sorted(rows, key=lambda r: _CASE_INFO[r["panel"]][1])
+
+
+def _plot_latency(rows: list[dict]) -> Path:
+    rows = _sorted_by_case(rows)
+    labels = [_CASE_INFO[r["panel"]][0] for r in rows]
+    lat_us = [r["latency_ns"] / 1e3 for r in rows]
+    colors = ["#888", "#888", "#888", "#3b6ea5"]  # Case 4 highlighted
+    fig, ax = plt.subplots(figsize=(8.0, 4.5))
+    bars = ax.bar(labels, lat_us, color=colors, width=0.6)
+    ax.set_ylabel("end-to-end latency (µs)")
+    ax.set_title(
+        "Long-context prefill 4-cases — end-to-end latency per case\n"
+        "LLaMA-3.1-70B single-KV-head group (8 cubes × 8 PEs)"
+    )
+    ax.bar_label(bars, fmt="%.1f", padding=3, fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    ax.set_ylim(0, max(lat_us) * 1.15)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_prefill_long_ctx_4cases_latency.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def _plot_traffic(rows: list[dict]) -> Path:
+    rows = _sorted_by_case(rows)
+    labels = [_CASE_INFO[r["panel"]][0] for r in rows]
+    x = list(range(len(rows)))
+    keys = ["ipcq_copy_count", "dma_read_count", "dma_write_count"]
+    disp = ["IPCQ copy", "DMA read", "DMA write"]
+    colors = ["#c0504d", "#9bbb59", "#8064a2"]
+    w = 0.25
+    fig, ax = plt.subplots(figsize=(9.0, 4.5))
+    for i, (k, d, c) in enumerate(zip(keys, disp, colors)):
+        vals = [r["op_log_summary"][k] for r in rows]
+        ax.bar([xi + (i - 1) * w for xi in x], vals, width=w, label=d, color=c)
+    ax.set_xticks(list(x))
+    ax.set_xticklabels(labels, fontsize=9)
+    ax.set_ylabel("op count")
+    ax.set_title("Long-context prefill 4-cases — op-count breakdown per case")
+    ax.legend(fontsize=9)
+    ax.grid(axis="y", ls=":", alpha=0.5)
+    fig.tight_layout()
+    out = _FIG_DIR / "gqa_prefill_long_ctx_4cases_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.
+
+    Encodes the cube/pe sharding axes from the panel name:
+      cube_sp_pe_tp (Case 1): S_kv / C       (pe=replicate within cube)
+      cube_repl_pe_tp (Case 2): S_kv         (full S_kv per active PE)
+      cube_repl_pe_sp (Case 3): S_kv / P     (pe=row_wise within cube)
+      cube_sp_pe_sp (Case 4): S_kv / (C·P)   (★ 64-way split)
+    """
+    S_per_cube = S_kv if "cube_repl" in panel else S_kv // C
+    return S_per_cube // P if "pe_sp" in panel else S_per_cube
+
+
+def _active_pe_count(panel: str, *, C: int, P: int) -> int:
+    """Number of PEs doing non-idle attention work.
+
+    Prefill T_q≫1 means PE-TP is *useful* (not wasted as in decode):
+      cube_sp_pe_tp (Case 1): C·P   (T_q sharded across all 64 ranks)
+      cube_repl_pe_tp (Case 2): P   (CUBE 0 only; T_q sharded across its P PEs)
+      cube_repl_pe_sp (Case 3): C·P (all PEs busy but cubes redundant)
+      cube_sp_pe_sp (Case 4): C·P   (all 64 PEs doing unique work)
+    """
+    if "cube_repl" in panel and "pe_tp" in panel:
+        return P
+    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)
+    return 2 * s_local * h_kv * d_head * 2
+
+
+def _plot_memory(rows: list[dict]) -> Path:
+    """Per-PE KV bytes — Case 4 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
+    ]
+    colors = ["#888", "#c0504d", "#888", "#3b6ea5"]
+    fig, ax = plt.subplots(figsize=(8.0, 4.5))
+    bars = ax.bar(labels, mib_per_pe, color=colors, width=0.6)
+    ax.set_ylabel("KV bytes per PE (MiB, K + V, f16)")
+    ax.set_title(
+        "Long-context prefill 4-cases — KV memory per PE\n"
+        "(one KV-head group; per-layer, full S_kv 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_prefill_long_ctx_4cases_memory.png"
+    fig.savefig(out, dpi=150)
+    plt.close(fig)
+    return out
+
+
+def _t_q_per_pe(panel: str, *, T_q: int, C: int, P: int) -> int:
+    """T_q row count each active PE computes attention for.
+
+      cube_sp_pe_tp (Case 1): T_q / (C·P)  (T_q sharded across all 64 ranks)
+      cube_repl_pe_tp (Case 2): T_q / P    (T_q sharded across CUBE 0's P PEs)
+      cube_repl_pe_sp (Case 3): T_q        (Q replicated on every PE)
+      cube_sp_pe_sp (Case 4): T_q / C      (Q sharded by cube, replicated within)
+    """
+    if "cube_sp" in panel and "pe_tp" in panel:
+        return T_q // (C * P)
+    if "cube_repl" in panel and "pe_tp" in panel:
+        return T_q // P
+    if "cube_repl" in panel and "pe_sp" in panel:
+        return T_q
+    return T_q // C  # cube_sp_pe_sp
+
+
+def _s_kv_processed_per_pe(panel: str, *, S_kv: int, C: int, P: int) -> int:
+    """S_kv tokens each PE actually processes attention over.
+
+    Differs from ``_s_local_per_pe`` (OWNED KV bytes): for cases with
+    a Ring (Case 1, 4) each PE sees C ring steps so processes more
+    tokens than it locally owns.
+
+      cube_sp_pe_tp (Case 1): S_kv          (Ring + pe=replicate within cube)
+      cube_repl_pe_tp (Case 2): S_kv        (full KV per PE)
+      cube_repl_pe_sp (Case 3): S_kv / P    (pe=row_wise; no Ring)
+      cube_sp_pe_sp (Case 4): S_kv / P      (Ring restores full S_kv/P per PE)
+    """
+    if "cube_sp" in panel and "pe_tp" in panel:
+        return S_kv
+    if "cube_repl" in panel and "pe_tp" in panel:
+        return S_kv
+    return S_kv // P  # both pe_sp cases
+
+
+def _plot_parallelism(rows: list[dict]) -> Path:
+    """Total compute work (PE × T_q × S_kv token-pairs) — exposes Case 3's
+    redundancy. Cases 1, 2, 4 all do the same total work (correct
+    attention over T_q × S_kv); Case 3 does C× more (cubes redundantly
+    repeat the same compute because K/V is cube-replicated).
+    """
+    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"])
+        * _t_q_per_pe(r["panel"], T_q=r["T_q"], C=r["C"], P=r["P"])
+        * _s_kv_processed_per_pe(
+            r["panel"], S_kv=r["S_kv"], C=r["C"], P=r["P"],
+        )
+        for r in rows
+    ]
+    colors = ["#888", "#888", "#c0504d", "#3b6ea5"]  # Case 3 red, Case 4 highlighted
+    fig, ax = plt.subplots(figsize=(8.0, 4.5))
+    bars = ax.bar(labels, total_work, color=colors, width=0.6)
+    ax.set_ylabel(
+        "total compute (PE × T_q × S_kv token-pairs; lower ⇒ less wasted work)"
+    )
+    ax.set_title(
+        "Long-context prefill 4-cases — total compute work across active PEs\n"
+        "(Case 3 replicates K/V across 8 cubes ⇒ 8× redundant compute)"
+    )
+    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_prefill_long_ctx_4cases_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_sip_architecture.py

@@ -0,0 +1,91 @@
+"""Re-emit sip_view.svg in an academic palette (white background, black
+strokes) and convert it to PDF for the 1H-codesign-paper Figure 1.
+
+Source of truth: docs/diagrams/sip_view.svg (generated by
+src/kernbench/topology/visualizer.py, dark-ish theme).
+
+Treatment per user request: keep the layout intact, just force a white
+canvas and turn every stroke black; promote faint label text to black so
+all annotations stay legible on white.
+"""
+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" / "sip_view.svg"
+OUT_DIR = REPO / "docs" / "report" / "1H-codesign-paper" / "figures"
+OUT_SVG = OUT_DIR / "sip_architecture.svg"
+OUT_PDF = OUT_DIR / "sip_architecture.pdf"
+
+# Drop the "SIP VIEW" title; the figure caption already names the level.
+TITLE_REMOVE: list[tuple[str, str]] = [
+    ('  <text x="324" y="18" text-anchor="middle" font-family="monospace" '
+     'font-size="14" font-weight="bold" fill="#1e293b">SIP VIEW</text>\n',
+     ''),
+]
+
+COLOR_MAP: list[tuple[str, str]] = [
+    # canvas background slate-50 -> pure white
+    ('fill="#f8fafc"', 'fill="#ffffff"'),
+    # all strokes -> black
+    ('stroke="#3b82f6"', 'stroke="#000000"'),  # UCIe mesh blue lines
+    ('stroke="#475569"', 'stroke="#000000"'),  # cube block borders
+    ('stroke="#0ea5e9"', 'stroke="#000000"'),  # I/O sky-blue lines
+    # link annotation text slate-500 -> black for readability
+    ('fill="#64748b"', 'fill="#000000"'),
+]
+
+# Font-size bumps so labels survive LaTeX \linewidth scaling at the
+# half-text-width subfigure. CUBE block labels overflow the 48px block
+# rects, which is acceptable here.
+FONT_MAP: dict[str, str] = {
+    "7":  "10",
+    "14": "17",
+}
+
+# Tighten whitespace: cube grid occupies x=[84,564], y=[128,520]; IO
+# chiplet sits around y~50. With the title removed, crop top to y=40 so
+# the IO chiplet keeps a small headroom. Crop 70px on each side and 113
+# px from bottom.
+LAYOUT_FIXUP: list[tuple[str, str]] = [
+    ('<svg xmlns="http://www.w3.org/2000/svg" width="648" height="648" '
+     'viewBox="0 0 648 648">',
+     '<svg xmlns="http://www.w3.org/2000/svg" width="508" height="495" '
+     'viewBox="70 40 508 495">'),
+]
+
+
+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 TITLE_REMOVE + COLOR_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()