From 927b60f015076854e53c8dc6ae55de12f698872c Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 17:40:41 +0000 Subject: [PATCH] =?UTF-8?q?record(quant):=20FLASH-AUDIT=20#68=20=E2=80=94?= =?UTF-8?q?=20-use=5Ffast=5Fmath=20REJECTED=20(measured=20flash=20regressi?= =?UTF-8?q?on);=20gap=20is=20occupancy/L2,=20not=20the=20SASS?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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] --- .agents/NOW.md | 2 +- .agents/benchmark-record.md | 73 +++++++++++++++++++ .agents/specs/mxfp4-flash-audit-2026-08-06.md | 42 +++++++++++ .agents/state.md | 35 +++++++++ CMakeLists.txt | 9 +++ docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- 7 files changed, 162 insertions(+), 3 deletions(-) create mode 100644 .agents/specs/mxfp4-flash-audit-2026-08-06.md diff --git a/.agents/NOW.md b/.agents/NOW.md index 06ad6637..3227199e 100644 --- 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 index 059324d4..17c51ad2 100644 --- 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 new file mode 100644 index 00000000..8efe1d9a --- /dev/null +++ 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) + + + +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 index 4b96c623..243babef 100644 --- 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 + + + +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 index bf771b3b..c202820e 100644 --- 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 index 7c20280a..a13fb3b9 100644 --- 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) | diff --git a/docs/STATUS.md b/docs/STATUS.md index 30d7552f..fffd9cf8 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -85,7 +85,7 @@ token-for-token correctness against the pinned oracle. | Safetensors loading | Supported | Both gate models plus every registered dense/MoE family | | GGUF loading (F32/F16/BF16/Q4_0/Q8_0/Q2_K/Q3_K/Q4_K/Q5_K/Q6_K/IQ2_XXS/IQ3_XXS/IQ2_S/MXFP4/NVFP4) | Supported; compute-in-quant (keep-quant) on CPU AND now CUDA for the six K-block encodings PLUS Q2_K/IQ2_XXS/IQ3_XXS (DeepSeek-V4 W8, 2026-07-29 - the FIRST CUDA keep-quant GGUF k-quant GEMM `KERNEL-QUANT-CIQ-GEMM-CUDA`, MMVQ-style dequant-in-kernel, GB10-gated 92401/92401 vs the CPU oracle, so a CUDA runner keeps blocks compressed and dots them on the GPU instead of the ARM cores); **NVFP4 now COMPUTES IN FP4 on CUDA for the dense-MLP and full-attention projections (2026-07-29, `CLAIM-GGUF-NVFP4-COMPUTE`), no longer materialize-only** | Weights in six block encodings stay compressed from file to matmul on CPU (no BF16 expansion). NVFP4 (ggml type 40) DEQUANTIZES, including the per-tensor (per-expert) `.scale` sidecar the container keeps outside the blocks; gated BIT-EXACT against the compressed-tensors NVFP4 path on real Qwen3.6-27B bytes from both containers. **It no longer expands to bf16 on CUDA:** an NVFP4 matmul/expert weight is REPACKED at load into the same (`weight_packed [N,K/2]`, `weight_scale [N,K/16]`) operand pair the compressed-tensors path produces - a pure byte permutation, gated BYTE-IDENTICAL against that container - and the existing `vt::MatmulNvfp4*` kernels run on it, so no new kernel exists and no numerics are re-derived. Covers the dense MLP + full-attention q/k/v/o and the MoE shared/routed experts; the GDN `in_proj_*` family and `ssm_out` still expand (the V-head reorder rewrites their layout) and a CPU build still expands everything - the documented `part` subset. **MEASURED on GB10 (2026-07-29), same-binary A/B, one `flock`, idle box, 2 reps per arm:** peak RSS **50.8 -> 25.7 GiB**, load-and-generate **1:58 -> 0:41**, and the 256 projections that move cost 35 840 MiB expanded against 10 080 MiB fp4-resident (3.56x). **The divergence against the safetensors sibling CLOSES:** the fp4 arm is token-IDENTICAL over 24 greedy tokens where the bf16 arm of the same binary diverges at index 4, which retires the reading that that divergence was permanent. It is REPORTED, not gated: the two containers are not the same model - the GGUF NVFP4-quantizes 192 GDN `in_proj` tensors the safetensors keeps BF16 (mean relative weight error ~0.18) and their activation global scales differ - so identity is not guaranteed and a cross-container throughput arm is not valid. SACRED gates unmoved: `test_qwen27_paged_engine` 235/235, `test_qwen36_paged_engine` 315/315. **The MoE (35B) stacked-expert arm is now HARDWARE-GATED too (2026-07-29)**, superseding the gap recorded here: the real 35B A3B NVFP4 GGUF loads and generates through the fp4 path, its 120 routed-expert stacks x 256 experts repack to the modelopt safetensors' own operands with ZERO differing bytes over 840 sampled (tensor, expert) slabs, and all 840 per-expert `.scale[e]` are bit-identical to that expert's `weight_scale_2` - the per-expert scale INDEXING, mutation-proved against both a `scales[0]`-for-every-expert and an expert-0-slab-for-every-expert mutant. Same-binary A/B: peak RSS 68.5 -> 22.7 GiB (3.01x), load-and-generate 1:51.9 -> 0:28.8, tokens IDENTICAL (correct here, since the 35B routed experts run the W4A16 grouped GEMM in both arms). Recorded as OPEN, not smoothed over: this case's 24-token greedy stream is NOT run-to-run stable (one of three `use_a16` runs and one of four safetensors-reference runs differed), so the binding results are the weight-level byte identity and the residency audit, not a token-exactness claim; `test_qwen36_paged_engine` is token-exact at ITS engine params, so the instability belongs to this case's configuration and attributing it is owed work. That run also found and FIXED a latent defect the MoE arm made reachable: the two fp4 fused MoE blocks issued the router GEMM assuming the safetensors `[K,N]` gate layout and threw `matmul: inner dims mismatch` on the GGUF's `[N,K]` one; `MoeRouterLogits` now branches on `nk` (inert for the safetensors path, SACRED gates unmoved). **Q2_K (id 10) + IQ2_XXS (id 16) DEQUANTIZE (2026-07-29, `CLAIM-DSV4-GGUF-LOADER`):** the ~2-bit types the single-Spark `DeepSeek-V4-Flash-GGUF UD-IQ2_XXS`/`UD-Q2_K_XL` vehicles use, ported 1:1 from llama.cpp `ggml-quants.c` (`iq2xxs_grid` codebook + signs; Q2_K nibble sub-scale/min), unit-gated on hand-derived bytes (`test_gguf_dequant` 15/15). Dequant-only (no vec_dot -> expand-bf16). A V4-GGUF model still cannot RUN: the V4-GGUF name map (tensor-manifest-blocked) + the V4 forward (W3-W8) remain. **Multi-shard split GGUF READING landed (2026-08-03, `CLAIM-GGUF-SPLIT-SHARDS`):** `GgufFile::Open` now transparently stitches llama.cpp `gguf-split` shards (`...-00001-of-00003.gguf`) — every shard mmap'd, tensor tables merged, KV metadata taken from shard `00001`, and the sibling shard mappings kept alive by the primary so keep-quant mmap-borrows stay valid across shards (`OwnsSpan` is shard-aware); `VT_GGUF_NO_SPLIT=1` opts out; unit-gated (`test_gguf` split-merge / no-split / count-mismatch cases, 33/33 local). This unblocks the real 3-shard `unsloth/DeepSeek-V4-Flash-0731 UD-IQ2_M` (~91 GiB), whose layout is the NATIVE `deepseek4` arch — per-block `ffn_gate_tid2eid` hash tables (hash layers 0/1/2) + `hc_*` MHC + DSA compressor/indexer are all PRESENT (name-map 1328/1328), `vocab_size` derives from `token_embd` — NOT a standard llama.cpp conversion, so no loader-layout change is owed. It now loads THROUGH 1324/1328 tensors; the sole remaining gap is 4 routed-expert slabs quantized with IQ2_S (id 22, ×2) + MXFP4 (id 39, ×2) — encodings we have GGUF block traits for but no keep-quant vec_dot, so they hit the expand→dequant path which lacks them. Dequant-expanding those 4 big expert tensors to bf16 would add ~17 GiB (~106 GiB total → GB10 OOM-reboot risk), so the memory-safe fix is an IQ2_S+MXFP4 keep-quant kernel (CPU dequant dispatch + the `iq2s_grid` codebook + a CUDA `DotSuperblock`), spec'd as the next brick **IQ2_S (id 22) + MXFP4 (id 39) DEQUANTIZE + KEEP-QUANT on CPU (2026-08-03, `CLAIM-DSV4-UDIQ2M-QUANT`, off-GPU):** the extra per-tensor "dynamic" encodings the `unsloth/DeepSeek-V4-Flash-GGUF UD-IQ2_M` checkpoint mixes into its last routed-expert slabs (IQ2_S `ffn_gate/up` dotting Q8_K, MXFP4 `ffn_down` dotting Q8_0) — ported 1:1 from llama.cpp `ggml-quants.c` @ 237ad9b96 (`iq2s_grid` 1024-entry codebook + DIRECT sign bytes; MXFP4 `kvalues_mxfp4` + `e8m0_to_fp32_half` micro-scaling, distinct from the compressed-tensors `E8M0ToF32` NVFP4 path). CPU dequant + keep-quant `vec_dot`, unit-gated on hand-derived golden bytes (`test_gguf_dequant` 17/17), an INDEPENDENT f64 dequant-then-dot + GEMM NMSE (`test_ops_quant_dot` 19/19), and keep-quant routing (`test_gguf_keep_quant` 37/37) — all CPU-green, so UD-IQ2_M's four previously-`unsupported ggml type 22/39` slabs now load COMPRESSED (no ~17 GiB bf16 expansion that OOM-reboots the box). CUDA: the IQ2_S device `DotSuperblock` is wired into the Q8_K grouped-MoE GEMM and now **CUDA-BUILT + LINKED on GB10 (sm_121a, CUDA 13.0, `-Werror`, 2026-08-03 integration)** — it compiles clean and the merged binary links; MXFP4's device dot (`DotMXFP4`) is written but NOT wired (Q8_0-activation needs a separate 32-block GEMM) so it is marked `[[maybe_unused]]` to keep the ready math without tripping nvcc #177-D, and on GPU MXFP4 CPU-fallbacks like Q4_0/Q8_0. The V4-GGUF forward + a real UD-IQ2_M GPU load/coherence run are owed | | AWQ / GPTQ quantization | W0 spike + W1 CPU INT4 dequant primitive; not yet loadable end to end | INT4 unpack+dequant-to-bf16 for BOTH community formats, mirroring vLLM 1:1 (AWQ reverse-order `awq_triton.py`; GPTQ `qdq_4.cuh` with zero_offset v1/v2 + act-order g_idx). Unit-gated RED-first (hand-computed known bytes + double-precision roundtrip). NOT wired to a loader, no GPU Marlin compute, no model run yet: config recognizer (W2), Marlin GPU GEMM riding the vendored NVFP4 Marlin (W4), CPU e2e (W3), GPTQ 8/2/3-bit (W5) and MoE (W6) are named next bricks. See [.agents/specs/awq-gptq-quant.md](../.agents/specs/awq-gptq-quant.md) | -| MXFP4 (compressed-tensors `mxfp4-pack-quantized`) | Compute PROVEN (#38); GQA-swap ON (#49); decode-graph+gate_up FUSION default-ON. `VT_MARLIN_DENSE` DEFAULT-ON (`KERNEL-MARLIN-DENSE-EXEC`): dense marlin 48-CTA byte-faithful (32B 0.000, 263/263), binding beats #51 every axis (c1 1.020, c8 0.969, mem 2.63x). **`QUANT-CT-MXFP4-FINAL-STACK` TERMINAL — both last levers exhausted: num_splits cap `VT_FA2_NSPLITS_CAP` gated-OFF (c1-only, self-corrects@c8; 32B strict char-identical); glue folds via `vt::FusedChain`; `FUSED-GLUE` W0: residual FLASH-dominant, not glue (vLLM does not fuse into extern marlin). c1 1.020x PASS, c2-c8 0.962-0.969.** state.md | Shared with DeepSeek-V4-Flash + Kimi-K3 MXFP4 paths. CPU E8M0 dequant 5/5·1142. GPU W4A4 + MoE-expert e2e later | +| MXFP4 (compressed-tensors `mxfp4-pack-quantized`) | Compute PROVEN (#38); GQA-swap ON (#49); decode-graph+gate_up FUSION default-ON. `VT_MARLIN_DENSE` DEFAULT-ON (`KERNEL-MARLIN-DENSE-EXEC`): dense marlin 48-CTA byte-faithful (32B 0.000, 263/263), binding beats #51 every axis (c1 1.020, c8 0.969, mem 2.63x). **`QUANT-CT-MXFP4-FINAL-STACK` TERMINAL — both last levers exhausted: num_splits cap `VT_FA2_NSPLITS_CAP` gated-OFF (c1-only, self-corrects@c8; 32B strict char-identical); glue folds via `vt::FusedChain`; `FLASH-AUDIT` #68: c8 flash gap +12.5us/call is occupancy/L2-bound; `-use_fast_math` TRIED, REGRESSES flash (168.8→189.8), rejected. c1 1.020x PASS, c2-c8 0.962-0.969.** state.md | Shared with DeepSeek-V4-Flash + Kimi-K3 MXFP4 paths. CPU E8M0 dequant 5/5·1142. GPU W4A4 + MoE-expert e2e later | | CPU backend vs llama.cpp | At or ahead on every axis (GGUF) | Prefill 1.18x ahead, decode at parity, peak memory 1.01x, byte-identical greedy tokens. Single-stream only; no concurrent-serving comparison has been measured | | Paged KV cache + prefix caching | Supported | Block-paged full attention, hybrid full-attention + GDN state groups, automatic prefix caching (APC) on by default for dense models (cache-ON gated end to end: token-identical output, cache hits, faster TTFT) | | fp8 KV cache (`cache_dtype=fp8`) | In progress (W1 CPU brick), not yet usable end-to-end | HIGH-priority memory/throughput lever (halves the KV footprint). W0 spike + W1 CPU brick landed (`KV-FP8` ACTIVE): fp8-e4m3 K/V STORE (`Quantize(hp/scale)`) + the paged-attention READ dequant (`Dequant(fp8)*scale`) + the `cache_dtype` config parse, all CPU-gated RED-first (`test_ops_fp8_kv_cache` 8/8·511; a wrong store direction fails 3/480). Storage is 1-byte fp8 (`DType::kI8`) + a `Fp8KVCacheDataType` interpretation enum, per-tensor k/v scales (mirroring vLLM `BaseKVCacheMethod`). The CUDA store + fp8 paged-attention read (the GPU memory-halving path, DGX-blocked), the runner/spec integration (half-sized KV blocks + checkpoint-scale threading + `--kv-cache-dtype`/`--calculate-kv-scales`), fp8_e5m2 and per-head scales are named W2-W5 in [.agents/specs/fp8-kv-cache.md](../.agents/specs/fp8-kv-cache.md). No model can run with an fp8 KV cache yet |