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>
Mid-pass through §2. Captured progress so far:
- §2 intro previews the three new threads (graph view, DES engine,
congestion) the rewritten latency section weaves together.
- §2.3 retitled to "Latency model: a graph traversed by events" and
restructured into bold-led paragraphs:
- The hardware as a graph (nodes = components, edges = links).
- From graph to discrete-event simulation (node/edge events,
deterministic ordering, correlation-ID trace).
- Latency contributions (per-node fixed, per-edge size-aware,
per-service occupancy).
- Congestion: per-edge FIFO BW occupancy, HBM per-PC parallelism,
component serial workers — the mechanisms that surface real
bottlenecks instead of peak-BW roofline.
- Control-plane cost model (FIXED + b·R) — unchanged.
- Accuracy: extended with a second cross-check from the all-reduce
study (torus vs. analytic startup-plus-per-packet model + FSIM
external single-device reference). Existing GEMM analytic-vs-
measured 10-20% check retained.
- §2.4 Hardware-configuration table unchanged.
Still TODO: cross-check accuracy numbers against a fresh
milestone-1h-gemm/ccl re-run; the current artifacts predate the
ADR-0064 Rev2 cost-model and IPCQ Phase-2 race fix and may need a
refresh before §2.3 final lock-in.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
- §1 Introduction rewritten with AHBM-first opening, in-paper KernBench
platform overview, GQA motivation and three architectural requirements,
three matching HW-SW co-design mechanisms (GQA-aware placement,
PE_IPCQ, composite-command GEMM under PE_SCHEDULER), broader-
applicability outlook (PE_IPCQ as collective substrate, composite
command for fused operator pipelines), and outline. Consistent with
the Executive Summary structure.
- Executive Summary: extended broader-applicability closing with a
forward-looking line previewing FFN/MoE integration.
- §5 section heading renamed from
"All-Reduce Acceleration via PE_IPCQ" to
"PE_IPCQ and Collective Communication".
- Title page date line refined to "2026 H1 Report".
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
test_bench_registry EXPECTED_NAMES: add the 3 milestone benches the pull registered (milestone-1h-ccl, milestone-1h-gemm, milestone-gqa-headline) — 11 benches, alphabetical. test_memory_store::test_shape_mismatch_raises: the pull made same-size reshapes byte-conserving (allowed) and an over-large read raise 'Out-of-bounds read' (not 'Shape mismatch') — match the new message. Full suite now 820 passed / 1 skipped / 0 failed.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The reconciliation table, corrected file/test lists, measured 5.22x ratio, and the new section 15 (data-mode numerics D8 + N1-N4) — the content the git-mv in 30a0451 left uncommitted.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Updated the detailed design to match the shipped implementation: an as-built reconciliation table (P1-first, pinned-based DMA, D7 hard-cap-error, D8 output handle, prologue 1-stage-tile feed, D6.7 prologue-scoped), corrected file/test lists, the measured 5.22x dispatch ratio, and a new section 15 documenting the data-mode numeric subsystem (D8 + N1-N4) the original DDD omitted. Status Draft->Accepted; moved docs/adr-proposed -> docs/ddd/.
Co-Authored-By: Claude Opus 4.8 (1M context) <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>
Two fixes make the recipe's P.V GEMM fold into the rescaled accumulator: (1) op_log marks the composite gemm record accumulate=True when an 'add' epilogue targets the GEMM's own output; data_executor._execute_gemm then does O = O_existing + a@b instead of overwriting. (2) PE_SCHEDULER records the composite's numeric op at COMPLETION (in _retire_on_done) instead of at dispatch, so the GEMM's t_start sorts AFTER its prologue MATH ops -> it reads the computed P and the rescaled O.
Verified end-to-end (DataExecutor): the full recipe produces O = O_old*corr + P@V matching a numpy flash-attention reference. Suite 816 pass / 3 pre-existing fail.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
data_executor._compute_math gains the 5 recipe ops (rmax/rsum as keepdims reductions, max_elem/exp_diff/mul_bcast as binary; numpy broadcasting covers bcast_axis). tiling._math_stage now carries operand+output addrs/shapes/spaces + axis on the recipe prologue MATH stages. op_log.record_end promotes a MATH stage to op_kind='math' ONLY when it carries input_addrs -> the DataExecutor runs it; legacy epilogue MATH stages (bias/relu, no addrs) stay op_kind-opaque -> byte-equal.
Fixed a latent P2 lowering bug exposed by data mode: _resolve_recipe_dst gave reductions shape (G,) (axis removed) but max_elem broadcasts them against the running m=(G,1); reductions now keep the reduced axis as size 1 (keepdims) -> (G,1). Verified end-to-end: running the recipe's 8 MATH ops through the DataExecutor produces m, l (fully updated online-softmax) and O (rescaled by corr) matching a numpy reference. Suite 815 pass / 3 pre-existing fail.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Two changes make a composite's matmul replay in Phase 2: (1) op_log _extract_op_info(CompositeCmd) now SCANS ops for the gemm op (it is ops[0] for a legacy composite but sits after the prologue MATH ops in a recipe), emitting one composite_gemm record with the gemm's a/b/out addrs+spaces; (2) PE_SCHEDULER._dispatch_composite records the composite as a zero-duration numeric op (no-op without an op_logger, so latency-only mode is unchanged). Verified: a composite GEMM with seeded inputs writes a@b to its HBM output.
DataExecutor._execute_gemm is now best-effort: if an operand's data was never produced (torch.empty bench input) or a shard read is out-of-bounds (column-wise sharded operand read at full shape), skip the matmul instead of crashing. This keeps composite-using benches (qkv-gemm, composite-epilogue) green now that their previously-uncomputed GEMM actually fires. Suite 813 pass / 3 pre-existing fail.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
tl.composite now returns the output TensorHandle (not a CompletionHandle) so its result chains like tl.dot's; the handle carries the completion in a CompositeFuture pending so downstream ops and tl.wait auto-await it. out is a handle: out=tl.ref(addr,shape) (HBM, DMA_WRITE inside the composite) or an in-place TCM handle (STORE only); omitted -> TLContext auto-allocates a TCM scratch. out_ptr kept as HBM shorthand (= out=tl.ref(out_ptr, shape)) to avoid churning ~30 existing call sites.
tl.ref now returns space=hbm (it references HBM data; operand-input DMA stays pinned-based per D4 so input streaming is unchanged). tiling: the tile loop's DMA_WRITE is gated on out.space==hbm (out analog of the operand pinned rule) — a TCM output stays on-chip (chainable) and its high-bit scratch address no longer hits the DMA PA decoder.
Fixes the opt2 data-mode crash: the recipe accumulator O is TCM -> no DMA_WRITE -> opt2 now RUNS end-to-end in data mode (enable_data=True). Numeric parity of the recipe MATH ops is the next step. Suite 812 pass / 3 pre-existing fail; existing composite benches use out_ptr->hbm->DMA_WRITE unchanged (byte-equal).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
New `/paper` slash-command skill that synthesizes ADR/SPEC content and live
KernBench benchmark results into a sectioned LaTeX technical paper compiled
to PDF with Tectonic (auto-installed). The skill negotiates a TOC, grounds
every number in committed artifacts or fresh bench runs, and keeps
report-only benches isolated.
This commit also includes the first generated report:
- docs/report/1H-codesign-paper/ — main.tex + per-section .tex, figures,
toc.md contract, and the built 8-page main.pdf. Covers the platform
(source-level kernels, latency model + accuracy, HW config from
topology.yaml), GEMM via composite command, All-Reduce via PE_IPCQ, and
fused GQA combining both, plus discussion/conclusion/2H future work.
- scripts/paper/ — isolated report harnesses (not registered benches):
paper_gqa_latency.py harvests per-panel GQA end-to-end latency + engine
occupancy (the milestone only emitted op-counts); paper_plot_gqa.py
renders the GQA figures.
GEMM/All-Reduce reuse committed milestone figures/CSVs; GQA results are
generated fresh. Honest flags retained: PE_CPU dispatch cost is 0 in this
config, and the proposed two-composite softmax_merge decode is marked
designed-not-measured.
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>
_RwHazardTracker (pe_scheduler): a composite registers its write set (rw_handles) as in-flight; a new composite whose read (op operands) or write handles intersect an in-flight write set waits on those composites' done events before admission (ADR-0065 D6.3 / DDD 8). _dispatch_composite calls admit() before feeding and retires on the done event.
Legacy composites (rw_handles=()) never register and never block -> existing benches untouched (byte-equal). Strict FIFO is conservative (partial overlap waits for full drain); RW-aware reorder is deferred (A4).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
tiling.generate_plan_from_ops: scan flat ops for the GEMM (<=1); pre-GEMM KERNEL ops become single-shot prologue MATH stages, post-GEMM ops split by scope (K_TILE/OUTPUT_TILE epilogue + KERNEL post-loop). MATH-only composite reproduces the legacy math-head plan. Prologue/post-loop stages fold into the first/last tile so the feeder + completion counting are untouched (existing benches have neither -> byte-equal op_log).
PipelinePlan gains prologue_stages/epilogue_stages. pe_scheduler._generate_plan delegates to generate_plan_from_ops. DMA keeps the existing pinned signal (NOT a space flip); recipe scratch/primary-out handles are pinned=True so the head GEMM's auto-bound a (=P) is consumed in place.
ADR-0065 D4/D6.7/Test#5 amended (EN+KO): DMA decision from pinned, not space; prologue recipe ops TCM-only (head op exempt -- it may stream from HBM).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
New triton_emu/tl_recipes.py: RecipeDescriptor/EngineOp/PrimaryOutSpec + RECIPE_DESCRIPTORS['softmax_merge'] (8-step engine_seq). PE_SCHEDULER does not import it (ADR-0065 D5 boundary).
TLContext.composite(): add prologue=[...] + out=TensorHandle kwargs, a optional. _expand_prologue lowers a recipe into flat MATH OpSpecs (scope=KERNEL), allocates TCM scratch, derives the primary-out slot 'P' and auto-binds it into the head GEMM a (D6.6 conflict check); rw_handles=(m,l,O). decode-opt2 #2 lowers to 10 ops [rmax,max_elem,exp_diff,exp_diff,rsum,fma,mul_bcast,copy,gemm,add].
D6.7 (MATH operand TCM-only) scoped to prologue recipe ops only — the head op (gemm or math) keeps existing DMA-staged-from-HBM behavior. D6.1 (GEMM count <=1) on the whole composite. Host-side lowering only; PE_SCHEDULER position-scan is P3.
Also commit the ADR-0064 Rev2 promotion content that the prior commit's git mv dropped: Status Proposed->Accepted + D7 amended to hard-cap ValueError (no segmentation), EN+KO.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Comment-only change. Adds a kernbench-only limitation note at the
Tile-0 / bootstrap block in all four GQA attention kernels:
- persistent (m, l, O) must live outside tl.scratch_scope so
subsequent tiles' merges can read them;
- kernbench has no scratch-backed initializer (tl.zeros / tl.full
return addr=0 handles), so we can't seed (-inf, 0, 0) and rely
on tl.copy_to;
- therefore Tile 0 must compute the initial running state directly.
prefill_long adds a third bullet noting the ring-step ordering reason
(k=0 must send W before any k>0 can recv E, so the (t=0, k=0) step has
to run outside the scoped loop where the send is conditional on C > 1).
Each block also notes that a real Triton port collapses Tile 0 into a
unified loop (SSA tensors stay live across iterations).
No behavior change; tests unchanged.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
ADR-0064 Revision 2 second-pass review fixes:
- D7 added 1024-byte cap rationale: explicitly framed as a *safe
engineering limit* (intentionally above all known composites at ~322
bytes, finite descriptor capacity placeholder), not a measured HW
number. Default is the discipline; topology override for real HW.
- D7 segment ordering: strict FIFO is the *ordering source*; rw_handles
is dependency metadata, not the ordering primitive. Explicit note for
the future RW-aware reorder migration path.
- Tests rewritten around the formula (D1) instead of specific numbers:
- Test #1 → "formula preservation" (parametrized across composites,
asserts dispatch == FIXED + bytes × R, not anchor 43 ns).
- Test #2 → robust "opt3 > 2 × opt2" qualitative gate; ≈ 4.0× is
informative only, not the gate, so calibration changes don't break.
- Test #9 → qualitative (opt2 < opt3 at all R), not a numeric ratio.
- Anchor description in D3 stays as *informative* (shows the model
produces ≈ 43 ns at defaults), but is no longer a test gate.
ADR-0065 second-pass review:
- D5 step 5 aligned with D6.6 invariant: explicit conflict check before
auto-binding (previously written as if auto-bind always happens).
Mirrors D6.6 wording: error if both kernel explicit `a` and prologue
primary_out are present.
- Test #7 made robust: `opt3 > 2 × opt2` gate (was ≈ 4.0×).
Default-calibration model expectation (≈ 4.0×) recorded as informative
reference, not the gate.
DDD-0065 follow-on robust thresholds:
- §1.4 success criteria: ratio condition rewritten as `opt3 > 2 × opt2`
with model-expected ≈ 4.0× as informative.
- §9 P6 phase gate: `opt3 > 2 × opt2` (was 4.0× ± 15%).
- §10 dispatch ratio test: assert `opt3 > 2 × opt2`; record observed
ratio for performance tracking but do not gate on it.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
ADR-0065 review fixes:
- D3 semantics clarified: position determines *phase* (pre-loop / head /
in-accumulation / per-output-tile / post-loop), scope determines
*repetition*. KERNEL = "once per CompositeCmd invocation", not "once
per kernel launch" — removes ambiguity around softmax_merge's KERNEL
scope at different positions before/after GEMM.
- D4 MATH-TCM invariant: Phase 1 limits DMA auto-insertion to GEMM
operands and head outputs; MATH operand with space="hbm" raises
validation error at TLContext emit. Catches kernel misuse early; a
MATH-with-HBM path is future work (explicit tl.load + MATH composite).
- D6 #6 NEW: auto-bind conflict — explicit GEMM `a=` + prologue
primary_out simultaneously → validation error (prevents "which value
wins" ambiguity).
- D6 #7 NEW: MATH operand TCM-only restatement.
- Consequences/Negative: strict-FIFO conservatism note — safe but may
under-expose composite-level overlap when handle-sharing composites
could in principle run in parallel.
- Test Requirements #10, #11: auto-bind conflict + MATH-HBM validation
error tests.
ADR-0064 Revision 2 clarification:
- D2 counting rule: per-op summation, no deduplication. Same handle
appearing in multiple OpSpecs or in rw_handles is counted
independently — matches HW reality (each descriptor field is a
distinct address slot, rw_handles is separate metadata block).
Example walked through (mul_bcast with O in src_a + out + rw_handles).
DDD-0065 follow-on:
- §3.2: new test file test_tl_composite_validation.py covers auto-bind
conflict, MATH-TCM, GEMM-count invariants.
- §4.3: phase vs repetition split aligned with ADR-0065 D3.
- §6 lowering pipeline: step 5 adds auto-bind conflict check; step 8
NEW — operand space validation; step 9 (size cap); step 10 (emit).
- §8 strict-FIFO conservatism note: explains why next-tile #1 cross-tile
pipelining still works under FIFO (no rw overlap), and where the
conservatism would hurt (future handle-sharing recipes).
- §10 verification: invariant-guard tests (auto-bind conflict, MATH
TCM, per-op summation accuracy).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
ADR-0064 Revision 2 review fixes:
- D7 NEW: composite size cap (MAX_COMPOSITE_LOGICAL_BYTES default 1024
bytes). Oversized recipes deterministically segmented into N
CompositeCmds; each segment incurs its own dispatch cost. Models real
HW limits (descriptor queue entry, scheduler parser buffer, command
SRAM) and prevents the model from rewarding pathologically-large
fused composites.
- D2: type-aware extra-field byte counting (int/float=4, bool=1,
tuple/list=1+4N, str=1) — replaces uniform 4 bytes per extra.
- D3: recalibrated defaults to FIXED=40 cycles, R=0.0625 cycles/byte
(16 B/cycle — typical on-die descriptor queue width); anchor stays at
~43 ns for typical 1-OpSpec composite. Clarified anchor description:
DMA stages do not appear in logical_bytes (auto-inserted by
PE_SCHEDULER from operand.space per ADR-0065 D4).
- D4: removed clock_freq_ghz from pe_cost_model: override block;
conversion uses the PE node's existing clock_freq_ghz attr. Added
max_composite_logical_bytes knob.
- Context: emphasized command-count reduction (FIXED) as the primary
signal; byte term as secondary refinement.
- Open review: added large-composite scheduler-cost stress test.
- Test req: added composite-size-cap (#8) and R-sensitivity sweep (#9).
ADR-0065 + DDD-0065 follow-on updates:
- opt2 vs opt3 dispatch ratio updated 2.4× → ≈4.0× under new defaults
(FIXED-dominated, reflecting the corrected framing).
- Test req #9: decode opt2 composite fits within 1024-byte cap; no
segmentation needed for the GQA workload.
- DDD §6: TLContext lowering checks logical_bytes against cap (step 8).
- DDD §11: performance model recomputed with new defaults + sensitivity
table across R ∈ {0.25, 0.0625, 0.03125} confirming opt2 < opt3 holds.
- DDD §9 P6 gate: ratio band 2.4×±10% → 4.0×±15%; sensitivity sweep added.
Co-Authored-By: Claude Opus 4.7 (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>
Amends ADR-0026 to reflect the cube_start field added in e2fe331.
The production code shipped without an ADR update; this fills that
gap. Documentation-only change (no production code, no test code).
Changes (mirrored in both EN and KO):
- Status: Revision 5 → Revision 6
- D1: add ``cube_start: int = 0`` to the canonical DPPolicy dataclass
- D3: ``cube = policy.cube_start + cube_id`` in resolve_dp_policy
- D8: new section explaining purpose (disjoint cube sub-meshes for
GQA Llama-70B 8-KV-group headline), semantics, default-0 backward
compatibility, intra-device constraint, design rationale
(scalar vs 2D origin vs cube_ids list), and the kernel-side
cube_start subtraction needed to compensate for ADR-0022's
physical-cube-id ``program_id(axis=1)`` semantics.
tools/verify_adr_lang_pairs.py passes (EN/KO Status keyword and
title in sync).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per review:
- Placement unified to *contiguous* C×P blocks throughout (S0/S2.1/S2.2/S0.5);
round-robin demoted to the rejected alternative. Adds driver SP-enable
threshold fallback (smaller C / C=P=1 for short/early decode).
- Ring-vs-reduce cost model in S5.5: reduce ~ G*T_q*log(C*P) (O dominant;
m,l scalars), ring ~ 2*S (total K+V bytes a CUBE injects over C-1
rotations; recv_async pipelines so latency ~ max-step) -> ring wins when
T_q > 2S/(G*log(C*P)).
- opt3 'removes the bubble' -> 'hides (subject to scheduler+engine balance)'
everywhere; table 'hidden*' with footnote.
- 'rank' defined (SP participant = a PE in a CUBE, KV shard in its HBM->TCM).
- out-proj handoff contract (S0.5.4); S11 gate-type note (absolute latency
deferred to ADR-0064; structural + relative-to-baseline gates now).
- greenlet-as-contrast tightened to 'primitive-op (tl.dot) path' (S1, SB).
KO mirror synced; DDD gets the SP-enable threshold fallback. Docs only.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
DDD-0060 was rewritten to the two-kernel composite-hybrid design, so SB
item 1 ('DDD not yet synced') is stale. Update EN + KO mirror to record the
DDD as synced (ADR authoritative, DDD = impl how-to). Docs only.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Supersede the old greenlet-primitive/load_async/single-kernel DDD. Align to
ADR-0060's composite-hybrid + hierarchical CUBE-Group SP with two kernels:
decode+SP (head-replicated, contiguous C×P static shard, 2-level
reduce-to-root) and prefill+SP (1 Q head/CUBE, Ring KV, no reduce). Updated
file plan (lazy tl.load, scratch_scope, broadcast-opt; evolve
_attention_mesh_mlo_2d/_kv; DPPolicy.cube_start; llama70b_4sip.yaml),
placement (contiguous shared KV), phase plan (P1..P8), verification, perf
model, risks, glossary. Defers design rationale to ADR-0060 (now
authoritative); open items point to ADR-0060 §B. ADR-0064 cost model noted.
Docs only.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Regenerate the KO mirror to match the current EN: two SP kernels
(decode=reduce / prefill=ring), TL;DR full code for the 3 decode variants
+ prefill, 'Q replicated / M-fold' and '1 Q head per CUBE' terminology,
contiguous shared KV layout, opt2 tl.wait, and all SB items. Docs only.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The ex_composite #2 updates acc=(m,l,O) asynchronously on the scheduler and
the kernel never reads its output, so no auto-wait fires; acc is only final
after the last #2 in the serial chain. Add tl.wait() (drain all composites)
before hierarchical_reduce_and_store reads acc. opt1/opt3 don't need it
(their composite outputs are consumed in-iteration -> auto-wait). Docs only.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The ring rotates K/V over IPCQ, so they must be kernel-held TCM handles
(tl.load), not HBM tl.ref streamed inside the composite. Correct the
prefill kernel: load K pre-transposed [d, S/C] and pass the TCM-resident
Kc/Vc directly to the composite (drop the bogus tl.ref(Kc,(d,TILE_S)) and
TILE_S); recv shapes match ([d,S/C] for K, [S/C,d] for V). Comments note
why K/V live in TCM. Docs only.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Move the 3 decode CPU-pipelining variants (opt1 current / opt3 sw-pipe /
opt2 ex_composite) up into the TL;DR as full standalone kernels alongside
the full prefill ring kernel, with the comparison table. S5.6 is reduced to
a brief anchor (still referenced by S8/SB) pointing to the TL;DR code +
keeping the recommend/cost-model linkage. Docs only.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The prior commit updated S0/S10/S5.5 terminology and added S5.6 (3 decode
variants) but left the TL;DR with the old 'all G heads replicated' wording
and no pointer to S5.6. Sync the TL;DR: 'Q replicated (G heads stacked into
the GEMM M-dim, M-fold)', '1 Q head per CUBE', and a one-line pointer to the
3 CPU-pipelining variants in S5.6. Docs only.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
- Terminology: 'Q replicated' (all G query heads stacked into the GEMM M-dim;
M-fold explained) for decode; 'one Q head per CUBE' precise for prefill.
- New S5.6: three decode CPU-pipelining variants — opt1 current CompositeCmd
(has GEMM-engine bubble), opt3 software pipelining (issue next Q.Kt before
this tile's softmax; Sj in persistent double buffer; ships now, no new cmd),
opt2 ex_composite split into two (#1 = existing GEMM+scale reads K first;
#2 = softmax+P.V+accumulator merge, the only new flash-epilogue machinery,
gives DMA K-before-V priority). MATH engine already has max/sum/exp — the
new part is the stateful flash accumulator, not the ops.
- S2.1/SB: shared prefill/decode KV layout = contiguous CxP blocks (prefill
causal-skip needs contiguous; avoids prefill->decode reshard; short-context
under-use caveat). S8 item 4 sizing note for the two-composite split.
Prefill note: opt2/opt3 give little for prefill (causal if can't enter a
composite; recv_async already overlaps). Docs only; KO mirror deferred.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Decode-SP and prefill-SP are structurally different and cannot share one
kernel (principle: move the smaller thing):
- Decode: O=[G,d] tiny, KV cache big -> keep KV statically sharded/resident,
G heads replicated (M-fold), move only (m,l,O) via the 2-level reduce (S4).
- Prefill: O=[S,d] big -> shard heads (1 query head per CUBE, C=G),
rotate KV (Ring KV, S5.5), no (m,l,O) reduce; each CUBE writes its own head.
Rewrites TL;DR (two kernels), S0 (head map differs by case), S0.5.4 (output
head distribution differs -> downstream out-proj impact), S4 (scoped to
decode; S4.1 = intra-CUBE KV-split + PE reduce, the only way decode uses P
PEs), S5.1 (decode skeleton), S5.5 (head-parallel Ring KV), S9/S10/S11, and
adds SB items (two head mappings, output asymmetry, prefill within-CUBE PE,
C=G coupling, reconcile with _attention_mesh_mlo_2d). KO mirror deferred
until the design stabilizes (adr-proposed is mirror-exempt). Docs only.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Placement pivot (user-approved): a CUBE Group = C CUBEs within one SIP that
jointly own one KV head + its G=8 query heads. KV sequence sharded 2 levels
(Level-1 inter-CUBE over C, Level-2 intra-CUBE PE over P=8 = C*P ranks); G
folded into matmul M. C is a knob (8/4; C=1 = single-CUBE). Reverses the old
'a query head never spans CUBEs' non-goal (held only because the baseline
was H_kv=1). device=SIP; the for-kv loop is gone (head picked by CUBE coord).
Reduction is a 2-level reduce-to-root (not all-reduce): Level-2 PE tree ->
Level-1 center-root CUBE-mesh reduce, adapting lrab_hierarchical_allreduce's
inter-CUBE pattern as reduce-only + log-sum-exp. Data-driven (send on local
P.V completion, no global barrier) + level-pipelined; per-level topology
configurable (tree for decode, ring for long prefill).
Rewrites SS0/2/4/5/0.5/8-11, pseudocode, and adds SSB items (4-SIP config,
mesh partition, C knob, invariant-reversal check, index-math test). KO mirror
updated. Topology grounded in topology.yaml (4x4 CUBE mesh/SIP, 8 PE/cube).
Docs only; no production code changed.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
mukesh
ywkang