mudler · localai-bot · Aug 6, 2026
diff --git a/.agents/NOW.md b/.agents/NOW.md
@@ -22,7 +22,7 @@ checkpoint on `upstream/main` at `59674cf1d`.
 | Kimi-Linear-48B (KDA+NoPE-MLA+MoE) | **e2e RUNS** (bf16-resident §13): 13/13·656. Token gate **NEAR-TIE 106/128** | device GDN/MLA islands; 1.59 tok/s; default OFF |
 | 35B fresh grid | **BOUND** @`1ea26427`: 0.93-1.03x, c16 0.93x. INTAKE + Option A both NEGATIVE | Lever left: prefill glue (#61) |
 | Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` |
-| MXFP4 parity | **c1 1.020 PASS, c2-c8 0.962-0.969, mem 2.63x** (`VT_MARLIN_DENSE` banks +951us marlin). `FUSED-GLUE` W0: glue-into-marlin REFUTED at source; residual FLASH-dominant | flash same-tool audit OWED |
+| MXFP4 parity | **c1 1.020, c2-c8 0.962-0.969, mem 2.63x**. `FLASH-AUDIT` #68: c8 flash +12.5us/call = occupancy/L2-bound; `-use_fast_math` TRIED→REGRESSES, rejected | real lever: flash occupancy/L2; no flip |
 | ROW-SERVE-ASYNC-DENSE-MIRROR | **LANDED+dgx-VERIFIED** (`f9c969ae`): async mirror on classic dense Qwen3; SACRED 184/184 | Residual: sibling scope one-liner |
 | CPU levers (`QUANT-GGUF-CIQ-GEMM`) | Op-dispatch profile DONE: decode **47% threadpool sync**, prefill **~39% paged attn**. **G5 not next** | Parakeet encoder; attn dtype hoist |
 

diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md
@@ -19,6 +19,79 @@ from relative link targets repointed for this file's location.
 
 # Benchmarks
 
+## QUANT-CT-MXFP4-FLASH-AUDIT — `--use_fast_math` on the FA2 TUs REJECTED (measured flash REGRESSION); the flash decode gap vs vLLM is a runtime memory/occupancy effect, not the SASS instruction count (2026-08-06, `row/QUANT-CT-MXFP4-FLASH-AUDIT`, base `origin/main` `4ce9fb74`, GB10 sm_121a, PR #69)
+
+The #67 refutation re-attributed the c2-c8 MXFP4 residual FLASH-dominant and OWED a
+STRUCTURAL-lens audit of why the IDENTICAL-grid `flash_fwd_splitkv` decode kernel runs
+slower on our side. This row runs that audit to a MEASURED verdict.
+
+**W1 — the flash term is REAL and fresh same-tool; on CURRENT main it is SMALLER than #57.**
+`analyze_decode2.py` (dense-marlin-aware) c8 decode-window, both `nsys --cuda-graph-trace=node`,
+modal marlin=144 / gridZ=64 / 231 steady steps:
+
+| build (current-main dense-direct default) | flash MAIN us/step (us/call) | marlin us/step | glue |
+|---|---|---|---|
+| ours (no fast-math) | 6075.7 (**168.8**) | 16203.8 | 776.0 |
+| vLLM oracle | 5628.5 (**156.3**) | 16286.2 | 670.6 |
+
+Gap = **+12.5 us/call (+450 us/step)** — NOT #57's +807. Current-main's leaner marlin/glue
+(16204/776 vs #57's 16512/869) drop flash 178.1→168.8, itself a hint the residual is L2/context
+(the flash kernel's cost moves with its NEIGHBOURS' L2 footprint), not the kernel body.
+
+**W2 lens 1 (cuobjdump, HYPOTHESIS) — vLLM's flash-attn is `--use_fast_math`, ours was not.**
+Kernel-version REFUTED: vLLM v0.25.0 pins vllm-flash-attn @ `2c839c33`, the exact commit we
+vendored — flash SOURCE byte-identical. vLLM ships sm_80 ELF + compute_80 PTX (driver-JIT to
+sm_121). Decode kernel `<128,64,128,4>`: ours 5448 instrs/REG246; +`-use_fast_math` → 4832/REG255
+(= vLLM's 4880/REG255), HMMA(768)/LDSM(408)/LDGSTS(120) byte-identical. The compile lens
+SUGGESTED the +11.6% scalar bloat was the gap.
+
+**W2 lens 2 (MEASURED, the arbiter) — the compile lens was WRONG; `-use_fast_math` REGRESSES
+flash.** Controlled SAME-BUILD nsys A/B (current-main dense-direct default, flash source
+identical, only the flag differs), same-tool c8 decode-window:
+
+| build | flash MAIN us/call | flash instrs / REG |
+|---|---|---|
+| ours **no fast-math** | **168.8** | 5448 / 246 |
+| ours **+`-use_fast_math`** | **189.8** | 4832 / 255 |
+| vLLM | 156.3 | 4880 / 255 |
+
+**`-use_fast_math` makes flash +21 us/call SLOWER (168.8 → 189.8), not faster.** The kernel is
+MEMORY-LATENCY-bound: the instruction count barely gates time, and fast-math's higher register
+count (246→255) LOWERS occupancy, which dominates. Matching vLLM's SASS profile (4832/255 ≈
+4880/255) did NOT close the gap — ours-fastmath is +33 us/call above vLLM at the SAME instr/reg
+profile. This is the STRUCTURAL-lens "same kernel NAME, different resolved cost" rule and the
+"MEASURED not inferred" rule: cuobjdump was necessary but NOT sufficient. Correctness footnote:
+fast-math stayed token-exact (#44 smoke 3/3 + coherent) — a faithful mirror of vLLM's build,
+just not a speed win.
+
+**W2 real mechanism (sudo ncu; OURS characterized, vLLM-side LOST to a box OOM-reboot,
+OWED).** OURS no-fastmath flash (c1): **Achieved Occupancy 8.3%** (register-limited — 4 of 12
+warps/SM at 246 regs; fast-math's 255 regs → even fewer warps = the measured regression),
+**L2 Hit 53%**, Memory/Compute throughput ~1-1.5% (pure latency-bound). Warp stalls: **~38%
+short-scoreboard (MMA waiting on LDSM K/V staged to shared memory) + ~37% CTA-barrier
+(`__syncthreads`)** = ~75% of cycles. The flash decode is OCCUPANCY-STARVED: too few warps to
+hide the LDSM→MMA shared-memory dependency and the barrier. vLLM runs the SAME-SOURCE kernel at
+255 regs (same-or-worse occupancy) yet 156.3 — so its ~12 us/call edge is NOT occupancy; the
+precise ours-vs-vLLM L2/scheduling delta needs the vLLM-side ncu, lost when the shared box
+OOM-rebooted mid-capture (twice this session, under 3-way `row/H3-FP4-GPU-E2E` + CPU-GGUF-gate
+contention) — re-run owed on an idle box. Counter-pointed NEXT lever: lift occupancy above
+8.3% (cut the flash kernel's register pressure / `__launch_bounds__`) or cut the
+barrier/smem-scoreboard stalls (ncu's own hint: smaller warp groups / fewer syncs).
+
+**VERDICT.** `-use_fast_math` REJECTED and REVERTED (measured +21 us/call regression; a
+non-byte-exact change with a negative speed effect). No functional code ships (a CMakeLists
+NOTE records why not, so it is not re-tried). The flash decode residual vs vLLM is **+12.5
+us/call (+450 us/step) at c8**, a runtime memory/occupancy effect (ncu above), NOT the
+instruction count. The full SACRED battery + c1..c8 binding a default-flip would need are
+therefore NOT owed (no throughput win to flip). NEXT lever = the occupancy/L2 difference
+(candidate axes a/c: L2 pollution from our marlin/glue neighbours, register pressure /
+`__launch_bounds__`) — the +450 us/step is ~40% of the c8 residual after the dense-direct
+marlin banked the rest.
+
+Evidence: `dgx:~/mxfp4-nsys/{ours_fm_c8,ours_nofm_c8}_cuda_gpu_trace.csv` + `analyze_decode2.py`
+(A/B); `ours_nofm_flash_ncu.ncu-rep` (mechanism); `smoke_fm.log` (#44 3/3); compile arbiter
+(superseded lens). Box OOM-rebooted twice under 3-way contention; left clean.
+
 ## QUANT-CT-MXFP4-FUSED-GLUE W0 — the funded glue-fusion-into-Marlin kernel target is SOURCE-REFUTED; vLLM does NOT fuse glue into the extern Marlin GEMM for W4A16; the c2-c8 residual is FLASH-dominant, not glue (2026-08-06, `row/QUANT-CT-MXFP4-FUSED-GLUE`, source-side on dev box; GPU dumps box-contended by `row/H3-FP4-GPU-E2E`, base `origin/main` `672fc760`)
 
 W0 of the funded MXFP4 kernel campaign. The chartered vehicle is `Yi30/Qwen3-8B-MXFP4`

diff --git a/.agents/specs/mxfp4-flash-audit-2026-08-06.md b/.agents/specs/mxfp4-flash-audit-2026-08-06.md
@@ -0,0 +1,42 @@
+# QUANT-CT-MXFP4-FLASH-AUDIT — the flash decode gap vs vLLM (audit → MEASURED verdict)
+
+<!-- spec -->
+
+Row: `QUANT-CT-MXFP4-FLASH-AUDIT` (helper, worktree `row/QUANT-CT-MXFP4-FLASH-AUDIT`,
+draft PR #69). Base: `origin/main` `4ce9fb74`. Vehicle: `Yi30/Qwen3-8B-MXFP4` (dense
+`Qwen3ForCausalLM`, compressed-tensors W4A16 Marlin keep-quant). Oracle: vLLM **0.25.0**,
+arm `VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel`.
+
+## Target
+The #67 refutation re-attributed the c2-c8 MXFP4 residual FLASH-dominant and OWED a
+STRUCTURAL-lens audit of why the IDENTICAL-grid `flash_fwd_splitkv` decode kernel runs slower
+on our side.
+
+## W1 — DONE. The flash term is REAL, and on CURRENT main SMALLER than #57.
+`analyze_decode2.py` c8 decode-window (both `nsys --cuda-graph-trace=node`, modal marlin=144 /
+gridZ=64): ours 168.8 vs vLLM 156.3 us/call = **+12.5 us/call (+450 us/step)**, not #57's +807.
+Current-main's leaner marlin/glue drop flash 178.1→168.8 — a hint the residual is L2/context.
+
+## W2 — DONE. The mechanism is occupancy/L2, NOT `-use_fast_math`.
+- **Lens 1 (cuobjdump, HYPOTHESIS):** vLLM v0.25.0 pins vllm-flash-attn @ `2c839c33` (exact
+  commit we vendored) → source byte-identical; kernel-version REFUTED. vLLM's flash-attn is
+  built `--use_fast_math`; ours was not (ours 5448 instrs/REG246 → 4832/REG255 with it, =
+  vLLM's 4880/REG255, HMMA/LDSM/LDGSTS identical). SUGGESTED the scalar bloat was the gap.
+- **Lens 2 (MEASURED, the arbiter):** controlled SAME-BUILD nsys A/B (dense-direct default,
+  flash source identical, only the flag differs): `-use_fast_math` makes flash **+21 us/call
+  SLOWER (168.8→189.8)**. Memory-latency-bound kernel; the higher reg count (246→255) LOWERS
+  occupancy, which dominates the instruction reduction. cuobjdump was necessary but NOT
+  sufficient (STRUCTURAL-lens "MEASURED not inferred").
+- **ncu (OURS):** occupancy **8.3%** (register-limited, 4/12 warps/SM), L2 hit **53%**, ~38%
+  short-scoreboard (LDSM→MMA smem) + ~37% CTA-barrier stalls = occupancy-starved latency
+  exposure. vLLM at 255 regs is still 156.3 → its edge is L2/scheduling, not occupancy. The
+  ours-vs-vLLM ncu diff was LOST to a shared-box OOM-reboot (twice, 3-way contention); OWED.
+
+## W3 — VERDICT
+`-use_fast_math` REJECTED + REVERTED (measured +21 us/call regression; a non-byte-exact change
+with a negative speed effect); a CMakeLists NOTE records why so it is not re-tried. No
+functional code ships. Token-exact footnote: fast-math passed #44 smoke 3/3 + coherent (a
+faithful mirror of vLLM's build, just not a speed win). The +450 us/step flash residual is
+occupancy/L2 — a default-flip's SACRED battery + c1..c8 binding are NOT owed (no win to flip).
+NEXT lever (counter-pointed): lift flash occupancy above 8.3% (register pressure /
+`__launch_bounds__`) or cut the barrier/smem stalls; and the vLLM-side ncu on an idle box.
diff --git a/.agents/state.md b/.agents/state.md
@@ -39357,3 +39357,38 @@ encoder of `nano_nemotron_vl.py`, which vllm.cpp already carries rows for, so it
 is mirror-vLLM work that is owed anyway). The transducer decode half of
 parakeet.cpp (RNN-T/TDT/CTC search, joint network) is NOT in vLLM and is a
 separate scope call.
+
+## QUANT-CT-MXFP4-FLASH-AUDIT — `-use_fast_math` on the FA2 TUs REJECTED (measured flash regression); the flash decode gap vs vLLM is occupancy/L2-bound, not the SASS instruction count
+
+<!-- state: 2026-08-06T17:38 -->
+
+Helper row `QUANT-CT-MXFP4-FLASH-AUDIT` (worktree `row/QUANT-CT-MXFP4-FLASH-AUDIT`,
+draft PR #69, base `origin/main` `4ce9fb74`). Runs the #67-OWED flash audit to a MEASURED
+verdict.
+
+W1 (fresh same-tool, `analyze_decode2.py` c8 decode-window): the flash term is REAL, and on
+CURRENT main it is SMALLER than #57 — ours 168.8 vs vLLM 156.3 us/call = **+12.5 us/call
+(+450 us/step)**, not #57's +807 (current-main's leaner marlin/glue drop flash 178.1→168.8,
+a hint the residual is L2/context).
+
+W2 lens 1 (cuobjdump HYPOTHESIS): vLLM's flash-attn is `--use_fast_math`, ours was not (same
+`2c839c33` source; +fast-math → 4832/REG255 = vLLM's 4880/REG255, HMMA/LDSM/LDGSTS identical).
+W2 lens 2 (MEASURED, the arbiter): a controlled SAME-BUILD nsys A/B REFUTES it — `-use_fast_math`
+makes flash **+21 us/call SLOWER (168.8→189.8)**. The kernel is memory-latency-bound; the
+instruction count barely gates time and fast-math's higher reg count (246→255) LOWERS
+occupancy, which dominates. cuobjdump was necessary but NOT sufficient — the STRUCTURAL-lens
+"MEASURED not inferred" rule, learned the hard way (I committed the flag on the compile
+inference before measuring; the runtime overturned it).
+
+W2 mechanism (sudo ncu, OURS): flash occupancy **8.3%** (register-limited, 4/12 warps/SM),
+L2 hit **53%**, ~38% short-scoreboard (LDSM→MMA smem) + ~37% CTA-barrier stalls = occupancy-
+starved latency exposure. vLLM at 255 regs (worse occupancy) is still 156.3, so its edge is
+L2/scheduling, not occupancy — the ours-vs-vLLM ncu diff was LOST when the shared box
+OOM-rebooted mid-capture (twice, under 3-way H3 + CPU-GGUF-gate contention); re-run OWED on an
+idle box.
+
+VERDICT: `-use_fast_math` REVERTED (a CMakeLists NOTE records why, so it is not re-tried); no
+functional code ships. The +450 us/step flash residual is occupancy/L2, NOT the instruction
+count; NEXT lever = lift flash occupancy above 8.3% (register pressure / `__launch_bounds__`)
+or cut the barrier/smem stalls. No default flip (no throughput win). Full tables in
+`.agents/benchmark-record.md` (QUANT-CT-MXFP4-FLASH-AUDIT) + the spec. Box left clean.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -1412,6 +1412,15 @@ if(VLLM_CPP_CUDA)
       src/vt/cuda/cuda_flash_attn_fa2.cu
       ${_FA2_KERNEL_SRCS})
     target_compile_definitions(vllm PUBLIC VLLM_CPP_FLASH_ATTN)
+    # NOTE (QUANT-CT-MXFP4-FLASH-AUDIT #68): -use_fast_math was TRIED here to mirror
+    # vLLM's own flash-attn build (its _vllm_fa2_C is fast-math; it makes our decode
+    # kernel's SASS match vLLM's 4832/REG255 vs 4880/REG255). But a controlled
+    # same-build nsys A/B REFUTED it as a speed lever: this decode kernel is
+    # memory-latency-bound (ncu occ ~10.7%), so the +11.6% instruction reduction does
+    # not help, while fast-math's higher register count (246->255) LOWERS occupancy
+    # and REGRESSED flash 168.8 -> 189.8 us/call. So we do NOT pass -use_fast_math.
+    # The residual flash gap to vLLM (168.8 vs 156.3 us/call) is a runtime memory/L2
+    # effect, not the SASS; see .agents/benchmark-record.md.
     set_source_files_properties(
       src/vt/cuda/cuda_flash_attn_fa2.cu
       ${_FA2_KERNEL_SRCS} PROPERTIES

diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md
@@ -300,7 +300,7 @@ built on it rather than keeping the flattering one.
 | Kimi-Linear-48B-A3B (KDA+MLA+MoE) | Full-model GB10 e2e RUNS (bf16-resident §13), NEAR-TIE 106/128, pool math CLOSES; default OFF | Full model RUNS on GB10 (bf16-resident, RSS peak 1.7 GiB, min-avail 21 GiB, no OOM). Token NEAR-TIE 106/128 (6/8 prompts exact, numerics vs deterministic oracle). 1.59 tok/s. Detail: spec §13 |
 | vLLM 0.26 re-benchmark | Pending | Re-run the binding grids on the advanced pin |
 | MiniMax-H3 FP4 speed (W-FP4a) | Pending. fp4-resident Marlin-W4A16 routing CPU-landed (62/62); GB10 delta + per-step unmeasured (disk window); real e2e disk-blocked; vLLM-Omni has no quantized H3 (BF16-only) | Build CUDA `test_minimax_h3` on dgx, run the NVFP4 case (Marlin via `marlin_gemms`), capture delta + s/step. Detail: benchmark-record + spec §8 |
-| MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** banks +951us marlin. `FUSED-GLUE` W0: glue-into-marlin REFUTED (vLLM doesn't fuse into extern marlin); c8 residual FLASH-dominant, not glue. Flash same-tool audit OWED; see record |
+| MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** banks +951us marlin. `FUSED-GLUE` W0: glue-into-marlin REFUTED (vLLM doesn't fuse into extern marlin); `FLASH-AUDIT` #68: c8 flash occupancy/L2-bound; `-use_fast_math` regresses, rejected |
 | SGLang floor arms | Never ran | Both arms of the SGLang comparison |
 | cuBLAS invocation-parity guard | CI guard landed (CPU); `kGemvHeuristicAlgos` refactor build-verify owed | `nvcc` rebuild + SACRED gate on dgx |
 | Ampere consumer (`sm_86`, RTX 3090 class) | **No number owed; no such board here.** 2026-08-06 build-verify: 7/7 FA2 TUs 0-warn, real `sm_86` SASS. [Detail](../.agents/benchmark-record.md) | External RTX 3090 report. Floor is llama.cpp on that card (GGUF, not our Blackwell-only NVFP4 grid) |