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>
Add 5 of the 6 figure renderers ADR-0057 D3 sub-cycle 4c specifies:
- gqa_op_log_{panel}.png × 4 — per-panel bar chart of the 5 op_log
counts (gemm, ipcq_send, ipcq_recv, dma_read, dma_write).
- gqa_comparison.png — cross-panel grouped bars over the same 5 series.
Sixth figure (gqa_scaling.png) depends on sub-cycle 4b's Q/cube ∈
{1, 2, 4} sweep on multi_user_* panels and is deferred until that
data exists; emit_all_gqa_plots returns just the 5 in-scope paths.
Add MILESTONE_FAST=1 mode to run(): skip the panel sweep, reuse the
committed sweep.json, render figures only. Validation mode unchanged.
The runtime errors clearly when neither env var is set, listing the
two supported modes.
Renderers live in the bench module (the milestone-1h-gemm pattern);
tests/gqa/_gqa_plot_helpers.py re-exports them for figure tests.
Tests: tests/gqa/test_plot_gqa_figures.py — 7 tests, all green:
- 4 parametrized per-panel emit assertions
- 1 comparison emit assertion
- 1 emit_all returns exactly 5 PNG paths
- 1 default out_dir matches the bench _OUTPUT_DIR
Commits the 5 PNG baselines under the bench output dir alongside
sweep.json, mirroring milestone-1h-gemm's committed-figures pattern.
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>
Phase 1 cannot snapshot math-output sources at outbound send time
because math executes only in Phase 2 — so token.data stays None and
PE_DMA inbound can't write the recv slot. For own-sends this is harmless
(Phase 2 replay reads the stable scratch addr after math runs). For
forwarded sends in mesh kernels (ADR-0059), src_addr is a recv slot
that gets wrapped by later inbounds before this read's Phase 2 turn,
yielding a shape mismatch on the fallback MemoryStore.read.
Fix: DataExecutor maintains a per-slot, time-ordered, shape-keyed
history. Every ipcq_copy write appends (t_write, value) to the slot's
history; _resolve_read falls back to the most recent shape-matching
entry with t_write <= the consuming op's t_start. Applied uniformly
to _execute_memory, _execute_gemm, and _execute_math.
Secondary: OpLogger.record_end for math ops now prefers
TensorHandle.data carried by the input handle over a MemoryStore
re-read, closing the smaller record-end race covered by the new
test_op_log_input_snapshot_race.py unit tests.
Tests: 4 new race tests + 6 existing op_log + mesh decode diag +
mesh kv/mlo spec — all green. Full repo sweep: 760 passed (3
pre-existing failures unrelated: bench-registry list drift +
Windows Tkinter env).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per request, the milestone bench output is now tracked in git instead of
gitignored, so the figures/results are viewable on the remote:
- src/kernbench/benches/1H_milestone_output/gemm/ (3 PNGs + gemm_sweep.json)
- src/kernbench/benches/1H_milestone_output/ccl/ (3 per-topology PNGs,
buffer-kind PNG+CSV, FSIM comparison PNG, topology.png, summary.csv)
Drop the .gitignore rule; update ADR-0054 D3 + Negative (EN+KO) to say the
output is committed (regenerable by rerunning the bench). Artifacts produced
by full bench runs (milestone-1h-gemm non-FAST, milestone-1h-ccl).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Move the GEMM + allreduce sweep/render logic out of scripts/ and tests/
into two self-contained eval benches so a user can regenerate every
result + figure with one command:
kernbench run --bench milestone-1h-gemm (MILESTONE_FAST=1 reuses JSON)
kernbench run --bench milestone-1h-ccl
- benches/milestone_1h_{gemm,ccl}.py: single home for each domain; the
run(torch) entry drives the sweeps and writes figures into
benches/1H_milestone_output/{gemm,ccl}/ (gitignored), then submits a
sentinel tensor to satisfy the run_bench contract.
- tests/gemm + tests/sccl helpers and scripts/gemm_sweep.py become thin
re-export/wrapper shims over the benches (single source preserved); the
pytest-only param builders + _run_distributed wrapper stay in the shim.
- eval-bench pattern: a bench may drive many configs + build its own
per-config engines (extends ADR-0045 D5; reverses ADR-0044 D1/D2).
ADR-0054 (EN+KO) records the design; ADR-0043/0044/0045 + CLAUDE.md CLI
Semantics amended; ADR INDEX regenerated. Verified: milestone benches run
clean (ok=True, all artifacts), full suite 67 passed, lang-pairs OK.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds a section-based table of contents for the 46-ADR corpus, mirroring
the /report skill's classification (Design Principles / High-level
Architecture / Detailed Architecture by component / Implementation
Decisions by topic). Generated for both docs/adr/ (EN titles) and
docs/adr-ko/ (KO titles) from one tool.
tools/generate_adr_index.py:
- Single CLASSIFICATION dict per ADR — add an entry when introducing a
new ADR; the script fails loud if any file is missing from the table.
- DETAILED_COMPONENTS lists each builtin component and the ADR(s) that
cover it (ADR-0014 appears under six PE engines; ADR-0023 under
pe_dma + pe_ipcq).
- Accepts both ":" and "—" title separators (matching ADR-0033's
existing format).
- --check mode for CI: exits 1 if INDEX.md is stale.
Also includes the docs/report/architecture-2026-1H.md generated by the
prior /report write (the public-facing architecture document; 836 lines,
76 source-attribution comments).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Documents four cross-cutting surfaces one layer deeper than the prior
G4 batch:
- 0050 par-ccl-algorithm-module-contract: how to author a new CCL
algorithm in src/kernbench/ccl/algorithms/. Pairs with ADR-0045's
bench-module contract. Pins the four required public symbols
(kernel, kernel_args, TOPO_NAME_TO_KIND constants, kernel alias),
the 9 + tl standardized kernel signature, the kernel_args tuple
format, sip_topo_kind dispatch, and the ccl.yaml entry workflow.
- 0051 lat-routing-helper-api: every public method of AddressResolver
(resolve, find_m_cpu, find_pcie_ep, find_io_cpu, find_all_pcie_eps)
and PathRouter (find_path, find_path_with_distance,
find_mcpu_dma_path, find_memory_path, find_node_path + 2 shims).
Pins the four adjacency graphs (_adj_all / _adj / _adj_mcpu_dma /
_adj_local) and the edge-kind exclusion sets they use, plus the
single-owner naming convention.
- 0052 dev-oplog-memory-store-schemas: OpRecord's 7 fields, the
per-op_name params matrix (dma_read, dma_write, gemm_*, math, math
reduction, composite_gemm, ipcq_copy, unknown), snapshot timing
rules (math = all inputs, dma_write = HBM-only — ADR-0027 race
avoidance), TileToken stage_type capture, and MemoryStore's
(space, addr) two-level dict with reference-store semantics.
- 0053 dev-topology-builder-algorithms: the 6-stage compile pipeline,
cube_mesh.yaml's source_hash cache and its 5 input fields, the
cube NoC auto-layout algorithm (row/col placement, HBM exclusion
zone, PE/M_CPU/SRAM attachment via nearest-router, UCIe N/S/E/W
distribution), the node naming convention (single-owner with
router.py), the edge-kind catalog, the 4 view projections, and a
table of spec-field changes vs mesh regeneration.
Bilingual pair verifier passes for all four EN/KO pairs.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Documents four cross-cutting surfaces that previously had no ADR backing,
each surfaced as a G4 candidate by /report:
- 0046 prog-tl-context-contract: the kernel-side tl.* API. Enumerates
all primitives (ref/load/store/dot/composite/math/reduction/IPCQ/...),
the two execution modes (command-list vs greenlet runner), scratch
allocator semantics, dispatch-overhead model, and the kernel registry.
- 0047 par-ahbm-ccl-backend: torch.distributed.init_process_group
(backend="ahbm") install path. world_size priority (algorithm >
defaults > topology), the 4-step init sequence (load ccl.yaml, import
algorithm module, derive world_size, install SFR + IPCQ), greenlet-
local rank registry, all_reduce dispatch via _defer_wait, barrier
no-op rationale, and the explicit list of unsupported dist.* APIs.
- 0048 mem-allocator-algorithms: VirtualAllocator + PEMemAllocator
free-list semantics. Offset-keyed first-fit with coalescing, the
no-validation trust model for free(), HBM/TCM channel separation,
page-aligned VA allocation, the page_size dual-default
(VirtualAllocator 2 MiB / _ensure_allocators 4 KiB fallback), and
one-allocator-per-sub-unit rule.
- 0049 ver-probe-subcommand: kernbench probe traffic-pattern catalog.
H2D / D2H / PE DMA categories with their exact cube-index choices,
the 32 KiB reference size, the 5-point utilization sweep, the
formula vs actual column meanings, automatic invariant checks
(monotonicity, D2H >= H2D, best < worst), per-case GraphEngine
isolation, and the human-readable (not machine-parsable) output
contract.
Bilingual pair verifier passes for all four EN/KO pairs.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Documents src/kernbench/benches/: how @bench registration + audit work,
how the CLI dispatches via run_bench/RuntimeContext, and the contract a
new bench module must satisfy.
Nine decisions (D1-D9) cover:
- @bench name/description rules and duplicate detection
- Module-file convention (_-prefixed helpers vs bench modules)
- def run(torch) signature; torch = RuntimeContext
- Minimum-one-submit rule (else NO_REQUESTS)
- Single-device convention + multi-SIP CCL exception (ADR-0024/0027)
- resolve() name/index decision tree; indices are not a stable API
- Exact RuntimeContext surface exposed to benches
- Env-var parameterization (matmul_composite / gemm_sweep.py pattern)
Four alternatives rejected with documented reasons (manifest YAML,
decorator entry= arg, @multi_device_bench split, stable indices).
Verifier (tools/verify_adr_lang_pairs.py) passes for EN/KO pair.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Document the allreduce + GEMM evaluation harnesses and bring the affected
allreduce ADRs in line with the refactored code.
New (Accepted, EN + KO):
- ADR-0043 — allreduce evaluation harness (tests/sccl/): distributed-driven
correctness, latency/buffer-kind sweeps, sessionfinish plot aggregators,
topology + FSIM-comparison figures. Verified against the implementation.
- ADR-0044 — GEMM evaluation harness (scripts/gemm_sweep.py + tests/gemm/):
heavy-script data gen vs. fast test-rendered figures, slow regenerator,
the 3-figure set. Records two limitations as open questions: the
theoretical-model constants are inherited (not yet traced to ADR-0033/
0014), and the *_measured figure is a naming misnomer.
Updated (EN + KO):
- ADR-0024 — add D5: SIP grid w/h resolution (explicit sips.w/h, square
fallback, fail-loud), documenting the AhbmCCLBackend fix.
- ADR-0032 — D4/D5/Non-goals reconciled: rectangular SIP grids (e.g. 6 SIPs
as 3x2) are supported via explicit w/h; the square requirement now
applies only to the fallback. Affected-files repointed to tests/sccl/.
Verification: ADR-0023 and ADR-0042 confirmed still matching the code (no
change). verify_adr_lang_pairs.py passes (EN/KO Status blocks byte-equal).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Mirror the sccl pattern for GEMM figures: a tests/gemm/ package renders the
GEMM bar charts as PNGs from the committed docs/diagrams/gemm_sweep.json, so
the figures are fast test artifacts (run by default) while the heavy sim sweep
stays a manual script (scripts/gemm_sweep.py, kept) wrapped by a slow
regenerator test.
tests/gemm/:
- _gemm_plot_helpers.py: matplotlib renderers (series logic mirrors the
GEMM _render_* functions in scripts/build_overview_slides.py).
- test_plot_gemm_stage_breakdown.py: gemm_stage_breakdown.png (load_ref).
- test_plot_gemm_mac_utilization.py: gemm_mac_utilization_measured.png +
gemm_mac_utilization_theoretical_vs_measured.png (load_ref).
- test_gemm_sweep.py: @pytest.mark.slow regenerator (runs scripts/gemm_sweep.py).
Chart set trimmed to three (stage breakdown, MAC util, theoretical-vs-measured);
"formula" relabeled to "theoretical" throughout the comparison chart.
Known follow-ups (not blocking):
- gemm_mac_utilization_measured.png currently plots the theoretical ideal-
pipeline model, not simulator-measured data; the name is a misnomer pending
a decision to repoint its content or retitle.
- The theoretical-model constants (HBM 256 GB/s, T_stage 16 ns, 3 stages) are
inherited verbatim from build_overview_slides.py and not yet verified against
ADR-0033 / ADR-0014 / topology.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Convert the multidevice allreduce correctness + latency/buffer-kind sweeps
to run through the real PyTorch-distributed path
(init_process_group(backend="ahbm") -> mp.spawn -> dist.all_reduce) instead
of direct ctx.launch, and reorganize the CCL/allreduce tests into a
tests/sccl/ package split one test per file.
Production change (required for the distributed path on non-square SIP grids):
- AhbmCCLBackend now reads explicit system.sips.w/h from the spec, with a
square-only sqrt fallback that raises on ambiguity, instead of silently
guessing round(sqrt(count)). This fixes the 2x3 / 3x2 torus + mesh cases,
which previously resolved to a wrong 2x2 grid. Mirrors the test helper's
_sip_topo_dims precedence (explicit w/h > square fallback > raise).
Test reorganization (tests/sccl/):
- _allreduce_helpers.py: shared plumbing (distributed driver, config writers,
direct-launch run_allreduce parity reference, sweep/buffer-kind constants,
plot aggregators, topology-diagram + FSIM-comparison emitters).
- test_allreduce_ring_torus_mesh.py: correctness across ring/torus/mesh.
- test_distributed_default_topology.py: full distributed path on topology.yaml.
- test_plot_latency_sweep.py / test_plot_buffer_kind_sweep.py: sweep rows.
- test_plot_topology_diagram.py / test_plot_comparison_fsim.py: plot emitters.
- test_intercube_root_center.py: moved in (ADR-0032 center-root latency guard).
Also:
- Move the FSIM comparison plot generator out of scripts/ into the sccl suite.
- Delete superseded test files (test_allreduce_multidevice,
test_distributed_lrab_hierarchical_allreduce, test_allreduce_buffer_kind_sweep)
and repoint conftest aggregators + the ipcq buffer-kind importers.
- Regenerate the allreduce_latency_plots derived artifacts from the full sweep.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Rename the intercube all-reduce identity to lrab_hierarchical_allreduce
(module, config key, distributed test) so the name reflects both levels
it implements: LRAB intra-SIP (local reduce to center root + broadcast)
and the hierarchical inter-SIP topology exchange (ring/torus/mesh).
ADR-0032 slug kept as the stable decision id; pure rename, no logic change.
Also in this batch:
- ADR-0032 (EN+KO): document the shipped center-root bidirectional reduce
(doc was stale corner-root); annotate ccl.yaml root_cube as a placeholder.
- Rename allreduce + pe2pe latency plots to descriptive, title-matching
filenames and retitle the in-plot headings; drop overview/overview_log.
- Point the PPTX image refs at the new plot names.
Doc + derived-artifact + rename only; no simulation behavior changed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replace BenchResult.engine: object | None with GraphEngine | None via
TYPE_CHECKING import (avoids circular import at runtime). Cast the
topology argument to TopologyGraph at the GraphEngine call site for
the duck-typed engine_factory. Fixes Pylance reportAttributeAccessIssue
warnings on op_log and topology arg. Type annotations only; no runtime
behavior change.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Fill component-model coverage gaps surfaced by /report's G4 analysis.
Each ADR documents the component's First action, latency model, and
honest notes on dormant code or implementation asymmetries discovered
during re-evaluation against current code.
- 0038 pcie_ep: thin protocol-overhead model; ComponentBase forwarding
worker as-is; named-node contract for router helpers
- 0039 pe_mmu: component + utility dual role; sub-page region stopgap;
D2.1 flags pipeline path missing mmu.overhead_ns timeout (asymmetric
with non-pipeline; not visible at default tlb_overhead_ns=0)
- 0040 pe_tcm: dual-channel BW serialization (read/write Resource cap=1);
TcmRequest schema owned by TCM; timing-only (no data store)
- 0041 sram: terminal scratchpad model + ResponseMsg on reverse path;
D1.1 flags _worker override as currently dormant (no Transaction
actually targets the SRAM node today)
- 0042 tiling: pure plan-generator module, not a component; corrects
the G4 misclassification; pins GEMM/Math stage sequences and
epilogue scope contract
Also: /report skill G3 refinement — only flag older->newer asymmetric
cross-references; newer->older (e.g., 0034-0037 citing infrastructure
ADRs) are expected one-way and no longer reported.
Bilingual pair verifier (tools/verify_adr_lang_pairs.py) passes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Move benches/ -> src/kernbench/benches/ and src/kernbench/cli/probe.py ->
src/kernbench/probes/probe.py. Each bench self-registers via
@bench(name=..., description=...); kernbench list enumerates benches
with auto-assigned indices, --bench accepts kebab-case name or numeric
index. Audit at package-import time fails if any non-underscore module
forgets the decorator. ADR-0010 (EN + KO) updated to reflect the new
resolver path, list subcommand, and probes package separation.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Follow-up to the bilingual-structure commit: docs/adr-ko/ now holds
only Korean versions (24 files translated from English placeholders),
ADR-0013 slug uses kebab-case in both folders, and the verify tool
allows translated parenthetical commentary in the Status block.
- Translate 24 English files in docs/adr-ko/ to Korean. The previous
bilingual-structure commit had left these as English copies because
their source content was already English; this commit fulfills the
policy that docs/adr-ko/ contains only Korean.
- Rename ADR-0013 in both adr/ and adr-ko/ from
ver-verification_strategy.md to ver-verification-strategy.md
(kebab-case consistency with other ADRs).
- CLAUDE.md (ADR Translation Discipline): clarify that only the
Status lifecycle keyword (Accepted / Proposed / Stub / Draft /
Superseded by ADR-NNNN / Merged into ADR-NNNN) must match across
EN and KO; parenthetical commentary and trailing list items may be
translated.
- tools/verify_adr_lang_pairs.py: replace byte-equal Status check
with normalize_status_keyword() which strips parenthetical
commentary and takes only the first non-empty line.
- tests/test_verify_adr_lang_pairs.py: update existing test names,
add coverage for translated parenthetical, translated trailing
list, and Superseded-by-NNNN keyword equality.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Establish English as the canonical ADR language with Korean translations
held in a parallel docs/adr-ko/ tree as derived artifacts (1:1 mirror).
Promotion from adr-proposed/ to adr/ now writes English to adr/ and the
Korean to adr-ko/; bidirectional sync rule documented in CLAUDE.md.
- Migrate 30 ADRs in docs/adr/: 28 Korean-only translated to English,
2 bilingual pairs (ADR-0020, ADR-0023) consolidated (.en.md suffix
dropped). ADR-0023 EN regenerated against KO source which had newer
HW Realization Notes (D16-D23) section.
- docs/adr-history/ left frozen by design (transitional state).
- CLAUDE.md (Part 2): update ADR Lifecycle for 4-folder layout, mark
docs/adr-ko/ as a Derived Artifact, add ADR Translation Discipline
section covering bidirectional sync, conflict resolution (EN wins),
and proposed-language freedom.
- tools/verify_adr_lang_pairs.py: new verification tool checking pair
completeness, filename mirroring, ADR-ID match, Status byte-equality.
Pre-commit hook intentionally not added; run on demand or in CI.
- tests/test_verify_adr_lang_pairs.py: 11 cases including CRLF/LF
normalization, em-dash title separator, underscore-slug edge case.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
- CLAUDE.md: add ADR Lifecycle subsection (superseded → docs/history/,
immutable numbering, no renumber)
- ADR-0011: merge ADR-0018 content as "Address Model: LA" section
alongside PA / VA; status notes VA model is currently implemented
- ADR-0018 / 0029 / 0031: moved to docs/history/ with status updates
(0018 merged into 0011, 0029 superseded by 0032, 0031 absorbed
into 0001 rev 2)
- ADR-0019: rewrite Context as PE-HBM connectivity decision
(self-contained, no LA model framing)
- ADR-0019/0020/0021/0023/0025/0027: Status Proposed → Accepted
(code verified) and prune Implementation Notes / Affected files /
Test strategy / "현재 상태" sub-sections describing pre-impl state
- ADR-0024/0026: same migration-flavor cleanup; 0026 also drops D6
Migration and D8 docs-update sub-decisions
- ADR-0030: status simplified (blocker ADR-0031 now superseded)
- SPEC.md: R10 + §0.2 reflect PA / VA / LA model names
- ADR-0008/0012/0013: refresh ADR-0011 subtitle in Links
21 files changed, 553 insertions(+), 1290 deletions(-).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
- Reorganize rules into reusable general behavior vs kernbench-specific
foundations + rules
- Add Surfacing Choices, Coding Style (Simplicity First, Surgical Changes),
Mental Model, Common Failure Modes
- Clarify Phase 1 forbidden vs permitted-for-discussion (pseudocode,
sketches allowed; final ready-to-apply diffs are Phase 2 only)
- Tighten dead-code handling: mention + options before deletion
- Drop redundant "SPEC.md and ADRs are the final authority" from
Enforcement Defaults (already in Authority & Scope)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
mukesh
ywkang
Yangwook Kang