diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index debdffb26..04f8907ad 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#L172) pins the reader and vt tables against each other. **`C` = `Y` since [#2247](https://github.com/mudler/vllm.cpp/issues/2247):** keep-quant [`VecDotIQ2_XSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L783) against the Q8_K activation, ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-cpu/quants.c:948` `ggml_vec_dot_iq2_xs_q8_K_generic`, with [`BlockIQ2_XS`](../src/vt/cpu/cpu_quant_blocks.h#L191) (`ggml-common.h:388-393`) and the [traits row](../src/vt/cpu/cpu_quant_traits.cpp#L111) (`ggml-cpu.c:342-347`), so `HasQuantDotKernel` is TRUE and the loader keeps the blocks. **Gated BIT FOR BIT against the oracle's OWN KERNEL on the same real artifact bytes the decoder gate uses** — 4 super-blocks of `blk.3.ffn_gate_exps.weight` dotted against a Q8_K activation the oracle itself encoded, total and per block ([goldens + provenance](../tests/vt/iq2xs_iq4xs_dot_golden.h#L268), [case](../tests/vt/test_ops_quant_dot.cpp#L857)); the comparison is against upstream's own f32 accumulation, because a cleaner f64 reference agrees with a reduction-order defect. The grid SEAL and the KERNEL are tied together by a mutation, not by inspection: swapping `kIq2xsGrid` for `kIq2xxsGrid` inside the kernel reds this case while the seal itself stays green ([coupling case](../tests/vt/test_ops_quant_dot.cpp#L879)). **Residency MEASURED, not inferred:** `RouteGgufTensor` driven over all 1412 tensors of the staged artifact's own headers routes 774 to `kKeepQuant` for **101.14 GiB resident** against 426.72 GiB before this row — the 82 IQ2_XS tensors stop expanding from 53.33 GiB to 369.00 GiB — which fits the ~119.63 GiB of `dgx:gpu0` with 18.49 GiB of headroom. [Routing table](../tests/vllm/test_gguf_keep_quant.cpp#L425) restated: the GEMM term moved 20 -> 24 and the GATHER term stayed 13. **The CUDA arm now HAS a keep-quant kernel for this dtype** ([#2260](https://github.com/mudler/vllm.cpp/issues/2260), spec [cuda-keepquant-iq2xs-iq4xs](specs/cuda-keepquant-iq2xs-iq4xs.md)): `src/vt/cuda/cuda_quant_dot.cu::DotIQ2XS` is a 1:1 port of `b10451 quants.c:948` reading the generated `src/vt/cuda/cuda_quant_iq_tables.cuh::d_iq2xs_grid`, and `src/vt/cuda/cuda_quant_dot.cu::IsCudaKeepQuantSupported` admits twelve Q8_K-family encodings with IQ2_XS among them, wired into all THREE dispatch switches. Before that the 82 tensors still kept — `src/vllm/model_executor/model_loader/gguf_keep_quant.cpp::DeviceKeepQuantSupported` returns `true` for CUDA on its `default:` arm regardless — so the artifact DID fit and the expert GEMM then took the CPU fallback inside `src/vt/cuda/cuda_quant_dot.cu::MatmulBTQuantGroupedKernelCuda` behind a full `cudaStreamSynchronize` per grouped call, while the fused `src/vt/cuda/cuda_quant_dot.cu::MoeGateUpSwiGLUGroupedCuda` seam THREW `gate/up must be the SAME CUDA keep-quant dtype`. **101.14 GiB remains a RESIDENCY result and no speed number is claimed**, and the kernel is not yet SELECTED by a running model: `Glm5NextHostForward` refuses a non-CPU queue by name, so the wiring is still owed by `MODEL-MM-glm5-next-glm5-next-for-conditional-generation`. NOT reached from this row today — `glm5_next_moe.cpp` is W5's host reference and `laguna.cpp` is the only model on the fused seam — but AGENTS.md routes mergeable MLP projections through `vt::MergedGemmGroup`, so W5b ([#2241](https://github.com/mudler/vllm.cpp/issues/2241)) and W5c ([#2242](https://github.com/mudler/vllm.cpp/issues/2242)) make the throw live on a 101 GiB-resident model. Contrast `QUANT-GGUF-IQ2_XXS` above, which records a CUDA compute landing; this row has none. Carried as O19 in [glm5-next-flash](specs/glm5-next-flash.md). | [glm5-next-flash](specs/glm5-next-flash.md) | - | -| `QUANT-GGUF-IQ3_XXS` | IQ3_XXS | 18 / output | Y | Y | Y | - | - | `READY` | reader trait [`:237`](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L237) (98 B block, ADDED W8); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L349) (`DequantIQ3_XXS` + the shared [`iq3xxs_grid`](../src/vt/cpu/cpu_quant_iq_tables.h), ported 1:1 from llama.cpp `ggml-quants.c:2503` + `ggml-common.h:1007`); vt block dtype `kIQ3_XXS` [geometry](../src/vt/dtype.cpp#L86). **`C` = `Y` (DeepSeek-V4 W8, 2026-07-29):** the keep-quant `vec_dot` [`VecDotIQ3_XXSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L1) (1:1 port of ggml `vec_dot_iq3_xxs_q8_K_generic`, `quants.c:999`) + the [Q8_K traits row](../src/vt/cpu/cpu_quant_traits.cpp#L1) keep IQ3_XXS blocks COMPRESSED — the MEMORY ENABLER for the `UD-IQ2_XXS` **down** routed experts (`ffn_down_exps`, which are IQ3_XXS, NOT IQ2_XXS; without keep-quant they alone OOM the box). Op gate [test_ops_quant_dot](../tests/vt/test_ops_quant_dot.cpp#L1) (vec_dot vs f64 dequant-dot ≤1e-5·L1, NMSE ≤5e-4). **W2b GGUF→tower materialization LANDED (2026-07-29, `CLAIM-DEEPSEEK-V4-W2B`):** `LoadDeepseekV4FromGguf` wires the name-map + keep-quant blocks into the `DeepseekV4` towers (`test_deepseek_v4_gguf_load` 5/5·149). `E`/`P` = `-`: the real DeepSeek-V4 e2e run stays W8-final. **CUDA compute LANDED 2026-07-29 (`KERNEL-QUANT-CIQ-GEMM-CUDA`, `CLAIM-CUDA-KEEPQUANT-GEMM`):** a native kCUDA `kMatmulBTQuant` provider ([cuda_quant_dot.cu](../src/vt/cuda/cuda_quant_dot.cu), MMVQ-style) dots these blocks ON the GPU (Q8_K activation quant + integer dot, dequant-in-kernel, weights stay COMPRESSED in the unified pool), GB10-gated ([test_cuda_quant_dot](../tests/vt/test_cuda_quant_dot.cpp)) 2/2 · 92401 vs the CPU oracle (NMSE ≤1e-6, int core bit-exact) + f64 dequant (≤5e-4), compute-sanitizer 0, RED-first proven — so on a CUDA runner these experts dispatch to the GPU, not the 20 ARM cores. | [CUDA keep-quant GEMM](specs/cuda-keepquant-gemm.md); [iquant DSV4 spike](specs/gguf-iquant-dsv4.md) | - | +| `QUANT-GGUF-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 | Y | Y | Y | - | - | `ACTIVE` | reader trait [`FindGgmlTraits` case 20](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L292) (18 B block, `QK4_NL = 32`: Q4_0's geometry with the 16-entry NON-LINEAR codebook `kvalues_iq4nl` of llama.cpp `b10451` `ggml/src/ggml-common.h:1120` in place of the affine step), ADDED by [#1989](https://github.com/mudler/vllm.cpp/issues/1989) and gated against the pinned oracle by [`DequantGgufRowToF32 IQ4_NL row matches the pinned oracle`](../tests/vllm/test_gguf_dequant.cpp#L528), with the keep-quant admission asserted BY NAME at [`test_gguf_keep_quant.cpp`](../tests/vllm/test_gguf_keep_quant.cpp#L598) and its device-routing case at [`:655`](../tests/vllm/test_gguf_keep_quant.cpp#L655), together with the dequantizer [`DequantIQ4_NL`](../src/vt/cpu/cpu_quant_dequant.cpp#L92) and a dequantizing gather that makes `GgufTensorRole::kEmbeddingTable` keep-quant eligible. **`C` = `Y` ON CPU, CUDA AND ROCm, and the activation pairing is `Q8_0` rather than the `Q8_K` every neighbouring i-quant uses** ([`VecDotIQ4_NLQ8_0`](../src/vt/cpu/cpu_quant_dot.cpp#L140), ported 1:1 from `b10451` `ggml/src/ggml-cpu/quants.c:1254` `ggml_vec_dot_iq4_nl_q8_0_generic`, traits row [`cpu_quant_traits.cpp`](../src/vt/cpu/cpu_quant_traits.cpp#L52) pairing `kIQ4_NL` -> `kQ8_0`, read off `ggml-cpu.c:379-384` and NOT inherited from IQ4_XS, which pairs `Q8_K`). That 32-element block on a Q8_0 activation is why NO accelerator served it for so long: the resident device GEMM on both ROCm and CUDA quantizes activations to `BlockQ8_K` over 256-element super-blocks, so each backend needed a SECOND path rather than a table entry. **CUDA HAS an arm** since [#2419](https://github.com/mudler/vllm.cpp/issues/2419) (`593b888b5`), on a SECOND templated GEMM beside the Q8_K one: [`IsCuda32BlockKeepQuantSupported`](../src/vt/cuda/cuda_quant_dot.cu#L2060) admits `kIQ4_NL`/`kQ5_0`/`kQ4_0` and is consulted by the single-matrix, grouped and fused gate/up seams alike. **Do not read `IsCudaKeepQuantSupported` and conclude otherwise** — that predicate is the 256-element Q8_K family only, and a 32-element Q8_0-activation encoding can never appear in it. MXFP4 is the one 32-block encoding still owed there. **ROCm HAS an arm** since QUANT-GGUF-IQ4_NL ([#3149](https://github.com/mudler/vllm.cpp/pull/3149)), on the same second-path shape: [`DotIQ4_NL`](../src/vt/rocm/rocm_grouped_gemm.hip) reads the sealed `d_kvalues_iq4nl` of [`rocm_quant_iq_tables.h`](../src/vt/rocm/rocm_quant_iq_tables.h) and feeds `IQ4NLGemmK` (single matrix) and `GroupedIQ4NLK` (expert towers), both against the `QuantizeQ8_0K` activation the Q8_0 arm already used; `DeviceKeepQuantSupported` admits `kIQ4_NL` for `kROCM` ([`gguf_keep_quant.cpp`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L146)) and the wrapper [`rocm_quant_dot.hip`](../src/vt/rocm/rocm_quant_dot.hip) delegates it to those kernels on BOTH the single and the grouped seam. **`IsRocmKeepQuantSupported` still excludes it and always will** — that predicate is the Q8_K-activation family, exactly as `IsCudaKeepQuantSupported` is on the other backend, and reading it and concluding "no arm" is the error this row made twice. Gated on `strix:gpu0` (`gfx1151`, ROCm 7.2.4) red-then-green, 2026-09-12, in [the spec's `## Gates` G1](specs/quant-gguf-iq4nl.md). The refusal it removes was TERMINAL, because the backend is discrete and an unsupported dtype cannot fall back to a host kernel that would follow device pointers. **The ROCm GATHER, by contrast, LANDED on 2026-09-12** via [#3097](https://github.com/mudler/vllm.cpp/pull/3097) (`82de418e8`), which registers `OpId::kEmbeddingQuant` for `kROCM` at [`rocm_ops.hip:191`](../src/vt/rocm/rocm_ops.hip#L191) and carries IQ4_NL in its codec list at [`rocm_embedding_quant.hip:88`](../src/vt/rocm/rocm_embedding_quant.hip#L88). `DeviceQuantGatherSupported` is exactly `vt::OpRegistered(kEmbeddingQuant, dev)` ([`gguf_keep_quant.cpp:220`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L220)), so it is now TRUE for ROCm and the [`qwen4_exp_weights.cpp:665`](../src/vllm/model_executor/models/qwen4_exp_weights.cpp#L665) refusal no longer fires on that device. **Gather and dot are now BOTH served on ROCm**: the gather reads the n-gram table and the dot multiplies the experts, so nothing in this encoding routes `kExpandBf16` on gfx1151 any more. What still stands between this row and a `qwen4_exp` forward on that box is the seven model operations with no ROCm arm, recorded under the spec's `## Owed` and owned elsewhere. **Both halves are load-bearing for a SHIPPED artifact and neither is optional:** every published `unsloth/Qwen3.8-Flash-Next-GGUF` quant stores **91,465,564,160 elements** in IQ4_NL, read from the three shard headers over HTTP range requests rather than from prose - the 20M-entry n-gram table, which is gathered and never multiplied, and the 48 `ffn_down_exps`, which are multiplied and never gathered. `UD-IQ1_S` (67.56 GiB) needs IQ4_NL ALONE; `UD-IQ1_M` additionally needs IQ1_M; `UD-Q2_K_XL` (73.45 GiB) additionally needs IQ2_XS, and those two formats are about 97% of its weight elements. **`E`/`P` = `-`:** this row claims no token and no number, and none is admissible while the `gfx1151` token gate reads `TOKEN_GATE=FAIL` at 3 of 6 ([`rocm-gfx1151-q4k-token-gate-v2.md`](specs/rocm-gfx1151-q4k-token-gate-v2.md)). **This cell previously read `INVENTORIED` with R, M, C, E and P all unset**, which understated the tree by three stages: the reader, the dequantizer and the CPU dot had all landed under #1989. | [quant-gguf-iq4nl](specs/quant-gguf-iq4nl.md) | `CLAIM-QUANT-GGUF-IQ4_NL` | | `QUANT-GGUF-IQ3_S` | IQ3_S | 21 / IQ3 S/XS/M storage | Y | Y | - | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 21](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L305) (110 B block, ADDED by [#2510](https://github.com/mudler/vllm.cpp/issues/2510)); [codebook dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L399) `DequantIQ3_S` + the 512-entry [`kIq3sGrid`](../src/vt/cpu/cpu_quant_iq_tables.h#L168), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2607` + `ggml/src/ggml-common.h:1052`; vt block dtype `kIQ3_S` [geometry](../src/vt/dtype.cpp#L150) and the CUDA gather codec [`DqIQ3_S`](../src/vt/cuda/cuda_quant_dequant.cuh#L476). **THE REFUSAL THIS ROW EXISTS FOR:** id 21 was the ONE hole in the reader's i-quant run 16..23, and `GgufFile::Open` therefore refused whole artifacts over it — `unsloth/Qwen3.8-27B-GGUF` `UD-Q4_K_M.gguf` stores 4 of its 866 tensors in IQ3_S and `UD-Q4_K_XL.gguf` stores 1 of 866, while every OTHER encoding in those files (F32, Q3_K, Q4_K, Q5_K, Q6_K, Q8_0, IQ4_NL, IQ4_XS) was already handled. That is the file the published Strix Halo rows ran, so [#2497](https://github.com/mudler/vllm.cpp/issues/2497) had to substitute the plain `Q4_K_M`. **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.11.ffn_gate.weight` at absolute offset 4,262,628,128 in the staged `Qwen3.8-27B-UD-Q4_K_M.gguf` (sha256 `322e194f…3482`), decoded by the pinned upstream's own `dequantize_row_iq3_s` and compared as raw f32 bit patterns ([goldens + provenance](../tests/vt/iq3s_golden_vectors.h#L69), [unit](../tests/vllm/test_gguf_dequant.cpp#L663)); the four blocks span 190 distinct 9-bit grid indices of which 73 are >= 256, so the `qh`-spliced ninth bit is exercised rather than left at zero. The 110-byte stride is the ORACLE's own `sizeof(block_iq3_s)`, printed by the harness, not a sum read off the struct. The table itself is sealed by FNV-1a digest and by its lane alphabet ([test_ops_quant_dot](../tests/vt/test_ops_quant_dot.cpp#L717)), because `kIq3sGrid` (512 u32) and `kIq3xxsGrid` (256 u32) are the same shape and a decoder pointed at the sibling still decodes. Reached through `GgufFile::Open` + `DequantGgufRowToF32` in the same file, which is the production path the real artifact took. **`C` = `-`, DELIBERATELY, and it is the one FILE encoding in this tree that decodes without dotting.** `KeepQuantDType` gates on `HasQuantDotKernel`, so IQ3_S routes `kExpandBf16` on the GEMM arm of EVERY tier (CPU, CUDA, ROCm, Metal, Vulkan) and `kKeepQuant` only on the gather, where CPU and CUDA both decode. That is asserted BY NAME rather than left to a count ([routing](../tests/vllm/test_gguf_keep_quant.cpp#L429), [decode-only class](../tests/vt/test_ops_quant_traits.cpp#L266)), so the day the `vec_dot` lands those cases red. **The CPU `VecDotIQ3_SQ8_K` and the CUDA `WType::kIQ3_S` are OWED AS ONE UNIT** (`ggml-cpu/quants.c:1094`, `.vec_dot_type = GGML_TYPE_Q8_K` at `ggml-cpu.c:355-360`): `IsCudaKeepQuantSupported` has no refusal arm, so landing the CPU half alone would flip the loader to keep IQ3_S blocks and send every CUDA IQ3_S GEMM through the host fallback that [#2260](https://github.com/mudler/vllm.cpp/issues/2260) MEASURED to segfault on a discrete card. On the measured artifact the expansion is 146.13 MiB of blocks against 680.00 MiB of bf16, 3.4 % of a 15.33 GiB file — bounded here, and NOT bounded on an artifact whose experts are IQ3_S, which is why it is recorded rather than waved at (compare [#1870](https://github.com/mudler/vllm.cpp/issues/1870)). `E`/`P` = `-`: this row makes the artifact OPENABLE and claims no token and no number; the `gfx1151` quant-matched decode number stays owed by [#2497](https://github.com/mudler/vllm.cpp/issues/2497). | [gguf-iq3s](specs/gguf-iq3s.md) | - | | `QUANT-GGUF-IQ2_S` | IQ2_S | 22 / IQ2_M storage | Y | - | - | - | - | `INVENTORIED` | reader `:229-234`; [explicit rejection](../tests/vllm/test_gguf_dequant.cpp#L223) | leaf open | - | -| `QUANT-GGUF-IQ4_XS` | IQ4_XS | 23 / output | Y | Y | Y | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 23](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L311) (136 B block; the trait predates the decoder, the decoder is [#2240](https://github.com/mudler/vllm.cpp/issues/2240)); [dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L488) `DequantIQ4_XS`, ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2743`; vt block dtype `kIQ4_XS` [geometry](../src/vt/dtype.cpp#L140). NOT a codebook delta from IQ4_NL — it reuses `kValuesIq4nl` unchanged and differs only in the SUPER-BLOCK SCALE LAYOUT (a 6-bit `ls` spliced from a `scales_l` nibble and a `scales_h` bit pair, then biased by -32). **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.11.ffn_down_exps.weight` from the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL arm, whose four super-blocks span `ls` 0..61 so both nibbles and both `scales_h` bit pairs and both signs of `dl` are exercised ([goldens + provenance](../tests/vt/iq2xs_iq4xs_golden_vectors.h#L276), [unit](../tests/vllm/test_gguf_dequant.cpp#L576)). [Geometry cross-check](../tests/vt/test_ops_quant_traits.cpp#L172) pins the reader and vt tables against each other. The former `explicit rejection` evidence is GONE by construction — that test now guards Q1_0 (41), the remaining tabulated-but-undecodable id. **`C` = `Y` since [#2247](https://github.com/mudler/vllm.cpp/issues/2247):** keep-quant [`VecDotIQ4_XSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L844), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-cpu/quants.c:1283` `ggml_vec_dot_iq4_xs_q8_K_generic`, with [`BlockIQ4_XS`](../src/vt/cpu/cpu_quant_blocks.h#L205) (`ggml-common.h:454-460`) and the [traits row](../src/vt/cpu/cpu_quant_traits.cpp#L121). **The activation pairing is `Q8_K`, READ OFF the oracle and not inherited from IQ4_NL:** `type_traits_cpu` at `ggml/src/ggml-cpu/ggml-cpu.c:385-390` carries `.vec_dot = ggml_vec_dot_iq4_xs_q8_K, `.vec_dot_type = GGML_TYPE_Q8_K`, against `:379-384`'s `GGML_TYPE_Q8_0` for IQ4_NL — the 16-entry codebook is shared, the block geometry is not, and a 256-element super-block pairs with the 256-element activation encoding. **Gated BIT FOR BIT against the oracle's OWN KERNEL** on 4 super-blocks of `blk.11.ffn_down_exps.weight`, total and per block ([goldens](../tests/vt/iq2xs_iq4xs_dot_golden.h#L269), [case](../tests/vt/test_ops_quant_dot.cpp#L869)). These 3 tensors stop expanding from 3.59 GiB to 13.50 GiB; the whole-artifact measurement is in the `QUANT-GGUF-IQ2_XS` row above. **It also unblocks a SECOND consumer:** the expert-tower streaming lane is all-or-nothing (`GgufExpertTowersReachSlotLane` returns false on the FIRST tower that does not keep), so the 4 IQ4_XS tensors of the GLM-5.3 `UD-IQ1_S` arm were dropping an arm whose other 221 towers all kept — one tower 6.375 -> 24.000 GiB, and a 4096-slot cache 25.5 -> 96 GiB. Asserted in the `kStackedExpertWeight` role both models store these towers in ([lane case](../tests/vllm/model_executor/test_gguf_device_fit.cpp#L672)). **The CUDA arm now HAS a keep-quant kernel for this dtype too** ([#2260](https://github.com/mudler/vllm.cpp/issues/2260), spec [cuda-keepquant-iq2xs-iq4xs](specs/cuda-keepquant-iq2xs-iq4xs.md)): `src/vt/cuda/cuda_quant_dot.cu::DotIQ4XS` is a 1:1 port of `b10451 quants.c:1283` over the generated `src/vt/cuda/cuda_quant_iq_tables.cuh::d_kvalues_iq4nl`, keeping upstream's EIGHT f32 accumulation steps per super-block rather than one integer core, because that association is what the oracle's golden numbers were produced with. **This completes the device admission of the GLM-5.3 non-flash `UD-IQ1_S` arm**, whose other five encodings (IQ1_S, IQ3_XXS, IQ2_XXS, Q2_K, Q3_K) were already there. Before it, these 3 tensors fit `dgx:gpu0` with their expert GEMM on the CPU fallback behind a `cudaStreamSynchronize` and the fused `src/vt/cuda/cuda_quant_dot.cu::MoeGateUpSwiGLUGroupedCuda` seam throwing. The fit remains RESIDENCY and no speed number is claimed; O19's premise in [glm5-next-flash](specs/glm5-next-flash.md) is discharged, and what still blocks a CUDA run of either model is that model's own host-f32 forward. | [glm5-next-flash](specs/glm5-next-flash.md) | - | +| `QUANT-GGUF-IQ4_XS` | IQ4_XS | 23 / output | Y | Y | Y | - | - | `PARTIAL` | reader trait [`FindGgmlTraits` case 23](../src/vllm/model_executor/model_loader/gguf_reader.cpp#L311) (136 B block; the trait predates the decoder, the decoder is [#2240](https://github.com/mudler/vllm.cpp/issues/2240)); [dequant kernel](../src/vt/cpu/cpu_quant_dequant.cpp#L488) `DequantIQ4_XS`, ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-quants.c:2743`; vt block dtype `kIQ4_XS` [geometry](../src/vt/dtype.cpp#L140). NOT a codebook delta from IQ4_NL — it reuses `kValuesIq4nl` unchanged and differs only in the SUPER-BLOCK SCALE LAYOUT (a 6-bit `ls` spliced from a `scales_l` nibble and a `scales_h` bit pair, then biased by -32). **`M` = `Y`, gated BYTE-FOR-BYTE against the oracle** over REAL checkpoint bytes: 4 blocks of `blk.11.ffn_down_exps.weight` from the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL arm, whose four super-blocks span `ls` 0..61 so both nibbles and both `scales_h` bit pairs and both signs of `dl` are exercised ([goldens + provenance](../tests/vt/iq2xs_iq4xs_golden_vectors.h#L276), [unit](../tests/vllm/test_gguf_dequant.cpp#L576)). [Geometry cross-check](../tests/vt/test_ops_quant_traits.cpp#L172) pins the reader and vt tables against each other. The former `explicit rejection` evidence is GONE by construction — that test now guards Q1_0 (41), the remaining tabulated-but-undecodable id. **`C` = `Y` since [#2247](https://github.com/mudler/vllm.cpp/issues/2247):** keep-quant [`VecDotIQ4_XSQ8_K`](../src/vt/cpu/cpu_quant_dot.cpp#L844), ported 1:1 from llama.cpp `b10451` `ggml/src/ggml-cpu/quants.c:1283` `ggml_vec_dot_iq4_xs_q8_K_generic`, with [`BlockIQ4_XS`](../src/vt/cpu/cpu_quant_blocks.h#L205) (`ggml-common.h:454-460`) and the [traits row](../src/vt/cpu/cpu_quant_traits.cpp#L121). **The activation pairing is `Q8_K`, READ OFF the oracle and not inherited from IQ4_NL:** `type_traits_cpu` at `ggml/src/ggml-cpu/ggml-cpu.c:385-390` carries `.vec_dot = ggml_vec_dot_iq4_xs_q8_K, `.vec_dot_type = GGML_TYPE_Q8_K`, against `:379-384`'s `GGML_TYPE_Q8_0` for IQ4_NL — the 16-entry codebook is shared, the block geometry is not, and a 256-element super-block pairs with the 256-element activation encoding. **Gated BIT FOR BIT against the oracle's OWN KERNEL** on 4 super-blocks of `blk.11.ffn_down_exps.weight`, total and per block ([goldens](../tests/vt/iq2xs_iq4xs_dot_golden.h#L269), [case](../tests/vt/test_ops_quant_dot.cpp#L869)). These 3 tensors stop expanding from 3.59 GiB to 13.50 GiB; the whole-artifact measurement is in the `QUANT-GGUF-IQ2_XS` row above. **It also unblocks a SECOND consumer:** the expert-tower streaming lane is all-or-nothing (`GgufExpertTowersReachSlotLane` returns false on the FIRST tower that does not keep), so the 4 IQ4_XS tensors of the GLM-5.3 `UD-IQ1_S` arm were dropping an arm whose other 221 towers all kept — one tower 6.375 -> 24.000 GiB, and a 4096-slot cache 25.5 -> 96 GiB. Asserted in the `kStackedExpertWeight` role both models store these towers in ([lane case](../tests/vllm/model_executor/test_gguf_device_fit.cpp#L672)). **The CUDA arm now HAS a keep-quant kernel for this dtype too** ([#2260](https://github.com/mudler/vllm.cpp/issues/2260), spec [cuda-keepquant-iq2xs-iq4xs](specs/cuda-keepquant-iq2xs-iq4xs.md)): `src/vt/cuda/cuda_quant_dot.cu::DotIQ4XS` is a 1:1 port of `b10451 quants.c:1283` over the generated `src/vt/cuda/cuda_quant_iq_tables.cuh::d_kvalues_iq4nl`, keeping upstream's EIGHT f32 accumulation steps per super-block rather than one integer core, because that association is what the oracle's golden numbers were produced with. **This completes the device admission of the GLM-5.3 non-flash `UD-IQ1_S` arm**, whose other five encodings (IQ1_S, IQ3_XXS, IQ2_XXS, Q2_K, Q3_K) were already there. Before it, these 3 tensors fit `dgx:gpu0` with their expert GEMM on the CPU fallback behind a `cudaStreamSynchronize` and the fused `src/vt/cuda/cuda_quant_dot.cu::MoeGateUpSwiGLUGroupedCuda` seam throwing. The fit remains RESIDENCY and no speed number is claimed; O19's premise in [glm5-next-flash](specs/glm5-next-flash.md) is discharged, and what still blocks a CUDA run of either model is that model's own host-f32 forward. **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 cdc0f7374..f95c345ea 100644 --- a/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md +++ b/.agents/specs/kernel-quant-ciq-gemm-rocm-iquant.md @@ -11,7 +11,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 @@ -112,26 +113,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.** IQ4_XS shares the existing `nsb = K/256` decomposition, so it compiles against the current launch shape. Re-measure the `nwarps=8` choice rather @@ -149,28 +155,88 @@ llama.cpp, pin `b10451` per `.agents/upstream-sync.md`. ## Tests -- Extend `test_ops_quant_dot.cpp`'s existing IQ4_XS `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`): bit-exact 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 IQ4_XS to the CPU-vs-ROCM cross-check. -- Rerun `ROCM-KQUANT-NWARPS-DECODE`'s own measurement recipe - (`rocprofv3 --kernel-trace` on a real quant-matched trace workload) for - IQ4_XS 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 @@ -178,10 +244,11 @@ llama.cpp, pin `b10451` per `.agents/upstream-sync.md`. MXFP4): tracked by #1940 and left for a follow-on row. - A WMMA/tensor-core tile for IQ4_XS, if the scalar tier's measured throughput warrants one (mirroring how `KERNEL-QUANT-CIQ-GEMM-ROCM` - followed the existing scalar formats): 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. + followed the existing four formats' scalar tier): not attempted here. +- 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 @@ -197,7 +264,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. The GFX1100 reconstruction supplies the former IQ3_XXS scope. -Next: W0 probes the FMA-contraction question on gfx1200, then W1 ports -`DotIQ4XS` with its 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 383ae4020..89796c603 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -346,7 +346,7 @@ expand at load instead of reaching a device provider that cannot execute them. | 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 grouped BF16 expert kernels pass component gates on gfx1100; full Qwen3 MoE token parity remains open ([spec](../.agents/specs/rocm-bf16-moe.md)) | 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 b72c953d5..2dea52f81 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -859,8 +859,8 @@ skips with that refusal quoted. [complete model gate](../.agents/specs/rocm-f16-weights.md) remains failing. - On ROCm, GGUF mixture-of-experts checkpoints compute on the quantized expert blocks (Q8_0, Q2_K, Q3_K, Q4_K, Q5_K, Q6_K, IQ2_XXS, IQ3_XXS, - IQ2_S, IQ1_S, IQ1_XXXS) instead of being dequantized to bf16 at load time. - IQ4_XS remains separate work in #3029. + IQ2_S, IQ1_S, IQ1_XXXS, IQ4_XS) 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 f5096530e..3fc5d6947 100644 --- a/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp +++ b/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp @@ -136,23 +136,21 @@ bool KeepNvfp4DType(uint32_t ggml_type) { return ggml_type == 40; } bool DeviceKeepQuantSupported(vt::DType dt, vt::DeviceType dev) { switch (dev) { case vt::DeviceType::kROCM: - // rocm_grouped_gemm.hip implements Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K, while - // rocm_quant_dot.hip adds the seven Q8_K-activation formats below on - // both grouped and non-grouped arms. IQ4_XS remains with #3029 and is - // not admitted by this row; Q4_0/Q5_0/IQ2_XS/IQ3_S/IQ4_XS/MXFP4 stay - // on the named expand-or-refuse path. + // rocm_grouped_gemm.hip implements Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K/IQ4_XS/IQ3_XXS, + // while rocm_quant_dot.hip adds the six Q8_K-activation formats below + // (IQ2_XXS, Q2_K, Q3_K, IQ2_S, IQ1_S, IQ1_XXXS) on both grouped and + // non-grouped arms. // // IQ4_NL is admitted on BOTH arms or neither. It is the only entry here - // whose activation encoding is Q8_0 rather than Q8_K, and it is served by - // DotIQ4_NL through IQ4NLGemmK (single) and GroupedIQ4NLK (expert - // towers). The grouped arm is the one that matters for the shipped - // Qwen3.8-Flash-Next checkpoints, whose 48 ffn_down_exps are IQ4_NL; - // admitting the encoding with only the single-matrix arm would throw at - // the first expert forward with the model already resident, which is the - // exact failure this predicate exists to prevent. + // whose activation encoding is Q8_0 rather than Q8_K, served by DotIQ4_NL + // through IQ4NLGemmK (single) and GroupedIQ4NLK (expert towers). + // + // IQ4_XS/IQ3_XXS added by KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940). + // Q4_0/Q5_0/IQ2_XS/IQ3_S/MXFP4 stay on the named expand-or-refuse path. return dt == vt::DType::kQ8_0 || dt == vt::DType::kIQ4_NL || dt == vt::DType::kQ4_K || dt == vt::DType::kQ5_K || dt == vt::DType::kQ6_K || + dt == vt::DType::kIQ4_XS || dt == vt::DType::kIQ2_XXS || dt == vt::DType::kIQ3_XXS || dt == vt::DType::kQ2_K || dt == vt::DType::kQ3_K || dt == vt::DType::kIQ2_S || dt == vt::DType::kIQ1_S || diff --git a/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index f1573a214..270c9ce84 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -43,11 +43,14 @@ #include "vt/rocm/rocm_f16_codec.h" #include "vt/rocm/rocm_mmvq_policy.h" #include "vt/rocm/rocm_norm_quant_bridge.h" -#include "vt/rocm/rocm_quant_iq_tables.h" +#include "vt/rocm/rocm_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 { @@ -59,6 +62,8 @@ using vt::cpu::BlockQ4_K; using vt::cpu::BlockQ5_K; using vt::cpu::BlockQ6_K; using vt::cpu::kQK4_NL; +using vt::cpu::BlockIQ4_XS; +using vt::cpu::BlockIQ3_XXS; using vt::cpu::kQK8_0; using vt::cpu::kQK_K; @@ -325,6 +330,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 @@ -852,6 +949,12 @@ __global__ void GroupedIQ4NLK(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, @@ -871,10 +974,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); @@ -957,10 +1063,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); @@ -1604,6 +1722,23 @@ const void* NormQuantLastScratchForTesting() { return g_norm_quant_last_scratch.load(std::memory_order_acquire); } +// 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. // @@ -1723,13 +1858,22 @@ void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, const Te 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; - 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); const detail::MmvqRoute mmvq_route = detail::SelectMmvqRoute( std::getenv("VT_GEMV_MMVQ"), std::getenv("VT_GEMV_MMVQ_FOLD_MAX"), m, @@ -1885,6 +2029,8 @@ void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, const Te 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{}); @@ -1893,8 +2039,7 @@ void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, const Te } throw std::runtime_error( "vt rocm: internal GDN matmul_bt_quant provider received an unsupported weight dtype " - "(GDN-owned: Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K; wrapper-owned: Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ2_S/IQ1_S/IQ1_XXXS; " - "unimplemented: Q4_0/Q5_0/IQ2_XS/IQ3_S/IQ4_XS/MXFP4 -- the ROCm wrapper must delegate only the GDN-owned set)"); + "(GDN-owned: Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K/IQ4_XS/IQ3_XXS; wrapper-owned: Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ2_S/IQ1_S/IQ1_XXXS; unimplemented: Q4_0/Q5_0/IQ2_XS/IQ3_S/MXFP4 -- the ROCm wrapper must delegate only the GDN-owned set)\n" } @@ -1959,25 +2104,34 @@ void MatmulBTQuantGroupedKernelRocmGdn(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{}); @@ -1987,8 +2141,7 @@ void MatmulBTQuantGroupedKernelRocmGdn(Queue& q, Tensor& out, const Tensor& act, throw std::runtime_error( "vt rocm: internal GDN matmul_bt_quant_grouped provider received an unsupported weight dtype " - "(GDN-owned: Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K; wrapper-owned: Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ2_S/IQ1_S/IQ1_XXXS; " - "unimplemented: Q4_0/Q5_0/IQ2_XS/IQ3_S/IQ4_XS/MXFP4 -- the ROCm wrapper must delegate only the GDN-owned set)"); + "(GDN-owned: Q8_0/IQ4_NL/Q4_K/Q5_K/Q6_K/IQ4_XS/IQ3_XXS; wrapper-owned: Q2_K/Q3_K/IQ2_XXS/IQ3_XXS/IQ2_S/IQ1_S/IQ1_XXXS; unimplemented: Q4_0/Q5_0/IQ2_XS/IQ3_S/MXFP4 -- the ROCm wrapper must delegate only the GDN-owned set)\n" } } // 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/tests/vllm/model_executor/test_gguf_device_fit.cpp b/tests/vllm/model_executor/test_gguf_device_fit.cpp index 291ee1a60..7fa06dc91 100644 --- a/tests/vllm/model_executor/test_gguf_device_fit.cpp +++ b/tests/vllm/model_executor/test_gguf_device_fit.cpp @@ -704,21 +704,32 @@ 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. ROCm now - // serves 11 formats, but the IQ2_XS and IQ4_XS tensors in this fixture remain - // unsupported (#1940 owns both gaps). 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 either gap lands. + // it RED on every ROCm build for a reason the case did not state -- + // `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,10 +943,15 @@ TEST_CASE( vllm::GgufLoadPolicy rocm = PolicyWith(true, false, false, false); rocm.device = vt::DeviceType::kROCM; - // NO PLAN: ROCm supports IQ2_XXS, but not this fixture's IQ4_XS tower - // (#1940), so that tower expands 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 938b1a157..ab5cad6b6 100644 --- a/tests/vllm/test_gguf_keep_quant.cpp +++ b/tests/vllm/test_gguf_keep_quant.cpp @@ -587,6 +587,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. @@ -643,8 +665,10 @@ 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. This row gives ROCm the - // seven Q8_K-activation formats in addition to its four prior formats. + // narrower than the loader's CPU-derived list — ROCm now serves 12 + // formats: Q8_0/Q4_K/Q5_K/Q6_K, the seven Q8_K-activation formats + // (IQ2_XXS, Q2_K, Q3_K, IQ2_S, IQ1_S, IQ1_XXXS, IQ3_XXS), and IQ4_XS + // via KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940). // 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. @@ -668,7 +692,8 @@ TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { !rocm || type == kQ8_0 || type == kQ2_K || type == kQ3_K || type == kQ4_K || type == kQ5_K || type == kQ6_K || type == kIQ2_XXS || type == kIQ3_XXS || type == kIQ2_S || - type == kIQ1_S || type == kIQ1_XXXS || type == kIQ4_NL; + type == kIQ1_S || type == kIQ1_XXXS || type == kIQ4_NL || + 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) @@ -757,12 +782,14 @@ TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { // Both outcomes are actually exercised (a table that never keeps anything // would pass every assertion above vacuously). The kept count is // device-dependent (review #523): 17 block-capable encodings x 2 keep-capable - // GEMM roles where the device covers the CPU list; 12 x 2 on ROCm. The - // GATHER role adds 19 more (the 17, plus Q8_K and IQ3_S) on a device that - // REGISTERS the block - // gather, and nothing on a device that does not. Written as named terms - // 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; 13 x 2 on ROCm (the + // 11 Q8_K-activation and base formats, plus IQ4_NL via QUANT-GGUF-IQ4_NL + // and IQ4_XS via KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT #1940). The GATHER + // role adds 19 more (the 17, plus Q8_K and IQ3_S) on a device that + // REGISTERS the block gather, and nothing on a device that does not. + // Written as named terms 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. @@ -785,11 +812,14 @@ TEST_CASE("routing table is TOTAL: every role x every encoding is explicit") { // per-tier result: IQ3_S stays compressed in a gather table and expands to // bf16 in a GEMM, on every device. // - // QUANT-GGUF-IQ4_NL moves the ROCm GEMM term 22 -> 24 and leaves the CPU/CUDA - // term at 34, because IQ4_NL was already in the CPU list and only the DEVICE - // set was narrower. That is the shape of a device-arm port: one encoding, two - // keep-capable GEMM roles, and no change to either gather term. - const int gemm_kept = kRouteDev == vt::DeviceType::kROCM ? 24 : 34; + // QUANT-GGUF-IQ4_NL and KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940) each + // move the ROCm GEMM term 22 -> 24 (one encoding, two keep-capable GEMM + // roles apiece) and together move it 22 -> 26. Both leave the CPU/CUDA + // term at 34, because IQ4_NL and IQ4_XS were already in the CPU list and + // only the DEVICE set was narrower. Neither changes either gather term. + // IQ3_XXS was already in ROCm's Q8_K-activation set, so it does not move + // this number. + const int gemm_kept = kRouteDev == vt::DeviceType::kROCM ? 26 : 34; const int gather_kept = vt::OpRegistered(vt::OpId::kEmbeddingQuant, kRouteDev) ? 19 : 0; CHECK(kept == gemm_kept + gather_kept); diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 2a5d45575..88f730f4f 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 { @@ -2541,13 +2551,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[] = { @@ -2560,6 +2577,8 @@ TEST_CASE("non-grouped keep-quant GEMM (Q8_0/Q4_K/Q5_K/Q6_K) matches the CPU ora // Qwen3.8-Flash-Next quant stores its ffn_down_exps and its n-gram table // in, so a ROCm arm that lacks it cannot multiply those experts at all. {vt::DType::kIQ4_NL, 18, 0, -1, "iq4_nl"}, + {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 = [&] { @@ -2865,6 +2884,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 @@ -3913,11 +3961,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) @@ -3935,6 +3989,8 @@ TEST_CASE("grouped quant expert GEMM (Q8_0/Q4_K/Q6_K) matches the CPU oracle") { // provider, so the single-matrix arm alone would still throw at the first // expert forward with the model already resident. {vt::DType::kIQ4_NL, 18, 0, -1, "iq4_nl"}, + {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 @@ -4182,6 +4238,134 @@ TEST_CASE("fused MoE gate+up+SwiGLU grouped GEMM matches the CPU oracle and is N } } +// ─── KERNEL-QUANT-CIQ-GEMM-ROCM-IQUANT (#1940): the risk this row's spec ──── +// named -- CUDA needed __fmul_rn/__fadd_rn in DotIQ4XS to stay bit-exact +// against the oracle, because nvcc's default -fmad=true silently contracted +// the two-rounding accumulation into a single-rounding FMA and two of eight +// real super-blocks then disagreed by 1-4 ULP (cuda_quant_dot.cu). This +// project's CMakeLists.txt already passes -ffp-contract=off to +// $ project-wide, unlike its CXX-only reach on CUDA, so +// DotIQ4XS on ROCm uses plain `*`/`+` rather than carrying that workaround +// over unexamined (see the function's own comment). This is the gate that +// proves whether that bet paid off, over the SAME real checkpoint bytes and +// the SAME expected oracle bits the CUDA gate uses -- reused, not +// re-derived, so a drift between the two device gates cannot hide. +namespace { +// test_cuda_quant_dot.cpp's MakeDotActivation, restated rather than shared +// (no cross-file dependency): every value is an integer in [-1024, 1023] over +// 64, exact in binary32 on any compiler, so the same bytes come out of any +// build and the golden's provenance carries over unchanged. +void MakeIq4xsDotActivation(int n, uint32_t seed, float* x) { + uint32_t st = seed; + for (int i = 0; i < n; ++i) { + st = st * 1664525U + 1013904223U; + const int32_t v = static_cast((st >> 16) & 0x7ffU) - 1024; + x[i] = static_cast(v) / 64.0F; + } +} +uint32_t FloatBitsOf(float f) { + uint32_t u = 0; + std::memcpy(&u, &f, sizeof(u)); + return u; +} +float BitsFloat(uint32_t u) { + float f = 0.0F; + std::memcpy(&f, &u, sizeof(f)); + return f; +} +} // namespace + +TEST_CASE("ROCm IQ4_XS dots the ORACLE's own numbers on REAL checkpoint bytes") { + const bool any_rocm = [&] { + for (DeviceType dt : RegisteredDevices()) if (dt == DeviceType::kROCM) return true; + return false; + }(); + if (!any_rocm) { + MESSAGE("no ROCm backend on this host; ROCm IQ4_XS oracle-dot gate skipped"); + return; + } + REQUIRE_MESSAGE(OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM), + "kMatmulBTQuant must be registered on ROCm"); + + constexpr int kBlocks = 4; + constexpr int kK = 256 * kBlocks; + REQUIRE(vt::cpu::QuantTraits(DType::kIQ4_XS).vec_dot_type == DType::kQ8_K); + const size_t wbytes = std::size(vllm_test::kIq4xsGoldenBlocks); + const size_t block_bytes = wbytes / kBlocks; + REQUIRE(block_bytes * kBlocks == wbytes); + + std::vector act(kK); + MakeIq4xsDotActivation(kK, 0x4247U, act.data()); + + vt::Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + const Device gd{DeviceType::kROCM, 0}; + DevBuf da(gpu, gq, act.size()); + DevBufBytes dwt(gpu, gq, wbytes); + DevBuf dout(gpu, gq, kBlocks); + da.Upload(act); + dwt.Upload(vllm_test::kIq4xsGoldenBlocks); + + // k=256, one super-block at a time: ONE contributing lane, so BIT FOR BIT + // against the oracle's own per-block number -- no reassociation exists to + // explain a difference away. + for (int b = 0; b < kBlocks; ++b) { + CAPTURE(b); + Tensor at = Tensor::Contiguous(static_cast(da.ptr()) + b * 256 * sizeof(float), + DType::kF32, gd, {1, 256}); + Tensor wt = Tensor::Contiguous(static_cast(dwt.ptr()) + b * block_bytes, + DType::kIQ4_XS, gd, {1, 256}); + Tensor ot = Tensor::Contiguous(static_cast(dout.ptr()) + b * sizeof(float), + DType::kF32, gd, {1, 1}); + vt::MatmulBTQuant(gq, ot, at, wt); + float got = 0.0F; + gpu.Synchronize(gq); + gpu.Copy(gq, &got, static_cast(dout.ptr()) + b * sizeof(float), sizeof(float)); + gpu.Synchronize(gq); + REQUIRE(std::isfinite(got)); + CHECK(FloatBitsOf(got) == vllm_test::kIq4xsDotPerBlockBits[b]); + } + + // k=1024, all four super-blocks in one warp: the SAME __shfl_down_sync + // offsets (16,8,4,2,1) as the CUDA gate, over four live lanes, giving the + // same (v0+v2)+(v1+v3) grouping the CUDA gate's comment derives -- so the + // PRIMARY assertion is bit equality against the oracle's four numbers + // recombined in that order, and a SECONDARY one bounds the difference from + // the oracle's own sequential total by the reassociation error. + { + Tensor at = T2(da.ptr(), gd, 1, kK); + Tensor wt = Tensor::Contiguous(dwt.ptr(), DType::kIQ4_XS, gd, {1, kK}); + Tensor ot = T2(dout.ptr(), gd, 1, 1); + vt::MatmulBTQuant(gq, ot, at, wt); + float got = 0.0F; + gpu.Synchronize(gq); + gpu.Copy(gq, &got, dout.ptr(), sizeof(float)); + gpu.Synchronize(gq); + REQUIRE(std::isfinite(got)); + + const float p0 = BitsFloat(vllm_test::kIq4xsDotPerBlockBits[0]); + const float p1 = BitsFloat(vllm_test::kIq4xsDotPerBlockBits[1]); + const float p2 = BitsFloat(vllm_test::kIq4xsDotPerBlockBits[2]); + const float p3 = BitsFloat(vllm_test::kIq4xsDotPerBlockBits[3]); + const float tree = (p0 + p2) + (p1 + p3); + CAPTURE(got); + CAPTURE(tree); + CHECK(FloatBitsOf(got) == FloatBitsOf(tree)); + + const float seq = BitsFloat(vllm_test::kIq4xsDotExpectedBits); + const double mag = static_cast(std::fabs(p0)) + std::fabs(p1) + + std::fabs(p2) + std::fabs(p3); + const double bound = 4.0 * 1.1920929e-7 * mag; + const double margin = std::fabs(static_cast(got) - seq); + CAPTURE(seq); + CAPTURE(margin); + CAPTURE(bound); + CHECK(margin <= bound); + } + + gpu.DestroyQueue(gq); +} + TEST_CASE("ReshapeAndCache->PagedAttention composition matches CPU (real dims, shuffled blocks)") { // The "paged attention" case above hand-builds a contiguous KV cache; the // real model path writes it with ReshapeAndCache and reads it back. This