QUANT-CT-MXFP4-FLASH-AUDIT: -use_fast_math REJECTED (measured flash regression); gap is occupancy/L2 - #69
Closed
localai-bot wants to merge 1 commit into
Closed
QUANT-CT-MXFP4-FLASH-AUDIT: -use_fast_math REJECTED (measured flash regression); gap is occupancy/L2#69localai-bot wants to merge 1 commit into
localai-bot wants to merge 1 commit into
Conversation
mudler
force-pushed
the
row/QUANT-CT-MXFP4-FLASH-AUDIT
branch
from
August 6, 2026 17:05
2e36821 to
2e56699
Compare
…ash regression); gap is occupancy/L2, not the SASS QUANT-CT-MXFP4-FLASH-AUDIT runs the #67-OWED flash audit to a MEASURED verdict. W1: fresh same-tool c8 decode-window — the flash gap on CURRENT main is +12.5us/call (+450us/step) ours 168.8 vs vLLM 156.3, SMALLER than #57's +807 (leaner marlin/glue). W2 lens 1 (cuobjdump HYPOTHESIS): vLLM's flash-attn is --use_fast_math, ours was not (same 2c839c33 src; +fast-math -> 4832/REG255 = vLLM's 4880/255). W2 lens 2 (MEASURED, the arbiter): a controlled same-build nsys A/B REFUTES it — -use_fast_math makes flash +21us/call SLOWER (168.8->189.8); the kernel is memory-latency-bound (ncu occ 8.3%, register-limited), so the higher reg count (246->255) lowers occupancy, which dominates. cuobjdump was necessary but NOT sufficient. So the flag is REVERTED (a CMakeLists NOTE records why); no functional code ships. Real residual = occupancy/L2 (ncu: 38% smem-scoreboard + 37% barrier stalls); vLLM-side ncu owed (box OOM-rebooted twice under 3-way contention). No default flip (no throughput win). Records + spec updated. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler
force-pushed
the
row/QUANT-CT-MXFP4-FLASH-AUDIT
branch
from
August 6, 2026 17:40
2e56699 to
927b60f
Compare
localai-bot
changed the title
mudler
added a commit
that referenced
this pull request
Aug 6, 2026
…idual is occupancy/L2 (#69) row/QUANT-CT-MXFP4-FLASH-AUDIT squash, records + a CMakeLists NOTE (no functional change ships). Fresh same-tool c8 decode-window on current main: flash +12.5us/call (+450us/step) — smaller than #57's +807 (the leaner dense-marlin context already recovered part). The compile-lens hypothesis was MEASURED and REFUTED: vLLM pins vllm-flash-attn at the exact commit we vendored (source byte-identical; the earlier ~5824e6e note was wrong); adding vLLM's -use_fast_math reproduces its SASS (4832 vs 4880 instrs, REG255, HMMA/LDSM/LDGSTS identical) yet runs +21us/call SLOWER — the 246->255 register bump cuts occupancy on a latency-bound kernel (ncu: 8.3% occupancy, L2 53%, smem-scoreboard + CTA-barrier stalls dominate). Rejected + reverted with a NOTE so it is not re-tried; cuobjdump is necessary but not sufficient — runtime arbitrates. vLLM at the same 255 regs is still 156.3us/call, so its edge is L2/scheduling: the next lever is flash occupancy (register pressure / launch bounds) or barrier/smem stalls. The precise ours-vs-vLLM ncu diff is OWED on an idle box — the shared GB10 OOM-rebooted twice under 3-way campaign contention during this run. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
Collaborator
Author
|
Landed on main as the squash (mudler authorship). The ncu ours-vs-vLLM diff re-runs on an idle box once the H3 campaign clears — single-campaign box discipline from here on. 🤖 Generated with Claude Code |
mudler
added a commit
that referenced
this pull request
Aug 6, 2026
…IDENTICAL, gap is an irreducible-for-us ptxas SASS-quality difference; no lever (#75) Runs the #69-owed ours-vs-vLLM flash decode ncu diff to a MEASURED verdict on an idle box, and CORRECTS #69's "occupancy/L2-bound" framing. W1 (matched-c8 ncu, both engines, grid 1x3x64, full section set): - vLLM runs at the SAME 8.33% occupancy as ours — both smem-limited to 1 CTA/SM by the byte-identical 81.92 KB kSmemSize (the 216-vs-241 reg delta is moot). - L2 hit ~1% on BOTH (KV streamed from DRAM, no adjacency/warmth to win), identical short-scoreboard stall structure. Only measured diff: +13% instr. - The prior #69 "8.3% occupancy" was the batch-1 short-context num_splits=1 kernel; the prior ncu pair was context-mismatched. Now measured on the real c8 kernel with everything else held identical. W2 (mirror-first levers, all refuted DECISIVELY by building vLLM's exact recipe): - arch-mirror (compute_80 PTX driver-JIT): neutral (no change from native). - compute_80 + -use_fast_math (vLLM's EXACT recipe): reproduces vLLM's SASS profile EXACTLY (241 reg, 17,008 instr vs vLLM 241/17,020) and is STILL ~167us = the same as native ~165us and ~10us slower than vLLM's ~157us. So matching vLLM's arch, fast-math, register AND instruction count does NOT close the gap: the kernel is not instruction-bound; the residual is vLLM's wheel-ptxas SASS-scheduling quality, un-reachable from nvcc 13.0. (native+fast-math REGRESSES to 189.8us per #68; compute_80+fast-math is neutral.) W3 VERDICT: MXFP4 stays BELOW-FLOOR at c2-c8 (binding unchanged 1.020/0.962/0.966/0.969). NO lever exists on our stack — occupancy, L2, num_splits, reg count, instruction count, __launch_bounds__, arch-mirror and fast-math ALL refuted by measurement. The flash term (+450us/step, ~40% of the c8 residual) is an irreducible-for-us ptxas quality gap. NO default flip owed; no functional code ships. CMakeLists NOTE + benchmark-record #75 record the closed levers so they are not re-tried. No byte-exact razor / SACRED battery owed (nothing shipped to gate). check-fusion-consistency (minimax_h3_video_vae_device) is a pre-existing red, untouched by this change. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler
added a commit
that referenced
this pull request
Aug 6, 2026
…as SASS scheduling (#75) row/QUANT-CT-MXFP4-FLASH-OCCUPANCY squash, records + CMakeLists NOTE. The owed matched-context ncu diff (the #69 pair was context-mismatched — its "8.3% ours" was the num_splits=1 batch-1 kernel): both engines run the IDENTICAL flash_fwd_splitkv at IDENTICAL 8.33% occupancy (smem-bound to 1 CTA/SM by the byte-identical 81.92 KB kSmemSize), L2 ~1% both (KV streamed, no warmth story). Rebuilding our TU with vLLM's EXACT recipe (compute_80 PTX + -use_fast_math) reproduces its SASS profile precisely (241 reg, 17k instr/sched) and is STILL ~10us slower — every nameable lever measured and refuted (occupancy, L2, launch bounds, arch, fast-math, registers, instructions). The residual is the wheel-ptxas SASS-SCHEDULING quality from a different toolkit lineage than nvcc 13.0. Binding unchanged: c1 1.020 / c2-c8 0.962-0.969. Fresh c8 map: flash-ptxas ~40% / glue ~18% (sub-parity) / marlin-host remainder. Next per the vendor-generated-kernels directive: ptxas-lineage arbitration (assemble our PTX with the wheel's ptxas; if it hits 157us, AOT-vendor the cubin like the GDN Triton cubins) — cross-model, FA2 decode is shared by every dense/GQA lane. One OOM-reboot root-caused: full-section ncu kernel-replay at util 0.5 blows the unified pool — util 0.15 is the ncu recipe. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler
added a commit
that referenced
this pull request
Aug 6, 2026
…ineage REFUTED, gap is engine context not codegen #75 owed obtaining vLLM's wheel ptxas (hypothesized older CUDA 12.x lineage) and A/B'ing it. Done. The lineage hypothesis dies three ways: (a) cuobjdump of vLLM 0.25.0 _vllm_fa2_C.abi3.so = 52 sm_80 cubins + 52 PTX .version 9.0 (= CUDA 13.0, our own major), NO sm_12x cubin. On GB10 the sm_80 SASS can't run, so vLLM's flash SASS is DRIVER-JIT'd from compute_80 PTX by the box CUDA-13.0 driver — the same assembler #75's Build C already used. There is no older-CUDA-12.x wheel ptxas anywhere in vLLM's path. (b) ptxas 12.8 (the only sub-13 ptxas on the box) tops out at sm_120a and cannot target sm_121a or read PTX 9.0 — an old-12.x sm_121a cubin is impossible. (c) same-params cuModule A/B: the byte-identical c8 decode kernel from our compute_80+fast-math PTX and vLLM's own PTX #30, via driver-JIT / ptxas 13.0 / ptxas 13.2 (all REG=241), ties at ~144us (module/native ratios 0.969-1.013, a +-1.3% tie; the lone 0.969 is box-drift, contradicted by its sisters). Corrected mechanism: flash-kernel CODEGEN is at PARITY across every reachable toolchain AND vs vLLM's own PTX; the +10us/call the engine showed (167 vs 157) is ENGINE CONTEXT (neighbour L2/orchestration, per #69), NOT ptxas. Retires #75's "ptxas SASS-scheduling quality" attribution. No vendor (nothing beat the driver JIT we already use). Binding UNCHANGED c1 1.020 / c2 0.962 / c4 0.966 / c8 0.969; MXFP4 TERMINAL below-floor at c2-c8. hd256 (27B/35B) projection: same structure (sm_80-only PTX driver-JIT'd), no hd256 vendor owed. No functional code shipped; CMakeLists NOTE + benchmark-record (#82) + spec CLOSED record the closed levers. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Draft PR = the helper claim for
row/QUANT-CT-MXFP4-FLASH-AUDIT. Runs the #67-OWED flash audit to a MEASURED verdict. No functional code ships (a CMakeLists NOTE documents a rejected lever).Verdict:
-use_fast_mathREJECTED; the flash gap is occupancy/L2, not the SASS instruction countW1 (fresh same-tool c8 decode-window,
analyze_decode2.py): the flash gap on CURRENT main is +12.5us/call (+450us/step) — ours 168.8 vs vLLM 156.3 — smaller than #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): kernel-version REFUTED (vLLM v0.25.0 pins vllm-flash-attn @
2c839c33, the exact commit we vendored → source byte-identical). vLLM's flash-attn is built--use_fast_math; ours was not. Adding it matches vLLM's SASS (ours 5448/REG246 → 4832/REG255 = vLLM's 4880/255, HMMA/LDSM/LDGSTS identical). This suggested the scalar bloat was the gap.W2 lens 2 (MEASURED — the arbiter): a controlled same-build nsys A/B REFUTES it:
-use_fast_math-use_fast_mathmakes flash +21us/call SLOWER. The kernel is memory-latency-bound (ncu: occ 8.3%, register-limited; L2 hit 53%; ~38% smem-scoreboard + ~37% barrier stalls), so the higher reg count (246→255) lowers occupancy, which dominates the instruction reduction. cuobjdump was necessary but not sufficient — the STRUCTURAL-lens "MEASURED not inferred" rule.What ships
-use_fast_mathreverted (CMakeLists NOTE records why it is not re-tried). Token-exact footnote: fast-math passed #44 smoke 3/3 + coherent — a faithful mirror of vLLM's build, just not a speed win. No default flip (no throughput win to flip), so the SACRED battery + c1..c8 binding are not owed for this lever.Owed
The precise ours-vs-vLLM ncu diff (vLLM at 255 regs is still 156.3, so its edge is L2/scheduling, not occupancy) — lost when the shared box OOM-rebooted twice under 3-way
row/H3-FP4-GPU-E2E+ CPU-GGUF-gate contention. Re-run on an idle box. Next lever the counters point to: lift flash occupancy above 8.3% (register pressure /__launch_bounds__) or cut the barrier/smem stalls.Note: branch based on
4ce9fb74; main has advanced — keyed records may need union-merge on integration.🤖 Generated with Claude Code
https://claude.ai/code/session_014fZAcg1WcU8V629k6HWKys