Complete the composite-command study across both attention regimes and
write up the result, resolving when the composite command helps latency.
Decode (memory-bound, T_q=1, M=8): the command form is latency-neutral —
the kernel is bound by KV streaming and the MAC array is near-idle, so the
composite's only benefit here is host-issue offload (O(n_tiles) -> O(1)
PE_CPU commands). Figure: gqa_decode_long_ctx_composite (latency flat,
command count saturates).
Prefill (compute-bound, large M=G*T_q tile-filling): add three command-form
variants of a single-rank FlashAttention prefill kernel
(_gqa_prefill_compute_bound: primitive / composite / composite_extended).
Here the composite's scheduler-internal per-tile DMA<->compute pipeline
keeps the MAC array fed while the primitive's blocking tl.dot starves it,
so composite wins on MAC utilization (68% flat -> 83%) and wall-clock
(19% faster at ctx=1024), with the margin growing with context (deeper
P.V reduction = more tiles to pipeline) — the compute-bound mirror of the
GEMM result in section 3. Figure: gqa_prefill_compute_bound (latency +
MAC util). Sweep wired as GQA_1H_SWEEPS=prefill_cb; tests cover structure,
e2e, and composite<primitive in the compute-bound regime.
The synthesis: composite has two benefits set by roofline position —
host-issue offload (always) and MAC-array feeding (compute-bound only).
Decode exercises the first, prefill both.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Add two command-form variants of the Case-6 (Cube-SP x PE-SP) long-context
decode kernel alongside the primitive baseline:
- composite: per-tile GEMMs issued as one coarse tl.composite over the
full S_local (PE_SCHEDULER tiles/streams K,V) -- O(1) PE_CPU commands
vs the primitive kernel's O(n_tiles).
- composite_extended: Q.K^T composite + softmax_merge recipe composite
(ADR-0065), folding the per-tile online merge + P.V.
Extract the shared two-level (m,l,O) reduce into _gqa_mlo_reduce and
refactor the baseline to use it (byte-equal, guarded by a 64-rank
command-stream digest test).
Engine fix: _make_compute_out omitted pinned=True, so an on-chip TCM
compute result consumed as a composite GEMM operand was scheduled as an
HBM DMA_READ of its bit-61 scratch address -> PhysAddrError in data mode.
Mark compute outputs pinned (resident), matching the composite
auto-output and recipe scratch handles. .pinned is read only by the
tiling DMA gate, so this is byte-equal for existing benches.
Add the composite sweep (emit-level PE_CPU dispatch to 1M + data-mode
latency to 128K), its plot, umbrella wiring (GQA_1H_SWEEPS=composite),
and tests (refactor guard, dispatch saturation, engine-fix, e2e).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
ADR-0064 D8 broadened from "DMA fast-path" to "single-op-cmd fast-path":
every single-op command (DmaRead/DmaWrite/Gemm/Math/Copy) now pays the
lighter FIXED=8; only CompositeCmd keeps the 40-cycle control-path FIXED
(it alone needs scheduler plan generation + per-tile RW-hazard tracking +
completion wiring). Renamed knob fixed_per_dma_cmd_cycles ->
fixed_per_single_op_cmd_cycles. dispatch_cycles now branches on
"is CompositeCmd" rather than enumerating DMA types.
Term choice: "single-op" (not "atomic", which read as sync/async) — the
axis is composition (one engine op vs fused multi-op plan), orthogonal to
timing. single-op <-> composite.
Tests: test_pe_cost_model.py updated to the single-op surface (defaults,
fast-path over all 5 single-op cmd types, composite general path, yaml
override). All green.
Recalibrated tests/attention/test_gqa_decode_opt2.py
::test_opt3_dispatch_exceeds_opt2 — NOT a regression: D8 makes single-op
cmds 5x cheaper, so opt2's two-composite fusion win over opt3's many
single-ops narrowed from pre-D8 ~3.7x to ~1.87x (opt3=224 > opt2=120).
The CPU-offload invariant (opt2 cheaper) still holds; only the model-
dependent ">2x" constant was over-fit to the old uniform-40 model. Gate
now: direction + >1.5x margin (matches sibling R-sweep test's stated
"absolute ratio informative-only" philosophy).
NOTE for review: ADR-0065's "2x CPU-offload win" headline may want a
refresh to reflect the post-D8 ~1.87x — left to user (architectural doc).
Full regression: 826 passed, 1 skipped (tests/ excl. tests/gemm).
--- Remaining work (resume here if interrupted) ---
5. Re-run scripts/paper/paper_plot_gemm_async_vs_composite.py with new
cost model; verify async-tiled dispatch overhead drops (~4576ns ->
~1536ns expected) and the composite-vs-async-tiled gap narrows from
the prior ~6.3x at K=3072.
6. Copy regenerated gemm_composite_vs_async_tflops.png to
docs/report/1H-codesign-paper/figures/.
7. Paper §3.4 (03-gemm.tex sec:gemm-vs-async): finish naive->async-full /
chunked->async-tiled rename AND reframe FIXED_DMA wording to single-op
vs composite (currently still says "lighter FIXED for DMA descriptors,
FIXED_DMA=8"). Table 2 (02-platform) + §2 dispatch prose already done.
8. Paper §3.4 K=3072 corner para + mechanism #3: update dispatch breakdown
to new model (96 DMA*8 + 95 single-op*8 ≈ 1.5us vs old 4.6us); update
headline ratio if it changed.
9. Rebuild docs/report/1H-codesign-paper/build/main.pdf (tectonic) +
verify via pdftotext.
10. Then this is the bench-harness + paper commits (Groups 2 & 3).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Mirrors the existing decode 4-cases study for prefill on the LLaMA-3.1-70B
single-KV-head group (h_q=8, h_kv=1, d_head=128) across 8 cubes × 8 PEs:
Case 1 Cube-SP × PE-TP → KV S_kv-Ring across cubes; Q/O T_q-split 64-way
Case 2 Cube-Repl × PE-TP → full KV per cube; only CUBE 0's P PEs split T_q
Case 3 Cube-Repl × PE-SP → full KV per cube; PEs SP on S_kv (intra-cube AR)
Case 4 Cube-SP × PE-SP → KV split 64-way + Q T_q-split across cubes
(Ring KV + per-cube intra-CUBE AR) ★ optimal
Cases 2 and 3 use Q-axis tiling (TILE_Q) so the per-PE scratch is bounded
by TILE_Q × TILE_S_KV regardless of T_q (the full Q tensor would be
G·T_q·d_head·2 bytes which scales linearly with T_q and blew the 1 MB
budget at T_q≥128 without tiling).
Per-panel try/except in run_sweep tolerates per-panel failures so
sweep.json always lands with whichever cases succeed plus a failures
list. Case 4 uses configure_sfr_intercube_ring with the snake submesh
(2,4) which installs both the Ring E/W lanes and the intra_* lanes
needed for the intra-CUBE reduce — single SFR config for all 4 cases.
The plot script generates 4 PNGs (latency, traffic, memory, parallelism)
into 1H_milestone_output/gqa/long_ctx/. The parallelism chart shows
total compute work (active_PE × T_q_per_PE × S_kv_processed) — exposes
Case 3's 8× redundancy vs. Cases 1/2/4 which all do unique work.
gqa_prefill_long_ctx_4cases.py exposes run_sweep() rather than a @bench
decorator — the umbrella bench in the next commit invokes it.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Splits the GQA helpers into a dedicated subpackage to make room for the
prefill 4-cases study (next commit) and a single umbrella bench
(milestone-1h-gqa, after that).
Layout:
benches/gqa_helpers/
long_ctx/ — decode 4-cases kernels + sweep runner
short_ctx/ — prefill/decode short-context kernels
shared/ — _gqa_panel_helpers + decode_opt2 (context-agnostic)
The registry audit now skips subpackages so gqa_helpers/ (without a
leading underscore) doesn't get audited for @bench decorators.
Also drops the legacy milestone-gqa-headline bench, its
_gqa_attention_prefill_long kernel, 6 dependent prefill tests, the
paper_gqa_latency.py report harness, and the 3 stale headline-derived
PNGs the §6 wire-up referenced (paper will re-pull from the new
1H_milestone_output/gqa/long_ctx/ once §6 is updated).
The _ccl_cfg and _summarize_op_log helpers used to live in the
headline bench; extracted them to gqa_helpers/shared/_gqa_panel_helpers.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Bench changes:
- new _end_to_end_ns(op_log) and _engine_occupancy_ns(op_log) helpers
in milestone_gqa_decode_long_ctx_4cases.py (mirror paper_gqa_latency.py)
- _run_panel return dict now carries latency_ns + engine_occupancy_ns
alongside op_log_summary, so sweep.json is the single source of truth
for the comparative figures
Plot script:
- new scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py reads
sweep.json and emits 3 PNGs to docs/report/1H-codesign-paper/figures/:
gqa_decode_long_ctx_4cases_latency.png (end-to-end latency / case)
gqa_decode_long_ctx_4cases_traffic.png (ipcq/dma op counts / case)
gqa_decode_long_ctx_4cases_memory.png (KV bytes per cube / case)
Test changes:
- 2 new tests verifying the helpers + _run_panel dict shape
- lower smoke S_kv from 8192 -> 2048 (4x faster Case 2; assertions
are S_kv-independent; one fold-loop iteration preserved)
18/18 tests pass in ~4 min.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Option B fold: the lrab math previously in
``_gqa_attention_decode_long.py`` (used only as a thin re-export by
the Case 4 wrapper and as the import source for ``_merge_running``
in Cases 1 / 3) now lives inline in the Case 4 kernel file. ``sub_w``
is hardcoded to 4 (the 4×2 cube sub-mesh geometry; root cube 6);
the legacy ``sub_w=0`` 1D-chain backward-compat path is removed
along with its dedicated tests — the dropped ``single_user_*`` /
``multi_user_*`` panels that exercised it are already gone.
Production changes:
- inline lrab math into _gqa_attention_decode_long_ctx_cube_sp_pe_sp.py
(drops sub_w param; _ROOT_CUBE=6 baked in)
- inline _merge_running into Cases 1 (cube_sp_pe_tp) and 3
(cube_repl_pe_sp) so they no longer depend on the deleted file
- delete src/kernbench/benches/_gqa_attention_decode_long.py
- remove dead _run_decode_panel / _DECODE_* constants /
decode-side _PANEL_DISPATCH / _make_bench_fn decode branch
from milestone_gqa_headline.py (only the prefill panel remains)
- update _gqa_attention_decode_opt2.py docstring reference
Test changes:
- delete 7 legacy test_gqa_*.py files that pre-dated the 4-cases
architectural split (coverage now subsumed by the 16 4-cases tests)
- remove test_opt2_matches_opt3_data_mode + _run_decode_data helper
from test_gqa_decode_opt2.py (the parity check required sub_w=0
which no longer exists; opt2's other 4 tests preserved)
20/20 tests pass (16 4-cases + 4 opt2 smoke/dispatch).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Cases 1-3 each have a dedicated _gqa_attention_decode_<case>.py
kernel file; Case 4 previously reached into _gqa_attention_decode_long.py
via the headline bench's _run_decode_panel helper, breaking the
one-file-per-case convention. Adds _gqa_attention_decode_cube_sp_pe_sp.py
as a 20-line wrapper that bakes in sub_w=4 (the C=8 lrab geometry)
and gives Case 4 its own kind ("decode_cube_sp_pe_sp") and helper
(_run_decode_panel_cube_sp_pe_sp). decode_long.py is unchanged
(still serves the legacy decode_long tests). 16 tests pass.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds the fourth and final 4-cases panel: KV split S_kv-wise across
the 8 cubes (cube=row_wise, S_local = S_kv/C), replicated within
each cube (pe=replicate). PE-TP at B=1 means only PE 0 of each cube
has work; PEs 1-7 early-return (slide-11 PE-TP-at-B=1 waste). No
intra-CUBE comm; inter-CUBE 8-way reduce reuses the lrab-adapted
center-root pattern (root cube 6) — same structural cost as Case 4's
inter-CUBE phase (21 ipcq_copy). 16 tests pass (4 per case).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds the third 4-cases panel: KV replicated per cube (8× memory
waste), PEs SP on S_kv within each cube, intra-CUBE 8-way reduce on
(m, ℓ, O), and no inter-CUBE comm (every cube ends with full answer;
designated writer = cube 0). Reuses the row-chain + col-bridge
intra-CUBE pattern that anchors Case 4 (21 ipcq_copy per cube × 8
cubes = 168 total). 12 tests pass (4 Case 4 + 4 Case 2 + 4 Case 3).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Second case in the GQA decode 4-cases comparative study per
GQA_full_deck.pptx slide 11. Case 2 replicates K, V across all 8
cubes × 8 PEs (the slide-11 8 KB/tok/PE memory waste) and has zero
inter-rank comm. For B=1 (single-user decode default), only PE 0
of CUBE 0 has work — the inherent PE-TP waste slide 11 calls out.
Changes:
- New kernel: src/kernbench/benches/_gqa_attention_decode_cube_repl_pe_tp.py
Simplest of the 4 cases. Active rank loads full Q/K/V from HBM,
computes attention via S_kv tile sweep with online-softmax merge,
writes O. All non-(0,0) ranks early-return. No tl.send/recv.
- src/kernbench/benches/milestone_gqa_decode_4cases.py:
- Add panel single_kv_group_decode_gqa_cube_repl_pe_tp (Case 2)
to _PANELS and _PANEL_DISPATCH.
- Add _run_decode_panel_cube_repl_pe_tp helper: DPPolicy K/V/Q/O
= cube=replicate, pe=replicate (models 8× memory waste).
- Extend _make_bench_fn to dispatch kind="decode_cube_repl_pe_tp"
to the new runner.
- tests/attention/test_milestone_gqa_decode_4cases.py:
4 new tests assert Case 2 contract: panel registered, smoke
completion, zero ipcq_copy (no comm), single dma_write from cube 0.
Verification: 8/8 tests pass (4 Case 4 anchor + 4 new Case 2).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The legacy panel names suggested batched serving semantics they never
had — all four modeled a single user with KV sharded differently
(C=1 single-cube; C=4 multi-cube Cube-SP), at toy dims (T_q=4, S_kv≤128).
The single-KV-group C=8 panel + the new milestone-gqa-decode-4cases
bench cover the meaningful comparisons; pytest regression already
covers C=1/C=4 configurations end-to-end at richer scale.
Changes:
- milestone_gqa_headline.py: drop the 4 legacy panels; _PANELS now
contains only single_kv_group_prefill_gqa_c8_p8. Update docstring.
- tests/attention/test_milestone_gqa_headline.py: drop the 3 legacy-
panel architectural tests (Ring-KV traffic, root-only decode write,
per-CUBE distributed output) and test_decode_panels_use_real_gqa
(no decode panels in this bench anymore). Equivalent properties
are asserted in test_milestone_gqa_single_kv_group_prefill_panel.py
(64 dma_writes, 896 ipcq_copy) and test_milestone_gqa_decode_4cases.py
(1 dma_write at cube 6, 189 ipcq_copy).
- tests/attention/test_milestone_gqa_single_kv_group_prefill_panel.py:
drop test_existing_prefill_panel_runner_backward_compat (it exercised
multi_user_prefill_gqa which no longer exists).
- scripts/paper/paper_plot_gqa.py: replace the 4 legacy _LABELS entries
with the single single_kv_group_prefill_gqa_c8_p8 label.
- Regenerate 1H_milestone_output/gqa_headline/sweep.json from the new
panel set.
Verification: 9/9 tests pass across the 3 affected test files.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
First milestone of the decode 4-cases comparative study per
GQA_full_deck.pptx slides 11-17. Case 4 (Cube-SP × PE-SP, the optimal
case per slide 11) is structurally the existing _gqa_attention_decode_long
at sub_w=4 (Increment 2's lrab-adapted center-root reduce). This commit
wires it into a dedicated bench so the remaining cases land alongside.
Changes:
- New bench: src/kernbench/benches/milestone_gqa_decode_4cases.py
Houses the 4 case panels under one milestone-gqa-decode-4cases entry
(gated by GQA_DECODE_4CASES_RUN=1; output to
1H_milestone_output/gqa_decode_4cases/sweep.json). Cases 1-3 are
TBD in subsequent sub-increments (5C.A/B/C).
- New panel: single_kv_group_decode_gqa_cube_sp_pe_sp
C=8, P=8, sub_w=4, T_q=1, S_kv=131_072, d_head=128, h_q=8, h_kv=1.
- src/kernbench/benches/milestone_gqa_headline.py: _run_decode_panel
extended with keyword-only sub_w/T_q/d_head/h_q/h_kv overrides
(defaults preserve existing-panel behaviour).
- tests/attention/test_milestone_gqa_decode_4cases.py: 4 new tests
asserting registration, smoke completion, reduce-to-root at the lrab
center cube (cube 6), and the predicted 189-ipcq Case-4 traffic
pattern (168 intra-CUBE + 21 inter-CUBE lrab Phase 1+2).
- tests/attention/test_milestone_gqa_headline.py: rename
test_sweep_json_has_four_panels -> test_sweep_json_has_expected_panels
and switch hardcoded 4 to len(PANELS) (the panel set grew to 5
with Increment 5's single_kv_group_prefill_gqa_c8_p8).
Deviation noted: slide 13 prescribes AllReduce on (m,ℓ,O); our kernel
does reduce-to-root (only the lrab center cube has the answer) per
ADR-0060 §4. Treated as the kernbench Case-4 baseline.
Verification: all 4 new tests pass; 90 regression tests pass; the
previously-failing test_sweep_json_has_four_panels now passes under
its renamed form.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The headline T_q=S_kv=32K target overflows the 1 MB per-PE scratch
pool: at T_q_local=4K and tile_s=1024 the scores matrix alone is 8 MB.
The prefill kernel's bootstrap section also leaves K_t/V_t/scores/
exp_scores persistent (outside tl.scratch_scope), inflating baseline.
Scale-down to T_q=S_kv=1K (T_q_local=128, fits comfortably) preserves
the C=8 + P=8 architecture demonstration; the true LLaMA 32K headline
awaits a future increment to add Q-axis tiling and tighten bootstrap
scratch discipline.
Verified end-to-end: kernbench run --bench milestone-gqa-headline now
produces sweep.json with all 5 panels. The new panel shows
ipcq_copy=896 (matches (C-1)·n_tiles·2·C·P = 7·1·2·8·8) and
dma_write=64 (one per PE, head-parallel + intra-CUBE PE-SP).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
End-to-end wires C=8 P=8 d_head=128 prefill with snake-ring inter-CUBE
SFR, intra-CUBE PE-SP (all 64 ranks active), and the milestone bench
panel. Decode kernel gains lrab-adapted center-root reduce for the
2×4 sub-mesh per ADR-0060 §4.2.
Increment 1 — SFR multi-row snake
src/kernbench/ccl/sfr_config.py: configure_sfr_intercube_ring gains
submesh_shape / submesh_origin kwargs; installs a Hamiltonian snake
ring through a rectangular sub-mesh (every hop is 1-hop physical
neighbour). Backward-compat: 1D-row behaviour preserved when
submesh_shape is None.
tests/test_intercube_snake_ring.py (12 tests)
Increment 2 — Decode lrab-adapted center-root reduce
src/kernbench/benches/_gqa_attention_decode_long.py: new sub_w param
(default 0 = existing 1D-chain). sub_w >= 2 selects the ADR-0060
§4.2 prescribed lrab-adapted Phase 1+2 reduce (bidirectional row +
bidirectional col converge to the center cube), with log-sum-exp
_merge_running replacing the plain + of lrab.
tests/attention/test_gqa_decode_long_2d_reduce.py (4 tests)
Increment 3 — Prefill kernel at C=8 (no production change)
Verified by inspection that the existing prefill_long kernel +
Increment 1's snake SFR already work at C=8 without any kernel
edit. The kernel speaks logical W/E; the snake routes it.
tests/attention/test_gqa_prefill_long_c8_snake.py (3 tests)
Increment 4 — Intra-CUBE PE-SP in prefill (all 64 ranks)
src/kernbench/benches/_gqa_attention_prefill_long.py: new P param
(default 1 = existing PE-0-only). P > 1 splits T_q query-axis-wise
across the P PEs of each CUBE; output rows are disjoint per PE so
no intra-CUBE reduce is needed; each PE drives its own same-lane
ring (P parallel rings).
tests/attention/test_gqa_prefill_long_pe_sp.py (5 tests)
Increment 5 — LLaMA-scale milestone bench panel
src/kernbench/benches/milestone_gqa_headline.py: new panel
single_kv_group_prefill_gqa_c8_p8 (C=8, P=8, T_q=S_kv=32K,
d_head=128). _run_prefill_panel extended with P/T_q/d_head
defaults; routes snake SFR when C > mesh_w.
tests/attention/test_milestone_gqa_single_kv_group_prefill_panel.py (3 tests)
Total: 4 production files modified, 5 new test files, 27 new tests.
Followed the Phase 1/2 protocol per CLAUDE.md throughout.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
composite() now awaits its prologue recipe operands (e.g. the score tile from a prior #1 composite) so the cross-composite RAW (#1 writes scores -> #2 reads them) is ordered: #2 waits for #1 before reading. With N1-N3, opt2 (the two-composite recipe kernel) now computes a correct attention output end-to-end in data mode and matches opt3 within fp tolerance.
Scope note: the e2e opt2<->opt3 parity holds in the near-uniform-attention regime (small inputs). kernbench's K reshape-as-transpose (opt3 kernel docstring: 'correct for zero/symmetric inputs') makes opt2's per-sub-tile K interpretation differ from opt3's single-tile one for high-contrast inputs — a pre-existing kernbench limitation shared by both kernels, not the recipe. The EXACT recipe math (online-softmax merge m/l + O = O*corr + P@V) is verified non-trivially and exactly by N2/N3 (test_recipe_data_mode.py). Suite 817 pass / 3 pre-existing fail.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Extend the opt2 R-sensitivity test to assert the opt3/opt2 dispatch ratio increases as R decreases (4.20x -> 5.22x -> 5.45x at R=0.25/0.0625/0.03125) — the cost is FIXED-dominated (command-count-driven), per ADR-0064 Test #9. Test-only; no production change (the P0 cost model + P2 recipe already make the ratio measurable).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
New _gqa_attention_decode_opt2.py: per-tile attention as #1 Q.Kt GEMM composite + #2 softmax_merge recipe composite (online merge + P.V + add). Runnable in op_log mode; full data-mode numeric parity is a deferred follow-up (P5-numerics).
Measured per-tile PE_CPU dispatch (ADR-0064 Rev2 default): opt3=960ns vs opt2=184ns = 5.22x (gate >2x), FIXED-dominated (command-count reduction) — the ADR-0060/0064 CPU-offload win. opt2<opt3 across R in {0.25,0.0625,0.03125}. K-before-V: softmax_merge prologue carries no DMA; V (ref) streams only in the GEMM tile loop.
Fix latent P3 bug surfaced by the first e2e recipe run: prologue MATH stages were FOLDED into the first GEMM tile, but a folded MATH->DMA_READ boundary needs a PE_MATH->PE_DMA token route the pipeline never wires (KeyError pe_dma). Now prologue/post-loop stages are fed as standalone 1-stage tiles (each completes on its component); completion count includes them. Existing benches have no prologue -> tiles + feed order unchanged (byte-equal); full suite 806 pass / 3 pre-existing fail.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Three logically distinct changes, bundled for atomic test green:
1. **P3c — prefill_long tile-granular Ring KV** (ADR-0060 §5.5.1 amendment).
Convert the ring from slice-granular (one full ``(d_head, S_local)``
KV slice per step) to tile-granular (``n_tiles`` tiles of
``TILE_S_KV`` per step). Nested loop with outer tile, inner ring step:
each tile propagates through all C ring positions before the next
tile starts, so IPCQ in-flight depth stays at 1 per direction.
Bootstrap at ``(t=0, k=0)`` outside the scratch_scope establishes the
persistent ``(m, ℓ, O)``; every other iteration scope-wraps + persists
via ``copy_to``. Per-rank persistent scratch shrinks to ~1 KB; per-tile
scope bounded by TILE_S_KV regardless of S_local. Headline:
prefill_long now completes at S_kv=128K (previously overflowed).
New: ``tests/attention/test_gqa_prefill_long_tile_ring.py``
(3 tests — ceiling-lift + tile-granular ipcq_copy count +
per-CUBE distributed output regression guard).
2. **Rename ``gqa_*`` → ``gqa_attention_*``** across kernel files,
function names, and importers. The "attention" name makes the role
explicit (GQA is grouped-query attention) and matches upstream Triton
FlashAttention naming conventions. Renames:
_gqa_decode_long.py -> _gqa_attention_decode_long.py
_gqa_decode_short.py -> _gqa_attention_decode_short.py
_gqa_prefill_long.py -> _gqa_attention_prefill_long.py
_gqa_prefill_short.py -> _gqa_attention_prefill_short.py
And function names ``gqa_<phase>_<context>_kernel`` →
``gqa_attention_<phase>_<context>_kernel``. Updated 1 bench file
(milestone_gqa_headline.py) and 10 test files.
3. **ADR-0060 / 0062 / 0063 / 0064: Proposed → Accepted**.
All four are reflected in production code and covered by tests:
- ADR-0060 (GQA fused attention): 4 kernels deployed; §5.5.1
amendment added for the tile-granular Ring KV introduced by P3c
(EN + KO mirror).
- ADR-0062 (lazy tl.load): LoadFuture + _await_pending live in
tl_context.py.
- ADR-0063 (tl.scratch_scope + tl.copy_to): used in every chain
reduce + tile sweep + ring step. EN-only previously; KO
translation authored as part of this commit (CLAUDE.md
bidirectional rule).
- ADR-0064 (per-op-type CPU issue cost): cpu_issue_cost.py +
issue_cost_table wiring in tl_context.py (Phase E).
Files git mv'd from docs/adr-proposed/ to docs/adr/ (EN) and
docs/adr-ko/ (KO). ADR-0061 (tl.broadcast) stays Proposed — no
implementation; documented as optional convenience primitive in
the ADR itself.
Tests: 88/88 focused regression green
(tests/attention/ + Phase E + TL discipline).
ADR pair verification: ``python tools/verify_adr_lang_pairs.py`` OK.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Replace each kernel's local one-shot partial with a Tile-0 bootstrap + a
``scratch_scope``-wrapped merge loop. Per-rank scratch is now bounded by
``TILE_S_KV = 1024`` regardless of ``S_local``, lifting the long-context S
ceiling. Backward compatible at ``S_local <= TILE_S_KV`` (loop body is
empty; op_log structurally identical to today).
Framework: extend ``memory_store.read`` to support partial reads at
offsets within a stored region. Models real Triton ``tl.load(ptr+offset,
shape=...)`` — needed because each tile loads ``k_ptr + tile_start*row_bytes``
for its sub-slice of the per-rank KV region.
prefill_long is intentionally left untiled. Its ring loop carries
full-slice ``Kc``/``Vc`` for ``tl.send`` between CUBEs, so tile-sweeping
step 0 wouldn't shrink the kernel's actual scratch footprint. Lifting
prefill_long's ceiling requires a tile-granular ring rewrite — separate
phase. Docstring + inline comments cleaned to reflect the current shape.
Docstrings across all 4 GQA kernels: drop P1a/P2a/P2b/P6a/P6b lineage
paragraphs and historical deviation lists; describe each kernel by what
it does, not how it got here. Add explicit ``# Local attention`` /
``# Communication`` section headers.
Tests: new ``tests/attention/test_gqa_tile_sweep.py`` with 5 tests
covering single-tile op_log stability, multi-tile ``copy_to`` emission
across all 3 refactored kernels, and the headline ceiling-lift
(decode_long at S_kv=256K). Full focused regression green: 85/85 across
``tests/attention/`` + Phase E + TL discipline tests.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Make the file + function naming symmetric:
_gqa_decode.py -> _gqa_decode_long.py
_gqa_prefill.py -> _gqa_prefill_long.py
gqa_decode_kernel -> gqa_decode_long_kernel
gqa_prefill_kernel -> gqa_prefill_long_kernel
Mirrors the existing _gqa_{decode,prefill}_short.py naming. Updates the
two imports + two call sites in milestone_gqa_headline.py and the 9
attention tests that import the kernels.
Tests: 72/72 focused regression green (tests/attention/ + Phase E + TL
discipline). milestone-gqa-headline bench passes its 7 panel/schema
assertions.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Land the new GQA fused-attention kernels (ADR-0060) for prefill/decode
across long and short context, the TL discipline primitives they depend
on (ADR-0062 lazy load, ADR-0063 scratch_scope + copy_to), and the
per-op-type CPU issue cost model (ADR-0064). Remove the pre-ADR-0060
mesh-attention baseline now that the unified kernels supersede it.
ADR-0060 (long context)
- _gqa_decode.py: M-fold + 2-level chain reduce-to-root (Level-2
intra-CUBE row-then-col + Level-1 inter-CUBE) — root-only output.
- _gqa_prefill.py: head-parallel + Ring KV rotation around C CUBEs,
online-softmax merge per ring step, per-CUBE distributed output.
- Each merge stage wraps in scratch_scope() and persists running
(m, l, O) via copy_to() to lift the 1 MiB scratch ceiling.
ADR-0060 §B.split.2 (short context, kv_per_cube in {1,2,4,8})
- _gqa_decode_short.py / _gqa_prefill_short.py: no cube-SP; each CUBE
owns whole KV heads; PE-parallel heads with intra-group chain
reduce. Prefill has no Ring KV (each head fully resident).
ADR-0062 (lazy tl.load): future-bearing TensorHandle, auto-wait at
first consuming op (dot/MATH/store/send/copy_to/composite).
ADR-0063 (tl.scratch_scope + tl.copy_to): scoped per-tile arena with
copy_to writeback primitive for persistent running state.
ADR-0064 (CPU issue cost model)
- common/cpu_issue_cost.py: per-op-type table (composite=40 ns,
primitives=5 ns); ratios are load-bearing per D1.
- TLContext: issue_cost_table param; _emit_dispatch_overhead(kind)
consults table with dispatch_cycles fallback (ADR-0046 §D6
back-compat).
- Live PE_CPU paths (greenlet + legacy) construct TLContext with
DEFAULT_CPU_ISSUE_COST so saturation lever (ADR-0060 §1) is
measurable end-to-end.
P7 headline bench: milestone-gqa-headline writes per-panel
op_log_summary to 1H_milestone_output/gqa_headline/sweep.json. No
figure renderers yet (deferred).
Removals (pre-ADR-0060 baseline now superseded):
- benches: _attention_mesh_kv.py, _attention_mesh_mlo.py,
_attention_mesh_mlo_2d.py, milestone_gqa_llama70b.py
- tests: test_attention_*, test_mesh_*, test_milestone_gqa_llama70b
- topology: llama70b_4sip.yaml (only consumer was the deleted diag)
- artifacts: 1H_milestone_output/gqa/ (sweep.json + 5 PNGs)
- tests/gqa/ plot helper + test (broken on Windows Tcl/Tkinter)
- ADR-0060/0061 references to deleted file paths cleaned up
(EN + KO kept in sync).
Tests: 124/124 focused regression green (attention + Phase E + TL
discipline + triton_emu + pe_components). Full regression: 764 pass,
2 pre-existing test_bench_registry failures (stale EXPECTED_NAMES
across multiple benches, not introduced here).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
New ``test_attention_8kv_groups_diag.py`` probes the path from
validation scale to the GQA Llama-70B 1-Q-head-per-cube headline
(64 cubes = 8 KV-groups x 8 cubes/group on a 4-SIP topology) in
three incremental steps:
step_1_single_kv_group_at_full_breadth
ONE 2x4 multi_user_decode launch (8 cubes) via the 2D mesh-mlo
kernel. Verifies the per-KV-group 2D AllReduce works at full
breadth.
step_2_four_kv_groups_one_per_sip
Four sequential 2x4 launches, one per SIP. Uses the CCL
milestone's set_device pattern (milestone_1h_ccl.py:283-292):
ONE run_bench with ``target_device="all"`` and
``ctx.ahbm.set_device(sip)`` between launches — not four
separate run_bench calls with ``DeviceSelector("sip:N")``
(which would misuse DeviceSelector and hit a
``DPPolicy x target_device`` allocator mismatch).
step_3_eight_kv_groups_two_per_sip
Two 2x4 launches per SIP x 4 SIPs = 8 KV-groups (64 cubes
total). Second launch per SIP uses ``cube_start=8`` to address
cubes 8..15 — the disjoint sub-mesh that was unreachable before
``DPPolicy.cube_start`` landed. Demonstrates the headline-enabling
use of cube_start end-to-end (DPPolicy + kernel kwarg).
All three steps pass at validation dims (S_q=1, S_kv=16, h=1,
d_head=64). Headline-scale dims, multi_user_prefill 2D, and the B=8
"Batch on batch" PE parallelism remain follow-up work.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
New ``_attention_mesh_mlo_2d.py`` decomposes a
``(mesh_rows x mesh_cols)`` cube sub-mesh into two stages of
bidirectional AllReduce-mlo:
Stage 1 — row reduce (E/W edges, mesh_cols-1 steps)
Stage 2 — col reduce (N/S edges, mesh_rows-1 steps)
After both stages every cube holds the same final ``(m, l, o)`` and
writes the normalized output. The online-softmax mlo merge is
associative, so row-then-col partitioning is mathematically
equivalent to a 1D ring AllReduce-mlo over all
``mesh_rows * mesh_cols`` cubes but uses fewer hops:
- 2x4 (8 cubes / KV-group): 4 steps vs 7 (1.75x faster)
- 4x4 (16 cubes / full SIP): 6 steps vs 15 (2.5x faster)
Motivation: the original 1D ring kernel ``_attention_mesh_mlo.py``
hit ``IpcqInvalidDirection`` at cube 4 when ``n_ranks=8`` on the
4x4 cube mesh — cube 4 has no W neighbor at the row 0/1 boundary.
N/S edges are already installed by ``configure_sfr_intercube_multisip``
so the 2D kernel runs on existing wiring without SFR changes.
The kernel accepts ``cube_start: int = 0`` and subtracts it from
``program_id(axis=1)`` so the ring math uses launch-local rank. This
matters because kernbench's ``program_id(axis=1)`` returns the
physical cube id (ADR-0022), so a launch starting at cube 8 would
otherwise compute ``my_row = 8//4 = 2`` (out of sub-mesh bounds) and
deadlock. Default ``cube_start=0`` keeps the existing
multi_user_decode validation behavior bit-for-bit.
Bench dispatch: ``multi_user_decode`` in milestone-gqa-llama70b now
uses the 2D kernel via a new ``mesh_shape`` column in
``_PANEL_DISPATCH``. At validation ``N_RANKS_MULTI_USER=4``, the
shape is ``(1, 4)`` — a degenerate single-row mesh, equivalent in
step count and op_log structure to the prior 1D ring at n_ranks=4.
The other three panels keep their 1D kernels.
Tests: 4 new unit tests in ``test_mesh_mlo_2d_correctness.py`` —
1x4 (degenerate row), 2x4 (8-KV-group target), 4x4 (full SIP), and
2x4 at cube_start=8 (the second sub-mesh per SIP). Existing
milestone (12 tests) and mesh-kernels-rank-axis (7 tests) suites
stay green — no regression.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Self-contained eval bench (ADR-0054) that drives the four GQA Llama-70B
panels through run_bench with enable_data=True at validation scale and
emits sweep.json with the v1 schema (ADR-0057 D7).
Panel dispatch table maps each panel to (kernel, SFR install, S_q,
n_ranks, rank_axis):
single_user_prefill mesh_kv_kernel, intracube_pe_ring, S_q=16, n=8, rank_axis=0
multi_user_prefill mesh_kv_kernel, intercube_multisip, S_q=16, n=4, rank_axis=1
single_user_decode mesh_mlo_kernel, intracube_pe_ring, S_q=1, n=8, rank_axis=0
multi_user_decode mesh_mlo_kernel, intercube_multisip, S_q=1, n=4, rank_axis=1
multi_user panels pass _auto_dim_remap=False (avoid d_head=64
colliding with K's global M=64) and rank_axis=1 (cube-level ring,
gates 7 of every 8 PEs to silence).
Each panel runs on a fresh per-config GraphEngine, then op_log is
summarized into gemm/dma/ipcq counts. Both decode panels emit exactly
2*n_ranks GEMMs (one-shot partial attention per rank, ADR-0056 D3).
v1 supports GQA_VALIDATION=1 only; headline mode + figures deferred to
sub-cycles 4b/4c. Sentinel tensor satisfies the run_bench
"at least one request" contract (ADR-0045 D4 / ADR-0054 D2 carve-out).
Tests: tests/attention/test_milestone_gqa_llama70b.py — all 12 pass.
Includes committed sweep.json baseline at the bench's _OUTPUT_DIR so
subsequent test runs reuse it instead of re-simulating.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
ADR-0059 single_user_* panels run the ring across PEs in one cube
(rank == tl.program_id(axis=0)). multi_user_* panels run the ring
across cubes — rank should be cube_id (axis=1), and 7 of every 8 PEs
in each cube must stay silent because the cube-level SFR install only
gives the cube-coordinate PE 0 an E/W neighbor.
Add ``rank_axis: int = 0`` kwarg to both ``attention_mesh_mlo_kernel``
and ``attention_mesh_kv_kernel``:
- 0 (default): rank == tl.program_id(axis=0). Existing single_user
behavior, all spec tests unchanged.
- 1: gate ``if tl.program_id(axis=0) != 0: return`` at kernel start,
then ``rank = tl.program_id(axis=1)``. multi_user_* panels pass
this to the kernel via ctx.launch positional arg.
Also brings in _attention_mesh_kv.py and _attention_mesh_mlo.py as
the committed home of the ADR-0059 kernels (previously living
uncommitted in the working tree from sub-cycle 4b).
Tests: 7-test rank_axis spec file (default-path + rank_axis=1 gating
and cube-id semantics, both kernels); 4-panel diag harness now green
end-to-end (single_user_prefill/decode + multi_user_prefill/decode);
763-test wider sweep clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two compounding bugs in ctx.launch's dim-translation path surfaced
by multi_user_* panels of milestone-gqa-llama70b (sub-cycle 4c step 2):
Bug A: _compute_local_shape divided by self._num_cubes (the topology's
cube count, 16 in default topology.yaml) instead of the DPPolicy's
effective num_cubes (4 for validation-scale multi_user). The tensor
allocator at context.py:471-484 already honored dp.num_cubes; the
parallel computation inside launch was out of sync. Fix mirrors the
allocator's eff_num_cubes precedence pattern.
Bug B: dim_map was keyed by value, so any scalar whose value
coincidentally equaled a global tensor dim got rewritten to that dim's
local value — e.g. d_head=64 colliding with K's global M=64 in
multi_user mode. Legacy bench kernels (va_offset etc.) rely on this
remap, so the fix is opt-out: ctx.launch(..., _auto_dim_remap=False)
preserves scalars exactly as passed. Default remains True.
Tests: 3 new dim-translation tests + 4-panel diag harness covers
single_user_* (PASS) and multi_user_* (advances to new SFR/axis layer
failure, tracked separately). va_offset + full attention spec suite
unchanged.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
eusonice
ywkang
mukesh