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>

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

@@ -252,6 +252,74 @@ composite (가장 큰 것이 decode opt2 의 `#2` ~322 bytes) 보다 훨씬 위
 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 표 유지

docs/adr/ADR-0014-dev-pe-pipeline-execution-model.md

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

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

@@ -270,6 +270,80 @@ 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

docs/diagrams/allreduce_latency_plots/AllReduce_LRAB_2Dtorus_6SiP_2x3_with_TCM_SRAM_HBM.csv

@@ -1,13 +1,13 @@
 buffer_kind,sip_topology,n_sips,n_elem,bytes_per_pe,latency_ns
-hbm,torus_2d,6,128,256,2120.040000000012
+hbm,torus_2d,6,128,256,2345.040000000012
-hbm,torus_2d,6,1024,2048,2717.2783333333473
+hbm,torus_2d,6,1024,2048,2942.2783333333473
-hbm,torus_2d,6,8192,16384,7315.184999999989
+hbm,torus_2d,6,8192,16384,7540.184999999989
-hbm,torus_2d,6,32768,65536,23081.26500000037
+hbm,torus_2d,6,32768,65536,23306.26500000037
-sram,torus_2d,6,128,256,2060.040000000012
+sram,torus_2d,6,128,256,2285.040000000012
-sram,torus_2d,6,1024,2048,2909.2783333333473
+sram,torus_2d,6,1024,2048,3134.2783333333527
-sram,torus_2d,6,8192,16384,9523.184999999869
+sram,torus_2d,6,8192,16384,9748.184999999869
-sram,torus_2d,6,32768,65536,32201.265000000385
+sram,torus_2d,6,32768,65536,32426.265000000385
-tcm,torus_2d,6,128,256,1964.040000000012
+tcm,torus_2d,6,128,256,2189.040000000012
-tcm,torus_2d,6,1024,2048,2477.2783333333473
+tcm,torus_2d,6,1024,2048,2702.2783333333473
-tcm,torus_2d,6,8192,16384,6403.185000000109
+tcm,torus_2d,6,8192,16384,6628.18500000005
-tcm,torus_2d,6,32768,65536,19865.265000000378
+tcm,torus_2d,6,32768,65536,20090.265000000378

docs/diagrams/allreduce_latency_plots/summary.csv

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

docs/diagrams/pe2pe_latency_plots/summary.csv

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

docs/report/1H-codesign-paper/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/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

@@ -13,6 +13,7 @@
 \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}

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

@@ -25,9 +25,12 @@ serves as the common evaluation platform for all mechanisms and kernels
 discussed in this study.
 
 This report focuses on Grouped-Query Attention (GQA), one of the most
-performance- and bandwidth-critical components of LLM inference. Modern
+performance- and bandwidth-critical components of LLM inference. GQA
-decoder-only models such as Llama~3 and Mistral have largely transitioned
+dominates inference-time memory traffic and KV-cache capacity in modern
-from GPT-3-style multi-head attention (MHA) to GQA, in which multiple
+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

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

@@ -1,3 +1,64 @@
+\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}
 
@@ -8,7 +69,7 @@ execution model it presents to a kernel writer, how its latency model
 turns that execution into a number, and the concrete hardware
 configuration used for every experiment that follows.
 
-\subsection{Why KernBench: source-level kernels without a software stack}
+\subsection{Why KernBench}
 \label{sec:why}
 
 In a production end-to-end (E2E) stack, kernel performance is entangled
@@ -39,66 +100,37 @@ KernBench models AHBM as a hierarchy of SIPs, CUBEs, and processing
 elements (PEs). At the system level, multiple SIPs are connected
 through inter-package links, while each SIP contains a collection of
 CUBEs joined by an on-package interconnect
-(Fig.~\ref{fig:sip-arch}).
+(Fig.~\ref{fig:hw-arch}\subref{fig:sip-arch}).
-
-\begin{figure}[t]
-\centering
-\includegraphics[width=\linewidth]{sip_architecture.pdf}
-\caption{Modeled hardware graph at the \emph{SIP} level (one
-example configuration; specific parameters in \S\ref{sec:hw} /
-Table~\ref{tab:hw}). The SIP holds a $4{\times}4$ mesh of CUBEs and
-an IO chiplet, with each line a directed link labelled by its
-physical distance and bandwidth.}
-\label{fig:sip-arch}
-\end{figure}
 
 Each CUBE contains eight PEs, shared SRAM, HBM controllers, an
 \textsf{M\_CPU} control processor, and an intra-CUBE router mesh
-(Fig.~\ref{fig:cube-arch}). Together these components form the
+(Fig.~\ref{fig:hw-arch}\subref{fig:cube-arch}). Together these
-execution substrate for all kernels evaluated in this report.
+components form the execution substrate for all kernels evaluated in
-
+this report. This organization reflects the memory-centric nature of
-\begin{figure}[t]
+AHBM: each PE is paired with a dedicated slice of HBM bandwidth and
-\centering
+local on-PE storage (TCM), so compute lives next to the data it
-\includegraphics[width=\linewidth]{cube_architecture.pdf}
+consumes rather than fetching it through a far-away memory
-\caption{Modeled hardware graph at the \emph{CUBE} level---a
+controller. The platform's performance question is therefore not
-zoom-in on one CUBE node from Fig.~\ref{fig:sip-arch}. The CUBE
+``how many FLOPs can the chip do'' but ``how well can a kernel keep
-holds 8 PEs (each with its own HBM channels), the \textsf{M\_CPU}
+each compute engine fed from its locally-attached memory while
-control processor, a shared SRAM, an intra-CUBE NoC router mesh, and
+moving the unavoidable traffic between PEs efficiently.''
-UCIe links to neighbouring CUBEs.}
-\label{fig:cube-arch}
-\end{figure}
 
 Within a PE, commands are dispatched by \textsf{PE\_CPU} to
 \textsf{PE\_SCHED}, which routes work to specialized execution
 engines---DMA, FETCH/STORE, GEMM, vector-math, and IPCQ
-(Fig.~\ref{fig:pe-arch}). Commands come in two flavours. \emph{Atomic}
+(Fig.~\ref{fig:hw-arch}\subref{fig:pe-arch}). Commands come in two
-commands target a single engine---a plain DMA read, a GEMM tile, a
+flavours. \emph{Atomic} commands target a single engine---a plain
-vector-math op, an IPCQ send/receive---and are the natural unit for
+DMA read, a GEMM tile, a vector-math op, an IPCQ send/receive---and
-short or special-purpose work; the \textsf{PE\_CPU} itself also runs
+are the natural unit for short or special-purpose work; the
-control-plane work directly. \emph{Composite} commands, by contrast,
+\textsf{PE\_CPU} itself also runs control-plane work directly.
-carry an ordered pipeline of operations across multiple engines
+\emph{Composite} commands, by contrast, carry an ordered pipeline of
-(\textsf{DMA\_READ} $\rightarrow$ \textsf{FETCH} $\rightarrow$
+operations across multiple engines (\textsf{DMA\_READ} $\rightarrow$
-\textsf{GEMM}/\textsf{MATH} $\rightarrow$ \textsf{STORE} $\rightarrow$
+\textsf{FETCH} $\rightarrow$ \textsf{GEMM}/\textsf{MATH} $\rightarrow$
-\textsf{DMA\_WRITE}) that the \textsf{PE\_SCHED} tiles and streams
+\textsf{STORE} $\rightarrow$ \textsf{DMA\_WRITE}) that the
-without per-tile redispatch. The composite form is the substrate for
+\textsf{PE\_SCHED} tiles and streams without per-tile redispatch. The
-the GEMM optimization (\S\ref{sec:gemm}) and, combined with on-PE
+composite form is the substrate for the GEMM optimization
-collectives, for fused attention (\S\ref{sec:gqa}).
+(\S\ref{sec:gemm}) and, combined with on-PE collectives, for fused
-
+attention (\S\ref{sec:gqa}).
-\begin{figure*}[t]
-\centering
-\includegraphics[width=\linewidth]{pe_architecture.png}
-\caption{Modeled PE architecture used in this report---one example
-configuration whose specific parameters are listed in
-\S\ref{sec:hw} (Table~\ref{tab:hw}); KernBench is not tied to this
-particular decomposition and supports arbitrary PE block layouts
-provided each component is given a port and bandwidth model.
-\textsf{PE\_CPU} dispatches commands to the \textsf{PE\_SCHED}, which
-routes tile-token streams through the \textsf{PE\_DMA},
-\textsf{PE\_FETCH\_STORE}, and \textsf{GEMM}/\textsf{MATH} engines
-along on-chip links; the \textsf{PE\_IPCQ} provides the control plane
-for on-device collective communication.}
-\label{fig:pe-arch}
-\end{figure*}
 
 KernBench is layered along the flow of a request:
 
@@ -114,8 +146,20 @@ KernBench is layered along the flow of a request:
   controllers, and the inter-chiplet links.
 \end{itemize}
 
-Data and timing are handled in two passes, so that a kernel's numeric
+Data and timing are handled in two passes, so that a kernel's
-results and its latency are computed consistently but independently.
+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}
@@ -135,37 +179,14 @@ wire, \textbf{drain} (per-flit service occupancy) at the destination,
 and \textbf{queuing delay} at the shared FIFOs that the wire and the
 destination share between concurrent transactions.
 
-\begin{figure*}[t]
-\centering
-\includegraphics[width=\textwidth]{latency_model.png}
-\caption{Conceptual schematic of the latency model. Two source nodes
-(Requester A/B) inject flits through a chain of routers into a
-destination node; on the shared edge between routers, flits from the
-two transactions are interleaved flit-by-flit by the wire's FIFO
-arrival order. End-to-end latency is the sum of four contributions:
-\textbf{per-node overhead} (the switch's fixed processing cost,
-shown in light yellow---the same colour as the destination's
-processing-logic block), \textbf{per-edge transmission}
-($\textit{flit\_size}/\textit{BW}$ on each wire),
-\textbf{drain} (per-flit service occupancy at the destination's
-channel), and \textbf{queuing delay} (waiting in a FIFO when a
-shared resource is busy). The places where queuing actually
-accumulates are highlighted in green: the router output queue and
-the destination's input queue. This is the model, not a
-measurement; specific bandwidths and overheads are listed in
-\S\ref{sec:hw} (Table~\ref{tab:hw}).}
-\label{fig:latency-model}
-\end{figure*}
-
 \paragraph{The hardware as a graph.} The topology is compiled once at
 configuration time into this graph and is never mutated during a run.
 There are no hidden shortcuts, implicit bypasses, or magic paths: if a
 request reaches its destination, the path it took is explicit in the
 graph, and the latency it incurred is the sum of the per-node and
 per-edge costs paid along that path. The same graph representation
-applies recursively at every hierarchy level---system, SIP
+applies recursively at every hierarchy level---system, SIP, CUBE, and
-(Fig.~\ref{fig:sip-arch}), CUBE (Fig.~\ref{fig:cube-arch}), and PE
+PE (Fig.~\ref{fig:hw-arch}).
-(Fig.~\ref{fig:pe-arch}).
 
 \paragraph{From graph to discrete-event simulation.} The graph is
 driven by a discrete-event engine. Two kinds of events advance
@@ -178,7 +199,7 @@ fires them one at a time, with ties broken under a deterministic
 policy so that the same kernel on the same topology always yields the
 same trace. Per-request correlation IDs are stamped at injection and
 carried through every hop, so the path from injection to completion
-is fully traceable. Every nanosecond in a reported latency
+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.
 
@@ -195,7 +216,73 @@ collective engines hold the request for their service time before
 releasing it downstream. Each of these is attached to a specific node
 or edge in the graph; together they make up the entire latency budget.
 
-\paragraph{Congestion: where bottlenecks emerge.} The simulator's
+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
@@ -215,70 +302,127 @@ it reveals where the real bottlenecks form and which hardware levers
 actually relieve them---which is exactly the question the codesign
 work in this report turns on.
 
-\paragraph{Control-plane (issue) cost model.} The cost of \emph{issuing}
-a command is modelled structurally rather than with a per-operation
-calibration table. The PE control processor charges, per command,
-\[
-  d_{\text{cmd}} = \textsf{FIXED} + b_{\text{logical}} \cdot R,
-\]
-where $b_{\text{logical}}$ is the command's hardware-logical byte size,
-\textsf{FIXED} captures the fixed per-command cost (queue-tail update,
-completion registration) and $R$ captures the per-byte cost of
-serializing the command descriptor into the scheduler queue. The
-default anchoring (\textsf{FIXED} $= 40$ cycles, $R = 0.0625$
-cycles/byte, i.e.\ \SI{16}{\byte\per\cycle}, at \SI{1}{\giga\hertz})
-places a typical composite at roughly \SI{43}{\nano\second}, and a
-hard cap on a composite's descriptor size prevents the model from
-rewarding arbitrarily large fused commands beyond what real descriptor
-queues accept. In the configurations measured here, command issue is
-not the bottleneck---data movement is---so this term stays small
-relative to DMA and collective time.
 
-\paragraph{Accuracy.} The model is precise about the effects that
+\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
+translates into physically reasonable kernel latencies.
-GEMM study (\S\ref{sec:gemm}), simulator-measured MAC efficiency
+
-tracks an analytic ideal-pipeline model within roughly
+First, in the GEMM study (\S\ref{sec:gemm}), simulator-measured MAC
-\SIrange{10}{20}{\percent} across a wide range of tile counts; the
+efficiency tracks an analytic ideal-pipeline model within
-residual gap is attributable to pipeline-fill and DMA effects the
+\SI{1.4}{ppt} across every swept shape---from a single-tile
-analytic model omits. Second, in the all-reduce study
+$M{=}K{=}N{=}32$ at $\sim\SI{7.7}{\percent}$ up to the deep-$K$
-(\S\ref{sec:allreduce}, Fig.~\ref{fig:allreduce-cmp}), simulator
+$K{=}3072$ case at $\sim\SI{90}{\percent}$. The residual gap is
-latency for a 2D-torus over six devices follows the expected
+fill/tail overhead the closed-form pipeline model omits.
-startup-plus-per-packet shape across the entire payload sweep---tight
+
-at small payloads where startup dominates, and within a single-digit
+Second, in the all-reduce study (\S\ref{sec:allreduce},
-multiplicative factor at the largest payloads, where the residual gap
+Fig.~\ref{fig:allreduce-cmp}), simulator latency for a 2D-torus over
-is explained by per-router switching the analytic shape elides. A
+six devices follows the expected startup-plus-per-packet shape across
-single-device point from an external full-system simulator (FSIM) at
+the entire payload sweep---tight at small payloads where startup
-the largest payload sits an order of magnitude above the KernBench
+dominates, and within a single-digit multiplicative factor at the
-multi-device torus, illustrating the well-known gap between an
+largest payloads, where the residual gap is explained by per-router
-achievable-kernel number and a full end-to-end-stack number rather
+switching the analytic shape elides.
-than a model error. The known simplifications---FIFO router
+
-arbitration (instead of round-robin), HBM scheduler without
+Third, the simulator's per-traversal behaviour can be probed
-write-buffer reordering, no bank conflict, no refresh or thermal
+directly. Running \texttt{kernbench probe} on the modelled topology
-effects, and no upstream backpressure---are the price of a
+issues a sequence of PE-to-HBM DMA reads at progressively greater
-deterministic, inspectable model. These simplifications bound the
+hop distances and reports the per-component overhead, per-edge
-absolute accuracy, but the agreement with both analytic models and the
+serialization, and per-PC drain that the model charges
-external full-system simulator data above indicates that KernBench is
+(Table~\ref{tab:probe-pe-dma}). Three properties stand out. First,
-sufficiently accurate for evaluating the \emph{relative}
+the reported latency increases \emph{monotonically} with hop
-hardware--software design trade-offs (tiling A vs.\ B, topology X
+count---from \SI{141}{\nano\second} at the local HBM slice to
-vs.\ Y, with vs.\ without composite command, mesh vs.\ torus) that
+\SI{677}{\nano\second} at the worst-case remote-CUBE slice---with the
-are the primary objective of this work.
+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 for
+Table~\ref{tab:hw} summarizes the hardware configuration
-every experiment in this report. It is read directly from the
+used throughout this report. The intent of this
-simulator's topology description; per-experiment workload parameters
+configuration is not to model a specific product, but to
-(matrix shapes, collective sizes, sequence lengths) are stated in
+represent a realistic memory-centric accelerator and to
-their respective sections rather than here.
+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 (shared by all experiments).}
+\caption{Modeled hardware configuration}
 \label{tab:hw}
 \small
 \begin{tabular}{@{}ll@{}}
@@ -296,11 +440,12 @@ GEMM engine peak & \SI{8}{\tera\flop\per\second} (f16) \\
 TCM (on-PE) & \SI{16}{\mega\byte}, \SI{512}{\giga\byte\per\second} R/W \\
 \quad kernel scratch & \SI{1}{\mega\byte} \\
 DMA engines & 1 read + 1 write \\
-CPU / scheduler overhead & \SI{2}{\nano\second} / \SI{1}{\nano\second} \\
+\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{1024}{\giga\byte\per\second} \\
+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} \\
@@ -311,7 +456,8 @@ Inter-CUBE (UCIe PHY) & \SI{512}{\giga\byte\per\second}, \SI{8}{\nano\second}, X
 Inter-SIP (PCIe) & \SI{768}{\giga\byte\per\second} per endpoint \\
 \midrule
 \multicolumn{2}{@{}l}{\emph{Command-issue cost model (defaults)}} \\
-FIXED per command & 40 cycles \\
+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

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

@@ -1,8 +1,6 @@
 \section{GEMM Acceleration via the Composite Command}
 \label{sec:gemm}
 
-\subsection{Why it is needed}
-
 GEMM is the compute core of every transformer block---the QKV
 projections, the attention score and context products, and the
 feed-forward matrices are all matrix multiplications. On a tiled
@@ -21,20 +19,17 @@ in command overhead?
 
 \subsection{Design}
 
-The answer is the \emph{composite command}. A single command carries the
+The answer is the \emph{composite command}. A single command carries
-ordered tile pipeline
+the ordered five-stage tile pipeline (\textsf{DMA\_READ},
-\[
+\textsf{FETCH}, \textsf{GEMM}, \textsf{STORE}, \textsf{DMA\_WRITE}),
-\textsf{DMA\_READ}\rightarrow\textsf{FETCH}\rightarrow\textsf{GEMM}
-\rightarrow\textsf{STORE}\rightarrow\textsf{DMA\_WRITE},
-\]
 and the PE scheduler splits the payload into hardware tiles
-(here $32\times64\times32$), emitting one tile token per tile. Subsequent
+(here $32\times64\times32$), emitting one tile token per tile.
-stages are reached by \emph{token self-routing} between the on-PE engines,
+Subsequent stages are reached by \emph{token self-routing} between
-so a tile flows DMA\,$\rightarrow$\,fetch\,$\rightarrow$\,GEMM\,$%
+the on-PE engines, so a tile flows through the chain without
-\rightarrow$\,store without returning to the scheduler between stages.
+returning to the scheduler between stages. Because the whole
-Because the whole pipeline is described by one command, the issue cost is
+pipeline is described by one command, the issue cost is paid once
-paid once per GEMM rather than once per tile-stage, and the scheduler is
+per GEMM rather than once per tile-stage, and the scheduler is free
-free to keep every stage busy on different tiles simultaneously---tile
+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
@@ -45,66 +40,362 @@ kernel to fuse its softmax work into the GEMM pipeline
 \subsection{Results}
 
 We sweep eight GEMM shapes spanning square, tall, wide, and deep-$K$
-geometries, under three operand-staging variants
+geometries under two operand-staging variants that bracket the
-(\textsf{ref\_ref}, both operands streamed from HBM; \textsf{load\_ref},
+realistic LLM cases. In \textsf{load\_ref}, the activation $A$ is
-one operand resident in TCM; \textsf{load\_load}, both resident).
+pre-staged on chip and only the weight $W$ streams from HBM during
-Figure~\ref{fig:gemm-util} reports MAC utilization and efficiency, and
+the composite (the ``activation in TCM, weights from HBM'' case
-Figure~\ref{fig:gemm-stages} breaks the kernel into per-stage engine
+typical of decoding with a small batch). In \textsf{ref\_ref}, both
-busy time.
+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.
+\caption{GEMM MAC utilization and efficiency, theoretical vs.\ measured
-Tile-fill sets the ceiling: under-tile shapes (marked $\ast$) such as
+(\textsf{load\_ref} staging). Tile-fill sets the ceiling: under-tile
-$M{=}K{=}N{=}32$, $M{=}8$, and $K{=}8$ cannot fill the MAC tile and cap at
+shapes (marked $\ast$) such as $M{=}K{=}N{=}32$, $M{=}8$, and $K{=}8$
-\SI{50}{\percent}, \SI{25}{\percent}, \SI{12.5}{\percent}. For
+cannot fill the MAC tile and cap at \SI{50}{\percent},
-tile-filling shapes, efficiency climbs with tile count---from
+\SI{25}{\percent}, \SI{12.5}{\percent}. For tile-filling shapes,
-\textasciitilde\SI{23}{\percent} at one tile to \textasciitilde%
+efficiency climbs with tile count---from
-\SI{78}{\percent} measured at 48 tiles ($K{=}3072$)---and the measured
+\textasciitilde\SI{15}{\percent} at one tile to \textasciitilde%
-bars track the analytic prediction within
+\SI{90}{\percent} measured at 48 tiles ($K{=}3072$)---and the
-\SIrange{10}{20}{\percent}.}
+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
+\caption{Per-stage engine wall-clock (DMA in, Fetch, GEMM, DMA out)
-\textsf{load\_ref} staging. For the deep-$K$ shape ($K{=}3072$, 48 tiles)
+under \textsf{load\_ref} staging. For the deep-$K$ shape ($K{=}3072$,
-the Fetch and GEMM stages are large and comparable
+48 tiles) DMA in, Fetch, and GEMM are all close to
-(\textasciitilde\SI{770}{} and \SI{785}{\nano\second}) while DMA-out is
+\textasciitilde\SI{785}{\nano\second}, running concurrently in a
-negligible---a compute-rich, well-pipelined regime. For the low-reuse
+tightly pipelined regime while DMA-out is negligible. For the
-shape ($M{=}128,K{=}8,N{=}128$) DMA-out grows to
+low-reuse shape ($M{=}128,K{=}8,N{=}128$, 16 output tiles) DMA-out
-\textasciitilde\SI{350}{\nano\second} and compute is small---a
+grows to \textasciitilde\SI{336}{\nano\second} while GEMM compute is
-data-movement-bound regime.}
+\textasciitilde\SI{262}{\nano\second}---a data-movement-bound regime
+where the output write becomes the largest stage.}
 \label{fig:gemm-stages}
 \end{figure}
 
-Two regularities stand out. (i) Utilization is governed by how completely
-the problem fills the MAC tile: the three under-tile shapes are hard-capped
-well below \SI{100}{\percent}, independent of how the kernel is issued.
-(ii) Among tile-filling shapes, efficiency is governed by tile count---more
-tiles amortize the one-time pipeline fill and issue cost, so the deep-$K$
-shape reaches \textasciitilde\SI{78}{\percent} of peak while a single-tile
-shape reaches only \textasciitilde\SI{23}{\percent}. The stage breakdown
-explains why: with 48 tiles the GEMM and Fetch stages overlap and stay
-busy, whereas the low-reuse shape spends most of its time moving data in
-and out.
-
 \subsection{Analysis and meaning}
 
-The composite command does not manufacture bandwidth or MAC throughput; it
+The composite command does not manufacture bandwidth or MAC
-removes the two software-shaped obstacles between a GEMM and its hardware
+throughput; it removes the two software-shaped obstacles between a
-roofline. By paying issue cost once per GEMM and chaining tile-stages
+GEMM and its hardware roofline. By paying the issue cost once per
-through token self-routing, it lets the scheduler keep the pipeline full,
+GEMM rather than once per tile-stage, and by chaining tile-stages
-so compute-rich shapes actually reach the efficiency their arithmetic
+through token self-routing, it lets the scheduler keep every stage
-intensity allows ($\sim$\SI{78}{\percent} measured at 48 tiles), and
+busy on a different tile, so compute-rich shapes actually reach the
-data-bound shapes actually reach their DMA bound instead of stalling on
+efficiency their arithmetic intensity allows
-command overhead. The close agreement between measured and theoretical
+($\sim$\SI{90}{\percent} of GEMM peak at the deep-$K$
-efficiency (Figure~\ref{fig:gemm-util}) is also the report's primary
+\textsf{load\_ref} corner) and data-bound shapes actually reach
-validation that KernBench's latency model is faithful in the regime that
+their HBM-BW ceiling instead of stalling on command overhead.
-matters. The hardware implication is concrete: a single-command,
+
-self-routing tile pipeline is the issue mechanism that makes the MAC array
+The split between \textsf{load\_ref} and \textsf{ref\_ref} also makes
-usable, and it is a prerequisite---not a luxury---for the fused attention
+the hardware-software boundary explicit: when the activation fits in
-kernel of \S\ref{sec:gqa}.
+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

@@ -1,46 +1,215 @@
 \section{PE\_IPCQ and Collective Communication}
 \label{sec:allreduce}
 
-\subsection{Why it is needed}
-
 Distributing a transformer across devices turns every tensor-parallel
 layer into a collective: partial results computed on different PEs, CUBEs,
-and SIPs must be summed and redistributed with an all-reduce. If that
+and SIPs must be summed and redistributed with an all-reduce. Underneath
-collective is handled by the host or by a generic DMA path, three problems
+the algorithm this is fundamentally a PE-to-PE problem---many short
-appear. The reduction traffic competes with the kernel's own compute DMA
+messages flowing between neighbors as the reduction proceeds. The natural
-on the same links, causing head-of-line blocking; there is no efficient
+software realization is a per-direction ring buffer whose head and tail
-peer-to-peer ring primitive, so data takes extra hops; and the ordering is
+pointers the producer and consumer update atomically and poll, but our
-hard to make deterministic. The hardware question is how to perform
+H2 2025 report measured this scheme end-to-end and found that the
-collectives \emph{on the device}, overlapped with compute and reproducible
+atomic-pointer traffic together with the consumer's polling loop dominate
-run-to-run.
+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}
 
-KernBench models a dedicated per-PE collective engine, \textbf{PE\_IPCQ}
+The proposed block is \textbf{PE\_IPCQ} (inter-PE communication queue),
-(inter-PE communication queue). It is a control-plane block: it owns the
+a small controller dropped into every PE next to PE\_DMA, PE\_GEMM,
-ring-buffer address arithmetic, head/tail pointers, peer-pointer caches,
+PE\_MATH, and PE\_TCM. It is a control-plane block---it holds no
-backpressure, and the four-direction (N/S/E/W) neighbor map, with eight
+payload data---and consists of three pieces: a per-direction
-ring buffers per PE (four directions $\times$ \{tx, rx\}). Crucially,
+\emph{QPair register file} (\textasciitilde\SI{576}{\byte} of flip-flops
-PE\_IPCQ does \emph{not} move data itself---it delegates the actual
+covering up to eight directions), a combinational slot-address
-transfer to PE\_DMA, keeping a clean control/data split. To stop
+generator and backpressure comparator, and a credit injector/receiver
-collective traffic from blocking compute, PE\_DMA is extended into a
+wired to the NoC. Each QPair holds the local pointers
-two-channel virtual-channel model: \texttt{vc\_compute} carries tile
+(\texttt{my\_head}, \texttt{my\_tail}), shadowed views of the peer's
-load/store for GEMM and vector math, \texttt{vc\_comm} carries IPCQ sends,
+(\texttt{peer\_head\_cache}, \texttt{peer\_tail\_cache}), the local and
-each with an independent state machine. The same physical link is shared
+peer rx-buffer physical bases, ring depth and slot size (both
-but progresses in chunks (\SI{256}{\byte}), so a large GEMM DMA does not
+power-of-two), and the peer's credit-target address. The ring data
-lock the link end-to-end against a pending reduction. On top of this
+itself lives in a reserved slot region of TCM (or PE-local HBM, or
-substrate the collective runs a hierarchical local-reduce / global
+cube-shared SRAM), addressed by the QPair registers; PE\_IPCQ never
-all-reduce-broadcast schedule across whatever inter-device topology the
+touches the bytes, only the pointers. The PE\_CPU sees the controller
-configuration specifies.
+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}
 
-Following the milestone-evaluation convention, the collective sweep builds
+All measurements in this section run on the PE\_IPCQ substrate
-its own six-device (six-SIP, $2\times3$) configurations---distinct from
+described above; the topology sweep is intended to characterize how
-the two-SIP default of Table~\ref{tab:hw}---and measures all-reduce
+effectively the proposed mechanism exposes the underlying interconnect's
-latency as a function of payload size for three inter-device topologies:
+properties, not to compare PE\_IPCQ against an alternative
-a 1D ring, a 2D mesh (no wrap), and a 2D torus. Table~\ref{tab:allreduce}
+communication primitive. Following the milestone-evaluation convention,
-and Figure~\ref{fig:allreduce-cmp} report the result.
+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
@@ -52,11 +221,11 @@ per-PE payload. Lower is better; the torus wins at every size.}
 \toprule
 \textbf{Bytes/PE} & \textbf{2D mesh} & \textbf{Ring 1D} & \textbf{2D torus} \\
 \midrule
-256     & 2667  & 2365  & 1701 \\
+256     & 4189  & 3883  & 2957 \\
-4{,}096  & 4450  & 4082  & 3038 \\
+4{,}096  & 5566  & 5240  & 4031 \\
-16{,}384 & 8900  & 8217  & 6403 \\
+16{,}384 & 10016 & 9376  & 7396 \\
-65{,}536 & 26705 & 24766 & 19865 \\
+65{,}536 & 27821 & 25925 & 20858 \\
-98{,}304 & 38574 & 35798 & 28840 \\
+98{,}304 & 39690 & 36957 & 29833 \\
 \bottomrule
 \end{tabular}
 \end{table}
@@ -64,13 +233,19 @@ per-PE payload. Lower is better; the torus wins at every size.}
 \begin{figure}[t]
 \centering
 \includegraphics[width=\linewidth]{allreduce_comparison.png}
-\caption{All-reduce latency vs.\ per-PE payload for the three topologies,
+\caption{All-reduce latency vs.\ per-PE payload for the three PE\_IPCQ
-against the analytic torus model and an external full-system simulator
+topologies, against the analytic torus model. The isolated point in
-(FSIM) reference. The measured torus tracks the analytic curve within a
+the top panel (\SI{366}{\micro\second}) is the only data point
-small constant factor; the FSIM single-device point
+available from the H2 2025 software-queue measurement campaign---an
-(\SI{366}{\micro\second}) sits an order of magnitude above the
+intra-device all-reduce across 16 CUBEs on a single device. A true
-KernBench algorithmic latency, illustrating the difference between an
+6-device inter-SIP measurement under the same software queue was not
-achievable-kernel number and a full end-to-end-stack number.}
+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}
 
@@ -78,9 +253,9 @@ achievable-kernel number and a full end-to-end-stack number.}
 \centering
 \includegraphics[width=\linewidth]{allreduce_buffer_kind.png}
 \caption{Effect of IPCQ staging-buffer placement (2D torus). At
-\SI{64}{\kibi\byte}/PE, TCM staging (\SI{19865}{\nano\second}) beats HBM
+\SI{64}{\kibi\byte}/PE, TCM staging (\SI{20858}{\nano\second}) beats HBM
-(\SI{23081}{\nano\second}) by \textasciitilde\SI{14}{\percent} and SRAM
+(\SI{24074}{\nano\second}) by \textasciitilde\SI{13}{\percent} and SRAM
-(\SI{32201}{\nano\second}) by \textasciitilde\SI{38}{\percent}; at small
+(\SI{33194}{\nano\second}) by \textasciitilde\SI{37}{\percent}; at small
 payloads the three are indistinguishable.}
 \label{fig:allreduce-buf}
 \end{figure}

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

@@ -1,8 +1,6 @@
 \section{Fused Grouped-Query Attention}
 \label{sec:gqa}
 
-\subsection{Why it is needed}
-
 Attention is the 1H focus, and it is where the two preceding optimizations
 have to come together. Grouped-Query Attention (GQA) shrinks the KV cache
 by sharing each KV head across a group of query heads (here $h_q=8$ query
@@ -16,12 +14,99 @@ realizing it as a fast \emph{fused} kernel needs both building blocks from
 this report: efficient GEMM issue (\S\ref{sec:gemm}) for the
 $Q\!\cdot\!K^{\top}$ and $P\!\cdot\!V$ products, and an efficient on-device
 reduction (\S\ref{sec:allreduce}) for the multi-user and
-sequence-parallel KV reductions. This section is the capstone: the fused
+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{Design}
+\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
@@ -37,74 +122,273 @@ overlaps score computation; and per-tile \emph{scratch recycling} that
 keeps the running softmax accumulators ($m,\ell,O$) in a persistent arena
 while freeing per-tile temporaries, so the kernel fits the
 \SI{1}{\mebi\byte} scratch budget across many tiles. A further refinement
-that restructures the decode step into two stateful composites (a named
+restructures the decode step so the per-tile matrix products and the
-\textsf{softmax\_merge} recipe) is designed but not yet wired into the
+online-softmax merge are issued as \emph{composite} commands rather than
-measured path; results below reflect the implemented kernel only.
+hand-tiled primitives; \S\ref{sec:gqa-composite} measures it against the
+primitive baseline at long context.
 
-\subsection{Results}
+% 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/
 
-We measure four headline panels that vary the user count $C$ and the phase:
+% TODO: prefill performance figure (latency, stage breakdown).
-single- and multi-user prefill ($T_q=4$, $S_{kv}=16$), and single- and
+% TODO: decode performance figure (latency, stage breakdown).
-multi-user decode ($P=8$ PEs, $S_{kv}=64$ and $128$), all at $d_{\text{head}}=64$
+%       Bench output for short_ctx to be generated.
-and $G=8$. For each panel we harvest end-to-end latency
-(max event end minus min event start, the same window convention as the
-GEMM study) together with the per-engine busy time and the operation mix.
-Figure~\ref{fig:gqa-lat} and Figure~\ref{fig:gqa-break} report the result;
-the underlying numbers are in Table~\ref{tab:gqa}.
 
-\begin{table}[t]
+\subsection{Inference with Long-Context Length}
-\centering
+\label{sec:gqa-long}
-\caption{Fused GQA per-panel latency and operation mix. Compute (GEMM,
+
-MATH) is a tiny fraction of DMA occupancy; IPCQ copies grow with users and
+% TODO: prefill long-context kernel implementation description
-PEs.}
+% (Sequence-Parallel partition of S_kv, per-case mechanics).
-\label{tab:gqa}
+% Bench code: src/kernbench/benches/gqa_helpers/long_ctx/
-\small
+
-\begin{tabular}{@{}lrrrr@{}}
+% TODO: prefill long-context performance figure.
-\toprule
+
-\textbf{Panel} & \textbf{Lat.\ (ns)} & \textbf{GEMM} & \textbf{IPCQ} & \textbf{DMA rd} \\
+Long-context decode is the regime where the KV cache, not attention
-\midrule
+compute, sets serving cost, so the placement question of
-prefill C=1        & 445   & 2  & 0  & 3  \\
+\S\ref{sec:gqa-placement} becomes decisive here. To pick the right
-prefill C=4 (Ring) & 4630  & 32 & 24 & 12 \\
+placement for this regime we run each of the six options as a fused
-decode C=1, P=8    & 3632  & 16 & 21 & 24 \\
+decode kernel on the simulator (one decode step on the LLaMA-3.1-70B
-decode C=4, P=8    & 6693  & 64 & 93 & 96 \\
+single-KV-head-group target, $C{=}8$ CUBEs $\times$ $P{=}8$ PEs) at a
-\bottomrule
+tractable $S_{kv}{=}8192$, and read off end-to-end latency, on-device op
-\end{tabular}
+traffic, and the redundant compute each one induces. The swept context is
-\end{table}
+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_latency_by_panel.png}
+\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_latency.png}
-\caption{Fused GQA end-to-end latency. Latency grows from
+\caption{Measured end-to-end decode latency per placement
-\SI{445}{\nano\second} (single-user prefill) to \SI{6693}{\nano\second}
+($S_{kv}{=}8192$; red = ruled out by the long-context memory budget,
-(four-user decode) as the KV history and the number of participating
+blue = the predicted Pareto choice). The fastest raw latency belongs to
-devices grow.}
+Case~3 (\SI{17.8}{\micro\second})---but Case~3 replicates the full KV
-\label{fig:gqa-lat}
+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_op_engine_breakdown.png}
+\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_parallelism.png}
-\caption{Where the work goes. Left: operation counts---GEMM and IPCQ-copy
+\caption{Redundant compute per placement, measured as active-PE
-volume both scale with users and PEs. Right: summed engine occupancy on a
+$\times$ $S_{\text{local}}$ (PE-tokens; lower means less wasted work).
-log scale---the DMA engine dominates by two to three orders of magnitude
+The minimum is \num{8192} PE-tokens---one pass over the sequence.
-over the GEMM and MATH engines in every panel.}
+Case~3 inflates this 8$\times$ to \num{65536} by replicating the KV
-\label{fig:gqa-break}
+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}
 
-The dominant observation is in Figure~\ref{fig:gqa-break}: the compute
+\begin{figure}[t]
-engines are almost idle. The GEMM engine accumulates only
+\centering
-\SIrange{2}{33}{\nano\second} of busy time across the panels and the
+\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_traffic.png}
-vector-math engine \SIrange{5}{688}{\nano\second}, while the DMA engine
+\caption{Measured on-device op traffic per placement. The unsharded
-accumulates \SIrange{72}{15920}{\nano\second}. Fused GQA, as modeled here,
+Case~1 issues no IPCQ copies (each PE has the full cache, nothing to
-is overwhelmingly data-movement bound. The operation mix shows why the
+reduce); the single-axis Case~3 charges 168. Among the 64-way splits,
-collective machinery matters: IPCQ-copy count rises from zero (single-user
+the $d_{\text{head}}$-split Cases~4--5 charge the most---280 IPCQ copies
-prefill) to 93 (four-user decode) as the kernel reduces partial outputs
+and 8 DMA writes each, the partial-score all-reduce that head-slicing
-across more PEs and CUBEs, and DMA-read count rises in step as more KV
+forces---while Case~6~$\star$ needs only 189 IPCQ copies and a single DMA
-shards are streamed. The PE control-processor dispatch cost registered as
+write, because it merges just the running softmax state $(m,\ell,O)$
-zero in this configuration---command issue is simply not on the critical
+rather than partial scores. This is the on-device collective traffic that
-path when data movement is this dominant.
+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}
 
-\subsection{Analysis and meaning}
+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
@@ -121,3 +405,9 @@ 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

@@ -40,7 +40,10 @@ hardware levers for this workload:
 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.
+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

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_gqa_latency.py

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

scripts/paper/paper_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_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

@@ -1,12 +1,25 @@
 """Comparative figures for milestone-gqa-decode-long-ctx-4cases.
 
-Reads sweep.json (emitted by ``kernbench run --bench
+Reads sweep_decode.json (emitted by the milestone-1h-gqa bench) and
-milestone-gqa-decode-long-ctx-4cases``) and writes three PNGs into
+writes four PNGs into the same bench-output dir
-``docs/report/1H-codesign-paper/figures/``:
+(src/kernbench/benches/1H_milestone_output/gqa/long_ctx/):
 
-  gqa_decode_long_ctx_4cases_latency.png   end-to-end latency per case
+  gqa_decode_long_ctx_6cases_latency.png       end-to-end latency per case
-  gqa_decode_long_ctx_4cases_traffic.png   ipcq/dma op-count breakdown
+  gqa_decode_long_ctx_6cases_traffic.png       ipcq/dma op-count breakdown
-  gqa_decode_long_ctx_4cases_memory.png    KV bytes per cube per case
+  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 \\
@@ -24,22 +37,34 @@ matplotlib.use("Agg")
 import matplotlib.pyplot as plt  # noqa: E402
 
 _REPO_ROOT = Path(__file__).resolve().parents[2]
-_FIG_DIR = _REPO_ROOT / "docs" / "report" / "1H-codesign-paper" / "figures"
+# Sweep JSON + PNGs live together under the bench output dir.
-_SWEEP_JSON = (
+_FIG_DIR = (
     _REPO_ROOT / "src" / "kernbench" / "benches"
-    / "1H_milestone_output" / "gqa_decode_long_ctx_4cases" / "sweep.json"
+    / "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"
 
-# Panel name → (short label, case ordinal for left-to-right plot order).
 _CASE_INFO = {
-    "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp": (
+    # panel name                                                       label                            ord  flag
-        "Case 1\nCube-SP × PE-TP", 1),
+    "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_repl_pe_tp": (
+    "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp":              ("Case 2\nCube-SP × PE-repl",     2, _OVERFLOW),
-        "Case 2\nCube-Repl × PE-TP", 2),
+    "single_kv_group_decode_long_ctx_gqa_cube_repl_pe_sp":            ("Case 3\nCube-Repl × PE-SP",     3, _OVERFLOW),
-    "single_kv_group_decode_long_ctx_gqa_cube_repl_pe_sp": (
+    "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp_dhead":        ("Case 4\nCube-SP × PE-TP-d_head", 4, _NORMAL),
-        "Case 3\nCube-Repl × PE-SP", 3),
+    "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": (
+    "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_sp":              ("Case 6 ★\nCube-SP × PE-SP",     6, _PARETO),
-        "Case 4 ★\nCube-SP × PE-SP", 4),
+}
+
+# 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
 }
 
 
@@ -51,23 +76,26 @@ 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]
-    colors = ["#888", "#888", "#888", "#3b6ea5"]  # Case 4 highlighted
+    fig, ax = plt.subplots(figsize=(12.0, 4.8))
-    fig, ax = plt.subplots(figsize=(8.0, 4.5))
+    bars = ax.bar(labels, lat_us, color=_bar_colors(rows), width=0.6)
-    bars = ax.bar(labels, lat_us, color=colors, width=0.6)
     ax.set_ylabel("end-to-end latency (µs)")
     ax.set_title(
-        "Long-context decode 4-cases — end-to-end latency per case\n"
+        "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_4cases_latency.png"
+    out = _FIG_DIR / "gqa_decode_long_ctx_6cases_latency.png"
     fig.savefig(out, dpi=150)
     plt.close(fig)
     return out
@@ -81,58 +109,123 @@ def _plot_traffic(rows: list[dict]) -> Path:
     disp = ["IPCQ copy", "DMA read", "DMA write"]
     colors = ["#c0504d", "#9bbb59", "#8064a2"]
     w = 0.25
-    fig, ax = plt.subplots(figsize=(9.0, 4.5))
+    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 4-cases — op-count breakdown per case")
+    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_4cases_traffic.png"
+    out = _FIG_DIR / "gqa_decode_long_ctx_6cases_traffic.png"
     fig.savefig(out, dpi=150)
     plt.close(fig)
     return out
 
 
-def _kv_bytes_per_cube(panel: str, *, S_kv: int, h_kv: int,
+def _s_local_per_pe(panel: str, *, S_kv: int, C: int, P: int) -> int:
-                       d_head: int, C: int) -> int:
+    """S_local (token count) each PE attends over locally.
-    """KV bytes a single cube's HBM holds (K + V together, f16)."""
+
-    # Cube-replicate ⇒ each cube holds full S_kv.
+      cube_repl_pe_tp  (Case 1): S_kv        (no sharding, KV-wise)
-    # Cube-SP        ⇒ each cube holds S_kv / C.
+      cube_sp_pe_tp    (Case 2): S_kv / C    (cube splits S_kv, PEs replicate)
-    # PE replicate vs row_wise share the cube's HBM and don't change
+      cube_repl_pe_sp  (Case 3): S_kv / P
-    # per-cube bytes.
+      cube_sp_pe_tp_dhead (Case 4): S_kv / C (cube splits S_kv, PE splits d_head)
-    S_per_cube = S_kv if "cube_repl" in panel else S_kv // C
+      cube_tp_dhead_pe_sp (Case 5): S_kv / P (cube splits d_head, PE splits S_kv)
-    return 2 * S_per_cube * h_kv * d_head * 2  # K + V, f16 (2 B/elem)
+      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_cube = [
+    mib_per_pe = [
-        _kv_bytes_per_cube(
+        _kv_bytes_per_pe(
             r["panel"], S_kv=r["S_kv"], h_kv=r["h_kv"],
-            d_head=r["d_head"], C=r["C"],
+            d_head=r["d_head"], C=r["C"], P=r["P"],
         ) / (1024 * 1024)
         for r in rows
     ]
-    # SP = blue (efficient); Repl = red (wasteful).
+    fig, ax = plt.subplots(figsize=(12.0, 4.8))
-    colors = ["#3b6ea5", "#c0504d", "#c0504d", "#3b6ea5"]
+    bars = ax.bar(labels, mib_per_pe, color=_bar_colors(rows), width=0.6)
-    fig, ax = plt.subplots(figsize=(8.0, 4.5))
+    ax.set_ylabel("KV bytes per PE (MiB, K + V, f16)")
-    bars = ax.bar(labels, mib_per_cube, color=colors, width=0.6)
-    ax.set_ylabel("KV bytes per cube (MiB, K + V, f16)")
     ax.set_title(
-        "Long-context decode 4-cases — KV memory per cube\n"
+        "Long-context decode 6-cases — KV memory per PE\n"
         "(one KV-head group; per-layer, per-token state)"
     )
-    ax.bar_label(bars, fmt="%.1f", padding=3, fontsize=9)
+    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_cube) * 1.15)
+    ax.set_ylim(0, max(mib_per_pe) * 1.15)
     fig.tight_layout()
-    out = _FIG_DIR / "gqa_decode_long_ctx_4cases_memory.png"
+    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
@@ -144,9 +237,11 @@ def main() -> None:
     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__":

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

scripts/verify_case4_dhead_tp.py

@@ -0,0 +1,148 @@
+"""Phase 1a smoke test for the Case 4 d_head-TP decode kernel.
+
+Runs `gqa_attention_decode_long_ctx_cube_sp_pe_tp_dhead_kernel` at small
+S_kv (2K) to verify:
+  1. it imports cleanly,
+  2. it runs without error under the same harness as Case 6,
+  3. captures op_log_summary (gemm/ipcq/dma counts),
+  4. compares vs analytical predictions and vs Case 6 (same memory tier).
+
+Usage:
+  python scripts/verify_case4_dhead_tp.py
+"""
+from __future__ import annotations
+
+import os
+import sys
+from pathlib import Path
+
+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_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, _summarize_op_log,
+)
+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
+
+# Small smoke-test params (S_kv=2K is enough to exercise tile-loop + AR).
+_PARAMS = dict(C=8, P=8, T_q=1, S_kv=2_048,
+               d_head=128, h_q=8, h_kv=1)
+
+
+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 main() -> int:
+    topology = os.environ.get("GQA_1H_TOPOLOGY", "topology.yaml")
+    topo = resolve_topology(topology)
+
+    print(f"Smoke params: {_PARAMS}")
+    print()
+
+    for label, bench_fn in (
+        ("Case 4 (Cube-SP × PE-TP d_head)",   _bench_fn_case4),
+        ("Case 5 (Cube-TP d_head × PE-SP)",   _bench_fn_case5),
+        ("Case 6 (Cube-SP × PE-SP S_kv)",     _bench_fn_case6),
+    ):
+        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}")
+            continue
+        if not res.completion.ok:
+            print(f"  {label:<42} ENGINE FAIL: {res.completion}")
+            continue
+        s = _summarize_op_log(res.engine.op_log)
+        lat = (res.engine.op_log[-1].t_end if res.engine.op_log else 0.0)
+        print(f"  {label:<42} "
+              f"gemm={s['gemm_count']:>4}  "
+              f"ipcq={s['ipcq_copy_count']:>4}  "
+              f"dma_r={s['dma_read_count']:>4}  "
+              f"dma_w={s['dma_write_count']:>3}  "
+              f"latency={lat:.1f} ns  "
+              f"(n_ops={len(res.engine.op_log)})")
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

src/kernbench/benches/1H_milestone_output/ccl/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
+buffer_kind,sip_topology,n_sips,n_elem,bytes_per_pe,latency_ns
-hbm,torus_2d,6,128,256,2120.040000000012
+hbm,torus_2d,6,128,256,3113.040000000012
-hbm,torus_2d,6,1024,2048,2717.2783333333473
+hbm,torus_2d,6,1024,2048,3710.2783333333527
-hbm,torus_2d,6,8192,16384,7315.184999999989
+hbm,torus_2d,6,8192,16384,8308.184999999929
-hbm,torus_2d,6,32768,65536,23081.26500000037
+hbm,torus_2d,6,32768,65536,24074.26500000037
-sram,torus_2d,6,128,256,2060.040000000012
+sram,torus_2d,6,128,256,3053.040000000012
-sram,torus_2d,6,1024,2048,2909.2783333333473
+sram,torus_2d,6,1024,2048,3902.2783333333573
-sram,torus_2d,6,8192,16384,9523.184999999869
+sram,torus_2d,6,8192,16384,10516.18499999987
-sram,torus_2d,6,32768,65536,32201.265000000385
+sram,torus_2d,6,32768,65536,33194.265000000385
-tcm,torus_2d,6,128,256,1964.040000000012
+tcm,torus_2d,6,128,256,2957.040000000012
-tcm,torus_2d,6,1024,2048,2477.2783333333473
+tcm,torus_2d,6,1024,2048,3470.2783333333527
-tcm,torus_2d,6,8192,16384,6403.185000000109
+tcm,torus_2d,6,8192,16384,7396.18499999999
-tcm,torus_2d,6,32768,65536,19865.265000000378
+tcm,torus_2d,6,32768,65536,20858.265000000378

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

@@ -1,37 +1,37 @@
-algorithm,sip_topology,n_sips,n_elem,bytes_per_pe,bytes_per_sip,latency_ns
+algorithm,sip_topology,n_sips,n_elem,bytes_per_pe,bytes_per_sip,latency_ns
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8,16,256,2666.552500000015
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8,16,256,3782.5525000000202
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,2747.7400000000152
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,3863.7400000000207
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,2855.990000000018
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,3971.9900000000216
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,3072.490000000019
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,4188.4900000000225
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,3337.1133333333582
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,4453.113333333365
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,3708.0333333333692
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,4824.033333333375
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,4449.873333333393
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,5565.873333333401
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,5933.020000000124
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,7049.020000000062
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,8900.379999999863
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,10016.379999999745
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,14835.099999999224
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,15951.099999999256
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,26704.540000000765
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,27820.540000000765
-lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,38573.97999999701
+lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,39689.97999999693
-lrab_hierarchical_allreduce,ring_1d,6,8,16,256,2365.255833333347
+lrab_hierarchical_allreduce,ring_1d,6,8,16,256,3524.2558333333504
-lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,2436.9433333333473
+lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,3595.943333333351
-lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,2532.526666666683
+lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,3691.5266666666857
-lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,2723.693333333349
+lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,3882.6933333333527
-lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,3048.635000000021
+lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,4207.101666666693
-lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,3393.4016666666957
+lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,4550.801666666699
-lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,4082.401666666714
+lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,5240.335000000055
-lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,5458.80166666677
+lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,6617.801666666763
-lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,8216.934999999943
+lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,9375.934999999827
-lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,13733.201666665835
+lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,14892.201666666035
-lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,24765.73500000064
+lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,25924.735000000648
-lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,35798.268333331536
+lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,36957.268333330976
-lrab_hierarchical_allreduce,torus_2d,6,8,16,256,1700.6025000000095
+lrab_hierarchical_allreduce,torus_2d,6,8,16,256,2693.6025000000113
-lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,1753.2900000000102
+lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,2746.290000000012
-lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,1823.540000000012
+lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,2816.5400000000127
-lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,1964.040000000012
+lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,2957.040000000012
-lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,2196.8183333333463
+lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,3189.81833333335
-lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,2477.2783333333473
+lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,3470.2783333333527
-lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,3038.1983333333583
+lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,4031.1983333333665
-lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,4159.5050000000665
+lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,5152.5050000000665
-lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,6403.185000000109
+lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,7396.18499999999
-lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,10890.5449999995
+lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,11883.544999999496
-lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,19865.265000000378
+lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,20858.265000000378
-lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,28839.98500000059
+lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,29832.98500000003

src/kernbench/benches/1H_milestone_output/gemm/gemm_composite_vs_async.json

@@ -0,0 +1,338 @@
+[
+  {
+    "M": 32,
+    "K": 32,
+    "N": 32,
+    "bench": "matmul-composite",
+    "variant": "load_ref",
+    "pe_window_ns": 171.394,
+    "engine_window_ns": 106.38400000000001,
+    "flops": 65536,
+    "tflops": 0.6160324860881335,
+    "n_records": 7
+  },
+  {
+    "M": 32,
+    "K": 32,
+    "N": 32,
+    "bench": "matmul-async",
+    "variant": null,
+    "pe_window_ns": 103.22199999999998,
+    "engine_window_ns": 71.71199999999999,
+    "flops": 65536,
+    "tflops": 0.9138777331548418,
+    "n_records": 4
+  },
+  {
+    "M": 32,
+    "K": 32,
+    "N": 32,
+    "bench": "matmul-async-chunked",
+    "variant": null,
+    "pe_window_ns": 103.22199999999998,
+    "engine_window_ns": 71.71199999999999,
+    "flops": 65536,
+    "tflops": 0.9138777331548418,
+    "n_records": 4
+  },
+  {
+    "M": 32,
+    "K": 32,
+    "N": 32,
+    "bench": "matmul-async-chunked-db",
+    "variant": null,
+    "pe_window_ns": 103.22199999999998,
+    "engine_window_ns": 71.71199999999999,
+    "flops": 65536,
+    "tflops": 0.9138777331548418,
+    "n_records": 4
+  },
+  {
+    "M": 32,
+    "K": 64,
+    "N": 32,
+    "bench": "matmul-composite",
+    "variant": "load_ref",
+    "pe_window_ns": 179.394,
+    "engine_window_ns": 106.38400000000001,
+    "flops": 131072,
+    "tflops": 1.232064972176267,
+    "n_records": 7
+  },
+  {
+    "M": 32,
+    "K": 64,
+    "N": 32,
+    "bench": "matmul-async",
+    "variant": null,
+    "pe_window_ns": 127.41399999999999,
+    "engine_window_ns": 87.904,
+    "flops": 131072,
+    "tflops": 1.4910811794685113,
+    "n_records": 4
+  },
+  {
+    "M": 32,
+    "K": 64,
+    "N": 32,
+    "bench": "matmul-async-chunked",
+    "variant": null,
+    "pe_window_ns": 127.41399999999999,
+    "engine_window_ns": 87.904,
+    "flops": 131072,
+    "tflops": 1.4910811794685113,
+    "n_records": 4
+  },
+  {
+    "M": 32,
+    "K": 64,
+    "N": 32,
+    "bench": "matmul-async-chunked-db",
+    "variant": null,
+    "pe_window_ns": 127.41399999999999,
+    "engine_window_ns": 87.904,
+    "flops": 131072,
+    "tflops": 1.4910811794685113,
+    "n_records": 4
+  },
+  {
+    "M": 32,
+    "K": 128,
+    "N": 32,
+    "bench": "matmul-composite",
+    "variant": "load_ref",
+    "pe_window_ns": 211.77800000000002,
+    "engine_window_ns": 122.76800000000003,
+    "flops": 262144,
+    "tflops": 2.135279551674703,
+    "n_records": 10
+  },
+  {
+    "M": 32,
+    "K": 128,
+    "N": 32,
+    "bench": "matmul-async",
+    "variant": null,
+    "pe_window_ns": 175.798,
+    "engine_window_ns": 120.28800000000001,
+    "flops": 262144,
+    "tflops": 2.1793030061186482,
+    "n_records": 4
+  },
+  {
+    "M": 32,
+    "K": 128,
+    "N": 32,
+    "bench": "matmul-async-chunked",
+    "variant": null,
+    "pe_window_ns": 185.81799999999998,
+    "engine_window_ns": 130.308,
+    "flops": 262144,
+    "tflops": 2.011726064401265,
+    "n_records": 8
+  },
+  {
+    "M": 32,
+    "K": 128,
+    "N": 32,
+    "bench": "matmul-async-chunked-db",
+    "variant": null,
+    "pe_window_ns": 185.81799999999998,
+    "engine_window_ns": 130.308,
+    "flops": 262144,
+    "tflops": 2.011726064401265,
+    "n_records": 8
+  },
+  {
+    "M": 32,
+    "K": 128,
+    "N": 128,
+    "bench": "matmul-composite",
+    "variant": "load_ref",
+    "pe_window_ns": 315.394,
+    "engine_window_ns": 226.38400000000001,
+    "flops": 1048576,
+    "tflops": 4.631846773623577,
+    "n_records": 34
+  },
+  {
+    "M": 32,
+    "K": 128,
+    "N": 128,
+    "bench": "matmul-async",
+    "variant": null,
+    "pe_window_ns": 394.102,
+    "engine_window_ns": 338.592,
+    "flops": 1048576,
+    "tflops": 3.096871751252245,
+    "n_records": 4
+  },
+  {
+    "M": 32,
+    "K": 128,
+    "N": 128,
+    "bench": "matmul-async-chunked",
+    "variant": null,
+    "pe_window_ns": 368.12200000000007,
+    "engine_window_ns": 312.6120000000001,
+    "flops": 1048576,
+    "tflops": 3.354241040011259,
+    "n_records": 8
+  },
+  {
+    "M": 32,
+    "K": 128,
+    "N": 128,
+    "bench": "matmul-async-chunked-db",
+    "variant": null,
+    "pe_window_ns": 368.12200000000007,
+    "engine_window_ns": 312.6120000000001,
+    "flops": 1048576,
+    "tflops": 3.354241040011259,
+    "n_records": 8
+  },
+  {
+    "M": 32,
+    "K": 3072,
+    "N": 32,
+    "bench": "matmul-composite",
+    "variant": "load_ref",
+    "pe_window_ns": 1701.4420000000007,
+    "engine_window_ns": 876.4320000000005,
+    "flops": 6291456,
+    "tflops": 7.178487321320988,
+    "n_records": 148
+  },
+  {
+    "M": 32,
+    "K": 3072,
+    "N": 32,
+    "bench": "matmul-async",
+    "variant": null,
+    "pe_window_ns": 2401.4620000000004,
+    "engine_window_ns": 1609.952,
+    "flops": 6291456,
+    "tflops": 3.9078531533859397,
+    "n_records": 4
+  },
+  {
+    "M": 32,
+    "K": 3072,
+    "N": 32,
+    "bench": "matmul-async-chunked",
+    "variant": null,
+    "pe_window_ns": 3264.8920000000217,
+    "engine_window_ns": 2473.3820000000214,
+    "flops": 6291456,
+    "tflops": 2.5436653133240017,
+    "n_records": 192
+  },
+  {
+    "M": 32,
+    "K": 3072,
+    "N": 32,
+    "bench": "matmul-async-chunked-db",
+    "variant": null,
+    "pe_window_ns": 3286.402000000022,
+    "engine_window_ns": 2494.8920000000217,
+    "flops": 6291456,
+    "tflops": 2.5217348085608298,
+    "n_records": 192
+  },
+  {
+    "M": 8,
+    "K": 128,
+    "N": 128,
+    "bench": "matmul-composite",
+    "variant": "load_ref",
+    "pe_window_ns": 291.394,
+    "engine_window_ns": 226.38400000000001,
+    "flops": 262144,
+    "tflops": 1.1579616934058943,
+    "n_records": 34
+  },
+  {
+    "M": 8,
+    "K": 128,
+    "N": 128,
+    "bench": "matmul-async",
+    "variant": null,
+    "pe_window_ns": 247.798,
+    "engine_window_ns": 216.288,
+    "flops": 262144,
+    "tflops": 1.212013611480988,
+    "n_records": 4
+  },
+  {
+    "M": 8,
+    "K": 128,
+    "N": 128,
+    "bench": "matmul-async-chunked",
+    "variant": null,
+    "pe_window_ns": 251.42399999999998,
+    "engine_window_ns": 214.914,
+    "flops": 262144,
+    "tflops": 1.2197623235340647,
+    "n_records": 8
+  },
+  {
+    "M": 8,
+    "K": 128,
+    "N": 128,
+    "bench": "matmul-async-chunked-db",
+    "variant": null,
+    "pe_window_ns": 251.42399999999998,
+    "engine_window_ns": 214.914,
+    "flops": 262144,
+    "tflops": 1.2197623235340647,
+    "n_records": 8
+  },
+  {
+    "M": 128,
+    "K": 8,
+    "N": 128,
+    "bench": "matmul-composite",
+    "variant": "load_ref",
+    "pe_window_ns": 462.01,
+    "engine_window_ns": 397.0,
+    "flops": 262144,
+    "tflops": 0.6603123425692695,
+    "n_records": 82
+  },
+  {
+    "M": 128,
+    "K": 8,
+    "N": 128,
+    "bench": "matmul-async",
+    "variant": null,
+    "pe_window_ns": 247.798,
+    "engine_window_ns": 216.288,
+    "flops": 262144,
+    "tflops": 1.212013611480988,
+    "n_records": 4
+  },
+  {
+    "M": 128,
+    "K": 8,
+    "N": 128,
+    "bench": "matmul-async-chunked",
+    "variant": null,
+    "pe_window_ns": 247.798,
+    "engine_window_ns": 216.288,
+    "flops": 262144,
+    "tflops": 1.212013611480988,
+    "n_records": 4
+  },
+  {
+    "M": 128,
+    "K": 8,
+    "N": 128,
+    "bench": "matmul-async-chunked-db",
+    "variant": null,
+    "pe_window_ns": 247.798,
+    "engine_window_ns": 216.288,
+    "flops": 262144,
+    "tflops": 1.212013611480988,
+    "n_records": 4
+  }
+]