QUANT-CT-MXFP4-FLASH-OCCUPANCY: the owed ours-vs-vLLM flash decode ncu diff + occupancy/L2/codegen lever - #75
Closed
localai-bot wants to merge 2 commits into
Closed
Conversation
…ys re-read reframes the premise Opens row/QUANT-CT-MXFP4-FLASH-OCCUPANCY (helper) for the ncu diff #69 owed after the box OOM-reboot lost the vLLM-side capture. A free re-read of the existing c8 nsys CSVs (gpu_trace_c8_dflt vs vllm_offline_trace_c8) reframes the premise BEFORE any GPU work: - ours-vs-vLLM c8 decode flash_fwd_splitkv grid is IDENTICAL: 1x3x64 (num_splits=3, gridZ=64=batch*kv_heads) — same GQA-pack, same split heuristic. - ours uses FEWER registers (216) than vLLM (241) => ours has MORE occupancy headroom, not less. The #69 "8.3% occupancy, register-limited" figure was the batch-1 short-context num_splits=1 kernel (Split=false); the prior ncu pair was context-MISMATCHED (ours ~5-token cli prompt vs vLLM lens=1024). - so at the REAL failing c8 kernel occupancy is NOT vLLM's advantage; the +18us (offline; +12.5 online) residual is codegen (our native sm_121a SASS vs vLLM's sm_80-PTX driver-JIT SASS) or L2 — the matched-workload ncu will disambiguate. Spec .agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md carries the spike contract (W1 matched-c8 ncu, W2 mirror-first codegen lever, W3 binding+verdict). STATUS/BENCHMARKS keyed rows updated (spike-open); STATUS ratchet lowered 3. 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]
…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]
Collaborator
Author
|
Landed on main as the squash (mudler authorship). Next: ptxas-lineage arbitration per the vendor-generated-kernels directive — cross-model payoff if it holds. 🤖 Generated with Claude Code |
Collaborator
Author
|
Correction: the earlier close comment fired before the push (a red NOW-budget gate stopped the chain). Now actually landed on main; the honest close stands. 🤖 Generated with Claude Code |
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]
mudler
added a commit
that referenced
this pull request
Aug 6, 2026
…RITY; MXFP4 terminal complete (#82) row/QUANT-CT-MXFP4-FLASH-PTXAS squash, docs-only. THE ARBITER: NO. (1) vLLM's wheel ships ONLY sm_80 cubins + sm_80 PTX at ISA 9.0 = CUDA 13.0 — the SAME lineage as our nvcc; on GB10 its flash SASS is driver-JIT'd from that PTX, never made by a "wheel ptxas". (2) A CUDA 12.x sm_121a cubin is impossible — 12.x ptxas cannot target sm_121a nor read ISA 9.0. (3) The six-arm cuModule A/B (our PTX and vLLM's own PTX, each under driver-JIT/ptxas-13.0/13.2; identical mangled kernel, grid, 81920B smem, REG=241) ties at [0.969, 1.013] — +/-1.3% box drift. Retires #75's "ptxas SASS-scheduling" attribution: in isolation every arm is ~144-151us, so the engine's +10us/call is ENGINE-CONTEXT adjacency (neighbour L2/orchestration), not kernel SASS. hd256 (27B/35B) is structurally identical — no vendor owed anywhere. MXFP4 TERMINAL, now complete to the assembler-lineage level: c1 1.020 PASS + mem 2.63x; c2-c8 0.962-0.969 with the final residual map = flash engine-context adjacency + portable glue (~18%, numerics- delicate) + marlin/host remainder — all sub-parity levers. No default flip; binding unchanged. 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.
Row
QUANT-CT-MXFP4-FLASH-OCCUPANCY(helper). Ran the #69-owed ours-vs-vLLM flash decode ncu diff to a MEASURED verdict on an idle box, and CORRECTED #69's "occupancy/L2-bound" framing. No functional code ships — this is a diagnosis that closes the flash lever search.W1 — matched-c8 ncu diff (the owed table)
Both engines, IDENTICAL grid
1×3×64, full section set. (Prior #69 ncu was context-mismatched: ours ran a ~5-token prompt = num_splits=1; the "8.3% occupancy" was that batch-1 kernel.)vLLM runs at the SAME 8.33% occupancy (both smem-limited to 1 CTA/SM by the byte-identical 81.92 KB smem). L2 ~1% both ⇒ no adjacency/warmth to win. Only diff: +13% instructions.
W2 — every mirror-first lever refuted (decisively, by building vLLM's exact recipe)
_vllm_fa2_CBuilding vLLM's EXACT recipe reproduces its SASS profile exactly (241 reg, 17,008 instr) and is STILL ~10us slower. Matching vLLM's arch + fast-math + register + instruction count does NOT close the gap: the kernel is not instruction-bound; the residual is vLLM's wheel-
ptxasSASS-scheduling quality, un-reachable from nvcc 13.0.W3 — THE PARITY 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, register count, instruction count,
__launch_bounds__, arch-mirror and fast-math ALL refuted by measurement. The flash term (+450 us/step, ~40% of the c8 residual) is an irreducible-for-us ptxas quality gap. No default flip owed. CMakeLists NOTE + benchmark-record #75 record the closed levers.check-fusion-consistency(minimax_h3) is a pre-existing red, untouched.🤖 Generated with Claude Code