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>
Adds the 4-SIP topology referenced by tests/attention/test_attention_8kv_groups_diag.py
(landed in 39fc2e9). Identical to the repo-root topology.yaml except
``system.sips.count: 2 -> 4``, matching the GQA Llama-70B sharding
study's 1 Q-head-per-cube baseline (64 cubes = 8 KV-groups x 8
cubes/group across 4 SIPs).
Opt-in only — only the diag harness and milestone-gqa-llama70b use
this file via env var. The repo-root topology.yaml is unchanged and
continues to drive milestone-1h-ccl / -gemm at their 2-SIP scale.
Without this commit the diag tests pytest.skip on a fresh checkout
("4-SIP topology missing"). With it, step 1 / step 2 / step 3 all run
and pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <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>
Adds an optional ``cube_start: int = 0`` field to ``DPPolicy``. In
``resolve_dp_policy`` the generated ``ShardSpec.cube`` is now
``policy.cube_start + cube_id`` instead of just ``cube_id``. With the
default value (0) every existing call site resolves to identical
``ShardSpec.cube`` values (cubes 0..num_cubes-1).
Why this is needed: the GQA Llama-70B 8-KV-group headline target lays
two 2x4 KV-groups per SIP — the first on cubes 0..7 (rows 0..1), the
second on cubes 8..15 (rows 2..3). Without ``cube_start``, ``DPPolicy``
can only address the first 8 row-major cubes, so a second launch on
the same SIP overlaps the first. ``cube_start=8`` selects the second
2x4 sub-mesh directly.
Design choice: scalar ``cube_start`` (vs a 2D ``cube_mesh_origin`` or
arbitrary ``cube_ids`` list) was picked because (a) the 2D mesh-mlo
kernel already assumes row-major contiguous cubes, (b) it pairs
naturally with ``num_cubes`` (range = ``[start, start+count)``), and
(c) zero migration churn — every existing call site stays unchanged.
Scope: 5 production LOC (one field + one expression in resolve). No
kernel changes here; kernels that consume ``program_id(axis=1)`` for
ring arithmetic need their own follow-up (see ADR-0022 contract).
Tests: 7 new unit tests in ``test_dppolicy_cube_start.py`` covering
default behavior preservation (cube_start=0), shifted ranges
(cube_start=8 → cubes 8..15), full-SIP CCL pattern, replicate cube
policy, shard-count preservation, and uniqueness. ADR-0026 regression
suite (12 tests) stays green.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The single_user_prefill/single_user_decode panels of milestone-gqa-llama70b
(landed in e748a62) import configure_sfr_intracube_pe_ring from
kernbench.ccl.sfr_config, but the function definition was left in the
local working tree and never made it into that commit. Because the
bench registry eager-imports every bench module at CLI startup, the
missing symbol breaks `kernbench` entirely and cascades into 7 test
failures + 18 collection errors for anyone on a fresh checkout.
Installs an 8-PE logical ring inside every cube on every SIP (E/W
direction labels, no intercube or inter-SIP edges). Mutually exclusive
with configure_sfr_intercube_multisip on the same engine — the bench
picks one per panel-run (ADR-0058 D2).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Both mesh kernels and the milestone-gqa bench described single_user vs
multi_user in ways that were architecturally misleading. Source of truth
is the GQA Llama-70B sharding study at
llm_paper_review/notes/GQA_MHA_sharding/scripts/_gen_llama70b_1M_4cases.py:
Single User (B=1) — TL prefill, BL decode:
KV split @ PEs intra-cube. Each cube does its own 8-PE ring,
independent of other cubes. Headline = 64 cubes serving one user.
Multi User (B=8) — TR prefill, BR decode:
KV split @ cubes inter-cube. 8 cubes/KV-group form the ring; inside
each cube, "8 PEs each handle 1 different user → Batch on batch."
Updates the rank_axis docstring on both mesh kernels and the bench's
module docstring to make this accurate, including the v1 simplification
of the multi_user kernel (gating pe_id != 0 collapses B=8 → B=1; the
per-cube batch parallelism is deferred to sub-cycle 4c headline).
No behavior change; existing tests unchanged.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
- ADR-0060: GEMMs (Q.Kt, P.V) via existing tl.composite (scheduler-managed
tiling + K/V DMA streaming); softmax merge + IPCQ tree reduction stay in
kernel. Front TL;DR pseudocode of the final composite kernel; new section
B lists open design items (DDD sync, K pre-transpose, dma_read lever,
kernel-vs-scheduler tiling, ring path).
- ADR-0062: redefined from a new load_async op to global lazy tl.load
(non-blocking + auto-wait on first use; API unchanged; goldens regenerate).
- ADR-0064 (new): per-op-type CPU issue cost model (composite ~40ns >>
primitive) so the hybrid's CPU-saturation win becomes measurable
(currently dispatch_cycles=0 hides it). Cost-model impl deferred.
- KO mirrors for ADR-0060/0062/0064 (-ko suffix, adr-proposed).
Rationale: non-blocking CompositeCmd offloads tiling to PE_SCHEDULER,
decoupling CPU issue-rate from execution so the CPU can saturate the
engines; the prior 'composite = no latency benefit' claim was an artifact
of dispatch_cycles=0. Docs only; no production code changed.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Verified against sim_engine data path (memory_store, data_executor):
- GQA reuse does NOT need a broadcast op. The baseline's h_q==h_kv limit
is its head-packing reshape hack, not a missing primitive. Correct fix:
per-KV-head loop with G folded into matmul M dim (byte-conserving
reshape) — runs today, timing correct (m=G*T_q), data mode runs.
- ADR-0061 broadcast demoted from 'the blocker' to optional convenience.
- Surfaced tl.trans = reshape-not-transpose (memory_store reshapes;
data_executor np.matmul on reshaped operands) -> numeric parity is
bounded; verification is structural/timing/determinism-first (matches
SPEC perf-model purpose). Optional tl.transpose deferred.
- Reordered DDD phase plan (P1 GQA needs no new feature; P3 scratch_scope
is the key scale feature); added open decisions 10.10 (transpose) and
10.11 (GQA-via-M-fold finding).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Reframes the proposed GQA FlashAttention design onto kernbench's actual
execution model (greenlet tl API + IPCQ), replacing the composite-centric
mechanism that does not match the simulator:
- Records relationship to existing baseline kernels (_attention_mesh_kv/mlo,
milestone-gqa-llama70b) and their 3 deliberate limitations.
- Mechanism is greenlet tl (per-op latency; no fusion benefit), not
composites; running (m,l,O) is Python handles; reduction is tl.send/recv.
- Tree reduction (log N) replaces baseline all-to-all fan-out (N-1).
- Pseudocode rewritten in real tl.* signatures; depends on ADR-0061/62/63.
- Rejects composite-IPCQ-push + composite-carried-state + flash-composite
with documented efficient alternatives.
- Adds verification plan.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Proposed prerequisites surfaced while evaluating the GQA fused-attention
ADR against the actual kernbench tl/sim_engine implementation:
- ADR-0061 tl.broadcast: data-faithful GQA head reuse (fixes the
MemoryStore nbytes check that forces h_q==h_kv==1 today).
- ADR-0062 tl.load_async: non-blocking HBM tile load for KV prefetch
(KV-load-bound decode/long-context overlap).
- ADR-0063 tl.scratch_scope: per-tile scratch recycling (removes the
1 MiB bump-allocator ceiling that caps context at S=16).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
ywkang
mukesh
Yangwook Kang