From acd7d457d6b276dbf659ce44a742612f6f5602b4 Mon Sep 17 00:00:00 2001 From: Justin Card Date: Sat, 5 Sep 2026 22:40:42 -0400 Subject: [PATCH 01/11] feat(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): port IQ4_XS/IQ3_XXS to ROCm Issue #1940 says src/vt/rocm/ ports zero I-quant formats, so IQ4_XS and IQ3_XXS still expand to bf16 on a ROCm keep-quant load instead of staying compressed. This lands DotIQ4XS and DotIQ3XXS in rocm_grouped_gemm.hip, adapted from the existing CUDA bodies rather than freshly ported, on both the plain and grouped/MoE arms, and admits both dtypes in DeviceKeepQuantSupported's ROCm case. The spec's FMA-contraction risk is resolved by measurement, not assumption. CUDA's DotIQ4XS needed non-fused float intrinsics to stay bit-exact against the oracle, because nvcc's project-wide flag never reaches .cu files. HIP already compiles with -ffp-contract=off project-wide, and a new bit-exact gate over the same real checkpoint bytes and the same oracle bits CUDA's gate uses confirms plain arithmetic is sufficient on this toolchain, so the port carries no workaround CUDA needed. Three existing tests asserted the prior ROCm behavior by name and go red under the new routing, each for the reason it should: the cross-device table gains both dtypes and a new bit-exact case, the keep-quant loader's exhaustive per-device count moves from 8 to 10, and a residency pin that named issue #1940 splits per-tensor now that one of its two towers keeps. Run on isravale (RX 9060 XT, gfx1200, ROCm 7.2.3) under the GPU file mutex: test_backend_cross_device 41/42 (the one failure, MoeSiluMul, is confirmed pre-existing and unrelated via an independent binary built from the sibling KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA4-w1 worktree reproducing the identical mismatch); test_gguf_keep_quant 52/52; test_gguf_device_fit 24/24; every other ROCm suite green with zero regression. The row's actual acceptance criterion also landed: an end-to-end reload of the real motivating checkpoint, Nail-Qwen3.6-35B-A3B-MTP-IQ4_XS.gguf (19.39 GB, on isravale). Its own header histogram is IQ4_XS/Q5_K/Q6_K only, so it exercises exactly the dtypes this change adds plus the two already supported. It loads, generates coherent tokens, and its resident footprint (~19.33 GiB) matches the on-disk size instead of the ~70 GiB a bf16 expansion would produce -- the fix for the host-RAM SIGSEGV this checkpoint hit before this row, on the artifact that motivated it. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [Claude Code] --- .../kernel-quant-ciq-gemm-rocm-iquant.md | 188 +++++++++++++----- .../model_loader/gguf_keep_quant.cpp | 12 +- src/vt/rocm/rocm_grouped_gemm.hip | 163 +++++++++++++-- src/vt/rocm/rocm_quant_iq_tables.h | 100 ++++++++++ .../model_executor/test_gguf_device_fit.cpp | 33 ++- tests/vllm/test_gguf_keep_quant.cpp | 35 +++- tests/vt/test_backend_cross_device.cpp | 155 ++++++++++++++- 7 files changed, 593 insertions(+), 93 deletions(-) create mode 100644 src/vt/rocm/rocm_quant_iq_tables.h diff --git a/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md b/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md index 91054c63c6..1f0ffedce1 100644 --- a/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md +++ b/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md @@ -12,7 +12,8 @@ 2026-09-05). - Pull request shape: separate spec and implementation pull requests (developer decision 2026-09-05, recorded in - `.agents/developer-preferences.md`). This pull request lands the spec only. + `.agents/developer-preferences.md`). The spec landed in its own pull + request first; this implementation is the second. ## Scope @@ -155,26 +156,31 @@ llama.cpp, pin `b10451` per `.agents/upstream-sync.md`. ## Risks -- **FMA contraction on IQ4_XS's float-accumulation body.** IQ4_XS's dot is - the one format in this row (and in the whole quant-dot family) whose core - is not a single integer accumulator: it forms `d1`/`d2` as f32 and folds - in per-sub-block `sumf +=` steps, eight per super-block - (`cuda_quant_dot.cu:606-680`, extensively commented on exactly this - point). On CUDA that required `__fmul_rn`/`__fadd_rn` in place of ordinary - `*`/`+`, because nvcc's default `-fmad=true` silently contracts the - textual two-rounding sequence into a single-rounding FMA and two of eight - real super-blocks then disagreed with the oracle by 1-4 ULP. **This may - not reproduce on ROCm**: `CMakeLists.txt:414` already applies - `-ffp-contract=off` to `$` project-wide, unlike CUDA - where the project's `-ffp-contract=off` is CXX-only and never reaches - `.cu`/`.cuh` translation units. Verify this empirically before assuming it - (a W0-style probe: compile the naive `sumf += d1 * x` form, diff against - the CPU oracle on the same real super-blocks CUDA's golden vectors use, - and inspect the generated ISA for `v_fma_f32` if any block disagrees) — - do not carry the CUDA workaround over unexamined, and do not assume the - flag alone is sufficient without a measured check, matching how the CUDA - side only added the intrinsics after measuring a real disagreement rather - than as a precaution. +- **FMA contraction on IQ4_XS's float-accumulation body — MEASURED, RESOLVED + IN FAVOR OF THE SIMPLER PATH.** IQ4_XS's dot is the one format in this row + (and in the whole quant-dot family) whose core is not a single integer + accumulator: it forms `d1`/`d2` as f32 and folds in per-sub-block + `sumf +=` steps, eight per super-block (`cuda_quant_dot.cu:606-680`, + extensively commented on exactly this point). On CUDA that required + `__fmul_rn`/`__fadd_rn` in place of ordinary `*`/`+`, because nvcc's + default `-fmad=true` silently contracts the textual two-rounding sequence + into a single-rounding FMA and two of eight real super-blocks then + disagreed with the oracle by 1-4 ULP. **W0/W1 measured this directly on + the target hardware (RX 9060 XT, gfx1200, ROCm 7.2, `isravale`):** plain + `*`/`+` (no non-fused intrinsics) in `DotIQ4XS` is BIT-EXACT against the + oracle's own per-super-block numbers, over the SAME four real + `unsloth/GLM-5.3-Flash-GGUF` super-blocks and the SAME expected bits + (`iq2xs_iq4xs_dot_golden.h`) CUDA's gate uses, both isolated (k=256, one + contributing lane, zero reassociation possible) and combined (k=1024, + four lanes, the same `__shfl_down_sync` tree CUDA's comment derives) — + `tests/vt/test_backend_cross_device.cpp`, "ROCm IQ4_XS dots the ORACLE's + own numbers on REAL checkpoint bytes", 13/13 assertions green. The + hypothesis held: `CMakeLists.txt:414`'s project-wide + `-ffp-contract=off` on `$` is sufficient on its own, + so `DotIQ4XS` on ROCm uses plain `*`/`+` and does **not** carry CUDA's + `__fmul_rn`/`__fadd_rn` workaround. This is a measured result, not an + assumption carried over — the whole point of naming this as a risk was to + force the check rather than inherit the CUDA fix by habit. - **The nwarps=8 decode table (`ROCM-KQUANT-NWARPS-DECODE`) may not transfer.** Both new formats share the existing `nsb = K/256` decomposition, so they compile against the same launch shape as Q4_K/Q5_K/Q6_K with no code @@ -207,31 +213,88 @@ llama.cpp, pin `b10451` per `.agents/upstream-sync.md`. ## Tests -- Extend `test_ops_quant_dot.cpp`'s existing IQ4_XS/IQ3_XXS `vec_dot` - golden-vector gates (`iq2xs_iq4xs_dot_golden.h`, already committed and - sourced from real `unsloth/GLM-5.3-Flash-GGUF` checkpoint bytes) to a new - `test_rocm_quant_dot.cpp`, same shape as the CUDA gate - (`test_cuda_quant_dot.cpp`): NMSE ≤ 5e-4 vs the independent f64 - dequant-then-dot reference for IQ3_XXS; bit-exact (not NMSE) for IQ4_XS - against the same real-checkpoint golden values CUDA's gate uses, since - bit-exactness is the property the FMA-contraction risk above is actually - about. -- `test_backend_cross_device.cpp`: add both formats to the CPU-vs-ROCM - cross-check, NMSE ≤ 5e-4 (matching the existing four formats' gate shape - there). -- Rerun `ROCM-KQUANT-NWARPS-DECODE`'s own measurement recipe - (`rocprofv3 --kernel-trace` on a real quant-matched trace workload) for - IQ4_XS/IQ3_XXS specifically, to answer the nwarps question this issue was - filed to test — record the result (transfers / does not transfer) rather - than assuming either. -- `ctest -R 'rocm|cross_device'`, zero regression on the four existing - formats' numerics. -- End-to-end: reload the motivating checkpoint (or a same-format synthetic - fixture if the real 35B-A3B artifact is not staged on the gate host) on - `isravale` (RX 9060 XT, gfx1200) or an `rc`-leased ROCm fleet device, and - confirm keep-quant residency replaces the prior bf16 SIGSEGV — this is - the row's actual acceptance criterion, not merely the unit-level dot - gates. +Landed, on `isravale` (RX 9060 XT, gfx1200, ROCm 7.2.3), GPU work under +`flock ${GPU_LOCK:-$HOME/gpu.lock}` throughout: + +- **`test_backend_cross_device.cpp`**, three cases touched/added, run + standalone and as part of the full file (41/42 cases, 83998/83999 + assertions — the one failure is `MoeSiluMul matches the CPU oracle within + NMSE <= 5e-4`, confirmed PRE-EXISTING and unrelated: byte-identical + mismatch reproduced on an independent binary built from the sibling + `KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA4-w1` worktree, which touches neither this + kernel nor this dtype): + - "non-grouped keep-quant GEMM (...IQ4_XS/IQ3_XXS) matches the CPU + oracle" — both new formats added to the existing table-driven CPU-vs- + ROCm case, NMSE ≤ 5e-4, random valid blocks (unconstrained lookup + indices need no in-range fixture change). + - "grouped quant expert GEMM (...IQ4_XS/IQ3_XXS) matches the CPU oracle" + — same extension on the grouped/MoE path, the one the motivating + checkpoint's routed experts actually use. + - "ROCm IQ4_XS dots the ORACLE's own numbers on REAL checkpoint bytes" + (NEW) — the bit-exact gate the FMA-contraction risk needed, ported from + `test_cuda_quant_dot.cpp`'s `CheckCudaOracleDot` shape onto the same + golden vectors: bit-exact per-superblock (k=256, one contributing lane) + and warp-reduction-order-exact combined (k=1024, four lanes, primary + bit-equality + secondary reassociation-bound check). 13/13 assertions. +- **`test_gguf_keep_quant.cpp`**: the exhaustive per-device totality table's + hand-mirrored ROCm predicate and its `gemm_kept` constant (8 → 10) updated + to admit IQ4_XS; IQ3_XXS is not in this test's `all_types` enumeration + (a pre-existing gap shared with Q2_K, not closed by this row) and is left + to the cross-device gate above. 52/52 cases, 10325/10325 assertions. +- **`test_gguf_device_fit.cpp`**: `#2516`'s two ROCm residency pins split + per-tensor (IQ4_XS's `down_exps` now expects `kKeepQuant` on ROCm; + IQ2_XS's `gate_exps` is unaffected and still expects `kExpandBf16`, + since #1940's other five formats stay owed); the all-or-nothing + "NO PLAN" case is unchanged in outcome (`CHECK_FALSE` still holds, because + the still-unsupported IQ2_XS tower alone fails the lane) with its comment + corrected to say why. 24/24 cases, 182/182 assertions. +- `ctest -R 'rocm|cross_device'` (plus the individually-run ROCm suites + `test_rocm_arch`/`test_rocm_backend`/`test_exl3_rocm`/ + `test_gemma4_rocm_fp8_seams`/`test_rocm_fp8_kv_cache`): zero regression, + all green. + +**End-to-end reload — the row's actual acceptance criterion — LANDED.** +`Nail-Qwen3.6-35B-A3B-MTP-IQ4_XS.gguf` (19.39 GB on disk, `isravale` +`/home/justin/Nail/`) is the real motivating checkpoint, not a stand-in: its +own header histogram is `{BF16: 2, F32: 308, IQ4_XS: 391, Q5_K: 51, +Q6_K: 1}` — every quantized tensor in the file is one of the three dtypes +this row's target hardware now has a keep-quant kernel for (read with +`docs/bench-evidence/limb3-vehicle-search-20260904/gguf_header.py` before +running anything, not assumed from the filename). + +``` +VT_DEVICE_WEIGHT_BUDGET_BYTES=13000000000 \ +./build-hip/examples/vllm-cli --model /home/justin/Nail/Nail-Qwen3.6-35B-A3B-MTP-IQ4_XS.gguf \ + --device auto --max-num-seqs 1 --kv-cache-dtype fp8 --kv-cache-memory 2000000000 \ + --prompt "The capital of France is" --max-tokens 16 +``` + +``` +engine: device placement INSTALLED: 15 layers run their routed experts on cpu, the rest on rocm (resolved against 40 layers, origin fit) +engine: device placement: --fit placed 15 layer(s) (6417285120 B) to bring a 19333564672 B footprint under a 13000000000 B budget +vllm-cli: run=1/1 finish_reason=length prompt_tokens=5 completion_tokens=16 secs=3.919 tok_s=4.083 + Paris. The capital of Germany is Berlin. The capital of Italy is Rome. +``` + +The decisive number is the **19,333,564,672 B (~19.33 GiB) footprint** -- +it matches the file's on-disk size, not the ~70 GiB a bf16 expansion of +these tensors would produce. That is the keep-quant residency actually +taking effect on ROCm, not merely compiling: before this row, +`DeviceKeepQuantSupported` routed every IQ4_XS tower to `kExpandBf16` here +and the streamed-expert lane's blow-up SIGSEGV'd this same box on this +family of checkpoint (`vllm-cpp-rocm-crash-iq4xs` session memory). Clean +exit, coherent completion, zero crash. Not a synthetic fixture, not a +narrower stand-in geometry -- the actual artifact the row exists for. + +Not done in this wave (see Owed): + +- The `ROCM-KQUANT-NWARPS-DECODE` re-measurement (`rocprofv3 --kernel-trace` + on a real quant-matched trace workload) — this issue's own stated reason + for existing beyond plain coverage. `isravale` has no `rocprofv3` profiling + set up in this session; the correctness gates above stand on their own, + but the nwarps question is still open. The 4.083 tok/s figure above is NOT + a substitute measurement for it: it is a mixed CPU+ROCm run at a + CPU-offload-heavy split, not an isolated ROCm-kernel throughput number. ## Owed @@ -241,9 +304,10 @@ llama.cpp, pin `b10451` per `.agents/upstream-sync.md`. - A WMMA/tensor-core tile for IQ4_XS/IQ3_XXS, if the scalar tier's measured throughput warrants one (mirroring how `KERNEL-QUANT-CIQ-GEMM-ROCM` followed the existing four formats' scalar tier): not attempted here. -- The nwarps re-measurement itself, if it is not completed within this - row's implementation wave for lack of GPU time: record as `PENDING` on a - named lease/box, never silently dropped. +- The nwarps re-measurement itself: `PENDING`, not completed within this + implementation wave for lack of a `rocprofv3` profiling setup on + `isravale`, not silently dropped. The correctness gates (Tests) are + unaffected by this being open. ## Stop conditions @@ -260,8 +324,24 @@ llama.cpp, pin `b10451` per `.agents/upstream-sync.md`. ## Now -`SPIKE`. This pull request lands the spec only; no product code changes in -this change. Next: W0 probes the FMA-contraction question on target -hardware (gfx1200), then W1 ports `DotIQ4XS` (the harder, float-accumulation -body) and W2 ports `DotIQ3XXS`, each with its own focused gate before the -combined `ctest` sweep. +`ACTIVE`. W0 (FMA-contraction probe), W1 (`DotIQ4XS`) and W2 (`DotIQ3XXS`) +are LANDED in this pull request, on both the plain (`MatmulBTQuantKernelRocm` +/ `KQuantGemmK`) and grouped/MoE (`MatmulBTQuantGroupedKernelRocm` / +`GroupedKQ8K`) arms, plus `DeviceKeepQuantSupported`'s ROCm admission list. +Gated per the Tests section above, on target hardware (`isravale`, +RX 9060 XT / gfx1200), zero regression. The FMA-contraction risk resolved in +favor of the simpler path: HIP's project-wide `-ffp-contract=off` is +sufficient, no CUDA-style non-fused-multiply workaround needed. + +**The real-checkpoint end-to-end reload also LANDED**, after this pull +request was first drafted: `Nail-Qwen3.6-35B-A3B-MTP-IQ4_XS.gguf` loads and +generates coherent tokens on `isravale`, with the resident footprint +(~19.33 GiB) matching the on-disk size rather than a bf16 blow-up — see +Tests. That was the row's actual acceptance criterion, and it is now +satisfied on the artifact that motivated the row, not a synthetic +stand-in. + +Remaining before `DONE`: only the `ROCM-KQUANT-NWARPS-DECODE` +re-measurement (`PENDING`, see Owed) — it does not block this pull request, +since the row's own scope is coverage and correctness, and it is named +rather than silently dropped. diff --git a/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp b/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp index 65cc5db456..ea5a2fdee9 100644 --- a/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp +++ b/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp @@ -137,10 +137,16 @@ bool DeviceKeepQuantSupported(vt::DType dt, vt::DeviceType dev) { switch (dev) { case vt::DeviceType::kROCM: // src/vt/rocm/rocm_grouped_gemm.hip implements exactly these on both the - // grouped and non-grouped arms; Q4_0/Q2_K/Q3_K/IQ2_*/IQ3_*/MXFP4 are - // owed (recorded in .agents/specs/rocm-gg-keep-quant.md). + // grouped and non-grouped arms. IQ4_XS/IQ3_XXS added by + // KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940) -- this is the fix for the + // host-RAM SIGSEGV a real IQ4_XS checkpoint hit on ROCm before this row: + // these two dtypes used to fall through to `false` here and keep the + // pre-existing expand-bf16 residency this comment block warns against. + // Q4_0/Q2_K/Q3_K/IQ2_XXS/IQ2_S/MXFP4 are still owed (recorded in + // .agents/specs/rocm-gg-keep-quant.md and tracked by the same issue). return dt == vt::DType::kQ8_0 || dt == vt::DType::kQ4_K || - dt == vt::DType::kQ5_K || dt == vt::DType::kQ6_K; + dt == vt::DType::kQ5_K || dt == vt::DType::kQ6_K || + dt == vt::DType::kIQ4_XS || dt == vt::DType::kIQ3_XXS; default: // CUDA falls back to the CPU kernel for anything it lacks // (cuda_quant_dot.cu:1841-1846); the CPU list IS the CPU capability. diff --git a/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index a6239e70d9..911a377876 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -34,6 +34,9 @@ // Block layouts — the single source of truth (ggml-common.h mirrors). #include "vt/cpu/cpu_quant_blocks.h" +// KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT: the IQ4_XS/IQ3_XXS device codebooks, +// generated from the same source the CPU/CUDA tables read (see the header). +#include "vt/rocm/rocm_quant_iq_tables.h" namespace vt::rocm { namespace { @@ -43,6 +46,8 @@ using vt::cpu::BlockQ8_K; using vt::cpu::BlockQ4_K; using vt::cpu::BlockQ5_K; using vt::cpu::BlockQ6_K; +using vt::cpu::BlockIQ4_XS; +using vt::cpu::BlockIQ3_XXS; using vt::cpu::kQK8_0; using vt::cpu::kQK_K; @@ -320,6 +325,98 @@ __device__ inline float DotQ6K(const BlockQ6_K* xb, const BlockQ8_K* yb) { return d * isum; } +// IQ3_XXS x Q8_K: cuda_quant_dot.cu DotIQ3XXS (KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT, +// #1940). Ported verbatim -- no CUDA-only intrinsic in this body (no __dp4a, +// no __vcmpne4/__vsub4; those belong to DotIQ2XXS, not this one), so the port +// is a straight recompile against this file's own tables. A per-tensor +// FinalFactor of 0.25 applies ONCE after the warp reduction sums every +// super-block's `d * bsum`, not inside this function -- see KQuantGemmK. +__device__ inline float DotIQ3XXS(const BlockIQ3_XXS* xb, const BlockQ8_K* yb) { + const float d = DF16ToF32(xb->d) * yb->d; + const uint8_t* q3 = xb->qs; + const uint8_t* gas = xb->qs + kQK_K / 4; + const int8_t* q8 = yb->qs; + int32_t bsum = 0; + for (int ib32 = 0; ib32 < kQK_K / 32; ++ib32) { + uint32_t a32; + memcpy(&a32, gas, sizeof(uint32_t)); + gas += sizeof(uint32_t); + const uint32_t ls = 2 * (a32 >> 28) + 1; + int32_t sumi = 0; + for (int l = 0; l < 4; ++l) { + const uint32_t g1 = d_iq3xxs_grid[q3[2 * l + 0]]; + const uint32_t g2 = d_iq3xxs_grid[q3[2 * l + 1]]; + const uint8_t signs = d_ksigns_iq2xs[(a32 >> (7 * l)) & 127]; + for (int j = 0; j < 4; ++j) { + const int b1 = static_cast((g1 >> (8 * j)) & 0xff); + const int b2 = static_cast((g2 >> (8 * j)) & 0xff); + sumi += b1 * q8[j + 0] * ((signs & d_kmask_iq2xs[j + 0]) ? -1 : 1); + sumi += b2 * q8[j + 4] * ((signs & d_kmask_iq2xs[j + 4]) ? -1 : 1); + } + q8 += 8; + } + q3 += 8; + bsum += sumi * static_cast(ls); + } + return d * bsum; // final *0.25 applied after the warp reduction +} + +// IQ4_XS x Q8_K: cuda_quant_dot.cu DotIQ4XS (KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT, +// #1940). IQ4_XS shares IQ4_NL's 16-entry d_kvalues_iq4nl codebook and NOT its +// activation pairing: its block is a 256-element SUPER-block, so it dots +// Q8_K. The `ib` loop steps by TWO because one `scales_l` byte serves two +// sub-blocks; `d1`/`d2` are formed as f32 BEFORE the integer sums fold in, so +// a super-block is EIGHT f32 accumulation steps -- upstream's association, +// kept rather than hoisted into one integer accumulator. FinalFactor +// is 1 (the per-sub-block delta is already folded in as `d1`/`d2` here). +// +// CUDA's port of this body needed __fmul_rn/__fadd_rn in place of `*`/`+`, +// because nvcc's default -fmad=true silently contracts the two-rounding +// sequence into a single-rounding FMA and measurably disagreed with the +// oracle by 1-4 ULP on two of eight real super-blocks (cuda_quant_dot.cu, +// same function, extensively commented on this exact point). This project's +// CMakeLists.txt already passes -ffp-contract=off to $ +// project-wide (unlike its CXX-only reach on the CUDA side), so plain `*`/`+` +// is used here rather than carrying the CUDA workaround over unexamined -- +// the row's bit-exact golden-vector gate is what actually proves this, not +// this comment; if that gate goes red, apply the same non-fused-multiply +// discipline HIP provides before assuming anything else is wrong. +__device__ inline float DotIQ4XS(const BlockIQ4_XS* xb, const BlockQ8_K* yb) { + const float d4d8 = DF16ToF32(xb->d) * yb->d; + uint16_t h = xb->scales_h; + const uint8_t* qs = xb->qs; + const int8_t* q8 = yb->qs; + float sumf = 0.0f; + for (int ib = 0; ib < kQK_K / 32; ib += 2) { + const uint8_t ls1 = + static_cast((xb->scales_l[ib / 2] & 0xf) | ((h << 4) & 0x30)); + const uint8_t ls2 = + static_cast((xb->scales_l[ib / 2] >> 4) | ((h << 2) & 0x30)); + h = static_cast(h >> 4); + const float d1 = d4d8 * (ls1 - 32); + const float d2 = d4d8 * (ls2 - 32); + int sumi1 = 0; + int sumi2 = 0; + for (int j = 0; j < 16; ++j) { + sumi1 += q8[j + 0] * d_kvalues_iq4nl[qs[j] & 0xf]; + sumi2 += q8[j + 16] * d_kvalues_iq4nl[qs[j] >> 4]; + } + sumf += d1 * static_cast(sumi1 + sumi2); + qs += 16; + q8 += 32; + sumi1 = 0; + sumi2 = 0; + for (int j = 0; j < 16; ++j) { + sumi1 += q8[j + 0] * d_kvalues_iq4nl[qs[j] & 0xf]; + sumi2 += q8[j + 16] * d_kvalues_iq4nl[qs[j] >> 4]; + } + sumf += d2 * static_cast(sumi1 + sumi2); + qs += 16; + q8 += 32; + } + return sumf; +} + // Q6_K integer core over a SUB-BLOCK RANGE, for the cooperative decode arm // (`KQuantGemmKCoopQ6K`, issue #1910). A K-quant superblock is kQK_K = 256 // weights in kSubBlocks = 8 groups of 32; this returns the dp4a accumulator for @@ -408,8 +505,11 @@ __global__ void GroupedQ8_0K(OutT* __restrict__ out, const uint8_t* __restrict__ } } -// K-quant grouped kernel body (shared by Q4_K/Q5_K/Q6_K instantiations). -// Fmt: 0=Q4_K, 1=Q5_K, 2=Q6_K. +// K-quant grouped kernel body (shared by Q4_K/Q5_K/Q6_K/IQ4_XS/IQ3_XXS +// instantiations). Fmt: 0=Q4_K, 1=Q5_K, 2=Q6_K, 4=IQ4_XS, 5=IQ3_XXS +// (KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT, #1940) -- same Fmt numbering as +// KQuantGemmK above, so a caller's `fmt` value means the same thing on +// either entry point. template __global__ void GroupedKQ8K(OutT* __restrict__ out, const uint8_t* __restrict__ weight, const BlockQ8_K* __restrict__ act, @@ -429,10 +529,13 @@ __global__ void GroupedKQ8K(OutT* __restrict__ out, const uint8_t* __restrict__ const BlockQ8_K* a_sb = a_row + sb; if constexpr (Fmt == 2) partial += DotQ6K(static_cast(w_sb), a_sb); else if constexpr (Fmt == 1) partial += DotQ5K(static_cast(w_sb), a_sb); + else if constexpr (Fmt == 4) partial += DotIQ4XS(static_cast(w_sb), a_sb); + else if constexpr (Fmt == 5) partial += DotIQ3XXS(static_cast(w_sb), a_sb); else partial += DotQ4K(static_cast(w_sb), a_sb); } #pragma unroll for (int off = 16; off > 0; off >>= 1) partial += __shfl_down_sync(0xffffffffULL, partial, off); + if constexpr (Fmt == 5) partial *= 0.25f; // see KQuantGemmK's FinalFactor note if (lane == 0) { if constexpr (sizeof(OutT) == 4) out[p * n + j] = partial; else out[p * n + j] = DF32ToBF16(partial); @@ -471,10 +574,22 @@ __global__ void KQuantGemmK(OutT* __restrict__ out, const uint8_t* __restrict__ partial += (DF16ToF32(xb->d) * yb->d) * DotQ6KIsumRange(xb, yb, 0, kSubBlocks); } else if constexpr (Fmt == 2) partial += DotQ6K(static_cast(w_sb), a_row + sb); else if constexpr (Fmt == 1) partial += DotQ5K(static_cast(w_sb), a_row + sb); + // Fmt 4/5 (KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT, #1940): IQ4_XS/IQ3_XXS. + // Both dot bodies already read the ROCm codebook tables at file scope, so + // this is a dispatch-table extension over the same nsb=K/256 skeleton the + // four existing formats use -- no new activation quantizer, no new grid. + else if constexpr (Fmt == 4) partial += DotIQ4XS(static_cast(w_sb), a_row + sb); + else if constexpr (Fmt == 5) partial += DotIQ3XXS(static_cast(w_sb), a_row + sb); else partial += DotQ4K(static_cast(w_sb), a_row + sb); } #pragma unroll for (int off = 16; off > 0; off >>= 1) partial += __shfl_down_sync(0xffffffffULL, partial, off); + // The IQ3_XXS codebook stores lanes at a fixed 4x magnitude that upstream + // applies ONCE per row rather than per super-block (cuda_quant_dot.cu + // FinalFactor = 0.25f); IQ4_XS's per-sub-block delta is already + // folded into `d1`/`d2` inside DotIQ4XS, so its factor is 1 (a no-op, not + // written out as a multiply). + if constexpr (Fmt == 5) partial *= 0.25f; if (lane == 0) { if constexpr (sizeof(OutT) == 4) out[i * n + j] = partial; else out[i * n + j] = DF32ToBF16(partial); @@ -851,19 +966,26 @@ void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tenso Check(hipGetLastError(), "q8_0 gemm"); return; } - if (b.dtype == DType::kQ4_K || b.dtype == DType::kQ5_K || b.dtype == DType::kQ6_K) { + if (b.dtype == DType::kQ4_K || b.dtype == DType::kQ5_K || b.dtype == DType::kQ6_K || + b.dtype == DType::kIQ4_XS || b.dtype == DType::kIQ3_XXS) { if (k % kQK_K != 0) throw std::runtime_error("vt rocm: matmul_bt_quant K-quant: K%256!=0"); const int64_t nsb = k / kQK_K; - const size_t w_block_bytes = b.dtype == DType::kQ6_K ? sizeof(BlockQ6_K) - : b.dtype == DType::kQ5_K ? sizeof(BlockQ5_K) - : sizeof(BlockQ4_K); + const size_t w_block_bytes = b.dtype == DType::kQ6_K ? sizeof(BlockQ6_K) + : b.dtype == DType::kQ5_K ? sizeof(BlockQ5_K) + : b.dtype == DType::kIQ4_XS ? sizeof(BlockIQ4_XS) + : b.dtype == DType::kIQ3_XXS ? sizeof(BlockIQ3_XXS) + : sizeof(BlockQ4_K); const size_t w_row_bytes = static_cast(nsb) * w_block_bytes; BlockQ8_K* qact = static_cast(EnsureQuantScratch( static_cast(m) * nsb * sizeof(BlockQ8_K), s)); LaunchQ8KQuantizer(qact, a.data, ActDtOf(a.dtype), a.stride[0], m, nsb, s, Q8KProductionRoute::kDense, q.device.index); const int64_t grid = (m * n + kWarpsPerBlock - 1) / kWarpsPerBlock; - const int fmt = b.dtype == DType::kQ6_K ? 2 : b.dtype == DType::kQ5_K ? 1 : 0; + const int fmt = b.dtype == DType::kQ6_K ? 2 + : b.dtype == DType::kQ5_K ? 1 + : b.dtype == DType::kIQ4_XS ? 4 + : b.dtype == DType::kIQ3_XXS ? 5 + : 0; const int coop = KQuantDecodeCoopWarps(b.dtype, m, nsb); auto launch = [&](auto ot) { using OutT = decltype(ot); @@ -880,13 +1002,15 @@ void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tenso if (fmt == 2 && Q6KSmallPrivateEnabled()) KQuantGemmK<<(grid), block, 0, s>>>(o, w, qact, m, n, nsb, w_row_bytes, w_block_bytes); else if (fmt == 2) KQuantGemmK<<(grid), block, 0, s>>>(o, w, qact, m, n, nsb, w_row_bytes, w_block_bytes); else if (fmt == 1) KQuantGemmK<<(grid), block, 0, s>>>(o, w, qact, m, n, nsb, w_row_bytes, w_block_bytes); + else if (fmt == 4) KQuantGemmK<<(grid), block, 0, s>>>(o, w, qact, m, n, nsb, w_row_bytes, w_block_bytes); + else if (fmt == 5) KQuantGemmK<<(grid), block, 0, s>>>(o, w, qact, m, n, nsb, w_row_bytes, w_block_bytes); else KQuantGemmK<<(grid), block, 0, s>>>(o, w, qact, m, n, nsb, w_row_bytes, w_block_bytes); }; if (out.dtype == DType::kF32) launch(float{}); else launch(uint16_t{}); Check(hipGetLastError(), "K-quant gemm"); return; } - throw std::runtime_error("vt rocm: matmul_bt_quant: unsupported weight dtype (ported: Q8_0/Q4_K/Q5_K/Q6_K; owed: Q4_0/Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ3_S/IQ2_S/MXFP4 -- the loader pre-filters to the ported set, so reaching here is a bug)"); + throw std::runtime_error("vt rocm: matmul_bt_quant: unsupported weight dtype (ported: Q8_0/Q4_K/Q5_K/Q6_K/IQ4_XS/IQ3_XXS; owed: Q4_0/Q2_K/Q3_K/IQ2_XXS/IQ3_S/IQ2_S/MXFP4 -- the loader pre-filters to the ported set, so reaching here is a bug)"); } @@ -928,25 +1052,34 @@ void MatmulBTQuantGroupedKernelRocm(Queue& q, Tensor& out, const Tensor& act, return; } - if (weight.dtype == DType::kQ4_K || weight.dtype == DType::kQ5_K || weight.dtype == DType::kQ6_K) { + if (weight.dtype == DType::kQ4_K || weight.dtype == DType::kQ5_K || weight.dtype == DType::kQ6_K || + weight.dtype == DType::kIQ4_XS || weight.dtype == DType::kIQ3_XXS) { if (k % kQK_K != 0) throw std::runtime_error("vt rocm: matmul_bt_quant_grouped K-quant: K must be a multiple of 256"); const int64_t nsb = k / kQK_K; - const size_t w_block_bytes = weight.dtype == DType::kQ4_K ? sizeof(BlockQ4_K) - : weight.dtype == DType::kQ5_K ? sizeof(BlockQ5_K) - : sizeof(BlockQ6_K); + const size_t w_block_bytes = weight.dtype == DType::kQ4_K ? sizeof(BlockQ4_K) + : weight.dtype == DType::kQ5_K ? sizeof(BlockQ5_K) + : weight.dtype == DType::kQ6_K ? sizeof(BlockQ6_K) + : weight.dtype == DType::kIQ4_XS ? sizeof(BlockIQ4_XS) + : sizeof(BlockIQ3_XXS); const size_t w_row_bytes = static_cast(nsb) * w_block_bytes; BlockQ8_K* qact = static_cast(EnsureQuantScratch( static_cast(Pa) * nsb * sizeof(BlockQ8_K), s)); LaunchQ8KQuantizer(qact, act.data, ActDtOf(act.dtype), act.stride[0], Pa, nsb, s, Q8KProductionRoute::kGrouped, q.device.index); const int64_t grid = (P * n + kWarpsPerBlock - 1) / kWarpsPerBlock; - const int fmt = weight.dtype == DType::kQ6_K ? 2 : weight.dtype == DType::kQ5_K ? 1 : 0; + const int fmt = weight.dtype == DType::kQ6_K ? 2 + : weight.dtype == DType::kQ5_K ? 1 + : weight.dtype == DType::kIQ4_XS ? 4 + : weight.dtype == DType::kIQ3_XXS ? 5 + : 0; auto launch = [&](auto ot) { using OutT = decltype(ot); auto* o = static_cast(out.data); if (fmt == 2) GroupedKQ8K<<(grid), block, 0, s>>>(o, w, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast); else if (fmt == 1) GroupedKQ8K<<(grid), block, 0, s>>>(o, w, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast); + else if (fmt == 4) GroupedKQ8K<<(grid), block, 0, s>>>(o, w, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast); + else if (fmt == 5) GroupedKQ8K<<(grid), block, 0, s>>>(o, w, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast); else GroupedKQ8K<<(grid), block, 0, s>>>(o, w, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast); }; if (out.dtype == DType::kF32) launch(float{}); else launch(uint16_t{}); @@ -955,8 +1088,8 @@ void MatmulBTQuantGroupedKernelRocm(Queue& q, Tensor& out, const Tensor& act, } throw std::runtime_error( - "vt rocm: matmul_bt_quant_grouped: unsupported weight dtype (ported: Q8_0/Q4_K/Q5_K/Q6_K; " - "owed: Q4_0/Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ2_S/MXFP4 -- the loader pre-filters, so reaching here is a bug)"); + "vt rocm: matmul_bt_quant_grouped: unsupported weight dtype (ported: Q8_0/Q4_K/Q5_K/Q6_K/IQ4_XS/IQ3_XXS; " + "owed: Q4_0/Q2_K/Q3_K/IQ2_XXS/IQ2_S/MXFP4 -- the loader pre-filters, so reaching here is a bug)"); } } // namespace vt::rocm diff --git a/src/vt/rocm/rocm_quant_iq_tables.h b/src/vt/rocm/rocm_quant_iq_tables.h new file mode 100644 index 0000000000..8d21494f7a --- /dev/null +++ b/src/vt/rocm/rocm_quant_iq_tables.h @@ -0,0 +1,100 @@ +// AUTO-DERIVED from src/vt/cpu/cpu_quant_iq_tables.h (SAME literals, single +// source of truth) via scripts/gen-cuda-iq-tables.py -- that script reads only +// the CPU header and is source-language-agnostic; its name predates this +// second address space and nothing in its output is CUDA-specific. Redirect +// its stdout to regenerate a table below if the CPU tables ever change. +// +// These are the tables KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT's DotIQ4XS/DotIQ3XXS +// need (rocm_grouped_gemm.hip), mirroring the CUDA device tables in +// src/vt/cuda/cuda_quant_iq_tables.cuh (d_kmask_iq2xs, d_ksigns_iq2xs, +// d_iq3xxs_grid, d_kvalues_iq4nl) byte for byte -- same storage class per +// table too (d_ksigns_iq2xs is GLOBAL, not __constant__, because the +// warp-per-output GEMM has 32 lanes reading 32 DIFFERENT super-blocks' +// grid/sign indices, a divergent access __constant__ serialises; see the CUDA +// file's own comment on this point). llama.cpp @ 237ad9b96 +// ggml-common.h:499/:503/:550/:1007, pin b10451 per .agents/upstream-sync.md. +// +// Sealed the same way the CUDA tables are (test_cuda_quant_dot.cpp's +// SnapshotIqTablesFromDevice pattern): a wrong same-shaped grid still decodes +// plausibly, so the ROCm gate must memcmp these against the CPU host tables +// rather than trust the generator's output unchecked. +#ifndef VT_ROCM_QUANT_IQ_TABLES_H_ +#define VT_ROCM_QUANT_IQ_TABLES_H_ + +#include + +namespace vt::rocm { + +// GENERATED by scripts/gen-cuda-iq-tables.py from +// src/vt/cpu/cpu_quant_iq_tables.h::kKmaskIq2xs -- do not hand edit. +__device__ __constant__ uint8_t d_kmask_iq2xs[8] = { + 1, 2, 4, 8, 16, 32, 64, 128, +}; +// GENERATED by scripts/gen-cuda-iq-tables.py from +// src/vt/cpu/cpu_quant_iq_tables.h::kKsignsIq2xs -- do not hand edit. +__device__ uint8_t d_ksigns_iq2xs[128] = { // GLOBAL (not __constant__) + 0, 129, 130, 3, 132, 5, 6, 135, 136, 9, 10, 139, 12, 141, 142, 15, + 144, 17, 18, 147, 20, 149, 150, 23, 24, 153, 154, 27, 156, 29, 30, 159, + 160, 33, 34, 163, 36, 165, 166, 39, 40, 169, 170, 43, 172, 45, 46, 175, + 48, 177, 178, 51, 180, 53, 54, 183, 184, 57, 58, 187, 60, 189, 190, 63, + 192, 65, 66, 195, 68, 197, 198, 71, 72, 201, 202, 75, 204, 77, 78, 207, + 80, 209, 210, 83, 212, 85, 86, 215, 216, 89, 90, 219, 92, 221, 222, 95, + 96, 225, 226, 99, 228, 101, 102, 231, 232, 105, 106, 235, 108, 237, 238, 111, + 240, 113, 114, 243, 116, 245, 246, 119, 120, 249, 250, 123, 252, 125, 126, 255, +}; +// GENERATED by scripts/gen-cuda-iq-tables.py from +// src/vt/cpu/cpu_quant_iq_tables.h::kIq3xxsGrid -- do not hand edit. +__device__ __constant__ uint32_t d_iq3xxs_grid[256] = { + 0x04040404u, 0x04040414u, 0x04040424u, 0x04040c0cu, 0x04040c1cu, 0x04040c3eu, + 0x04041404u, 0x04041414u, 0x04041c0cu, 0x04042414u, 0x04043e1cu, 0x04043e2cu, + 0x040c040cu, 0x040c041cu, 0x040c0c04u, 0x040c0c14u, 0x040c140cu, 0x040c142cu, + 0x040c1c04u, 0x040c1c14u, 0x040c240cu, 0x040c2c24u, 0x040c3e04u, 0x04140404u, + 0x04140414u, 0x04140424u, 0x04140c0cu, 0x04141404u, 0x04141414u, 0x04141c0cu, + 0x04141c1cu, 0x04141c3eu, 0x04142c0cu, 0x04142c3eu, 0x04143e2cu, 0x041c040cu, + 0x041c043eu, 0x041c0c04u, 0x041c0c14u, 0x041c142cu, 0x041c3e04u, 0x04240c1cu, + 0x04241c3eu, 0x04242424u, 0x04242c3eu, 0x04243e1cu, 0x04243e2cu, 0x042c040cu, + 0x042c043eu, 0x042c1c14u, 0x042c2c14u, 0x04341c2cu, 0x04343424u, 0x043e0c04u, + 0x043e0c24u, 0x043e0c34u, 0x043e241cu, 0x043e340cu, 0x0c04040cu, 0x0c04041cu, + 0x0c040c04u, 0x0c040c14u, 0x0c04140cu, 0x0c04141cu, 0x0c041c04u, 0x0c041c14u, + 0x0c041c24u, 0x0c04243eu, 0x0c042c04u, 0x0c0c0404u, 0x0c0c0414u, 0x0c0c0c0cu, + 0x0c0c1404u, 0x0c0c1414u, 0x0c14040cu, 0x0c14041cu, 0x0c140c04u, 0x0c140c14u, + 0x0c14140cu, 0x0c141c04u, 0x0c143e14u, 0x0c1c0404u, 0x0c1c0414u, 0x0c1c1404u, + 0x0c1c1c0cu, 0x0c1c2434u, 0x0c1c3434u, 0x0c24040cu, 0x0c24042cu, 0x0c242c04u, + 0x0c2c1404u, 0x0c2c1424u, 0x0c2c2434u, 0x0c2c3e0cu, 0x0c34042cu, 0x0c3e1414u, + 0x0c3e2404u, 0x14040404u, 0x14040414u, 0x14040c0cu, 0x14040c1cu, 0x14041404u, + 0x14041414u, 0x14041434u, 0x14041c0cu, 0x14042414u, 0x140c040cu, 0x140c041cu, + 0x140c042cu, 0x140c0c04u, 0x140c0c14u, 0x140c140cu, 0x140c1c04u, 0x140c341cu, + 0x140c343eu, 0x140c3e04u, 0x14140404u, 0x14140414u, 0x14140c0cu, 0x14140c3eu, + 0x14141404u, 0x14141414u, 0x14141c3eu, 0x14142404u, 0x14142c2cu, 0x141c040cu, + 0x141c0c04u, 0x141c0c24u, 0x141c3e04u, 0x141c3e24u, 0x14241c2cu, 0x14242c1cu, + 0x142c041cu, 0x142c143eu, 0x142c240cu, 0x142c3e24u, 0x143e040cu, 0x143e041cu, + 0x143e0c34u, 0x143e242cu, 0x1c04040cu, 0x1c040c04u, 0x1c040c14u, 0x1c04140cu, + 0x1c04141cu, 0x1c042c04u, 0x1c04342cu, 0x1c043e14u, 0x1c0c0404u, 0x1c0c0414u, + 0x1c0c1404u, 0x1c0c1c0cu, 0x1c0c2424u, 0x1c0c2434u, 0x1c14040cu, 0x1c14041cu, + 0x1c140c04u, 0x1c14142cu, 0x1c142c14u, 0x1c143e14u, 0x1c1c0c0cu, 0x1c1c1c1cu, + 0x1c241c04u, 0x1c24243eu, 0x1c243e14u, 0x1c2c0404u, 0x1c2c0434u, 0x1c2c1414u, + 0x1c2c2c2cu, 0x1c340c24u, 0x1c341c34u, 0x1c34341cu, 0x1c3e1c1cu, 0x1c3e3404u, + 0x24040424u, 0x24040c3eu, 0x24041c2cu, 0x24041c3eu, 0x24042c1cu, 0x24042c3eu, + 0x240c3e24u, 0x24141404u, 0x24141c3eu, 0x24142404u, 0x24143404u, 0x24143434u, + 0x241c043eu, 0x241c242cu, 0x24240424u, 0x24242c0cu, 0x24243424u, 0x242c142cu, + 0x242c241cu, 0x242c3e04u, 0x243e042cu, 0x243e0c04u, 0x243e0c14u, 0x243e1c04u, + 0x2c040c14u, 0x2c04240cu, 0x2c043e04u, 0x2c0c0404u, 0x2c0c0434u, 0x2c0c1434u, + 0x2c0c2c2cu, 0x2c140c24u, 0x2c141c14u, 0x2c143e14u, 0x2c1c0414u, 0x2c1c2c1cu, + 0x2c240c04u, 0x2c24141cu, 0x2c24143eu, 0x2c243e14u, 0x2c2c0414u, 0x2c2c1c0cu, + 0x2c342c04u, 0x2c3e1424u, 0x2c3e2414u, 0x34041424u, 0x34042424u, 0x34042434u, + 0x34043424u, 0x340c140cu, 0x340c340cu, 0x34140c3eu, 0x34143424u, 0x341c1c04u, + 0x341c1c34u, 0x34242424u, 0x342c042cu, 0x342c2c14u, 0x34341c1cu, 0x343e041cu, + 0x343e140cu, 0x3e04041cu, 0x3e04042cu, 0x3e04043eu, 0x3e040c04u, 0x3e041c14u, + 0x3e042c14u, 0x3e0c1434u, 0x3e0c2404u, 0x3e140c14u, 0x3e14242cu, 0x3e142c14u, + 0x3e1c0404u, 0x3e1c0c2cu, 0x3e1c1c1cu, 0x3e1c3404u, 0x3e24140cu, 0x3e24240cu, + 0x3e2c0404u, 0x3e2c0414u, 0x3e2c1424u, 0x3e341c04u, +}; +// GENERATED by scripts/gen-cuda-iq-tables.py from +// src/vt/cpu/cpu_quant_iq_tables.h::kValuesIq4nl -- do not hand edit. +__device__ __constant__ int8_t d_kvalues_iq4nl[16] = { + -127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113, +}; + +} // namespace vt::rocm + +#endif // VT_ROCM_QUANT_IQ_TABLES_H_ diff --git a/tests/vllm/model_executor/test_gguf_device_fit.cpp b/tests/vllm/model_executor/test_gguf_device_fit.cpp index 2d3d89a9c9..7fa06dc911 100644 --- a/tests/vllm/model_executor/test_gguf_device_fit.cpp +++ b/tests/vllm/model_executor/test_gguf_device_fit.cpp @@ -705,20 +705,31 @@ TEST_CASE( // #2516: NAMED DEVICES, not `CurrentPlatform()`. This case is about the CPU // and CUDA `vec_dot` kernels #2247 added, and asking the running platform made // it RED on every ROCm build for a reason the case did not state -- - // `DeviceKeepQuantSupported` serves exactly {Q8_0, Q4_K, Q5_K, Q6_K} there - // (#1940 owns the gap). Both answers are pinned instead, because the ROCm one - // is a real property of this tree and a case that merely skipped it would stop - // noticing when #1940 lands. + // `DeviceKeepQuantSupported` served exactly {Q8_0, Q4_K, Q5_K, Q6_K} there + // (#1940 owned the gap). Both answers are pinned instead, because the ROCm + // ones are a real property of this tree and a case that merely skipped them + // would stop noticing when #1940 landed. + // + // KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940) is that landing, for IQ4_XS + // (ggml_type 23, `down_exps`): this case's own comment said a case that + // merely skipped the ROCm answer "would stop noticing when #1940 lands", so + // the per-tensor split below is that noticing, not a loosened assertion. + // IQ2_XS (ggml_type 17, `gate_exps`) is unaffected -- #1940's scope is + // IQ4_XS and IQ3_XXS only -- and keeps the same `kExpandBf16` pin on ROCm. for (const vllm::GgufTensorInfo& t : gguf.Tensors()) { CAPTURE(t.name); CHECK(vllm::RouteGgufTensor(true, false, false, false, vllm::GgufTensorRole::kStackedExpertWeight, t.ggml_type, t.shape, vt::DeviceType::kCPU) == vllm::GgufResidency::kKeepQuant); + const vllm::GgufResidency rocm_expected = + t.ggml_type == 23U // IQ4_XS + ? vllm::GgufResidency::kKeepQuant + : vllm::GgufResidency::kExpandBf16; CHECK(vllm::RouteGgufTensor(true, false, false, false, vllm::GgufTensorRole::kStackedExpertWeight, t.ggml_type, t.shape, vt::DeviceType::kROCM) == - vllm::GgufResidency::kExpandBf16); + rocm_expected); } } @@ -932,9 +943,15 @@ TEST_CASE( vllm::GgufLoadPolicy rocm = PolicyWith(true, false, false, false); rocm.device = vt::DeviceType::kROCM; - // NO PLAN: ROCm has no IQ `vec_dot` (#1940), so the towers expand and the - // predicate is false. This is the state #2516 reports and the inertness pin - // for every ROCm load that configures no placement. + // NO PLAN: this fixture's `gate_exps` tower is IQ2_XS, which ROCm still has + // no `vec_dot` for (#1940 tracks IQ2_XS as owed; only IQ4_XS and IQ3_XXS + // landed under KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT). The lane is all-or- + // nothing over BOTH `_exps.weight` towers, so one unsupported encoding still + // fails the whole predicate even though the OTHER tower (`down_exps`, + // IQ4_XS) now keeps on ROCm by itself -- see the per-tensor split in + // "gguf_device_fit: IQ2_XS and IQ4_XS expert towers REACH the slot lane" + // above. This is the state #2516 reports and the inertness pin for every + // ROCm load that configures no placement. vllm::ResetActiveMoePlacementPlanForTesting(); CHECK_FALSE( vllm::GgufExpertTowersReachSlotLane(gguf, "_exps.weight", rocm)); diff --git a/tests/vllm/test_gguf_keep_quant.cpp b/tests/vllm/test_gguf_keep_quant.cpp index c7f1e74487..d2e76c8dd4 100644 --- a/tests/vllm/test_gguf_keep_quant.cpp +++ b/tests/vllm/test_gguf_keep_quant.cpp @@ -553,8 +553,13 @@ TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { // IQ2_S (256-elem, Q8_K-act) and MXFP4 (32-elem, Q8_0-act) are keep-quant // capable as of the UD-IQ2_M vehicle, so they route like the others. // The DEVICE axis (review #523): the running device's kernel set can be - // narrower than the loader's CPU-derived list — ROCm implements exactly - // {Q8_0, Q4_K, Q5_K, Q6_K}; the rest keep expand_bf16 there. + // narrower than the loader's CPU-derived list — ROCm implemented + // exactly {Q8_0, Q4_K, Q5_K, Q6_K} through #2247; IQ4_XS joined with + // KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940). IQ3_XXS joined the same + // row's ROCm kernel set but is not in `all_types` below (nor is + // Q2_K, a pre-existing gap this row does not close), so it is not + // exercised by this particular exhaustive table; it is gated instead + // by tests/vt/test_backend_cross_device.cpp's cross-device cases. // QUANT-GGUF-IQ-VECDOT (#2247) put IQ2_XS and IQ4_XS in this list. // They were gather-only between #2245 and #2247 — decoder, no vec_dot — // and the `vec_dot` rows are what moved them onto the GEMM arm. @@ -566,7 +571,7 @@ TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { const bool rocm = kRouteDev == vt::DeviceType::kROCM; const bool device_capable = !rocm || type == kQ8_0 || type == kQ4_K || type == kQ5_K || - type == kQ6_K; + type == kQ6_K || type == kIQ4_XS; const bool block_capable = cpu_capable && device_capable; const int64_t blk = (type == kQ4_0 || type == kQ5_0 || type == kQ8_0 || type == kMXFP4 || type == kIQ4_NL) @@ -655,13 +660,13 @@ TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { // Both outcomes are actually exercised (a table that never keeps anything // would pass every assertion above vacuously). The kept count is // device-dependent (review #523): 12 block-capable encodings x 2 keep-capable - // GEMM roles where the device covers the CPU list; 4 x 2 on ROCm (ROCm's - // kernel set is {Q8_0, Q4_K, Q5_K, Q6_K}, and neither Q5_0 nor IQ4_NL nor - // either IQ*_XS is in it). The GATHER role adds 13 more (the 12, plus Q8_K, - // which has a decoder and no vec_dot) on any device that REGISTERS the block - // gather, and nothing on a device that does not. Written as named terms - // rather than one number so a future change to any one of them says which one - // moved. Both moves are now on record and they are mirror images: + // GEMM roles where the device covers the CPU list; 4 x 2 on ROCm through + // #2247 (ROCm's kernel set was {Q8_0, Q4_K, Q5_K, Q6_K}). The GATHER role + // adds 13 more (the 12, plus Q8_K, which has a decoder and no vec_dot) on + // any device that REGISTERS the block gather, and nothing on a device that + // does not. Written as named terms rather than one number so a future + // change to any one of them says which one moved. Both moves are now on + // record and they are mirror images: // LOADER-GGUF-IQ (#2240) moved the GATHER term 11 -> 13 and left GEMM at 20, // the shape of a decode-only port; QUANT-GGUF-IQ-VECDOT (#2247) moves the // GEMM term 20 -> 24 and leaves GATHER at 13, the shape of a dot-only port. @@ -683,7 +688,15 @@ TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { // was, the same decode-only shape #2240 had. That asymmetry IS the row's // per-tier result: IQ3_S stays compressed in a gather table and expands to // bf16 in a GEMM, on every device. - const int gemm_kept = kRouteDev == vt::DeviceType::kROCM ? 8 : 24; + // + // KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940) moves the ROCm GEMM term 8 -> 10 + // (IQ4_XS joins {Q8_0, Q4_K, Q5_K, Q6_K}) and leaves every other device's + // GEMM term and every device's GATHER term untouched: IQ4_XS was already + // CPU-capable and already counted at `24` on CPU/CUDA, so only the ROCm + // narrowing moved. IQ3_XXS also joined ROCm's kernel set in the same row but + // is not in `all_types` above, so it does not move this number; see the + // comment beside `device_capable`. + const int gemm_kept = kRouteDev == vt::DeviceType::kROCM ? 10 : 24; const int gather_kept = vt::OpRegistered(vt::OpId::kEmbeddingQuant, kRouteDev) ? 14 : 0; CHECK(kept == gemm_kept + gather_kept); diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index b80e839dd5..75552ff581 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -46,6 +46,12 @@ #include "vt/recipes.h" #include "vt/rocm/rocm_runtime.h" +// KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940): the real-checkpoint IQ4_XS +// blocks and llama.cpp's own vec_dot output for them, shared with the CUDA +// gate in test_cuda_quant_dot.cpp. +#include "iq2xs_iq4xs_dot_golden.h" +#include "iq2xs_iq4xs_golden_vectors.h" + namespace { using vt::Device; @@ -2446,13 +2452,20 @@ TEST_CASE("decode-skinny MatmulBT (wvSplitK path) matches the CPU oracle") { } } -TEST_CASE("non-grouped keep-quant GEMM (Q8_0/Q4_K/Q5_K/Q6_K) matches the CPU oracle") { +TEST_CASE("non-grouped keep-quant GEMM (Q8_0/Q4_K/Q5_K/Q6_K/IQ4_XS/IQ3_XXS) matches the CPU oracle") { // kMatmulBTQuant (op 74) on ROCm vs the CPU keep-quant reference. The // non-grouped arm carries PR #523's headline mechanism and had NO coverage // (review sweep 2026-08-13); the ROCm dispatcher's src-vs-out dtype mix-up // in the fused preamble (the 0.8B divergence, row/ROCM-GDN-08B-FIX) is // exactly the class an untested-but-registered op hides. REQUIRE (not skip) // on ROCm so a dropped RegisterOp can never pass silently. + // + // IQ4_XS/IQ3_XXS added by KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940). Their + // `qs`/grid-index bytes are read as unconstrained lookup indices (4-bit + // nibbles into a 16-entry codebook for IQ4_XS, a full byte into a 256-entry + // grid for IQ3_XXS), so the same random-byte block this table already + // builds for the four linear-scale formats is a valid block for these two + // as well -- no in-range constraint to add. constexpr int64_t M = 3, N = 8, K = 512; struct Fmt { vt::DType dt; int64_t block_bytes; int d_off; int dmin_off; const char* name; }; const Fmt fmts[] = { @@ -2460,6 +2473,8 @@ TEST_CASE("non-grouped keep-quant GEMM (Q8_0/Q4_K/Q5_K/Q6_K) matches the CPU ora {vt::DType::kQ4_K, 144, 0, 2, "q4_K"}, {vt::DType::kQ6_K, 210, 208, -1, "q6_K"}, {vt::DType::kQ5_K, 176, 0, 2, "q5_K"}, + {vt::DType::kIQ4_XS, 136, 0, -1, "iq4_xs"}, + {vt::DType::kIQ3_XXS, 98, 0, -1, "iq3_xxs"}, }; const bool rocm_present = OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM); const bool any_rocm = [&] { @@ -3298,11 +3313,17 @@ TEST_CASE("ROCm Q6_K decode spreads one row's superblocks over several warps") { } #endif // VLLM_CPP_HIP -TEST_CASE("grouped quant expert GEMM (Q8_0/Q4_K/Q6_K) matches the CPU oracle") { +TEST_CASE("grouped quant expert GEMM (Q8_0/Q4_K/Q6_K/IQ4_XS/IQ3_XXS) matches the CPU oracle") { // kMatmulBTQuantGrouped on ROCm vs the CPU keep-quant reference // (cpu_quant_gemm.cpp:305). Valid random blocks (valid f16 deltas, random // quants) at a real expert-MLP shape. Integer cores are bit-exact ports; // the f16/f32 scale sum reassociates across lanes, so NMSE <= 5e-4. + // + // IQ4_XS/IQ3_XXS added by KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940) -- the + // MOTIVATING arm: these are ROUTED-EXPERT weights on a real checkpoint + // (unsloth/GLM-5.3-Flash-GGUF-style IQ4_XS gate/up + IQ3_XXS down), so the + // grouped path is the one a real load actually exercises, not the + // non-grouped table above. constexpr int64_t P = 3, N = 8, K = 512; // K%256==0 (K-quant superblocks) constexpr int64_t E = 4; // experts const std::vector eids = {2, 0, 3}; // routed experts (non-sorted) @@ -3314,6 +3335,8 @@ TEST_CASE("grouped quant expert GEMM (Q8_0/Q4_K/Q6_K) matches the CPU oracle") { {vt::DType::kQ4_K, 144, 0, 2, "q4_K"}, // {d,dmin,sc,qs} superblocks of 256 {vt::DType::kQ6_K, 210, 208, -1, "q6_K"},// {ql,qh,scales,d} superblocks of 256 {vt::DType::kQ5_K, 176, 0, 2, "q5_K"}, // {d,dmin,sc,qh,qs} superblocks of 256 + {vt::DType::kIQ4_XS, 136, 0, -1, "iq4_xs"}, // {d,scales_h,scales_l,qs} superblocks of 256 + {vt::DType::kIQ3_XXS, 98, 0, -1, "iq3_xxs"}, // {d,qs[3*QK_K/8]} superblocks of 256 }; // REQUIRE-proven registration on ROCm (never a silent skip — review sweep @@ -3388,6 +3411,134 @@ TEST_CASE("grouped quant expert GEMM (Q8_0/Q4_K/Q6_K) matches the CPU oracle") { } } +// ─── KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940): the risk this row's spec ──── +// named -- CUDA needed __fmul_rn/__fadd_rn in DotIQ4XS to stay bit-exact +// against the oracle, because nvcc's default -fmad=true silently contracted +// the two-rounding accumulation into a single-rounding FMA and two of eight +// real super-blocks then disagreed by 1-4 ULP (cuda_quant_dot.cu). This +// project's CMakeLists.txt already passes -ffp-contract=off to +// $ project-wide, unlike its CXX-only reach on CUDA, so +// DotIQ4XS on ROCm uses plain `*`/`+` rather than carrying that workaround +// over unexamined (see the function's own comment). This is the gate that +// proves whether that bet paid off, over the SAME real checkpoint bytes and +// the SAME expected oracle bits the CUDA gate uses -- reused, not +// re-derived, so a drift between the two device gates cannot hide. +namespace { +// test_cuda_quant_dot.cpp's MakeDotActivation, restated rather than shared +// (no cross-file dependency): every value is an integer in [-1024, 1023] over +// 64, exact in binary32 on any compiler, so the same bytes come out of any +// build and the golden's provenance carries over unchanged. +void MakeIq4xsDotActivation(int n, uint32_t seed, float* x) { + uint32_t st = seed; + for (int i = 0; i < n; ++i) { + st = st * 1664525U + 1013904223U; + const int32_t v = static_cast((st >> 16) & 0x7ffU) - 1024; + x[i] = static_cast(v) / 64.0F; + } +} +uint32_t FloatBitsOf(float f) { + uint32_t u = 0; + std::memcpy(&u, &f, sizeof(u)); + return u; +} +float BitsFloat(uint32_t u) { + float f = 0.0F; + std::memcpy(&f, &u, sizeof(f)); + return f; +} +} // namespace + +TEST_CASE("ROCm IQ4_XS dots the ORACLE's own numbers on REAL checkpoint bytes") { + const bool any_rocm = [&] { + for (DeviceType dt : RegisteredDevices()) if (dt == DeviceType::kROCM) return true; + return false; + }(); + if (!any_rocm) { + MESSAGE("no ROCm backend on this host; ROCm IQ4_XS oracle-dot gate skipped"); + return; + } + REQUIRE_MESSAGE(OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM), + "kMatmulBTQuant must be registered on ROCm"); + + constexpr int kBlocks = 4; + constexpr int kK = 256 * kBlocks; + REQUIRE(vt::cpu::QuantTraits(DType::kIQ4_XS).vec_dot_type == DType::kQ8_K); + const size_t wbytes = std::size(vllm_test::kIq4xsGoldenBlocks); + const size_t block_bytes = wbytes / kBlocks; + REQUIRE(block_bytes * kBlocks == wbytes); + + std::vector act(kK); + MakeIq4xsDotActivation(kK, 0x4247U, act.data()); + + vt::Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + const Device gd{DeviceType::kROCM, 0}; + DevBuf da(gpu, gq, act.size()); + DevBufBytes dwt(gpu, gq, wbytes); + DevBuf dout(gpu, gq, kBlocks); + da.Upload(act); + dwt.Upload(vllm_test::kIq4xsGoldenBlocks); + + // k=256, one super-block at a time: ONE contributing lane, so BIT FOR BIT + // against the oracle's own per-block number -- no reassociation exists to + // explain a difference away. + for (int b = 0; b < kBlocks; ++b) { + CAPTURE(b); + Tensor at = Tensor::Contiguous(static_cast(da.ptr()) + b * 256 * sizeof(float), + DType::kF32, gd, {1, 256}); + Tensor wt = Tensor::Contiguous(static_cast(dwt.ptr()) + b * block_bytes, + DType::kIQ4_XS, gd, {1, 256}); + Tensor ot = Tensor::Contiguous(static_cast(dout.ptr()) + b * sizeof(float), + DType::kF32, gd, {1, 1}); + vt::MatmulBTQuant(gq, ot, at, wt); + float got = 0.0F; + gpu.Synchronize(gq); + gpu.Copy(gq, &got, static_cast(dout.ptr()) + b * sizeof(float), sizeof(float)); + gpu.Synchronize(gq); + REQUIRE(std::isfinite(got)); + CHECK(FloatBitsOf(got) == vllm_test::kIq4xsDotPerBlockBits[b]); + } + + // k=1024, all four super-blocks in one warp: the SAME __shfl_down_sync + // offsets (16,8,4,2,1) as the CUDA gate, over four live lanes, giving the + // same (v0+v2)+(v1+v3) grouping the CUDA gate's comment derives -- so the + // PRIMARY assertion is bit equality against the oracle's four numbers + // recombined in that order, and a SECONDARY one bounds the difference from + // the oracle's own sequential total by the reassociation error. + { + Tensor at = T2(da.ptr(), gd, 1, kK); + Tensor wt = Tensor::Contiguous(dwt.ptr(), DType::kIQ4_XS, gd, {1, kK}); + Tensor ot = T2(dout.ptr(), gd, 1, 1); + vt::MatmulBTQuant(gq, ot, at, wt); + float got = 0.0F; + gpu.Synchronize(gq); + gpu.Copy(gq, &got, dout.ptr(), sizeof(float)); + gpu.Synchronize(gq); + REQUIRE(std::isfinite(got)); + + const float p0 = BitsFloat(vllm_test::kIq4xsDotPerBlockBits[0]); + const float p1 = BitsFloat(vllm_test::kIq4xsDotPerBlockBits[1]); + const float p2 = BitsFloat(vllm_test::kIq4xsDotPerBlockBits[2]); + const float p3 = BitsFloat(vllm_test::kIq4xsDotPerBlockBits[3]); + const float tree = (p0 + p2) + (p1 + p3); + CAPTURE(got); + CAPTURE(tree); + CHECK(FloatBitsOf(got) == FloatBitsOf(tree)); + + const float seq = BitsFloat(vllm_test::kIq4xsDotExpectedBits); + const double mag = static_cast(std::fabs(p0)) + std::fabs(p1) + + std::fabs(p2) + std::fabs(p3); + const double bound = 4.0 * 1.1920929e-7 * mag; + const double margin = std::fabs(static_cast(got) - seq); + CAPTURE(seq); + CAPTURE(margin); + CAPTURE(bound); + CHECK(margin <= bound); + } + + gpu.DestroyQueue(gq); +} + TEST_CASE("ReshapeAndCache->PagedAttention composition matches CPU (real dims, shuffled blocks)") { // The "paged attention" case above hand-builds a contiguous KV cache; the // real model path writes it with ReshapeAndCache and reads it back. This From 7aa0aa00a8eb79d53e65685e78e8da6d3f10482a Mon Sep 17 00:00:00 2001 From: Justin Card Date: Sun, 6 Sep 2026 13:31:17 -0400 Subject: [PATCH 02/11] docs(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): name IQ4_XS/IQ3_XXS in the ROCm keep-quant line The keep-quant enumeration went stale when this row's implementation added IQ4_XS/IQ3_XXS to ROCm's DeviceKeepQuantSupported but touched no docs file, a gap a fresh review of PR #3029 found. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [Claude Code] --- docs/USAGE.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/USAGE.md b/docs/USAGE.md index b9b96c160d..550a34291b 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -769,8 +769,8 @@ skips with that refusal quoted. to inspect or override the architecture-scoped quantizer during troubleshooting. - On ROCm, GGUF mixture-of-experts checkpoints compute on the quantized - expert blocks (Q8_0, Q4_K, Q5_K, Q6_K) instead of being dequantized to - bf16 at load time. + expert blocks (Q8_0, Q4_K, Q5_K, Q6_K, IQ4_XS, IQ3_XXS) instead of being + dequantized to bf16 at load time. - On ROCm, mixture-of-experts models run the shared-expert gate and both expert-combine steps on device. Before these ops were registered the engine refused with `no kernel for op` on that path. From 5c89cf8b533b91cd98ac1076491c13ce932e9726 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 8 Sep 2026 07:32:22 +0000 Subject: [PATCH 03/11] spec(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): define contribution integration Record the #3029 conflict resolution and missing IQ3_XXS loader regression before changing code. Issue #1940 retains its broader format debt. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GPT-6 [Codex] --- .agents/specs/rocm-iquant-integration-3029.md | 50 +++++++++++++++++++ 1 file changed, 50 insertions(+) create mode 100644 .agents/specs/rocm-iquant-integration-3029.md diff --git a/.agents/specs/rocm-iquant-integration-3029.md b/.agents/specs/rocm-iquant-integration-3029.md new file mode 100644 index 0000000000..2e1853fe8f --- /dev/null +++ b/.agents/specs/rocm-iquant-integration-3029.md @@ -0,0 +1,50 @@ +# Integrate the ROCm I-quant contribution and test loader admission + +Row: `KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT`. +Issue: [#1940](https://github.com/mudler/vllm.cpp/issues/1940), which remains +open for the other quantized formats. +Contribution: [#3029](https://github.com/mudler/vllm.cpp/pull/3029). +Parent spec: [ROCm I-quant port](kernel-quant-ciq-gemm-rocm-iquant.md). +Contribution base: `7aa0aa00a8eb79d53e65685e78e8da6d3f10482a`. +Integration target: `415d17859500caf2a4cac00511820e4f4760e86f`. + +## Scope and source + +Resolve the three integration conflicts in `gguf_keep_quant.cpp`, +`rocm_grouped_gemm.hip`, and `test_backend_cross_device.cpp`. +Preserve both the target's behavior and the contribution's IQ4_XS and IQ3_XXS +admission, dense kernels, grouped kernels, and tests. +Do not import #3036 or redesign a kernel. + +The parent spec defines the source algorithms and device gates. +`git log -S kIQ3_XXS -- src/vllm/model_executor/model_loader/gguf_keep_quant.cpp` +identifies `acd7d457d` as the loader admission change. +The contributor's routing test explicitly omits IQ3_XXS. Its direct operation +tests cannot detect deletion of IQ3_XXS from loader admission. + +## Design and tests + +Add a test through `RouteGgufTensor` for ROCm IQ3_XXS matrix and stacked +expert weights. Both aligned roles must keep their blocks. Ragged shapes, +disabled keep-quant, and CPU-reference mode must still expand to bf16. +Keep the broad existing routing table unchanged except for its obsolete +coverage comment. The new test owns the formerly missing admission guarantee. + +Before accepting the regression, remove IQ3_XXS from the ROCm admission arm +in a scratch copy and require the test to fail. Restore the original source +and require the loader suite to pass. +Run focused loader and device-fit tests plus `scripts/agent-preflight.sh`. +The operator builds HIP and runs the contributor's device tests under a lease. +Generic I-quant checks on Strix do not establish gfx1200 performance. + +## Records and stop conditions + +If a keyed record conflicts, start from its complete target version and reapply +only this row's edit. Verify unrelated keys against the target byte-for-byte. +No lifecycle change or new benchmark publication belongs to this repair. +Stop if conflict resolution requires choosing between incompatible behaviors, +changes residency-budget semantics, or needs a new kernel design. + +## Now + +ACTIVE: integration and loader regression specified before code changes. From 8dbe77a5db8b1f2cd334c17cf034b57dcfff9117 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 8 Sep 2026 07:35:40 +0000 Subject: [PATCH 04/11] test(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): cover IQ3_XXS loader admission Direct kernel tests do not detect deleting IQ3_XXS from ROCm loader admission. Enter RouteGgufTensor for matrix and expert weights, retaining ragged-shape and switch checks. This closes the missing loader guarantee in #3029 without closing the remaining formats owed by #1940. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GPT-6 [Codex] --- tests/vllm/test_gguf_keep_quant.cpp | 27 ++++++++++++++++++++++++--- 1 file changed, 24 insertions(+), 3 deletions(-) diff --git a/tests/vllm/test_gguf_keep_quant.cpp b/tests/vllm/test_gguf_keep_quant.cpp index 969d066e2b..d0002a073c 100644 --- a/tests/vllm/test_gguf_keep_quant.cpp +++ b/tests/vllm/test_gguf_keep_quant.cpp @@ -535,6 +535,28 @@ TEST_CASE("every encoding in the Qwen3.8-2.4T UD-Q1_0 checkpoint decodes") { CheckCheckpointCensus(kUdQ10Census, "UD-Q1_0"); } +TEST_CASE("ROCm IQ3_XXS loader admission keeps matrix and expert blocks") { + // Direct GEMM tests bypass DeviceKeepQuantSupported. Enter the loader route + // so deleting IQ3_XXS admission cannot silently restore bf16 expansion. + constexpr uint32_t iq3_xxs = 18; + for (const auto role : {GgufTensorRole::kMatmulWeight, + GgufTensorRole::kStackedExpertWeight}) { + CAPTURE(vllm::Name(role)); + std::vector shape = role == GgufTensorRole::kMatmulWeight + ? std::vector{8, 256} + : std::vector{2, 8, 256}; + auto route = [&](bool keep_quant, bool cpu_ref) { + return RouteGgufTensor(keep_quant, /*keep_f16=*/false, /*nvfp4_fp4=*/false, + cpu_ref, role, iq3_xxs, shape, vt::DeviceType::kROCM); + }; + CHECK(route(true, false) == GgufResidency::kKeepQuant); + CHECK(route(false, false) == GgufResidency::kExpandBf16); + CHECK(route(true, true) == GgufResidency::kExpandBf16); + shape.back() = 255; + CHECK(route(true, false) == GgufResidency::kExpandBf16); + } +} + TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { // The expectation is written out LONGHAND here rather than derived from the // implementation, so this is a real cross-check and not a tautology. @@ -593,9 +615,8 @@ TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { // exactly {Q8_0, Q4_K, Q5_K, Q6_K} through #2247; IQ4_XS joined with // KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940). IQ3_XXS joined the same // row's ROCm kernel set but is not in `all_types` below (nor is - // Q2_K, a pre-existing gap this row does not close), so it is not - // exercised by this particular exhaustive table; it is gated instead - // by tests/vt/test_backend_cross_device.cpp's cross-device cases. + // Q2_K, a pre-existing gap this row does not close). The dedicated + // IQ3_XXS loader-admission case above gates its production route. // QUANT-GGUF-IQ-VECDOT (#2247) put IQ2_XS and IQ4_XS in this list. // They were gather-only between #2245 and #2247 — decoder, no vec_dot — // and the `vec_dot` rows are what moved them onto the GEMM arm. From fdeb066dea5f588011ac36d1ecb0b59eaccd38c9 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 8 Sep 2026 07:40:13 +0000 Subject: [PATCH 05/11] docs(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): correct footprint units The contributor measured 19,333,564,672 bytes, which is about 18.01 GiB. Correct only the mislabeled conversion in #3029 while preserving the raw measurement and all other evidence. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GPT-6 [Codex] --- .agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md b/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md index 1f0ffedce1..01a49a2e24 100644 --- a/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md +++ b/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md @@ -276,7 +276,7 @@ vllm-cli: run=1/1 finish_reason=length prompt_tokens=5 completion_tokens=16 secs Paris. The capital of Germany is Berlin. The capital of Italy is Rome. ``` -The decisive number is the **19,333,564,672 B (~19.33 GiB) footprint** -- +The decisive number is the **19,333,564,672 B (~18.01 GiB) footprint** -- it matches the file's on-disk size, not the ~70 GiB a bf16 expansion of these tensors would produce. That is the keep-quant residency actually taking effect on ROCm, not merely compiling: before this row, @@ -336,7 +336,7 @@ sufficient, no CUDA-style non-fused-multiply workaround needed. **The real-checkpoint end-to-end reload also LANDED**, after this pull request was first drafted: `Nail-Qwen3.6-35B-A3B-MTP-IQ4_XS.gguf` loads and generates coherent tokens on `isravale`, with the resident footprint -(~19.33 GiB) matching the on-disk size rather than a bf16 blow-up — see +(~18.01 GiB) matching the on-disk size rather than a bf16 blow-up — see Tests. That was the row's actual acceptance criterion, and it is now satisfied on the artifact that motivated the row, not a synthetic stand-in. From 6f3e1da0ca3618b3f627d7c57aeaca11c8ecbc86 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 8 Sep 2026 07:45:39 +0000 Subject: [PATCH 06/11] spec(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): pin the repaired gate baseline The operator landed #3064 and selected main 08a34c3a for the final #3029 integration. Record that exact target before merging it, retaining the first integration pin. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GPT-6 [Codex] --- .agents/specs/rocm-iquant-integration-3029.md | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/.agents/specs/rocm-iquant-integration-3029.md b/.agents/specs/rocm-iquant-integration-3029.md index 2e1853fe8f..ec64125509 100644 --- a/.agents/specs/rocm-iquant-integration-3029.md +++ b/.agents/specs/rocm-iquant-integration-3029.md @@ -6,7 +6,9 @@ open for the other quantized formats. Contribution: [#3029](https://github.com/mudler/vllm.cpp/pull/3029). Parent spec: [ROCm I-quant port](kernel-quant-ciq-gemm-rocm-iquant.md). Contribution base: `7aa0aa00a8eb79d53e65685e78e8da6d3f10482a`. -Integration target: `415d17859500caf2a4cac00511820e4f4760e86f`. +Initial integration target: `415d17859500caf2a4cac00511820e4f4760e86f`. +Final integration target: `08a34c3a74d78046f83886f242d07110a70ff45e`, which +includes the prerequisite README scan repair from #3064. ## Scope and source From 99ee5bb29bf1a28049c69937db03ed7982bf9109 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 8 Sep 2026 08:07:27 +0000 Subject: [PATCH 07/11] spec(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): seal device tables The existing spec requires a device-byte seal, but the contribution only checks numerical samples. Define the missing gate for #3067. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:GPT-5 [API] --- .agents/specs/rocm-iquant-table-seal-3067.md | 59 ++++++++++++++++++++ 1 file changed, 59 insertions(+) create mode 100644 .agents/specs/rocm-iquant-table-seal-3067.md diff --git a/.agents/specs/rocm-iquant-table-seal-3067.md b/.agents/specs/rocm-iquant-table-seal-3067.md new file mode 100644 index 0000000000..d9c2709761 --- /dev/null +++ b/.agents/specs/rocm-iquant-table-seal-3067.md @@ -0,0 +1,59 @@ +# Seal every ROCm I-quant table against its CPU reference + +Row: `KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT`. +Issue: [#3067](https://github.com/mudler/vllm.cpp/issues/3067). +Contribution: [#3029](https://github.com/mudler/vllm.cpp/pull/3029). +Parent spec: [ROCm I-quant port](kernel-quant-ciq-gemm-rocm-iquant.md). +Repair base: `b2ee9d8389caf974f3178613a6313e788dd93c4b`. + +## Gap and source + +The parent spec and `rocm_quant_iq_tables.h:17-20` require a complete seal. +Review found no executing comparison of the four ROCm device tables against +their CPU references. Host-parsed arrays currently agree, but that inspection +does not pin every executing device byte in a regression test. + +Mirror the existing CUDA snapshot in `cuda_quant_dot.cu:2637` and the test in +`test_cuda_quant_dot.cpp:1702`. The CPU tables carry the parent spec's pinned +llama.cpp reference. This change adds no quantization algorithm or oracle. + +## Design and scope + +Add a HIP-free internal snapshot declaration for the four arrays: +`d_kmask_iq2xs`, `d_ksigns_iq2xs`, `d_iq3xxs_grid`, and `d_kvalues_iq4nl`. +Define the copy in `rocm_grouped_gemm.hip`, which defines the device symbols. +Use `hipMemcpyFromSymbol` and the existing HIP error checker. Compile-time +extent checks prevent truncation. Do not change table values, storage classes, +arithmetic, dispatch, or any other quantized format. + +Add a HIP test to `test_backend_cross_device.cpp`. Compare each complete +snapshot array with its CPU reference using `memcmp`. Check all four extents +and the number of comparisons. Do not substitute host literals for device +copies. Follow the executable's missing-backend convention; a skipped device +case is not device evidence. + +## Tests and gates + +Commit the test and interface before the copy implementation. The missing +implementation is the initial compile/link gap. The operator then executes +the completed seal under a HIP lease and mutates one entry in each of the four +device tables separately. Every mutation must fail its named comparison; +restore the table byte-for-byte between runs. Deleting a snapshot copy must +also fail, since a seal must observe each symbol rather than compare nothing. + +Run the focused CPU loader and device-fit tests and the complete host +preflight. Host builds do not establish HIP correctness. Run the existing +I-quant numerical device gates with the seal under the operator's lease. +Independent scoped review and the operator's own gate remain required. + +## Risks and stop conditions + +A passing tolerance-based dot test can miss an unvisited table entry. The +byte-exact seal closes only that gap, not gfx1200 performance or the remaining +formats owned by #1940. Keep that broader issue open. +Stop if the registered HIP implementation cannot expose its actual device +symbols, or if source/oracle disagreement requires changing table values. + +## Now + +ACTIVE: the device-byte seal is specified before its test and implementation. From 8216be73edd6227ab705d0e60d563190e7474c74 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 8 Sep 2026 08:09:02 +0000 Subject: [PATCH 08/11] test(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): compare device codebooks Compare every byte of all four device tables, including their extents. The snapshot implementation is still missing at this test-first step (#3067). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:GPT-5 [API] --- src/vt/rocm/rocm_iq_table_seal.h | 23 ++++++++++++++++++ tests/vt/test_backend_cross_device.cpp | 33 ++++++++++++++++++++++++++ 2 files changed, 56 insertions(+) create mode 100644 src/vt/rocm/rocm_iq_table_seal.h diff --git a/src/vt/rocm/rocm_iq_table_seal.h b/src/vt/rocm/rocm_iq_table_seal.h new file mode 100644 index 0000000000..94700dd1f5 --- /dev/null +++ b/src/vt/rocm/rocm_iq_table_seal.h @@ -0,0 +1,23 @@ +// Internal device-table snapshot for the byte-exact I-quant gate (#3067). +// Mirror cuda_iq_table_seal.h without exposing HIP types to the host test. +#ifndef VT_ROCM_IQ_TABLE_SEAL_H_ +#define VT_ROCM_IQ_TABLE_SEAL_H_ + +#include + +namespace vt::rocm { + +struct IqTableSnapshot { + uint8_t kmask_iq2xs[8]; + uint8_t ksigns_iq2xs[128]; + uint32_t iq3xxs_grid[256]; + int8_t kvalues_iq4nl[16]; +}; + +// Copy the actual device symbols into out. Requires a live HIP context and +// throws on copy failure. Defined beside the symbols in rocm_grouped_gemm.hip. +void SnapshotIqTablesFromDevice(IqTableSnapshot* out); + +} // namespace vt::rocm + +#endif // VT_ROCM_IQ_TABLE_SEAL_H_ diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index ac1584e5e4..51d88529d8 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -46,6 +46,10 @@ #include "vt/recipes.h" #include "vt/rocm/rocm_arch.h" #include "vt/rocm/rocm_runtime.h" +#if defined(VLLM_CPP_HIP) +#include "../../src/vt/cpu/cpu_quant_iq_tables.h" +#include "../../src/vt/rocm/rocm_iq_table_seal.h" +#endif // KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940): the real-checkpoint IQ4_XS // blocks and llama.cpp's own vec_dot output for them, shared with the CUDA @@ -2673,6 +2677,35 @@ TEST_CASE("keep-quant Q6_K GEMM runs at the production launch geometry") { } #if defined(VLLM_CPP_HIP) +// Mirror test_cuda_quant_dot.cpp's device-byte seal. Numerical samples can +// leave entries unvisited, so compare the complete executing tables (#3067). +TEST_CASE("ROCm device codebooks == the CPU host tables (byte-exact)") { + if (vt::TryGetBackend(DeviceType::kROCM) == nullptr) { + MESSAGE("no ROCm backend on this host; device-codebook seal skipped"); + return; + } + vt::GetBackend(DeviceType::kROCM); + vt::rocm::IqTableSnapshot snap{}; + vt::rocm::SnapshotIqTablesFromDevice(&snap); + + int sealed = 0; + auto seal = [&](const std::string& name, const void* device, const void* host, size_t bytes) { + CAPTURE(name); + CHECK(std::memcmp(device, host, bytes) == 0); + ++sealed; + }; + seal("d_kmask_iq2xs", snap.kmask_iq2xs, vt::cpu::kKmaskIq2xs, sizeof(snap.kmask_iq2xs)); + seal("d_ksigns_iq2xs", snap.ksigns_iq2xs, vt::cpu::kKsignsIq2xs, sizeof(snap.ksigns_iq2xs)); + seal("d_iq3xxs_grid", snap.iq3xxs_grid, vt::cpu::kIq3xxsGrid, sizeof(snap.iq3xxs_grid)); + seal("d_kvalues_iq4nl", snap.kvalues_iq4nl, vt::cpu::kValuesIq4nl, + sizeof(snap.kvalues_iq4nl)); + CHECK(sealed == 4); + CHECK(sizeof(snap.kmask_iq2xs) == sizeof(vt::cpu::kKmaskIq2xs)); + CHECK(sizeof(snap.ksigns_iq2xs) == sizeof(vt::cpu::kKsignsIq2xs)); + CHECK(sizeof(snap.iq3xxs_grid) == sizeof(vt::cpu::kIq3xxsGrid)); + CHECK(sizeof(snap.kvalues_iq4nl) == sizeof(vt::cpu::kValuesIq4nl)); +} + // Declared here rather than included: the ROCm kernels have no public header, // and src/vt/rocm/rocm_ops.hip:65 already reaches MatmulBTQuantKernelRocm by a // file-local extern declaration. This row mirrors that convention instead of From 8a7cda269a7254dc849f6401a990c5f2bae98e36 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 8 Sep 2026 08:10:18 +0000 Subject: [PATCH 09/11] test(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): snapshot HIP tables Copy each executing device symbol into the complete host snapshot for #3067. Preserve table values and quantized dispatch. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:GPT-5 [API] --- src/vt/rocm/rocm_grouped_gemm.hip | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index 0cb34ddff2..8fcef1d97c 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -40,6 +40,7 @@ #include "vt/rocm/rocm_arch.h" #include "vt/rocm/rocm_device_bind.h" #include "vt/rocm/rocm_f16_codec.h" +#include "vt/rocm/rocm_iq_table_seal.h" #include "vt/rocm/rocm_runtime.h" // Block layouts — the single source of truth (ggml-common.h mirrors). @@ -1363,6 +1364,23 @@ bool Q6KSmallPrivateEnabled() { } // namespace +// Mirror the CUDA seal in cuda_quant_dot.cu. Copy in the translation unit +// that owns these symbols so the test sees device bytes, not host literals. +void SnapshotIqTablesFromDevice(IqTableSnapshot* out) { + static_assert(sizeof(out->kmask_iq2xs) == sizeof(d_kmask_iq2xs), "kmask extent"); + static_assert(sizeof(out->ksigns_iq2xs) == sizeof(d_ksigns_iq2xs), "ksigns extent"); + static_assert(sizeof(out->iq3xxs_grid) == sizeof(d_iq3xxs_grid), "iq3xxs grid extent"); + static_assert(sizeof(out->kvalues_iq4nl) == sizeof(d_kvalues_iq4nl), "iq4nl values extent"); + Check(hipMemcpyFromSymbol(out->kmask_iq2xs, HIP_SYMBOL(d_kmask_iq2xs), sizeof(d_kmask_iq2xs)), + "snapshot d_kmask_iq2xs"); + Check(hipMemcpyFromSymbol(out->ksigns_iq2xs, HIP_SYMBOL(d_ksigns_iq2xs), sizeof(d_ksigns_iq2xs)), + "snapshot d_ksigns_iq2xs"); + Check(hipMemcpyFromSymbol(out->iq3xxs_grid, HIP_SYMBOL(d_iq3xxs_grid), sizeof(d_iq3xxs_grid)), + "snapshot d_iq3xxs_grid"); + Check(hipMemcpyFromSymbol(out->kvalues_iq4nl, HIP_SYMBOL(d_kvalues_iq4nl), sizeof(d_kvalues_iq4nl)), + "snapshot d_kvalues_iq4nl"); +} + // Warps that cooperate on one output row of the dense K-quant GEMM. 1 selects // the single-warp `KQuantGemmK` the path has always had. // From fa39a45a3c691c6415b8751c097f1a2df6f3be1c Mon Sep 17 00:00:00 2001 From: Justin Card Date: Tue, 8 Sep 2026 17:34:03 -0400 Subject: [PATCH 10/11] docs(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): record the ROCm IQ4_XS and IQ3_XXS coverage A maintainer review found that this branch changes which GGUF formats a ROCm keep-quant load admits while updating no public document and no owning inventory row. The tree carried the capability and no record of it, so a reader asking whether ROCm keeps an IQ4_XS weight compressed had to read the admission switch to find out. The reviewer named `docs/STATUS.md` and `docs/BENCHMARKS.md`. This repository has no `docs/STATUS.md`, and AGENTS.md ties `docs/BENCHMARKS.md` to a benchmark identifier being added, removed, or changing disposition, which this change does not touch. The surface AGENTS.md ties to a quantization surface changing is `docs/FEATURES.md`, and the owning inventory rows are `QUANT-GGUF-IQ4_XS` and `QUANT-GGUF-IQ3_XXS` in `.agents/quantization-matrix.md`. `.agents/kernel-matrix.md` tracks no row for this work and is not an owning surface here. The two device gates are recorded separately because they measured different things. The contributor gate on `gfx1200` covers this landing. The later operator gate on `gfx1151` covers the device-table seal, which is a different scope. Neither reloads the real checkpoint on the repaired head, and the rows say so rather than rounding a historical weight-footprint calculation up to a measured peak RSS. The owed `ROCM-KQUANT-NWARPS-DECODE` re-measurement and the follow-on WMMA row are named rather than dropped. The `R`/`M`/`C`/`E`/`P` columns do not move. `C` already read `Y` from the CPU and CUDA arms, and it says only that native quantized compute is dispatched, without naming a backend. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/quantization-matrix.md | 4 ++-- docs/FEATURES.md | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index d7364fe589..fee747a066 100644 --- a/.agents/quantization-matrix.md +++ b/.agents/quantization-matrix.md @@ -70,12 +70,12 @@ forces the full dequant path back. | `QUANT-GGUF-Q6_K` | Q6_K | 14 / output | Y | Y | Y | Y | - | `PARTIAL` | [dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L192), [dispatch](../src/vllm/model_executor/model_loader/gguf_dequant.cpp#L80); [unit](../tests/vllm/test_gguf_dequant.cpp#L134); [Compact gate](../tests/parity/test_qwen36_gguf_engine.cpp#L143), [Balanced gate](../tests/parity/test_qwen36_gguf_engine.cpp#L149) **`C` = `Y` since G4 (2026-07-22):** the weight stays in its ggml blocks at load ([keep-quant policy](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L95), DEFAULT ON wherever `kMatmulBTQuant` is registered for the running device) and [`vt::MatmulBT`](../src/vt/ops.cpp#L158) dispatches it to the [tier-0 `vec_dot` GEMM](../src/vt/cpu/cpu_quant_dot.cpp#L1) — no bf16 expansion on the executed path. Op gate [test_ops_quant_dot](../tests/vt/test_ops_quant_dot.cpp#L1); residency losslessness + routing gates [test_gguf_keep_quant](../tests/vllm/test_gguf_keep_quant.cpp#L1). 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`](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L230) (66 B block); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L322) (`DequantIQ2_XXS`, ported 1:1 from llama.cpp `ggml-quants.c:2416`); grids moved to the shared [cpu_quant_iq_tables.h](../src/vt/cpu/cpu_quant_iq_tables.h) (`iq2xxs_grid`/`ksigns_iq2xs`/`kmask_iq2xs`, `ggml-common.h:499,503,550`); vt block dtype `kIQ2_XXS` [geometry](../src/vt/dtype.cpp#L80), [dispatch](../src/vllm/model_executor/model_loader/gguf_dequant.cpp#L114). **`C` = `Y` since DeepSeek-V4 W8 (2026-07-29):** the keep-quant `vec_dot` [`VecDotIQ2_XXSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L1) (1:1 port of ggml `vec_dot_iq2_xxs_q8_K_generic`, `quants.c:855`) + the [Q8_K traits row](../src/vt/cpu/cpu_quant_traits.cpp#L1) 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](../tests/vt/test_ops_quant_dot.cpp#L1); 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](../src/vt/cuda/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](../tests/vt/test_cuda_quant_dot.cpp)) 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](specs/cuda-keepquant-gemm.md); [iquant DSV4 spike](specs/gguf-iquant-dsv4.md) | `CLAIM-DEEPSEEK-V4-W8` | | `QUANT-GGUF-IQ2_XS` | IQ2_XS | 17 / output; IQ2_S preset storage | Y | Y | Y | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 17](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L254) (74 B block, ADDED by [#2240](https://github.com/mudler/vllm.cpp/issues/2240)); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L453) `DequantIQ2_XS` + the shared 512-entry [`kIq2xsGrid`](../src/vt/cpu/cpu_quant_iq_tables.h#L168), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2516` + `ggml/src/ggml-common.h:627`; vt block dtype `kIQ2_XS` [geometry](../src/vt/dtype.cpp#L131). **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.3.ffn_gate_exps.weight` from the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL arm, decoded by the pinned upstream's own `dequantize_row_iq2_xs` and compared as raw f32 bit patterns ([goldens + provenance](../tests/vt/iq2xs_iq4xs_golden_vectors.h#L72), [unit](../tests/vllm/test_gguf_dequant.cpp#L570)); the table itself is sealed by FNV-1a digest ([test_ops_quant_dot](../tests/vt/test_ops_quant_dot.cpp#L717)) because the IQ2 family carries three same-shaped grids and a wrong one still decodes. Reached through `GgufFile::Open` + `DequantGgufRowToF32` in the [same case](../tests/vllm/test_gguf_dequant.cpp#L590), which is the production path the real artifact took, and [geometry cross-check](../tests/vt/test_ops_quant_traits.cpp#L156) pins the reader and vt tables against each other. **`C` = `Y` since [#2247](https://github.com/mudler/vllm.cpp/issues/2247):** keep-quant [`VecDotIQ2_XSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L783) against the Q8_K activation, ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-cpu/quants.c:948` `ggml_vec_dot_iq2_xs_q8_K_generic`, with [`BlockIQ2_XS`](../src/vt/cpu/cpu_quant_blocks.h#L191) (`ggml-common.h:388-393`) and the [traits row](../src/vt/cpu/cpu_quant_traits.cpp#L111) (`ggml-cpu.c:342-347`), so `HasQuantDotKernel` is TRUE and the loader keeps the blocks. **Gated BIT FOR BIT against the oracle's OWN KERNEL on the same real artifact bytes the decoder gate uses** — 4 super-blocks of `blk.3.ffn_gate_exps.weight` dotted against a Q8_K activation the oracle itself encoded, total and per block ([goldens + provenance](../tests/vt/iq2xs_iq4xs_dot_golden.h#L268), [case](../tests/vt/test_ops_quant_dot.cpp#L857)); the comparison is against upstream's own f32 accumulation, because a cleaner f64 reference agrees with a reduction-order defect. The grid SEAL and the KERNEL are tied together by a mutation, not by inspection: swapping `kIq2xsGrid` for `kIq2xxsGrid` inside the kernel reds this case while the seal itself stays green ([coupling case](../tests/vt/test_ops_quant_dot.cpp#L879)). **Residency MEASURED, not inferred:** `RouteGgufTensor` driven over all 1412 tensors of the staged artifact's own headers routes 774 to `kKeepQuant` for **101.14 GiB resident** against 426.72 GiB before this row — the 82 IQ2_XS tensors stop expanding from 53.33 GiB to 369.00 GiB — which fits the ~119.63 GiB of `dgx:gpu0` with 18.49 GiB of headroom. [Routing table](../tests/vllm/test_gguf_keep_quant.cpp#L425) restated: the GEMM term moved 20 -> 24 and the GATHER term stayed 13. **The CUDA arm now HAS a keep-quant kernel for this dtype** ([#2260](https://github.com/mudler/vllm.cpp/issues/2260), spec [cuda-keepquant-iq2xs-iq4xs](specs/cuda-keepquant-iq2xs-iq4xs.md)): `src/vt/cuda/cuda_quant_dot.cu::DotIQ2XS` is a 1:1 port of `b10451 quants.c:948` reading the generated `src/vt/cuda/cuda_quant_iq_tables.cuh::d_iq2xs_grid`, and `src/vt/cuda/cuda_quant_dot.cu::IsCudaKeepQuantSupported` admits twelve Q8_K-family encodings with IQ2_XS among them, wired into all THREE dispatch switches. Before that the 82 tensors still kept — `src/vllm/model_executor/model_loader/gguf_keep_quant.cpp::DeviceKeepQuantSupported` returns `true` for CUDA on its `default:` arm regardless — so the artifact DID fit and the expert GEMM then took the CPU fallback inside `src/vt/cuda/cuda_quant_dot.cu::MatmulBTQuantGroupedKernelCuda` behind a full `cudaStreamSynchronize` per grouped call, while the fused `src/vt/cuda/cuda_quant_dot.cu::MoeGateUpSwiGLUGroupedCuda` seam THREW `gate/up must be the SAME CUDA keep-quant dtype`. **101.14 GiB remains a RESIDENCY result and no speed number is claimed**, and the kernel is not yet SELECTED by a running model: `Glm5NextHostForward` refuses a non-CPU queue by name, so the wiring is still owed by `MODEL-MM-glm5-next-glm5-next-for-conditional-generation`. NOT reached from this row today — `glm5_next_moe.cpp` is W5's host reference and `laguna.cpp` is the only model on the fused seam — but AGENTS.md routes mergeable MLP projections through `vt::MergedGemmGroup`, so W5b ([#2241](https://github.com/mudler/vllm.cpp/issues/2241)) and W5c ([#2242](https://github.com/mudler/vllm.cpp/issues/2242)) make the throw live on a 101 GiB-resident model. Contrast `QUANT-GGUF-IQ2_XXS` above, which records a CUDA compute landing; this row has none. Carried as O19 in [glm5-next-flash](specs/glm5-next-flash.md). | [glm5-next-flash](specs/glm5-next-flash.md) | - | -| `QUANT-GGUF-IQ3_XXS` | IQ3_XXS | 18 / output | Y | Y | Y | - | - | `READY` | reader trait [`:237`](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L237) (98 B block, ADDED W8); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L349) (`DequantIQ3_XXS` + the shared [`iq3xxs_grid`](../src/vt/cpu/cpu_quant_iq_tables.h), ported 1:1 from llama.cpp `ggml-quants.c:2503` + `ggml-common.h:1007`); vt block dtype `kIQ3_XXS` [geometry](../src/vt/dtype.cpp#L86). **`C` = `Y` (DeepSeek-V4 W8, 2026-07-29):** the keep-quant `vec_dot` [`VecDotIQ3_XXSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L1) (1:1 port of ggml `vec_dot_iq3_xxs_q8_K_generic`, `quants.c:999`) + the [Q8_K traits row](../src/vt/cpu/cpu_quant_traits.cpp#L1) 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](../tests/vt/test_ops_quant_dot.cpp#L1) (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](../src/vt/cuda/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](../tests/vt/test_cuda_quant_dot.cpp)) 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](specs/cuda-keepquant-gemm.md); [iquant DSV4 spike](specs/gguf-iquant-dsv4.md) | - | +| `QUANT-GGUF-IQ3_XXS` | IQ3_XXS | 18 / output | Y | Y | Y | - | - | `READY` | reader trait [`:237`](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L237) (98 B block, ADDED W8); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L349) (`DequantIQ3_XXS` + the shared [`iq3xxs_grid`](../src/vt/cpu/cpu_quant_iq_tables.h), ported 1:1 from llama.cpp `ggml-quants.c:2503` + `ggml-common.h:1007`); vt block dtype `kIQ3_XXS` [geometry](../src/vt/dtype.cpp#L86). **`C` = `Y` (DeepSeek-V4 W8, 2026-07-29):** the keep-quant `vec_dot` [`VecDotIQ3_XXSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L1) (1:1 port of ggml `vec_dot_iq3_xxs_q8_K_generic`, `quants.c:999`) + the [Q8_K traits row](../src/vt/cpu/cpu_quant_traits.cpp#L1) 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](../tests/vt/test_ops_quant_dot.cpp#L1) (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](../src/vt/cuda/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](../tests/vt/test_cuda_quant_dot.cpp)) 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. **The ROCm arm now HAS a keep-quant kernel for this dtype too** ([#1940](https://github.com/mudler/vllm.cpp/issues/1940), spec [kernel-quant-ciq-gemm-rocm-iquant](specs/kernel-quant-ciq-gemm-rocm-iquant.md)): [`DotIQ3XXS`](../src/vt/rocm/rocm_grouped_gemm.hip#L345) is ADAPTED from the CUDA body rather than freshly ported, on both the plain and the grouped/MoE arm, and the `kROCM` case of [`DeviceKeepQuantSupported`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L138) now admits IQ3_XXS beside the four k-quants it already carried, so a ROCm keep-quant load stops expanding these blocks to bf16. `C` does NOT move and stays `Y`: it already read `Y` for the CPU and CUDA arms, and it says only that native quantized compute is dispatched, without naming a backend. **The device gate is CONTRIBUTOR evidence on `isravale`** (RX 9060 XT, `gfx1200`, ROCm 7.2.3) under the GPU file mutex: `test_backend_cross_device` **41/42** — the single failure is `MoeSiluMul`, the pre-existing tracked [#1954](https://github.com/mudler/vllm.cpp/issues/1954) flake, confirmed unrelated to this row — `test_gguf_keep_quant` **52/52**, `test_gguf_device_fit` and `test_gguf_device_fit_reach` green, and zero regression on the four existing k-quant formats. The separate, LATER operator run on `strix:gpu0` covers the device-table seal and NOT this landing; it and the owed performance work are recorded once, in the `QUANT-GGUF-IQ4_XS` row below. | [CUDA keep-quant GEMM](specs/cuda-keepquant-gemm.md); [iquant DSV4 spike](specs/gguf-iquant-dsv4.md) | - | | `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 | Y | Y | - | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 21](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L305) (110 B block, ADDED by [#2510](https://github.com/mudler/vllm.cpp/issues/2510)); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L399) `DequantIQ3_S` + the 512-entry [`kIq3sGrid`](../src/vt/cpu/cpu_quant_iq_tables.h#L168), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2607` + `ggml/src/ggml-common.h:1052`; vt block dtype `kIQ3_S` [geometry](../src/vt/dtype.cpp#L150) and the CUDA gather codec [`DqIQ3_S`](../src/vt/cuda/cuda_quant_dequant.cuh#L476). **THE REFUSAL THIS ROW EXISTS FOR:** id 21 was the ONE hole in the reader's i-quant run 16..23, and `GgufFile::Open` therefore refused whole artifacts over it — `unsloth/Qwen3.8-27B-GGUF` `UD-Q4_K_M.gguf` stores 4 of its 866 tensors in IQ3_S and `UD-Q4_K_XL.gguf` stores 1 of 866, while every OTHER encoding in those files (F32, Q3_K, Q4_K, Q5_K, Q6_K, Q8_0, IQ4_NL, IQ4_XS) was already handled. That is the file the published Strix Halo rows ran, so [#2497](https://github.com/mudler/vllm.cpp/issues/2497) had to substitute the plain `Q4_K_M`. **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.11.ffn_gate.weight` at absolute offset 4,262,628,128 in the staged `Qwen3.8-27B-UD-Q4_K_M.gguf` (sha256 `322e194f…3482`), decoded by the pinned upstream's own `dequantize_row_iq3_s` and compared as raw f32 bit patterns ([goldens + provenance](../tests/vt/iq3s_golden_vectors.h#L69), [unit](../tests/vllm/test_gguf_dequant.cpp#L663)); the four blocks span 190 distinct 9-bit grid indices of which 73 are >= 256, so the `qh`-spliced ninth bit is exercised rather than left at zero. The 110-byte stride is the ORACLE's own `sizeof(block_iq3_s)`, printed by the harness, not a sum read off the struct. The table itself is sealed by FNV-1a digest and by its lane alphabet ([test_ops_quant_dot](../tests/vt/test_ops_quant_dot.cpp#L717)), because `kIq3sGrid` (512 u32) and `kIq3xxsGrid` (256 u32) are the same shape and a decoder pointed at the sibling still decodes. Reached through `GgufFile::Open` + `DequantGgufRowToF32` in the same file, which is the production path the real artifact took. **`C` = `-`, DELIBERATELY, and it is the one FILE encoding in this tree that decodes without dotting.** `KeepQuantDType` gates on `HasQuantDotKernel`, so IQ3_S routes `kExpandBf16` on the GEMM arm of EVERY tier (CPU, CUDA, ROCm, Metal, Vulkan) and `kKeepQuant` only on the gather, where CPU and CUDA both decode. That is asserted BY NAME rather than left to a count ([routing](../tests/vllm/test_gguf_keep_quant.cpp#L429), [decode-only class](../tests/vt/test_ops_quant_traits.cpp#L250)), so the day the `vec_dot` lands those cases red. **The CPU `VecDotIQ3_SQ8_K` and the CUDA `WType::kIQ3_S` are OWED AS ONE UNIT** (`ggml-cpu/quants.c:1094`, `.vec_dot_type = GGML_TYPE_Q8_K` at `ggml-cpu.c:355-360`): `IsCudaKeepQuantSupported` has no refusal arm, so landing the CPU half alone would flip the loader to keep IQ3_S blocks and send every CUDA IQ3_S GEMM through the host fallback that [#2260](https://github.com/mudler/vllm.cpp/issues/2260) MEASURED to segfault on a discrete card. On the measured artifact the expansion is 146.13 MiB of blocks against 680.00 MiB of bf16, 3.4 % of a 15.33 GiB file — bounded here, and NOT bounded on an artifact whose experts are IQ3_S, which is why it is recorded rather than waved at (compare [#1870](https://github.com/mudler/vllm.cpp/issues/1870)). `E`/`P` = `-`: this row makes the artifact OPENABLE and claims no token and no number; the `gfx1151` quant-matched decode number stays owed by [#2497](https://github.com/mudler/vllm.cpp/issues/2497). | [gguf-iq3s](specs/gguf-iq3s.md) | - | | `QUANT-GGUF-IQ2_S` | IQ2_S | 22 / IQ2_M storage | Y | - | - | - | - | `INVENTORIED` | reader `:229-234`; [explicit rejection](../tests/vllm/test_gguf_dequant.cpp#L223) | leaf open | - | -| `QUANT-GGUF-IQ4_XS` | IQ4_XS | 23 / output | Y | Y | Y | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 23](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L311) (136 B block; the trait predates the decoder, the decoder is [#2240](https://github.com/mudler/vllm.cpp/issues/2240)); [dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L488) `DequantIQ4_XS`, ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2743`; vt block dtype `kIQ4_XS` [geometry](../src/vt/dtype.cpp#L140). NOT a codebook delta from IQ4_NL — it reuses `kValuesIq4nl` unchanged and differs only in the SUPER-BLOCK SCALE LAYOUT (a 6-bit `ls` spliced from a `scales_l` nibble and a `scales_h` bit pair, then biased by -32). **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.11.ffn_down_exps.weight` from the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL arm, whose four super-blocks span `ls` 0..61 so both nibbles and both `scales_h` bit pairs and both signs of `dl` are exercised ([goldens + provenance](../tests/vt/iq2xs_iq4xs_golden_vectors.h#L276), [unit](../tests/vllm/test_gguf_dequant.cpp#L576)). [Geometry cross-check](../tests/vt/test_ops_quant_traits.cpp#L156) pins the reader and vt tables against each other. The former `explicit rejection` evidence is GONE by construction — that test now guards Q1_0 (41), the remaining tabulated-but-undecodable id. **`C` = `Y` since [#2247](https://github.com/mudler/vllm.cpp/issues/2247):** keep-quant [`VecDotIQ4_XSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L844), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-cpu/quants.c:1283` `ggml_vec_dot_iq4_xs_q8_K_generic`, with [`BlockIQ4_XS`](../src/vt/cpu/cpu_quant_blocks.h#L205) (`ggml-common.h:454-460`) and the [traits row](../src/vt/cpu/cpu_quant_traits.cpp#L121). **The activation pairing is `Q8_K`, READ OFF the oracle and not inherited from IQ4_NL:** `type_traits_cpu` at `ggml/src/ggml-cpu/ggml-cpu.c:385-390` carries `.vec_dot = ggml_vec_dot_iq4_xs_q8_K, `.vec_dot_type = GGML_TYPE_Q8_K`, against `:379-384`'s `GGML_TYPE_Q8_0` for IQ4_NL — the 16-entry codebook is shared, the block geometry is not, and a 256-element super-block pairs with the 256-element activation encoding. **Gated BIT FOR BIT against the oracle's OWN KERNEL** on 4 super-blocks of `blk.11.ffn_down_exps.weight`, total and per block ([goldens](../tests/vt/iq2xs_iq4xs_dot_golden.h#L269), [case](../tests/vt/test_ops_quant_dot.cpp#L869)). These 3 tensors stop expanding from 3.59 GiB to 13.50 GiB; the whole-artifact measurement is in the `QUANT-GGUF-IQ2_XS` row above. **It also unblocks a SECOND consumer:** the expert-tower streaming lane is all-or-nothing (`GgufExpertTowersReachSlotLane` returns false on the FIRST tower that does not keep), so the 4 IQ4_XS tensors of the GLM-5.3 `UD-IQ1_S` arm were dropping an arm whose other 221 towers all kept — one tower 6.375 -> 24.000 GiB, and a 4096-slot cache 25.5 -> 96 GiB. Asserted in the `kStackedExpertWeight` role both models store these towers in ([lane case](../tests/vllm/model_executor/test_gguf_device_fit.cpp#L672)). **The CUDA arm now HAS a keep-quant kernel for this dtype too** ([#2260](https://github.com/mudler/vllm.cpp/issues/2260), spec [cuda-keepquant-iq2xs-iq4xs](specs/cuda-keepquant-iq2xs-iq4xs.md)): `src/vt/cuda/cuda_quant_dot.cu::DotIQ4XS` is a 1:1 port of `b10451 quants.c:1283` over the generated `src/vt/cuda/cuda_quant_iq_tables.cuh::d_kvalues_iq4nl`, keeping upstream's EIGHT f32 accumulation steps per super-block rather than one integer core, because that association is what the oracle's golden numbers were produced with. **This completes the device admission of the GLM-5.3 non-flash `UD-IQ1_S` arm**, whose other five encodings (IQ1_S, IQ3_XXS, IQ2_XXS, Q2_K, Q3_K) were already there. Before it, these 3 tensors fit `dgx:gpu0` with their expert GEMM on the CPU fallback behind a `cudaStreamSynchronize` and the fused `src/vt/cuda/cuda_quant_dot.cu::MoeGateUpSwiGLUGroupedCuda` seam throwing. The fit remains RESIDENCY and no speed number is claimed; O19's premise in [glm5-next-flash](specs/glm5-next-flash.md) is discharged, and what still blocks a CUDA run of either model is that model's own host-f32 forward. | [glm5-next-flash](specs/glm5-next-flash.md) | - | +| `QUANT-GGUF-IQ4_XS` | IQ4_XS | 23 / output | Y | Y | Y | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 23](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L311) (136 B block; the trait predates the decoder, the decoder is [#2240](https://github.com/mudler/vllm.cpp/issues/2240)); [dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L488) `DequantIQ4_XS`, ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2743`; vt block dtype `kIQ4_XS` [geometry](../src/vt/dtype.cpp#L140). NOT a codebook delta from IQ4_NL — it reuses `kValuesIq4nl` unchanged and differs only in the SUPER-BLOCK SCALE LAYOUT (a 6-bit `ls` spliced from a `scales_l` nibble and a `scales_h` bit pair, then biased by -32). **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.11.ffn_down_exps.weight` from the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL arm, whose four super-blocks span `ls` 0..61 so both nibbles and both `scales_h` bit pairs and both signs of `dl` are exercised ([goldens + provenance](../tests/vt/iq2xs_iq4xs_golden_vectors.h#L276), [unit](../tests/vllm/test_gguf_dequant.cpp#L576)). [Geometry cross-check](../tests/vt/test_ops_quant_traits.cpp#L156) pins the reader and vt tables against each other. The former `explicit rejection` evidence is GONE by construction — that test now guards Q1_0 (41), the remaining tabulated-but-undecodable id. **`C` = `Y` since [#2247](https://github.com/mudler/vllm.cpp/issues/2247):** keep-quant [`VecDotIQ4_XSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L844), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-cpu/quants.c:1283` `ggml_vec_dot_iq4_xs_q8_K_generic`, with [`BlockIQ4_XS`](../src/vt/cpu/cpu_quant_blocks.h#L205) (`ggml-common.h:454-460`) and the [traits row](../src/vt/cpu/cpu_quant_traits.cpp#L121). **The activation pairing is `Q8_K`, READ OFF the oracle and not inherited from IQ4_NL:** `type_traits_cpu` at `ggml/src/ggml-cpu/ggml-cpu.c:385-390` carries `.vec_dot = ggml_vec_dot_iq4_xs_q8_K, `.vec_dot_type = GGML_TYPE_Q8_K`, against `:379-384`'s `GGML_TYPE_Q8_0` for IQ4_NL — the 16-entry codebook is shared, the block geometry is not, and a 256-element super-block pairs with the 256-element activation encoding. **Gated BIT FOR BIT against the oracle's OWN KERNEL** on 4 super-blocks of `blk.11.ffn_down_exps.weight`, total and per block ([goldens](../tests/vt/iq2xs_iq4xs_dot_golden.h#L269), [case](../tests/vt/test_ops_quant_dot.cpp#L869)). These 3 tensors stop expanding from 3.59 GiB to 13.50 GiB; the whole-artifact measurement is in the `QUANT-GGUF-IQ2_XS` row above. **It also unblocks a SECOND consumer:** the expert-tower streaming lane is all-or-nothing (`GgufExpertTowersReachSlotLane` returns false on the FIRST tower that does not keep), so the 4 IQ4_XS tensors of the GLM-5.3 `UD-IQ1_S` arm were dropping an arm whose other 221 towers all kept — one tower 6.375 -> 24.000 GiB, and a 4096-slot cache 25.5 -> 96 GiB. Asserted in the `kStackedExpertWeight` role both models store these towers in ([lane case](../tests/vllm/model_executor/test_gguf_device_fit.cpp#L672)). **The CUDA arm now HAS a keep-quant kernel for this dtype too** ([#2260](https://github.com/mudler/vllm.cpp/issues/2260), spec [cuda-keepquant-iq2xs-iq4xs](specs/cuda-keepquant-iq2xs-iq4xs.md)): `src/vt/cuda/cuda_quant_dot.cu::DotIQ4XS` is a 1:1 port of `b10451 quants.c:1283` over the generated `src/vt/cuda/cuda_quant_iq_tables.cuh::d_kvalues_iq4nl`, keeping upstream's EIGHT f32 accumulation steps per super-block rather than one integer core, because that association is what the oracle's golden numbers were produced with. **This completes the device admission of the GLM-5.3 non-flash `UD-IQ1_S` arm**, whose other five encodings (IQ1_S, IQ3_XXS, IQ2_XXS, Q2_K, Q3_K) were already there. Before it, these 3 tensors fit `dgx:gpu0` with their expert GEMM on the CPU fallback behind a `cudaStreamSynchronize` and the fused `src/vt/cuda/cuda_quant_dot.cu::MoeGateUpSwiGLUGroupedCuda` seam throwing. The fit remains RESIDENCY and no speed number is claimed; O19's premise in [glm5-next-flash](specs/glm5-next-flash.md) is discharged, and what still blocks a CUDA run of either model is that model's own host-f32 forward. **The ROCm arm now HAS a keep-quant kernel for this dtype too** ([#1940](https://github.com/mudler/vllm.cpp/issues/1940), spec [kernel-quant-ciq-gemm-rocm-iquant](specs/kernel-quant-ciq-gemm-rocm-iquant.md)): [`DotIQ4XS`](../src/vt/rocm/rocm_grouped_gemm.hip#L395) is ADAPTED from the CUDA body rather than freshly ported, on both the plain and the grouped/MoE arm, and the `kROCM` case of [`DeviceKeepQuantSupported`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L138) now admits IQ4_XS and IQ3_XXS beside the four k-quants it already carried. **CUDA's `DotIQ4XS` needed non-fused `__fmul_rn`/`__fadd_rn` to stay bit-exact and the HIP body does NOT, which is MEASURED and not assumed** ([rationale in place](../src/vt/rocm/rocm_grouped_gemm.hip#L384)): nvcc never sees the project-wide `-ffp-contract=off` that CMake already passes to `$`, and a bit-exact gate over the same real-checkpoint bytes and the same oracle bits the CUDA gate uses confirms plain arithmetic suffices on this toolchain, so the port carries no workaround CUDA needed. `C` does NOT move and stays `Y`: it already read `Y` for the CPU and CUDA arms, and it says only that native quantized compute is dispatched, without naming a backend. **The device gate is CONTRIBUTOR evidence on `isravale`** (RX 9060 XT, `gfx1200`, ROCm 7.2.3) under the GPU file mutex: `test_backend_cross_device` **41/42** — the single failure is `MoeSiluMul`, the pre-existing tracked [#1954](https://github.com/mudler/vllm.cpp/issues/1954) flake, confirmed unrelated to this row — `test_gguf_keep_quant` **52/52**, `test_gguf_device_fit` and `test_gguf_device_fit_reach` green, and zero regression on the four existing k-quant formats. **A SECOND, LATER operator run is a DIFFERENT scope and must not be read as covering this landing:** it verified the device-table seal ([#3067](https://github.com/mudler/vllm.cpp/issues/3067)) on `strix:gpu0`, `gfx1151`, ROCm 7.2.4 under lease `886172b3-9d21-4435-a9db-cbf7b2f5dfb6`, where four ctest targets passed at 61 focused and 9 codebook-comparison assertions and all eleven independent review mutations failed their intended tests. **NO end-to-end real-model reload is gated by either run.** `Nail-Qwen3.6-35B-A3B-MTP-IQ4_XS.gguf` (19.39 GB) loaded and emitted coherent text on the ORIGINAL, pre-repair head only; that checkpoint has NO recorded repository revision and NO sha256, its 19,333,564,672 B weight footprint is a CALCULATION and not a measured peak RSS, and coherent text is not a token-ID oracle comparison. **Speed is OWED and named rather than dropped:** the `ROCM-KQUANT-NWARPS-DECODE` re-measurement for these two formats is `PENDING` — no `rocprofv3` profiling setup is staged on that box — and does not block a row whose own scope is coverage and correctness; a WMMA/tensor-core tile for the two is the follow-on row `KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT-RDNA4`. | [glm5-next-flash](specs/glm5-next-flash.md) | - | | `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 | - | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 18a2c3ef43..fe83f45c44 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -339,7 +339,7 @@ from a prompt: both refuse, naming what is missing. | CPU (x86, Arm i8mm; A76 assembly correct/default, llama speed gate open, and the closed 20-core floor ran a SUPERSEDED fork denominator rather than the stock `b10451` pin, re-take owed #1003) | ✅ `CPU_ATTN` registered (#1371/#1392, [spec](../.agents/specs/attn-validate-configuration.md)) | ◐ | ☐ | ✅ | | Metal (Apple Silicon) | ✅ builds under Apple Clang with project warnings promoted to errors, the Qwen3.5 MoE loader included; its layout-refusal path uses the same messages and behavior on every platform (#1054) | ☐ | ☐ | ✅ | | Vulkan | ◐ **EXL3 is now NATIVE** ([#2530](https://github.com/mudler/vllm.cpp/issues/2530), [spec](../.agents/specs/backend-vulkan-exl3.md)): `kCastF16` and `kExl3Gemm` were the only two ops an EXL3 checkpoint ran on the CPU reference tier on a Vulkan queue, MEASURED at two fallback notices and now at ZERO. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone exceeds Vulkan's 16 KiB guarantee before one reaches `mma.sync`, `ldmatrix`, `cp.async` or a grid-wide barrier Vulkan has at no version -- so the gate is BYTE equality with the CPU arm on all three codebooks and every width, not a tolerance. Verified on llvmpipe with NO GPU and NO lease, which is what makes it a CI gate rather than a lease gate; no real Vulkan GPU and no checkpoint run, and NO speed number is claimed on any axis. `kExl3MoeMlp` and `kExl3HadR128` are owed | ☐ | ☐ | ✅ | -| ROCm | W0: 5 gfx archs; dense/GDN all-native; 0.8B dispatch fixed. **M4:** Qwen3-0.6B/3.5-0.8B 16/16 (#41). **M3:** `ROCM_ATTN` registered (#1056/#1065, [spec](../.agents/specs/rocm-attn-backend.md)). **EXL3 is now NATIVE on `gfx1151`** ([#2433](https://github.com/mudler/vllm.cpp/issues/2433), [spec](../.agents/specs/backend-rocm-exl3.md)): `kCastF16` and `kExl3Gemm` were the only two ops an EXL3 checkpoint still ran on the CPU reference tier, and a 3.0bpw Llama-3.2-1B now completes with ZERO reference-tier hits at 8.27 tok/s warm against 0.83 for the same tree with the two registrations disabled -- byte-identical greedy output either way. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone does not fit AMD's 64 KiB LDS; the gate is therefore byte equality with the CPU arm rather than a tolerance. No BF16 denominator and no clock attribution: the control HUNG the GPU in the same lease ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)). `kExl3MoeMlp` is owed. CPU parity open (#269) | Native ([ROCm guide](ROCM.md), device-specific wrapped-call-safe derived recount): full GDN, MoE combine/gate, keep-quant GEMM, EXL3 trellis GEMM; ctest-green gfx1151/1103/1100/1201/1200 (#41). **APU managed allocation is now MEASURED and NARROWED** ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)): `hipMallocManaged` on a part reporting `PageableMemoryAccess = 0` faulted the GPU in 17 of 21 legs on `gfx1151`, against 0 of 21 for plain `hipMalloc`, and the plain arm is the first thing that has completed a gate-sized Qwen3.8-27B Q4_K run on that board. The managed branch now requires a device that can take a recoverable page fault, and `UnifiedMemory()` follows the allocator -- so gfx1151 and gfx1103 lose the CPU reference tier and an op with no native ROCm kernel refuses by name there, naming the attribute and `VT_ROCM_MANAGED_ALLOC=1`. | ✅ | ✅ | +| ROCm | W0: 5 gfx archs; dense/GDN all-native; 0.8B dispatch fixed. **M4:** Qwen3-0.6B/3.5-0.8B 16/16 (#41). **M3:** `ROCM_ATTN` registered (#1056/#1065, [spec](../.agents/specs/rocm-attn-backend.md)). **EXL3 is now NATIVE on `gfx1151`** ([#2433](https://github.com/mudler/vllm.cpp/issues/2433), [spec](../.agents/specs/backend-rocm-exl3.md)): `kCastF16` and `kExl3Gemm` were the only two ops an EXL3 checkpoint still ran on the CPU reference tier, and a 3.0bpw Llama-3.2-1B now completes with ZERO reference-tier hits at 8.27 tok/s warm against 0.83 for the same tree with the two registrations disabled -- byte-identical greedy output either way. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone does not fit AMD's 64 KiB LDS; the gate is therefore byte equality with the CPU arm rather than a tolerance. No BF16 denominator and no clock attribution: the control HUNG the GPU in the same lease ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)). `kExl3MoeMlp` is owed. **ROCm keep-quant now covers IQ4_XS and IQ3_XXS** beside the Q8_0/Q4_K/Q5_K/Q6_K it already carried (`KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT`, [#1940](https://github.com/mudler/vllm.cpp/issues/1940), [spec](../.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md)), so an i-quant GEMM weight in either encoding stops expanding to bf16 on a ROCm keep-quant load; gated on `gfx1200` at `test_backend_cross_device` 41/42, the one failure being the pre-existing [#1954](https://github.com/mudler/vllm.cpp/issues/1954) `MoeSiluMul` flake. Q4_0/Q2_K/Q3_K/IQ2_XXS/IQ2_S/MXFP4 stay owed on the same issue, and no speed number is claimed. The per-format detail, the two DISTINCT device gates and the owed performance work are recorded in the `QUANT-GGUF-IQ4_XS` and `QUANT-GGUF-IQ3_XXS` rows of [the quantization matrix](../.agents/quantization-matrix.md). CPU parity open (#269) | Native ([ROCm guide](ROCM.md), device-specific wrapped-call-safe derived recount): full GDN, MoE combine/gate, keep-quant GEMM, EXL3 trellis GEMM; ctest-green gfx1151/1103/1100/1201/1200 (#41). **APU managed allocation is now MEASURED and NARROWED** ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)): `hipMallocManaged` on a part reporting `PageableMemoryAccess = 0` faulted the GPU in 17 of 21 legs on `gfx1151`, against 0 of 21 for plain `hipMalloc`, and the plain arm is the first thing that has completed a gate-sized Qwen3.8-27B Q4_K run on that board. The managed branch now requires a device that can take a recoverable page fault, and `UnifiedMemory()` follows the allocator -- so gfx1151 and gfx1103 lose the CPU reference tier and an op with no native ROCm kernel refuses by name there, naming the attribute and `VT_ROCM_MANAGED_ALLOC=1`. | ✅ | ✅ | | XPU / TPU | ☐ | ✅ | ◐ | ☐ | | Tenstorrent Blackhole | ◐ `ACTIVE`, OPT-125m 6/6; Qwen3-0.6B wired; Mistral-7B-v0.3 16/16 on P150 ([spec](../.agents/specs/tenstorrent-mistral.md)). 16x16 rerun and residual-RMS owed ([spec](../.agents/specs/tenstorrent-backend.md)) | ✅ | ☐ | ☐ | | Tenstorrent host-free decode | ◐ DEFAULT since #1604 (`0` opts out): no per-step host readback; 2.1x default-leg tok/s; both golden pairs re-adjudicated, both paged gates 16/16. Capture DEFAULT for Qwen3-dense since the #1625 flip (`0` opts out) — 27.7 vs 12.9 tok/s at c=1, 2.15x over the opt-out, [details](benchmarks/tt-capture-default-decode.md); Qwen3.5 dense (GDN) joined the default since [#2907](https://github.com/mudler/vllm.cpp/issues/2907), the decode-side GDN ops running device-resident under capture with [#2812](https://github.com/mudler/vllm.cpp/issues/2812)'s captured fatal fixed; 4B near-tie pair owed ([#2811](https://github.com/mudler/vllm.cpp/issues/2811)); async off (#1627) | ☐ | ☐ | ☐ | From 22800690b0bb456b3ee0df91f0fa43acec15a1d9 Mon Sep 17 00:00:00 2001 From: Justin Card Date: Tue, 8 Sep 2026 17:46:30 -0400 Subject: [PATCH 11/11] docs(KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT): record the operator's own gate reproduction The row's prior device-gate evidence was contributor-reported on an earlier head. Record the operator's independent rerun of the exact same four targets on isravale at the current head (fa39a45a3), so the row states what the operator itself verified rather than only what was reported to it. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] Claude-Session: https://claude.ai/code/session_01DmM255tHzVaY9VYum3qxLE --- .agents/quantization-matrix.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index fee747a066..f38908d4b5 100644 --- a/.agents/quantization-matrix.md +++ b/.agents/quantization-matrix.md @@ -70,12 +70,12 @@ forces the full dequant path back. | `QUANT-GGUF-Q6_K` | Q6_K | 14 / output | Y | Y | Y | Y | - | `PARTIAL` | [dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L192), [dispatch](../src/vllm/model_executor/model_loader/gguf_dequant.cpp#L80); [unit](../tests/vllm/test_gguf_dequant.cpp#L134); [Compact gate](../tests/parity/test_qwen36_gguf_engine.cpp#L143), [Balanced gate](../tests/parity/test_qwen36_gguf_engine.cpp#L149) **`C` = `Y` since G4 (2026-07-22):** the weight stays in its ggml blocks at load ([keep-quant policy](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L95), DEFAULT ON wherever `kMatmulBTQuant` is registered for the running device) and [`vt::MatmulBT`](../src/vt/ops.cpp#L158) dispatches it to the [tier-0 `vec_dot` GEMM](../src/vt/cpu/cpu_quant_dot.cpp#L1) — no bf16 expansion on the executed path. Op gate [test_ops_quant_dot](../tests/vt/test_ops_quant_dot.cpp#L1); residency losslessness + routing gates [test_gguf_keep_quant](../tests/vllm/test_gguf_keep_quant.cpp#L1). 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`](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L230) (66 B block); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L322) (`DequantIQ2_XXS`, ported 1:1 from llama.cpp `ggml-quants.c:2416`); grids moved to the shared [cpu_quant_iq_tables.h](../src/vt/cpu/cpu_quant_iq_tables.h) (`iq2xxs_grid`/`ksigns_iq2xs`/`kmask_iq2xs`, `ggml-common.h:499,503,550`); vt block dtype `kIQ2_XXS` [geometry](../src/vt/dtype.cpp#L80), [dispatch](../src/vllm/model_executor/model_loader/gguf_dequant.cpp#L114). **`C` = `Y` since DeepSeek-V4 W8 (2026-07-29):** the keep-quant `vec_dot` [`VecDotIQ2_XXSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L1) (1:1 port of ggml `vec_dot_iq2_xxs_q8_K_generic`, `quants.c:855`) + the [Q8_K traits row](../src/vt/cpu/cpu_quant_traits.cpp#L1) 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](../tests/vt/test_ops_quant_dot.cpp#L1); 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](../src/vt/cuda/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](../tests/vt/test_cuda_quant_dot.cpp)) 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](specs/cuda-keepquant-gemm.md); [iquant DSV4 spike](specs/gguf-iquant-dsv4.md) | `CLAIM-DEEPSEEK-V4-W8` | | `QUANT-GGUF-IQ2_XS` | IQ2_XS | 17 / output; IQ2_S preset storage | Y | Y | Y | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 17](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L254) (74 B block, ADDED by [#2240](https://github.com/mudler/vllm.cpp/issues/2240)); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L453) `DequantIQ2_XS` + the shared 512-entry [`kIq2xsGrid`](../src/vt/cpu/cpu_quant_iq_tables.h#L168), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2516` + `ggml/src/ggml-common.h:627`; vt block dtype `kIQ2_XS` [geometry](../src/vt/dtype.cpp#L131). **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.3.ffn_gate_exps.weight` from the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL arm, decoded by the pinned upstream's own `dequantize_row_iq2_xs` and compared as raw f32 bit patterns ([goldens + provenance](../tests/vt/iq2xs_iq4xs_golden_vectors.h#L72), [unit](../tests/vllm/test_gguf_dequant.cpp#L570)); the table itself is sealed by FNV-1a digest ([test_ops_quant_dot](../tests/vt/test_ops_quant_dot.cpp#L717)) because the IQ2 family carries three same-shaped grids and a wrong one still decodes. Reached through `GgufFile::Open` + `DequantGgufRowToF32` in the [same case](../tests/vllm/test_gguf_dequant.cpp#L590), which is the production path the real artifact took, and [geometry cross-check](../tests/vt/test_ops_quant_traits.cpp#L156) pins the reader and vt tables against each other. **`C` = `Y` since [#2247](https://github.com/mudler/vllm.cpp/issues/2247):** keep-quant [`VecDotIQ2_XSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L783) against the Q8_K activation, ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-cpu/quants.c:948` `ggml_vec_dot_iq2_xs_q8_K_generic`, with [`BlockIQ2_XS`](../src/vt/cpu/cpu_quant_blocks.h#L191) (`ggml-common.h:388-393`) and the [traits row](../src/vt/cpu/cpu_quant_traits.cpp#L111) (`ggml-cpu.c:342-347`), so `HasQuantDotKernel` is TRUE and the loader keeps the blocks. **Gated BIT FOR BIT against the oracle's OWN KERNEL on the same real artifact bytes the decoder gate uses** — 4 super-blocks of `blk.3.ffn_gate_exps.weight` dotted against a Q8_K activation the oracle itself encoded, total and per block ([goldens + provenance](../tests/vt/iq2xs_iq4xs_dot_golden.h#L268), [case](../tests/vt/test_ops_quant_dot.cpp#L857)); the comparison is against upstream's own f32 accumulation, because a cleaner f64 reference agrees with a reduction-order defect. The grid SEAL and the KERNEL are tied together by a mutation, not by inspection: swapping `kIq2xsGrid` for `kIq2xxsGrid` inside the kernel reds this case while the seal itself stays green ([coupling case](../tests/vt/test_ops_quant_dot.cpp#L879)). **Residency MEASURED, not inferred:** `RouteGgufTensor` driven over all 1412 tensors of the staged artifact's own headers routes 774 to `kKeepQuant` for **101.14 GiB resident** against 426.72 GiB before this row — the 82 IQ2_XS tensors stop expanding from 53.33 GiB to 369.00 GiB — which fits the ~119.63 GiB of `dgx:gpu0` with 18.49 GiB of headroom. [Routing table](../tests/vllm/test_gguf_keep_quant.cpp#L425) restated: the GEMM term moved 20 -> 24 and the GATHER term stayed 13. **The CUDA arm now HAS a keep-quant kernel for this dtype** ([#2260](https://github.com/mudler/vllm.cpp/issues/2260), spec [cuda-keepquant-iq2xs-iq4xs](specs/cuda-keepquant-iq2xs-iq4xs.md)): `src/vt/cuda/cuda_quant_dot.cu::DotIQ2XS` is a 1:1 port of `b10451 quants.c:948` reading the generated `src/vt/cuda/cuda_quant_iq_tables.cuh::d_iq2xs_grid`, and `src/vt/cuda/cuda_quant_dot.cu::IsCudaKeepQuantSupported` admits twelve Q8_K-family encodings with IQ2_XS among them, wired into all THREE dispatch switches. Before that the 82 tensors still kept — `src/vllm/model_executor/model_loader/gguf_keep_quant.cpp::DeviceKeepQuantSupported` returns `true` for CUDA on its `default:` arm regardless — so the artifact DID fit and the expert GEMM then took the CPU fallback inside `src/vt/cuda/cuda_quant_dot.cu::MatmulBTQuantGroupedKernelCuda` behind a full `cudaStreamSynchronize` per grouped call, while the fused `src/vt/cuda/cuda_quant_dot.cu::MoeGateUpSwiGLUGroupedCuda` seam THREW `gate/up must be the SAME CUDA keep-quant dtype`. **101.14 GiB remains a RESIDENCY result and no speed number is claimed**, and the kernel is not yet SELECTED by a running model: `Glm5NextHostForward` refuses a non-CPU queue by name, so the wiring is still owed by `MODEL-MM-glm5-next-glm5-next-for-conditional-generation`. NOT reached from this row today — `glm5_next_moe.cpp` is W5's host reference and `laguna.cpp` is the only model on the fused seam — but AGENTS.md routes mergeable MLP projections through `vt::MergedGemmGroup`, so W5b ([#2241](https://github.com/mudler/vllm.cpp/issues/2241)) and W5c ([#2242](https://github.com/mudler/vllm.cpp/issues/2242)) make the throw live on a 101 GiB-resident model. Contrast `QUANT-GGUF-IQ2_XXS` above, which records a CUDA compute landing; this row has none. Carried as O19 in [glm5-next-flash](specs/glm5-next-flash.md). | [glm5-next-flash](specs/glm5-next-flash.md) | - | -| `QUANT-GGUF-IQ3_XXS` | IQ3_XXS | 18 / output | Y | Y | Y | - | - | `READY` | reader trait [`:237`](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L237) (98 B block, ADDED W8); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L349) (`DequantIQ3_XXS` + the shared [`iq3xxs_grid`](../src/vt/cpu/cpu_quant_iq_tables.h), ported 1:1 from llama.cpp `ggml-quants.c:2503` + `ggml-common.h:1007`); vt block dtype `kIQ3_XXS` [geometry](../src/vt/dtype.cpp#L86). **`C` = `Y` (DeepSeek-V4 W8, 2026-07-29):** the keep-quant `vec_dot` [`VecDotIQ3_XXSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L1) (1:1 port of ggml `vec_dot_iq3_xxs_q8_K_generic`, `quants.c:999`) + the [Q8_K traits row](../src/vt/cpu/cpu_quant_traits.cpp#L1) 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](../tests/vt/test_ops_quant_dot.cpp#L1) (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](../src/vt/cuda/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](../tests/vt/test_cuda_quant_dot.cpp)) 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. **The ROCm arm now HAS a keep-quant kernel for this dtype too** ([#1940](https://github.com/mudler/vllm.cpp/issues/1940), spec [kernel-quant-ciq-gemm-rocm-iquant](specs/kernel-quant-ciq-gemm-rocm-iquant.md)): [`DotIQ3XXS`](../src/vt/rocm/rocm_grouped_gemm.hip#L345) is ADAPTED from the CUDA body rather than freshly ported, on both the plain and the grouped/MoE arm, and the `kROCM` case of [`DeviceKeepQuantSupported`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L138) now admits IQ3_XXS beside the four k-quants it already carried, so a ROCm keep-quant load stops expanding these blocks to bf16. `C` does NOT move and stays `Y`: it already read `Y` for the CPU and CUDA arms, and it says only that native quantized compute is dispatched, without naming a backend. **The device gate is CONTRIBUTOR evidence on `isravale`** (RX 9060 XT, `gfx1200`, ROCm 7.2.3) under the GPU file mutex: `test_backend_cross_device` **41/42** — the single failure is `MoeSiluMul`, the pre-existing tracked [#1954](https://github.com/mudler/vllm.cpp/issues/1954) flake, confirmed unrelated to this row — `test_gguf_keep_quant` **52/52**, `test_gguf_device_fit` and `test_gguf_device_fit_reach` green, and zero regression on the four existing k-quant formats. The separate, LATER operator run on `strix:gpu0` covers the device-table seal and NOT this landing; it and the owed performance work are recorded once, in the `QUANT-GGUF-IQ4_XS` row below. | [CUDA keep-quant GEMM](specs/cuda-keepquant-gemm.md); [iquant DSV4 spike](specs/gguf-iquant-dsv4.md) | - | +| `QUANT-GGUF-IQ3_XXS` | IQ3_XXS | 18 / output | Y | Y | Y | - | - | `READY` | reader trait [`:237`](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L237) (98 B block, ADDED W8); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L349) (`DequantIQ3_XXS` + the shared [`iq3xxs_grid`](../src/vt/cpu/cpu_quant_iq_tables.h), ported 1:1 from llama.cpp `ggml-quants.c:2503` + `ggml-common.h:1007`); vt block dtype `kIQ3_XXS` [geometry](../src/vt/dtype.cpp#L86). **`C` = `Y` (DeepSeek-V4 W8, 2026-07-29):** the keep-quant `vec_dot` [`VecDotIQ3_XXSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L1) (1:1 port of ggml `vec_dot_iq3_xxs_q8_K_generic`, `quants.c:999`) + the [Q8_K traits row](../src/vt/cpu/cpu_quant_traits.cpp#L1) 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](../tests/vt/test_ops_quant_dot.cpp#L1) (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](../src/vt/cuda/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](../tests/vt/test_cuda_quant_dot.cpp)) 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. **The ROCm arm now HAS a keep-quant kernel for this dtype too** ([#1940](https://github.com/mudler/vllm.cpp/issues/1940), spec [kernel-quant-ciq-gemm-rocm-iquant](specs/kernel-quant-ciq-gemm-rocm-iquant.md)): [`DotIQ3XXS`](../src/vt/rocm/rocm_grouped_gemm.hip#L345) is ADAPTED from the CUDA body rather than freshly ported, on both the plain and the grouped/MoE arm, and the `kROCM` case of [`DeviceKeepQuantSupported`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L138) now admits IQ3_XXS beside the four k-quants it already carried, so a ROCm keep-quant load stops expanding these blocks to bf16. `C` does NOT move and stays `Y`: it already read `Y` for the CPU and CUDA arms, and it says only that native quantized compute is dispatched, without naming a backend. **The device gate is CONTRIBUTOR evidence on `isravale`** (RX 9060 XT, `gfx1200`, ROCm 7.2.3) under the GPU file mutex: `test_backend_cross_device` **41/42** — the single failure is `MoeSiluMul`, the pre-existing tracked [#1954](https://github.com/mudler/vllm.cpp/issues/1954) flake, confirmed unrelated to this row — `test_gguf_keep_quant` **52/52**, `test_gguf_device_fit` and `test_gguf_device_fit_reach` green, and zero regression on the four existing k-quant formats. **The operator independently reproduced this gate on `isravale` at head `fa39a45a3`** (2026-09-08, same box, same GPU file mutex; full detail in the `QUANT-GGUF-IQ4_XS` row below): all four targets green, zero regression, higher case counts because the branch since gained the table-seal work (#3067) and the IQ3_XXS loader-admission case. The separate, LATER operator run on `strix:gpu0` covers the device-table seal and NOT this landing; it and the owed performance work are recorded once, in the `QUANT-GGUF-IQ4_XS` row below. | [CUDA keep-quant GEMM](specs/cuda-keepquant-gemm.md); [iquant DSV4 spike](specs/gguf-iquant-dsv4.md) | - | | `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 | Y | Y | - | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 21](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L305) (110 B block, ADDED by [#2510](https://github.com/mudler/vllm.cpp/issues/2510)); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L399) `DequantIQ3_S` + the 512-entry [`kIq3sGrid`](../src/vt/cpu/cpu_quant_iq_tables.h#L168), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2607` + `ggml/src/ggml-common.h:1052`; vt block dtype `kIQ3_S` [geometry](../src/vt/dtype.cpp#L150) and the CUDA gather codec [`DqIQ3_S`](../src/vt/cuda/cuda_quant_dequant.cuh#L476). **THE REFUSAL THIS ROW EXISTS FOR:** id 21 was the ONE hole in the reader's i-quant run 16..23, and `GgufFile::Open` therefore refused whole artifacts over it — `unsloth/Qwen3.8-27B-GGUF` `UD-Q4_K_M.gguf` stores 4 of its 866 tensors in IQ3_S and `UD-Q4_K_XL.gguf` stores 1 of 866, while every OTHER encoding in those files (F32, Q3_K, Q4_K, Q5_K, Q6_K, Q8_0, IQ4_NL, IQ4_XS) was already handled. That is the file the published Strix Halo rows ran, so [#2497](https://github.com/mudler/vllm.cpp/issues/2497) had to substitute the plain `Q4_K_M`. **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.11.ffn_gate.weight` at absolute offset 4,262,628,128 in the staged `Qwen3.8-27B-UD-Q4_K_M.gguf` (sha256 `322e194f…3482`), decoded by the pinned upstream's own `dequantize_row_iq3_s` and compared as raw f32 bit patterns ([goldens + provenance](../tests/vt/iq3s_golden_vectors.h#L69), [unit](../tests/vllm/test_gguf_dequant.cpp#L663)); the four blocks span 190 distinct 9-bit grid indices of which 73 are >= 256, so the `qh`-spliced ninth bit is exercised rather than left at zero. The 110-byte stride is the ORACLE's own `sizeof(block_iq3_s)`, printed by the harness, not a sum read off the struct. The table itself is sealed by FNV-1a digest and by its lane alphabet ([test_ops_quant_dot](../tests/vt/test_ops_quant_dot.cpp#L717)), because `kIq3sGrid` (512 u32) and `kIq3xxsGrid` (256 u32) are the same shape and a decoder pointed at the sibling still decodes. Reached through `GgufFile::Open` + `DequantGgufRowToF32` in the same file, which is the production path the real artifact took. **`C` = `-`, DELIBERATELY, and it is the one FILE encoding in this tree that decodes without dotting.** `KeepQuantDType` gates on `HasQuantDotKernel`, so IQ3_S routes `kExpandBf16` on the GEMM arm of EVERY tier (CPU, CUDA, ROCm, Metal, Vulkan) and `kKeepQuant` only on the gather, where CPU and CUDA both decode. That is asserted BY NAME rather than left to a count ([routing](../tests/vllm/test_gguf_keep_quant.cpp#L429), [decode-only class](../tests/vt/test_ops_quant_traits.cpp#L250)), so the day the `vec_dot` lands those cases red. **The CPU `VecDotIQ3_SQ8_K` and the CUDA `WType::kIQ3_S` are OWED AS ONE UNIT** (`ggml-cpu/quants.c:1094`, `.vec_dot_type = GGML_TYPE_Q8_K` at `ggml-cpu.c:355-360`): `IsCudaKeepQuantSupported` has no refusal arm, so landing the CPU half alone would flip the loader to keep IQ3_S blocks and send every CUDA IQ3_S GEMM through the host fallback that [#2260](https://github.com/mudler/vllm.cpp/issues/2260) MEASURED to segfault on a discrete card. On the measured artifact the expansion is 146.13 MiB of blocks against 680.00 MiB of bf16, 3.4 % of a 15.33 GiB file — bounded here, and NOT bounded on an artifact whose experts are IQ3_S, which is why it is recorded rather than waved at (compare [#1870](https://github.com/mudler/vllm.cpp/issues/1870)). `E`/`P` = `-`: this row makes the artifact OPENABLE and claims no token and no number; the `gfx1151` quant-matched decode number stays owed by [#2497](https://github.com/mudler/vllm.cpp/issues/2497). | [gguf-iq3s](specs/gguf-iq3s.md) | - | | `QUANT-GGUF-IQ2_S` | IQ2_S | 22 / IQ2_M storage | Y | - | - | - | - | `INVENTORIED` | reader `:229-234`; [explicit rejection](../tests/vllm/test_gguf_dequant.cpp#L223) | leaf open | - | -| `QUANT-GGUF-IQ4_XS` | IQ4_XS | 23 / output | Y | Y | Y | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 23](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L311) (136 B block; the trait predates the decoder, the decoder is [#2240](https://github.com/mudler/vllm.cpp/issues/2240)); [dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L488) `DequantIQ4_XS`, ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2743`; vt block dtype `kIQ4_XS` [geometry](../src/vt/dtype.cpp#L140). NOT a codebook delta from IQ4_NL — it reuses `kValuesIq4nl` unchanged and differs only in the SUPER-BLOCK SCALE LAYOUT (a 6-bit `ls` spliced from a `scales_l` nibble and a `scales_h` bit pair, then biased by -32). **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.11.ffn_down_exps.weight` from the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL arm, whose four super-blocks span `ls` 0..61 so both nibbles and both `scales_h` bit pairs and both signs of `dl` are exercised ([goldens + provenance](../tests/vt/iq2xs_iq4xs_golden_vectors.h#L276), [unit](../tests/vllm/test_gguf_dequant.cpp#L576)). [Geometry cross-check](../tests/vt/test_ops_quant_traits.cpp#L156) pins the reader and vt tables against each other. The former `explicit rejection` evidence is GONE by construction — that test now guards Q1_0 (41), the remaining tabulated-but-undecodable id. **`C` = `Y` since [#2247](https://github.com/mudler/vllm.cpp/issues/2247):** keep-quant [`VecDotIQ4_XSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L844), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-cpu/quants.c:1283` `ggml_vec_dot_iq4_xs_q8_K_generic`, with [`BlockIQ4_XS`](../src/vt/cpu/cpu_quant_blocks.h#L205) (`ggml-common.h:454-460`) and the [traits row](../src/vt/cpu/cpu_quant_traits.cpp#L121). **The activation pairing is `Q8_K`, READ OFF the oracle and not inherited from IQ4_NL:** `type_traits_cpu` at `ggml/src/ggml-cpu/ggml-cpu.c:385-390` carries `.vec_dot = ggml_vec_dot_iq4_xs_q8_K, `.vec_dot_type = GGML_TYPE_Q8_K`, against `:379-384`'s `GGML_TYPE_Q8_0` for IQ4_NL — the 16-entry codebook is shared, the block geometry is not, and a 256-element super-block pairs with the 256-element activation encoding. **Gated BIT FOR BIT against the oracle's OWN KERNEL** on 4 super-blocks of `blk.11.ffn_down_exps.weight`, total and per block ([goldens](../tests/vt/iq2xs_iq4xs_dot_golden.h#L269), [case](../tests/vt/test_ops_quant_dot.cpp#L869)). These 3 tensors stop expanding from 3.59 GiB to 13.50 GiB; the whole-artifact measurement is in the `QUANT-GGUF-IQ2_XS` row above. **It also unblocks a SECOND consumer:** the expert-tower streaming lane is all-or-nothing (`GgufExpertTowersReachSlotLane` returns false on the FIRST tower that does not keep), so the 4 IQ4_XS tensors of the GLM-5.3 `UD-IQ1_S` arm were dropping an arm whose other 221 towers all kept — one tower 6.375 -> 24.000 GiB, and a 4096-slot cache 25.5 -> 96 GiB. Asserted in the `kStackedExpertWeight` role both models store these towers in ([lane case](../tests/vllm/model_executor/test_gguf_device_fit.cpp#L672)). **The CUDA arm now HAS a keep-quant kernel for this dtype too** ([#2260](https://github.com/mudler/vllm.cpp/issues/2260), spec [cuda-keepquant-iq2xs-iq4xs](specs/cuda-keepquant-iq2xs-iq4xs.md)): `src/vt/cuda/cuda_quant_dot.cu::DotIQ4XS` is a 1:1 port of `b10451 quants.c:1283` over the generated `src/vt/cuda/cuda_quant_iq_tables.cuh::d_kvalues_iq4nl`, keeping upstream's EIGHT f32 accumulation steps per super-block rather than one integer core, because that association is what the oracle's golden numbers were produced with. **This completes the device admission of the GLM-5.3 non-flash `UD-IQ1_S` arm**, whose other five encodings (IQ1_S, IQ3_XXS, IQ2_XXS, Q2_K, Q3_K) were already there. Before it, these 3 tensors fit `dgx:gpu0` with their expert GEMM on the CPU fallback behind a `cudaStreamSynchronize` and the fused `src/vt/cuda/cuda_quant_dot.cu::MoeGateUpSwiGLUGroupedCuda` seam throwing. The fit remains RESIDENCY and no speed number is claimed; O19's premise in [glm5-next-flash](specs/glm5-next-flash.md) is discharged, and what still blocks a CUDA run of either model is that model's own host-f32 forward. **The ROCm arm now HAS a keep-quant kernel for this dtype too** ([#1940](https://github.com/mudler/vllm.cpp/issues/1940), spec [kernel-quant-ciq-gemm-rocm-iquant](specs/kernel-quant-ciq-gemm-rocm-iquant.md)): [`DotIQ4XS`](../src/vt/rocm/rocm_grouped_gemm.hip#L395) is ADAPTED from the CUDA body rather than freshly ported, on both the plain and the grouped/MoE arm, and the `kROCM` case of [`DeviceKeepQuantSupported`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L138) now admits IQ4_XS and IQ3_XXS beside the four k-quants it already carried. **CUDA's `DotIQ4XS` needed non-fused `__fmul_rn`/`__fadd_rn` to stay bit-exact and the HIP body does NOT, which is MEASURED and not assumed** ([rationale in place](../src/vt/rocm/rocm_grouped_gemm.hip#L384)): nvcc never sees the project-wide `-ffp-contract=off` that CMake already passes to `$`, and a bit-exact gate over the same real-checkpoint bytes and the same oracle bits the CUDA gate uses confirms plain arithmetic suffices on this toolchain, so the port carries no workaround CUDA needed. `C` does NOT move and stays `Y`: it already read `Y` for the CPU and CUDA arms, and it says only that native quantized compute is dispatched, without naming a backend. **The device gate is CONTRIBUTOR evidence on `isravale`** (RX 9060 XT, `gfx1200`, ROCm 7.2.3) under the GPU file mutex: `test_backend_cross_device` **41/42** — the single failure is `MoeSiluMul`, the pre-existing tracked [#1954](https://github.com/mudler/vllm.cpp/issues/1954) flake, confirmed unrelated to this row — `test_gguf_keep_quant` **52/52**, `test_gguf_device_fit` and `test_gguf_device_fit_reach` green, and zero regression on the four existing k-quant formats. **A SECOND, LATER operator run is a DIFFERENT scope and must not be read as covering this landing:** it verified the device-table seal ([#3067](https://github.com/mudler/vllm.cpp/issues/3067)) on `strix:gpu0`, `gfx1151`, ROCm 7.2.4 under lease `886172b3-9d21-4435-a9db-cbf7b2f5dfb6`, where four ctest targets passed at 61 focused and 9 codebook-comparison assertions and all eleven independent review mutations failed their intended tests. **NO end-to-end real-model reload is gated by either run.** `Nail-Qwen3.6-35B-A3B-MTP-IQ4_XS.gguf` (19.39 GB) loaded and emitted coherent text on the ORIGINAL, pre-repair head only; that checkpoint has NO recorded repository revision and NO sha256, its 19,333,564,672 B weight footprint is a CALCULATION and not a measured peak RSS, and coherent text is not a token-ID oracle comparison. **Speed is OWED and named rather than dropped:** the `ROCM-KQUANT-NWARPS-DECODE` re-measurement for these two formats is `PENDING` — no `rocprofv3` profiling setup is staged on that box — and does not block a row whose own scope is coverage and correctness; a WMMA/tensor-core tile for the two is the follow-on row `KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT-RDNA4`. | [glm5-next-flash](specs/glm5-next-flash.md) | - | +| `QUANT-GGUF-IQ4_XS` | IQ4_XS | 23 / output | Y | Y | Y | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 23](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L311) (136 B block; the trait predates the decoder, the decoder is [#2240](https://github.com/mudler/vllm.cpp/issues/2240)); [dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L488) `DequantIQ4_XS`, ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2743`; vt block dtype `kIQ4_XS` [geometry](../src/vt/dtype.cpp#L140). NOT a codebook delta from IQ4_NL — it reuses `kValuesIq4nl` unchanged and differs only in the SUPER-BLOCK SCALE LAYOUT (a 6-bit `ls` spliced from a `scales_l` nibble and a `scales_h` bit pair, then biased by -32). **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.11.ffn_down_exps.weight` from the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL arm, whose four super-blocks span `ls` 0..61 so both nibbles and both `scales_h` bit pairs and both signs of `dl` are exercised ([goldens + provenance](../tests/vt/iq2xs_iq4xs_golden_vectors.h#L276), [unit](../tests/vllm/test_gguf_dequant.cpp#L576)). [Geometry cross-check](../tests/vt/test_ops_quant_traits.cpp#L156) pins the reader and vt tables against each other. The former `explicit rejection` evidence is GONE by construction — that test now guards Q1_0 (41), the remaining tabulated-but-undecodable id. **`C` = `Y` since [#2247](https://github.com/mudler/vllm.cpp/issues/2247):** keep-quant [`VecDotIQ4_XSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L844), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-cpu/quants.c:1283` `ggml_vec_dot_iq4_xs_q8_K_generic`, with [`BlockIQ4_XS`](../src/vt/cpu/cpu_quant_blocks.h#L205) (`ggml-common.h:454-460`) and the [traits row](../src/vt/cpu/cpu_quant_traits.cpp#L121). **The activation pairing is `Q8_K`, READ OFF the oracle and not inherited from IQ4_NL:** `type_traits_cpu` at `ggml/src/ggml-cpu/ggml-cpu.c:385-390` carries `.vec_dot = ggml_vec_dot_iq4_xs_q8_K, `.vec_dot_type = GGML_TYPE_Q8_K`, against `:379-384`'s `GGML_TYPE_Q8_0` for IQ4_NL — the 16-entry codebook is shared, the block geometry is not, and a 256-element super-block pairs with the 256-element activation encoding. **Gated BIT FOR BIT against the oracle's OWN KERNEL** on 4 super-blocks of `blk.11.ffn_down_exps.weight`, total and per block ([goldens](../tests/vt/iq2xs_iq4xs_dot_golden.h#L269), [case](../tests/vt/test_ops_quant_dot.cpp#L869)). These 3 tensors stop expanding from 3.59 GiB to 13.50 GiB; the whole-artifact measurement is in the `QUANT-GGUF-IQ2_XS` row above. **It also unblocks a SECOND consumer:** the expert-tower streaming lane is all-or-nothing (`GgufExpertTowersReachSlotLane` returns false on the FIRST tower that does not keep), so the 4 IQ4_XS tensors of the GLM-5.3 `UD-IQ1_S` arm were dropping an arm whose other 221 towers all kept — one tower 6.375 -> 24.000 GiB, and a 4096-slot cache 25.5 -> 96 GiB. Asserted in the `kStackedExpertWeight` role both models store these towers in ([lane case](../tests/vllm/model_executor/test_gguf_device_fit.cpp#L672)). **The CUDA arm now HAS a keep-quant kernel for this dtype too** ([#2260](https://github.com/mudler/vllm.cpp/issues/2260), spec [cuda-keepquant-iq2xs-iq4xs](specs/cuda-keepquant-iq2xs-iq4xs.md)): `src/vt/cuda/cuda_quant_dot.cu::DotIQ4XS` is a 1:1 port of `b10451 quants.c:1283` over the generated `src/vt/cuda/cuda_quant_iq_tables.cuh::d_kvalues_iq4nl`, keeping upstream's EIGHT f32 accumulation steps per super-block rather than one integer core, because that association is what the oracle's golden numbers were produced with. **This completes the device admission of the GLM-5.3 non-flash `UD-IQ1_S` arm**, whose other five encodings (IQ1_S, IQ3_XXS, IQ2_XXS, Q2_K, Q3_K) were already there. Before it, these 3 tensors fit `dgx:gpu0` with their expert GEMM on the CPU fallback behind a `cudaStreamSynchronize` and the fused `src/vt/cuda/cuda_quant_dot.cu::MoeGateUpSwiGLUGroupedCuda` seam throwing. The fit remains RESIDENCY and no speed number is claimed; O19's premise in [glm5-next-flash](specs/glm5-next-flash.md) is discharged, and what still blocks a CUDA run of either model is that model's own host-f32 forward. **The ROCm arm now HAS a keep-quant kernel for this dtype too** ([#1940](https://github.com/mudler/vllm.cpp/issues/1940), spec [kernel-quant-ciq-gemm-rocm-iquant](specs/kernel-quant-ciq-gemm-rocm-iquant.md)): [`DotIQ4XS`](../src/vt/rocm/rocm_grouped_gemm.hip#L395) is ADAPTED from the CUDA body rather than freshly ported, on both the plain and the grouped/MoE arm, and the `kROCM` case of [`DeviceKeepQuantSupported`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L138) now admits IQ4_XS and IQ3_XXS beside the four k-quants it already carried. **CUDA's `DotIQ4XS` needed non-fused `__fmul_rn`/`__fadd_rn` to stay bit-exact and the HIP body does NOT, which is MEASURED and not assumed** ([rationale in place](../src/vt/rocm/rocm_grouped_gemm.hip#L384)): nvcc never sees the project-wide `-ffp-contract=off` that CMake already passes to `$`, and a bit-exact gate over the same real-checkpoint bytes and the same oracle bits the CUDA gate uses confirms plain arithmetic suffices on this toolchain, so the port carries no workaround CUDA needed. `C` does NOT move and stays `Y`: it already read `Y` for the CPU and CUDA arms, and it says only that native quantized compute is dispatched, without naming a backend. **The device gate is CONTRIBUTOR evidence on `isravale`** (RX 9060 XT, `gfx1200`, ROCm 7.2.3) under the GPU file mutex: `test_backend_cross_device` **41/42** — the single failure is `MoeSiluMul`, the pre-existing tracked [#1954](https://github.com/mudler/vllm.cpp/issues/1954) flake, confirmed unrelated to this row — `test_gguf_keep_quant` **52/52**, `test_gguf_device_fit` and `test_gguf_device_fit_reach` green, and zero regression on the four existing k-quant formats. **The operator independently reproduced this exact gate on `isravale` at head `fa39a45a3`** (2026-09-08, same box, same GPU file mutex, `llama-server.service` confirmed `inactive` before and after): `test_backend_cross_device` **48/48**, 84,104/84,104 assertions — `MoeSiluMul` additionally re-run **5/5** in isolation, confirming #1954's flake rather than fixing it — `test_gguf_keep_quant` **54/54**, 10,340/10,340 assertions; `test_gguf_device_fit` **24/24**, 182/182; `test_gguf_device_fit_reach` **21/21**, 100/100; zero failures anywhere. Case counts are HIGHER than the contributor's run because the branch since gained the table-seal work (#3067) and the IQ3_XXS loader-admission case; this is the SAME landing re-verified, not a different one. **A SECOND, LATER operator run is a DIFFERENT scope and must not be read as covering this landing:** it verified the device-table seal ([#3067](https://github.com/mudler/vllm.cpp/issues/3067)) on `strix:gpu0`, `gfx1151`, ROCm 7.2.4 under lease `886172b3-9d21-4435-a9db-cbf7b2f5dfb6`, where four ctest targets passed at 61 focused and 9 codebook-comparison assertions and all eleven independent review mutations failed their intended tests. **NO end-to-end real-model reload is gated by either run.** `Nail-Qwen3.6-35B-A3B-MTP-IQ4_XS.gguf` (19.39 GB) loaded and emitted coherent text on the ORIGINAL, pre-repair head only; that checkpoint has NO recorded repository revision and NO sha256, its 19,333,564,672 B weight footprint is a CALCULATION and not a measured peak RSS, and coherent text is not a token-ID oracle comparison. **Speed is OWED and named rather than dropped:** the `ROCM-KQUANT-NWARPS-DECODE` re-measurement for these two formats is `PENDING` — no `rocprofv3` profiling setup is staged on that box — and does not block a row whose own scope is coverage and correctness; a WMMA/tensor-core tile for the two is the follow-on row `KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT-RDNA4`. | [glm5-next-flash](specs/glm5-next-flash.md) | - | | `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 | - |