merge+ledger: fused preamble +0.8% (cumulative +1.9% with GDN cluster), stacking validated; prefill levers additive

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
Claude-Session: https://claude.ai/code/session_01UJyFKcK62CcR3imhgbiBnW

Author: mudler

.agents/parity-ledger.md

@@ -246,3 +246,5 @@ Columns:
 | 2026-07-08 (PREFILL gap-scan vs vLLM+deps — 44-agent workflow, 20 CONFIRMED gaps; the 0.77x prefill gap is FUSION + bf16-intermediates, NOT cudagraph/fp8) | Per user: re-scan gaps vs vLLM + deps (flashinfer/cutlass/FLA) for the confirmed PREFILL frontier. 10 subsystems × find→adversarial-verify. Full results: .agents/prefill-gap-scan-2026-07-08.json. | (44-agent scan+verify) | **20 CONFIRMED gaps. NEGATIVES (ruled out): prefill-cudagraph (vLLM runs large-prefill EAGER too — graphing is the WRONG lever) + fp8 GEMM (0/5, at-parity). DOMINANT THEME (6+ findings): vLLM fuses the full-attn PREAMBLE (split+q/k-RMSNorm+partial-RoPE+gate) into ONE bf16 kernel reading a precomputed cos_sin_cache; we run 4-5 SEPARATE f32 kernels that recompute RoPE transcendentals in DOUBLE + round-trip f32 intermediates through HBM (tasks #14+#15). RANKED: LOW-effort quick wins — (a) silu+quant fusion is DEFAULT-OFF, flip FuseSiluQuantEnabled ON (2-4%, ZERO code, fused kernel exists); (b) v_new full-buffer memset ~97% dead → zero only slack rows (cuda_gdn.cu:1854, ~1-2%); (c) GDN recurrence output core_attn_out + z → bf16 (1.5-2.5%); (d) SwizzleBlockscale redundant per-projection re-run → swizzle shared act-SF once (0.5-1%). MEDIUM — (e) full-attn preamble fusion 4→1 bf16 + cos_sin_cache (3-5%, the biggest coherent lever); (f) GDN in_proj mixed_qkv + causal-conv1d f32→bf16 (widest tensors, 3-4%); (g) fp4 GEMM per-shape autotune 8 flashinfer SM120 configs + StreamK for large-K down/o_proj (2-4%+1-3%); (h) swizzle-in-quant emit swizzled SF directly (1-2%); (i) CausalConv1d retile 2D + register window (1-2%). HIGH — (j) DeltaH/ChunkO V-split + register state + cp.async (2-4%); (k) prefill-attn cp.async (FA-2 addresses). Stackable low+medium ~ enough to close a large chunk of the 23% prefill gap. NEXT: drive the low-effort cluster first, then the preamble-fusion + bf16 cluster.** |
 
 | 2026-07-08 (PREFILL lever #1 driven from the gap-scan: GDN bandwidth cluster — v_new slack-memset + recurrence-output bf16, +0.8% token-exact, MERGED default-OFF) | First levers from the 44-agent scan. VT_GDN_SLACK_MEMSET: replace full v_new cudaMemsetAsync (~97% dead) with a slack-only zero kernel (cuda_gdn.cu:1705/1901; mirrors FLA boundary_check, no memset). VT_GDN_OUT_BF16: GDN recurrence output dcore + z gate bf16 (halves [T,Hv,Dv] core traffic; the f32 dcore VT_BF16_GDN left untouched). Merged to main (default OFF toggles, code available to stack on). | **All gates token-exact: OFF 9/9, slack-memset 9/9, out-bf16 9/9, both 9/9.** | **MEASURED A/B (in1024/out128 conc32 np192, baseline 829.90): slack-memset 831.73 (+0.2%, marginal — memset overlaps compute), out-bf16 836.77 (+0.8%), BOTH 836.84 (+0.8%). Small real win, driven by the bf16 core (bandwidth cut). Near the ~1.2% noise floor → kept default-OFF pending a clean multi-sample cumulative A/B once more levers stack (vLLM's +22% prefill edge is the SUM of many such small fusions/traffic-cuts). LESSON: bandwidth levers give small (~0.8%) but REAL positive wins (unlike the launch-overhead fusions that were neutral). Continuing to stack: preamble fusion 4→1 (biggest, 3-5%), fp4 GEMM per-shape autotune, swizzle-in-quant, GDN in_proj/conv bf16.** |
+
+| 2026-07-08 (PREFILL lever #2: fused full-attn preamble 4->1 + cos_sin_cache — +0.8%, CUMULATIVE +1.9%, token-exact, MERGED default-OFF; STACKING VALIDATED) | Biggest confirmed gap-scan lever (tasks #14+#15). VT_FUSE_ATTN_PREAMBLE: one vt::AttnQkNormRopeGate kernel (split+gemma-qk-RMSNorm+partial-NeoX-RoPE+gate) replaces 4-5 separate f32 kernels in the 11 full-attn layers; + vt::RopeCosSinCache (per-step [T,rot] fill from sdi.positions, kills RopeNeoxKernel's DOUBLE-precision per-element cos/sin/pow recompute). Outputs f32 (value-exact; bf16-q is a follow-up). Byte-identical CPU parity for q/k/gate proven. | **Gates token-exact: OFF 9/9, preamble 9/9, cumulative 9/9.** | **MEASURED A/B (in1024/out128 conc32 np192, base ~832.8): preamble 839.48 (+0.8%), CUMULATIVE (preamble+VT_GDN_OUT_BF16+VT_GDN_SLACK_MEMSET) 848.26 (+1.9%). STACKING VALIDATED — 2 lever-clusters stack roughly ADDITIVELY (+0.8% + ~0.8% ≈ +1.9%), confirming vLLM's +22% prefill edge is the SUM of many small fusions/traffic-cuts, each individually small (why single levers measured 'neutral' — they're ~0.8% at the noise floor, only visible cumulatively). Merged default-OFF (flip all ON after a clean full-cumulative A/B). NEXT: stack fp4 GEMM per-shape autotune (43% bucket, biggest remaining), swizzle-in-quant (re-profile showed 2856 SwizzleBlockscale launches), GDN in_proj/conv bf16, CausalConv1d retile → target ~0.85x prefill.** |