Three coupled fixes so enable_data=False now reports byte-identical
kernel-body sim latency (max(t_end) - min(t_start) over op_log) as
enable_data=True, and skips setup-write sim events that were pure
wall-clock overhead.
Before this change, at Case 4 (Cube-SP x PE-SP) decode, S_kv=8K:
enable_data=False -> latency_ns = 0 (op_log empty)
enable_data=True -> latency_ns = 30.646 us (correct)
After:
enable_data=False -> latency_ns = 30.646 us (byte-equal)
enable_data=True -> latency_ns = 30.646 us (unchanged)
engine.py: always create OpLogger. Previously OpLogger was gated on
enable_data=True together with MemoryStore, so op_log stayed empty when
data mode was off. Decouple: MemoryStore is gated (Phase 2 DataExecutor
still needs it) but OpLogger runs unconditionally, receiving
memory_store=None when data mode is off. OpLogger already guards its
arr.copy() snapshot paths on `if self._memory_store is not None`.
pe_cpu.py: always use the ADR-0020 greenlet execution path. The
if store is not None: _execute_greenlet(); else: _execute_legacy branch
gated the *execution model* on data-mode presence. The legacy
command-list path predates ADR-0020 and doesn't route IpcqSendCmd
through PE_IPCQ (it goes to PE_SCHEDULER instead), so all 189 Case-4
ipcq_copy fabric transfers were silently no-op'd in that mode. Forcing
greenlet gives identical sim behavior in both modes. KernelRunner
already guards its store reads on `self._store is not None`.
context.py: gate setup MemoryWriteMsg on memory_store presence. Under
enable_data=False there is no MemoryStore to populate and no Phase 2
replay, so the per-shard sim events for Q/K/V deploy are pure wall-clock
overhead with no effect on reported kernel latency (Yangwook's max-min
formula excludes pre-kernel ops). Handle count for Case 4 at S_kv=8K
drops 229 -> 37 under enable_data=False.
Verified:
- tests/attention/test_milestone_gqa_decode_long_ctx_4cases.py: 18/18 pass
- Parity probe: enable_data=False/True both report 30.646 us kernel sim
time, 1649 op_log records, 189 ipcq_copy events
Known follow-up: pe_cpu._execute_legacy is now dead code but left in
place; separate cleanup once we confirm no downstream caller.
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>
- KernelLaunchMsg gains target_start_ns: IO_CPU stamps a global barrier
(max path latency across every target PE), M_CPU passes it through,
PE_CPU yields until it before recording pe_exec_start. Every PE in a
launch begins kernel execution at the same env.now regardless of its
dispatch path length — eliminates per-PE dispatch-offset artifact in
cross-PE and cross-cube latency measurements.
- PE_DMA._handle_ipcq_inbound now pays Transaction.drain_ns at the top,
matching the terminal-drain behavior of ComponentBase._forward_txn for
every non-IPCQ Transaction. SRC-side tl.send stays fire-and-forget
(sender doesn't yield on sub_done); tl.recv now blocks until bytes
have actually drained into its inbox.
- ComponentContext: new compute_path_latency_ns helper + node_overhead_ns
field populated by GraphEngine.
- tests/test_kernel_launch_sync.py: asserts all PEs in one launch
produce identical pe_exec_ns for a no-op kernel (zero spread).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
tl.program_id(axis=0) returns local PE id within cube,
tl.program_id(axis=1) returns cube id. Enables cube-aware
sharding in benchmark kernels.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
mukesh
ywkang