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..f682b5590 --- /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 — 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/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..411c5c01a --- /dev/null +++ b/.agents/issues/QUANT-GGUF-IQ4_NL/ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD.md @@ -0,0 +1,91 @@ +ID: ISSUE-LOCAL-01M29ECVDVH5A89TBKA8K72YRD +Title: IQ4_NL has no native quantized compute on ROCm +Row: QUANT-GGUF-IQ4_NL +State: CLOSED +Kind: bug +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-11 +Updated: 2026-09-12 +Closed: 2026-09-12 + +## 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 + +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 + +**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. + +## 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. + +**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 939763ca1..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 | - | - | - | - | - | `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, 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 new file mode 100644 index 000000000..a20f5516a --- /dev/null +++ b/.agents/specs/quant-gguf-iq4nl.md @@ -0,0 +1,550 @@ +# 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: `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. + +## 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_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. | +| 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.**~~ **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`. + +### 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. + +**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 +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. **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 + +`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. **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 + 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 + +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 + +- **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. + +- **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`. + +- **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 + 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 + 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. +- **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. +- **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 + +- #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. 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 +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 +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. 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/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/src/vt/rocm/rocm_quant_dot.hip b/src/vt/rocm/rocm_quant_dot.hip index 1dd358992..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; } } @@ -847,12 +850,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 +922,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; 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/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 7ca7d5e12..ab39e561b 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; @@ -2620,6 +2626,113 @@ 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"); + + // 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)); + 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 @@ -3816,6 +3929,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 +3951,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; @@ -3928,6 +4048,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 = [&] { @@ -3955,7 +4082,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;