CUDA Inference Path — Optimization History and Current State

STATUS (2026-09-06, post-r60): history-organized reference. This document was restructured from a part-based roadmap into a history-ordered record: §0 is the master history table (the outline — every landed, reverted, or measured lever with its commit and perf delta), §1 is the current state, §2 is one chapter per table row, and §3 holds the appendices (env-gate reference, methodology, and the pre-Phase-7 legacy history). Single-sourced implementation records: docs/CUDA-BACKEND-DESIGN.md (Phase 7a–7e) and the per-step documents of §2 (Phase 8: docs 01–11 plus the supplementary records 78–79 — the former CUDA-FOLLOWUP-PLAN.md was consolidated into them and retired on 2026-09-10); the per-round MMQ redesign records are mirrored in docs/LLAMA-CPP-MMQ-ANALYSIS.md §11. Default env = the verified 1.080×-vs-llama path; MINFER_MMQ=0 = the legacy f16 escape (§1, Appendix A).

§0 Master history table — the complete optimization record

One row per optimization step: every lever that landed, every lever that was reverted, and the measurement-only rounds that directed the campaign. Chapters in §2 follow this table row by row.

Reading conventions.

  • Perf column: whole-prefill tok/s for the 7B q4_k_m model on DGX Spark GB10 unless noted. The absolute anchor moves with the prompt length used by each session (2630/2659 → 3325/3314/3354 tokens ≈ "pp2K/3314-eq") and with machine state (co-tenant load); every row's numbers are interleaved same-binary A/B medians within one session window — cross-session absolutes are not comparable (the r59b lesson). "—" = not measured at whole prefill (kernel-level or decode-level metric instead).
  • vs-llama column: whole-prefill vs llama-bench at the 3325-eq (later pp3314) anchor — 3324.42 tok/s (clean machine) and 3323.29 (r59b window). The 8m–8p/P5 rows' early multipliers are vs the 3401 @2K llama-bench figure instead. Before r37 the campaign ran on the opt-in MMQ path, so no whole-prefill vs-llama was recorded at the 3325-eq anchor for those rows ("—"); the default f16 path sat at 1.43× (2340–2370 vs 3401) from P5 until the q6_K/FA lines moved it.
  • Commit(s): code commit first, record/docs commit second where both exist. Twelve hashes quoted in the older session records are pre-amend duplicates that no longer resolve (e.g. the r59 record's feb37de); this table cites their reachable twins (same subject — see the r59 chapter for the one case worth naming).
  • Cross-reference: the r28→r59 MMQ-redesign rounds are mirrored, round for round, in docs/LLAMA-CPP-MMQ-ANALYSIS.md §11.8–§11.37 (11.8 = r28 Phase-2 outcome, 11.26 = r47, 11.34 = r55, 11.36 = r58, 11.37 = r59).
  • r26–r27 do not appear in the record — the P6 round numbering skips from r25 to r28 (the intermediate commits are the MMQ-analysis §11 design/doc work, kept in docs/LLAMA-CPP-MMQ-ANALYSIS.md).
StepLeverCommit(s)Perf before→after (anchor)Δvs-llamaStatusOne-line lesson
Phase 7CUDA backend: raw FFI device layer + graph backend (7a–7e)0dc2a54baseline 7B @2K 30.7 tok/s—~110×LANDEDresident weights + per-op dispatch ended the Part-IV ping-pong
8m/8m②tiled wmma f16 prefill GEMM + cp.async tile stagingba3f317, cdc659930.7 → 294 → 1204 tok/s @2K39×~2.8×LANDEDone tensor-core GEMM over all 8 types replaces per-token weight re-streaming
8nFA-style tiled prefill attentioncb66fca176 → 8.5 ms/layer @2K20×—LANDEDonline softmax + register O accumulator; 256 B P stride avoids a score-clobber race
8odecode-start CPU stalls killed65b686cfirst decode step 724 → 35 ms20×—LANDEDCow::Owned clone + eager concat probe cost ~1.6 s per graph rebuild
8ppersistent f16 weight cache + fused dequant-in-GEMM2992f57 (+b9e7a91 docs)7B @2K prefill → ~1400–1500~4.7× vs 8m~2.3×LANDEDdequant once per weight at load (≥2 GB gate); exposed a latent Q5_0 misaligned load
8bKV f16 on CUDA (store_kv_f16 + f16-KV attention mirror)f7b00367B @2K decode +~11%+11%—LANDEDauto-f16 at n_layers×n_kv_embd ≥ 8192 (MINFER_CACHE_TYPE override); f32 accumulation
8cprefill Q8_0-activation GEMM, shape-gated69a27c50.5B @3.6K 1005 → 1246 tok/s+24%—LANDEDthe shape gate is load-bearing: −63% at 7B ffn_down (weight-bound shapes stream slower)
8dsplit-K flash-decoding decode attentiona5af60f7B @2K decode 10.1 → 13.7 tok/s+36%—LANDED28 warps → an 8-way KV-split scan; superseded by R4's dim-parallel rewrite
8fQ5_K + Q5_1 f32-activation kernelsb959ec90.5B q5_k_m admitted to CUDA (was CPU wholesale)——LANDEDthe all-or-nothing gate needs a kernel for EVERY matmul weight type
8e/8e②decode MMVQ (dp4a, per-type kernels, shape gate)b7b8e73, 1298cb2, 1d282357B decode +37% (q4_K), then q6_K/q5_K+37%—LANDEDinteger dp4a dots + llama.cpp MMVQ_PARAMETERS_GB10 launch table
8lllama.cpp parity benchmark (the decode + prefill gap sheet)acca28fdecode 1.15–1.76×, prefill 18–110×——MEAS-ONLYfound the Q5_K registration gap (51.6 → 246.3 tok/s, 4.8×); the sheet ranked 8m–8p / R1 / the r-campaign
8qQ5_0 CUDA enablement + u16 qh loads9f419f90.5B q4_k_m 148.7 → ~1200 prefill / 56.9 → ~306 decode——LANDEDa 22-byte block's qh word is not 4-byte aligned — two u16 loads; the CPU fallback eliminated
R3-A1single-split prefill (tail_ids input at graph head)029a9a44 splits → 1 per prefill forward——LANDEDa mid-graph declared input forced 2 extra full-stream syncs per forward
R3-A2pinned D2H logits readback, no redundant clonea213c89parity-to-slightly-ahead under load——LANDEDpageable readback paid a driver-internal pinned bounce
R3-Bprefill capture defaults ON (3-run protocol)761e236repeated identical-nt prefills capture automatically——LANDEDone-shot CLI prefill never reaches 3 runs and pays nothing
R1int8 MMQ prefill GEMM, opt-in MINFER_MMQ=140e97c9155 (co-tenant) / 412 quiet vs f16 630–880 / 1460; 441 in the r7–r8 window——LANDED (opt-in; superseded by raw line)parity-clean but ~2.9 TMAC/s vs llama ~24 — the 8× gap was unprofiled
R2MMVQ weight-streaming rework (per-thread sub-pairs, uint4)6df3245tg128 42.2 → 45.1; @2K 36.7 → 38.8+6.9% / +5.7%1.05× / 1.16×LANDEDper-sub-block nibble re-reads doubled load instructions; L1 hid the bytes, not the issue stream
R4decode split-attention dim-parallel rewrite70f57db@2K 39.2 → 43.2–45.1; tg128 45.1 → 47.5+10–15%~0–4% / aheadLANDEDLOCAL-memory float4 oc[32] accumulator = ~80 MB/layer of local traffic
P5·0FA prefill P·V on tensor cores86ca78c10.06 → 4.24 ms/layer; @2K prefill +15%+15%2.37×LANDEDwmma for P·V halves the FA pass
P5·1elementwise vectorization (store_kv, convert)d713e6e1435 → 1493+4%2.28×LANDED1-elem kernels left 15/16 of every transaction unused
P5·2128-wide GEMM tiles (TM=128)725e3071493 → 2267+30%1.50×LANDEDhalves B-panel L2 re-reads and barriers per FLOP; fb[1] offsets +16 elements, not rows
P5·3FA softmax on all 8 warps + padded smem rowsfc07c042267 → 2365–2371+4–5%1.44×LANDED256 B rows ≡ 0 mod 32 banks = 8-way ldmatrix conflicts; +8-half stride fixes
P5·4GEMM k-step KS=641365c821464 vs 2345−38%—REVERTED56 KB footprint halves resident blocks — depth vs occupancy inverted
P5·negTM=256 GEMM tilea189837−3%−3%—REVERTEDwider tile, same wall
P5·negin-kernel f32→f16 A staging (AF32 mirror)b254c22 (WIP a3b0dcd, 69c3933, aa40ed3)−8% end-to-end; parity hole open−8%—REVERTEDthe convert pass is cheaper off the hot path (r23 later quantified it at 6%)
r5–r6MMQ re-rank + structural rewrite spec1e0673f, 491eb5cKD=4 re-negative (427 vs 438); 4-warp 32×32 tile 399 + parity hole; dequant pass is at LOAD, not in the wall——MEAS-ONLY + REVERTEDdepth (KD=8) beats occupancy for the word-staging kernel; spec: raw-byte smem, dequant at mma time
r7–r8raw-byte MMQ kernel + 128-token wide tile; FA KV L2-prefetch probed440d16, d9d626a, a41eac0, ef9d5b4, 87bade0raw 472 vs 441 (+7%), quantize 129→74 ms; FA probe null (2319–2345 vs 2345)+7%—LANDED (raw) + REVERTED (probe)wide KD=8 first measured 2124 = phantom (silent smem-cap failure) — guard the cap; GQA already keeps KV L2-hot
r9llama.cpp MMQ reference decoded; shape axis closed84831d2narrow cp.async KD=8 481 = local optimum (6-shape matrix)——MEAS-ONLY~0.018 inst/MAC/thread vs our 0.133 — the ratio, not tile shape, is their speed
r10reference inner-loop decomposition ported025a69f462–468 vs 470 (flat, ~6.4 TMAC/s)~0—REVERTED (edits lost to a post-checkout hook; measurements valid)same decomposition as llama still 5× slower — residual is ILP depth/ldsm/tile
r11ILP-chain reading verified (tile ne = I·J/32)83e5580, 6f29e65———MEAS-ONLYthe I·J/64 ne was the AMD MFMA branch; NVIDIA needs 4 C regs per m16n8k32
r1216-chain warp tile + ldmatrix774a116wide-16 KD=4 1020–1058 vs narrow 441–481 (~2.3×); KD=8 973–9952.3×—LANDEDaccumulator depth wall passed; 4-warp 32×32 and 48B A-padding both negative
x-tile256-token wide-MMQ blockf061cb8~942 vs 1035–1058; kernel +27%−9%—REVERTEDA re-reads (327 MB) already ≈2× B re-reads (152 MB) — trading B for A is strictly negative
j-tile128tok × 256od A-reuse (outer jh loop)4993804~1067 (+2.6%, bar 1150 missed)+2.6%—REVERTEDL2-byte savings do not convert to time at 1 block/SM — latency-bound, not byte-bound
cp.async-dbcp.async double-buffered raw staging784786d~1034 (neutral; baseline band 1035–1050)~0—REVERTEDkernel is L2-throughput-bound, not MLP-starved; re-timing identical bytes wins nothing
r13ncu counter forensics vs llama.cpp + FULL/MINIMAL staging fixes5ca037dFULL 865–871 (−16%); MINIMAL +0.3% (noise)——MEAS-ONLYthe 3× gap is the per-MAC warp-instruction stream (10.14 vs 6.06 M/GMAC), not bytes
r14B-fragments via one ldmatrix.x4 + widened scale loadsc64cd991225 @KD=4 / 1273 @KD=8; kernel 3.632 → 2.378 ms+18.5% / +23–30%—LANDEDslot-major 48B qb8 tiling is conflict-free; fewer/wider smem ops, zero staging-ALU growth
r15f32-accum s8 mma probe (dead) + rank-1 term2 rescaleb999e9a1295 @KD=8; inst 9.45 → 8.69 M/GMAC+1.9%—LANDEDf32-accumulate integer mma does not exist (ptxas probe); merged-chunk rescale is mathematically invalid
r16narrow kernel gets the rank-1 fold151fa97480 vs 447–473+1.7% (noise-band)—LANDED3-site port of r15; narrow is not the perf path
r17wide warp remap 32od × 64tok9d09a81inst −5.8%; wall +0.9/+1.0% (noise)~0—REVERTEDpure per-MAC instruction cuts pay ~0 wall while SM% sits at ~30 — stall-bound
r18load-time B pre-expansion (bulk-copy staging)0a26b35KD=8 +0.9% (noise); KD=4 −19% median~0—REVERTED+5.8 GB for a −4.5% kernel-time that does not reach the wall; EB/SB machinery preserved for future L2 experiments
r19weight L2 residency (__ldg, persisting window)072dd9a__ldg +2% (noise); L2WIN −50%~0 / −50%—REVERTEDweight tiles already re-read from L2; a persisting carveout starves C stores/activations/KV
r20split-phase A staging5ac89171230.4 → 1317.7 @KD=4; 1275.8 → 1319.9 @KD=8+7.1% / +3.5%—LANDEDthe gap carrier is long_scoreboard (97% of named-stall excess) in the LDG→STS chains; freed stalls re-saturate on lg_throttle
r21coalesced block-linear A staging3c009ccsectors −28.6%, lg_throttle −90%, wall −2.3%/−0.7%−2%—REVERTEDstall mass is conserved: sectors/queue are not the binder, warp-instruction count is
r22qa8 XOR swizzle (+ d/ssum fold negative)5b40058KD=8 1329.6 vs 1311.8 (+1.4%, 3/3); op_ld 16.86 M → 0+1.4%—LANDED (fold REVERTED)precompute all 8 A-frag offsets once — per-ldmatrix address ALU eats freed wavefronts; d/ssum loads are L1 hits
r23f16-path full-graph wall decomposition + FA_TKV=32 occupancy lifte8c348dlift: occ 16.7 → 32.68%, kernel −6.7%, wall −0.3%−0.3%1.43× (f16 path)MEAS-ONLY + REVERTEDFA's 2.5×/layer gap is structural (llama keeps 128-wide KV tiles); GEMM is 74% of the f16 wall
r24scheduling ladder: tile-order swizzle + persistent blocksd90b3e9swizzle −2.3/−4.7%; persistent −3.3%negative—REVERTEDdefault B-hot x-fastest order is best; no wave-quantization tail exists to remove
r25SASS opcode-class census + kd-unroll attempt8658f1binst −9.7% (surplus halved to +46.6k/tile); wall +0.37/+0.49%~0—MEAS-ONLY (census is the deliverable) + REVERTED100% of the +25% inst surplus is support instructions (int ALU 77%) — but it is wall-inert at 1 block/SM
r28Direction-A raw-nibble NB kernel (2 blocks/SM)0957a08 (design 2f783a3)1375.2 → 1410.4 @KD=8; 45,056 B smem, 123 regs+2.56%—LANDEDoccupancy was the binding resource; unsigned-nibble + rank-1 rescale is the B-frag contract
r29NB kd-loop unrollbfe6bba1387.9 → 1426.8; int ALU −25%, inst −6.5%+2.80%—LANDEDr25's wall-inert int-ALU cut becomes real at 2 blocks/SM — occupancy unlocks instruction cuts
r30SWAR word-granular B-nibble unpack0071b31+0.54% (noise); SASS byte-identical~0—REVERTEDr29's unroll already induced the exact CSE — check SASS before writing the lever
r31q-major sda scale-read repack851a896 (+76d495a docs)1424.10 → 1439.40; LDS.64 32→0, LDS.128 16→32+1.07%—LANDED (sub-bar)naive q-major repack was 2-way conflicted — region-split layout is conflict-free
r32finite-lever sweep (staging hoist, epilogue widening)153d28cepilogue +0.46% (structurally capped ~0.4%)~0—REVERTEDstaging addressing already hoisted by ptxas; run-once epilogue cannot clear a bar
r33hybrid inner-loop port (llama j0/n order)697ef04median −0.25%; SASS byte-identical~0—REVERTED (hypothesis falsified)ptxas already schedules the 8-mma + rescale loop optimally — source reorders are SASS no-ops
r34quantize-transpose prepass (A-side layout transform out of the kernel)ba977bf1364.2 → 1496.8 @3354 tok; prepass 0.908×; 103 regs+9.72%—LANDEDthe residual was layout-transformation locality (llama's quantize_mmq_q8_1 design), not instruction composition
r35sda d/ssum scale pre-decode6112db3−0.46%; SHF 64→0 but wall flat~0—REVERTEDdecode ALU hides in the IMMA shadow — removing int/fp ops that fill idle slots frees nothing
r36A-frag wavefront economics (H1/H2 endpoint)f44fc441.76× shared wavefronts/IMMA but 1.85× wavefronts/s at equal IMMA rate——MEAS-ONLY (H1 refuted)the MIO pipe is not scarce; llama's edge is A-fragment reuse (0.125 vs 0.5 LDSM/IMMA), a tiling property
r37post-parity whole-prefill attributionea234f11521 tok/s; q6_K GEMM 1094.7 ms = 51.2% of wall at 6.38×/GMAC—2.15×MEAS-ONLYMMQ made q4_K fast and left q6_K on a slower-than-f16 path — the next lever is a different kernel
r38q6_K BT-style raw-byte mma kernel (KSPLIT=2, KDR=4)75aabb91518.4 → 1561.9; q6_K 368.9 → 221.8 µs/GMAC (1.66×)+2.87%2.13×LANDEDq6_K is 16 sub-blocks of 16 (not 8×32) → two m16n8k16 with separate dsc; KDR=8 regressed to 1097.8
r39q6_K KDR=2 double-buffer (A+B pipelined)f2b9e541568.7 → 1777.5; attn_v kernel −19.7%+13.3%1.87×LANDEDdoubling every plane at KDR=4 = the 1-block/SM trap; KDR=2 hits the same 29,696 B with real overlap
r403rd resident block via __launch_bounds__(256,3)65ecef71784.0 → 2015.6; kernel −23%+13.0%1.65×LANDEDthe 0-spill gate is disproven-immaterial: 80 regs + 4 B spill beats 87 regs at 2 blocks
r41q6_K B-expand widened to uint4 groupsb891e1b (+aa82e8f docs)1979.9 → 2605.2; kernel 1.70 → 0.654 ms (−61.5%)+30.7%1.27×LANDED32 per-byte ql/qh LDGs per thread-kt = the L1TEX scoreboard (85.5% → 33.6%)
r42q6_K stage-wide dsc scale reada1421e6−0.19%; L1TEX traffic down, stall share unchanged (33.6%)~01.27×REVERTEDcutting dsc bytes does not cut dsc latency — the stall is at the I2F consumer
r43PC-sampling attribution + pre-expand-B (parity FAIL)b7fa305attribution: B-expand recomb 45% + A-sts 28% + dsc I2F 26%; W_exp byte-correct but diff 448—1.27×MEAS-ONLY + REVERTEDattribute stalls to the consuming instruction; byte-correct data at wrong offsets = stride mismatch next door
r44W_exp stride mismatch root-caused; dense-index fix6d02017parity green; kernel −10.9% but wall −0.42%~01.27×REVERTEDdense W_exp indexed with the padded raw-W stride = the paradox; removing recomb only transforms the latency
r45cp.async the q6_K A-side staging9825ffdkernel −10.2%, longsb −18%, wall −0.34%~01.27×REVERTEDafter r41 the q6_K GEMM is no longer the bottleneck — a faster kernel that is not the wall does not reach it
r46 (FAP1)FA audit + FA_TKV 64→32 + S/P row paddinga186f51FA kernel 5.16 → 4.58 ms (−11%); wall +0.27%~01.27×REVERTEDFA was already wmma + online-softmax — occupancy-starved and S/P-conflicted; but not wall-critical yet
r47converged-regime wall decomposition (r37 table stale)11e36402585–2623 tok/s; q6_K 1094.7 → 196.4 ms; FA = #1 residual 5.72× (125.8 ms, 10.2%)—1.27×MEAS-ONLYq4_K 1.06× and q6_K 1.13× both at parity — recommend FAP2 register-resident softmax (2× → −4.9% wall)
r48 (FAP2)register-resident softmax in fa_prefill_f16kvd38744d (+7e2ee62 docs)2603.5 → 2749.9; FA kernel 5.16 → 2.12 ms (2.43×)+5.6%1.21×LANDEDsoftmax on the QK^T accumulator fragments; P built in-register as the P·V A-operand; K col_major vs V row_major is the trap
r49A-quantize prepass shared-A dedup (window memoization)87a75a32734.1 → 2797.5; prepass 193 → 110 launches, 118.4 → 83.9 ms+2.32%1.18×LANDEDq/k/v and gate/up share one A — the prepass was re-quantizing it per GEMM; cache keyed on (src ptr, nt, id), cleared at any non-MatMul node
r50FA_TKV 32→16 occupancy trial9128468−0.5% / −0.01% (3 blocks/SM reached); greedy identity lost~01.18×REVERTEDoccupancy gain cancelled by doubled per-tile sync/softmax overhead — FA_TKV reduction is a dead lever that also breaks byte-identity
r51producer-fused A-quantize, mode 1 (rms/swiglu emit pad40_t)cf1ed4b (+bf0c986 docs)2803.4 → 2856.4; prepass 110 → 28 launches, 83.0 → 10.1 ms+1.89%1.16×LANDEDthe quantize input is L2-hot in the producer; register-resident swiglu quantize was rejected (uncoalesced f32 stores)
r52fused-producer phase 2: skip-write mode (MINFER_MMQ_A_FUSE=2)910d967 (+fb659f7 docs)2855.7 → 3011.3; fused producers 151.9 → 86.5 ms+5.45%1.09×LANDEDf32 output is provably dead under the window enumeration; dead-write backstop turns violations into loud errors
r53q6_K bundle: pre-expanded-B dense W_exp + cp.async B staging83fee77 (+4907d9f docs)3024.7 → 3176.9; ffn_down kernel −20.5%; +1.52 GB device+5.03%1.05×LANDEDr44 (removes WORK) + r45 (removes WAIT) are individually wall-neutral and compose — the basket thesis
r54MINFER_MMQ_Q6K_EXP opt-out of the W_exp plane3252e96 (+b860b7e docs)default 3181 (unchanged) / EXP=0 3020.7; 7636 vs 6182 MiB−5.04% for 1.52 GB1.04×LANDED (gate)memory-for-speed knob with a three-way liveness label (exp=off vs fallback!)
r55fused-swiglu roofline audit + one-shot prefill CUDA-Graph decision83c3c67swiglu at 242 GB/s = 89% roofline (cap +0.74%); capture ≤ +0.1% + capture-illegal malloc—1.05×MEAS-ONLY (both documented skips; campaign CONVERGED)bound the roofline before coding — no implementation of this kernel can clear the bar
r56q6_K A-side bundle: A cp.async + W_dsc f32 plane4cf7c74 (+29084de docs)3138.6 → 3212.5; ffn_down −5.9%, attn_v −4.2% kernel; +363 MB+2.35%1.035×LANDEDpost-r53 the A-side wait and dsc consumer became the wall — r45's mechanism finally lands in a bundle
r57FA KV staging double-buffer (FA_TQ=48)c3268ccgreedy-32 identity lost at token 19 (×2 attempts)—1.035×REVERTEDFA_TQ is a tile size too — the r50 rounding caveat applies to any FA tile change
r58q4_K BT structural spec + cp.async-db2 transplant093ae41 (code reverted)2819.3 vs 3227.6 (−12.6%); spec: gate/up −37% potential, ceil-waves 2.6%−12.6%1.03×REVERTEDa mechanism whose COST depends on the granularity of what it replaces is amortization-bound (the r45 mirror)
r59q4_K W_dsc f32-pair plane + pre-warm/pre-grow riders36a481f (+15c04ba docs)recorded +26.2% (2843.2 → 3588.8, co-tenant) — superseded by r59b; bt kernel busy −30.9%; +1456 MB+11.1% (true)~aheadLANDED (Δ corrected)gate/up (−37%) and q/o (−18%) carried it, not ffn_down — the r58 "staging scales with kt" reading confounded decode cost with A-plane DRAM traffic
r59bclean-machine re-measure + baseline-poisoning correction074ca94definitive 3590.8 (HEAD) vs 3232.0 (fresh baseline rebuild) = 1.080× llama-bench 3323.29 @pp3314+11.1%1.080× (ahead)MEAS-ONLY (correction)the r59 "baseline" binary was the stale r58-delta build (−12.5%) — anchor every A/B baseline behaviorally in the same window
r60PROMOTION: the verified MMQ gate set flips DEFAULT-ON57edcf6 (+7029ee4 docs)default ≈3578–3599 (~3581); MINFER_MMQ=0 legacy f16 ~2226–2353-class; planes +3.27 GB—1.080×LANDEDpromotion = default-on with "0" opt-outs (r54 pattern); the bisect caught the mode-2 multiturn break → NB-BT-only guard
D1decode @1641 KV attribution: gqa_attn_split_partial is 100% of the KV-scaling wall (34.1 µs/launch, 76.5% long_scoreboard); ATTN_SPLITS sweep = dead endmeasurement-only (/tmp/d1/)tg128 49.3 (KV~1) / 47.2 (@1641) vs llama 49.41 (tg128)—0.956× (tg128)MEASUREDprobe-verified: staging-depth changes bitwise-safe (ndiff=0); ATTN_SPLITS changes reorder the float sum
D2explicit K+V register staging in gqa_attn_split_partial (4-row window staged before the softmax chain); cp.async smem pipe + pair-lookahead measured worseD2 commit (this row)decode @1641 47.2 → 48.2 (+2.0%); kernel 34.1 → 19.4 µs/launch; tg128 49.4 flat—0.975× (tg@1641)LANDEDbitwise-identical end-to-end (greedy-32/256 byte-identical); probe −42%, nsys −43%, wall +2.0% agree
D3b-1bdown-q6K pipelined MMVQ q6_k_q8_mmvq_v2_pf (npair>256: both serial units' weight+q8 loads issue up front)f1825b57B decode tg128 48.03→49.47 (+3.0% SEP), @1641 46.66→47.94 (+2.7% SEP); 14B tg128 22.80→22.90 (+0.44%), @3254 21.01→21.06 (+0.24% SEP)+2.7–3.0% (7B decode)0.970× (7B tg@1641, this window)LANDEDbitwise-identical (114/114 dump memcmp, greedy-256 byte-identical, suite 169/0/3); for npair>256 the second serial unit's exposed load latency WAS the 198.9-vs-220 GB/s gap — D4-2 CORRECTION: the 7B numbers are void (the dispatch dropped units 512..591 on npair-592 rows; the "gain" was mostly the missing work — see the D4-2 chapter); 14B numbers stand
D3b-1aattn_v-q6K off the padded-f32 kernel: (a) MMVQ routing via a lowered od*id>=24M gate — NOT bitwise (MMVQ quantizes activations to q8, different accumulation semantics); (b) NSG 2→1 row→warp re-map — bitwise-green but kernel 36.4→39.9 µs (2× warps = 2× y re-read L2 traffic)reverted (both routes)14B tg128 −0.74%, @3254 −0.33%——REVERTEDthe padded kernel is not warp-starved; y re-read traffic scales 1:1 with warp count — rows-per-warp is the only bitwise-free knob and 2 is already the sweet spot
D3b-1coutput-head dynamic block size (npair 160 → 160-thread blocks, warp-count-bounded mmvq_block_reduce)reverted (patch /tmp/d3/patch_1c.py)7B @1641 +0.26% (SEP); 14B tg128 +0.04%, @3254 +0.09%——REVERTEDGB10's 1536-thread/SM limit: 9 blocks×160 live threads ≈ 6×256 allocated (960 live) — the idle-thread win does not exist at 14B shapes
D3b-2short-KV combine skip (single-split path for nkv ≤ threshold)not implemented———ANALYSIS-NEGATIVEsingle-split ≠ 32-split partial+combine bitwise for ANY nkv>1 (the merge reorders the float sum — D1's split-count evidence: ndiff 3.6e-3 of outputs, max|Δ|~3e-9); the split grid is frozen by CUDA-graph replay capture; the bitwise-safe residual (combine early-out of empty splits — exact +0.0 terms) is ≤ ~15 µs/step, below every bar
D3a4-warp fattn-vec-style split-attention rewrite (gqa_attn_split_partial_h4w, hd=128: 128 threads, K/V streamed, Q in registers, 8-lane subgroups, 32-row windows/warp, smem LSE merge; grid unchanged, replay-safe)reverted (patch /tmp/d3/d3a_kernel_patch.diff, findings /tmp/d3/D3A_FINDINGS.md)kernel 14B @3254 73.4 → 68.4 µs (−6.9%) but 7B @1641 21.1 → 34.7 µs (+64%); wall 14B @3254 −0.95% (noise), 7B @1641 −3.4% (real), tg128 noise-level——REVERTEDrows-per-warp pathology: rpw = ceil(ceil(nkv/32)/4) = 26/13/1 at @3254/@1641/tg128 — the 32-row window idles 59–75% of lanes below rpw≈16 and per-block fixed costs amortize over rpw; kernel −6.9% at the best shape is only ~+0.5% wall (attention = 7.4% of the step), under the +1.5% bar and the ±2% A/B noise; numerics fully green (probe ≤1.3e-7 vs CPU, argmax hard-gated, greedy 0/10 diverged) — the session's durable output is the tolerance-gate calibration: end-to-end max|Δlogits| is 0.38/0.39 (14B/7B) for ANY accumulation-order change, so the D3-1 ≤1e-3 logits gate is unsatisfiable; argmax+greedy+A/B are the operative gate set
D3-4 L1hybrid rpw dispatch: dual-kernel self-gating split attention (f16 KV, hd==128) — 4-warp h4w kernel when rpw = ceil(ceil(nkv/32)/4) ≥ 16 (nkv ≥ 1921), incumbent 32-thread kernel below; BOTH launch per layer with static grids, each re-reads positions[0] per replay, exactly one is live per nkv (nkv-uniform branch → replay-safe)22336b214B @3254 split 72.1 → 62.1 µs (−13.9%) + 1.5 µs dud launch; wall 14B @3254 21.20 → 21.33 (+0.61%, SEP), tg128 22.96 → 22.94, 7B tg128 50.28 → 50.20, @1641 48.78 → 48.68 (guards hold)+0.61% (14B @3254)0.944× / 0.877× (14B tg128/@3254 vs llama 24.31/24.32)LANDED1-warp path bitwise (7B @1845 dump: all gated files identical; node{3,5,8}_prefill diffs = pre-calibrated slot aliasing); h4w tolerance class (max|Δlogits| 0.309, argmax identical margin 0.716, 1 greedy flip at the regime entry = 1/256 < 2%, temp-0.8 controls identical); suite 169/0/3 incl. the hd=128/n_ctx-4200 parity shape sweeping the rpw 15/16 boundary; an in-kernel 1-warp-fallback form was REJECTED pre-commit: inside 128-thread blocks the 1-warp body caps at 12 working warps/SM (1536/128) = +78% kernel at 7B @1641 (35.4 vs 19.8 µs nsys) — geometry, not math
D3-4 L2window-level K/V prefetch pipelining in the h4w body (K-pass software pipeline +2 uint4, 4-deep V-bulk ring +8 uint4; issue-point-only → bitwise vs h4w by construction)reverted (patch /tmp/d3/patch_l2.py)14B @3254 h4w 62.1 → 66.5 µs (+7%)−7% kernel—REVERTEDthe kernel is bytes+tail-bound at 79% of the 48.9 µs floor, not chain-bound enough: funding the pipeline buffers needs __launch_bounds__ minBlocks 8→4 (64→128 regs) → occupancy 32→16 warps/SM and 3.33→4.44 waves — the occupancy/wave-tail tax outweighs the shorter load chains; the D2 4-row-scale lesson (issue-point moves are free) does NOT transplant to window scale under a 64-reg budget
D3-4 findingspre-existing behaviors calibrated this session: (a) at long prompts (≥2.8K tokens) MINFER_GRAPH_DUMP PREFILL-phase files (all kv*_prefill, logits_prefill, prefill nodes) are non-deterministic pre-vs-pre (wholesale, garbage-magnitude — aliased dump reads); decode-phase dumps stay deterministic; (b) CLI prompts longer than the n_ctx default 4096 leave zero generation headroom (position N exceeds n_ctx N panic, graph.rs:367); bench unaffectedmeasurement-only———RECORDEDdump gates at long prompts must anchor pre-vs-pre at the EXACT shape and gate only decode-phase files; long-prompt CLI greedy needs prompt + n ≤ 4096 until n_ctx sizing is fixed
D3-5 1afused-producer decode A-quantize: rms_norm_quant_pad40 (rms body + the standalone per-block quantize body, both verbatim) and swiglu_quant_pad40 write the pad40 q8 plane beside their f32 output; decode matmuls consult decode_quantize_native (MmqCache, native form) and skip the standalone quantize launch on a hit; FusedFFN joins the cache-clear preserve set; MINFER_NO_DECODE_A_FUSE=1 opt-out3230b2bnsys 14B @3254 (NO_CUDA_GRAPH): standalone quantize_q8_0_pad40 4448 → 964 launches (−78%; the ~50/step remainder = the attn_o class), total kernels 29402 → 25918, sub-2µs 15171 → 11723; wall 14B tg128 23.05 → 23.35 (+1.30% SEP), @3254 21.36 → 21.68 (+1.50% SEP), 7B tg128 49.86 → 50.69 (+1.66% SEP), @1641 48.47 → 49.16 (+1.43% SEP)+1.30% (14B tg128)0.961× / 0.891× (14B tg128/@3254 vs llama 24.31/24.32); 7B 1.026× / 0.995×LANDEDq8 bytes bit-identical by construction (max is exact for any association; rintf/clamp elementwise — the epilogue IS the standalone body) and probe-verified bitwise (rms+swiglu q8 buffers, f32 producer outputs, MMVQ outputs through the cache-hit path); suite 170/0/3; 14B −n 4 dump gate: logits both phases + all KV byte-identical, the 3 node-dump diffs are same-binary pool-slot aliasing reproduced pre-vs-pre AND post-vs-post; greedy −n 256 token streams byte-identical both models; fused epilogues add ~0.2–0.25 µs/launch (swiglu 2.06 → 2.31 µs), priced into the wall
D3-5 1boutput-head od-split / 512-thread re-map (lm_head q6_K od 152064, id 5120, npair 160)not implemented———ANALYSIS-NEGATIVEall three candidate mechanisms are measured or computed dead at this shape: (a) idle-thread removal (96 of 256 idle at npair 160) = D3b-1c, measured neutral; (b) rows-in-flight: 288 resident rows either way (6×256-thread blocks/SM vs 3×512 dual-row), and D3b-1c's 9×160 = 432-row form was ALSO neutral — occupancy is not the limiter; (c) block-scheduling rate: the head sustains 47.6 blocks/µs while ffn_gu demonstrates 76/µs — not the limiter. A dual-row 512-thread form is bitwise-capable (per-row half-block reduce with the same 8-warp tree) but carries no mechanism → not built per the "measured mechanism, don't guess" rule
D3-5 1cffn_down-q6K 512-thread single-unit variant (id 13824 → npair 432)not implemented———ANALYSIS-NEGATIVENOT bitwise vs the landed v2_pf: in the 256-thread form thread t accumulates fma(u_t) then += fma(u_{t+256}) into ONE float acc before the block reduce; at 512 threads those units live in different threads and their sum happens in the reduce tree (16-warp cross-warp serial order) — a different float sum. A bitwise emulation (smem pair-exchange so thread t still sums u_t+u_{t+256} first) adds a barrier for zero resident-parallelism gain (5120 rows = 18 waves either way), and the exposure mechanism 1c targets was already fixed by v2_pf's up-front load issue (D3b-1b)
D3-6 2aGQA q-head batching in the decode split attention (grid (ATTN_SPLITS, n_kv_heads), 32*gqa threads, warp w = q head hk*gqa+w, full split stripe per warp, shared h4w_warp_windows window pass, per-warp 8/16-butterfly epilogue; same rpw ≥ 16 dispatch slot, static grids → replay-safe)reverted (patch /tmp/d3/patch_2a.py)nsys 14B @3254: h4w 63.76 → batched 64.79 µs (+1.6%); ncu lts__t_sectors: 2,121,671 (4.63× analytic 1×) → 472,583 (1.03×) — traffic ÷4.5 with time flat−1.6% kernel—REVERTEDmechanism-nailed: the 5× L2 re-read is fully HIDDEN under the latency roofline in the live regime (D1's latency-bound attribution stands; D3-4's L2-composition re-attribution revised) — ncu's −28% appears only serialized/cold; ALL gates were green first: parity ≤1e-4 at gqa 5/7 incl. exact nkv 2808 + outlier calibration (shapes kept as permanent h4w coverage), dump argmax HARD gate (margins 2.187/0.557), greedy byte-identical with repeat-penalty 1.0 both models, default-penalty flips = 1/256 sampler knife-edges (14B step-63 raw top-2 probgap 0.0167; 7B step-8 penalized rank-6 winner), temp-0.8 controls identical, suite 170/0/3; new gate rule: attribute greedy flips to sampler vs kernel via the penalty-free stream + per-step logits_top trace
D3-7 2battn_v-q6K decode MMVQ routing: Q6_K decode dispatch gate lowered od*id >= 24M → >= 4M (the only affected shape in the supported set is the 14B attn_v, od 1024 × id 5120 = 5.24M, 11 layers; GGUF census; 7B attn_v od 512 × id 3584 = 1.8M stays padded-f32)this commitnsys 14B @3254: attn_v kernel 33.16 → 24.32 µs (−26.7%, ~177 GB/s — short of the 220 class, as the 8e small-shape crossover data warned for od≈1024 but still −26.7%); ×11 layers ≈ 97 µs/step; quantize launch count UNCHANGED (attn_v joins attn_o's MmqCache hit — same src buffer + id)+0.42% (14B @3254, 8-pair median)see D3-7LANDEDTolerance-gated per the D3a package: logits_decode max|Δ| 0.254 (calibrated 0.39-class), logits_prefill byte-identical, argmax HARD gate green (margin 1.915), kv1+ decode-side f16-boundary drift (kv0 bitwise — earlier onset than D3-6's sub-ULP class, expected for input-quantization noise); penalty-free (rp=1.0) greedy streams byte-identical both models (the D3-6 clean kernel gate); default-penalty flips = 1 knife-edge event/256 steps (5/5 seeds, coherent text, no degeneracy); temp-0.8 controls: 7B identical, 14B reorders (sampled reordering expected at 0.22-logit drift); wall: @3254 21.605 → 21.72 (+0.42%, 7/7 clean pairs positive, sign-test p≈0.008; strict SEP missed by 0.05% — min-new-excl-outlier 21.66 vs max-base 21.67 — medians carried per the D3-5 outlier precedent); tg128 clean-window +0.26% (sub-bar; the extension window was co-tenant-contaminated post-side)
D3-7 2crms/elementwise-launch consolidation, two bitwise sub-levers: (i) rms_norm_quant_pad40 wide-block geometry (launch 32 → 128 threads; the reduction keeps lanes 0..31 exactly — same element→lane map, serial per-lane chains, warp_reduce_sum tree; scale broadcasts via smem; write/quantize loops are element/per-32-block independent so their wider mapping cannot move a bit; reduce loop #pragma unroll 8 deepens load pipelining), (ii) positions_i32 one-execution-window memo (every Rope/KvcacheStore/Attn node re-converted the same positions buffer: 240 launches/step at 14B; key (buf id, pool_gen), cleared in synchronize next to the MmqCache clear; capture-safe: only the first consumer's conversion is recorded and replay re-executes it)this commitnsys 14B @3254 decode census: rms_norm_quant_pad40 9.43 → 5.66 µs (−40%), 94.6–96/step; f32_bits_to_i32 239.6 → 1.2 launches/step; wall-effective ≈ −0.62 ms/step (rms −0.348 + bits −0.275)+1.76% (14B @3254 cumulative with 2b, SEP)see D3-7LANDEDBitwise end-to-end: dump gate (both sides under MINFER_NO_KQ_MMVQ=1) 109/114 files byte-identical, the 5 diffs = the documented slot-aliasing class; logits both phases + all KV byte-identical; 7B greedy streams byte-identical 5/5 seeds + temp-0.8 control; suite green. GATE GOTCHA recorded: MINFER_NO_KQ_MMVQ=1 also reverts the Q5_K decode arm (pre-existing), so a 2b-off control must set it on BOTH sides — a one-sided control shows a fake 0.22-logit drift from the Q5_K f32-activation fallback
D3-8FusedQKV decode fusion ported to CUDA (Stage-3 Tier A; the G4 Metal fusion): (1) attn_bias_rope_store_f32 kernel + launch_attn_bias_rope_store — one launch replaces the per-layer add_bias×3 + rope×2 + store_kv×2 chain, pointer-form (serves concat sections q=base/k=base+nqt/v=base+2nkt AND three separate buffers), positions read device-side (positions[0]) so the launch is capture-safe, math verbatim add_bias_f32+rope_f32+store_kv_f32/f16; (2) class 1 (wq|wk|wv same quant type): Op::FusedQKV — one concat matmul (blk.{i}.attn_qkv loader-registered wq|wk|wv rows) + the fused epilogue (MMVQ is per-row, dispatch on (ttype,id,nt) only → concat bitwise-equal to 3 separate matmuls, probe-proven); (3) class 2 (mixed quant, e.g. Q6_K attn_v among Q4_K q/k — 24/48 layers at 14B, 14/28 at 7B): new Op::QkvBiasRopeStore — the three SEPARATE matmuls (bias-free) + one epilogue launch (CUDA-only; Metal keeps the unfused chain for these layers), builder wires attention to the epilogue node so q's matmul buffer has exactly one consumer and the §5 in-place alias applies; gated by nt==1 && gpu && fuse_qkv (part of the reuse identity — MINFER_NO_FUSE_QKV=1 reverts both classes for A/B)this commitnsys 14B @3254: total launches −4968/trace (−22.5%); per decode step −310 (add_bias −144, rope −96, store_kv −96, fused +48, mmvq −48 (24 concat layers 3→1), quantize +24) ≈ the D3-7 §2-listed 0.45 ms/step qkv-chain item+1.63% (14B @3254; tg128 +3.11%; 7B +1.23%/+1.05%; isolation A/B post-vs-post NO_FUSE_QKV: +2.28%/+3.15%/+1.00%/+1.03% — all 3/3 pairs clean-separated)see D3-8LANDEDBitwise: probe tests (epilogue vs the 7-launch chain bitwise on q/k/v sections + KV rows, f32+f16 KV, both pointer forms, 14B+7B shapes; concat matmul vs 3 separate bitwise) + dump gate logits_prefill/decode + ALL kv*.f32 byte-identical both models (98/98 + 58/58; the informational node* dumps are a documented instrument limitation — recycled pool slots, binary-layout-dependent) + greedy −n 256 byte-identical 5/5 seeds × both models + rp=1.0 + MINFER_NO_FUSE_QKV=1 control + temp-0.8 controls (the only diffs are the perf-banner tok/s numbers); prefill DOT graph byte-identical (prefill untouched); suite 172/0/3 (D3-7's 170 + 2 probes); prefill graph topology unchanged (nt>1 gate) so prefill perf untouched (pp3254 1833 t/s pre-vs-post)
D4-2Decode-GEMM Tier B session (design /tmp/d4/D4_DESIGN.md): (B0) correctness — v2_pf dispatch bounded to npair ≤ 512 (see the D4-2 chapter; D3b-1b's 7B gains were mostly dropped units); (A) llama L2-prefetch port closed PRE-BUILD: prefetch distance 2·bpi requires bpr > 32 blocks (QI4_K=16/VDR=2, QI6_K=8/VDR=1 → bpi = 4·nwarps) — at 14B only ffn_down (bpr 54) qualifies, and our kernels map one 64-elem unit per thread over 256 threads → exactly ONE K-loop iteration at every decode shape (npair 80/216/160; v2_pf's 432 are unrolled u0/u1): there is no "2 iterations ahead" to prefetch, and where llama's prefetch does fire its ffn_down-q4K runs 224.1 GB/s vs our 228.6; (B1) __launch_bounds__(256,6) on v2_pf: 48→40 regs + 40 B stack spill, probe tg128 +0.25% / @3254 −0.75% → killed; (B1c) v2-loop at npair 432 via MINFER_Q6K_PF=0: v2_pf wins/ties (the 5-block × 2-unit-MLP form beats 6-block × 1-unit) → default kept, env kept as opt-out; (B2) 160-thread v2 right-size (= D3b-1c repeat, re-measured with per-kernel isolation): bitwise 98/98 but nsys lm_head +1.51%, attn_v +3.04% → killedb31084c (B0) + docs commitper-kernel nsys deltas above; walls ≈ 0 as expected for a fix-only treecorrectness fix; perf-neutral—see D4-2Dump gates 107/107 (14B pre-vs-post) + 98/98 (B2 bitwise check); greedy byte-identical 14B pre-vs-post; 7B fixed-vs-v1 first-step logits at the v1-vs-v2 rounding class (max|Δ| 0.254, argmax same) vs 4.72-4.79 pre-bug; suite 173/0/3
D4-3Attention-structure attempt 2 (probe /tmp/d4/probe_attn2.cu, 690-row sweep, 0 skips): llama-fattn-geometry split-attention kernel vec_attn over pb/T/R/minb/STG (load-scheduling axis); NO-GO per the pre-registered bar — best 41.07 µs kernel-total @14B (bar ≤~32; 1.64× vs current 67.4) and 16.22 @7B@1641 (bar ≤~10.35; 1.53×) → projected wall +1.6–1.8% < the +2% bar → no integration. Headline: the D4-1 llama attention target (14.02 µs/layer @14B/@3254) is a llama-bench artifact — ncu: the bench decode fattn-vec (grid (1,2,40)) loads a constant 5,427,200 B ≈ one 128-row KV iteration per block = 256 of 3255 rows covered, byte-identical at KV 1024/2474/3255, while llama-cli's decode (grid (1,7,40)) loads 53.2/142.7 MB scaling with context (mid-prompt recall A/B confirms). Honest llama full-context decode attention ≈ 2.0–2.1 TB/s ≈ 1.7 ms/step @14B — minfer's 3.32 ms is ~1.9× off, not 4.9×; the honest @3254 gap is ~10% wall (~3.5% attention)docs commitsweep table in the D4-3 chapterline closed (measurement-corrected)—see D4-3probe gate 0.05 abs w/ adversarial outliers, CPU ref in double; SASS-level LDG counts + recall A/B + reductio (9.6 TB/s impossible) all consistent
D4-4Final decode-kernel session (three levers, /tmp/d4/probe_l1_dpl.cu + /tmp/d4/probe_l3_fuse.cu): (L1) dense split-plane (dpl) q6_K decode MMVQ LANDED — the padded 256-elem/224B row layout streams 14 dead bytes per super-block (215/256 useful = 84%); dpl repacks to [ql: nbe×128][qh: nbe×64][sc: nbe×16][d: nbe×2] = 210B content/row at row stride (nbe·210+15)&~15 (16B-aligned uint4, zero pad sectors), same per-unit values + accumulation order → bitwise; sibling plane (+2.0 GB 14B, +0.9 GB 7B) under MINFER_Q6K_DPL ("0" opt-out), only id % 256 == 0 shapes, padded plane retained (prefill MMQ block_stride 224 + W_exp/W_dsc derivation + dequant/embed fallback). Probe: ffn_down 176.5→212.9 GB/s content (−17.1%), lm_head 208.3→250.0 (−16.7%); the group-split (gs) probe variant (−7.9/−10.7%, tolerance) dominated → dropped. (L2) PDL on the decode chain NO-GO in-situ: standalone probe green (graph capture+instantiate with cudaLaunchAttributeProgrammaticStreamSerialization OK on driver 580.173.02, 200 replays stable, PDL-graph vs plain-eager bitwise; but a compute-bound chain probe ran +2.8% slower — co-residency tax warning), full integration (PSS launch attribute + cudaGridDependencySynchronize() on 13 decode-chain kernels, MINFER_PDL gate) passed all bitwise gates, then the same-binary env-flip isolation A/B read 14B tg128 −2.6%/−1.8%, 7B ≈ 0%, @3254 within noise → below the +0.3% bar → reverted; mechanism: PSS early-launch co-residency taxes the compute-tail kernels (attention h4w, lm_head) more than the ~2 µs/launch graph-gap pool it recovers. (L3) fused gate+up+SwiGLU+q8 (Form B) NO-GO: 32-row-block fused q4_K kernel (grid nf/32 = 432 blocks, 64 serial per-row dots each, in-kernel silu + quantize_pad40_block) measured +28.2% vs the incumbent gu-matmul + swiglu_quant pair (412.1 vs 321.4 µs at 27648×5120; 193.2 vs 247.8 GB/s content) — grid = 1.5 waves at 6 blocks/SM (wave quantization) + exposed per-row latency; Form A (fused gu+swiglu f32, separate quantize) saves ~2–5 µs/step by arithmetic = sub-bar, not probedcode commit + docs commitwall deltas this row (3× interleaved same-window A/B, medians of 3)14B tg128 23.28→24.57 (+5.53%) / @3254 22.00→22.94 (+4.27%); 7B tg128 47.55→51.20 (+7.68%) / @1641 46.44→50.18 (+8.05%) — +5.53%/+4.27% (14B)vs-llama: 14B tg128 1.018×, @3254 0.950×; 7B tg128 1.074×, @1641 1.052× (llama 24.14/24.14/47.65/47.69)LANDED (L1); L2/L3 closed with mechanismL1: dpl kernels bitwise vs padded (probe + unit test both forms: od 512/id 8960 pf-form, od 4096/id 1024 loop-form); 14B −n 1 dumps 107 identical + 7 node{N} diffs (= the documented D4-2 pool-slot aliasing class), 7B 72 + 2; greedy rp=1.0 byte-identical both models; suite 174/0/3 (incl. the new dpl bitwise test; one earlier full-suite run flaked 2 pool/parity tests under the sglang co-tenant window — both pass in isolation and the rerun is green)
D5-0Speculative-decoding cost model (gate, no engine change): measured per-token costs (7B q4_k_m CUDA 54.3 tok/s; 0.5B q4_0 CUDA 342.2 / CPU 73.3; q4_k_m 365.0; q5_k_m 321.8) + real greedy acceptance via llama.cpp speculative-simple (0.5B-on-7B: aggregate 58.8/42.6/25.0% at d=2/4/8 → conditional p ≈ 0.68–0.70, stable across prose/code and draft quant) → break-even requires p* = 0.73/0.81/0.90 at d=2/4/8 → conditional go at d=2 only, gate = minfer measured nt=3 verify-batch amortization ≥ 2.5× (D4-1 anchor 2.7× at nt=4; interpolation 2.28× vs tile-step 2.7× disagree — D5-1 re-ordered primitive-first to measure); projected 1.04–1.05× at the anchor, ceiling ~1.2×; CPU-draft cross-device dead (1.35×)docs commitbattery: 3× interleaved -p 0 -n 128 medians; acceptance n=256 greedy, prose+code——MEAS-ONLY (gate open)measure the gate variable with someone else's binary; the cross-device fallback died by measurement not argument; interpolation is not measurement — nt=3 lands either side of the 2.5× line and only C_T(3) arbitrates
D5-1aThe gate measured end-to-end — FAILED, D5 closed: new minfer specverify instrument (src/spec_verify.rs; fixed-depth protocol, warmups absorb the graph rebuild, two-pass drift check) drives the generic forward_graph_cached at nt>1/n_out=nt; 7B q4_k_m CUDA @KV512: C_T(1)=18.33 ms, C_T(3)=106.0 ms → per-token amortization 0.52× vs ≥2.5× (needed C_T(3) ≤ 22.1 ms). Full curve: nt=2–8 costs a flat ~35 ms/token (batched path re-streams weights per row — 1.9× worse per token than the nt=1 MMVQ path; zero amortization anywhere), the real tile-regime step sits at M≥16 (nt=16 = 56.7 ms total, 3.1× the weight floor; nt=64 = 64.9 ms) — unreachable for verify (nt=d+1 ≤ 9), and even padding nt=3→16 caps at 0.97×. Five probes eliminated alternatives (graph-launch asymmetry, KV depth, FA kernel, n_out path, drift). The D4-1 2.7×@nt=4 anchor was a kernel micro-bench that never existed at graph levelthis commit + docs addendummedians of 15–40 reps × 2 passes, ±2% spread; probes + llama-cli battery in doc 81gate FAIL 4.8× (0.52× vs 2.5×); external check: llama-cli -md same pair lands 0.99–1.00× (doc 81 §4.1) — VOID, see doc 81 §4.3 errata: the battery never engaged the draft (missing --spec-type); corrected = 1.53–2.43×—see D5-1aLANDED (instrument) · D5 CLOSED per the pre-registered stop rule
82small-M dispatch fix: multi-token MMVQ (K-quants, nt 2–8, in-block token loop) + token-looped legacy kernels (grid.y=nt removed) + GEMM gate 16→9this commit7B batched decode nt=3 105.9 → 29.4 ms (3.60×), nt=8 279.3 → 48.6 ms (5.75×); marginal 34.4 → 4.3 ms/token; nt=1 paths bitwise-unchanged (tg128/pp512 clean)——LANDEDthe pre-registered 2.5× amortization bar was mis-derived (marginal ≈ nt×(attention+compute), not ε) — measured 1.87×; weight traffic is nt-independent now, D5 stays closed (C_T(3)=29.4 > 22.1); small models gain ~1.0× only (per-layer weights already L2-buffered)
D5-R ①speculative loop LANDED (reopen per doc 81 §4.3 + doc 82): src/spec.rs SpecEngine — d×nt==1 draft forwards + one nt=d+1 verify through forward_graph_cached, lazy accept loop (unit-tested), full-accept draft-KV repair, CLI --spec-draft/--spec-draft-n; two-model process fixes: namespaced GPU weight registries (load_model_ns), nb_bt_only global-mix semantics (q4_0 draft degrades mode-2→mode-1)this commit + doc 8314B+0.5B q4_0 d=2 greedy n=128: 34.0/40.6 tok/s = 1.34×/1.58× vs serial 25.5/25.5; 7B 1.18×; suite 179 greenG1 re-scoped by measurement: batched-verify vs decode kernels differ ~0.01–0.05 logits → near-tie flaps only (first-flap margin 0.043); d=0 fallback == non-spec to one exact-tie flap—LANDED"single-model-per-process" was load-bearing in three places (registry, dispatch flags, prewarm); all-or-nothing CUDA failure is silent CPU — watch tok/s, not errors; greedy equivalence across kernel paths is a numerics statement, not a logic statement
D5-R ②same-window dual-engine battery: 3 interleaved reps × prose/code × 4 cells (minfer off/d2 × llama base/d2), doc 81 §4.3 protocolthis commit + doc 84minfer 1.33×/1.59× (34.0/40.6) vs llama 1.64×/2.08× (38.8/49.0) in one window; minfer = 88%/83% of llama's absolute spec speed; gate ≥1.2× PASS——LANDEDsame-window interleaving beats rep count (llama's spec cell moved 13% between windows, base <1%); the whole gap to llama is the verify row marginal (minfer 8.8 vs llama 2.5 ms/row) — closing it prices at 1.62×; acceptance prose 51.6% (near-tie dilution) / code 73.1%

Footnotes.

  1. r59 correction (visible in-table). The r59 record originally reported +26.2% interleaved under a co-tenant and attributed −12% to a "co-tenant tax". r59b proved the r59 baseline binary was itself the stale r58-delta build (~−12.5% deficient), making the true clean delta +11.1% (3232.0 → 3590.8). The +26.2% and the co-tenant-tax claim are superseded; the corrected value is what the table's Δ column carries.
  2. Measurement contexts. Rows r12–r25 interleaved on a box that drifted session to session (−9% to +38% vs neighbors) — only intra-session deltas are meaningful. r59's interleaved series ran with a 46 GB sglang co-tenant (later shown irrelevant: idle residency taxes nothing, r59b §1). r55's baseline sanity was 3144.4–3151.4 under a live co-tenant vs the quiet-box 3181.
  3. Hash substitutions. The older records quote HEAD/revert anchors that were later amended away (e.g. r22's "HEAD 9819410", r18's "HEAD 1e0dded", r55's "HEAD dd6d842", r58's "HEAD 2105b08", r59's "commit feb37de"). Each has a reachable twin with an identical subject; the table cites the twins. The P5 record's session range start b8568cd does not resolve to any commit — the P5 code commits are 86ca78c, d713e6e, 725e307, fc07c04, 1365c82. llama.cpp-side hashes (ca3d5a3e1 bench build) are upstream identifiers, not minfer commits.
  4. Campaign arc: R1 MMQ 441 tok/s (first parity-clean MMQ measurement, r7–r8 window) → 3590.8 tok/s (r59b definitive) = 8.1×.

§1 Current state (post-D4-4, 2026-09-09)

1.1 Performance summary (DGX Spark GB10, 7B q4_k_m unless noted)

llama.cpp reference: llama-bench @ ca3d5a3e1 (upstream build), 8 threads, -ngl 99.

Config (all default env unless noted)7B q4_k_m whole-prefillMemory (peak, per-PID)Note
Default (promoted MMQ gate set, mode 2)~3581 (r59b definitive median; r60 A/B window 3578.0–3598.7)9484 MiB= the verified 1.080× path
MINFER_MMQ_Q6K_EXP=0 MINFER_MMQ_Q4K_DSC=0 (planes off)−5%-class (r54: −5.04%; r59-class)6217 MiB (planes cost +3.27 GB)fine-grained opt-out
MINFER_MMQ=0 (legacy f16 w16-cache escape)~2353 clean-class (~2226 in the r60 window)~20.5 GBthe escape is ~11 GB HEAVIER, not lighter
llama-bench pp3314 (r59b window)3323.29 ± 3.08—minfer 3590.8 / 3323.29 = 1.080×

Decode (nt==1) is untouched by the MMQ campaign (r60 evidence: no MINFER_MMQ* read on the nt==1 path; decode -n 16 --greedy byte-identical). Final decode-campaign state (D1→D4-4, 2026-09, all measured on the B0-fixed engine — the D4-2 correction found the v2_pf dispatch dropping units on npair-592 rows and re-anchored every 7B claim):

modeltg128long-ctxvs-llama (same-window)
7B q4_k_m51.20@1641 50.181.074× / 1.052× (ahead)
14B (48L)24.57@3254 22.941.018× tg128 / 0.950× @3254

Small models (pre-MMQ-campaign numbers, Part-I record): 0.6B q8_0 prefill @2K 4792 (llama 23909), decode tg128 ~195 (290); 0.5B q4_0 prefill ~3020 (30550), decode ~257 (453).

The per-session narratives (D2/D3-4/D3-5/D3-7/D3-8 updates), the D4-2/D4-3 correction chain, and the standing measurement rules live in the §0 rows D1–D4-4 and step docs 65–76. Two rules worth surfacing: never quote llama-bench long-ctx tg rates as attention targets without an ncu byte-count or llama-cli recall cross-check (D4-3: the bench decode fattn-vec covers only ~8% of KV rows; honest llama decode attention ≈ 1.7 ms/step, minfer ~1.9× off), and an end-to-end max|Δlogits| ≈ 0.38/0.39 is the inherent class of ANY accumulation-order change — argmax + greedy-divergence + A/B are the operative gates (D3a calibration).

1.2 Wall decomposition (converged regime, r55/r58/r59-era records)

Production nsys at nt=3314, kernel busy ≈ 984 ms (r58 census, pre-r59):

SliceSharevs-llamaState
q4_K BT GEMM (mmq_raw_nb_bt)63.2% (622.4 ms; r59: −30.9% → 526.8 ms)1.06×closed absent a llama.cpp-style q8_1 GEMM-prologue rewrite
q6_K BT GEMM (mmq_raw_nb_bt_q6k)15.4% (125.5–148.8 ms)~1.1× kernel-sideclosed (r53/r56: both staging planes precomputed, all stagings async)
fused A-producers (swiglu+quant, rms+quant)8.8% (swiglu 64.0 ms)swiglu at 89% of the 273 GB/s DRAM rooflineswiglu closed (cap +0.74%); rms at 56% roofline = last incremental lead (ideal +0.98%)
FA prefill attention5.3% (51.8 ms)2.43× taken in r48 (5.16 → 2.12 ms)closed for tile levers (r46/r50/r57)
elementwise/rope/kv/store~6%—bandwidth-bound
standalone wo quantize + mmvq tail~2%—tail section runs the MMVQ/native path
host/launch gaps~1%—one-time stalls 0.6% + recurring gaps ~0.1% + tail malloc 0.1%

Campaign verdict (r55, re-validated by r59/r60): no identified lever ≥ +1.5% remains within the current architecture; the next meaningful step is the step-function q8_1 GEMM-prologue fusion, not incremental optimization.

1.3 What the engine does now (dispatch shape)

All in src/cuda/kernels/*.cu + src/cuda.rs, dispatched by src/graph/cuda_backend.rs:

  • Weights resident at load: every matmul weight uploaded once and registered by name; q6_K optionally 224-byte-padded (7e②). Registration additionally builds the q6_K W_exp (1.52 GB) and q4_K/q6_K W_dsc (363.2 MB + 1456 MB) planes when their gates are on (Appendix A).
  • Prefill (nt ≥ 16): the promoted MMQ path — quantize_q8_0_pad40_t pre-transposed A planes (r34), fused producers (r51/r52), NB-BT raw-byte int8-tensor-core GEMMs for q4_K (mmq_raw_nb_bt_kernel, 2 blocks/SM) and q6_K (mmq_raw_nb_bt_q6k_kernel, KSPLIT=2, 3 blocks/SM, cp.async A/B/dsc staging), FA-style tiled prefill attention (8n, FAP2 register softmax). Fallbacks compile in and are byte-identical (EXP=false, DSC=false, mode-1 producers, generic mmq_nt/f16 arms).
  • Decode (nt == 1): per-type MMVQ over q8_0 activations (8e/8e② + R2 v2), fused bias+rope+store, whole-step CUDA-graph capture/replay (7d); pinned D2H logits readback (R3-A2). Repeated identical-nt prefills capture after the 3-run protocol (R3-B); one-shot prefills never capture (r55: measured no-win + capture-illegal mid-window malloc).
  • Memory etiquette (shared box): no raw allocation probes; check free -g before suite runs (the suite transiently reserves up to ~100 GB of the overcommitted pool); benches stay at single-process 7B scale while sglang serves.

1.4 Remaining roadmap (post-campaign)

  • D5 speculative decoding — CLOSED 2026-09-10, REOPENED as D5-R 2026-09-12 (docs 81 §4.3, 82, 83, 84): the original closure's bar was mis-derived and its llama-cli anchor was a measurement artifact (--spec-type silently defaults to none); doc 82's small-M dispatch fix made the verify amortization 2.14×. D5-R stage ① landed the loop (1.34×/1.58× at 14B d=2), stage ② priced the gap to llama (verify row marginal 8.8 vs 2.5 ms/row → 1.62× recoverable). Current plan: SPECULATIVE-DECODING-PLAN.md (the old plan is an appendix there).
  • Not planned (revisit with a concrete need): cuBLAS/cublasLt (closed as 8k — 8m's wmma GEMM covered the f16 path), VMM pool, multi-GPU, node reordering, Windows, IQ/Q2/Q3 quants.
  • Open leads, all sub-bar or step-function: q8_1 GEMM-prologue fusion (the step change); rms_nw roofline (+0.5–1%); wave re-tile for small-od classes (+0.3–0.8%, needs ≤85 regs); fused ffn_gu concat (needs the G5 nf≤16384 gate re-measured); FA deep-opt only with numerics-order-preserving structure (r50/r57 caveat).
  • Open Phase-8 ledger items (inherited from the retired CUDA-FOLLOWUP-PLAN.md; records in step docs 78/79): 8e② follow-up — llama.cpp's shape-dependent halve_iters idle-tail rule, not started; 8h② — self-hosted CUDA CI runner, DEFERRED (needs standing runner infrastructure); 8h③ — the Phase-7 /tmp/minfer_phase7/ ledger cleanup, awaiting user decision.
  • Closed Phase-8 ledger item: 8a① macOS Metal regression run (fuse_ffn decoupling + the MINFER_NO_FUSE_FFN A/B on 0.5B + 7B) — was BLOCKED on hardware; DONE 2026-09-10 on an Apple M4 Pro, all three checks green (pre/post greedy byte-identity, fused-vs-unfused byte-identity, 0.5B decode graph still emitting 24 × fused_ffn on Metal). Record: doc 78 §3.2.

§2 Step documents — one doc per history row

The full per-step chapters (process narrative, principle explanations, real code excerpts, verification gates, lessons — previously inlined here) now live in docs/cuda_optimization_steps/ as 109 standalone documents (01–109, including the Phase-8 supplementary records 78–79 and the verification-methodology capstone 77). The §0 master table above remains the one-row-per-step index; the tables below link each row to its step document. Appendix B points at the cross-cutting methodology.

Part I · Era A — Phase 7/8 foundations (rows 1–6 + 78–79)

#doc
01Phase 7 — CUDA backend: raw FFI device layer + graph backend 7a–7e (LANDED)
028m/8m② — Tiled wmma f16 prefill GEMM (LANDED)
038n — FA-style tiled prefill attention (LANDED)
048o — Killing the CPU stall at decode start (LANDED)
058p — Persistent f16 weight cache + fused dequant-in-GEMM (LANDED)
068e/8e② — decode MMVQ: dp4a integer dot products + the llama.cpp launch table (LANDED)
78Phase-8 correctness & engineering-debt batch — 8a/8h①/8i (LANDED; 8a① closed 2026-09-10)
79Phase-8 coverage & first-measurement batch — 8b/8c/8d/8f/8l/8q (LANDED)

Part II · Era B — R and P5 sessions (rows 7–13)

#doc
07R3 — Small-model per-token overhead: single-split prefill (LANDED)
08R1 — int8 MMQ prefill GEMM (opt-in): the parity-first strategy (LANDED)
09R2 — MMVQ weight-streaming rework (LANDED)
10R4 — decode split-attention dim-parallel rewrite (LANDED)
11P5 — prefill gap session: TM=128 big tiles + FA rewrite (LANDED, with three REVERTED probes)

Part III · Era C — the P6 q4_K MMQ line, r5–r37 (rows 14–51)

#doc
12r5–r6 re-ranking + structural rewrite spec (MEAS-ONLY + REVERTED)
13r7–r8 raw-byte MMQ kernel + wide tile + FA probe (LANDED (raw) + REVERTED (probe))
14r9 llama.cpp MMQ reference decode; shape axis closed (MEAS-ONLY)
15r10–r11 — reference inner-loop decomposition port; ILP reading verification (REVERTED / MEAS-ONLY)
16r12 — 16-chain warp tile + ldmatrix (LANDED)
17x-tile / j-tile / cp.async-db — the staging-shape family, closed (REVERTED / closed)
18r13 — counter forensics against llama.cpp (MEAS-ONLY, closed)
19r14 — B fragments via ldmatrix + widened scale reads (LANDED)
20r15 — f32-accumulate mma probe (dead end) + rank-1 term2 rescale (LANDED)
21r16 — Narrow kernel gets the rank-1 fold (LANDED)
22r17 — Wide warp remap 32od × 64tok (REVERTED)
23r18: Load-time B pre-expansion — staging becomes a bulk copy (W_exp's debut, REVERTED)
24r19: Weight L2 residency — __ldg imperceptible, persisting window catastrophic (REVERTED)
25r20: Split-phase A staging — attribute first, shoot second: the first hit (LANDED)
26r21 — Coalesced block-linear A staging: stall-mass conservation (REVERTED)
27r22 — qa8 XOR swizzle (LANDED); d/ssum fold reverted separately
28r23 — f16-path whole-graph wall decomposition + FA_TKV occupancy raise (MEAS-ONLY + REVERTED)
29r24 — The scheduling-structure ladder (REVERTED; the +1.5% whole-prefill landing bar calibrated here)
30r25 — SASS opcode census; the unroll is wall-inert (MEAS-ONLY + REVERTED)
31r28 — Direction-A raw-nibble NB kernel, 2 blocks/SM (LANDED)
32r29 — NB kd-loop unroll: 2 blocks/SM lets integer-ALU pruning move the wall clock for the first time (LANDED)
33r30 — SWAR unpack: the compiler already did it (REVERTED)
34r31 — q-major sda scale-read repack: a sub-bar positive gain caught by conflict analysis (LANDED)
35r32 — The finite lever sweep: two regions fenced off (REVERTED)
36r33 — Hybrid inner-loop port: SASS fully identical, hypothesis falsified (REVERTED)
37r34 — The quantize-transpose prepass: layout-transform locality (LANDED, +9.72%)
38r35 — sda scale predecode: a total SASS win, a wall-clock tie (REVERTED)
39r36 — A-frag wavefront economics: H1 falsified (MEAS-ONLY, no code change)
40r37 — Post-parity whole-prefill attribution: the wall clock re-decomposed (MEAS-ONLY, no code change)

Part IV · Era D — q6_K, FA, prepass, promotion, r38–r60 (rows 52–75)

#doc
41r38 — q6_K BT-style raw-byte mma kernel (LANDED)
42r39 — q6_K KDR=2 double-buffer (LANDED)
43r40 — __launch_bounds__(256,3) third resident block (LANDED)
44r41 — q6_K B-expand widened to uint4 group loads (LANDED)
45r42 — q6_K stage-wide dsc scale reads (REVERTED)
46r43 — PC-sampling attribution + pre-expand-B parity FAIL (MEAS-ONLY + REVERTED)
47r44 — W_exp stride mismatch root cause: fix goes parity all-green but wall-neutral (REVERTED)
48r45 — cp.async for the q6_K A-side staging: mechanism confirmed, wall-neutral (REVERTED)
49r46 (FAP1) — FA audit + occupancy/bank-conflict levers: kernel −11% but wall-neutral (REVERTED)
50r47 — converged-era whole-wall re-decomposition (MEAS-ONLY)
51r48 (FAP2) — register-resident softmax: deleting the S/P smem round trip outright (LANDED)
52r49 — A-quantize prepass shared-A dedup: consecutive-window memoization (LANDED)
53r50 — FA_TKV 32→16 occupancy experiment (REVERTED)
54r51 — producer-fused A-quantize mode 1 (LANDED)
55r52 — skip-write mode 2: skipping the f32 intermediate write-out (LANDED)
56r53 — q6_K bundle: W_exp pre-expansion plane + cp.async B staging (LANDED)
57r54 — MINFER_MMQ_Q6K_EXP: an exit valve for the 1.52 GB W_exp plane (LANDED)
58r55 — swiglu roofline audit + one-shot prefill CUDA-Graph: both closed on the record (CLOSED)
59r56 — q6_K A-side bundle: A cp.async + W_dsc f32 plane (LANDED)
60r57 — FA KV staging double buffering (REVERTED)
61r58 — q4_K BT spec + cp.async-db2 transplant (REVERTED)
62r59 — q4_K W_dsc plane + riders (LANDED, Δ corrected by r59b)
63r59b — clean re-measurement + baseline-contamination correction (measurement round)
64r60 — the coronation: flipping the verified gate set to default-on (PROMOTION, LANDED)

Part V · The decode campaign, D1→D4-4 (rows 76–87)

#doc
65D1 decode attribution: split-attention staging depth is the only wall that grows with KV (measurement round, CLOSED)
66D2: explicit K+V register staging (LANDED, +2.0% @1641) and the cp.async negative result
67D3: 14B decode attribution (D3-1) + the D3b bitwise MMVQ triple (1b LANDED; 1a/1c REVERTED)
68D3a — the 4-warp fattn-vec-style split-attention rewrite (REVERTED) + tolerance-gate calibration
69D3-4 — L1 hybrid rpw dispatch (LANDED) + L2 window-prefetch pipelining (REVERTED) + long-prompt dump calibration
70D3-5 — decode-alignment plan Stage 1: fused-producer decode A-quantize (LANDED) + negative analysis of the output-head/ffn_down geometry levers (1b/1c)
71D3-6 — GQA q-head batched attention: all gates green, still reverted — the 5× L2 re-read was not the residual (REVERTED)
72D3-7 — attn_v-q6K MMVQ routing (2b) + rms wide-block / positions memo (2c): the Stage-2 closing ledger (LANDED ×2)
73D3-8 — G4 FusedQKV ported to CUDA: both layer classes covered, 14B short-KV breaks through parity (LANDED)
74D4-2 — B0 latent correctness fix + all bitwise occupancy/prefetch axes closed (LANDED)
75D4-3 — attention structure rewrite attempt 2 NO-GO + D4-1's llama target was a llama-bench artifact (CLOSED)
76D4-4 — the endgame kernel session: dpl dense split-plane q6_K decode MMVQ lands (bitwise); PDL and fused-FFN closed with mechanism (LANDED)
80D5-0 — speculative-decoding cost model: measured baseline, acceptance, and the d=2 gate (MEAS-ONLY)
81D5-1a — the verify-batch gate measured end-to-end: no amortization at any nt, D5 closed (LANDED · gate FAIL)
82small-M dispatch fix — multi-token MMVQ + token-looped legacy kernels: the batching invariant restored, D5 verdict unchanged (LANDED)
83D5-R stage 1 — speculative decode loop: two-model process fixes, accept-loop unit tests, greedy-identity investigation (LANDED)
84D5-R stage 2 — same-window dual-engine battery vs llama.cpp: 1.33×/1.59× vs 1.64×/2.08×, gap = verify row marginal (LANDED)
D5-R ③verify marginal priced with an nsys per-kernel ledger (specverify per-nt runs, exact forward spans)
D5-R ④aattention for the verify shapes: fa_prefill gate nt≥64 → nt≥2 (one line; hd==128, kill switch, rc fallback kept; nt=1 keeps split-KV)
D5-R ④bmulti-MMVQ nt 9–16: token-groups-of-8 (parity 85.0 vs 84.9 — group re-streams weights from DRAM) and acc[16] single-pass (111 ms, register spill) both measured; doc-82 GEMM boundary at nt ≥ 9 stands, kernels keep the group structure (bitwise at nt ≤ 8)
D5-R ⑤final dual-engine battery + capture A/B: the doc-85 capture prize was already banked by R3-B prefill capture (48.08 captured vs 50.15 eager)
D5-R ⑤+row-marginal localization (no ncu): cold-L2 real-kernel bench + chain nsys + ablation
D5-R ⑤++doc-89 menu item 1 (R-rows-per-block) implemented → measured → reverted
D5-R ⑤+++mma path audit: BT kernel already tensor-core; small-M floor = block starvation (40 blocks < SMs); conditional double-buffer shipped
D5-R ⑤++++K-split shipped for both BT kernels (grid.z + deterministic reduce), gated; flip condition C_T(9)≤55 NOT met (72.9)
D5-R ⑤+++++auto-ksplit enabled on the DEFAULT path (user decision, doc 92 §3b)
D5-R follow-updraft-quant sweep (mixed knob: q4_k_m code +3.0%/prose −2%) + greedy identity test: spec output ≠ sequential — flips originate in batched verify attention/softmax, nt-invariance campaign proposed
D5-R ⑥++++++greedy identity ACHIEVED: spec output byte-identical to sequential (4-prompt battery) — batched verify attention (bitwise position-invariant, nkv<1921) + spec penalty window capped to repeat_last_n; cost ≈1% C_T(3); d=8 door CROSSES with q4_k_m draft (code 76.5% ≥ 0.755: 46.6 tok/s vs d=2 43.6, 96% of llama; prose stays d=2)
D5-R ⑦adaptive draft depth (--spec-draft-adaptive): per-round d from beta-smoothed per-depth acceptance + min-of-4 verify/draft costs, 10% switch hysteresis, unobserved depths inherit the depth-1 rate (emergent exploration); identity boundary discovered and pinned — verify nt ≤ 8 (single+multi MMVQ) bitwise vs decode, nt=9 (BT-MMQ) lm_head is tolerance-class while its KV stays bitwise (48-layer dump proof) → adaptive cap d=7
D5-R ⑦P0nt=9 verify profiled (Phase 0, pre-registered stop-gate): nsys — BT-MMQ kernels = 84% of GPU time (attention ~1%, ksplit reduce 1.3%, quantize 1.7%); ncu (sudo; ERR_NVGPUCTRPERM workaround) — both BT kernels at ~15-24% SM / ~15-26% memory throughput, smem scoreboard stall = 40–59% of warp cycles → the <30% stop-gate says PROCEED; Phase 1 = cp.async double-buffered staging (bitwise-preserving), EV narrow (prefill lever / identity-relaxed d=8 — adaptive d≈3.5 already beats d=8 identity-safely) → re-priced separately below 战役 97
D5-R ⑦✦speculative decoding wired into ALL frontends: --cnv --spec-draft (Engine trait hooks + SpecAwareEngine + a spec sibling decode loop mirroring the plain one token-for-token) and serve --spec-draft (per-slot draft engines, seed carry across rounds, mid-batch stop/EOG termination); pre-existing server bug fixed in passing — slot GraphCache reuse across different-prompt requests leaked stale KV rows into the new attention window (the plain path was contaminated too; identical back-to-back requests hid it) → per-request slot-cache + draft reset
D5-R ⑦P196 Phase 1 lever (cp.async double-buffered staging for the q4_K BT kernel) implemented in full — A/B planes + one-group-per-tile (r56 pattern) + the dbuf regime extended to the K-split path — and measured NULL on GB10: q4_K BT 188.5/146.1 µs vs baseline 187.6/142.4, dbuf on/off within noise of each other, C_T ladder and pp512 unchanged; bitwise-preservation gate passed (pre == post == dbuf-off, 4/4 prompts). The 40–59% Short-Scoreboard stall is compute-side smem dependency (ldmatrix→mma operand chains), not tile staging — doc 92's staging attribution corrected. Remaining lever class reorders fp32 accumulation → tolerance-class → excluded on the identity-claimed path → patch reverted, campaign closed; C_T(9) ≈ 73 ms stands as the identity-safe floor
D5-R ⑦P1b99: fast-verify P0/P1 — doc 98's tolerance-class blanket corrected (int mma is exact; only the per-kd fp32 fold carries order → the bitwise-safe set is wider), then four interventions measured: cp.async staging NULL, fragment prefetch NULL, non-volatile mma NULL, B-plane XOR swizzle kept (−4.3% instance / e2e noise; the 40% excessive shared wavefronts eliminated). pc-sampling fixes the root-cause chain: the stall is L1TEX (global) latency at a register-file-capped 16 warps/SM — not smem, not staging, not scheduling. MINFER_FAST_VERIFY not built: tolerance-class levers act on scheduling and are re-priced to single-digit %; weight repack (−27% staging sectors) is the only remaining priced lever (~7–10%)
D5-R ⑦P2100: the last priced lever (32-B-aligned qs plane for the BT B staging, +0.89x q4_K memory) implemented and measured timing-NULL under interleaved A/B (141.9 vs 141.6 ms kernel totals; sequential "−12.4%" was clock-ramp drift). Discovery: dgxspark's short benchmarks carry a ±7–12% SM-clock-ramp band (208 MHz idle → 3 GHz) — sequential comparisons across docs 92–99 all carry it; only interleaved A/B is valid. Traffic levers don't convert in a latency-bound regime → repack family closed; doc 99's swizzle stays on its mechanism evidence with honest error bars
D5-R ⑦P3101: doc 100's method rule made executable — specverify/bench warm on a time budget (default 2000 ms, MINFER_SPECVERIFY_WARMUP_MS / MINFER_BENCH_WARMUP_MS); headline table re-measured in one hot session: pp512 2083.0 ± 7.2 (was ±25), e2e prose 25.3→36.6 (d2) / 35.9 (adaptive), code 25.3→44.8 (adaptive, +7.7% over best static); doc-95 structure reproduces, doc-95 absolutes confirmed as drift-band artifacts
D5-R ⑦P4102: draft-scale sweep — bigger drafts lose (Qwen3-0.6B Q8_0: 39.7 vs incumbent 44.8 code-adaptive; 1.7B: 34.2; acceptance bounded by target predictability at 76.5% same-family vs 63-67% cross-family) → default draft unchanged. Flushed a latent bug: qwen3 loader registered Q6_K padded weights under the raw GGUF name, silently replacing the target's registry entry under the namespaced draft load and dropping BOTH models to CPU — fixed to reg_name; post-EOS token-text divergence downgraded to a warning (cross-family drafts legitimate; identity re-proven 4/4 vs fresh sequential)
⑧ legacy-quant decode103: the 8e MMVQ structure extended to q4_0/q8_0 decode (nt=1 + multi, NEW CODE ONLY — every landed K-quant kernel/arm byte-identical; size-floored gates id≥2048, q4_0 multi only above the 8c territory id>8192; MINFER_NO_Q40_MMVQ/MINFER_NO_Q80_MMVQ fallbacks). 7B q4_0 tg128 48.1→56.8 (+18.1%), 7B q8_0 26.7→28.8 (+7.9%), 7B q8_0 spec e2e 26.8→68.0 tok/s (+154%) — q8_0's verify batch previously ran the f32-activation kernel and re-read f32 activations per token. Same-file llama.cpp: q4_0 decode 1.045× FASTER, q8_0 closed 0.87×→0.94×; pp512 unchanged (prefill gap 0.21–0.24× is pre-existing, out of scope)
⑧ q8_0 p32 planes104: nsys shows 97% of q8_0 decode GPU time inside the MMVQ kernels; cold-L2 microbench isolates the 34B-stride 2B-load scatter (~2x q4_0's wavefront cost/byte) → p32 split planes (uint4 x2 payload + dense d, traffic unchanged, raw registration untouched, byte-equal outputs, MINFER_NO_Q80_P32 fallback). 7B Q8_0 tg128 → 31.9-32.1 tok/s = llama.cpp parity (0.87x→0.94x→1.00x across docs 103-104); spec e2e 79.0 tok/s = 2.47x sequential; both engines now bounded by the same ~253-266 GB/s DRAM ceiling — decode gains for any quant now need a higher streaming ceiling, not better kernels
T1 device tier tables105 (T-series, plan: DEVICE-ADAPTATION-PLAN.md): cc-keyed tier table + selector (device_tier.rs, pure data) — GB10 row Measured (docs 94-104), consumer rows Adopted from llama.cpp (1200 Blackwell q4_K→5/q5_K→6/q6_K→7, 870 Orin K-quants→1, 890/860 generic-8, 750 Turing mmq=false per ruling #4), GENERIC fallback (mmq = cc≥800); MINFER_DEVICE_TIER forced-key override for soak tests; mmq_active() = resolved tier flag; batch caps tabled + unit-tested but unwired (R8: no destination for the vacated nt range until small-nt BT / field A/B). Encoding fix en passant: runtime cc is 1201 (major*100+minor), stale "1210" comments corrected; table keys use llama.cpp encoding via llama_key() conversion
T2 query formula gates106 (T-series T2): (1) doc-92 auto-ksplit extracted to auto_ksplit(), block target = max(256, 2*SM) — GB10 (48 SMs) keeps the calibrated 256 exactly, larger SM counts scale; (2) BT smem feasibility: the tile demand is a single-source C formula exposed as cuda_mmq_smem_bytes(), init folds demand > per-block-optin → tier_mmq=false (f16 GEMM serves) — R2 fixed (externs queried device 0 unconditionally, now the current device); (3) plane_budget_ok() (free > extra +25%) gates every optional plane upload (p32 pair +100%, q6_K W_exp/W_dsc, q4_K W_dsc) — Orin-Nano-class 8 GB devices self-disable planes, raw paths serve. Tile candidate search + batch-cap activation deferred per plan §6.4/R8 (Orin Nano A/B first); T3 stays later/independent
C4 #144 packed Q8_0107: the packed cache's two missing tuned routes. (1) Packed fused decode epilogue — attn_bias_rope_store_q8_0 gives one thread a whole (head, 32-element K block) and one V block, quantizing with the same q8_0_quantize_block store_kv_q8_0 uses; the Qwen2 builders dropped && !packed, and the fused metas gained row_elems so the allocator sizes a packed region the way the store node does. (2) Packed FA prefill — fa_prefill_f16kv became fa_prefill_kv<CAUSAL,MAP,LAYOUT>; only the staging is layout-dependent (kv8_q8_0 dequantizes each packed block into the same f16 tile), the tensor-core QK^T/softmax/P·V are untouched, and the general layout-tagged kernel stays the documented fallback. (3) A dp4a packed K dot was deliberately NOT taken (numerics change, own accuracy statement) and is filed separately
C4 #186 dp4a packed K dot108: the packed Q8_0 decode K dot accumulates in int (__dp4a) against a per-(head, block)-quantized query and scales by d_q*d_k once per block — K is never converted to float (V still is); MINFER_NO_DP4A_Q8_KV=1 is the same-binary control; the nt > 1 prefill/verify kernel is deliberately not taken (per token and head query, no shared block scale)
C4 #202 packed KV L1 loads109: the packed cell's four s8 K/V quant loads become two u16 loads (34k + 2 + 4m is always 2-byte aligned even when only odd blocks are 4-byte aligned) — no layout, CPU, copy-stride, session-format or map_q8_0_cells change; MINFER_NO_Q8_KV_WIDE=1 is the control

Methodology

§3 Appendices

Appendix A — Env-gate reference (post-r60 semantics)

Promoted gates (r60): absent / any non-"0" value = ON (the verified best path); explicit "0" = opt-out (the pre-r60 default behavior). All reads single-sourced in CudaState::mmq_gate_on.

GateDefault"0" effectIntroduced
MINFER_MMQonwhole MMQ path off → legacy f16 w16-cache prefill (~20.5 GB, ~2353-class)R1 (opt-in), r60 (default-on)
MINFER_MMQ_RAWonraw-byte kernels off → generic mmq_nt armsr7–r8, r60
MINFER_MMQ_RAW_NBonNB kernels off → wide raw kernelr28, r60
MINFER_MMQ_A_TRANSPOSEonpre-transposed prepass off → native quantize + NB kernelr34, r60
MINFER_MMQ_Q6K_NBonq6_K BT kernel off → generic q6_K pathr38, r60
MINFER_MMQ_A_FUSEabsent = mode 2 (skip-write)off (mode 1 = "1": fused producers write plane AND f32; "2" = skip-write override)r51/r52, r60
MINFER_MMQ_Q6K_EXPonq6_K W_exp plane not built (−1.52 GB, ~−5% prefill) → EXP=false r41 pathr54
MINFER_MMQ_Q4K_DSConq4_K/q6_K W_dsc planes not built (−1.46 GB) → in-kernel decoder59
MINFER_Q6K_DPLondense split-plane q6_K decode planes not built (−2.0 GB 14B / −0.9 GB 7B) → padded-224B MMVQ path (bitwise)D4-4

Overrides / debug (opt-in "1"): MINFER_MMQ_RAW_KD (default 8), MINFER_MMQ_RAW_WIDE (wide 128×128 kernel), MINFER_MMQ_RAW_NB_DEBUG (prints the dispatch label — B=W_exp-cp.async / DSC=f32-plane / exp=off / fallback! — the liveness instrument from r53/r54).

Legacy f16-path A/B gates (unchanged): MINFER_NO_PREFILL_GEMM, MINFER_NO_W16CACHE, MINFER_NO_FA_PREFILL, MINFER_NO_CUDA_GRAPH, MINFER_NO_PREFILL_CAPTURE (MINFER_CAPTURE_PREFILL=1 accepted, redundant), MINFER_NO_PINNED_READBACK, MINFER_MMVQ_V1, MINFER_NO_KQ_MMVQ, MINFER_GEMM_TM (64), MINFER_GEMM_K64, MINFER_FUSED_B.

Dispatch guards (unchanged by r60, now protect the default path): MMQ entry nt >= 16 && id % 32 == 0 && !no_prefill_gemm; NB-BT (id / 32) % 8 == 0; plane registration id % 256 == 0 (+ od % 2 == 0 for dsc); fused producers rows >= 16 && dim % 256 == 0; mode-2 auto-degrade under MINFER_GRAPH_DUMP / MINFER_DUMP_DIR / MINFER_TRACE / viz capture; the r60 nb_bt_only flag degrades mode 2 → mode 1 on mixed-quant models.

Post-r60 gates (D3–D5-R and the spec/bench harness) — added after the promotion round, so they are not in the table above:

GateDefaultEffect
MINFER_NO_DECODE_A_FUSEoff1 skips the decode fused-producer A-quantize, restoring the standalone quantize launch (D3-5)
MINFER_NO_Q40_MMVQ / MINFER_NO_Q80_MMVQoff1 forces Q4_0/Q8_0 decode off the MMVQ path (f32 kernels)
MINFER_NO_Q80_P32off1 reverts the q8_0 p32 split planes and their dispatch (doc 104)
MINFER_Q6K_PFon0 disables the q6_K prefetching MMVQ form (D4-2)
MINFER_SMALL_M_GEMMoff1 routes nt 2..8 into the mma BT path (doc 91; measured ~1.7× worse)
MINFER_MMQ_KSPLIT_TARGETmax(256, 2×SM)target resident-block count for the auto-K-split (doc 92; doc 106 SM-parameterized — GB10 keeps 256; an explicit value overrides the formula)
MINFER_DEVICE_TIERoffllama.cpp-style tier key forced onto the selector, e.g. 870 (Orin) / 750 (Turing, MMQ off) / -1 (GENERIC) — soak-test override for the doc-105 device tier tables; the banner logs it as FORCED

Instrument / harness (not backend gates): the specverify instrument reads MINFER_SPECVERIFY_WARMUP_MS / _NTS / _NOUT, spec round tracing uses MINFER_SPEC_DEBUG, bench takes its steady-clock warmup budget from MINFER_BENCH_WARMUP_MS, and MINFER_BENCH_ROW_MARGINAL gates the cold-L2 row-marginal device bench test.

Appendix B — Verification methodology (summary)

The full capstone — the five-gate chain (① parity ×3 via cuda_prefill_mmq / cuda_prefill / cuda_fa_prefill_attention_parity, ② greedy-32 token identity, ③ interleaved A/B medians with the +1.5% whole-prefill bar, ④ the device suite, ⑤ the ncu/nsys/SASS protocol incl. the GB10 metric gaps and the sudo-LD_LIBRARY_PATH gotcha) and the campaign's transferable lessons (baseline anchoring, liveness labels, tile-size vs greedy-identity, the pipeline-value formula, roofline-before-coding, occupancy-before-instructions, SASS-first, stall-mass conservation, layout-transformation locality, mechanism composition, consumer attribution, expiring wall decompositions, phantom results, shared-box memory etiquette) — lives in cuda_optimization_steps/77-verification-methodology.md. Read it before running any A/B on this engine.

Appendix C — Part IV legacy: the pre-Phase-7 roadmap (2026-08-29) and where it ended

Everything below described the deleted imperative path (layer_gpu, forward.rs) on an RTX 4080 Laptop with Qwen2-0.5B Q4_0: prefill 40, decode 20 tok/s (CPU 18/15). Kept for the record; outcomes annotated.

Root cause as diagnosed then: per-op CPU↔GPU ping pong.

CPU path → quantize f32→Q8_0 (CPU) → cudaMemcpy H2D → CUDA kernel → sync → cudaMemcpy D2H → CPU path

~6 PCIe round trips × 24 layers ≈ 144 DMA operations per decode step, 2–7 ms of pure overhead. (Correct for that path; Phase 7's resident-weight graph backend eliminated it structurally.)

Original P0–P5 and actual outcomes:

ItemClaim thenOutcome
P0 full-layer GPU offloadadd Q4_1/Q8_0 kernels, kill 144 DMAs, 3–4×Absorbed by Phase 7 graph backend (weights resident, per-op dispatch, split syncs)
P1 GPU-side activation quantizeGPU q8_0 kernel unused, 1.2×Landed as 8c with a measure-first gate (nt>1 && id≤8192); the q8_0 path LOSES 63% at 7B ffn_down (weight-bound) — a blind wire would have regressed
P2 fused GQA on GPUwire gqa_attn_f32, 1.5×Landed in Phase 7a/7e (gqa_attn_f32_f16kv); prefill attention replaced by FA tiling (8n: 20×); decode attention is 0.12 ms/token at 2K — no longer material
P3 cuBLAS for output projectioncublasSgemm "leverages tensor cores", 2×Closed as 8k (not planned). Two errors: cublasSgemm is FP32 SGEMM — tensor cores require cublasGemmEx with f16/int8; and the need disappeared once 8m's custom wmma GEMM covered large matmuls
P4 tiled quantized matmulllama.cpp MMQ "shared-memory tiling with Stream-K decomposition", 1.5×First judged negative, then REVERSED same-day (8e): the real design is integer __dp4a dots over q8_0 activations with a per-CC launch table — ported in-tree as the decode MMVQ win; "Stream-K" was never part of llama.cpp's MMQ. The prefill int8 version became R1
P5 CUDA graph for launch overheadcapture decode, 1.2×Landed as Phase 7d decode capture/replay (one ~57 µs graph launch per token) + opt-in prefill capture (8g②, default since R3-B). The original "2,000+ launches (…× ~14 heads)" miscounted — heads don't multiply launches; the true figure is ~95 nodes/layer × 28 layers ≈ 2.7K, same order

Implementation order as drawn then:

P0 (layer_gpu) ─→ P1 (GPU quantize) ─→ P2 (GPU attention)
                                      ↘
                                       P3 (cuBLAS) ─→ P4 (tiled MMQ) ─→ P5 (CUDA Graph)

All six landed in some form by 2026-08-31 — none via its original mechanism except P0's idea. Measured budgets recorded at the Part-III era: f16 wmma GEMM ~35 TFLOPS (llama.cpp int8 MMQ ≈ 52 equivalent); MMVQ weight streaming 130–147 GB/s effective vs the 252.7 GB/s read-only probe (93% of the 273 GB/s theoretical); llama.cpp ~197 GB/s on the same decode shape.