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