This is the canonical per-scheme support inventory for roadmap C4/D1. The
coverage spike is specs/quantization-coverage.md.
Sources are pinned to vLLM 555967922 (0.26.0.dev0 + transformers 5.14.1;
advanced 2026-07-26 from the prior e24d1b24/0.25.0 pin, see
specs/pin-advance.md), llama.cpp 237ad9b plus the
documented local type-40/41 fork, MLX 4367c73b, and MLX-LM a790972f.
Maintenance checkpoint for QUANT-GGUF-KEEPQ-LOADER (2026-07-31): GCC 12
-Werror portability is restored for the L7 volatile prefault accumulator
with an expression-only assignment. Prefault reads, defaults, behavior, and
the row's ACTIVE lifecycle are unchanged.
| Column | Meaning |
|---|---|
R |
serialization/config is recognized and validated |
M |
dequantization/repack/materialization exists and is tested |
C |
native quantized compute is dispatched; bf16 expansion is not C |
E |
representative real-model correctness gate passed |
P |
reference-engine performance gate passed on the named backend |
Y means grounded support, - means absent, and part means a documented
subset. A row is DONE only when its declared scope has all required stages;
otherwise it remains PARTIAL or INVENTORIED even if parsing works.
| ID | Item | Upstream | Our code | Tests/evidence | Spike/spec | State | Owner |
|---|---|---|---|---|---|---|---|
QUANT-GGUF-COMPUTE |
Block row (claim the three leaves below, not this row): quantized vt:: storage plus direct GGUF compute kernels; bf16 expansion does not count |
llama.cpp ggml/src/ggml-cpu/ggml-cpu.c:211-406,1245-1443 tensor traits + ggml-cpu/quants.c, ggml-cpu/repack.cpp at 237ad9b96 |
the direct path is LIVE since CIQ G4: vt::MatmulBT routes block dtypes to kMatmulBTQuant and keep-quant is the production default wherever that op is registered; bf16 expansion remains for f16/f32 tensors, unported encodings, non-CPU devices and VT_CPU_REF=1 |
dequant units and same-file APEX gates cover only materialization; CPU perf floor vs same-file llama.cpp RE-MEASURED 2026-07-22 (binding arm = idle dgx.casa aarch64): decode 11.6×, prefill 33.5×, peak RSS 2.65× behind (floor re-measurement), superseding the B4 54–75× / ≈1,480× (ledger B4 row) — the threadpool moved decode/prefill, RSS is UNMOVED. Gap was ATTRIBUTED to ONE op (kMatmul = 95.4 % of wall time, tier-0 quant GEMM 14–44× faster but UNROUTED). G4 ROUTED it 2026-07-22 and re-measured on the same binding recipe: decode 3.38×, prefill 8.20×, peak RSS 2.29× behind (from 11.6/33.5/2.65), i.e. same-binary A/B gains of 3.45×/4.16×/1.16× with byte-identical output tokens. The projected 9–17× did NOT hold: on that mixed file only 1.062 GiB of weight bytes are q8_0 while 1.615 GiB are f16 (incl. the 970 MiB tied token_embd/lm_head), which no block encoding covers, so 60 % of the mass still runs the elementwise kernel. The residual gap is now the ELEMENTWISE bf16/f16 GEMM, not the quant one — that is the next lever, ahead of the SIMD/repack tiers (G4 result). THAT G4 POSITION IS SUPERSEDED and this cell must not be quoted as current: E1-E4, L5, L7, G6 (Arm i8mm) and G7 (q8_0 repack-at-load) have since taken the 20-core Arm/i8mm arm to llama.cpp parity or better on every axis — see BACKEND-GATE-CPU-LLAMACPP in the backend matrix, which is the one place this gate's live position lives. The ISA split matters: G6/G7 are Arm-only, so on x86_64 the quantized weights still take the portable tier (G5 open, no x86 consumer for the G7 repack layout) — #433, x86 arm |
leaves: threadpool, CIQ GEMM, keep-quant loader; umbrella coverage spike; floor re-measurement | READY |
- |
QUANT-GGUF-CPU-THREADPOOL |
vt CPU threadpool + chunked parallel op dispatch (GEMM first, then row/batch ops); bit-identical to single-thread; prerequisite leaf. W1-W3 complete; W4 idle-host speed/RSS gate pending | llama.cpp ggml/src/ggml-cpu/ggml-cpu.c:471-610 (pool/barrier/chunk), :3024-3390 (workers), :1155-1443 (mul_mat chunking) at 237ad9b96 |
pool core, GEMM chunking, row/batch dispatch | barrier/chunk/epoch/concurrency tests, 1/3/20 determinism battery; full CPU ctest 94/94 at each thread count + TSAN clean; checkpoint-gated model tests absent locally; performance run correctly deferred on a contended host (ledger recovery row). EXTENDED 2026-07-23 (CLAIM-CPU-THREAD-GDN-PAGED-1): the last two serial non-GEMM prefill kernels are now threaded — GdnPrefillKernel over (sequence, value-head) and PagedAttentionKernel over query-token rows, both via the same ParallelForRows, both bit-identical (determinism battery extended: single-seq Hv=20 GdnPrefill + 37-token causal PagedAttention, byte-identical at 1/3/20; qwen35 output-token md5 d235db12f2cd304007530286a1755c95 unchanged at threads 1/4/20 + VT_CPU_REF=1; CPU ctest 158/158). Binding dgx aarch64 (idle): prefill 1.382× same-binary (73.0→100.9 t/s), 2.43×→1.76× behind llama.cpp; op-scaling 1→20 GdnPrefill 7.08× / PagedAttention 8.96×; fresh profile re-ranks the NEW bottleneck to the GEMMs (80%) (two-kernel threading, BENCHMARKS). The row's remaining OPEN item is the W4 decode-scaling gate (8.05× < 10× at M=1 decode shapes), unchanged by this extension |
threadpool leaf, two-kernel threading | ANCHOR-BACKFILL |
CLAIM-CPU-THREAD-GDN-PAGED-1 |
QUANT-GGUF-CIQ-GEMM |
Compute-in-quant GEMM: activation quant (Q8_0/Q8_K) + per-type vec_dot dispatch for Q8_0/Q4_K/Q5_K/Q6_K/Q3_K/Q4_0; portable C++ tier, then x86/Arm SIMD + repack tiers. G1-G4 landed — the portable tier-0 path is complete, gated at the OP level, and ROUTED end to end: vt::MatmulBT dispatches a block-dtype weight to kMatmulBTQuant, keep-quant is the production DEFAULT wherever that op is registered, and the six routed encodings compute in quant with no token movement. G6 (2026-07-23) added the Arm i8mm mmla nrc==2 tier for q8_0/q4_0/q4_K/q6_K (q3_K/q5_K have no upstream mmla → stay portable), 2x2-tiled into kMatmulBTQuant at even M,N: op-level q4_K 7–8.4× / q6_K 3.8–4.5× / q8_0 ~1.2× over portable, e2e prefill +8.4 % on the q8_0-dominant bench file (1.44× behind llama.cpp), tokens byte-identical. G7 (2026-07-23) added q8_0 repack-at-load (the q8_0_4x8 tier ggml_repack_get_optimal_repack_type picks on NEON+i8mm): the loader repacks each q8_0 weight once into the block_q8_0x4 interleave and kMatmulBTQuant dispatches a pre-shuffled i8mm gemm/gemv with no per-block register shuffles — op-level q8_0 3.7–5.9× over the mmla tier, E2E prefill 1.92× same-binary → 223.8 t/s vs llama.cpp pp128 177.3 = at/beyond parity (was ~1.5× behind), decode at parity, tokens byte-identical. CPU prefill parity reached; the prefill-lever search is closed (remaining gap = peak RSS 1.39×, loader-bound). G5 (x86) + G8 open. The FRESH op-dispatch profile this row owed is DONE (2026-08-06, dgx aarch64, main @dfd29060, same bench file; see .agents/benchmark-record.md 'FRESH op-dispatch profile'), and it does NOT support starting G5 next: QuantRepackMatmul is 5.06 % of prefill and 15.99 % of decode on aarch64 where the i8mm tier already landed. The profile re-ranks the CPU levers to (1) threadpool synchronisation at 47 % of decode (ThreadReady+PollForWork+Barrier; M=1 cannot amortise the barrier) and (2) CPU paged attention at ~39 % of prefill, of which 20.68 % is a per-ELEMENT dtype switch in the attention dot loop (cpu_paged_attn.cpp:29 called from :143), the same defect class E1 already removed from the elementwise GEMM. G5 stays a real x86 gap worth closing for x86 users, but it is not the top lever, and the x86 box is VOID for timing so it cannot be speed-gated here |
llama.cpp ggml/src/ggml-cpu/ggml-cpu.c:211-406 traits table, ggml-cpu/quants.c:174-860 generic vec_dot, arch/{x86,arm}/quants.c, ggml-cpu/repack.cpp:4153-4830 at 237ad9b96 |
G1: block dtypes + geometry, quant traits table, shared block decoders, op surface. G2: activation quant + scratch sizing (quantize_row_q8_0/q8_K). G3: the six generic vec_dot, block-struct mirror, kMatmulBTQuant quantized path + composite fallback. G4: the routing point — vt::MatmulBT sends a block-dtype b to MatmulBTQuant and is otherwise unchanged, which is sufficient because every model matmul helper already routes an nk=true weight there (qwen3_5.cpp:1067); plus the default flip + expand_nk and the untransposed expand path. G6: Arm i8mm mmla tier (q8_0/q4_0/q4_K/q6_K vmmlaq_s32, HWCAP2_I8MM-probed, VT_CPU_QUANT_MMLA defeat) + 2x2 tile in kMatmulBTQuant, per-file +i8mm in CMakeLists |
G1 traits cross-check + fallback units — 8 cases / 5,615 assertions green (was 5,694; its composite case now covers Q8_K alone because the six weight types legitimately no longer take that path): vt geometry vs the reader's GgmlTraits vs ggml-common.h arithmetic all agree, and the composite equals the loader dequant byte-for-byte. G2/G3 units — 16 cases / 78,052 assertions green: every vec_dot gated against an INDEPENDENT f64 dequantize-then-dot reference (tolerance relative to the dot's L1 magnitude, actual agreement ~1e-6) over nblocks {1,2,3,5,7,16} incl. single-block and odd multiples; ragged K throws at every layer; upstream thresholds ported unwidened (test-quantize-fns:17-28, test-backend-ops:4277 NMSE ≤ 5e-4 at M {1,4,32,512} × N {1,7,16}); bit-exact run-to-run and across threads 1/2/4; byte-exact encoder gate pins the rounding rules; 14-mutant battery, 13 caught, the 1 uncaught mutant provably unreachable. dequant units still green after the decoder move. DGX (G2/G3 re-confirmed, each gate STANDALONE, goldens md5 identical before/after): clean CUDA -Werror build 0 warnings + full regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16 on both 0.6B and 4B, OPT 6/6, DeepSeek-V2 8/8) + test_qwen36_gguf_engine 28/28 with 16/16 tokens on both APEX files + the new CPU units green on aarch64 with identical counts. G4 (2026-07-22): test_qwen36_gguf_engine PASSES STANDALONE on a CPU-only dgx build (where keep-quant is live) — 2/2 cases, 16/16 greedy tokens on APEX-Compact AND APEX-Balanced vs the same-file llama.cpp oracle, exercising 5 of the 6 routed encodings end to end; the CUDA regression set is UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, DeepSeek-V2 8/8, gguf 28/28 incl. VT_CPU_REF=1), goldens md5 identical. Binding CPU A/B (idle dgx aarch64, one flock, same binary, 3 reps, Qwen3.5-2B-UD-Q8_K_XL): decode 2.216 -> 7.650 t/s (3.45x), prefill 5.149 -> 21.44 t/s (4.16x), peak RSS 7.428 -> 6.401 GiB, output tokens byte-identical across the pre-G4, post-G4 and VT_CPU_REF=1 arms. Still 3.38x / 8.20x / 2.29x behind llama.cpp — the projected 9-17x did NOT hold because 60 % of that file's weight bytes are f16, which no block encoding covers. That gap is now CLOSED by KERNEL-GEMM-CPU-ELEM (2026-07-22, same box/recipe/binary discipline): the elementwise kernel went 18-24 -> 69-351 GFLOP/s bit-exactly, taking the CPU position to decode 1.03x behind (parity within 3.1 %) and prefill 2.34x behind, tokens unchanged (same md5). Its measured NEGATIVE re-ranks G5-G8 once more: M-blocking the elementwise GEMM bought 1.63x op-level and 0.0 % end-to-end, so the 95.37 % kMatmul attribution these G-rows were ranked against is STALE and a FRESH op-dispatch profile is owed before G5/G6/G7 are started. G6 (2026-07-23): Arm i8mm mmla tier landed against the refreshed profile (kMatmulBTQuant 50 % + kMatmul 16 % + kMatmulBT 14 % = 80 % of prefill). test_ops_quant_dot G6 cross-check — 19 cases / 78,162 assertions on dgx aarch64: q8_0/q4_0 mmla BIT-IDENTICAL to the portable/scalar tier (vmlaq_f32 non-fused under -ffp-contract=off), q4_K/q6_K within NMSE ≤ 5e-4, mmla GEMM bit-identical across threads 1/2/4/20. test_qwen36_gguf_engine 2/2 · 16/16 on both APEX files with mmla live (q8_0/q4_K/q6_K at prefill), bench-file token md5 d235db12f2cd304007530286a1755c95 byte-identical across mmla-OFF/ON/VT_CPU_REF=1. Op-level portable→i8mm: q8_0 -Werror 0-warn, regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 138, Qwen3-dense 184, OPT, DeepSeek-V2 223), goldens untouched. G7 (2026-07-23): q8_0 repack transform + i8mm repack gemm/gemv dispatched from kMatmulBTQuant on b.repacked; loader repacks via OwnGgufQuantBlocks under GgufLoadPolicy::quant_repack, flag carried to the kernel through ResidentWeight. test_ops_quant_repack — 305 assertions on dgx aarch64: repacked gemm/gemv memcmp-equal to plain kMatmulBTQuant across decode/leftover/prefill, f32+bf16 out, strided activations, threads 1/2/4/20; interleave matches make_block_q8_0x4 byte-for-byte (110 on x86, numeric skip). test_qwen36_gguf_engine STANDALONE 2/2·16/16 on APEX Compact+Balanced (repack live), token md5 d235db12f2cd304007530286a1755c95 byte-identical across repack-ON/OFF/VT_CPU_REF=1. Binding dgx aarch64 (idle, one flock, 6 interleaved reps): op-level q8_0 3.7–5.9× (518→2401/583→3456/514→1902 GFLOP/s); E2E prefill 1.92× (1096→572 ms), 223.8 t/s vs llama.cpp pp128 177.3 = 1.26× at/beyond parity, decode at parity, RSS unchanged; fresh profile q8_0 GEMM 55%→-Werror 0-warn, regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, DeepSeek-V2 8/8, Llama 16/16), goldens content-hash identical . P0 REGRESSION FOUND + FIXED (2026-08-06, CLAIM-QUANT-GGUF-CIQ-GROUPED-DTYPE): the GROUPED provider MatmulBTQuantGroupedKernel was f32-ONLY — it advanced a float* by act.stride[0] and declared the row kF32 whatever act.dtype said, so a bf16/f16 activation was mis-strode 2x AND mis-decoded. Every prior caller/test passed f32; qwen3_5 W3b KqGrouped (bf16 act, b4f5610a) was the first non-f32 caller, so CPU-only GGUF 35B decode became all-token-0 while the CUDA gate stayed byte-exact (CUDA always honoured act.dtype). Fixed at cpu_quant_gemm.cpp:220-268 (rows addressed by SizeOf(act.dtype)/SizeOf(out.dtype); repacked/q8_0_aligned now propagate onto the per-expert slice — the CIQ-G7 all-zero mode). Gated per activation dtype + bf16-out by 2 NEW cases in test_ops_quant_dot.cpp (RED pre-fix on f16+bf16 for all 12 weight encodings, GREEN after; f32 unaffected either way) |
CIQ GEMM leaf | ANCHOR-BACKFILL |
CLAIM-QUANT-GGUF-CIQ-G7-1 |
QUANT-GGUF-KEEPQ-LOADER |
Keep-quantized GGUF loader: block-resident 2-D matmul weights ([N,K], no transpose), per-tensor routing, VT_CPU_REF dequant-oracle switch, bench-branch 7c91a42 merge. L1+L2+L3 landed — block residency, the TOTAL per-tensor routing policy and the VT_CPU_REF oracle switch all exist and are gated. Keep-quant is DEFAULT ON since CIQ G4 wherever the running device has a registered kMatmulBTQuant (CPU, and since 2026-07-29 also CUDA for the Q8_K family via the KERNEL-QUANT-CIQ-GEMM-CUDA kCUDA provider — a CUDA runner now keeps k-quant/i-quant blocks COMPRESSED instead of expanding), with VT_GGUF_KEEP_QUANT=0 as the opt-out. L4 measured; L5 LANDED (mmap in-place residency + tied-head sharing + read-once page release) — peak RSS 6.401 -> 3.884 GiB, 2.29x -> 1.39x llama.cpp, byte-identical |
llama.cpp src/llama-model-loader.cpp:1047,1385 (file-typed residency), :1676 + ggml/src/llama-mmap.cpp:490 (unmap_fragment), ggml/src/ggml-cpu/repack.cpp:4727 (repack-at-load hook) at 237ad9b96 |
L1: dense-arch (qwen35) GGUF path on main via the registry — dense GGUF load, arch->registered-ID map, F16/BF16 row dequant. L2: block residency OwnGgufQuantBlocks — raw ggml blocks into an OwnedTensor with a block vt::DType, file [N,K] orientation, nk=true, no transpose; stacked experts split by byte range. L3: routing policy + VT_CPU_REF/VT_GGUF_KEEP_QUANT (6 roles, no default: label so an unrouted role is a -Werror=switch build failure) wired at every loader call site via OwnMatmulWeight/RequireExpand. Default now device-derived (CIQ G4): GgufQuantComputeAvailable gates it on vt::OpRegistered(kMatmulBTQuant, CurrentPlatform().device_type()), and the same condition drives expand_nk, which stops transposing a weight that must expand. expand_nk now also covers the GDN split projections (2026-07-23, CLAIM-CPU-GDN-ORIENT-1): a fresh op-dispatch profile found LoadGdnGguf's in_proj_qkv/z/b/a + out_proj were the ONE expanded weight family still transposed to [K,N] (nk=false → slow kMatmul, 17.9 % of prefill); the new gdn_expand_nk field + MakeGdnProj keep them [N,K] nk=true (V-head reorder applied first, orthogonal to orientation; VT_GGUF_GDN_NK=0 A/B opt-out) → M-blocked kMatmulBT, same-binary prefill 1.090× / decode 1.09×, byte-identical (kMatmul 72→0 calls in prefill) |
L2/L3 units — 17 cases / 5,574 assertions green. Gate 1 (losslessness) proven PER ENCODING, one case each for Q4_0/Q8_0/Q3_K/Q4_K/Q5_K/Q6_K: resident bytes memcmp-equal to the file span and resident-block dequant BYTE-IDENTICAL to the direct-from-file expansion (f32 and bf16), over pseudo-random block bytes constrained only to finite f16 scales; at loader level the kept weight rehydrates to the expanded [K,N] bf16 tensor byte for byte, per weight and per expert, on dense and MoE fixtures. Totality: the audit hook proves routed == the file's complete tensor list on both fixtures, plus 6 roles × 12 encodings × 6 shapes against a LONGHAND expectation (12 keep / 420 expand, so neither outcome is vacuous). Gate 2 (oracle stability): VT_CPU_REF=1 keeps nothing quantized and every weight is bit-identical to the historical load; on dgx test_qwen36_gguf_engine under VT_CPU_REF=1 is 28/28 assertions, 16/16 tokens on both APEX files — same as without. 10-mutant battery, 10 caught (the expert-slice-offset mutant survived the first pass, exposed a real coverage hole, and drove the MoE fixture). DGX (each gate STANDALONE, production flags, goldens md5 identical before/after 2965ef5772b556d3f3f86fedf4221b2f): clean CUDA -Werror 0 warnings + regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16 on both, OPT 6/6, DeepSeek-V2 8/8) + gguf units green on aarch64 with identical counts; full CPU ctest 154/154. RSS at G4 was 6.401 GiB (2.29x); L5 took it to 3.884 GiB (1.39x) — binding, idle dgx aarch64, same-binary 3-rep A/B: mmap in-place residency (borrow kept q8_0 blocks out of the mapping, refcounted, -0.998 GiB), tied-head sharing (one bf16 vocab matrix for embed+lm_head, -0.946 GiB), read-once page release (MADV_DONTNEED the expanded tensors' file pages, port of llama.cpp unmap_fragment, -0.573 GiB). Decode TPOT 41.7 ms UNCHANGED, prefill TTFT +4% (first-touch faults move into the timed window), output md5 d235db12f2cd304007530286a1755c95 identical across BEFORE/AFTER/ORACLE. Lifetime safety tested explicitly (borrow outlives the GgufFile AND the on-disk file; shared head freed once either order). L6 (2026-07-23, CLAIM-QUANT-GGUF-KEEPF16-L6-1) implemented keep-f16 residency and REFUTED the "remaining gap is the f16 expansion" attribution above. New kKeepF16 residency + OwnGgufF16/OwnGgufKeptSlice keep F16 matmul weights (+ F16 embed/tied head, one shared f16 vocab matrix via OwnedBytes::KeepAlive) resident as F16, consumed by the elementwise f16 GEMM. Binding A/B: peak RSS 3.884 → 3.832 GiB (−52 MB, RSS-NEUTRAL) — L5's page-release ALREADY dropped the f16 file pages, so keep-f16 only swaps an anonymous bf16 buffer for equal-size file-backed f16 pages. smaps attribution: keep-f16 file-backed 2.634 GiB ≈ llama.cpp's 2.68 file (weight residency AT PARITY), anon 1.20 GiB; the remaining ~1.08 GiB gap is the engine's ANONYMOUS activation/KV workspace, NOT weights — the real, separate CPU RSS lever. Also regresses prefill (TTFT 577 → ~1000 ms, first-touch faults into the timed window; decode at parity). Tokens byte-identical (md5 d235db1…). Ships DEFAULT OFF at L6. L7 (2026-07-23, CLAIM-QUANT-GGUF-RSS-L7-1) REVERSED L6's refutation and CLOSED the CPU RSS gap to 1.01× llama.cpp. The profile disproved the "workspace" attribution — DevicePool 20 MiB, whole KV 115 MiB, both ≤ llama.cpp. The 1 GiB residual was a q8_0 repack-source DOUBLE-COUNT: on aarch64 the G7 repack COPIES q8_0 into an anonymous buffer while the f16 borrows keep the mapping alive, so the DEAD source blocks stay file-backed. OwnGgufQuantBlocks now DropSpanResidencyes the repack source (port of llama.cpp unmap_fragment), and PrefaultBorrowedSpan faults borrowed weights at load (port of llama.cpp mmap prefetch), removing L6's prefill regression — so keep-f16 flips DEFAULT ON (VT_GGUF_KEEP_F16=0 opt-out). Binding A/B (idle dgx aarch64, base-vs-L7 same-binary): peak RSS 3.884 → 2.832 GiB = 1.39× → 1.01× llama.cpp (File 2.632 → 1.629, the released q8_0 source; anon 1.200 unchanged), prefill 1.18× AHEAD (204 vs pp128 173.2), decode ~parity (24.4 vs 25.09, ~1.4% native-f16 cost), tokens BYTE-IDENTICAL (md5 809f2d0… base/L7/oracle). Anon 1.200 GiB is IRREDUCIBLE (repacked q8_0 1.06 + KV 0.115 + pool 0.02). Regressions UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, DeepSeek-V2 8/8, Llama 16/16, GGUF engine 28/28); test_gguf_keep_quant 36/36 (+1 L7 prefault byte-transparency case, x86+aarch64) |
keep-quant loader leaf | ANCHOR-BACKFILL |
CLAIM-QUANT-GGUF-RSS-L7-1 |
QUANT-GGUF-PRESETS |
Representative mixed-file gates for every llama.cpp output preset family | llama.cpp tools/quantize/quantize.cpp:34-74 |
only custom APEX mixed files are executable; no general preset dispatch | APEX gates do not prove llama.cpp preset breadth | coverage spike; split exact preset IDs before READY |
INVENTORIED |
- |
Upstream encodings: ggml/include/ggml.h:389-432. Our recognition is
gguf_reader.cpp:191;
materialization is gguf_dequant.cpp:246.
Since CIQ G4 (2026-07-22) a production load KEEPS the six routed encodings
in their ggml blocks and computes on them: block residency + the total
per-tensor routing policy
(gguf_keep_quant.cpp)
are DEFAULT ON wherever kMatmulBTQuant is registered for the running device,
and vt::MatmulBT dispatches a block-dtype weight to
it. Everything else still expands to bf16 —
f16/f32 tensors, unported encodings, value-rewritten tensors, the embedding
table, and every non-CPU device — and VT_GGUF_KEEP_QUANT=0 or VT_CPU_REF=1
forces the full dequant path back.
| ID | Encoding | ggml ID / upstream role | R | M | C | E | P | State | Our code + tests/evidence | Spike/spec | Owner |
|---|---|---|---|---|---|---|---|---|---|---|---|
QUANT-GGUF-F32 |
F32 | 0 / output | Y | Y | - | Y | - | PARTIAL |
dequant dispatch; unit; APEX gate | coverage spike; leaf open | - |
QUANT-GGUF-F16 |
F16 | 1 / output | Y | Y | - | Y | - | PARTIAL |
f16 row dequant; unit (exact half decode, plus a BF16 sibling case) (landed with keep-quant loader L1); executed end to end by the Qwen3.5-2B-UD-Q8_K_XL bench file, 56 of whose 335 tensors are f16 — including the 970 MiB tied token_embd/lm_head and whole ffn layers. C stays -: f16 is not a block encoding, so it never runs the QUANTIZED GEMM. Its speed gap (elementwise bf16 kernel at 17-25 GFLOP/s) was CLOSED by the elementwise-GEMM vectorization (KERNEL-GEMM-CPU-ELEM, 69-351 GFLOP/s). keep-f16 residency is now DEFAULT ON (keep-quant loader L6+L7): an F16 matmul weight (+ F16 embed/tied head) stays F16 resident and the elementwise f16 GEMM computes on it directly (mirrors llama.cpp ggml_vec_dot_f16), byte-faithful (VT_GGUF_KEEP_F16=0 opt-out). L7 CLOSED the CPU RSS gap to 1.01× llama.cpp (2.832 vs 2.798 GiB): keeping f16 file-backed removes the anonymous bf16 expansion, and releasing the dead q8_0 repack-source pages removes the double-count that made L6 look neutral; a load-time prefault removes the prefill regression (now 1.18× ahead). The L6 "gap is engine workspace" claim was WRONG — the profile measured DevicePool 20 MiB + KV 115 MiB (both ≤ llama.cpp); the residual was weight residency |
keep-f16 default-on residency L6/L7 | - |
QUANT-GGUF-Q4_0 |
Q4_0 | 2 / output | Y | Y | Y | - | - | PARTIAL |
dequant kernel, dispatch; unit C = Y since G4 (2026-07-22): the weight stays in its ggml blocks at load (keep-quant policy, DEFAULT ON wherever kMatmulBTQuant is registered for the running device) and vt::MatmulBT dispatches it to the tier-0 vec_dot GEMM — no bf16 expansion on the executed path. Op gate test_ops_quant_dot; residency losslessness + routing gates test_gguf_keep_quant. Binding CPU A/B on the mixed Qwen3.5-2B-UD-Q8_K_XL file (idle dgx aarch64, same binary, 3 reps): decode 3.45x, prefill 4.16x, peak RSS 1.16x less, output tokens byte-identical to both the pre-G4 arm and the VT_CPU_REF=1 oracle. P stays -: still 3.38x/8.20x behind llama.cpp on the same file, because 60% of ITS weight bytes are f16 and take the elementwise kernel. |
leaf open | - |
QUANT-GGUF-Q4_1 |
Q4_1 | 3 / output | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-Q5_0 |
Q5_0 | 6 / output | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-Q5_1 |
Q5_1 | 7 / output | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-Q8_0 |
Q8_0 | 8 / output | Y | Y | Y | Y | - | PARTIAL |
dequant kernel, dispatch; unit; APEX gate C = Y since G4 (2026-07-22): the weight stays in its ggml blocks at load (keep-quant policy, DEFAULT ON wherever kMatmulBTQuant is registered for the running device) and vt::MatmulBT dispatches it to the tier-0 vec_dot GEMM — no bf16 expansion on the executed path. Op gate test_ops_quant_dot; residency losslessness + routing gates test_gguf_keep_quant. Binding CPU A/B on the mixed Qwen3.5-2B-UD-Q8_K_XL file (idle dgx aarch64, same binary, 3 reps): decode 3.45x, prefill 4.16x, peak RSS 1.16x less, output tokens byte-identical to both the pre-G4 arm and the VT_CPU_REF=1 oracle. P stays -: still 3.38x/8.20x behind llama.cpp on the same file, because 60% of ITS weight bytes are f16 and take the elementwise kernel. |
leaf open | - |
QUANT-GGUF-Q2_K |
Q2_K | 10 / output | Y | Y | Y | - | - | ACTIVE |
dequant kernel (DequantQ2_K, ported 1:1 from llama.cpp ggml-quants.c:903), vt block dtype kQ2_K geometry, dispatch. C = Y since DeepSeek-V4 W8 (2026-07-29): the keep-quant vec_dot VecDotQ2_KQ8_K (1:1 port of ggml vec_dot_q2_K_q8_K_generic, quants.c:514) + the Q8_K traits row make HasQuantDotKernel TRUE, so the loader keeps Q2_K blocks COMPRESSED and dots them directly (never expand-to-bf16) — the UD-Q2_K_XL DeepSeek-V4 sibling vehicle. Op gate test_ops_quant_dot (added to kWeightCases: vec_dot vs f64 dequant-dot ≤1e-5·L1, NMSE ≤5e-4, RED-first proven). E/P = -: the DeepSeek-V4 e2e run is W8-residual (W2b tower materialization pending). CUDA compute LANDED 2026-07-29 (KERNEL-QUANT-CIQ-GEMM-CUDA, CLAIM-CUDA-KEEPQUANT-GEMM): a native kCUDA kMatmulBTQuant provider (cuda_quant_dot.cu, MMVQ-style) dots these blocks ON the GPU (Q8_K activation quant + integer dot, dequant-in-kernel, weights stay COMPRESSED in the unified pool), GB10-gated (test_cuda_quant_dot) 2/2 · 92401 vs the CPU oracle (NMSE ≤1e-6, int core bit-exact) + f64 dequant (≤5e-4), compute-sanitizer 0, RED-first proven — so on a CUDA runner these experts dispatch to the GPU, not the 20 ARM cores. |
CUDA keep-quant GEMM; iquant DSV4 spike | CLAIM-DEEPSEEK-V4-W8 |
QUANT-GGUF-Q3_K |
Q3_K | 11 / output + S/M/L presets | Y | Y | Y | Y | - | PARTIAL |
dequant kernel, dispatch; unit; APEX gate C = Y since G4 (2026-07-22): the weight stays in its ggml blocks at load (keep-quant policy, DEFAULT ON wherever kMatmulBTQuant is registered for the running device) and vt::MatmulBT dispatches it to the tier-0 vec_dot GEMM — no bf16 expansion on the executed path. Op gate test_ops_quant_dot; residency losslessness + routing gates test_gguf_keep_quant. Binding CPU A/B on the mixed Qwen3.5-2B-UD-Q8_K_XL file (idle dgx aarch64, same binary, 3 reps): decode 3.45x, prefill 4.16x, peak RSS 1.16x less, output tokens byte-identical to both the pre-G4 arm and the VT_CPU_REF=1 oracle. P stays -: still 3.38x/8.20x behind llama.cpp on the same file, because 60% of ITS weight bytes are f16 and take the elementwise kernel. |
leaf open | - |
QUANT-GGUF-Q4_K |
Q4_K | 12 / output + S/M presets | Y | Y | Y | Y | - | PARTIAL |
dequant kernel, dispatch; unit; APEX gate C = Y since G4 (2026-07-22): the weight stays in its ggml blocks at load (keep-quant policy, DEFAULT ON wherever kMatmulBTQuant is registered for the running device) and vt::MatmulBT dispatches it to the tier-0 vec_dot GEMM — no bf16 expansion on the executed path. Op gate test_ops_quant_dot; residency losslessness + routing gates test_gguf_keep_quant. Binding CPU A/B on the mixed Qwen3.5-2B-UD-Q8_K_XL file (idle dgx aarch64, same binary, 3 reps): decode 3.45x, prefill 4.16x, peak RSS 1.16x less, output tokens byte-identical to both the pre-G4 arm and the VT_CPU_REF=1 oracle. P stays -: still 3.38x/8.20x behind llama.cpp on the same file, because 60% of ITS weight bytes are f16 and take the elementwise kernel. |
leaf open | - |
QUANT-GGUF-Q5_K |
Q5_K | 13 / output + S/M presets | Y | Y | Y | Y | - | PARTIAL |
dequant kernel, dispatch; unit; APEX gate C = Y since G4 (2026-07-22): the weight stays in its ggml blocks at load (keep-quant policy, DEFAULT ON wherever kMatmulBTQuant is registered for the running device) and vt::MatmulBT dispatches it to the tier-0 vec_dot GEMM — no bf16 expansion on the executed path. Op gate test_ops_quant_dot; residency losslessness + routing gates test_gguf_keep_quant. Binding CPU A/B on the mixed Qwen3.5-2B-UD-Q8_K_XL file (idle dgx aarch64, same binary, 3 reps): decode 3.45x, prefill 4.16x, peak RSS 1.16x less, output tokens byte-identical to both the pre-G4 arm and the VT_CPU_REF=1 oracle. P stays -: still 3.38x/8.20x behind llama.cpp on the same file, because 60% of ITS weight bytes are f16 and take the elementwise kernel. |
leaf open | - |
QUANT-GGUF-Q6_K |
Q6_K | 14 / output | Y | Y | Y | Y | - | PARTIAL |
dequant kernel, dispatch; unit; Compact gate, Balanced gate C = Y since G4 (2026-07-22): the weight stays in its ggml blocks at load (keep-quant policy, DEFAULT ON wherever kMatmulBTQuant is registered for the running device) and vt::MatmulBT dispatches it to the tier-0 vec_dot GEMM — no bf16 expansion on the executed path. Op gate test_ops_quant_dot; residency losslessness + routing gates test_gguf_keep_quant. Binding CPU A/B on the mixed Qwen3.5-2B-UD-Q8_K_XL file (idle dgx aarch64, same binary, 3 reps): decode 3.45x, prefill 4.16x, peak RSS 1.16x less, output tokens byte-identical to both the pre-G4 arm and the VT_CPU_REF=1 oracle. P stays -: still 3.38x/8.20x behind llama.cpp on the same file, because 60% of ITS weight bytes are f16 and take the elementwise kernel. |
leaf open | - |
QUANT-GGUF-IQ2_XXS |
IQ2_XXS | 16 / output | Y | Y | Y | - | - | ACTIVE |
reader trait :230 (66 B block); codebook dequant kernel (DequantIQ2_XXS, ported 1:1 from llama.cpp ggml-quants.c:2416); grids moved to the shared cpu_quant_iq_tables.h (iq2xxs_grid/ksigns_iq2xs/kmask_iq2xs, ggml-common.h:499,503,550); vt block dtype kIQ2_XXS geometry, dispatch. C = Y since DeepSeek-V4 W8 (2026-07-29): the keep-quant vec_dot VecDotIQ2_XXSQ8_K (1:1 port of ggml vec_dot_iq2_xxs_q8_K_generic, quants.c:855) + the Q8_K traits row make HasQuantDotKernel TRUE ⇒ the loader keeps IQ2_XXS blocks COMPRESSED and dots them directly — the MEMORY ENABLER for the single-Spark unsloth/DeepSeek-V4-Flash-GGUF UD-IQ2_XXS gate/up routed experts (ffn_gate_exps/ffn_up_exps). Op gate test_ops_quant_dot; RED-first proven (perturb the 0.125 fold → fail, revert → green). W2b GGUF→tower materialization LANDED (2026-07-29, CLAIM-DEEPSEEK-V4-W2B): LoadDeepseekV4FromGguf wires the blk.N.* name-map + these keep-quant blocks into the DeepseekV4 weight towers (MW/SEW stay COMPRESSED, gated test_deepseek_v4_gguf_load 5/5·149 tiny synthetic). E/P = -: the real DeepSeek-V4 91 GB UD-IQ2_XXS e2e run stays W8-final (download + DGX). CUDA compute LANDED 2026-07-29 (KERNEL-QUANT-CIQ-GEMM-CUDA, CLAIM-CUDA-KEEPQUANT-GEMM): a native kCUDA kMatmulBTQuant provider (cuda_quant_dot.cu, MMVQ-style) dots these blocks ON the GPU (Q8_K activation quant + integer dot, dequant-in-kernel, weights stay COMPRESSED in the unified pool), GB10-gated (test_cuda_quant_dot) 2/2 · 92401 vs the CPU oracle (NMSE ≤1e-6, int core bit-exact) + f64 dequant (≤5e-4), compute-sanitizer 0, RED-first proven — so on a CUDA runner these experts dispatch to the GPU, not the 20 ARM cores. |
CUDA keep-quant GEMM; iquant DSV4 spike | CLAIM-DEEPSEEK-V4-W8 |
QUANT-GGUF-IQ2_XS |
IQ2_XS | 17 / output; IQ2_S preset storage | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-IQ3_XXS |
IQ3_XXS | 18 / output | Y | Y | Y | - | - | READY |
reader trait :237 (98 B block, ADDED W8); codebook dequant kernel (DequantIQ3_XXS + the shared iq3xxs_grid, ported 1:1 from llama.cpp ggml-quants.c:2503 + ggml-common.h:1007); vt block dtype kIQ3_XXS geometry. C = Y (DeepSeek-V4 W8, 2026-07-29): the keep-quant vec_dot VecDotIQ3_XXSQ8_K (1:1 port of ggml vec_dot_iq3_xxs_q8_K_generic, quants.c:999) + the Q8_K traits row keep IQ3_XXS blocks COMPRESSED — the MEMORY ENABLER for the UD-IQ2_XXS down routed experts (ffn_down_exps, which are IQ3_XXS, NOT IQ2_XXS; without keep-quant they alone OOM the box). Op gate test_ops_quant_dot (vec_dot vs f64 dequant-dot ≤1e-5·L1, NMSE ≤5e-4). W2b GGUF→tower materialization LANDED (2026-07-29, CLAIM-DEEPSEEK-V4-W2B): LoadDeepseekV4FromGguf wires the name-map + keep-quant blocks into the DeepseekV4 towers (test_deepseek_v4_gguf_load 5/5·149). E/P = -: the real DeepSeek-V4 e2e run stays W8-final. CUDA compute LANDED 2026-07-29 (KERNEL-QUANT-CIQ-GEMM-CUDA, CLAIM-CUDA-KEEPQUANT-GEMM): a native kCUDA kMatmulBTQuant provider (cuda_quant_dot.cu, MMVQ-style) dots these blocks ON the GPU (Q8_K activation quant + integer dot, dequant-in-kernel, weights stay COMPRESSED in the unified pool), GB10-gated (test_cuda_quant_dot) 2/2 · 92401 vs the CPU oracle (NMSE ≤1e-6, int core bit-exact) + f64 dequant (≤5e-4), compute-sanitizer 0, RED-first proven — so on a CUDA runner these experts dispatch to the GPU, not the 20 ARM cores. |
CUDA keep-quant GEMM; iquant DSV4 spike | - |
QUANT-GGUF-IQ1_S |
IQ1_S | 19 / output | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-IQ4_NL |
IQ4_NL | 20 / output | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-IQ3_S |
IQ3_S | 21 / IQ3 S/XS/M storage | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-IQ2_S |
IQ2_S | 22 / IQ2_M storage | Y | - | - | - | - | INVENTORIED |
reader :229-234; explicit rejection |
leaf open | - |
QUANT-GGUF-IQ4_XS |
IQ4_XS | 23 / output | Y | - | - | - | - | INVENTORIED |
reader :235-240; explicit rejection |
leaf open | - |
QUANT-GGUF-IQ1_M |
IQ1_M | 29 / output | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-BF16 |
BF16 | 30 / output | Y | - | - | - | - | INVENTORIED |
reader :261-264; executable path absent |
leaf open | - |
QUANT-GGUF-TQ1_0 |
TQ1_0 | 34 / output | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-TQ2_0 |
TQ2_0 | 35 / output | - | - | - | - | - | INVENTORIED |
- | leaf open | - |
QUANT-GGUF-MXFP4 |
MXFP4 | 39 / MoE output | Y | - | - | - | - | INVENTORIED |
reader :265-270; executable path absent |
leaf open | - |
QUANT-GGUF-NVFP4 |
NVFP4 | 40 / NVIDIA NVFP4 in a GGUF container (fork/toolchain type id; the real files are compressed-tensors nvfp4-pack-quantized repacked) |
Y | Y | part | - | - | PARTIAL |
block decode reusing the already-gated safetensors NVFP4 numerics (kE2M1Lut / F8E4M3ToF32) so the two containers cannot drift; sidecar + per-expert slab plumbing GgufGlobalScales / DqSlabs. The container is MEASURED, not assumed (spec Sec 5): d[4] fp8-e4m3 sub-block scales then qs[32] in the ggml SPLIT-HALF nibble order (byte j = element j low / j+8 high), and a per-tensor <stem>.scale f32 sidecar TENSOR that is bit-identical to float32(1)/float32(weight_global_scale) on the safetensors side, i.e. these blocks alone do NOT determine the weight. Stacked experts carry ONE scale per expert ([256] on the 35B). The scale-less entry point THROWS for type 40 rather than defaulting to 1.0. Gate = BIT-EXACT cross-format equivalence vs DequantNvfp4ToBf16 test_gguf_nvfp4 on real Qwen3.6-27B bytes from BOTH containers (goldens, regenerable with gen-gguf-nvfp4-goldens.py): 6/6 cases, 2141/2141 assertions. The plausible-but-wrong torch pairwise nibble order was implemented and run FIRST and the gate rejected it on 1686/2048, 1935/2048 and 1678/2048 VALUES, so the gate discriminates the layout rather than merely the symbol. Regression sweep UNCHANGED on the same build (test_gguf 30/103, test_gguf_dequant 13/215, test_gguf_keep_quant 36/5958, test_gguf_qwen36_loader 3/99, test_qwen3_5_gguf_mtp 2/2, test_ops_gdn 58/1825, test_capi 33/232, test_nvfp4_dequant 4/47, test_ct_nvfp4_emulation 6/81, test_model_loader_gguf 3/3, test_ops_quant_traits 8/5615). C stays -: NVFP4 has no vt::DType block encoding and no vec_dot, so RouteGgufTensor correctly expands it. E stays -: dequant alone does not load the 27B/35B models end to end Column C LANDED 2026-07-29 (CLAIM-GGUF-NVFP4-COMPUTE, spec gguf-nvfp4-native-compute), as part. An NVFP4 matmul/expert weight in a verbatim role now takes the new GgufResidency::kNvfp4Fp4 (policy) and is REPACKED at load (RepackGgufNvfp4Rows) into the (weight_packed [N,K/2] torch-pairwise, weight_scale [N,K/16] linear fp8) operand pair, filling the SAME Nvfp4Weight fields the safetensors loader fills (OwnGgufNvfp4), so the ALREADY-GATED vt::MatmulNvfp4 / kMatmulNvfp4Fp4 / kMatmulNvfp4Cutlass kernels run unchanged: NO kernel was written and nothing under src/vt/ was touched. NVFP4 deliberately gets NO vt::DType member - every vt block dtype is self-contained and consumed through a vec_dot, and NVFP4's value needs a scalar from a different TENSOR - so the residency's product is an Nvfp4Weight, not a block-typed tensor. Gate = BYTE-IDENTITY of the repacked operands against the compressed-tensors container's own weight_packed/weight_scale (test_gguf_nvfp4): 11/11 cases, 2207/2207 assertions, ZERO differing bytes; MUTATION-proved (the straight-through packing that preserves every value and scale is REJECTED, and flipping the nibble halves takes the suite to 9/11). Routing totality (test_gguf_keep_quant) 37/37, 5986/5986. Wiring + residency gate test_qwen27_gguf_nvfp4_compute. part, not Y, and the subset is exact: the dense MLP + full-attention q/k/v/o go native and are HARDWARE-GATED on the 27B; the MoE STACKED-EXPERT arm is now HARDWARE-GATED too (2026-07-29, spec native compute Sec D), superseding the UNVERIFIED gap this row previously recorded. Real 35B A3B asset ~/bench/q36-35b-a3b-nvfp4.gguf: its NVFP4 subset is DISJOINT from the 27B's - 120 3-D [256,out,in] routed-expert stacks + 120 2-D shared-expert projections, while its GDN, full-attention and whole MTP (blk.40/nextn.*) families are BF16 - so the two files together cover the column. Weight gate, asset-gated test_gguf_nvfp4: 840 (tensor, expert) slabs (every NVFP4 expert tensor of every layer x experts {0,1,2,7,128,254,255}) repack to the safetensors' own weight/weight_scale with ZERO differing bytes, and all 840 per-expert <stem>.scale[e] are bit-identical to expert e's weight_scale_2 - the per-expert SCALE INDEXING, proven on the real file. Container caveat: this checkpoint is MODELOPT, whose weight_scale_2 is already the multiply form, so the equality is DIRECT here where the 27B's compressed-tensors sibling needs the reciprocal. MUTATION-proved non-vacuous: the scales[0]-for-every-expert mutant AND the expert-0-SLAB-for-every-expert mutant are both constructed and REJECTED; the real file's own scales vary (117/117/138 distinct of 256 on layer 0's three stacks). Loader gate test_qwen36_gguf_nvfp4_compute: 240 of 733 routed tensors take the fp4 residency (120 stacked-expert), 120 fp4 expert stacks over 40 layers with 0 bf16 stacks, 17 280 MiB expert residency, 154 008/154 008 assertions rc=0. That run also FOUND A LATENT DEFECT, and it is fixed in the same change: the first 35B fp4 forward THREW vt: matmul: inner dims mismatch from MoeBlockFusedMarlinCuda - the two fp4 fused MoE blocks issued the ROUTER GEMM as a bare vt::Matmul, assuming the safetensors [K=H,N=E] layout, while the GGUF loader keeps the gate [N=E,K=H] nk=true under expand_nk (default ON now that CUDA registers kMatmulBTQuant). PRE-EXISTING and newly REACHABLE: a GGUF load never produced fp4-resident experts before, so a GGUF router gate never entered those blocks; test_qwen36_paged_engine cannot see it (safetensors, nk=false). Fix = MoeRouterLogits branching on router_gate.nk exactly as MatmulBf16D does, PROVABLY INERT for nk=false. Deterministic once attributed (3/3 throw; 3/3 pass with VT_GGUF_KEEP_QUANT=0, which restores the transposed gate). 35B same-binary A/B post-fix (one flock, idle box, fp4 arm 2 reps identical): peak RSS 68.50 -> 22.72 GiB (3.01x), load-and-generate 1:51.9 -> 0:28.8 (3.9x), and the 24 greedy tokens are IDENTICAL between the arms - which is the CORRECT outcome here and not a no-op: the 35B routed experts run the Marlin W4A16 grouped GEMM in both arms (it consumes scale2, ignores alpha), so token identity is the LOSSLESSNESS signal and 'the fp4 path ran' is carried by the residency, the 3.9x delta, and the fact that only the fp4 arm ever entered the fused block that threw. OPEN and recorded, not explained away: this case's greedy stream is NOT run-to-run stable - the VT_GGUF_NVFP4_W4A4=0 arm gave the canonical stream 2 of 3 runs and a different coherent one once, and the safetensors reference gave one stream 3 of 4 runs and another once (cross-container divergence index 7 three times, 16 once). So the BINDING results of this arm are the weight-level byte identity and the residency/routing audit; the token agreement is reported with that instability attached, never as token-exactness. The fp4 arm did reproduce 3 of 3. test_qwen36_paged_engine is token-exact 315/315 at ITS engine params, so this is a property of this case's block_size 32 / num_blocks 256 / max_num_seqs 1 configuration and NOT of column C; attributing it is OWED WORK and is not claimed. Also: the GDN in_proj_{qkv,z,a,b} family and ssm_out still expand (their V-head reorder rewrites the layout, and only ssm_out even has an fp4 field), and CPU keeps expanding everything (kMatmulNvfp4 is CUDA-only). VT_GGUF_NVFP4_FP4=0 is the same-binary opt-out, VT_GGUF_NVFP4_W4A4=0 selects vLLM's use_a16 W4A16 mode instead of true W4A4. The cross-container framing recorded on this row is CORRECTED: the two 27B 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 their token divergence is a WEIGHT delta and no compute change closes it; that is measured, not assumed (spec Sec A). GB10 sm_121a VERIFIED 2026-07-29 on a build proven production-configured three ways (0 cutlass not found; cuobjdump -lelf 41 cubins ALL sm_121a; SACRED test_qwen27_paged_engine 235/235 exit 0, and test_qwen36_paged_engine 315/315), one flock $HOME/gpu.lock across the series on an idle box: asset-gated test_gguf_nvfp4 11/11, 2784/2784 (full-file sweep compares 192 NVFP4 tensors across both containers), same-binary A/B VT_GGUF_NVFP4_FP4 0 vs 1 with 2 reps each, interleaved - peak RSS 50.8 -> 25.7 GiB, load-and-generate 1:58 -> 0:41, and the 256 moved projections 35 840 -> 10 080 MiB (3.56x, exactly 2 bytes/element against 0.5625); the routing audit shows 256 of 851 routed tensors taking the fp4 residency with 0 left bf16. The cross-container divergence CLOSES: 4 -> none. The fp4 arm is token-IDENTICAL to the safetensors container over 24 greedy tokens (divergence index 24 of 24) where the bf16 arm of the SAME binary diverges at index 4 - so the compute delta WAS the operative cause even though the weight delta above is real. Reported, never gated: identity is not guaranteed by construction and a cross-container throughput arm is not valid. E stays - (this is one prompt against our own engine on a sibling container, not a token-exact oracle for THIS container) and P stays - (the load/residency numbers are ACCEPTED in BENCHMARKS; a serving-throughput arm is PENDING). |
gguf NVFP4 notes, native compute | - |
QUANT-GGUF-Q1_0 |
Q1_0 | 41 / local fork output | Y | - | - | - | - | INVENTORIED |
reader :280-285; executable path absent |
leaf open | - |
| Internal dot/scalar tensor types Q8_1, Q8_K, I8/I16/I32/I64/F64 remain parser | |||||||||||
| concerns, not output-quantization completion rows. Removed GGUF types 4/5, | |||||||||||
| 31-33 and 36-38 are intentionally not roadmap targets. |
This is a non-claimable roll-up under QUANT-GGUF-PRESETS; it prevents
encoding support from being mistaken for mixed-preset coverage. Source:
tools/quantize/quantize.cpp:34-74. The umbrella spike assigns exact stable
preset IDs before any preset enters READY or implementation.
| Preset block | Exact accepted names | Required encodings | Gate state | Spike |
|---|---|---|---|---|
| Scalars | F32, F16, BF16 | matching scalar type | PARTIAL: F32 only, custom APEX files |
leaf specs open |
| Legacy | Q4_0, Q4_1, Q5_0, Q5_1, Q8_0 | matching legacy types | INVENTORIED: no representative llama-quantize preset gate |
leaf specs open |
| K2 | Q2_K, Q2_K_S | Q2_K plus mixed tensor choices | INVENTORIED |
leaf specs open |
| K3 | Q3_K, Q3_K_S, Q3_K_M, Q3_K_L | Q3_K plus mixed tensor choices | INVENTORIED |
leaf specs open |
| K4 | Q4_K, Q4_K_S, Q4_K_M | Q4_K plus mixed tensor choices | INVENTORIED |
leaf specs open |
| K5 | Q5_K, Q5_K_S, Q5_K_M | Q5_K plus mixed tensor choices | INVENTORIED |
leaf specs open |
| K6 | Q6_K | Q6_K plus mixed tensor choices | INVENTORIED |
leaf spec open |
| I1 | IQ1_S, IQ1_M | IQ1_S/IQ1_M plus mixed choices | INVENTORIED |
leaf specs open |
| I2 | IQ2_XXS, IQ2_XS, IQ2_S, IQ2_M | IQ2 variants plus mixed choices | INVENTORIED |
leaf specs open |
| I3 | IQ3_XXS, IQ3_XS, IQ3_S, IQ3_M | IQ3 variants plus mixed choices | INVENTORIED |
leaf specs open |
| I4 | IQ4_NL, IQ4_XS | IQ4 variants plus mixed choices | INVENTORIED |
leaf specs open |
| Ternary/new | TQ1_0, TQ2_0, Q1_0, MXFP4_MOE | matching types | INVENTORIED |
leaf specs open |
COPY is an operation, not a quantization format. Q3_K/Q4_K/Q5_K without a
suffix are aliases for the corresponding _M preset.
Registry source:
${VLLM_SOURCE}/vllm/model_executor/layers/quantization/__init__.py:12-46.
DONE below is deliberately limited to a named gate-model slice.
| ID | Upstream method / scheme | W/A/KV | Backend scope | R | M | C | E | P | State | Our code + tests/evidence | Spike/spec | Owner |
|---|---|---|---|---|---|---|---|---|---|---|---|---|
QUANT-NVFP4-MO-W4A16 |
ModelOpt NVFP4 W4A16 | W4/A16 | CUDA sm121a, 35B experts | Y | Y | Y | Y | Y | DONE |
weights, Marlin dispatch, kernel tests, 35B gate, final gate ledger | Marlin spike; implementation sequence 240d4a1..d6ed102, accepted at MVP closure 83010c7 |
83010c7 |
QUANT-NVFP4-CT-W4A16 |
compressed-tensors NVFP4 nvfp4-pack-quantized (A16 / weight-only) |
W4/A16 | CUDA sm121a, Qwen3-32B dense | Y | Y | Y | Y | - | ACTIVE |
NEW dispatcher dense_nvfp4_gemm.h (MatmulNvfp4W4A16D, resident Marlin repack, fused merged gate_up, Nvfp4W4A16Stats counters) reusing the ALREADY-VALIDATED vendored Marlin W4A16 GEMM vt::MoeGroupedGemmNvfp4Marlin + marlin_repack.h with ZERO new kernel code; append-only loader dense_weight_loaders.h (LoadCtNvfp4W4A16/LoadMergedCtNvfp4W4A16/IsCtNvfp4Projection); probe + fp4 fields in qwen3_weights.cpp / qwen3.h; dispatch in qwen3.cpp + dense_attn_block.h. Tests: loader gate (1603 tensors mapped, BF16 arms provably empty, alpha==0, zero input_global_scale, merged scale2 == 1/max(divisors)), forward doctest (synthetic NVFP4 vs BF16-on-dequantized, max abs delta-logit = 0, bit-identical), SACRED near-tie-robust gate + goldens qwen3_32b_nvfp4a16_greedy (incl. the our_ids.npy / neartie_gap_mnats.npy teacher-forcing evidence) + isolation script qwen3-32b-nvfp4a16-neartie-gap.py. ON-DISK (measured): weight_packed U8 [N,K/2] + weight_scale F8_E4M3 [N,K/16] + weight_global_scale F32 DIVISOR, and NO input_global_scale — that absence is what selects A16. vLLM has NO separate W4A16 scheme class: compressed_tensors.py:696-698 returns CompressedTensorsW4A4Fp4(use_a16=True), and kernels/linear/__init__.py:879-881 then FORCES Marlin, bypassing the capability registry. OBSERVED on sm_121: Using MarlinNvFp4LinearKernel for NVFP4 GEMM. GATE CLOSED 2026-07-21 with PROOF, in two measured stages. Stage 1: vLLM's own greedy is DETERMINISTIC (K=5, 0 multi-valued cells), so STRICT was run first and scored 4/6 prompts / 67/96 tokens — and was NOT loosened on that evidence alone. Stage 2, the ratified TEACHER-FORCING isolation (qwen3-32b-nvfp4a16-neartie-gap.py, + a VT_DUMP_IDS=1 bootstrap in the gate): all 29 divergent positions gap <= 0.0625 nats and 28/29 are EXACTLY 0.0 (our token IS vLLM's own argmax given OUR prefix). The prompt[5] root flip is an EXACT bf16 tie (" moon"/" Moon" bit-identical logprobs) where vLLM's teacher-forced argmax is OUR token and vLLM contradicts its own incremental greedy; the prompt[2] root flip gaps 0.0625 nats against 0.125 nats of vLLM's OWN batch-composition jitter at that same position, and is smaller than the 0.25-nat gap already ratified on the UNQUANTIZED dense Qwen3-4B row; the other 27 are downstream cascade. VERDICT: NOT a W4A16 defect — pre-existing dense-forward bf16 near-tie drift run 64 layers deep, recorded against MODEL-TEXT-qwen3-qwen3-for-causal-lm. The quant path is exonerated four ways (bit-exact CPU proof, fallback_gemms=0, invariance across both quantized GEMMs via the VT_NVFP4_MARLIN=0 A/B which scores 3/6 with the SAME divergences, and the <=0.0625-nat gaps). E = Y under the ratified near-tie-robust bar (kNearTieMnats = 500, identical to the dense Qwen3 / Qwen3-Coder gates) with the evidence COMMITTED as goldens and a hard anchor REQUIRE pinning the band to our exact sequence: 6/6 prompts PASS — strict 4/6 + band 2/6, max gap 0.062 nats, 0 forward-divergent, 142/142 assertions; path proven exercised (marlin_gemms=18432, fused_gate_up=6144, fallback_gemms=0). Loader gate 5144/5144; forward doctest 1031/1031; test_ops_nvfp4_fp4 27002/27002; CUDA -Werror 0-warn; memcheck 0 errors; REGRESSION 27B 235/235 + 35B 315/315 + Coder 138/138 + dense 664/664 + OPT 36/36 ALL UNCHANGED. P is - — SPEED explicitly PENDING, so the row stays ACTIVE, never DONE — anchor tests/vllm/models/test_qwen3_32b_nvfp4a16_load.cpp:115 |
NVFP4A16 spike; breadth-sweep rank 3 (plan B.3) | CLAIM-QUANT-NVFP4-CT-W4A16 |
QUANT-NVFP4-CT-W4A4 |
compressed-tensors NVFP4 | W4/A4 | CUDA sm121a, 27B dense | Y | Y | Y | Y | Y | DONE |
Existing weights, dispatch, quant, tests, and support ledger remain. 3f256ab stays 55/124; W3-I stays default-off and async is neutral. Finalized 179a0fc proves all 12 local ranges and 1,522 steady oracle windows resolve 128 Stream-K 128x64x256 + 80 static-persistent 128x32x256. Diagnostic local/oracle FP4 medians are 52.508720 / 52.734326 ms, so this quantized GEMM family is not the positive residual. Status 9e0143fa…7b57 is complete-diagnostic; support and binding performance are unchanged |
W4A4 notes; MVP closure 83010c7; W3-C spike; W3-F spike; W3-G spike; W3-H spike; W3-I spike |
83010c7 |
QUANT-FP8-MO-STATIC |
ModelOpt static per-tensor FP8 | W8/A8 | CUDA sm121a, 35B projections | Y | Y | Y | Y | Y | DONE |
weights, dispatch, kernel tests, 35B gate, final gate ledger | forward notes; native path 7b682cc/fd2a575, accepted at MVP closure 83010c7 |
83010c7 |
QUANT-FP8-GENERIC |
fp8 + ModelOpt FP8 breadth | static/dynamic, tensor/channel/token/block | all vLLM backends/SM dispatch | part | part | part | part | - | PARTIAL |
gate slice only: loader, kernel tests, 35B gate | leaf spec open | - |
QUANT-FP8-PB-WO |
ModelOpt FP8 per-block weight-only | W8/A16 | capability selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-FP8-PCPT |
ModelOpt FP8 per-channel/per-token | W8/A8 | capability selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-MXFP8-MODELOPT |
ModelOpt MXFP8 | W8/A8 | CUDA/ROCm/XPU dispatch | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-MIXED-MODELOPT |
ModelOpt mixed precision | FP8/NVFP4/MXFP8 groups | per-layer | part | part | part | Y | Y | PARTIAL |
35B FP8+NVFP4 slice: loader, FP8 tests, NVFP4 tests, gate | leaf spec open | - |
QUANT-CT-MXFP4 |
compressed-tensors MXFP4 mxfp4-pack-quantized (group 32, E8M0 block scales, NO global) |
W4/A16 native Marlin mxf4 keep-quant landed (GB10 target; W4A4 cute-dsl crashes sm_121) | CUDA GB10 Marlin W4A16 (E8M0, group_blocks=2); CPU dequant fallback | Y | Y | Y | Y | - | ANCHOR-BACKFILL |
W2 native compute + W3 gates + e2e (row/QUANT-CT-MXFP4, 1c5ee09e): W0 vehicle Yi30/Qwen3-8B-MXFP4 runs on the 0.25.0 oracle; W1 traced FlashInfer-W4A4-selected-but-crashes-on-sm_121 -> Marlin W4A16 is the GB10 target. Native path: generate_kernels.py MXFP4 config (kFE8M0fnu, group_blocks 2) + regenerated instances; MarlinProcessExpertScalesMxfp4 (byte-exact vs vLLM at all shapes); MoeMarlinArgs.{group_size,mxfp4} launcher branch; Nvfp4Weight.{group_size,is_mxfp4} + dense_nvfp4_gemm.h branch + MatmulMxfp4W4A16D; dense_weight_loaders.h MXFP4 loaders; qwen3_weights.cpp detect+load. Gates GREEN: op-level GEMM vs independent CPU dequant 0.36% M=1/M=8 all real shapes (test_ops_moe_grouped.cpp); model-facing MakeLinearMethod->Apply->BuildMarlinDenseResident bad=0 K=4096+12288 (test_linear_method.cpp); e2e 3/4 token-exact vs oracle golden (async-off) (evidence). Default-async degeneration RESOLVED (ROW-SERVE-ASYNC-DENSE-MIRROR, f9c969ae): the pre-existing classic-dense-Qwen3 async bug (device-mirror not wired for qwen3.cpp) is fixed; DEFAULT-config (async ON) e2e on dgx is now coherent + 3/4 token-exact vs the golden (p0/p1/p3 exact; p2 story = oracle-ratified near-tie, teacher-forced max gap 0.0000 nats). p3 formal near-tie gate now RATIFIED. W4 THROUGHPUT bench RAN on GB10 (row/QUANT-CT-MXFP4-BENCH 33e93608, c1/c2/c4/c8x3, production graphed vLLM 0.25.0, oracle VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel -> Marlin W4A16): BELOW-FLOOR (gate NO). Total tok/s ours/vLLM 0.99x@c1 but ~0.91x@c2-c8 (c2 0.911, c4 0.919, c8 0.913); driven by TPOT +10-12%@c2-c8 (parity@c1); TTFT at parity; peak GPU mem 2.6x LESS (28.3 vs 73.7 GiB). ATTRIBUTION (grounded, per-concurrency curve): the gap is a BATCHED-decode cost absent at batch-1 -> the grouped Marlin W4A16 decode GEMM at M=2..8, NOT batch-1 GEMV nor prefill; NOT a ceiling. Refined by nsys c1-vs-c2 (row/QUANT-CT-MXFP4-GLUE): the c2-c8 STEP is the FLASH decode launch (ours grid (batch,hq) over-waves at batch>=2; vLLM applies the FA2 seqlenq_ngroups_swapped decode swap → (batch,kv_heads)), NOT the marlin GEMM (per-shape parity) nor glue. Lever FLIPPED DEFAULT-ON (row/KERNEL-FA2-GQA-SWAP-FLIP, VT_FA2_DECODE_GQA_SWAP, =0 opts out): vLLM's decode group-swap in the d128 varlen launcher. Flip campaign on GB10 (HEAD 1f446fd7): op 280 assn GREEN + memcheck 0-err; SACRED Qwen3-0.6B/4B greedy near-tie gate 16/16 both, TOKEN-IDENTICAL swap-ON vs plain at default/=1/=0 + #44 smoke token-exact at the new default; nsys ours c2 decode-flash grid (1,3,64)=b×q_heads → (1,5,16)=b×kv_heads, per-call 63.7→45.3us (default no-env reproduces the swap grid). Binding q3mxfp4 grid swap-ON total tok/s c1 0.990 / c2 0.922 / c4 0.930 / c8 0.942 (was 0.989/0.911/0.919/0.913; c2-c8 gains OUTSIDE per-rep noise — each ours-rep beats every swap-OFF rep — c1 flat, NO regression), median TTFT at/above parity, peak GPU mem 2.614x LESS. Flipped per parity-enablers (improves + no regression + correctness holds). Still <1.0x on tput/TPOT: flash closed ~12%@c2 / ~30%@c8 of the gap (below #47's ~28/55% projection since actual flash drop was −29% not −34%); residual = grouped-Marlin decode +7-9% (E1 indirect gather vs vLLM dense direct-A) + ~0.7ms host/sched. No single lever ≥1.0x. Harness (all additive; 27/35 byte-unchanged; CPU tests 45/45): online_gate q3mxfp4 key + POINTS_BY_MODEL/points_for, mxfp4_smoke_gate.py (#44 battery), dgx-online-serving.sh q3mxfp4 branches, mxfp4-online-serving-grid.sh. Full table + build contract (RelWithDebInfo + oracle flashinfer-cutlass + DISK build) in benchmark-record. Earlier CPU weight unpack + E8M0 dequant: NEW mxfp4_dequant.h + .cpp (E8M0ToF32 = 2^(byte-127), DequantMxfp4ToBf16/ToF32, group 32, no global; reuses kE2M1Lut). Unit gate test_mxfp4_dequant.cpp — E8M0 known-byte decode, hand-computed 32-group dequant (bf16+f32), the E8M0-vs-fp8 + group-32-vs-16 RED traps, multi-row/group offsets, randomized rel-error vs a double-precision port of dq_mxfp4_torch with bf16==f32 exactness. CPU -Werror 0-warn. Ports FROM compressed_tensors_w4a4_mxfp4.py:20-97 + mxfp8_utils.py:61-65,222 + golden tests/quantization/reference_mxfp4.py:28-117. C/E/P PENDING: GPU W4A4 fp4 GEMM + Marlin W4A16 fallback + MoE expert path + e2e are NAMED later bricks; DeepSeek-V4 + Kimi-K3 loaders consume this once wired |
MXFP4 spike | CLAIM-QUANT-MXFP4 |
QUANT-CT-W4A8-FP8 |
compressed-tensors W4A8 FP8 | W4/A8 | CUTLASS | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-CT-W4A8-INT8 |
compressed-tensors W4A8 INT8 | W4/A8 | platform selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-CT-W8A8-FP8 |
compressed-tensors W8A8 FP8 | W8/A8 | platform selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-CT-W8A8-INT8 |
compressed-tensors W8A8 INT8 | W8/A8 | platform selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-CT-W8A8-MXFP8 |
compressed-tensors W8A8 MXFP8 | W8/A8 | platform selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-CT-W8A16-FP8 |
compressed-tensors FP8 weight-only | W8/A16 | platform selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-CT-WNA16 |
compressed-tensors WnA16 | symmetric 2-8b; asymmetric 4/8b | Marlin/Machete/etc. | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-CT-WNA8O8 |
compressed-tensors WnA8O8 | W2/4/8, static int8 A/O | platform selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-AWQ |
AWQ / AutoAWQ / AWQ-Marlin | W4/A16 | Marlin/CUDA/CPU variants | - | part | - | - | - | ANCHOR-BACKFILL |
W0 spike + W1 CPU dequant landed. INT4 unpack+dequant to bf16 mirroring awq_triton.py:11-105 (reverse-AWQ order [0,4,1,5,2,6,3,7], (w-z)*s): DequantAwq4ToBf16 / header. Unit gate test_awq_gptq_dequant: hand-computed known int32 (independent arithmetic oracle, col7=15 MSB sign-safety) + randomized double-precision layout roundtrip. M=part (4-bit dequant primitive only; loader R, Marlin C, e2e E are named later bricks — GPU rides the vendored Marlin, W4). |
spike | CLAIM-QUANT-AWQ-GPTQ |
QUANT-GPTQ |
GPTQ / AutoGPTQ / GPTQ-Marlin | W2-8/A16 | Marlin/CUDA/CPU variants | - | part | - | - | - | ANCHOR-BACKFILL |
W0 spike + W1 CPU dequant landed. INT4 unpack+dequant to bf16 mirroring qdq_4.cuh (dequant_4bit_8_gptq) + q_gemm.cu:201-202 (zero_offset v1=1/v2=0), optional act-order g_idx: DequantGptq4ToBf16 / header. Unit gate test_awq_gptq_dequant: hand-computed zero_offset v1/v2 + act-order g_idx group-selection + randomized double-precision roundtrip. M=part (4-bit only; 8/2/3-bit W5, loader R, Marlin C, e2e E are named later bricks). |
spike | CLAIM-QUANT-AWQ-GPTQ |
QUANT-BNB |
bitsandbytes | FP4/NF4/INT8 | CUDA library path | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-EXPERTS-I8 |
experts-int8 | MoE int8 | platform selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-MOE-WNA16 |
MoE WnA16 | W4/8-A16 | Marlin/Machete/ROCm | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-MXFP4-GPTOSS |
MXFP4 / GPT-OSS MXFP4 | W4/A16 or W4/A4 | model/kernel selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-QUARK |
AMD Quark | mixed int/fp | ROCm | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-TORCHAO |
TorchAO | multiple | compatibility path | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-INC |
Intel INC / AutoRound | WnA16 etc. | CPU/XPU | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-HUMMING |
Humming | weight-only | CUDA | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-DSV4-FP8 |
DeepSeek-V4 FP8 | model-specific | CUDA | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-ONLINE |
online FP8/MXFP8/int8 shorthands | runtime W/A | platform selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open | - |
QUANT-DEPRECATED |
FBGEMM-FP8, FPQuant | compatibility only | platform selected | - | - | - | - | - | INVENTORIED |
- | leaf spec open; do not prioritize | - |
Pinned vLLM source: vllm/config/cache.py:19-36.
| ID | Item | Upstream | Our code | Tests/evidence | Spike/spec | State | Owner |
|---|---|---|---|---|---|---|---|
QUANT-KV-FP8 |
fp8, fp8_e4m3, fp8_e5m2 | vllm/config/cache.py:19-25; vllm/model_executor/layers/quantization/kv_cache.py:42-191; store cache_kernels.cu:241-252; scale convention quant_utils.cuh:296-308 |
W1 CPU fp8-e4m3 store+read LANDED: codec, store kernel, read dequant, config parse. e5m2 CPU compute + per-head scales + CUDA + runner integration are named later bricks (see spec) | test_ops_fp8_kv_cache — 8 cases / 511 assertions, round-trip within the e4m3 band + fp8-vs-bf16 NMSE<1% + paged-attention e2e; RED-first (wrong store direction fails 3/480) | fp8-kv-cache | PARTIAL |
- |
QUANT-KV-FP8-VENDOR |
fp8_inc, fp8_ds_mla | vllm/config/cache.py:24-25; vendor KV implementations selected by attention backend |
- | no quantized KV cache | planned: specs/vendor-fp8-kv-cache.md |
INVENTORIED |
- |
QUANT-KV-TURBO |
k8v4, 4bit_nc, k3v4_nc, 3bit_nc | vllm/config/cache.py:28-33; TurboQuant dependency path |
- | no quantized KV cache | planned: specs/turboquant-kv-cache.md |
INVENTORIED |
- |
QUANT-KV-PER-HEAD |
int4/int8/fp8 per-token-head | vllm/config/cache.py:34; quantized cache kernels selected by backend |
- | no quantized KV cache | planned: specs/per-head-kv-cache.md |
INVENTORIED |
- |
QUANT-KV-NVFP4 |
NVFP4 KV | vllm/config/cache.py:35; CUDA FP4 cache kernel family |
- | no quantized KV cache | planned: specs/nvfp4-kv-cache.md |
INVENTORIED |
- |
These rows are required for the Apple backend, independent of CUDA/vLLM serialization. Correctness is against MLX/MLX-LM; performance is against oMLX and MLX-LM on the same Mac, model, quant and concurrency.
| ID | Item / required matrix | Upstream | Our code | Tests/evidence | Spike/spec | State | Owner |
|---|---|---|---|---|---|---|---|
QUANT-MLX-AFFINE-Q2 |
affine 2-bit; groups 32/64/128 | MLX quantized.py:11-95; ops.cpp:4624-4693 |
- | upstream test_quantized.py:11-170; local absent |
planned: specs/mlx-affine-quant.md |
INVENTORIED |
- |
QUANT-MLX-AFFINE-Q3 |
affine 3-bit; groups 32/64/128 | same MLX primitive/module path | - | upstream test_quantized.py:11-170; local absent |
planned: specs/mlx-affine-quant.md |
INVENTORIED |
- |
QUANT-MLX-AFFINE-Q4 |
affine 4-bit; groups 32/64/128; Q4-G64 first | same MLX primitive/module path | - | upstream test_quantized.py:11-170; local absent |
planned: specs/mlx-affine-quant.md |
INVENTORIED |
- |
QUANT-MLX-AFFINE-Q5 |
affine 5-bit; groups 32/64/128 | same MLX primitive/module path | - | upstream test_quantized.py:11-170; local absent |
planned: specs/mlx-affine-quant.md |
INVENTORIED |
- |
QUANT-MLX-AFFINE-Q6 |
affine 6-bit; groups 32/64/128 | same MLX primitive/module path | - | upstream test_quantized.py:11-170; local absent |
planned: specs/mlx-affine-quant.md |
INVENTORIED |
- |
QUANT-MLX-AFFINE-Q8 |
affine 8-bit; groups 32/64/128 | same MLX primitive/module path | - | upstream test_quantized.py:11-170; local absent |
planned: specs/mlx-affine-quant.md |
INVENTORIED |
- |
QUANT-MLX-MXFP4 |
E2M1 + E8M0; group 32 | MLX MX quantization mode in ops.cpp:4624-4693 |
- | upstream test_quantized.py MX cases; local absent |
planned: specs/mlx-mx-quant.md |
INVENTORIED |
- |
QUANT-MLX-MXFP8 |
E4M3 + E8M0; group 32 | same MLX MX primitive path | - | upstream test_quantized.py MX cases; local absent |
planned: specs/mlx-mx-quant.md |
INVENTORIED |
- |
QUANT-MLX-NVFP4 |
E2M1 + E4M3 scale; group 16; optional global scale | MLX NVFP4 mode in ops.cpp:4624-4693 |
- | upstream test_quantized.py NVFP4 cases; local absent |
planned: specs/mlx-nvfp4-quant.md |
INVENTORIED |
- |
QUANT-MLX-QQ-NVFP4 |
quantized activations + weights; tensor-global restriction separate | MLX quantized-quantized matmul implementation | - | upstream test_quantized.py:227-251; local absent |
planned: specs/mlx-qq-quant.md |
INVENTORIED |
- |
QUANT-MLX-QQ-MXFP8 |
quantized activations + weights; supported QQ shapes | MLX quantized-quantized matmul implementation | - | upstream test_quantized.py QQ cases; local absent |
planned: specs/mlx-qq-quant.md |
INVENTORIED |
- |
QUANT-MLX-MIXED |
mixed_2_6, mixed_3_4, mixed_3_6, mixed_4_6 per-layer recipes | MLX-LM convert.py:20-227 at a790972f |
- | upstream MLX-LM conversion tests; local absent | planned: specs/mlx-mixed-quant.md |
INVENTORIED |
- |
QUANT-MLX-FINE |
per-layer bits/group/mode plus skip rules | MLX-LM convert.py:20-227 at a790972f |
- | upstream MLX-LM conversion tests; local absent | planned: specs/mlx-fine-grained-quant.md |
INVENTORIED |
- |
QUANT-MLX-IMPORT |
AWQ/GPTQ 4b, legacy MXFP4, compressed-tensors, BitNet transforms | MLX-LM compatibility transforms at a790972f |
- | upstream MLX-LM conversion tests; local absent | planned: specs/mlx-quant-import.md |
INVENTORIED |
- |
- Preserve the three gate-specific CUDA rows and generalize their dispatch.
- Make GGUF compute-in-quant the speed foundation; loader-only breadth cannot satisfy llama.cpp parity.
- Close common GGUF and APEX blocker encodings, then IQ/TQ/Q1 breadth.
- Close generic vLLM FP8/MX/KV formats.
- Bring up MLX affine Q4-G64, expand the native matrix, then oMLX gates.
- Port AWQ/GPTQ/CT integer and the remaining registry methods.