From 5f137895b02f5e08fe61e14256331d20a4236d03 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 12 Sep 2026 00:12:45 +0000 Subject: [PATCH 1/9] spec(QUANT-GGUF-IQ4_NL): native IQ4_NL compute on ROCm and CUDA IQ4_NL reaches the GGUF reader, the CPU dequantizer and the CPU keep-quant dot, and no accelerator multiplies it. On ROCm the refusal is terminal, because the backend is discrete and an unsupported dtype cannot fall back to a host kernel that would follow device pointers. This blocks a shipped artifact rather than a hypothetical one. Every published unsloth/Qwen3.8-Flash-Next-GGUF quant stores 91,465,564,160 elements in IQ4_NL, read from the three shard headers over HTTP range requests: the 20M-entry n-gram table and the 48 ffn_down_exps. #3097 has since made the table readable on ROCm. The experts are still not multipliable, and that is now the whole gap for this encoding on gfx1151. The CUDA half is a variant and not a table entry. IQ4_NL is a 32-element block paired with a Q8_0 activation, where the resident device GEMM quantizes activations to Q8_K over 256-element super-blocks. cuda_quant_dot.cu already says so: DotMXFP4 sits there marked maybe_unused, awaiting the same Q8_0-activation GEMM variant that Q5_0 and Q4_0 also queue behind. The row understated the tree and now says what it does. The reader arm, the dequantizer and the CPU dot all landed under #1989 while the cell read INVENTORIED with every stage unset. Adds strix:gpu0 to the fleet table, which predates the box, with its memory measured under two rc leases. hipMallocManaged there is bounded by HOST memory rather than by the firmware carve: a bounded probe that stops at its first failure reached 76 GiB with plain hipMalloc and 27 GiB with hipMallocManaged, against 29.3 GiB host-available. That resolves an ambiguity #2518 could not, because its 58.000 GiB ceiling sat below both the 64.00 GiB carve and the 62 GiB host RAM of the time and so never said which bound it hit. Raising the carve to 96 GiB LOWERED the managed ceiling, because host RAM fell with it. The spec records two corrections against itself rather than applying them silently, because a scope that shrinks without a trace cannot be audited. It first planned a ROCm gather arm and first recorded #3029 as a hard dependency. #3097 landed both the gather and the ROCm codebook header while this was being drafted, so the arm is struck and the dependency is downgraded to a conflict surface. No product code lands here, and no throughput, latency or memory number is claimed. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md | 5 + .agents/environment.md | 48 +++ .../ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md | 62 +++ .agents/quantization-matrix.md | 2 +- .agents/specs/quant-gguf-iq4nl.md | 364 ++++++++++++++++++ 5 files changed, 480 insertions(+), 1 deletion(-) create mode 100644 .agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md create mode 100644 .agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md create mode 100644 .agents/specs/quant-gguf-iq4nl.md diff --git a/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md b/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md new file mode 100644 index 000000000..d63f40c33 --- /dev/null +++ b/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md @@ -0,0 +1,5 @@ +# CLAIM-QUANT-GGUF-IQ4_NL + +| Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | +|---|---|---|---|---|---|---|---| +| `CLAIM-QUANT-GGUF-IQ4_NL` | `QUANT-GGUF-IQ4_NL` (`ACTIVE`) | Claude Code (opus-5) as helper; implementation, review and repair to run as separate fresh agents per `AGENTS.md` "How work gets done" | local worktree `/home/mudler/.cache/sdd/mudler-vllm.cpp/quant-gguf-iq4nl`; ROCm arms gate on `strix:gpu0` (`gfx1151`), a fleet device reached only through an `rc` lease, never `ssh` | `row/QUANT-GGUF-IQ4_NL`, based `18f39771c`; issue `ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD` | Owns: the IQ4_NL keep-quant dot and its device admission on ROCm (`src/vt/rocm/rocm_quant_dot.hip`, `src/vt/rocm/rocm_grouped_gemm.hip`); the CUDA Q8_0-activation GEMM variant plus `DotIQ4_NL` and its `IsCudaKeepQuantSupported` arm (`src/vt/cuda/cuda_quant_dot.cu`); this row's spec, its `.agents/quantization-matrix.md` cell, and the `strix:gpu0` memory rows in `.agents/environment.md`. EXCLUDES the ROCm quantized GATHER, which #3097 landed on 2026-09-12 and which this row therefore does not reimplement; every other missing ROCm format (IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S, MXFP4 — #1940), any tensor-core tile for IQ4_NL, the existing Q8_K-activation dispatch and its kernels (untouched), the seven `qwen4_exp` operations with no ROCm arm (owed, separate row), and every throughput, latency or memory number (gate-blocked) | `ACTIVE` | 2026-09-12 — spec committed, no product code, base `18f39771c`. **Scope shrank during drafting and the record says so:** [#3097](https://github.com/mudler/vllm.cpp/pull/3097) landed the ROCm quantized gather (arm 2) and the ROCm IQ codebook header carrying `d_kvalues_iq4nl`, which together removed the #3029 dependency this spec first recorded. #3029 remains a conflict surface only, and its two `sanitize-cpu` reds are a repository-wide pre-existing failure in `dots3` tests it does not touch. Next: W1, the ROCm `DotIQ4_NL`, which needs no other pull request to land first | diff --git a/.agents/environment.md b/.agents/environment.md index 950d81db0..6d7d70804 100644 --- a/.agents/environment.md +++ b/.agents/environment.md @@ -68,6 +68,54 @@ The fleet, read from `rc devices` and `rc describe` on 2026-08-17: | `dgx:gpu0` | `gpu_model=GB10`, `class=train`, `k8s=true`, driver 580.173.02, `cpus=20`, 128 GB | the house NAS | | `thor:gpu0` | `gpu_model=NVIDIA-Thor`, `class=train`, `k8s=true`, driver 595.78, `cpus=14`, 132 GB | the house NAS, the SAME folder as `dgx` | | `orin:gpu0` | `gpu_model=AGX-Orin`, `class=train`, `k8s=true`, `cpus=12`, 32 GB, L4T R36.4.7 (JetPack 6), and NO detected GPU labels because Jetson carries no `nvidia-smi` | the house NAS, the SAME folder as `dgx` and `thor` | +| `strix:gpu0` | `gpu_model=Radeon-8060S`, `vendor=amd`, `class=train`, `k8s=true`, `cpus=32`, `gfx1151` (RDNA 3.5, AMD RYZEN AI MAX+ 395), ROCm 7.2.4 / HIP 7.2.53211, `mem_total_bytes=33270497280` | the house NAS, the SAME folder as `dgx`, `thor` and `orin` | + +### `strix:gpu0` was missing from the table above until 2026-09-11 + +**It is the only AMD device on this fleet, and the table predates it.** The +table's own header says it was read on 2026-08-17; `strix:gpu0` is leasable, +every `BACKEND-ROCM` row leases it, and a reader who trusted the table would +conclude this fleet has no AMD hardware. Claim it with `rc run` or `rc hold` +like any other fleet device, and never by `ssh`. + +**Its memory split is configurable in firmware and it was CHANGED on +2026-09-11**, so every figure recorded against this box before that date +describes a different machine. Measured inside a lease on 2026-09-11, `rc` +jobs `a8111ff8-3ce8-42f6-9034-36bdd2cacfe4` and +`c30dc437-bf12-4a23-ab8b-89b88fe767dd`: + +| Probe | Value | +|---|---| +| `mem_info_vram_total` | 103,079,215,104 B = **96.00 GiB** | +| `hipMemGetInfo` total / free | 96.000 GiB / 95.848 GiB idle | +| `mem_info_gtt_total` | 16,635,248,640 B = 15.49 GiB | +| host RAM total / available | 33,270,497,280 B / 29,304,037,376 B | +| device properties | `integrated=1 managedMemory=1 pageableMemoryAccess=0 gcn=gfx1151` | + +**`hipMallocManaged` on this board is bounded by HOST memory, not by the +carve.** A bounded probe that stops at its first failure reached **76 GiB with +plain `hipMalloc`** (its target, so the real ceiling is at least that and is +NOT measured above it) and **27 GiB with `hipMallocManaged`, which returned +`out of memory`**. 27 GiB against 29.3 GiB host-available is the match that +identifies which bound was hit. + +This **resolves an ambiguity [#2518](https://github.com/mudler/vllm.cpp/issues/2518) +could not**. Its 58.000 GiB managed ceiling was measured when the carve was +64.00 GiB and host RAM was 62 GiB; 58 sits below both, so that number never +said which one it was. The current split separates them, and the answer is +host. + +Three consequences for anyone sizing work here: + +1. **Raising the carve LOWERED the managed ceiling**, 58 GiB to 27 GiB, because + host RAM fell from 62 GiB to 31 GiB. More VRAM is not more of everything. +2. **`VT_ROCM_MANAGED_ALLOC=1` is actively harmful on this board.** The default + is already correct: `ResolveMemoryPolicy` (`include/vt/rocm/rocm_arch.h`) + sets `managed_alloc = pageable_memory_access` under `kUnset`, and this board + reports that 0, so the #2511 narrowing selects plain `hipMalloc`. +3. **A large CPU arm no longer fits this box.** The `qwen4_exp` CPU arm peaked + at 73.9 GiB `VmHWM`; the host side is now 31 GiB total. Run a CPU comparison + on `thor` or `dgx`, or on-box against an oracle instead. ### `orin:gpu0` needs L4T CUDA 12.6, and the DGX recipe breaks it diff --git a/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md b/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md new file mode 100644 index 000000000..ec9fbe3d2 --- /dev/null +++ b/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md @@ -0,0 +1,62 @@ +ID: ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD +Title: IQ4_NL has no native quantized compute on ROCm or CUDA +Row: QUANT-GGUF-IQ4_NL +State: OPEN +Kind: bug +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-11 +Updated: 2026-09-12 +Closed: - + +## Problem + +IQ4_NL (ggml id 20) reaches the GGUF reader, the CPU dequantizer and the CPU keep-quant dot, but has no native quantized compute on any device and no ROCm gather arm. + +Measured on the tree at b7fb4e51f: + +- ROCm: src/vt/rocm/rocm_quant_dot.hip declares ten WType entries (IQ2_XXS, IQ3_XXS, Q2_K, Q3_K, Q4_K, Q5_K, Q6_K, IQ2_S, IQ1_S, IQ1_XXXS). IQ4_NL is absent. src/vt/rocm/rocm_grouped_gemm.hip:1783 and :1855 throw naming it. The backend is discrete, so an unsupported dtype cannot fall back to the CPU kernel and the throw is terminal. +- CUDA: IsCudaKeepQuantSupported has no IQ4_NL arm. The blocker is structural rather than per-format: IQ4_NL is a 32-element block paired with Q8_0 activation, not a 256-element Q8_K super-block, and the CUDA file has no Q8_0-activation GEMM variant. DotMXFP4 already sits in cuda_quant_dot.cu marked [[maybe_unused]] with a comment stating it awaits exactly that variant. Q5_0, Q4_0 and MXFP4 are queued behind the same gap. +- Gather: OpId::kEmbeddingQuant is registered for kCPU (cpu_ops.cpp:4282) and kCUDA (cuda_ops.cu:4135) only. DeviceQuantGatherSupported is therefore false for kROCM, and qwen4_exp_weights.cpp:665 refuses the load by name before any tensor I/O. + +Why this blocks a shipped artifact. Every published unsloth/Qwen3.8-Flash-Next-GGUF quant stores 91,465,564,160 elements in IQ4_NL, read from the three shard headers over HTTP range requests: the 20M-entry n-gram embedding table and the 48 ffn_down_exps. The table needs the gather and the experts need the dot, so both halves are required and neither is optional for this checkpoint. UD-IQ1_S needs IQ4_NL alone; UD-IQ1_M additionally needs IQ1_M; UD-Q2_K_XL additionally needs IQ2_XS. + +Record defect in the same area. .agents/quantization-matrix.md carries QUANT-GGUF-IQ4_NL as INVENTORIED with R, M, C, E and P all unset, but the reader arm (gguf_reader.cpp case 20), the dequantizer (DequantIQ4_NL) and the CPU keep-quant dot (VecDotIQ4_NLQ8_0, cpu_quant_dot.cpp:140) all landed under W6a (#1989). The row understates what the tree does. + +## Resolution + +- + +## Reconciliation 2026-09-12: the gather half landed while this issue was open + +**The problem statement above is kept verbatim because it was accurate when +written, and one of its three bullets is now false.** + +[#3097](https://github.com/mudler/vllm.cpp/pull/3097) +(`feat(BACKEND-ROCM-QUANT-GATHER): gather packed embeddings on ROCm`, +`82de418e8`) landed on `main` and: + +- registers `OpId::kEmbeddingQuant` for `kROCM` at `src/vt/rocm/rocm_ops.hip:191`; +- carries IQ4_NL in the ROCm gather codec list, + `X(kIQ4_NL, DqIQ4_NL)` at `src/vt/rocm/rocm_embedding_quant.hip:88`; +- adds `src/vt/rocm/rocm_quant_iq_tables.h`, which carries + `d_kvalues_iq4nl[16]` at `:1198` from stock `ggml-common.h:1120`. + +`DeviceQuantGatherSupported` is exactly +`vt::OpRegistered(vt::OpId::kEmbeddingQuant, dev)` +(`gguf_keep_quant.cpp:220`), so it is **now true for ROCm** and the +`qwen4_exp_weights.cpp:665` refusal no longer fires on that device. + +**What is still open, and is what this issue now means:** the ROCm IQ4_NL +keep-quant DOT, and the CUDA IQ4_NL keep-quant GEMM. `WType` in +`rocm_quant_dot.hip` still declares ten entries without IQ4_NL, and +`rocm_grouped_gemm.hip:1783` still throws naming it. The gather can read the +n-gram table and the expert GEMM still cannot multiply `ffn_down_exps`. + +This also removed the hard dependency on +[#3029](https://github.com/mudler/vllm.cpp/pull/3029) that the spec first +recorded: the ROCm codebook header IQ4_NL needs is now on `main` by another +route. #3029 remains a conflict surface over eight files, and its two +`sanitize-cpu` reds are a repository-wide pre-existing failure in `dots3` +tests it does not touch. diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index 939763ca1..d3ca5609e 100644 --- a/.agents/quantization-matrix.md +++ b/.agents/quantization-matrix.md @@ -72,7 +72,7 @@ forces the full dequant path back. | `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#L172) 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-IQ1_S` | IQ1_S | 19 / output | - | - | - | - | - | `INVENTORIED` | - | leaf open | - | -| `QUANT-GGUF-IQ4_NL` | IQ4_NL | 20 / output | - | - | - | - | - | `INVENTORIED` | - | leaf open | - | +| `QUANT-GGUF-IQ4_NL` | IQ4_NL | 20 / output | Y | Y | Y | - | - | `ACTIVE` | reader trait [`FindGgmlTraits` case 20](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L292) (18 B block, `QK4_NL = 32`: Q4_0's geometry with the 16-entry NON-LINEAR codebook `kvalues_iq4nl` of llama.cpp `b10451` `ggml/src/ggml-common.h:1120` in place of the affine step), ADDED by [#1989](https://github.com/mudler/vllm.cpp/issues/1989) and gated against the pinned oracle by [`DequantGgufRowToF32 IQ4_NL row matches the pinned oracle`](../tests/vllm/test_gguf_dequant.cpp#L528), with the keep-quant admission asserted BY NAME at [`test_gguf_keep_quant.cpp`](../tests/vllm/test_gguf_keep_quant.cpp#L598) and its device-routing case at [`:655`](../tests/vllm/test_gguf_keep_quant.cpp#L655), together with the dequantizer [`DequantIQ4_NL`](../src/vt/cpu/cpu_quant_dequant.cpp#L92) and a dequantizing gather that makes `GgufTensorRole::kEmbeddingTable` keep-quant eligible. **`C` = `Y` ON CPU ONLY, and the activation pairing is `Q8_0` rather than the `Q8_K` every neighbouring i-quant uses** ([`VecDotIQ4_NLQ8_0`](../src/vt/cpu/cpu_quant_dot.cpp#L140), ported 1:1 from `b10451` `ggml/src/ggml-cpu/quants.c:1254` `ggml_vec_dot_iq4_nl_q8_0_generic`, traits row [`cpu_quant_traits.cpp`](../src/vt/cpu/cpu_quant_traits.cpp#L52) pairing `kIQ4_NL` -> `kQ8_0`, read off `ggml-cpu.c:379-384` and NOT inherited from IQ4_XS, which pairs `Q8_K`). That 32-element block on a Q8_0 activation is exactly why NO accelerator serves it: the resident device GEMM on both ROCm and CUDA quantizes activations to `BlockQ8_K` over 256-element super-blocks. **CUDA has no arm** (`IsCudaKeepQuantSupported` omits it; [`DotMXFP4`](../src/vt/cuda/cuda_quant_dot.cu#L698) already sits there `[[maybe_unused]]` saying in its own comment that it awaits the Q8_0-activation GEMM variant, which Q5_0, Q4_0 and MXFP4 also queue behind). **ROCm has no arm and its refusal is TERMINAL** (ten `WType` entries at [`rocm_quant_dot.hip`](../src/vt/rocm/rocm_quant_dot.hip#L580) exclude it; [`rocm_grouped_gemm.hip:1783`](../src/vt/rocm/rocm_grouped_gemm.hip#L1783) throws naming it, and the backend is discrete so an unsupported dtype cannot fall back to a host kernel that would follow device pointers). **The ROCm GATHER, by contrast, LANDED on 2026-09-12** via [#3097](https://github.com/mudler/vllm.cpp/pull/3097) (`82de418e8`), which registers `OpId::kEmbeddingQuant` for `kROCM` at [`rocm_ops.hip:191`](../src/vt/rocm/rocm_ops.hip#L191) and carries IQ4_NL in its codec list at [`rocm_embedding_quant.hip:88`](../src/vt/rocm/rocm_embedding_quant.hip#L88). `DeviceQuantGatherSupported` is exactly `vt::OpRegistered(kEmbeddingQuant, dev)` ([`gguf_keep_quant.cpp:220`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L220)), so it is now TRUE for ROCm and the [`qwen4_exp_weights.cpp:665`](../src/vllm/model_executor/models/qwen4_exp_weights.cpp#L665) refusal no longer fires on that device. **The gather reads the table and the dot still cannot multiply the experts**, which is the whole remaining gap for this encoding on gfx1151. **Both halves are load-bearing for a SHIPPED artifact and neither is optional:** every published `unsloth/Qwen3.8-Flash-Next-GGUF` quant stores **91,465,564,160 elements** in IQ4_NL, read from the three shard headers over HTTP range requests rather than from prose - the 20M-entry n-gram table, which is gathered and never multiplied, and the 48 `ffn_down_exps`, which are multiplied and never gathered. `UD-IQ1_S` (67.56 GiB) needs IQ4_NL ALONE; `UD-IQ1_M` additionally needs IQ1_M; `UD-Q2_K_XL` (73.45 GiB) additionally needs IQ2_XS, and those two formats are about 97% of its weight elements. **`E`/`P` = `-`:** this row claims no token and no number, and none is admissible while the `gfx1151` token gate reads `TOKEN_GATE=FAIL` at 3 of 6 ([`rocm-gfx1151-q4k-token-gate-v2.md`](specs/rocm-gfx1151-q4k-token-gate-v2.md)). **This cell previously read `INVENTORIED` with R, M, C, E and P all unset**, which understated the tree by three stages: the reader, the dequantizer and the CPU dot had all landed under #1989. | [quant-gguf-iq4nl](specs/quant-gguf-iq4nl.md) | `CLAIM-QUANT-GGUF-IQ4_NL` | | `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#L266)), 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#L172) 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) | - | diff --git a/.agents/specs/quant-gguf-iq4nl.md b/.agents/specs/quant-gguf-iq4nl.md new file mode 100644 index 000000000..d187daf4f --- /dev/null +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -0,0 +1,364 @@ +# Spec: QUANT-GGUF-IQ4_NL — native IQ4_NL compute on ROCm and CUDA + +- Issue: `ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD` +- Row: `QUANT-GGUF-IQ4_NL` (`.agents/quantization-matrix.md`). Owning feature + rows: `BACKEND-ROCM` ([#41](https://github.com/mudler/vllm.cpp/issues/41)) + and `QUANT-CUDA-GATES`. +- Claim: `CLAIM-QUANT-GGUF-IQ4_NL` +- Base: `18f39771c` (`origin/main`, 2026-09-12) +- Pull request shape: **one pull request** carrying the spec and the + implementation, spec committed first. Developer decision, 2026-09-11. + +## Contract + +| Field | Value | +|---|---| +| Scope | IN: the IQ4_NL keep-quant dot and device admission on ROCm; the CUDA Q8_0-activation GEMM variant plus `DotIQ4_NL` and its `IsCudaKeepQuantSupported` arm; the `QUANT-GGUF-IQ4_NL` record repair and the `strix:gpu0` rows in `.agents/environment.md`. OUT: the ROCm quantized GATHER, which #3097 landed on 2026-09-12; every other missing ROCm format (IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S, MXFP4 -- #1940); any tensor-core tile for IQ4_NL; the seven `qwen4_exp` operations with no ROCm arm; every throughput, latency and memory number. | +| Upstream chain | Numerical authority is the pinned [`llama-cpp`](../oracles/llama-cpp.md) oracle at `10bf611e5` (`b10451`), because vLLM defines no GGUF k-quant dot kernel and therefore has nothing to mirror here: `ggml/src/ggml-cpu/quants.c:1254` `ggml_vec_dot_iq4_nl_q8_0_generic` (the dot), `ggml/src/ggml-common.h:447-452` `block_iq4_nl` (18 B, `QK4_NL = 32`), `ggml/src/ggml-common.h:1120` `kvalues_iq4nl` (the 16-entry codebook), `ggml/src/ggml-cpu/ggml-cpu.c:379-384` (`.vec_dot_type = GGML_TYPE_Q8_0`, the activation pairing, read off upstream and NOT inherited from IQ4_XS which pairs `Q8_K`). | +| Our baseline | Reader, dequantizer and CPU dot all landed under #1989 and are the transcription source: `gguf_reader.cpp` case 20; `src/vt/cpu/cpu_quant_dequant.cpp:92` `DequantIQ4_NL`; `src/vt/cpu/cpu_quant_dot.cpp:140` `VecDotIQ4_NLQ8_0`; `src/vt/cpu/cpu_quant_blocks.h:62` `BlockIQ4_NL`; traits `src/vt/cpu/cpu_quant_traits.cpp:52` pairing `kIQ4_NL` -> `kQ8_0`. Absent everywhere else: ten `WType` entries in `rocm_quant_dot.hip:580` exclude it, `rocm_grouped_gemm.hip:1783` and `:1855` throw naming it, `IsCudaKeepQuantSupported` omits it, and `kEmbeddingQuant` is registered for `kCPU` and `kCUDA` only. | +| Port map | `cpu_quant_dot.cpp:140` `VecDotIQ4_NLQ8_0` -> a `DotIQ4_NL` device body in `src/vt/rocm/rocm_quant_dot.hip` (scalar, transcribed from the CPU body and NOT from CUDA, per `rocm_quant_dot.hip:226`: gfx1100 has no `__dp4a` so the CUDA integer-core shapes do not port) and a matching `DotIQ4_NL` in `src/vt/cuda/cuda_quant_dot.cu`. `kvalues_iq4nl` -> the ROCm IQ codebook header #3029 introduces, sealed. `cpu_ops.cpp:4282` `kEmbeddingQuant` registration -> a `kROCM` registration decoding IQ4_NL block rows, mirroring the CUDA codec `cuda_quant_dequant.cuh:156` `DqIQ4_NL`. The ROCm Q8_0 activation quantizer already exists (`rocm_grouped_gemm.hip` `QuantizeQ8_0Kernel`); CUDA needs one. | +| Tests to port | vLLM has no test to port here, because it has no GGUF k-quant dot. Goldens come from the pinned oracle's OWN kernel over real `UD-IQ1_S` checkpoint bytes, in the shape `tests/vllm/test_gguf_dequant.cpp:528` already uses for the IQ4_NL dequantizer and `tests/vt/test_ops_quant_dot.cpp:869` uses for the IQ4_XS dot: per block and in total, never a synthetic tensor. Newly authored: the IQ4_NL cases in `tests/vt/test_rocm_quant_dot.cpp`, `tests/vt/test_ops_quant_traits.cpp`, `tests/vllm/test_gguf_keep_quant.cpp`, `tests/vt/test_backend_cross_device.cpp`, and a device gather case. | +| Gates | See `## Gates`. G1 unit suites on CPU, on `strix:gpu0` for the ROCm arms and on a CUDA box for arm 3; G2 admission of `UD-IQ1_S` with zero reference-tier hits for the IQ4_NL GEMM and gather. NO throughput, latency or memory number is admissible from this row. | +| Dependencies | **None blocking.** #3029 was recorded as a hard dependency and is not one: #3097 landed `src/vt/rocm/rocm_quant_iq_tables.h` on `main` carrying `d_kvalues_iq4nl[16]` at `:1198`. #3029 remains a CONFLICT surface (eight files, measured by `git merge-tree`), and its two `sanitize-cpu` reds are a repository-wide pre-existing failure in `dots3` tests it does not touch. Hardware: `strix:gpu0` (`gfx1151`), a fleet device reachable only through an `rc` lease. No CI lane has an AMD runner, so a green CI is not evidence for arms 1 and 2. | +| Work breakdown | `W0` this spec -> `W1` ROCm `DotIQ4_NL` + `WType` entry + dispatch + device admission -> `W2` CUDA Q8_0-activation variant + `DotIQ4_NL` + `IsCudaKeepQuantSupported` arm -> `W3` record repair. The gather wave this spec originally planned is struck: #3097 landed it. W1 alone is what stands between `UD-IQ1_S` and a fully device-resident weight set on gfx1151. | +| Risks/decisions | R1 a codebook decoder pointed at a sibling table still decodes (IQ4_NL shares its codebook with IQ4_XS, so this is live); R2 reassociation passes a small golden; R3 the CUDA Q8_0-activation variant is a new dispatch path beside the Q8_K one it can regress; R4 `main` moves and #3029 also edits `.agents/quantization-matrix.md`; R5 no AMD runner in CI. D1 keep upstream's association order; D2 IQ4_NL is a 32-element block on a Q8_0 activation, which is why arm 3 is a variant rather than a table entry; D3 the gather and the dot are both required and are not the same code. Detail in `## Design` and `## Risks`. | + +## Why this row exists, in one paragraph + +IQ4_NL is the one encoding that every published `Qwen3.8-Flash-Next` GGUF +stores its two largest structures in, and it is the only encoding in those +files that no device can compute on. The reader opens it, the CPU dequantizes +it and the CPU dots it. Every accelerator refuses it. On ROCm the refusal is +terminal because the backend is discrete and cannot fall back to a host kernel +that would follow device pointers. + +## Scope + +Three arms and one record repair, in one change: + +1. **ROCm keep-quant dot.** An `IQ4_NL` arm for the ROCm quantized GEMM, with + device admission, so `MatmulBTQuant` and the grouped expert GEMM stop + throwing on it. +2. ~~**ROCm quantized gather.**~~ **LANDED ELSEWHERE ON 2026-09-12, while this + spec was being written, and struck rather than deleted so the change of + scope is auditable.** + [#3097](https://github.com/mudler/vllm.cpp/pull/3097) + (`feat(BACKEND-ROCM-QUANT-GATHER)`, `82de418e8`) registers + `OpId::kEmbeddingQuant` for `kROCM` at `src/vt/rocm/rocm_ops.hip:191` and + lists `X(kIQ4_NL, DqIQ4_NL)` in `src/vt/rocm/rocm_embedding_quant.hip:88`. + `DeviceQuantGatherSupported` is exactly + `vt::OpRegistered(vt::OpId::kEmbeddingQuant, dev)` + (`gguf_keep_quant.cpp:220`), so it is now true for ROCm and the + `qwen4_exp_weights.cpp:665` refusal no longer fires on this device. This row + does not reimplement it. +3. **CUDA keep-quant dot.** The Q8_0-activation GEMM variant the CUDA file is + already written against, plus `DotIQ4_NL` and its entry in + `IsCudaKeepQuantSupported`. +4. **Record reconciliation** of the `QUANT-GGUF-IQ4_NL` row, which understates + the tree: `R`, `M` and CPU `C` all landed under #1989 while the row still + read `INVENTORIED` with every stage unset. It moves to `ACTIVE`. + +### Out of scope, named rather than dropped + +- **IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S and MXFP4 on ROCm.** Tracked by + [#1940](https://github.com/mudler/vllm.cpp/issues/1940). `DotMXFP4` already + exists on CUDA and only waits on arm 3; wiring it is still a separate row + because it changes what `IsCudaKeepQuantSupported` admits. +- **Any tensor-core tile for IQ4_NL.** The scalar tier lands first, exactly as + `KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT` decided for IQ4_XS. A 16-entry codebook + format does not map onto the linear-scale k-quant WMMA tile. +- **The seven `qwen4_exp` operations that have no ROCm arm.** See `## Owed`. + This row makes the model's weights computable on ROCm. It does not make the + model run, and no part of this spec should be read as claiming that. +- **Any throughput, latency or memory number.** See `## Gates`. + +## Dependency on #3029: DOWNGRADED on 2026-09-12 + +**This spec first recorded #3029 as a hard dependency. It is not one any more, +and the correction is kept rather than quietly applied.** The premise was that +`src/vt/rocm/rocm_quant_iq_tables.h` existed only on that branch, so IQ4_NL's +codebook had nowhere to live. That was true when written and false hours later: +#3097 landed the header on `main`, and it carries `d_kvalues_iq4nl[16]` at +`:1198`, transcribed from stock `ggml-common.h:1120`. The codebook this row +needs is therefore already on `main`. + +What remains is a **conflict surface, not a dependency**. #3029 edits +`src/vt/rocm/rocm_grouped_gemm.hip` and `.agents/quantization-matrix.md`, which +this row also edits. Measured with `git merge-tree` against `origin/main` +`18f39771c` on 2026-09-12, #3029 carries **eight conflicts**: +`.agents/quantization-matrix.md`, +`.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md`, `docs/USAGE.md`, +`src/vllm/model_executor/model_loader/gguf_keep_quant.cpp`, +`src/vt/rocm/rocm_grouped_gemm.hip`, `src/vt/rocm/rocm_quant_iq_tables.h` +(add/add, because #3097 landed the same path), and two test files. + +**Its two `sanitize-cpu` reds are NOT its own**, and that was measured rather +than assumed: the failing cases are `test_dots3_note_vision`, +`test_dots3_note_audio` and `test_openai_api_server_dots3_mm_forward`, #3029 +touches no `dots3` file, and the same two lanes are red on +[#3140](https://github.com/mudler/vllm.cpp/pull/3140), +[#3133](https://github.com/mudler/vllm.cpp/pull/3133) and +[#3036](https://github.com/mudler/vllm.cpp/pull/3036) as well. They are a +repository-wide pre-existing red and they do not belong to this row either. +Its `windows-msvc-*` reds are the known +[#584](https://github.com/mudler/vllm.cpp/issues/584) crash baseline. + +## Upstream anchors + +Primary oracle is vLLM wherever it defines behavior. It does not define a GGUF +k-quant dot kernel, so the numerical authority here is the pinned llama.cpp +oracle [`llama-cpp`](../oracles/llama-cpp.md) at `10bf611e5` (`b10451`), exactly +as every sibling format in this tree used. + +| What | Where | +|---|---| +| `ggml_vec_dot_iq4_nl_q8_0_generic` | `b10451` `ggml/src/ggml-cpu/quants.c:1254` | +| `block_iq4_nl` layout | `b10451` `ggml/src/ggml-common.h:447-452` | +| `kvalues_iq4nl` codebook | `b10451` `ggml/src/ggml-common.h:1120` | +| activation pairing | `b10451` `ggml/src/ggml-cpu/ggml-cpu.c:379-384`, `.vec_dot_type = GGML_TYPE_Q8_0` | + +Our existing ports, which are the transcription source for the device arms: + +| What | Where | +|---|---| +| CPU dot | `src/vt/cpu/cpu_quant_dot.cpp:140` `VecDotIQ4_NLQ8_0` | +| CPU block | `src/vt/cpu/cpu_quant_blocks.h:62` `BlockIQ4_NL`, 18 bytes | +| CPU dequant | `src/vt/cpu/cpu_quant_dequant.cpp:92` `DequantIQ4_NL` | +| reader trait | `src/vllm/model_executor/model_loader/gguf_reader.cpp` case 20 | +| CUDA gather codec | `src/vt/cuda/cuda_quant_dequant.cuh:156` `DqIQ4_NL` | + +**Transcribe from the CPU bodies, not from CUDA.** That is what +`rocm_quant_dot.hip:226` records for the existing formats, and its reason still +holds: gfx1100 has no `__dp4a`, so the CUDA integer-core shapes do not port. + +## Design + +### D1. The association order is load-bearing + +`VecDotIQ4_NLQ8_0` forms the scale product **before** folding in the integer +sum, `d * (sumi1 + sumi2)`, which is the opposite association from the adjacent +`q4_0` kernel. `cpu_quant_dot.cpp:136-139` already says so in a comment, and +says why: it is what makes the GEMM bit-reproducible against upstream. Both +device arms keep that order. A kernel that reassociates is not this kernel, and +the goldens below will not detect the difference on small inputs, which is +precisely why this is stated as a design decision rather than left to review. + +### D2. IQ4_NL is a 32-element block on a Q8_0 activation + +This is the structural fact that makes arm 3 a variant rather than a table +entry. The resident device path on both backends quantizes activations to +`BlockQ8_K` over 256-element super-blocks. IQ4_NL pairs with `BlockQ8_0` over +32. On ROCm the activation quantizer already exists: +`rocm_grouped_gemm.hip` carries `QuantizeQ8_0Kernel` beside `QuantizeQ8KKernel`. +On CUDA it does not, and `cuda_quant_dot.cu:695-698` records the gap in its own +words, marking `DotMXFP4` `[[maybe_unused]]` because it "awaits the +Q8_0-activation GEMM variant above". Arm 3 is that variant. Q5_0, Q4_0 and +MXFP4 all queue behind it, which is this row's leverage and also the reason its +CUDA half is larger than one kernel. + +### D3. The gather and the dot are both required, and they are not the same code + +`unsloth/Qwen3.8-Flash-Next-GGUF` stores **91,465,564,160 elements** in IQ4_NL +in every published quant. Two structures account for it: the 20M-entry n-gram +embedding table, which is read by `kEmbeddingQuant` and never multiplied, and +the 48 `ffn_down_exps`, which are multiplied and never gathered. Landing only +the dot leaves `DeviceQuantGatherSupported(kROCM)` false and the loader still +refuses at `qwen4_exp_weights.cpp:665`, before any tensor I/O. Landing only the +gather leaves the expert GEMM throwing at `rocm_grouped_gemm.hip:1783`. This is +why the row's scope is both. + +### D4. Reachability + +Per `AGENTS.md` "Nothing lands dead", each arm enters through a production +entry point at this commit, and the smallest failing test enters the same way: + +| Arm | Production entry | Mutation that must red the gate | +|---|---|---| +| ROCm dot | `vt::MatmulBTQuant` / `kMatmulBTQuantGrouped` dispatch | delete the `IQ4_NL` dispatch case; the focused gate must fail rather than fall back | +| CUDA dot | `IsCudaKeepQuantSupported` plus the Q8_0-activation GEMM | delete the `IsCudaKeepQuantSupported` arm; the keep-quant routing case must red | + +A unit test that constructs the kernel by hand proves the arithmetic and not +the capability. Each arm therefore also carries a routing assertion by name, in +the shape `tests/vllm/test_gguf_keep_quant.cpp` already uses for IQ3_S and +IQ4_XS. + +## Evidence measured for this spec + +All read on 2026-09-11, recorded here because three of them contradict what the +tree currently says. + +### E1. What the shipped artifacts actually store + +Read from the three GGUF shard headers of each quant over HTTP range requests, +not from prose. `general.architecture = qwen4exp`, `split.tensors.count = 1224`. + +| Quant | Bytes | Formats with no ROCm arm | +|---|---|---| +| `UD-IQ1_S` | 72,546,461,344 (67.56 GiB) | **IQ4_NL only** | +| `UD-IQ1_M` | 74,538,755,776 (69.42 GiB) | IQ4_NL, IQ1_M | +| `UD-Q2_K_XL` | 78,869,128,864 (73.45 GiB) | IQ4_NL, IQ2_XS | + +`UD-Q2_K_XL` type histogram: IQ4_NL 49 tensors / 91,465,564,160 elements; +IQ2_XS 94 / 78,852,915,200; Q5_K 189 / 2,823,946,240; IQ3_XXS 2 / +1,677,721,600; Q8_0 244 / 747,110,400; Q6_K 64 / 642,908,160; Q4_K 1 / +635,699,200; F32 and BF16 581 / 98,034,560. The two formats ROCm lacks are +about 97 percent of its weight elements. + +`UD-IQ1_S` is therefore the cheapest ROCm target of the family, on two +independent counts: one missing format instead of two, and the smallest file. + +### E2. `strix:gpu0` carve, measured under an `rc` lease + +Job `a8111ff8-3ce8-42f6-9034-36bdd2cacfe4`, `rc run -d strix:gpu0`. + +| Probe | Value | +|---|---| +| `mem_info_vram_total` | 103,079,215,104 B = **96.00 GiB** | +| `mem_info_vram_used` | 154,816,512 B, idle | +| `mem_info_gtt_total` | 16,635,248,640 B = 15.49 GiB | +| host RAM total / available | 33,270,497,280 / 29,304,037,376 B | +| device | `gfx1151`, AMD RYZEN AI MAX+ 395 w/ Radeon 8060S | + +### E3. `hipMallocManaged` on gfx1151 is bounded by HOST memory, not the carve + +Job `c30dc437-bf12-4a23-ab8b-89b88fe767dd`, same box, HIP 7.2.53211, a bounded +probe that stops at the first failure and never walks past its target. + +```text +hipMemGetInfo: free=95.848 GiB total=96.000 GiB +integrated=1 managedMemory=1 pageableMemoryAccess=0 gcn=gfx1151 +hipMalloc, target 76 GiB: => 76 GiB (target reached, ceiling is >= 76) +hipMallocManaged, target 76 GiB: => 27 GiB (out of memory) +``` + +27 GiB against 29.3 GiB host-available is the match that identifies the source. +This **resolves the ambiguity [#2518](https://github.com/mudler/vllm.cpp/issues/2518) +could not**: its 58.000 GiB managed ceiling was measured when the carve was +64.00 GiB and host RAM was 62 GiB, and 58 sits below both, so that measurement +alone never said which bound it hit. The current split separates them. + +Three consequences, each of which outlives this row: + +1. **The shipped artifact fits, on the plain path.** 67.56 GiB and 73.45 GiB + both sit inside the proven 76 GiB of `hipMalloc`. Neither fits the 27 GiB + managed ceiling. +2. **`VT_ROCM_MANAGED_ALLOC=1` is now actively harmful on this board**, and + more so than before the carve change: it moves the ceiling from >= 76 GiB to + 27 GiB. The default is already correct without it. `ResolveMemoryPolicy` + (`include/vt/rocm/rocm_arch.h:164`) sets `managed_alloc = + pageable_memory_access` under `kUnset`, and this board reports that 0, so + the #2511 narrowing already selects plain `hipMalloc`. +3. **`.agents/environment.md` does not carry this box at all.** Its fleet + table is dated 2026-08-17 and lists `dgx:gpu0`, `thor:gpu0` and + `orin:gpu0` only, while `strix:gpu0` has been leasable for weeks and every + ROCm row leases it. The absence is worse than a stale number, because a + reader sizing a model for this board finds nothing and falls back to the + 64 GiB and 62 GiB figures scattered through + `q4km-limb3-kquant-vehicle.md` and `rocm-expert-lane-guard.md`, both of + which the carve change has now falsified. Repaired in this change by adding + the box with the values measured above. + +The `hipMalloc` ceiling above 76 GiB is **not measured**. The probe stopped at +its target by design. Do not quote 96 GiB as an allocatable figure. + +## Risks + +- **R1. A codebook decoder pointed at a sibling table still decodes.** This is + the exact failure `rocm_iq_table_seal.h` exists for, and IQ4_NL shares its + 16-entry codebook with IQ4_XS, so it is live here rather than theoretical. + Mitigation: seal `kValuesIq4nl` by digest and by its lane alphabet, in the + shape `tests/vt/test_ops_quant_dot.cpp:717` already uses. +- **R2. Reassociation passes a small golden.** See D1. Mitigation: goldens + taken from the oracle's own kernel over real checkpoint bytes, per block and + in total, not a synthetic tensor. +- **R3. The CUDA Q8_0-activation variant is a new dispatch path, not an + entry.** It can regress the Q8_K path it sits beside. Mitigation: the Q8_K + formats' existing cases must stay byte-identical; assert that rather than + assume it. +- **R4. `main` moves under this row.** #3029 touches + `.agents/quantization-matrix.md`, which this row also edits. Per + `AGENTS.md` "Records", take the complete target-branch version and re-apply + the scoped edit; never accept a three-way merge of a keyed record. +- **R5. No AMD runner in CI.** Every ROCm device test self-skips, so a green CI + is not evidence for arms 1 and 2. Their evidence is a named lease. + +## Tests + +Red first, in this order, each failing for the intended reason before any +kernel exists: + +1. `tests/vt/test_rocm_quant_dot.cpp` — IQ4_NL dot against oracle-produced + goldens over real `UD-IQ1_S` `ffn_down_exps` bytes, per block and total. +2. `tests/vt/test_ops_quant_traits.cpp` — reader and vt geometry agree for + IQ4_NL, 18-byte block, `QK4_NL = 32`, against the oracle's own + `sizeof(block_iq4_nl)` printed by the harness. +3. `tests/vllm/test_gguf_keep_quant.cpp` — routing by name: IQ4_NL routes + `kKeepQuant` on ROCm and CUDA GEMM arms, and on the gather. +4. `tests/vt/test_backend_cross_device.cpp` — IQ4_NL GEMM within NMSE <= 5e-4 + of the CPU oracle, on a real launch shape rather than a toy one. The Q6_K + precedent is explicit about why: its gate ran `M=3, N=8, K=512` against a + production launch of 6400 blocks and never tested a launch. +5. A gather case that decodes an IQ4_NL block row on device and matches + `DequantIQ4_NL`. +6. The three D4 mutations, each restoring the tree byte for byte afterwards. + +## Gates + +- **G1 (unit).** The suites above, green on CPU, on `strix:gpu0` for the ROCm + arms, and on a CUDA box for arm 3. +- **G2 (admission).** `UD-IQ1_S` opens and its IQ4_NL tensors keep their blocks + on ROCm, with `VT_OP_PROVIDER_STATS` showing zero reference-tier hits for the + IQ4_NL GEMM and gather. +- **G3 (correctness).** Not owned by this row. See `## Owed`. +- **No throughput, latency or memory number is admissible from this row.** + `AGENTS.md` Gates admits no performance result from an arm whose token gate + has not passed, and the `gfx1151` token gate currently reads + `TOKEN_GATE=FAIL` at 3 of 6 on a simpler model + (`rocm-gfx1151-q4k-token-gate-v2.md`). Numbers captured while this stands are + recorded as inadmissible, with that word, or they repeat the withdrawn 2.71x. + +## Owed + +- **The seven `qwen4_exp` operations with no ROCm arm**, which are what stands + between this row and a model that runs: `kQwen4ExpPleConv`, + `kQwen4ExpPleGate`, `kQwen4ExpGatedResidual`, + `kQwen4ExpGatedResidualWriteBack`, `kQwen4ExpQsaCompress`, + `kQwen4ExpQsaGatherAttention`, `kRmsNormGroup`. `grep -rn Qwen4Exp + src/vt/rocm/` returns nothing. Needs its own row and spec. +- **The ROCm token gate for `qwen4_exp`**, and the oracle it runs against. + vLLM at the current pin `e126687a9a` ships a first-class AMD backend for this + model (`vllm/models/qwen4_exp/amd/`), whose divergence from `nvidia/` is + about 840 lines confined to the Triton op layer, with `model.py`, + `model_state.py` and `mtp.py` byte-identical and no AITER or MFMA dependency. + That makes it the natural denominator on gfx1151, and it is unmeasured. +- **The CPU comparison arm no longer fits this box.** It needed 73.9 GiB + `VmHWM` and the host side is now 31 GiB. A CPU-versus-ROCm token comparison + has to run on `thor` or `dgx`, or on-box against an oracle instead. +- IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S and MXFP4 on ROCm (#1940). Wiring the + already-written CUDA `DotMXFP4` onto arm 3's variant. +- The `hipMalloc` ceiling above 76 GiB on this board, unmeasured by design. + +## Stop conditions + +- #3029 cannot be landed or its ROCm IQ table headers change shape. Stop and + re-scope arms 1 and 2; arm 3 and the record repair remain independent. +- The oracle goldens cannot be produced from real checkpoint bytes. Stop rather + than substitute a synthetic tensor, per R2. +- A device arm disagrees with the CPU oracle beyond NMSE 5e-4 and the cause is + not reassociation. Stop and report; do not widen the bound. +- `strix:gpu0` is unavailable or unhealthy. Arms 1 and 2 stay `PENDING` on a + named lease. Never convert an unrun gate into a pass. + +## Now + +`ACTIVE`, 2026-09-12. Spec committed, no product code, base pinned to +`origin/main` `18f39771c`. + +**The scope shrank between drafting and committing, and the record says so +rather than pretending it was always this size.** #3097 landed the ROCm +quantized gather, which was arm 2, and it also landed the ROCm IQ codebook +header that made #3029 look like a hard dependency. What remains is the ROCm +IQ4_NL dot, the CUDA IQ4_NL dot, and the record repair. + +Next action is W1, the ROCm dot, which needs no other pull request to land +first. From b4fa49117003d479d66eb107d391914acf4ac84e Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 12 Sep 2026 00:39:00 +0000 Subject: [PATCH 2/9] feat(QUANT-GGUF-IQ4_NL): IQ4_NL keep-quant GEMM on ROCm The ROCm backend refused IQ4_NL terminally: ten WType entries excluded it and rocm_grouped_gemm.hip threw naming it, with no host fallback available on a discrete backend. Every published Qwen3.8-Flash-Next quant stores its 48 ffn_down_exps in IQ4_NL, so those experts could not be multiplied at all. Two kernels, not one, because an expert tower does not reach the single-matrix provider. GroupedIQ4NLK is the arm the shipped checkpoints need; IQ4NLGemmK covers the non-grouped path for completeness. Admitting the encoding with only the latter would have thrown at the first expert forward with the model already resident, which is the failure DeviceKeepQuantSupported exists to prevent, so the predicate and both kernels move together. DotIQ4_NL is a 1:1 port of cpu_quant_dot.cpp VecDotIQ4_NLQ8_0, itself llama.cpp b10451 quants.c:1254. It keeps upstream's association, d * (sumi1 + sumi2), where the scale product is formed before the integer sum is folded in. That is the opposite order from DotQ8_0 beside it, and it is what makes the GEMM bit-reproducible against the oracle. The activation encoding is Q8_0, not the Q8_K every neighbouring i-quant here uses. ROCm already had that path (QuantizeQ8_0K, DotQ8_0, both GEMM shapes), so this is a new dot in existing machinery rather than new machinery. The spec assumed otherwise and is corrected in the same change; CUDA, which genuinely lacks the path, is unchanged and still owed. Deliberately not Dp4a-vectorized: the codebook lookup is per nibble, so a four-way byte dot has nothing to multiply until the values are gathered, and gfx1100 has no hardware dp4a in any case. Both cross-device cases gain IQ4_NL against the CPU oracle at NMSE <= 5e-4. They REQUIRE rather than skip when ROCm is registered, so a dropped dispatch cannot pass silently. NOT YET COMPILED OR RUN. This host has no ROCm toolchain and check-tree-compiles does not reach a .hip file, so the 721 translation units it checked exclude every line of the kernel work here. The build and the device gate are owed on a strix:gpu0 lease and are recorded as PENDING until then. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/specs/quant-gguf-iq4nl.md | 22 ++- .../model_loader/gguf_keep_quant.cpp | 18 ++- src/vt/rocm/rocm_grouped_gemm.hip | 144 +++++++++++++++++- tests/vt/test_backend_cross_device.cpp | 20 ++- 4 files changed, 185 insertions(+), 19 deletions(-) diff --git a/.agents/specs/quant-gguf-iq4nl.md b/.agents/specs/quant-gguf-iq4nl.md index d187daf4f..a17ba3399 100644 --- a/.agents/specs/quant-gguf-iq4nl.md +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -16,7 +16,7 @@ | Scope | IN: the IQ4_NL keep-quant dot and device admission on ROCm; the CUDA Q8_0-activation GEMM variant plus `DotIQ4_NL` and its `IsCudaKeepQuantSupported` arm; the `QUANT-GGUF-IQ4_NL` record repair and the `strix:gpu0` rows in `.agents/environment.md`. OUT: the ROCm quantized GATHER, which #3097 landed on 2026-09-12; every other missing ROCm format (IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S, MXFP4 -- #1940); any tensor-core tile for IQ4_NL; the seven `qwen4_exp` operations with no ROCm arm; every throughput, latency and memory number. | | Upstream chain | Numerical authority is the pinned [`llama-cpp`](../oracles/llama-cpp.md) oracle at `10bf611e5` (`b10451`), because vLLM defines no GGUF k-quant dot kernel and therefore has nothing to mirror here: `ggml/src/ggml-cpu/quants.c:1254` `ggml_vec_dot_iq4_nl_q8_0_generic` (the dot), `ggml/src/ggml-common.h:447-452` `block_iq4_nl` (18 B, `QK4_NL = 32`), `ggml/src/ggml-common.h:1120` `kvalues_iq4nl` (the 16-entry codebook), `ggml/src/ggml-cpu/ggml-cpu.c:379-384` (`.vec_dot_type = GGML_TYPE_Q8_0`, the activation pairing, read off upstream and NOT inherited from IQ4_XS which pairs `Q8_K`). | | Our baseline | Reader, dequantizer and CPU dot all landed under #1989 and are the transcription source: `gguf_reader.cpp` case 20; `src/vt/cpu/cpu_quant_dequant.cpp:92` `DequantIQ4_NL`; `src/vt/cpu/cpu_quant_dot.cpp:140` `VecDotIQ4_NLQ8_0`; `src/vt/cpu/cpu_quant_blocks.h:62` `BlockIQ4_NL`; traits `src/vt/cpu/cpu_quant_traits.cpp:52` pairing `kIQ4_NL` -> `kQ8_0`. Absent everywhere else: ten `WType` entries in `rocm_quant_dot.hip:580` exclude it, `rocm_grouped_gemm.hip:1783` and `:1855` throw naming it, `IsCudaKeepQuantSupported` omits it, and `kEmbeddingQuant` is registered for `kCPU` and `kCUDA` only. | -| Port map | `cpu_quant_dot.cpp:140` `VecDotIQ4_NLQ8_0` -> a `DotIQ4_NL` device body in `src/vt/rocm/rocm_quant_dot.hip` (scalar, transcribed from the CPU body and NOT from CUDA, per `rocm_quant_dot.hip:226`: gfx1100 has no `__dp4a` so the CUDA integer-core shapes do not port) and a matching `DotIQ4_NL` in `src/vt/cuda/cuda_quant_dot.cu`. `kvalues_iq4nl` -> the ROCm IQ codebook header #3029 introduces, sealed. `cpu_ops.cpp:4282` `kEmbeddingQuant` registration -> a `kROCM` registration decoding IQ4_NL block rows, mirroring the CUDA codec `cuda_quant_dequant.cuh:156` `DqIQ4_NL`. The ROCm Q8_0 activation quantizer already exists (`rocm_grouped_gemm.hip` `QuantizeQ8_0Kernel`); CUDA needs one. | +| Port map | `cpu_quant_dot.cpp:140` `VecDotIQ4_NLQ8_0` -> a `DotIQ4_NL` device body in `src/vt/rocm/rocm_grouped_gemm.hip` (scalar, transcribed from the CPU body and NOT from CUDA: gfx1100 has no hardware `__dp4a`, and the codebook lookup is per nibble so there is nothing for a four-way byte dot to multiply until the values are gathered) feeding TWO kernels, `IQ4NLGemmK` (single matrix) and **`GroupedIQ4NLK` (expert towers, the load-bearing one, because the 48 `ffn_down_exps` reach the GROUPED provider)**, plus a matching `DotIQ4_NL` in `src/vt/cuda/cuda_quant_dot.cu`. `kvalues_iq4nl` -> the ROCm IQ codebook header #3029 introduces, sealed. `cpu_ops.cpp:4282` `kEmbeddingQuant` registration -> a `kROCM` registration decoding IQ4_NL block rows, mirroring the CUDA codec `cuda_quant_dequant.cuh:156` `DqIQ4_NL`. The ROCm Q8_0 activation quantizer already exists (`rocm_grouped_gemm.hip` `QuantizeQ8_0Kernel`); CUDA needs one. | | Tests to port | vLLM has no test to port here, because it has no GGUF k-quant dot. Goldens come from the pinned oracle's OWN kernel over real `UD-IQ1_S` checkpoint bytes, in the shape `tests/vllm/test_gguf_dequant.cpp:528` already uses for the IQ4_NL dequantizer and `tests/vt/test_ops_quant_dot.cpp:869` uses for the IQ4_XS dot: per block and in total, never a synthetic tensor. Newly authored: the IQ4_NL cases in `tests/vt/test_rocm_quant_dot.cpp`, `tests/vt/test_ops_quant_traits.cpp`, `tests/vllm/test_gguf_keep_quant.cpp`, `tests/vt/test_backend_cross_device.cpp`, and a device gather case. | | Gates | See `## Gates`. G1 unit suites on CPU, on `strix:gpu0` for the ROCm arms and on a CUDA box for arm 3; G2 admission of `UD-IQ1_S` with zero reference-tier hits for the IQ4_NL GEMM and gather. NO throughput, latency or memory number is admissible from this row. | | Dependencies | **None blocking.** #3029 was recorded as a hard dependency and is not one: #3097 landed `src/vt/rocm/rocm_quant_iq_tables.h` on `main` carrying `d_kvalues_iq4nl[16]` at `:1198`. #3029 remains a CONFLICT surface (eight files, measured by `git merge-tree`), and its two `sanitize-cpu` reds are a repository-wide pre-existing failure in `dots3` tests it does not touch. Hardware: `strix:gpu0` (`gfx1151`), a fleet device reachable only through an `rc` lease. No CI lane has an AMD runner, so a green CI is not evidence for arms 1 and 2. | @@ -148,13 +148,19 @@ precisely why this is stated as a design decision rather than left to review. This is the structural fact that makes arm 3 a variant rather than a table entry. The resident device path on both backends quantizes activations to `BlockQ8_K` over 256-element super-blocks. IQ4_NL pairs with `BlockQ8_0` over -32. On ROCm the activation quantizer already exists: -`rocm_grouped_gemm.hip` carries `QuantizeQ8_0Kernel` beside `QuantizeQ8KKernel`. -On CUDA it does not, and `cuda_quant_dot.cu:695-698` records the gap in its own -words, marking `DotMXFP4` `[[maybe_unused]]` because it "awaits the -Q8_0-activation GEMM variant above". Arm 3 is that variant. Q5_0, Q4_0 and -MXFP4 all queue behind it, which is this row's leverage and also the reason its -CUDA half is larger than one kernel. +32. **The two backends are asymmetric here, and W1 measured the asymmetry rather +than inheriting this section's first guess.** ROCm turned out to carry a +COMPLETE Q8_0-activation path already: `QuantizeQ8_0K` beside +`QuantizeQ8KKernel`, `DotQ8_0(BlockQ8_0*, BlockQ8_0*)`, and both a single-matrix +and a grouped GEMM consuming `const BlockQ8_0* act`. The ROCm arm is therefore a +new dot slotted into existing machinery, not new machinery, and it is +correspondingly small. + +CUDA has no such path, and `cuda_quant_dot.cu:695-698` records the gap in its +own words, marking `DotMXFP4` `[[maybe_unused]]` because it "awaits the +Q8_0-activation GEMM variant above". The CUDA arm is that variant. Q5_0, Q4_0 +and MXFP4 all queue behind it, which is this row's leverage and also why the +CUDA half is the larger of the two. ### D3. The gather and the dot are both required, and they are not the same code 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 7b97e100b..f5096530e 100644 --- a/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp +++ b/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp @@ -136,12 +136,22 @@ bool KeepNvfp4DType(uint32_t ggml_type) { return ggml_type == 40; } bool DeviceKeepQuantSupported(vt::DType dt, vt::DeviceType dev) { switch (dev) { case vt::DeviceType::kROCM: - // rocm_grouped_gemm.hip implements Q8_0/Q4_K/Q5_K/Q6_K, while + // rocm_grouped_gemm.hip implements Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K, while // rocm_quant_dot.hip adds the seven Q8_K-activation formats below on // both grouped and non-grouped arms. IQ4_XS remains with #3029 and is - // not admitted by this row; Q4_0/Q5_0/IQ2_XS/IQ4_NL/IQ3_S/IQ4_XS/ - // MXFP4 stay on the named expand-or-refuse path. - return dt == vt::DType::kQ8_0 || dt == vt::DType::kQ4_K || + // not admitted by this row; Q4_0/Q5_0/IQ2_XS/IQ3_S/IQ4_XS/MXFP4 stay + // on the named expand-or-refuse path. + // + // IQ4_NL is admitted on BOTH arms or neither. It is the only entry here + // whose activation encoding is Q8_0 rather than Q8_K, and it is served by + // DotIQ4_NL through IQ4NLGemmK (single) and GroupedIQ4NLK (expert + // towers). The grouped arm is the one that matters for the shipped + // Qwen3.8-Flash-Next checkpoints, whose 48 ffn_down_exps are IQ4_NL; + // admitting the encoding with only the single-matrix arm would throw at + // the first expert forward with the model already resident, which is the + // exact failure this predicate exists to prevent. + return dt == vt::DType::kQ8_0 || dt == vt::DType::kIQ4_NL || + dt == vt::DType::kQ4_K || dt == vt::DType::kQ5_K || dt == vt::DType::kQ6_K || dt == vt::DType::kIQ2_XXS || dt == vt::DType::kIQ3_XXS || dt == vt::DType::kQ2_K || dt == vt::DType::kQ3_K || diff --git a/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index 27f91ccf6..f1573a214 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -43,6 +43,7 @@ #include "vt/rocm/rocm_f16_codec.h" #include "vt/rocm/rocm_mmvq_policy.h" #include "vt/rocm/rocm_norm_quant_bridge.h" +#include "vt/rocm/rocm_quant_iq_tables.h" #include "vt/rocm/rocm_runtime.h" // Block layouts — the single source of truth (ggml-common.h mirrors). @@ -51,11 +52,13 @@ namespace vt::rocm { namespace { +using vt::cpu::BlockIQ4_NL; using vt::cpu::BlockQ8_0; using vt::cpu::BlockQ8_K; using vt::cpu::BlockQ4_K; using vt::cpu::BlockQ5_K; using vt::cpu::BlockQ6_K; +using vt::cpu::kQK4_NL; using vt::cpu::kQK8_0; using vt::cpu::kQK_K; @@ -191,6 +194,34 @@ __device__ inline float DotQ8_0(const BlockQ8_0* wb, const BlockQ8_0* ab) { return sumi * (DF16ToF32(wb->d) * DF16ToF32(ab->d)); } +// IQ4_NL x Q8_0: cpu_quant_dot.cpp VecDotIQ4_NLQ8_0, itself a 1:1 port of +// llama.cpp b10451 ggml/src/ggml-cpu/quants.c:1254 +// ggml_vec_dot_iq4_nl_q8_0_generic. IQ4_NL is Q4_0's geometry with the 16-entry +// NON-LINEAR codebook (d_kvalues_iq4nl, stock ggml-common.h:1120) in place of +// the affine step, and it pairs with a Q8_0 activation rather than the Q8_K +// every neighbouring i-quant here uses (ggml-cpu.c:379-384). +// +// KEEP THE ASSOCIATION. The scale product is formed BEFORE the integer sum is +// folded in, `d * (sumi1 + sumi2)`, which is the OPPOSITE order from DotQ8_0 +// directly above. That is upstream's order and it is what makes this GEMM +// bit-reproducible against the oracle; a kernel that reassociates is a +// different kernel and a small golden will not tell you so. +// +// Deliberately NOT Dp4a-vectorized. The codebook lookup is per nibble, so the +// four-way byte dot has nothing to multiply until the values are gathered, and +// gfx1100 has no hardware dp4a anyway (see the note above Dp4a). +__device__ inline float DotIQ4_NL(const BlockIQ4_NL* wb, const BlockQ8_0* ab) { + const float d = DF16ToF32(ab->d) * DF16ToF32(wb->d); + int sumi1 = 0; + int sumi2 = 0; +#pragma unroll + for (int j = 0; j < kQK4_NL / 2; ++j) { + sumi1 += ab->qs[j] * d_kvalues_iq4nl[wb->qs[j] & 0x0F]; + sumi2 += ab->qs[j + kQK4_NL / 2] * d_kvalues_iq4nl[wb->qs[j] >> 4]; + } + return d * static_cast(sumi1 + sumi2); +} + // Q4_K x Q8_K: cuda_quant_dot.cu DotQ4K. dp4a-vectorized, one scale per 32. __device__ inline float DotQ4K(const BlockQ4_K* xb, const BlockQ8_K* yb) { const uint32_t kmask1 = 0x3f3f3f3f, kmask2 = 0x0f0f0f0f, kmask3 = 0x03030303; @@ -787,6 +818,38 @@ __global__ void GroupedQ8_0K(OutT* __restrict__ out, const uint8_t* __restrict__ } } +// IQ4_NL grouped expert GEMM against a Q8_0 activation. This is the arm the +// shipped Qwen3.8-Flash-Next checkpoints actually need: all 48 ffn_down_exps +// are IQ4_NL, and an expert tower reaches the grouped provider rather than the +// single-matrix one. Lane ownership, stride, expert indexing and reduction +// order are copied verbatim from GroupedQ8_0K above. +template +__global__ void GroupedIQ4NLK(OutT* __restrict__ out, const uint8_t* __restrict__ weight, + const BlockQ8_0* __restrict__ act, + const int32_t* __restrict__ expert_ids, int64_t P, int64_t n, + int64_t nb, size_t w_row_bytes, bool bcast) { + const int64_t warp = static_cast(blockIdx.x) * blockDim.y + threadIdx.y; + if (warp >= P * n) return; + const int64_t p = warp / n; + const int64_t j = warp % n; + const int lane = threadIdx.x; + const int64_t e = expert_ids[p]; + const uint8_t* w_row = weight + static_cast(e * n + j) * w_row_bytes; + const BlockQ8_0* a_row = act + (bcast ? 0 : p) * nb; + float partial = 0.0f; + for (int64_t b = lane; b < nb; b += 32) { + const BlockIQ4_NL* wb = reinterpret_cast( + w_row + static_cast(b) * sizeof(BlockIQ4_NL)); + partial += DotIQ4_NL(wb, a_row + b); + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) partial += __shfl_down_sync(0xffffffffULL, partial, off); + if (lane == 0) { + if constexpr (sizeof(OutT) == 4) out[p * n + j] = partial; + else out[p * n + j] = DF32ToBF16(partial); + } +} + // K-quant grouped kernel body (shared by Q4_K/Q5_K/Q6_K instantiations). // Fmt: 0=Q4_K, 1=Q5_K, 2=Q6_K. template @@ -1021,6 +1084,36 @@ __global__ void Q8_0GemmK(OutT* __restrict__ out, const uint8_t* __restrict__ we } } +// IQ4_NL weights against a Q8_0 activation. A SEPARATE kernel rather than a +// weight-templated Q8_0GemmK on purpose: Q8_0GemmK is on the GDN decode path, +// and templating it would re-codegen a kernel this row has no evidence about. +// The lane ownership, the stride and the reduction order are copied from it +// verbatim so the two stay comparable. +template +__global__ void IQ4NLGemmK(OutT* __restrict__ out, const uint8_t* __restrict__ weight, + const BlockQ8_0* __restrict__ act, int64_t m, int64_t n, + int64_t nb) { + const int64_t warp = static_cast(blockIdx.x) * blockDim.y + threadIdx.y; + if (warp >= m * n) return; + const int64_t i = warp / n; + const int64_t j = warp % n; + const int lane = threadIdx.x; + const uint8_t* w_row = weight + static_cast(j * nb) * sizeof(BlockIQ4_NL); + const BlockQ8_0* a_row = act + i * nb; + float partial = 0.0f; + for (int64_t bb = lane; bb < nb; bb += 32) { + const BlockIQ4_NL* wb = reinterpret_cast( + w_row + static_cast(bb) * sizeof(BlockIQ4_NL)); + partial += DotIQ4_NL(wb, a_row + bb); + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) partial += __shfl_down_sync(0xffffffffULL, partial, off); + if (lane == 0) { + if constexpr (sizeof(OutT) == 4) out[i * n + j] = partial; + else out[i * n + j] = DF32ToBF16(partial); + } +} + // The MMVQ decode arm assigns each superblock to an eight-lane octet. The // integer phase can reassociate exactly. Each octet leader then reconstructs // the baseline float term, and the warp replays the baseline lane ownership @@ -1608,6 +1701,27 @@ void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, const Te Check(hipGetLastError(), "q8_0 gemm"); return; } + if (b.dtype == DType::kIQ4_NL) { + NormQuantInvalidate(q); + if (k % kQK4_NL != 0) throw std::runtime_error("vt rocm: matmul_bt_quant IQ4_NL: K%32!=0"); + const int64_t nb = k / kQK4_NL; + // The activation encoding is Q8_0, NOT Q8_K: IQ4_NL is a 32-element block + // and the quantizer below is the one the Q8_0 arm already uses. + BlockQ8_0* qact = static_cast(EnsureQuantScratch( + static_cast(m) * nb * sizeof(BlockQ8_0), s)); + QuantizeQ8_0K<<((m * nb + 127) / 128), 128, 0, s>>>( + qact, a.data, ActDtOf(a.dtype), a.stride[0], m, nb); + Check(hipGetLastError(), "iq4_nl activation quant"); + const int64_t grid = (m * n + kWarpsPerBlock - 1) / kWarpsPerBlock; + if (out.dtype == DType::kF32) + IQ4NLGemmK<<(grid), block, 0, s>>>( + static_cast(out.data), w, qact, m, n, nb); + else + IQ4NLGemmK<<(grid), block, 0, s>>>( + static_cast(out.data), w, qact, m, n, nb); + Check(hipGetLastError(), "iq4_nl gemm"); + return; + } if (b.dtype == DType::kQ4_K || b.dtype == DType::kQ5_K || b.dtype == DType::kQ6_K) { 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; @@ -1779,8 +1893,8 @@ void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, const Te } throw std::runtime_error( "vt rocm: internal GDN matmul_bt_quant provider received an unsupported weight dtype " - "(GDN-owned: Q8_0/Q4_K/Q5_K/Q6_K; wrapper-owned: Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ2_S/IQ1_S/IQ1_XXXS; " - "unimplemented: Q4_0/Q5_0/IQ2_XS/IQ4_NL/IQ3_S/IQ4_XS/MXFP4 -- the ROCm wrapper must delegate only the GDN-owned set)"); + "(GDN-owned: Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K; wrapper-owned: Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ2_S/IQ1_S/IQ1_XXXS; " + "unimplemented: Q4_0/Q5_0/IQ2_XS/IQ3_S/IQ4_XS/MXFP4 -- the ROCm wrapper must delegate only the GDN-owned set)"); } @@ -1822,6 +1936,28 @@ void MatmulBTQuantGroupedKernelRocmGdn(Queue& q, Tensor& out, const Tensor& act, Check(hipGetLastError(), "q8_0 grouped"); return; } + if (weight.dtype == DType::kIQ4_NL) { + if (k % kQK4_NL != 0) + throw std::runtime_error("vt rocm: matmul_bt_quant_grouped IQ4_NL: K must be a multiple of 32"); + const int64_t nb = k / kQK4_NL; + const size_t w_row_bytes = static_cast(nb) * sizeof(BlockIQ4_NL); + // Q8_0 activation, not Q8_K — see DotIQ4_NL. + BlockQ8_0* qact = static_cast(EnsureQuantScratch( + static_cast(Pa) * nb * sizeof(BlockQ8_0), s)); + constexpr int kQBlock = 128; + QuantizeQ8_0K<<((Pa * nb + kQBlock - 1) / kQBlock), kQBlock, 0, s>>>( + qact, act.data, ActDtOf(act.dtype), act.stride[0], Pa, nb); + Check(hipGetLastError(), "iq4_nl activation quant (grouped)"); + const int64_t grid = (P * n + kWarpsPerBlock - 1) / kWarpsPerBlock; + if (out.dtype == DType::kF32) + GroupedIQ4NLK<<(grid), block, 0, s>>>( + static_cast(out.data), w, qact, eids, P, n, nb, w_row_bytes, bcast); + else + GroupedIQ4NLK<<(grid), block, 0, s>>>( + static_cast(out.data), w, qact, eids, P, n, nb, w_row_bytes, bcast); + Check(hipGetLastError(), "iq4_nl grouped"); + return; + } if (weight.dtype == DType::kQ4_K || weight.dtype == DType::kQ5_K || weight.dtype == DType::kQ6_K) { if (k % kQK_K != 0) @@ -1851,8 +1987,8 @@ void MatmulBTQuantGroupedKernelRocmGdn(Queue& q, Tensor& out, const Tensor& act, throw std::runtime_error( "vt rocm: internal GDN matmul_bt_quant_grouped provider received an unsupported weight dtype " - "(GDN-owned: Q8_0/Q4_K/Q5_K/Q6_K; wrapper-owned: Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ2_S/IQ1_S/IQ1_XXXS; " - "unimplemented: Q4_0/Q5_0/IQ2_XS/IQ4_NL/IQ3_S/IQ4_XS/MXFP4 -- the ROCm wrapper must delegate only the GDN-owned set)"); + "(GDN-owned: Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K; wrapper-owned: Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ2_S/IQ1_S/IQ1_XXXS; " + "unimplemented: Q4_0/Q5_0/IQ2_XS/IQ3_S/IQ4_XS/MXFP4 -- the ROCm wrapper must delegate only the GDN-owned set)"); } } // namespace vt::rocm diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 7ca7d5e12..b5b67a2c4 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -2555,6 +2555,11 @@ 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"}, + // IQ4_NL: 18-byte, 32-element block, and the ONLY entry here whose + // activation encoding is Q8_0 rather than Q8_K. It is what every published + // Qwen3.8-Flash-Next quant stores its ffn_down_exps and its n-gram table + // in, so a ROCm arm that lacks it cannot multiply those experts at all. + {vt::DType::kIQ4_NL, 18, 0, -1, "iq4_nl"}, }; const bool rocm_present = OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM); const bool any_rocm = [&] { @@ -2569,7 +2574,8 @@ TEST_CASE("non-grouped keep-quant GEMM (Q8_0/Q4_K/Q5_K/Q6_K) matches the CPU ora } for (const Fmt& f : fmts) { CAPTURE(f.name); - const int64_t elems_per_block = (f.dt == vt::DType::kQ8_0) ? 32 : 256; + const int64_t elems_per_block = + (f.dt == vt::DType::kQ8_0 || f.dt == vt::DType::kIQ4_NL) ? 32 : 256; const int64_t blocks_per_row = K / elems_per_block; const size_t row_bytes = static_cast(blocks_per_row) * f.block_bytes; const size_t wn = static_cast(N) * row_bytes; @@ -3816,6 +3822,12 @@ 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 + // IQ4_NL {d; qs[16]} — 32-element block on a Q8_0 activation. THIS is the + // arm the shipped Qwen3.8-Flash-Next checkpoints need: all 48 + // ffn_down_exps are IQ4_NL and an expert tower reaches the GROUPED + // provider, so the single-matrix arm alone would still throw at the first + // expert forward with the model already resident. + {vt::DType::kIQ4_NL, 18, 0, -1, "iq4_nl"}, }; // REQUIRE-proven registration on ROCm (never a silent skip — review sweep @@ -3832,7 +3844,8 @@ TEST_CASE("grouped quant expert GEMM (Q8_0/Q4_K/Q6_K) matches the CPU oracle") { } for (const Fmt& f : fmts) { CAPTURE(f.name); - const int64_t elems_per_block = (f.dt == vt::DType::kQ8_0) ? 32 : 256; + const int64_t elems_per_block = + (f.dt == vt::DType::kQ8_0 || f.dt == vt::DType::kIQ4_NL) ? 32 : 256; const int64_t blocks_per_row = K / elems_per_block; const size_t row_bytes = static_cast(blocks_per_row) * f.block_bytes; const size_t wn = static_cast(E) * N * row_bytes; @@ -3955,7 +3968,8 @@ TEST_CASE("fused MoE gate+up+SwiGLU grouped GEMM matches the CPU oracle and is N // As std::string: doctest stringifies a bare `const char*` as `1`, so the // capture in the case above cannot name the format that failed. CAPTURE(std::string(f.name)); - const int64_t elems_per_block = (f.dt == vt::DType::kQ8_0) ? 32 : 256; + const int64_t elems_per_block = + (f.dt == vt::DType::kQ8_0 || f.dt == vt::DType::kIQ4_NL) ? 32 : 256; const int64_t blocks_per_row = K / elems_per_block; const size_t row_bytes = static_cast(blocks_per_row) * f.block_bytes; const size_t wn = static_cast(E) * N * row_bytes; From 50a69cac7c99bdccfc8a146590c7c5bde660eddf Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 12 Sep 2026 00:50:50 +0000 Subject: [PATCH 3/9] fix(QUANT-GGUF-IQ4_NL): route IQ4_NL past the ROCm wrapper to its kernels The kernels landed in the GDN provider and nothing could reach them. The wrapper in rocm_quant_dot.hip delegates an explicit allowlist to the *Gdn kernels and refuses everything its own IsRocmKeepQuantSupported does not admit. IQ4_NL was in neither set, so both seams threw on a dtype the tree could now multiply. That is the refusal and its route predicate disagreeing, and the file already warned about this exact class beside the grouped list: Q8_0 has no Q8_K-superblock arm here because it dots a Q8_0 activation, so omitting it turns a served format into a throw. IQ4_NL is the second member of that class and the comment now says so. Measured on strix:gpu0 before the fix, which is the red this change turns green. Both new cross-device cases THREW for the intended reason and every other case passed: cross_device.cpp:2544 no keep-quant kernel for dtype iq4_nl cross_device.cpp:3809 matmul_bt_quant_grouped: no keep-quant kernel iq4_nl 47 cases | 45 passed | 2 failed | 84040 assertions | 0 failed Zero failed assertions against two failed cases is the signature of a route gap rather than a numerical one: nothing computed a wrong answer, the work never started. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .../ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md | 27 ++++++++++++++++++- .agents/quantization-matrix.md | 2 +- .agents/specs/quant-gguf-iq4nl.md | 17 +++++++++--- src/vt/rocm/rocm_quant_dot.hip | 21 ++++++++++----- 4 files changed, 56 insertions(+), 11 deletions(-) diff --git a/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md b/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md index ec9fbe3d2..afabc0f71 100644 --- a/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md +++ b/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md @@ -1,5 +1,5 @@ ID: ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD -Title: IQ4_NL has no native quantized compute on ROCm or CUDA +Title: IQ4_NL has no native quantized compute on ROCm Row: QUANT-GGUF-IQ4_NL State: OPEN Kind: bug @@ -60,3 +60,28 @@ recorded: the ROCm codebook header IQ4_NL needs is now on `main` by another route. #3029 remains a conflict surface over eight files, and its two `sanitize-cpu` reds are a repository-wide pre-existing failure in `dots3` tests it does not touch. + +## Reconciliation 2026-09-12 (second): the CUDA half was never missing + +**The problem statement's CUDA bullet was WRONG when written, not overtaken.** +Recording that distinction, because the two failures need different fixes. + +[#2419](https://github.com/mudler/vllm.cpp/issues/2419) (`593b888b5`, +`feat(QUANT-CUDA-KEEPQUANT-32B): give IQ4_NL, Q5_0 and Q4_0 a device GEMM +instead of a host drain`) had already landed. On CUDA: + +- `IsCuda32BlockKeepQuantSupported` (`cuda_quant_dot.cu:2060`) admits + `kIQ4_NL`, `kQ5_0` and `kQ4_0`; +- the single-matrix (`:2416`), grouped (`:2533`) and fused gate/up (`:2785`) + seams all consult it; +- `docs/FEATURES.md:111` states it in prose, naming this exact checkpoint. + +**How the wrong reading happened, so it does not happen again.** This issue +read `IsCudaKeepQuantSupported`, saw no IQ4_NL, and generalised to "CUDA has no +arm". That predicate governs the **256-element Q8_K family** only. IQ4_NL is a +32-element block on a Q8_0 activation and therefore CANNOT appear in it by +construction; CUDA keeps a second predicate and a second templated GEMM for +exactly that class. One predicate is not the backend. + +**What remains, and it is now a single arm:** the ROCm IQ4_NL keep-quant dot. +Nothing else in this issue's original three bullets is still true. diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index d3ca5609e..2bf9c8e64 100644 --- a/.agents/quantization-matrix.md +++ b/.agents/quantization-matrix.md @@ -72,7 +72,7 @@ forces the full dequant path back. | `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#L172) 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-IQ1_S` | IQ1_S | 19 / output | - | - | - | - | - | `INVENTORIED` | - | leaf open | - | -| `QUANT-GGUF-IQ4_NL` | IQ4_NL | 20 / output | Y | Y | Y | - | - | `ACTIVE` | reader trait [`FindGgmlTraits` case 20](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L292) (18 B block, `QK4_NL = 32`: Q4_0's geometry with the 16-entry NON-LINEAR codebook `kvalues_iq4nl` of llama.cpp `b10451` `ggml/src/ggml-common.h:1120` in place of the affine step), ADDED by [#1989](https://github.com/mudler/vllm.cpp/issues/1989) and gated against the pinned oracle by [`DequantGgufRowToF32 IQ4_NL row matches the pinned oracle`](../tests/vllm/test_gguf_dequant.cpp#L528), with the keep-quant admission asserted BY NAME at [`test_gguf_keep_quant.cpp`](../tests/vllm/test_gguf_keep_quant.cpp#L598) and its device-routing case at [`:655`](../tests/vllm/test_gguf_keep_quant.cpp#L655), together with the dequantizer [`DequantIQ4_NL`](../src/vt/cpu/cpu_quant_dequant.cpp#L92) and a dequantizing gather that makes `GgufTensorRole::kEmbeddingTable` keep-quant eligible. **`C` = `Y` ON CPU ONLY, and the activation pairing is `Q8_0` rather than the `Q8_K` every neighbouring i-quant uses** ([`VecDotIQ4_NLQ8_0`](../src/vt/cpu/cpu_quant_dot.cpp#L140), ported 1:1 from `b10451` `ggml/src/ggml-cpu/quants.c:1254` `ggml_vec_dot_iq4_nl_q8_0_generic`, traits row [`cpu_quant_traits.cpp`](../src/vt/cpu/cpu_quant_traits.cpp#L52) pairing `kIQ4_NL` -> `kQ8_0`, read off `ggml-cpu.c:379-384` and NOT inherited from IQ4_XS, which pairs `Q8_K`). That 32-element block on a Q8_0 activation is exactly why NO accelerator serves it: the resident device GEMM on both ROCm and CUDA quantizes activations to `BlockQ8_K` over 256-element super-blocks. **CUDA has no arm** (`IsCudaKeepQuantSupported` omits it; [`DotMXFP4`](../src/vt/cuda/cuda_quant_dot.cu#L698) already sits there `[[maybe_unused]]` saying in its own comment that it awaits the Q8_0-activation GEMM variant, which Q5_0, Q4_0 and MXFP4 also queue behind). **ROCm has no arm and its refusal is TERMINAL** (ten `WType` entries at [`rocm_quant_dot.hip`](../src/vt/rocm/rocm_quant_dot.hip#L580) exclude it; [`rocm_grouped_gemm.hip:1783`](../src/vt/rocm/rocm_grouped_gemm.hip#L1783) throws naming it, and the backend is discrete so an unsupported dtype cannot fall back to a host kernel that would follow device pointers). **The ROCm GATHER, by contrast, LANDED on 2026-09-12** via [#3097](https://github.com/mudler/vllm.cpp/pull/3097) (`82de418e8`), which registers `OpId::kEmbeddingQuant` for `kROCM` at [`rocm_ops.hip:191`](../src/vt/rocm/rocm_ops.hip#L191) and carries IQ4_NL in its codec list at [`rocm_embedding_quant.hip:88`](../src/vt/rocm/rocm_embedding_quant.hip#L88). `DeviceQuantGatherSupported` is exactly `vt::OpRegistered(kEmbeddingQuant, dev)` ([`gguf_keep_quant.cpp:220`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L220)), so it is now TRUE for ROCm and the [`qwen4_exp_weights.cpp:665`](../src/vllm/model_executor/models/qwen4_exp_weights.cpp#L665) refusal no longer fires on that device. **The gather reads the table and the dot still cannot multiply the experts**, which is the whole remaining gap for this encoding on gfx1151. **Both halves are load-bearing for a SHIPPED artifact and neither is optional:** every published `unsloth/Qwen3.8-Flash-Next-GGUF` quant stores **91,465,564,160 elements** in IQ4_NL, read from the three shard headers over HTTP range requests rather than from prose - the 20M-entry n-gram table, which is gathered and never multiplied, and the 48 `ffn_down_exps`, which are multiplied and never gathered. `UD-IQ1_S` (67.56 GiB) needs IQ4_NL ALONE; `UD-IQ1_M` additionally needs IQ1_M; `UD-Q2_K_XL` (73.45 GiB) additionally needs IQ2_XS, and those two formats are about 97% of its weight elements. **`E`/`P` = `-`:** this row claims no token and no number, and none is admissible while the `gfx1151` token gate reads `TOKEN_GATE=FAIL` at 3 of 6 ([`rocm-gfx1151-q4k-token-gate-v2.md`](specs/rocm-gfx1151-q4k-token-gate-v2.md)). **This cell previously read `INVENTORIED` with R, M, C, E and P all unset**, which understated the tree by three stages: the reader, the dequantizer and the CPU dot had all landed under #1989. | [quant-gguf-iq4nl](specs/quant-gguf-iq4nl.md) | `CLAIM-QUANT-GGUF-IQ4_NL` | +| `QUANT-GGUF-IQ4_NL` | IQ4_NL | 20 / output | Y | Y | Y | - | - | `ACTIVE` | reader trait [`FindGgmlTraits` case 20](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L292) (18 B block, `QK4_NL = 32`: Q4_0's geometry with the 16-entry NON-LINEAR codebook `kvalues_iq4nl` of llama.cpp `b10451` `ggml/src/ggml-common.h:1120` in place of the affine step), ADDED by [#1989](https://github.com/mudler/vllm.cpp/issues/1989) and gated against the pinned oracle by [`DequantGgufRowToF32 IQ4_NL row matches the pinned oracle`](../tests/vllm/test_gguf_dequant.cpp#L528), with the keep-quant admission asserted BY NAME at [`test_gguf_keep_quant.cpp`](../tests/vllm/test_gguf_keep_quant.cpp#L598) and its device-routing case at [`:655`](../tests/vllm/test_gguf_keep_quant.cpp#L655), together with the dequantizer [`DequantIQ4_NL`](../src/vt/cpu/cpu_quant_dequant.cpp#L92) and a dequantizing gather that makes `GgufTensorRole::kEmbeddingTable` keep-quant eligible. **`C` = `Y` ON CPU ONLY, and the activation pairing is `Q8_0` rather than the `Q8_K` every neighbouring i-quant uses** ([`VecDotIQ4_NLQ8_0`](../src/vt/cpu/cpu_quant_dot.cpp#L140), ported 1:1 from `b10451` `ggml/src/ggml-cpu/quants.c:1254` `ggml_vec_dot_iq4_nl_q8_0_generic`, traits row [`cpu_quant_traits.cpp`](../src/vt/cpu/cpu_quant_traits.cpp#L52) pairing `kIQ4_NL` -> `kQ8_0`, read off `ggml-cpu.c:379-384` and NOT inherited from IQ4_XS, which pairs `Q8_K`). That 32-element block on a Q8_0 activation is exactly why NO accelerator serves it: the resident device GEMM on both ROCm and CUDA quantizes activations to `BlockQ8_K` over 256-element super-blocks. **CUDA HAS an arm** since [#2419](https://github.com/mudler/vllm.cpp/issues/2419) (`593b888b5`), on a SECOND templated GEMM beside the Q8_K one: [`IsCuda32BlockKeepQuantSupported`](../src/vt/cuda/cuda_quant_dot.cu#L2060) admits `kIQ4_NL`/`kQ5_0`/`kQ4_0` and is consulted by the single-matrix, grouped and fused gate/up seams alike. **Do not read `IsCudaKeepQuantSupported` and conclude otherwise** — that predicate is the 256-element Q8_K family only, and a 32-element Q8_0-activation encoding can never appear in it. MXFP4 is the one 32-block encoding still owed there. **ROCm has no arm and its refusal is TERMINAL** (ten `WType` entries at [`rocm_quant_dot.hip`](../src/vt/rocm/rocm_quant_dot.hip#L580) exclude it; [`rocm_grouped_gemm.hip:1783`](../src/vt/rocm/rocm_grouped_gemm.hip#L1783) throws naming it, and the backend is discrete so an unsupported dtype cannot fall back to a host kernel that would follow device pointers). **The ROCm GATHER, by contrast, LANDED on 2026-09-12** via [#3097](https://github.com/mudler/vllm.cpp/pull/3097) (`82de418e8`), which registers `OpId::kEmbeddingQuant` for `kROCM` at [`rocm_ops.hip:191`](../src/vt/rocm/rocm_ops.hip#L191) and carries IQ4_NL in its codec list at [`rocm_embedding_quant.hip:88`](../src/vt/rocm/rocm_embedding_quant.hip#L88). `DeviceQuantGatherSupported` is exactly `vt::OpRegistered(kEmbeddingQuant, dev)` ([`gguf_keep_quant.cpp:220`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L220)), so it is now TRUE for ROCm and the [`qwen4_exp_weights.cpp:665`](../src/vllm/model_executor/models/qwen4_exp_weights.cpp#L665) refusal no longer fires on that device. **The gather reads the table and the dot still cannot multiply the experts**, which is the whole remaining gap for this encoding on gfx1151. **Both halves are load-bearing for a SHIPPED artifact and neither is optional:** every published `unsloth/Qwen3.8-Flash-Next-GGUF` quant stores **91,465,564,160 elements** in IQ4_NL, read from the three shard headers over HTTP range requests rather than from prose - the 20M-entry n-gram table, which is gathered and never multiplied, and the 48 `ffn_down_exps`, which are multiplied and never gathered. `UD-IQ1_S` (67.56 GiB) needs IQ4_NL ALONE; `UD-IQ1_M` additionally needs IQ1_M; `UD-Q2_K_XL` (73.45 GiB) additionally needs IQ2_XS, and those two formats are about 97% of its weight elements. **`E`/`P` = `-`:** this row claims no token and no number, and none is admissible while the `gfx1151` token gate reads `TOKEN_GATE=FAIL` at 3 of 6 ([`rocm-gfx1151-q4k-token-gate-v2.md`](specs/rocm-gfx1151-q4k-token-gate-v2.md)). **This cell previously read `INVENTORIED` with R, M, C, E and P all unset**, which understated the tree by three stages: the reader, the dequantizer and the CPU dot had all landed under #1989. | [quant-gguf-iq4nl](specs/quant-gguf-iq4nl.md) | `CLAIM-QUANT-GGUF-IQ4_NL` | | `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#L266)), 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#L172) 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) | - | diff --git a/.agents/specs/quant-gguf-iq4nl.md b/.agents/specs/quant-gguf-iq4nl.md index a17ba3399..9e7f4b149 100644 --- a/.agents/specs/quant-gguf-iq4nl.md +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -51,9 +51,20 @@ Three arms and one record repair, in one change: (`gguf_keep_quant.cpp:220`), so it is now true for ROCm and the `qwen4_exp_weights.cpp:665` refusal no longer fires on this device. This row does not reimplement it. -3. **CUDA keep-quant dot.** The Q8_0-activation GEMM variant the CUDA file is - already written against, plus `DotIQ4_NL` and its entry in - `IsCudaKeepQuantSupported`. +3. ~~**CUDA keep-quant dot.**~~ **ALREADY LANDED, and this spec was WRONG about + it. Struck rather than deleted.** #2419 (`593b888b5`, + `feat(QUANT-CUDA-KEEPQUANT-32B)`) gave IQ4_NL, Q5_0 and Q4_0 a CUDA device + GEMM on a second templated path beside the Q8_K one. + `IsCuda32BlockKeepQuantSupported` (`cuda_quant_dot.cu:2060`) admits + `kIQ4_NL`, and the single-matrix (`:2416`), grouped (`:2533`) and fused + gate/up (`:2785`) seams all consult it. + + **The error is worth naming because it is repeatable.** This spec read + `IsCudaKeepQuantSupported`, found no IQ4_NL, and concluded CUDA had none. + That predicate is the 256-element **Q8_K family** only; a 32-element + Q8_0-activation encoding cannot ever appear in it, and CUDA keeps a second + table for exactly those. Reading one predicate and generalising to "the + backend" is what produced a wrong scope twice in this row. 4. **Record reconciliation** of the `QUANT-GGUF-IQ4_NL` row, which understates the tree: `R`, `M` and CPU `C` all landed under #1989 while the row still read `INVENTORIED` with every stage unset. It moves to `ACTIVE`. diff --git a/src/vt/rocm/rocm_quant_dot.hip b/src/vt/rocm/rocm_quant_dot.hip index 1dd358992..cbf4a6481 100644 --- a/src/vt/rocm/rocm_quant_dot.hip +++ b/src/vt/rocm/rocm_quant_dot.hip @@ -847,12 +847,18 @@ void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const int64_t m = a.shape[0]; const int64_t k = a.shape[1]; const int64_t n = b.shape[0]; - // Delegate Q8_0/Q4_K/Q5_K/Q6_K to the optimized kernels in + // Delegate Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K to the optimized kernels in // rocm_grouped_gemm.hip (renamed *Gdn). These have cooperative-warp dispatch // and format-specific templates tuned for gfx1100; the new CUDA-ported // kernels in this file are slower for these types. The new kernels handle // IQ types and Q2_K/Q3_K. - if (b.dtype == DType::kQ8_0 || b.dtype == DType::kQ4_K || + // IQ4_NL is in the SAME CLASS as Q8_0 here and must be delegated for the same + // reason: it dots a 32-element block against a Q8_0 activation, so it has no + // Q8_K-superblock arm in this file and IsRocmKeepQuantSupported will never + // admit it. Leaving it out of this list turns a format the *Gdn kernel serves + // into a throw — which is exactly what the grouped cross-device case caught. + if (b.dtype == DType::kQ8_0 || b.dtype == DType::kIQ4_NL || + b.dtype == DType::kQ4_K || b.dtype == DType::kQ5_K || b.dtype == DType::kQ6_K) { MatmulBTQuantKernelRocmGdn(q, out, a, b); @@ -913,10 +919,13 @@ void MatmulBTQuantGroupedKernelRocm(Queue& q, Tensor& out, const Tensor& act, const int64_t k = act.shape[1]; if (P == 0 || n == 0) return; - // Delegate Q8_0/Q4_K/Q5_K/Q6_K to the optimized grouped kernels. Q8_0 has no - // Q8_K-superblock arm in this file (it dots a Q8_0 activation), so leaving it - // out of this list turns a format the *Gdn kernel serves today into a throw. - if (weight.dtype == DType::kQ8_0 || weight.dtype == DType::kQ4_K || + // Delegate Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K to the optimized grouped kernels. Q8_0 + // and IQ4_NL have no + // Q8_K-superblock arm in this file (both dot a Q8_0 activation), so leaving + // either out of this list turns a format the *Gdn kernel serves today into a + // throw. + if (weight.dtype == DType::kQ8_0 || weight.dtype == DType::kIQ4_NL || + weight.dtype == DType::kQ4_K || weight.dtype == DType::kQ5_K || weight.dtype == DType::kQ6_K) { MatmulBTQuantGroupedKernelRocmGdn(q, out, act, weight, expert_ids); return; From 266180142d5e5dc7fbba1ae33497f1cb962bdd99 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 12 Sep 2026 00:53:13 +0000 Subject: [PATCH 4/9] spec(QUANT-GGUF-IQ4_NL): record the strix gate, red and green The row's unit gate ran on the hardware it is for. Both runs are recorded, because the red is the evidence that the green means something. RED at 3c529c87e: 47 cases / 45 passed / 2 failed, 84040 assertions / 0 failed GREEN at 6f77c7081: 47 cases / 47 passed / 0 failed / 0 skipped, 84044 assertions / 0 failed, build -Werror with 0 warnings Zero failed assertions against two failed cases is a route gap and not a numerical one, and both failing cases were the two new IQ4_NL ones throwing by name. The assertion count rising by four across the pair is what shows the format executed rather than being skipped; in the red run those cases threw before their checks ran. Also records that the test binary does not find libamdhip64.so.7 by itself in the strix worker. Without LD_LIBRARY_PATH the gate exits 127 having measured nothing, which is a result that looks like a failure and is not one. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md | 2 +- .agents/specs/quant-gguf-iq4nl.md | 42 ++++++++++++++++++++--- 2 files changed, 39 insertions(+), 5 deletions(-) diff --git a/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md b/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md index d63f40c33..f682b5590 100644 --- a/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md +++ b/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-QUANT-GGUF-IQ4_NL` | `QUANT-GGUF-IQ4_NL` (`ACTIVE`) | Claude Code (opus-5) as helper; implementation, review and repair to run as separate fresh agents per `AGENTS.md` "How work gets done" | local worktree `/home/mudler/.cache/sdd/mudler-vllm.cpp/quant-gguf-iq4nl`; ROCm arms gate on `strix:gpu0` (`gfx1151`), a fleet device reached only through an `rc` lease, never `ssh` | `row/QUANT-GGUF-IQ4_NL`, based `18f39771c`; issue `ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD` | Owns: the IQ4_NL keep-quant dot and its device admission on ROCm (`src/vt/rocm/rocm_quant_dot.hip`, `src/vt/rocm/rocm_grouped_gemm.hip`); the CUDA Q8_0-activation GEMM variant plus `DotIQ4_NL` and its `IsCudaKeepQuantSupported` arm (`src/vt/cuda/cuda_quant_dot.cu`); this row's spec, its `.agents/quantization-matrix.md` cell, and the `strix:gpu0` memory rows in `.agents/environment.md`. EXCLUDES the ROCm quantized GATHER, which #3097 landed on 2026-09-12 and which this row therefore does not reimplement; every other missing ROCm format (IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S, MXFP4 — #1940), any tensor-core tile for IQ4_NL, the existing Q8_K-activation dispatch and its kernels (untouched), the seven `qwen4_exp` operations with no ROCm arm (owed, separate row), and every throughput, latency or memory number (gate-blocked) | `ACTIVE` | 2026-09-12 — spec committed, no product code, base `18f39771c`. **Scope shrank during drafting and the record says so:** [#3097](https://github.com/mudler/vllm.cpp/pull/3097) landed the ROCm quantized gather (arm 2) and the ROCm IQ codebook header carrying `d_kvalues_iq4nl`, which together removed the #3029 dependency this spec first recorded. #3029 remains a conflict surface only, and its two `sanitize-cpu` reds are a repository-wide pre-existing failure in `dots3` tests it does not touch. Next: W1, the ROCm `DotIQ4_NL`, which needs no other pull request to land first | +| `CLAIM-QUANT-GGUF-IQ4_NL` | `QUANT-GGUF-IQ4_NL` (`ACTIVE`) | Claude Code (opus-5) as helper; implementation, review and repair to run as separate fresh agents per `AGENTS.md` "How work gets done" | local worktree `/home/mudler/.cache/sdd/mudler-vllm.cpp/quant-gguf-iq4nl`; ROCm arms gate on `strix:gpu0` (`gfx1151`), a fleet device reached only through an `rc` lease, never `ssh` | `row/QUANT-GGUF-IQ4_NL`, based `18f39771c`; issue `ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD` | Owns: the IQ4_NL keep-quant dot and its device admission on ROCm (`src/vt/rocm/rocm_quant_dot.hip`, `src/vt/rocm/rocm_grouped_gemm.hip`); the CUDA Q8_0-activation GEMM variant plus `DotIQ4_NL` and its `IsCudaKeepQuantSupported` arm (`src/vt/cuda/cuda_quant_dot.cu`); this row's spec, its `.agents/quantization-matrix.md` cell, and the `strix:gpu0` memory rows in `.agents/environment.md`. EXCLUDES the ROCm quantized GATHER, which #3097 landed on 2026-09-12 and which this row therefore does not reimplement; every other missing ROCm format (IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S, MXFP4 — #1940), any tensor-core tile for IQ4_NL, the existing Q8_K-activation dispatch and its kernels (untouched), the seven `qwen4_exp` operations with no ROCm arm (owed, separate row), and every throughput, latency or memory number (gate-blocked) | `ACTIVE` | 2026-09-12 — implemented and GATED ON HARDWARE. PR [#3149](https://github.com/mudler/vllm.cpp/pull/3149) (draft). `strix:gpu0` lease, `gfx1151`, HIP 7.2.53211: RED at `3c529c87e` (47 cases / 45 passed / 2 failed, both the new IQ4_NL cases THROWING for the intended reason, 84040 assertions / 0 failed) -> GREEN at `6f77c7081` (47 / 47 / 0 failed / 0 skipped, 84044 assertions / 0 failed, build `-Werror` with 0 warnings). **Scope collapsed from four arms to one during the row and the records say so**: the ROCm gather landed as [#3097](https://github.com/mudler/vllm.cpp/pull/3097), and the CUDA dot was already landed as [#2419](https://github.com/mudler/vllm.cpp/issues/2419) — this row's first reading of CUDA was WRONG, not overtaken, because it consulted `IsCudaKeepQuantSupported` (the 256-element Q8_K family) for a 32-element Q8_0-activation encoding that can never appear in it. Next: fresh scoped review, then mark ready | diff --git a/.agents/specs/quant-gguf-iq4nl.md b/.agents/specs/quant-gguf-iq4nl.md index 9e7f4b149..6523d9923 100644 --- a/.agents/specs/quant-gguf-iq4nl.md +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -321,8 +321,38 @@ kernel exists: ## Gates -- **G1 (unit).** The suites above, green on CPU, on `strix:gpu0` for the ROCm - arms, and on a CUDA box for arm 3. +- **G1 (unit). PASSED on `strix:gpu0`, 2026-09-12.** `rc run -d strix:gpu0`, + HIP 7.2.53211, ROCm 7.2.4, `gfx1151`. Built in the lease from a `--depth 1` + clone of this row's branch with `git rev-parse HEAD` asserted equal to the + commit under test and `git status --porcelain` at 0 bytes. + + | Run | Head | Result | + |---|---|---| + | RED | `3c529c87e` | `47 cases / 45 passed / 2 failed`, `84040 assertions / 0 failed` | + | GREEN | `6f77c7081` | `47 cases / 47 passed / 0 failed / 0 skipped`, `84044 assertions / 0 failed`, `Status: SUCCESS!` | + + **The red is the evidence, not a mishap.** Both failing cases were the two new + IQ4_NL ones and both THREW for the intended reason: + `cross_device.cpp:2544 no keep-quant kernel for dtype iq4_nl` and + `cross_device.cpp:3809 matmul_bt_quant_grouped: no keep-quant kernel for dtype + iq4_nl`. **Zero failed assertions against two failed cases is the signature of + a ROUTE gap rather than a numerical one**: nothing computed a wrong answer, + the work never started. The cause was the wrapper in `rocm_quant_dot.hip`, + whose delegation allowlist and whose `IsRocmKeepQuantSupported` both excluded + IQ4_NL, so the kernels in the GDN provider were unreachable. + + The assertion count RISING by 4 across the pair is what shows the new format + executed rather than being skipped: in the red run those cases threw before + their checks ran. `0 skipped` on the green run is asserted, not assumed. + + Build was clean at `-Werror` with **0 warnings**. + + **One environment fact for the next runner:** the test binary does not find + `libamdhip64.so.7` on its own in the strix worker. Export + `LD_LIBRARY_PATH=$(dirname $(find /opt -name libamdhip64.so.7 | head -1))` + before running it, or the gate exits 127 having measured nothing. + +- **G1 (CUDA).** Not owed by this row: arm 3 was already landed by #2419. - **G2 (admission).** `UD-IQ1_S` opens and its IQ4_NL tensors keep their blocks on ROCm, with `VT_OP_PROVIDER_STATS` showing zero reference-tier hits for the IQ4_NL GEMM and gather. @@ -368,8 +398,12 @@ kernel exists: ## Now -`ACTIVE`, 2026-09-12. Spec committed, no product code, base pinned to -`origin/main` `18f39771c`. +`ACTIVE`, 2026-09-12. Base `18f39771c`. **The ROCm arm is implemented and its +unit gate PASSED on `strix:gpu0`** (see `## Gates` G1). Of the four arms this +spec opened with, three are gone: the gather landed as #3097, the CUDA dot was +already landed as #2419 and this spec was simply wrong about it, and what +remains is the ROCm dot, which is done, plus the record repair, which is in +this change. **The scope shrank between drafting and committing, and the record says so rather than pretending it was always this size.** #3097 landed the ROCm From 547f3318f37b3dda6bd18889b7f233406da54f7d Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 12 Sep 2026 01:48:25 +0000 Subject: [PATCH 5/9] fix(QUANT-GGUF-IQ4_NL): pin the association order and reconcile the records A fresh review returned FAIL on seven findings against the ROCm IQ4_NL arm. This change repairs each of them. ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD. The blocking one was a red gate the first wave never ran. `gguf_keep_quant.cpp` admitted `kIQ4_NL` on ROCm, and `tests/vllm/test_gguf_keep_quant.cpp` holds the ROCm admitted set INDEPENDENTLY, plus a total. The two are deliberately two descriptions of one rule, so moving only the first reds type 20 on both shapes and both totals: `60 cases | 59 passed | 1 failed`. The predicate is NOT narrowed, because admitting IQ4_NL on ROCm is exactly what puts a real `ffn_down_exps` tower on the new kernels; the test's expectation moves instead, and its GEMM term goes 22 -> 24. The second one is the more interesting failure. The spec's D1 and R2 call the association order `d * (sumi1 + sumi2)` load-bearing, and the reviewer mutated `DotIQ4_NL` to `(d*sumi1) + (d*sumi2)`, proved the binary changed, and watched 47 of 47 cases stay green. Every gate this row had was an NMSE band at `5e-4` and reassociation moves the result by about `1e-7` relative. The guarantee was claimed and unpinned, which is worse than not claiming it. It is now pinned bit for bit, by a new case in `test_backend_cross_device.cpp` that runs on every device registering `kMatmulBTQuant`, the CPU included. Three fixture properties make it able to see a one-ulp defect, and each is a way a lazier fixture stops being able to: ONE block at `M = N = 1`, so the warp reduction adds only zeros and the device result is the dot itself; an activation of `k/64` with `max |k| = 127`, so the Q8_0 quantizer is exact and the test KNOWS the quantized bytes rather than asking a kernel for them; and a full-mantissa f16 weight scale, because with a short-mantissa scale both products are exact and the two orders AGREE. A `REQUIRE` asserts the two orders differ on these operands, so the case fails rather than degenerating. Four smaller findings. `test_backend_cross_device.cpp` taught the fused gate+up+SwiGLU case's `elems_per_block` about IQ4_NL without extending its `fmts` list, so that branch was unreachable and read like coverage; the fused seam does serve the format (`rocm_moe_gate_up_swiglu.hip` delegates to the grouped GEMM), so the row is added rather than the line deleted. `docs/FEATURES.md` enumerates the ROCm keep-quant set and omitted IQ4_NL. The `quantization-matrix.md` cell still said ROCm has no arm and its refusal is terminal, contradicting the code in the same commit. The local issue was still OPEN with an empty `Resolution` and a closing line naming the arm that had just landed; it closes here with the dated `strix:gpu0` red-and-green evidence, and its two reconciliation sections are restored verbatim after `agent-issue.py close` dropped them, because deleting evidence to close a record is not closing it. One re-pin that no finding named, and that preflight found: this row entered the runnable population of `check-gate-commands.py` when its spec landed a `## Gates` section naming commands that can fail, and the ratchet requires the baseline to move in the same change. `test_check_gate_commands` was red on this branch before this commit. EXCEPTION ARGUED, not assumed: the 48-line `strix:gpu0` block in `.agents/environment.md` rides in this pull request although it is not this row's own code. AGENTS.md says a record edit rides in the pull request whose change made the record stale. The change that made it stale is the firmware memory carve on this box, which moved the VRAM split to 96.00 GiB and host RAM to 31 GiB, and this row is what measured it, under leases `a8111ff8-3ce8-42f6-9034-36bdd2cacfe4` and `c30dc437-bf12-4a23-ab8b-89b88fe767dd`. The fleet table also never carried the only AMD device on the fleet at all, so a reader sizing a model for it finds nothing and falls back to figures the carve has falsified. Splitting it would cost a branch, a gate run and a fresh review to move text whose measurement lives in this spec. A reviewer who rejects that reason does not merge this change. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .../ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md | 10 +- .agents/quantization-matrix.md | 2 +- .agents/specs/quant-gguf-iq4nl.md | 105 ++++++++++++++---- docs/FEATURES.md | 8 +- scripts/check-gate-commands.py | 7 ++ src/vt/rocm/rocm_quant_dot.hip | 5 +- tests/vllm/test_gguf_keep_quant.cpp | 19 +++- tests/vt/test_backend_cross_device.cpp | 99 +++++++++++++++++ 8 files changed, 224 insertions(+), 31 deletions(-) diff --git a/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md b/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md index afabc0f71..411c5c01a 100644 --- a/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md +++ b/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md @@ -1,14 +1,14 @@ ID: ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD Title: IQ4_NL has no native quantized compute on ROCm Row: QUANT-GGUF-IQ4_NL -State: OPEN +State: CLOSED Kind: bug GitHub: - Mirror: PENDING Availability: FULL Created: 2026-09-11 Updated: 2026-09-12 -Closed: - +Closed: 2026-09-12 ## Problem @@ -26,7 +26,7 @@ Record defect in the same area. .agents/quantization-matrix.md carries QUANT-GGU ## Resolution -- +Fixed by row QUANT-GGUF-IQ4_NL (PR #3149), 2026-09-12. The last remaining arm, the ROCm IQ4_NL keep-quant dot, landed: DotIQ4_NL in src/vt/rocm/rocm_grouped_gemm.hip reads the sealed d_kvalues_iq4nl of rocm_quant_iq_tables.h and feeds IQ4NLGemmK (single matrix) and GroupedIQ4NLK (expert towers) against the QuantizeQ8_0K activation the Q8_0 arm already used; the rocm_quant_dot.hip wrapper delegates kIQ4_NL to those kernels on both the single and the grouped seam, and DeviceKeepQuantSupported admits kIQ4_NL for kROCM in gguf_keep_quant.cpp. Gated on strix:gpu0 (gfx1151, HIP 7.2.53211, ROCm 7.2.4) under an rc lease, red then green on the same box: RED at 3c529c87e, 47 cases / 45 passed / 2 failed, 84040 assertions / 0 failed, both failures the new IQ4_NL cases throwing 'no keep-quant kernel for dtype iq4_nl'; GREEN at 6f77c7081, 47 cases / 47 passed / 0 failed / 0 skipped, 84044 assertions / 0 failed. Full evidence in .agents/specs/quant-gguf-iq4nl.md section Gates, G1. The other two bullets of the problem statement were already reconciled above: the ROCm gather landed as #3097 and the CUDA GEMM had landed as #2419. IsRocmKeepQuantSupported still excludes IQ4_NL by construction, because that predicate is the Q8_K-activation family; the delegation lists are the route. ## Reconciliation 2026-09-12: the gather half landed while this issue was open @@ -85,3 +85,7 @@ exactly that class. One predicate is not the backend. **What remains, and it is now a single arm:** the ROCm IQ4_NL keep-quant dot. Nothing else in this issue's original three bullets is still true. + +**That single arm landed on 2026-09-12 and this issue is CLOSED.** The two +reconciliations above are kept as written, so the order in which the three +bullets fell stays auditable. `## Resolution` carries the dated evidence. diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index 2bf9c8e64..debdffb26 100644 --- a/.agents/quantization-matrix.md +++ b/.agents/quantization-matrix.md @@ -72,7 +72,7 @@ forces the full dequant path back. | `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#L172) 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-IQ1_S` | IQ1_S | 19 / output | - | - | - | - | - | `INVENTORIED` | - | leaf open | - | -| `QUANT-GGUF-IQ4_NL` | IQ4_NL | 20 / output | Y | Y | Y | - | - | `ACTIVE` | reader trait [`FindGgmlTraits` case 20](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L292) (18 B block, `QK4_NL = 32`: Q4_0's geometry with the 16-entry NON-LINEAR codebook `kvalues_iq4nl` of llama.cpp `b10451` `ggml/src/ggml-common.h:1120` in place of the affine step), ADDED by [#1989](https://github.com/mudler/vllm.cpp/issues/1989) and gated against the pinned oracle by [`DequantGgufRowToF32 IQ4_NL row matches the pinned oracle`](../tests/vllm/test_gguf_dequant.cpp#L528), with the keep-quant admission asserted BY NAME at [`test_gguf_keep_quant.cpp`](../tests/vllm/test_gguf_keep_quant.cpp#L598) and its device-routing case at [`:655`](../tests/vllm/test_gguf_keep_quant.cpp#L655), together with the dequantizer [`DequantIQ4_NL`](../src/vt/cpu/cpu_quant_dequant.cpp#L92) and a dequantizing gather that makes `GgufTensorRole::kEmbeddingTable` keep-quant eligible. **`C` = `Y` ON CPU ONLY, and the activation pairing is `Q8_0` rather than the `Q8_K` every neighbouring i-quant uses** ([`VecDotIQ4_NLQ8_0`](../src/vt/cpu/cpu_quant_dot.cpp#L140), ported 1:1 from `b10451` `ggml/src/ggml-cpu/quants.c:1254` `ggml_vec_dot_iq4_nl_q8_0_generic`, traits row [`cpu_quant_traits.cpp`](../src/vt/cpu/cpu_quant_traits.cpp#L52) pairing `kIQ4_NL` -> `kQ8_0`, read off `ggml-cpu.c:379-384` and NOT inherited from IQ4_XS, which pairs `Q8_K`). That 32-element block on a Q8_0 activation is exactly why NO accelerator serves it: the resident device GEMM on both ROCm and CUDA quantizes activations to `BlockQ8_K` over 256-element super-blocks. **CUDA HAS an arm** since [#2419](https://github.com/mudler/vllm.cpp/issues/2419) (`593b888b5`), on a SECOND templated GEMM beside the Q8_K one: [`IsCuda32BlockKeepQuantSupported`](../src/vt/cuda/cuda_quant_dot.cu#L2060) admits `kIQ4_NL`/`kQ5_0`/`kQ4_0` and is consulted by the single-matrix, grouped and fused gate/up seams alike. **Do not read `IsCudaKeepQuantSupported` and conclude otherwise** — that predicate is the 256-element Q8_K family only, and a 32-element Q8_0-activation encoding can never appear in it. MXFP4 is the one 32-block encoding still owed there. **ROCm has no arm and its refusal is TERMINAL** (ten `WType` entries at [`rocm_quant_dot.hip`](../src/vt/rocm/rocm_quant_dot.hip#L580) exclude it; [`rocm_grouped_gemm.hip:1783`](../src/vt/rocm/rocm_grouped_gemm.hip#L1783) throws naming it, and the backend is discrete so an unsupported dtype cannot fall back to a host kernel that would follow device pointers). **The ROCm GATHER, by contrast, LANDED on 2026-09-12** via [#3097](https://github.com/mudler/vllm.cpp/pull/3097) (`82de418e8`), which registers `OpId::kEmbeddingQuant` for `kROCM` at [`rocm_ops.hip:191`](../src/vt/rocm/rocm_ops.hip#L191) and carries IQ4_NL in its codec list at [`rocm_embedding_quant.hip:88`](../src/vt/rocm/rocm_embedding_quant.hip#L88). `DeviceQuantGatherSupported` is exactly `vt::OpRegistered(kEmbeddingQuant, dev)` ([`gguf_keep_quant.cpp:220`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L220)), so it is now TRUE for ROCm and the [`qwen4_exp_weights.cpp:665`](../src/vllm/model_executor/models/qwen4_exp_weights.cpp#L665) refusal no longer fires on that device. **The gather reads the table and the dot still cannot multiply the experts**, which is the whole remaining gap for this encoding on gfx1151. **Both halves are load-bearing for a SHIPPED artifact and neither is optional:** every published `unsloth/Qwen3.8-Flash-Next-GGUF` quant stores **91,465,564,160 elements** in IQ4_NL, read from the three shard headers over HTTP range requests rather than from prose - the 20M-entry n-gram table, which is gathered and never multiplied, and the 48 `ffn_down_exps`, which are multiplied and never gathered. `UD-IQ1_S` (67.56 GiB) needs IQ4_NL ALONE; `UD-IQ1_M` additionally needs IQ1_M; `UD-Q2_K_XL` (73.45 GiB) additionally needs IQ2_XS, and those two formats are about 97% of its weight elements. **`E`/`P` = `-`:** this row claims no token and no number, and none is admissible while the `gfx1151` token gate reads `TOKEN_GATE=FAIL` at 3 of 6 ([`rocm-gfx1151-q4k-token-gate-v2.md`](specs/rocm-gfx1151-q4k-token-gate-v2.md)). **This cell previously read `INVENTORIED` with R, M, C, E and P all unset**, which understated the tree by three stages: the reader, the dequantizer and the CPU dot had all landed under #1989. | [quant-gguf-iq4nl](specs/quant-gguf-iq4nl.md) | `CLAIM-QUANT-GGUF-IQ4_NL` | +| `QUANT-GGUF-IQ4_NL` | IQ4_NL | 20 / output | Y | Y | Y | - | - | `ACTIVE` | reader trait [`FindGgmlTraits` case 20](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L292) (18 B block, `QK4_NL = 32`: Q4_0's geometry with the 16-entry NON-LINEAR codebook `kvalues_iq4nl` of llama.cpp `b10451` `ggml/src/ggml-common.h:1120` in place of the affine step), ADDED by [#1989](https://github.com/mudler/vllm.cpp/issues/1989) and gated against the pinned oracle by [`DequantGgufRowToF32 IQ4_NL row matches the pinned oracle`](../tests/vllm/test_gguf_dequant.cpp#L528), with the keep-quant admission asserted BY NAME at [`test_gguf_keep_quant.cpp`](../tests/vllm/test_gguf_keep_quant.cpp#L598) and its device-routing case at [`:655`](../tests/vllm/test_gguf_keep_quant.cpp#L655), together with the dequantizer [`DequantIQ4_NL`](../src/vt/cpu/cpu_quant_dequant.cpp#L92) and a dequantizing gather that makes `GgufTensorRole::kEmbeddingTable` keep-quant eligible. **`C` = `Y` ON CPU, CUDA AND ROCm, and the activation pairing is `Q8_0` rather than the `Q8_K` every neighbouring i-quant uses** ([`VecDotIQ4_NLQ8_0`](../src/vt/cpu/cpu_quant_dot.cpp#L140), ported 1:1 from `b10451` `ggml/src/ggml-cpu/quants.c:1254` `ggml_vec_dot_iq4_nl_q8_0_generic`, traits row [`cpu_quant_traits.cpp`](../src/vt/cpu/cpu_quant_traits.cpp#L52) pairing `kIQ4_NL` -> `kQ8_0`, read off `ggml-cpu.c:379-384` and NOT inherited from IQ4_XS, which pairs `Q8_K`). That 32-element block on a Q8_0 activation is why NO accelerator served it for so long: the resident device GEMM on both ROCm and CUDA quantizes activations to `BlockQ8_K` over 256-element super-blocks, so each backend needed a SECOND path rather than a table entry. **CUDA HAS an arm** since [#2419](https://github.com/mudler/vllm.cpp/issues/2419) (`593b888b5`), on a SECOND templated GEMM beside the Q8_K one: [`IsCuda32BlockKeepQuantSupported`](../src/vt/cuda/cuda_quant_dot.cu#L2060) admits `kIQ4_NL`/`kQ5_0`/`kQ4_0` and is consulted by the single-matrix, grouped and fused gate/up seams alike. **Do not read `IsCudaKeepQuantSupported` and conclude otherwise** — that predicate is the 256-element Q8_K family only, and a 32-element Q8_0-activation encoding can never appear in it. MXFP4 is the one 32-block encoding still owed there. **ROCm HAS an arm** since QUANT-GGUF-IQ4_NL ([#3149](https://github.com/mudler/vllm.cpp/pull/3149)), on the same second-path shape: [`DotIQ4_NL`](../src/vt/rocm/rocm_grouped_gemm.hip) reads the sealed `d_kvalues_iq4nl` of [`rocm_quant_iq_tables.h`](../src/vt/rocm/rocm_quant_iq_tables.h) and feeds `IQ4NLGemmK` (single matrix) and `GroupedIQ4NLK` (expert towers), both against the `QuantizeQ8_0K` activation the Q8_0 arm already used; `DeviceKeepQuantSupported` admits `kIQ4_NL` for `kROCM` ([`gguf_keep_quant.cpp`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L146)) and the wrapper [`rocm_quant_dot.hip`](../src/vt/rocm/rocm_quant_dot.hip) delegates it to those kernels on BOTH the single and the grouped seam. **`IsRocmKeepQuantSupported` still excludes it and always will** — that predicate is the Q8_K-activation family, exactly as `IsCudaKeepQuantSupported` is on the other backend, and reading it and concluding "no arm" is the error this row made twice. Gated on `strix:gpu0` (`gfx1151`, ROCm 7.2.4) red-then-green, 2026-09-12, in [the spec's `## Gates` G1](specs/quant-gguf-iq4nl.md). The refusal it removes was TERMINAL, because the backend is discrete and an unsupported dtype cannot fall back to a host kernel that would follow device pointers. **The ROCm GATHER, by contrast, LANDED on 2026-09-12** via [#3097](https://github.com/mudler/vllm.cpp/pull/3097) (`82de418e8`), which registers `OpId::kEmbeddingQuant` for `kROCM` at [`rocm_ops.hip:191`](../src/vt/rocm/rocm_ops.hip#L191) and carries IQ4_NL in its codec list at [`rocm_embedding_quant.hip:88`](../src/vt/rocm/rocm_embedding_quant.hip#L88). `DeviceQuantGatherSupported` is exactly `vt::OpRegistered(kEmbeddingQuant, dev)` ([`gguf_keep_quant.cpp:220`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L220)), so it is now TRUE for ROCm and the [`qwen4_exp_weights.cpp:665`](../src/vllm/model_executor/models/qwen4_exp_weights.cpp#L665) refusal no longer fires on that device. **Gather and dot are now BOTH served on ROCm**: the gather reads the n-gram table and the dot multiplies the experts, so nothing in this encoding routes `kExpandBf16` on gfx1151 any more. What still stands between this row and a `qwen4_exp` forward on that box is the seven model operations with no ROCm arm, recorded under the spec's `## Owed` and owned elsewhere. **Both halves are load-bearing for a SHIPPED artifact and neither is optional:** every published `unsloth/Qwen3.8-Flash-Next-GGUF` quant stores **91,465,564,160 elements** in IQ4_NL, read from the three shard headers over HTTP range requests rather than from prose - the 20M-entry n-gram table, which is gathered and never multiplied, and the 48 `ffn_down_exps`, which are multiplied and never gathered. `UD-IQ1_S` (67.56 GiB) needs IQ4_NL ALONE; `UD-IQ1_M` additionally needs IQ1_M; `UD-Q2_K_XL` (73.45 GiB) additionally needs IQ2_XS, and those two formats are about 97% of its weight elements. **`E`/`P` = `-`:** this row claims no token and no number, and none is admissible while the `gfx1151` token gate reads `TOKEN_GATE=FAIL` at 3 of 6 ([`rocm-gfx1151-q4k-token-gate-v2.md`](specs/rocm-gfx1151-q4k-token-gate-v2.md)). **This cell previously read `INVENTORIED` with R, M, C, E and P all unset**, which understated the tree by three stages: the reader, the dequantizer and the CPU dot had all landed under #1989. | [quant-gguf-iq4nl](specs/quant-gguf-iq4nl.md) | `CLAIM-QUANT-GGUF-IQ4_NL` | | `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#L266)), 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#L172) 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) | - | diff --git a/.agents/specs/quant-gguf-iq4nl.md b/.agents/specs/quant-gguf-iq4nl.md index 6523d9923..aad8f7246 100644 --- a/.agents/specs/quant-gguf-iq4nl.md +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -154,6 +154,37 @@ device arms keep that order. A kernel that reassociates is not this kernel, and the goldens below will not detect the difference on small inputs, which is precisely why this is stated as a design decision rather than left to review. +**Stating it was not enough, and a review proved it.** A fresh reviewer mutated +`DotIQ4_NL` to `(d*sumi1) + (d*sumi2)`, proved the binary changed, and watched +all 47 cases stay green on `strix:gpu0`. Every gate this row had was an NMSE +band at `5e-4`; reassociation moves the result by about `1e-7` relative, which +is four orders of magnitude inside it. The guarantee was claimed and unpinned. + +**It is now pinned BIT-EXACTLY** by +`tests/vt/test_backend_cross_device.cpp` `IQ4_NL keeps upstream's association +order d*(s1+s2), bit for bit`, which runs on every device that registers +`kMatmulBTQuant`, the CPU included, so CI carries half of it without an AMD +runner. Three properties make it able to see the defect, and each one is a way +a lazier fixture stops being able to: + +- **One block, `M = N = 1`.** With `nb == 1` the warp reduction adds only zeros + to lane 0's value, so the device result is the dot itself rather than a + reassociated sum over blocks. A multi-block case cannot be compared bit for + bit against anything, because the device sums blocks in warp order and the CPU + sums them sequentially. +- **An exactly-quantizable activation.** The values are `k/64` with + `max |k| = 127`, so `amax/127 = 1/64` is f16-exact and every `roundf` returns + the integer it started from. The test therefore KNOWS the quantized + activation and computes the expected value itself, instead of asking a kernel. +- **A full-mantissa weight scale** (`0x2123` as raw f16 bits). This is the part + that is easy to get wrong: with a short-mantissa scale such as a power of two, + `d*sumi1` and `d*sumi2` are both exact and the two associations AGREE. The + case carries a `REQUIRE` that the two orders differ on its own operands, so a + fixture that degenerates fails instead of passing vacuously. + +The difference it detects is one ulp. That is the true size of the guarantee, +and a test that needs an ulp is the only kind that can hold it. + ### D2. IQ4_NL is a 32-element block on a Q8_0 activation This is the structural fact that makes arm 3 a variant rather than a table @@ -285,9 +316,14 @@ its target by design. Do not quote 96 GiB as an allocatable figure. 16-entry codebook with IQ4_XS, so it is live here rather than theoretical. Mitigation: seal `kValuesIq4nl` by digest and by its lane alphabet, in the shape `tests/vt/test_ops_quant_dot.cpp:717` already uses. -- **R2. Reassociation passes a small golden.** See D1. Mitigation: goldens - taken from the oracle's own kernel over real checkpoint bytes, per block and - in total, not a synthetic tensor. +- **R2. Reassociation passes a small golden.** See D1. **This risk MATERIALISED + and was caught by review, not by a gate.** The planned mitigation (oracle + goldens over real checkpoint bytes) was never delivered, and it would not have + been sufficient on its own either: a golden compared under any tolerance band + cannot see a one-ulp reassociation, and a multi-block golden cannot be + compared without one. The delivered mitigation is the single-block bit-exact + case named in D1, which is the only shape in which the two orders are + distinguishable. - **R3. The CUDA Q8_0-activation variant is a new dispatch path, not an entry.** It can regress the Q8_K path it sits beside. Mitigation: the Q8_K formats' existing cases must stay byte-identical; assert that rather than @@ -301,23 +337,46 @@ its target by design. Do not quote 96 GiB as an allocatable figure. ## Tests -Red first, in this order, each failing for the intended reason before any -kernel exists: - -1. `tests/vt/test_rocm_quant_dot.cpp` — IQ4_NL dot against oracle-produced - goldens over real `UD-IQ1_S` `ffn_down_exps` bytes, per block and total. -2. `tests/vt/test_ops_quant_traits.cpp` — reader and vt geometry agree for - IQ4_NL, 18-byte block, `QK4_NL = 32`, against the oracle's own - `sizeof(block_iq4_nl)` printed by the harness. -3. `tests/vllm/test_gguf_keep_quant.cpp` — routing by name: IQ4_NL routes - `kKeepQuant` on ROCm and CUDA GEMM arms, and on the gather. -4. `tests/vt/test_backend_cross_device.cpp` — IQ4_NL GEMM within NMSE <= 5e-4 - of the CPU oracle, on a real launch shape rather than a toy one. The Q6_K - precedent is explicit about why: its gate ran `M=3, N=8, K=512` against a - production launch of 6400 blocks and never tested a launch. -5. A gather case that decodes an IQ4_NL block row on device and matches - `DequantIQ4_NL`. -6. The three D4 mutations, each restoring the tree byte for byte afterwards. +This list was written as a plan and **four of its six items were not delivered +in the first wave**. A fresh review found that, and the list below now records +what the tree has rather than what the plan wanted. Each undelivered item is +struck with its reason or moved to `## Owed`; none is quietly dropped. + +1. ~~`tests/vt/test_rocm_quant_dot.cpp` — IQ4_NL dot against oracle-produced + goldens over real `UD-IQ1_S` `ffn_down_exps` bytes.~~ **NOT DELIVERED.** + That file compiles the Q8_K-activation family only, and the goldens need the + 67.56 GiB artifact staged beside the oracle. Owed below, and it is a + correctness widening rather than the association guarantee, which item 7 + now holds. +2. ~~`tests/vt/test_ops_quant_traits.cpp` — reader and vt geometry for + IQ4_NL.~~ **NOT NEEDED: already present.** IQ4_NL's 18-byte block and + `QK4_NL = 32` landed under #1989 with the reader arm, and the geometry + cross-check covers it. Re-asserting it would have been a second reading of + the same table. +3. `tests/vllm/test_gguf_keep_quant.cpp` — **DELIVERED, and it was the defect + the review found.** The device admission set in `gguf_keep_quant.cpp` and the + independent set the test holds at its `device_capable` term are two + descriptions of one rule, and the first wave moved only the first: the ROCm + leg red on type 20 and on both totals. The test now admits `kIQ4_NL` on ROCm + and its GEMM term moves 22 -> 24. +4. `tests/vt/test_backend_cross_device.cpp` — **DELIVERED** on the single-matrix + and grouped arms, within NMSE <= 5e-4 of the CPU oracle, and **extended in + this repair to the fused gate+up+SwiGLU arm**, whose `fmts` table the first + wave left unextended while teaching its `elems_per_block` about IQ4_NL. That + branch was unreachable, which is a gate that reads like coverage and measures + nothing. The fused seam does serve the format + (`rocm_moe_gate_up_swiglu.hip:159-160` delegates to the grouped GEMM), so the + row was added rather than the line deleted. +5. ~~A gather case that decodes an IQ4_NL block row on device.~~ **LANDED + ELSEWHERE:** #3097 carries the ROCm gather and its case. This row does not + reimplement it. +6. The three D4 mutations — **PERFORMED BY THE FRESH REVIEW**, which is where + they belong, and one of them survived. See item 7. +7. **NEW, and the reason this section was rewritten:** + `tests/vt/test_backend_cross_device.cpp` `IQ4_NL keeps upstream's association + order d*(s1+s2), bit for bit`. It is the only case here that can fail on a + reassociated kernel. See D1 for why its three fixture properties are each + load-bearing. ## Gates @@ -384,6 +443,12 @@ kernel exists: - IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S and MXFP4 on ROCm (#1940). Wiring the already-written CUDA `DotMXFP4` onto arm 3's variant. - The `hipMalloc` ceiling above 76 GiB on this board, unmeasured by design. +- **An IQ4_NL dot golden from the oracle's own kernel over real `UD-IQ1_S` + `ffn_down_exps` bytes** (the struck test item 1). The delivered gates are an + NMSE band against our own CPU kernel plus the bit-exact association case; what + is missing is an independent numerical authority for the CODEBOOK and the + nibble order, which only upstream's own output can supply. The CPU dequantizer + is gated that way (`test_gguf_dequant.cpp:528`) and the dot is not. ## Stop conditions diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 925bcab36..383ae4020 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -333,10 +333,12 @@ from a prompt: both refuse, naming what is missing. ## Backends and hardware -ROCm GGUF keep-quant supports Q8_0, Q2_K, Q3_K, Q4_K, Q5_K, Q6_K, +ROCm GGUF keep-quant supports Q8_0, IQ4_NL, Q2_K, Q3_K, Q4_K, Q5_K, Q6_K, IQ2_XXS, IQ3_XXS, IQ2_S, IQ1_S, and IQ1_XXXS on matrix and grouped expert -weights. IQ4_XS remains separate work in #3029; unsupported formats expand at -load instead of reaching a device provider that cannot execute them. +weights. IQ4_NL is the one entry served against a Q8_0 activation rather than a +Q8_K one, and it is what the published Qwen3.8-Flash-Next quants store their +`ffn_down_exps` in. IQ4_XS remains separate work in #3029; unsupported formats +expand at load instead of reaching a device provider that cannot execute them. | Backend | vllm.cpp | vLLM | SGLang | llama.cpp | |---|---|---|---|---| diff --git a/scripts/check-gate-commands.py b/scripts/check-gate-commands.py index 88d175c32..075cf6f29 100755 --- a/scripts/check-gate-commands.py +++ b/scripts/check-gate-commands.py @@ -463,6 +463,13 @@ def audit() -> list[dict]: # reachable on this fleet. The credit returns when the row reaches a gate-obliged # state, which its W2 does. RUNNABLE_BASELINE = frozenset({ + # 2026-09-12: +QUANT-GGUF-IQ4_NL. GROWTH, and re-pinned in the change that + # caused it. The row is new (#3149) and its spec's `## Gates` section names + # a `ctest` invocation and the `strix:gpu0` ROCm suite, both of which can + # genuinely fail -- the section records this row's own RED run at + # `3c529c87e` (2 of 47 cases) before the arm existed. A row enters this + # population by what its Gates section can RUN, and these run. + "QUANT-GGUF-IQ4_NL", # 2026-09-05: +BACKEND-TENSTORRENT-KEEPQUANT. GROWTH, and re-pinned in the # same change that caused it. The row is new (#2959, W0 staging) and its # spec's `## Gates` section names `scripts/agent-preflight.sh` per wave, diff --git a/src/vt/rocm/rocm_quant_dot.hip b/src/vt/rocm/rocm_quant_dot.hip index cbf4a6481..949554b8f 100644 --- a/src/vt/rocm/rocm_quant_dot.hip +++ b/src/vt/rocm/rocm_quant_dot.hip @@ -779,7 +779,10 @@ bool IsRocmKeepQuantSupported(DType dt, WType* out) { case DType::kIQ2_S: *out = WType::kIQ2_S; return true; case DType::kIQ1_S: *out = WType::kIQ1_S; return true; case DType::kIQ1_XXXS: *out = WType::kIQ1_XXXS; return true; - // Q4_0 / Q8_0 / MXFP4 dot a Q8_0 activation and have no native arm here. + // Q4_0 / Q8_0 / IQ4_NL / MXFP4 dot a Q8_0 activation and have no native arm + // here. IQ4_NL is in that class since QUANT-GGUF-IQ4_NL, and it is why the + // delegation lists above name it: this predicate can never admit it, so the + // wrapper must hand it to the GDN kernels instead of falling through. default: return false; } } diff --git a/tests/vllm/test_gguf_keep_quant.cpp b/tests/vllm/test_gguf_keep_quant.cpp index 594370d26..938b1a157 100644 --- a/tests/vllm/test_gguf_keep_quant.cpp +++ b/tests/vllm/test_gguf_keep_quant.cpp @@ -656,11 +656,19 @@ TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { type == kIQ2_S || type == kIQ4_XS || type == kMXFP4 || type == kIQ1_XXXS; const bool rocm = kRouteDev == vt::DeviceType::kROCM; + // QUANT-GGUF-IQ4_NL adds kIQ4_NL to the ROCm set. It is the one entry + // here whose activation encoding is Q8_0 rather than Q8_K, and it is + // served on BOTH device arms by `DotIQ4_NL` through `IQ4NLGemmK` + // (single matrix) and `GroupedIQ4NLK` (expert towers). The grouped arm + // is the load-bearing one: the shipped Qwen3.8-Flash-Next checkpoints + // store all 48 `ffn_down_exps` in IQ4_NL, and an expert tower reaches + // the GROUPED provider, so this row is what stops those experts from + // expanding to bf16 on a ROCm box. const bool device_capable = !rocm || type == kQ8_0 || type == kQ2_K || type == kQ3_K || type == kQ4_K || type == kQ5_K || type == kQ6_K || type == kIQ2_XXS || type == kIQ3_XXS || type == kIQ2_S || - type == kIQ1_S || type == kIQ1_XXXS; + type == kIQ1_S || type == kIQ1_XXXS || type == kIQ4_NL; 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) @@ -749,7 +757,7 @@ 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): 17 block-capable encodings x 2 keep-capable - // GEMM roles where the device covers the CPU list; 11 x 2 on ROCm. The + // GEMM roles where the device covers the CPU list; 12 x 2 on ROCm. The // GATHER role adds 19 more (the 17, plus Q8_K and IQ3_S) on a device that // REGISTERS the block // gather, and nothing on a device that does not. Written as named terms @@ -776,7 +784,12 @@ 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 ? 22 : 34; + // + // QUANT-GGUF-IQ4_NL moves the ROCm GEMM term 22 -> 24 and leaves the CPU/CUDA + // term at 34, because IQ4_NL was already in the CPU list and only the DEVICE + // set was narrower. That is the shape of a device-arm port: one encoding, two + // keep-capable GEMM roles, and no change to either gather term. + const int gemm_kept = kRouteDev == vt::DeviceType::kROCM ? 24 : 34; const int gather_kept = vt::OpRegistered(vt::OpId::kEmbeddingQuant, kRouteDev) ? 19 : 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 b5b67a2c4..050ef0349 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -2626,6 +2626,98 @@ TEST_CASE("non-grouped keep-quant GEMM (Q8_0/Q4_K/Q5_K/Q6_K) matches the CPU ora } } +// QUANT-GGUF-IQ4_NL D1/R2: the IQ4_NL association order, pinned BIT-EXACTLY. +// +// `VecDotIQ4_NLQ8_0` and its two device ports form the scale product BEFORE the +// integer sum is folded in -- `d * (sumi1 + sumi2)` -- which is the OPPOSITE +// association from the neighbouring q4_0/Q8_0 kernels. The spec calls that +// load-bearing. The NMSE gates above cannot see it: reassociating moves the +// result by about 1e-7 relative, which is four orders of magnitude inside their +// 5e-4 band, and a fresh review duly mutated `DotIQ4_NL` to +// `(d*sumi1) + (d*sumi2)` and watched 47 of 47 cases stay green. +// +// This case is built so the two orders DIFFER and so the correct one is +// reproducible to the last bit: +// +// * ONE block (K = 32), M = N = 1. With a single block the warp reduction +// adds only zeros to lane 0's value, so the device result is the dot itself +// rather than a reassociated sum of dots. +// * The activation is k/64 with max |k| = 127, so the Q8_0 quantizer is +// EXACT: amax/127 = 1/64 is representable in f16 and every `roundf` lands +// on the integer it started from. `qs` is therefore known here, not guessed. +// * The weight scale carries a FULL f16 mantissa (0x2123), written as raw +// bits. A scale with a short mantissa (a power of two, say) makes `d*sumi1` +// exact and the two orders agree -- which is exactly how a toy fixture +// silently stops discriminating. +// +// The REQUIRE below is the anti-degeneracy guard: if the two orders ever agree +// on these operands the case FAILS rather than passing vacuously. +TEST_CASE("IQ4_NL keeps upstream's association order d*(s1+s2), bit for bit") { + constexpr int64_t M = 1, N = 1, K = 32; + // kvalues_iq4nl, llama.cpp b10451 ggml/src/ggml-common.h:1120. Transcribed + // here rather than included, so the test holds the codebook INDEPENDENTLY of + // whatever the CPU and device kernels read. + static const int kValues[16] = {-127, -104, -83, -65, -49, -35, -22, -10, + 1, 13, 25, 38, 53, 69, 89, 113}; + static const uint8_t kWeightQs[16] = {0xa8, 0x0f, 0xed, 0x48, 0x16, 0x2b, + 0xd2, 0x4b, 0x68, 0x07, 0xa2, 0xb1, + 0x5f, 0x4b, 0xd5, 0x2f}; + // max |k| == 127 on lane 16, which pins amax and therefore the Q8_0 scale. + static const int kActQ[32] = {-18, -50, -71, -52, 63, -115, 126, 124, + -81, -67, 93, 77, -58, 126, 118, 4, + 127, 64, 14, -2, 82, 29, -29, -32, + -19, 105, -84, -95, -56, -81, -81, -95}; + constexpr uint16_t kWeightDBits = 0x2123; // f16, full mantissa + constexpr uint16_t kActDBits = 0x2400; // f16 1/64, what the quantizer must emit + + std::vector wt(18); + std::memcpy(wt.data(), &kWeightDBits, 2); + std::memcpy(wt.data() + 2, kWeightQs, 16); + + std::vector act(K); + for (int64_t j = 0; j < K; ++j) act[j] = static_cast(kActQ[j]) / 64.0f; + + int32_t s1 = 0, s2 = 0; + for (int j = 0; j < 16; ++j) { + s1 += kActQ[j] * kValues[kWeightQs[j] & 0x0F]; + s2 += kActQ[j + 16] * kValues[kWeightQs[j] >> 4]; + } + // `act_d * weight_d`, in that order, because that is the order both kernels + // write and float multiplication is not associative across a rounding. + const float d = vt::F16ToF32(kActDBits) * vt::F16ToF32(kWeightDBits); + const float upstream_order = d * static_cast(s1 + s2); + const float reassociated = d * static_cast(s1) + d * static_cast(s2); + auto bits = [](float f) { uint32_t u; std::memcpy(&u, &f, 4); return u; }; + CAPTURE(s1); + CAPTURE(s2); + REQUIRE_MESSAGE(bits(upstream_order) != bits(reassociated), + "the fixture must be able to SEE a reassociation: these " + "operands make the two orders differ by one ulp, and a " + "fixture where they agree gates nothing"); + + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kMatmulBTQuant, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d_id{dt, 0}; + DevBuf da(dev, q, static_cast(K)); + DevBufBytes dwt(dev, q, wt.size()); + DevBuf dout(dev, q, static_cast(M) * N); + da.Upload(act); + dwt.Upload(wt.data()); + Tensor tact = T2(da.ptr(), d_id, M, K); + Tensor twt = Tensor::Contiguous(dwt.ptr(), vt::DType::kIQ4_NL, d_id, {N, K}); + Tensor tout = T2(dout.ptr(), d_id, M, N); + vt::MatmulBTQuant(q, tout, tact, twt); + const std::vector got = dout.Download(); + REQUIRE(got.size() == 1u); + CHECK(bits(got[0]) == bits(upstream_order)); + CHECK(bits(got[0]) != bits(reassociated)); + dev.DestroyQueue(q); + } +} + // Issue #2511: the keep-quant Q6_K arm at the launch geometry PRODUCTION uses. // // The gate above runs kMatmulBTQuant at M=3, N=8, K=512 -- nsb = 2 and a grid @@ -3941,6 +4033,13 @@ TEST_CASE("fused MoE gate+up+SwiGLU grouped GEMM matches the CPU oracle and is N {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"}, + // IQ4_NL {d; qs[16]} — 32-element block on a Q8_0 activation. The fused + // seam DOES serve it: `rocm_moe_gate_up_swiglu.hip` composes the epilogue + // over two `MatmulBTQuantGroupedKernelRocm` calls, which delegate IQ4_NL to + // `GroupedIQ4NLK`. Listing it here is what makes the `elems_per_block` + // branch below reachable — without the row the branch is dead code that + // reads like coverage. + {vt::DType::kIQ4_NL, 18, 0, -1, "iq4_nl"}, }; const bool rocm_built = [&] { From 50b13c090e550c4c2b07204cf5b60412f7050c8b Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 12 Sep 2026 02:18:40 +0000 Subject: [PATCH 6/9] spec(QUANT-GGUF-IQ4_NL): record the repair gates, CPU and strix The repair commit before this one landed the bit-exact association case and the record fixes. This commit writes down what the gates measured, because a claim without its numbers is the defect the review found in the first place. G1b, `strix:gpu0` under `rc` job 7a757f2a-d078-44f9-998d-2ef504b74cc6, built in the lease from a shallow clone of this branch at 0593ea477 with a clean tree: `test_backend_cross_device` 48 cases / 48 passed / 0 failed / 0 skipped and 84061 assertions, `test_gguf_keep_quant` 61 cases / 61 passed and 12626 assertions. The association mutation `(d*sumi1) + (d*sumi2)` was then applied in the lease, the binary was proved changed by md5, and the new case FAILED on `3240805058 == 3240805059`. One ulp is the whole size of the guarantee, and it is why an NMSE band at 5e-4 could never see it. The tree was restored byte for byte and the suite returned to 48 of 48. G1 CPU: `ctest --test-dir build`, 751 tests, 740 passed. The other 11 are 10 Skipped device or checkpoint arms and one pre-existing failure, `test_rocm_f16_contract`, which asserts that `OwnedTensor::View` carries the repacked markers and does not. The three files that decide it are byte-identical between this row's base and origin/main, and this row touches none of them. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/specs/quant-gguf-iq4nl.md | 40 +++++++++++++++++++++++++++++++ 1 file changed, 40 insertions(+) diff --git a/.agents/specs/quant-gguf-iq4nl.md b/.agents/specs/quant-gguf-iq4nl.md index aad8f7246..43dfc9db1 100644 --- a/.agents/specs/quant-gguf-iq4nl.md +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -411,6 +411,46 @@ struck with its reason or moved to `## Owed`; none is quietly dropped. `LD_LIBRARY_PATH=$(dirname $(find /opt -name libamdhip64.so.7 | head -1))` before running it, or the gate exits 127 having measured nothing. +- **G1b (unit, REVIEW REPAIR). PASSED on `strix:gpu0`, 2026-09-12.** Job + `7a757f2a-d078-44f9-998d-2ef504b74cc6`, `rc run -d strix:gpu0`, built in the + lease from a `--depth 1` clone of this branch at `0593ea477` with + `git status --porcelain` at 0 bytes, `-DVLLM_CPP_HIP=ON + -DVLLM_CPP_HIP_ARCHITECTURES=gfx1151`, `LD_LIBRARY_PATH=/opt/rocm-7.2.4/lib`. + + | Run | `test_backend_cross_device` | `test_gguf_keep_quant` | + |---|---|---| + | GREEN | `48 cases / 48 passed / 0 failed / 0 skipped`, `84061 assertions / 0 failed` | `61 cases / 61 passed / 0 failed / 0 skipped`, `12626 assertions / 0 failed` | + + **The association mutation now REDS, and that is the point of this run.** + `DotIQ4_NL` was rewritten in the lease to `(d*sumi1) + (d*sumi2)`, the binary + was proved changed (`md5` `d2ec91c1...` -> `2d0d4667...`), and the new case + failed: + + ```text + TEST CASE: IQ4_NL keeps upstream's association order d*(s1+s2), bit for bit + ERROR: CHECK( bits(got[0]) == bits(upstream_order) ) is NOT correct! + values: CHECK( 3240805058 == 3240805059 ) + 1 case | 0 passed | 1 failed | 47 skipped + ``` + + **One ulp.** `3240805058` against `3240805059` is the entire size of the + guarantee, and it is why every NMSE gate on this row was blind to it. The + tree was restored byte for byte afterwards (`git status --porcelain` 0 bytes, + `md5` back to `d2ec91c1...`) and the suite returned to + `48 cases / 48 passed / 0 failed`. + +- **G1 (CPU). PASSED, 2026-09-12.** `ctest --test-dir build`, 751 tests, **740 + passed**. The remaining 11 are 10 `Skipped` (CUDA, ROCm, `modelopt`, + `voxtral`, the two `minimax_music3` real-device arms, `capi` device arms) and + one pre-existing failure, `test_rocm_f16_contract`, which asserts that + `OwnedTensor::View` carries the `repacked` / `q8_0_aligned` / + `elem_kn_repacked` markers and does not. The three files that decide it + (`tests/vt/test_rocm_f16_contract.cpp`, `include/vt/tensor.h`, + `include/vllm/model_executor/models/qwen3_5_weights.h`) are byte-identical + between this row's base and `origin/main`, and this row touches none of them. + **This is the gate the first wave omitted**, and it is what found the + `test_gguf_keep_quant` red that the review reported. + - **G1 (CUDA).** Not owed by this row: arm 3 was already landed by #2419. - **G2 (admission).** `UD-IQ1_S` opens and its IQ4_NL tensors keep their blocks on ROCm, with `VT_OP_PROVIDER_STATS` showing zero reference-tier hits for the From fd2a12310ec6f4654ac67c05ee9714130014d38f Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 12 Sep 2026 03:02:45 +0000 Subject: [PATCH 7/9] test(QUANT-GGUF-IQ4_NL): pin the runnable baseline entry and stop the ROCm skip The branch moves `RUNNABLE_BASELINE` in `scripts/check-gate-commands.py` but never touched that checker's test, so `check-pr-size.py` reported the checker change as unevidenced. The move is load-bearing -- dropping the entry reds 17 cases including `test_the_baseline_matches_the_shipped_record` -- so the repair is to write the mutation down rather than to drop the entry. `Iq4nlRunnablePopulationTests` follows the two precedents already in that file: the row is in the pinned population, its audit verdict is earned by the record, and removing the entry breaks the exact pin. The bit-exact association case guarded its device loop with `if (!OpAvailable(kMatmulBTQuant, dt)) continue`, which is the opposite of the convention this same file states twice with explicit comments. On a ROCm build whose IQ4_NL admission regressed, the case would have passed having measured only the CPU. It now REQUIREs the registration whenever the ROCm backend is built, and skips only where the backend is absent. The spec's base was still the pre-rebase `18f39771c`; the merge-base is `97cb6964b`. `## Owed` now records what the bit-exact case does NOT reach: it pins one operand point and one block, so a kernel that reassociated only the reduction ACROSS blocks stays NMSE-gated. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/specs/quant-gguf-iq4nl.md | 16 +++++++-- tests/scripts/test_check_gate_commands.py | 41 +++++++++++++++++++++++ tests/vt/test_backend_cross_device.cpp | 15 +++++++++ 3 files changed, 70 insertions(+), 2 deletions(-) diff --git a/.agents/specs/quant-gguf-iq4nl.md b/.agents/specs/quant-gguf-iq4nl.md index 43dfc9db1..0acf21ac1 100644 --- a/.agents/specs/quant-gguf-iq4nl.md +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -5,7 +5,7 @@ rows: `BACKEND-ROCM` ([#41](https://github.com/mudler/vllm.cpp/issues/41)) and `QUANT-CUDA-GATES`. - Claim: `CLAIM-QUANT-GGUF-IQ4_NL` -- Base: `18f39771c` (`origin/main`, 2026-09-12) +- Base: `97cb6964b` (`origin/main`, 2026-09-12; the merge-base after the review rebase) - Pull request shape: **one pull request** carrying the spec and the implementation, spec committed first. Developer decision, 2026-09-11. @@ -489,6 +489,18 @@ struck with its reason or moved to `## Owed`; none is quietly dropped. is missing is an independent numerical authority for the CODEBOOK and the nibble order, which only upstream's own output can supply. The CPU dequantizer is gated that way (`test_gguf_dequant.cpp:528`) and the dot is not. +- **The ACROSS-block association order, which no bit-exact case pins.** The + case `IQ4_NL keeps upstream's association order d*(s1+s2), bit for bit` + (`tests/vt/test_backend_cross_device.cpp`) is built at ONE operand point and + ONE block (M = N = 1, K = 32) precisely so the warp reduction adds only zeros + and the device result IS the single-block dot. That is what makes it + discriminating WITHIN a block, and it is also its limit: a kernel that kept + `d*(sumi1+sumi2)` inside each block but reassociated the reduction ACROSS + blocks (`nb > 1`) would pass this case unchanged. That order remains gated + only by the NMSE band, which a reassociation of about 1e-7 relative sits four + orders of magnitude inside. Closing it needs a multi-block fixture whose + per-block partials are themselves chosen so the two reduction orders differ + in the last bit. ## Stop conditions @@ -503,7 +515,7 @@ struck with its reason or moved to `## Owed`; none is quietly dropped. ## Now -`ACTIVE`, 2026-09-12. Base `18f39771c`. **The ROCm arm is implemented and its +`ACTIVE`, 2026-09-12. Base `97cb6964b`. **The ROCm arm is implemented and its unit gate PASSED on `strix:gpu0`** (see `## Gates` G1). Of the four arms this spec opened with, three are gone: the gather landed as #3097, the CUDA dot was already landed as #2419 and this spec was simply wrong about it, and what diff --git a/tests/scripts/test_check_gate_commands.py b/tests/scripts/test_check_gate_commands.py index 368d962c2..9581920b7 100644 --- a/tests/scripts/test_check_gate_commands.py +++ b/tests/scripts/test_check_gate_commands.py @@ -1187,3 +1187,44 @@ def test_the_keepquant_row_is_credited_for_real_commands(self): gate_section = spec.split("\n## Gates\n", 1)[1].split("\n## ", 1)[0] self.assertIn("scripts/agent-preflight.sh", gate_section) self.assertNotIn("git diff", gate_section) + + +class Iq4nlRunnablePopulationTests(unittest.TestCase): + """The QUANT-GGUF-IQ4_NL entry is the whole of what the IQ4_NL ROCm dot + row changed in this checker, so the credit has to be checkable rather than + plausible: the row must be IN the pinned population, its audit verdict must + be earned by the record (not merely asserted), and removing the entry must + break the exact pin -- which is what proves the row was pinned because it + ENTERED the runnable population and not to quiet a gate.""" + + ROW = "QUANT-GGUF-IQ4_NL" + + def test_the_row_is_pinned_and_its_verdict_is_runnable(self): + self.assertIn(self.ROW, gates.RUNNABLE_BASELINE) + verdicts = {r["id"]: r["verdict"] for r in gates.audit()} + self.assertEqual(verdicts.get(self.ROW), "runnable") + + def test_dropping_the_iq4nl_entry_breaks_the_pin(self): + # MUTATION, in the direction this pin actually moved: remove the entry + # and the exact-population equality inside the checker has to go red, + # because the audit still reports the row as runnable from its spec. + reduced = set(gates.RUNNABLE_BASELINE) - {self.ROW} + self.assertNotEqual(reduced, set(gates.RUNNABLE_BASELINE)) + runnable = {r["id"] for r in gates.audit() if r["verdict"] == "runnable"} + self.assertNotEqual(runnable, reduced) + self.assertEqual(runnable - reduced, {self.ROW}) + + def test_the_iq4nl_row_is_credited_for_real_commands(self): + # A record row can only be credited for the checker that reads the + # record, so the load-bearing assertion is the last one: the `## Gates` + # section of the row's own spec must name a command that genuinely + # fails on a broken tree -- the unit suite this row's ROCm arm reded + # and then greened on -- not a command that exits 0 in any tree + # (`git diff`). + self.assertIn(self.ROW, gates.RUNNABLE_BASELINE) + spec = ( + gates.ROOT / ".agents/specs/quant-gguf-iq4nl.md" + ).read_text(encoding="utf-8") + gate_section = spec.split("\n## Gates\n", 1)[1].split("\n## ", 1)[0] + self.assertIn("ctest --test-dir build", gate_section) + self.assertNotIn("git diff", gate_section) diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 050ef0349..ab39e561b 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -2695,6 +2695,21 @@ TEST_CASE("IQ4_NL keeps upstream's association order d*(s1+s2), bit for bit") { "operands make the two orders differ by one ulp, and a " "fixture where they agree gates nothing"); + // REQUIRE-proven registration on ROCm (never a silent skip -- review sweep + // on #523: an OpAvailable-guarded case passes green with the registration + // deleted). ROCm is the arm this row added, so a build that carries the + // backend and has lost the IQ4_NL admission must FAIL here rather than fall + // through to a pass that measured only the CPU. + const bool rocm_built = [&] { + for (DeviceType dt : RegisteredDevices()) if (dt == DeviceType::kROCM) return true; + return false; + }(); + if (rocm_built) { + REQUIRE_MESSAGE(OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM), + "kMatmulBTQuant must be registered on ROCm -- a missing " + "registration is a failure, never a skip"); + } + for (DeviceType dt : RegisteredDevices()) { if (!OpAvailable(vt::OpId::kMatmulBTQuant, dt)) continue; CAPTURE(DeviceName(dt)); From 4353b5c32ababdfc02e93a39faf7dc67b0bfdb73 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 12 Sep 2026 03:22:50 +0000 Subject: [PATCH 8/9] spec(QUANT-GGUF-IQ4_NL): record the second review repair's gates G1c carries the numbers for the repair commit before it: `check-pr-size.py` now prints OK where it printed ERROR, the gate-commands suite is green at 70 cases, the CPU `ctest` is 757 of 758 with the pre-existing `test_rocm_f16_contract` red G1 already accounted for, and the `strix:gpu0` suites are green at `fd2a12310`. The strix assertion count rose from G1b's 84061 to 84062, which is the single REQUIRE the anti-skip repair adds. That difference is the evidence the new guard executed rather than being skipped. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/specs/quant-gguf-iq4nl.md | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/.agents/specs/quant-gguf-iq4nl.md b/.agents/specs/quant-gguf-iq4nl.md index 0acf21ac1..a20f5516a 100644 --- a/.agents/specs/quant-gguf-iq4nl.md +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -439,6 +439,24 @@ struck with its reason or moved to `## Owed`; none is quietly dropped. `md5` back to `d2ec91c1...`) and the suite returned to `48 cases / 48 passed / 0 failed`. +- **G1c (unit, SECOND REVIEW REPAIR). PASSED, 2026-09-12**, at `fd2a12310`. + + | Gate | Result | + |---|---| + | `check-pr-size.py --base origin/main --head HEAD --branch row/QUANT-GGUF-IQ4_NL` | `OK: every explicit path class is within its review budget.` (it printed `ERROR: checker change 'scripts/check-gate-commands.py' requires semantic mutation evidence` before this commit, and it exits 0 either way -- the OUTPUT is the result) | + | `tests/scripts/test_check_gate_commands.py` | `70 passed, 4 subtests passed` | + | `ctest --test-dir build`, CPU | `99% tests passed, 1 tests failed out of 758`; the one failure is `test_rocm_f16_contract`, the pre-existing red G1 above already proved this row does not own; 14 skipped | + | `strix:gpu0`, job `2d6f8de9-8e2f-412d-a514-f8db5512de0c` | `test_backend_cross_device` `48 cases / 48 passed / 0 failed / 0 skipped`, `84062 assertions / 0 failed`; `test_gguf_keep_quant` `61 cases / 61 passed / 0 failed / 0 skipped`, `12626 assertions / 0 failed`. `--depth 1` clone at `fd2a12310`, `git status --porcelain` 0 bytes, `LD_LIBRARY_PATH=/opt/rocm-7.2.4/lib` | + + **The baseline entry is written down as a mutation, not as a claim.** + `Iq4nlRunnablePopulationTests` was proved load-bearing by deleting + `"QUANT-GGUF-IQ4_NL"` from `RUNNABLE_BASELINE` with `__pycache__` cleared: all + three of its cases went red, and restoring the line left the working tree + byte-identical and returned the three to green. + + The assertion count on `strix` rose from G1b's `84061` to `84062`, which is + the single REQUIRE the anti-skip repair adds. It EXECUTED; it was not skipped. + - **G1 (CPU). PASSED, 2026-09-12.** `ctest --test-dir build`, 751 tests, **740 passed**. The remaining 11 are 10 `Skipped` (CUDA, ROCm, `modelopt`, `voxtral`, the two `minimax_music3` real-device arms, `capi` device arms) and From f34460fb9ef7d5c2ca100944ba1c8939692e339d Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 12 Sep 2026 03:59:49 +0000 Subject: [PATCH 9/9] spec(QUANT-GGUF-IQ4_NL): correct the branch base, the check-pr-size claim and the NMSE arithmetic The claim record still named `18f39771c` as this branch's base. The branch was rebased and `git merge-base HEAD origin/main` is `97cb6964b`. G1c said `check-pr-size.py` "exits 0 either way -- the OUTPUT is the result". That was wrong. `main()` returns 1 whenever its `errors` list is non-empty, so the gate exits non-zero on failure like every other gate here. The zero came from reading `$?` after a pipe, which reports `tail`'s status and not the checker's. The spec now says so in the open rather than quietly dropping the sentence, and it keeps the half of the advice that survives: read the output, because the exit status alone does not name the failing path class. The NMSE band is `5e-4` and the association effect is about `1e-7` relative. That ratio is `5e3`, so the band is about 3.7 orders of magnitude wide, not four. The same phrase had been copied into the R2 case header, and only the number changes there. The test edit is a comment. Preprocessing `tests/vt/test_backend_cross_device.cpp` with its own `compile_commands.json` flags before and after gives a byte-identical translation unit (`md5 2096681a4bf7714c7b03c5ef10650a1b`, `cmp` exit 0), so the `strix:gpu0` suites recorded at `fd2a12310` are not invalidated and were not re-run. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md | 2 +- .agents/specs/quant-gguf-iq4nl.md | 20 ++++++++++++++++---- tests/vt/test_backend_cross_device.cpp | 4 ++-- 3 files changed, 19 insertions(+), 7 deletions(-) diff --git a/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md b/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md index f682b5590..d9f98f8f9 100644 --- a/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md +++ b/.agents/claims/CLAIM-QUANT-GGUF-IQ4_NL.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-QUANT-GGUF-IQ4_NL` | `QUANT-GGUF-IQ4_NL` (`ACTIVE`) | Claude Code (opus-5) as helper; implementation, review and repair to run as separate fresh agents per `AGENTS.md` "How work gets done" | local worktree `/home/mudler/.cache/sdd/mudler-vllm.cpp/quant-gguf-iq4nl`; ROCm arms gate on `strix:gpu0` (`gfx1151`), a fleet device reached only through an `rc` lease, never `ssh` | `row/QUANT-GGUF-IQ4_NL`, based `18f39771c`; issue `ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD` | Owns: the IQ4_NL keep-quant dot and its device admission on ROCm (`src/vt/rocm/rocm_quant_dot.hip`, `src/vt/rocm/rocm_grouped_gemm.hip`); the CUDA Q8_0-activation GEMM variant plus `DotIQ4_NL` and its `IsCudaKeepQuantSupported` arm (`src/vt/cuda/cuda_quant_dot.cu`); this row's spec, its `.agents/quantization-matrix.md` cell, and the `strix:gpu0` memory rows in `.agents/environment.md`. EXCLUDES the ROCm quantized GATHER, which #3097 landed on 2026-09-12 and which this row therefore does not reimplement; every other missing ROCm format (IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S, MXFP4 — #1940), any tensor-core tile for IQ4_NL, the existing Q8_K-activation dispatch and its kernels (untouched), the seven `qwen4_exp` operations with no ROCm arm (owed, separate row), and every throughput, latency or memory number (gate-blocked) | `ACTIVE` | 2026-09-12 — implemented and GATED ON HARDWARE. PR [#3149](https://github.com/mudler/vllm.cpp/pull/3149) (draft). `strix:gpu0` lease, `gfx1151`, HIP 7.2.53211: RED at `3c529c87e` (47 cases / 45 passed / 2 failed, both the new IQ4_NL cases THROWING for the intended reason, 84040 assertions / 0 failed) -> GREEN at `6f77c7081` (47 / 47 / 0 failed / 0 skipped, 84044 assertions / 0 failed, build `-Werror` with 0 warnings). **Scope collapsed from four arms to one during the row and the records say so**: the ROCm gather landed as [#3097](https://github.com/mudler/vllm.cpp/pull/3097), and the CUDA dot was already landed as [#2419](https://github.com/mudler/vllm.cpp/issues/2419) — this row's first reading of CUDA was WRONG, not overtaken, because it consulted `IsCudaKeepQuantSupported` (the 256-element Q8_K family) for a 32-element Q8_0-activation encoding that can never appear in it. Next: fresh scoped review, then mark ready | +| `CLAIM-QUANT-GGUF-IQ4_NL` | `QUANT-GGUF-IQ4_NL` (`ACTIVE`) | Claude Code (opus-5) as helper; implementation, review and repair to run as separate fresh agents per `AGENTS.md` "How work gets done" | local worktree `/home/mudler/.cache/sdd/mudler-vllm.cpp/quant-gguf-iq4nl`; ROCm arms gate on `strix:gpu0` (`gfx1151`), a fleet device reached only through an `rc` lease, never `ssh` | `row/QUANT-GGUF-IQ4_NL`, based `97cb6964b`; issue `ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD` | Owns: the IQ4_NL keep-quant dot and its device admission on ROCm (`src/vt/rocm/rocm_quant_dot.hip`, `src/vt/rocm/rocm_grouped_gemm.hip`); the CUDA Q8_0-activation GEMM variant plus `DotIQ4_NL` and its `IsCudaKeepQuantSupported` arm (`src/vt/cuda/cuda_quant_dot.cu`); this row's spec, its `.agents/quantization-matrix.md` cell, and the `strix:gpu0` memory rows in `.agents/environment.md`. EXCLUDES the ROCm quantized GATHER, which #3097 landed on 2026-09-12 and which this row therefore does not reimplement; every other missing ROCm format (IQ2_XS, IQ1_M, Q5_0, Q4_0, IQ3_S, MXFP4 — #1940), any tensor-core tile for IQ4_NL, the existing Q8_K-activation dispatch and its kernels (untouched), the seven `qwen4_exp` operations with no ROCm arm (owed, separate row), and every throughput, latency or memory number (gate-blocked) | `ACTIVE` | 2026-09-12 — implemented and GATED ON HARDWARE. PR [#3149](https://github.com/mudler/vllm.cpp/pull/3149) (draft). `strix:gpu0` lease, `gfx1151`, HIP 7.2.53211: RED at `3c529c87e` (47 cases / 45 passed / 2 failed, both the new IQ4_NL cases THROWING for the intended reason, 84040 assertions / 0 failed) -> GREEN at `6f77c7081` (47 / 47 / 0 failed / 0 skipped, 84044 assertions / 0 failed, build `-Werror` with 0 warnings). **Scope collapsed from four arms to one during the row and the records say so**: the ROCm gather landed as [#3097](https://github.com/mudler/vllm.cpp/pull/3097), and the CUDA dot was already landed as [#2419](https://github.com/mudler/vllm.cpp/issues/2419) — this row's first reading of CUDA was WRONG, not overtaken, because it consulted `IsCudaKeepQuantSupported` (the 256-element Q8_K family) for a 32-element Q8_0-activation encoding that can never appear in it. Next: fresh scoped review, then mark ready | diff --git a/.agents/specs/quant-gguf-iq4nl.md b/.agents/specs/quant-gguf-iq4nl.md index a20f5516a..0ad303d3f 100644 --- a/.agents/specs/quant-gguf-iq4nl.md +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -158,7 +158,7 @@ precisely why this is stated as a design decision rather than left to review. `DotIQ4_NL` to `(d*sumi1) + (d*sumi2)`, proved the binary changed, and watched all 47 cases stay green on `strix:gpu0`. Every gate this row had was an NMSE band at `5e-4`; reassociation moves the result by about `1e-7` relative, which -is four orders of magnitude inside it. The guarantee was claimed and unpinned. +is about 3.7 orders of magnitude inside it (`5e-4 / 1e-7 = 5e3`). The guarantee was claimed and unpinned. **It is now pinned BIT-EXACTLY** by `tests/vt/test_backend_cross_device.cpp` `IQ4_NL keeps upstream's association @@ -443,11 +443,23 @@ struck with its reason or moved to `## Owed`; none is quietly dropped. | Gate | Result | |---|---| - | `check-pr-size.py --base origin/main --head HEAD --branch row/QUANT-GGUF-IQ4_NL` | `OK: every explicit path class is within its review budget.` (it printed `ERROR: checker change 'scripts/check-gate-commands.py' requires semantic mutation evidence` before this commit, and it exits 0 either way -- the OUTPUT is the result) | + | `check-pr-size.py --base origin/main --head HEAD --branch row/QUANT-GGUF-IQ4_NL` | `OK: every explicit path class is within its review budget.` (it printed `ERROR: checker change 'scripts/check-gate-commands.py' requires semantic mutation evidence` before this commit; read the OUTPUT, because it names which class failed) | | `tests/scripts/test_check_gate_commands.py` | `70 passed, 4 subtests passed` | | `ctest --test-dir build`, CPU | `99% tests passed, 1 tests failed out of 758`; the one failure is `test_rocm_f16_contract`, the pre-existing red G1 above already proved this row does not own; 14 skipped | | `strix:gpu0`, job `2d6f8de9-8e2f-412d-a514-f8db5512de0c` | `test_backend_cross_device` `48 cases / 48 passed / 0 failed / 0 skipped`, `84062 assertions / 0 failed`; `test_gguf_keep_quant` `61 cases / 61 passed / 0 failed / 0 skipped`, `12626 assertions / 0 failed`. `--depth 1` clone at `fd2a12310`, `git status --porcelain` 0 bytes, `LD_LIBRARY_PATH=/opt/rocm-7.2.4/lib` | + **CORRECTION, third review, 2026-09-12: an earlier version of the row above + said `check-pr-size.py` "exits 0 either way -- the OUTPUT is the result". + That was false.** `scripts/check-pr-size.py` `main()` returns 1 when its + `errors` list is non-empty (`scripts/check-pr-size.py:1043-1048`), so the + gate exits non-zero on failure like every other gate here. The zero came + from reading `$?` after a PIPE, which reports the status of the last stage + (`tail`) and not of the checker. Measured again at `547f3318f`: piped into + `tail -2` the shell reported `0`; redirected to a file the same command + reported `1`. The surviving advice is the other half of that row and it + still holds -- read the output, because the exit status alone does not say + which path class failed. + **The baseline entry is written down as a mutation, not as a claim.** `Iq4nlRunnablePopulationTests` was proved load-bearing by deleting `"QUANT-GGUF-IQ4_NL"` from `RUNNABLE_BASELINE` with `__pycache__` cleared: all @@ -515,8 +527,8 @@ struck with its reason or moved to `## Owed`; none is quietly dropped. discriminating WITHIN a block, and it is also its limit: a kernel that kept `d*(sumi1+sumi2)` inside each block but reassociated the reduction ACROSS blocks (`nb > 1`) would pass this case unchanged. That order remains gated - only by the NMSE band, which a reassociation of about 1e-7 relative sits four - orders of magnitude inside. Closing it needs a multi-block fixture whose + only by the NMSE band, which a reassociation of about 1e-7 relative sits + about 3.7 orders of magnitude inside (`5e-4 / 1e-7 = 5e3`). Closing it needs a multi-block fixture whose per-block partials are themselves chosen so the two reduction orders differ in the last bit. diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index ab39e561b..2a5d45575 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -2632,8 +2632,8 @@ TEST_CASE("non-grouped keep-quant GEMM (Q8_0/Q4_K/Q5_K/Q6_K) matches the CPU ora // integer sum is folded in -- `d * (sumi1 + sumi2)` -- which is the OPPOSITE // association from the neighbouring q4_0/Q8_0 kernels. The spec calls that // load-bearing. The NMSE gates above cannot see it: reassociating moves the -// result by about 1e-7 relative, which is four orders of magnitude inside their -// 5e-4 band, and a fresh review duly mutated `DotIQ4_NL` to +// result by about 1e-7 relative, which is about 3.7 orders of magnitude inside +// their 5e-4 band, and a fresh review duly mutated `DotIQ4_NL` to // `(d*sumi1) + (d*sumi2)` and watched 47 of 47 cases stay green. // // This case is built so the two orders DIFFER and so the correct one is