From f445fa8cbf58842b1db4cca7bb256ab5cb23df1e Mon Sep 17 00:00:00 2001 From: ghazni Date: Fri, 21 Aug 2026 11:29:04 +0000 Subject: [PATCH 1/2] spec(KERNEL-QUANT-CIQ-GEMM-ROCM): commit the keep-quant provider spec The ROCm backend registers no quantized-weight GEMM provider, so every GGUF k-quant weight on an AMD card computes off device. The spec scopes W1 as kMatmulBTQuant and kMatmulBTQuantGrouped providers mirroring the CUDA sibling, which the GGUF loader reaches automatically once they are registered, and owes the upstream csrc/rocm W4A16 family behind a loader consumer that does not exist yet. The issue index gains the three 2026-08-21 campaign issues: #1586 toolchain adoption and optimization, #1587 this row, #1588 the Qwen3.5-0.8B numerics characterization. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- .agents/issue-index.md | 3 + .agents/specs/kernel-quant-ciq-gemm-rocm.md | 160 ++++++++++++++++++++ 2 files changed, 163 insertions(+) create mode 100644 .agents/specs/kernel-quant-ciq-gemm-rocm.md diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 9a66f7d73..cd77064de 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -527,3 +527,6 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1563](https://github.com/mudler/vllm.cpp/issues/1563) | `GATE-SQUASH-SEPARATOR` | **A markdown `---` horizontal rule anywhere in a pull request body silently voids the trailer block, and `check-commit-trailers.py` blames the trailers instead of the framing.** Found 2026-08-21 writing the body for PR [#1550](https://github.com/mudler/vllm.cpp/pull/1550) ([#1542](https://github.com/mudler/vllm.cpp/issues/1542)). `parsed_trailers()` shells out to git's trailer parser, and **git treats a line of exactly `---` as the start of the patch section**, so everything after the first one is not part of the message and a trailer block below it is invisible. Reproduced with no repository state: a body of `subject / prose / --- / more prose / FOLLOWING_AGENTS_PROTOCOL / the three trailers` reports `[trailers] Following-Agents-Protocol must appear exactly once` and `[attribution] AI-Assisted must appear exactly once`; `sed -i '/^---$/d'` on that same file reports `OK: commit trailer contract`, and the `---` is the only difference. **The MESSAGE is the defect, not only the behaviour**: `Following-Agents-Protocol` appears EXACTLY ONCE in the body while the checker says it must appear exactly once, so a reader counts occurrences, finds one, counts again and dumps bytes before thinking to test the parser's own framing. `_strict_errors` already computes `_paragraphs(body)[-1]` correctly as the three trailers verbatim, so the checker holds the information needed to say "the trailer paragraph is present but git could not parse it; a `---` line at line N ends the message". Worse, the neighbouring `FOLLOWING_AGENTS_PROTOCOL must appear exactly once as a separate paragraph before the trailer paragraph` check stays SILENT, so the two errors that fire both point away from the cause. **Beyond one confusing message**: the repository sets `squash_merge_commit_message = PR_BODY`, so the body IS the landed commit message, and a body carrying a `---` lands a commit whose trailers `git interpret-trailers` cannot see, on a branch that is never force-pushed. Same permanent-damage shape AGENTS.md records for the `---------` separator GitHub wrote under `COMMIT_MESSAGES`, arriving from the AUTHOR side rather than the forge side. `scripts/agent-pr-body.py --pr ` DOES catch it and caught it here before the merge; the exposure is a body never passed through that command, which AGENTS.md notes is not a gate and cannot be one because it reaches the network, while the CI guard reads the frozen `pull_request` payload and so does not re-read a body edited after the final push. NOT FIXED HERE: it changes a checker's semantics and its message, so under `## Changing the rules or a checker` it needs its own row, a red-before test and green-after evidence. Two candidate repairs, neither chosen: name the `---` line, or strip patch-section framing before parsing so a markdown rule is inert -- the second changes what the contract accepts and is the larger decision. Suggested minimum: `tests/scripts/test_check_commit_trailers.py` gains a case pinning the reproduction above | bug | | [#1454](https://github.com/mudler/vllm.cpp/issues/1454) | `SPEC-MTP-GGUF` | **`test_qwen3_5_gguf_mtp.cpp` reported `Status: SUCCESS!` with `assertions: 0` on every CI run, and its one arithmetic guarantee was a tautology.** Both cases opened `if (path == nullptr) return;` on `VLLM_MTP_GGUF_MODEL`, and a bare `return` from a doctest case is a PASS: re-derived on a clean Release build at `947e5f648`, unset, the file printed `test cases: 2 \| 2 passed \| 0 failed \| 0 skipped`, `assertions: 0`, `Status: SUCCESS!`, exit 0, and printed nothing else. The variable is set nowhere in `.github/workflows/`, so that was the state of every run. Second defect in the same file: the comment at `:52` stated `num_hidden_layers + depth == block_count` and the line under it asserted `CHECK(c.num_hidden_layers > 0)`, true of every valid model. MEASURED, not argued: mutating `src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp:889` to `c.num_hidden_layers = block_count;` compiled clean and left the file at 2/2 cases, 0 assertions, `SUCCESS!`, exit 0. FIXED IN FLOW. The invariant is now pinned **HERMETICALLY** on KV-only synthetic GGUFs carrying no weight bytes, so CI checks it every run rather than never - 65/1 (the shipped Qwen3.8-27B pair), 25/1 (the Qwen3.5-2B reference this suite was developed against) and 28/3, the third arm separating `- nextn` from `- 1` - plus a head-less arm asserting the key is NOT published, which is the half `NumMtpLayers` cannot express because it answers 1 for an absent key. The two env-gated cases stay, now skipping with a `MESSAGE` naming the variable as `test_gguf_mmproj_reach.cpp` does, and the live one re-derives the invariant from the file's own `block_count` kv. Unset 4 cases / 18 assertions / `SUCCESS!` / rc 0; live on `Qwen3.8-27B-Q4_K_M.gguf` 4 / 38 / `SUCCESS!` / rc 0. Both mutants now red (9/18 and 5/18, exit 1), compiled clean, restored against a pre-taken sha256. **The production line is CORRECT and was not touched**: `block_count - nextn` landed `1a4db5c3c`, the `mtp_num_hidden_layers` republication `493327b4e`. Related but distinct: [#821](https://github.com/mudler/vllm.cpp/issues/821) W2 (`0adeb8b0e`) pins the same arithmetic for the 27B artifact on a committed manifest in `tests/vllm/models/test_qwen38_27b_gguf_manifest.cpp`, and that gate DOES catch both mutants - so the invariant was not globally unpinned, it was unpinned in this row's own file | bug | | [#1434](https://github.com/mudler/vllm.cpp/issues/1434) | `GATE-DOC-CHECKPOINT-STATES` | **`scripts/check-doc-checkpoint.py` could not see `PARTIAL`, so 118 state cells could move with no gate observing them.** `STATES` (`:56-66`) is the whole definition of what a lifecycle state IS for the gate that enforces AGENTS.md's `docs/STATUS.md` / `docs/BENCHMARKS.md` / spec `## Now` triple, and `row_states` drops any row it cannot match. `lifecycle_moves` and `moved_rows` then iterate the AFTER map, so leaving the matched set is silent by construction. Re-derived at `947e5f648` (the report measured `63d87805c`): `PARTIAL` **118** cells and `ANCHOR-BACKFILL` **73**, against `DONE` 77 and `BLOCKED` 9 — `PARTIAL` is the second most used state in the matrices and the gate was blind to it. Over the seven tables `ROW_TABLES` actually reads, the resolved population goes from **153 rows to 226**, a 47.7 % widening. Two of the transitions the report names behave differently from its description, measured with scratch commits at `947e5f648` on an unmodified checker: `READY -> PARTIAL` rc **0** and `PARTIAL -> READY` rc **0** are the real blind spots, while the report's suggested `PARTIAL -> ACTIVE` already reds — by accident, reporting **`added as ACTIVE`** for a row that has existed for months, because it is absent from the BEFORE map. FIXED IN FLOW for `PARTIAL` only. **`ANCHOR-BACKFILL` is deliberately excluded**: `.agents/feature-matrix.md:14-17` defines it as a property of the RECORD (*a legacy implemented row without exact code, test and real-spec anchors*), `docs/STATUS.md` carries no such term and would have nothing true to write on a `DONE <-> ANCHOR-BACKFILL` move, and `REQUIRED["lifecycle"]` cannot demand the spec's `## Now` alone — so admitting it would demand a public-document edit with nothing to say, which is the exact shape `check-doc-checkpoint.py:4-17` records as the reason the file was rewritten (16 of 20 red CI runs, six hardcoded escape hatches). One row's resolved state moves and the move is a REPAIR: `KV-BLOCK-POOL` says `` `PARTIAL` (not `DONE`) `` in its prose and the last-match heuristic believed the parenthesis, resolving `DONE`. No pinned counter moves — `check-gate-commands.py` has its own `GATED_STATES` and `RUNNABLE_BASELINE` is keyed on matrix rows, `UNOWNED_HIGH_WATER` is unmoved because this row names an owner, and no matrix row or public document changes — which was measured, not assumed, because this is the [#1376](https://github.com/mudler/vllm.cpp/issues/1376) ratchet shape. Remainder listed under `## Owed` in [doc-checkpoint-lifecycle-states.md](specs/doc-checkpoint-lifecycle-states.md): `ANCHOR-BACKFILL` moves, `.agents/sglang-matrix.md` never entering `ROW_TABLES`, a row that leaves the matched set entirely, and a new row added directly as `PARTIAL` | bug | +| [#1586](https://github.com/mudler/vllm.cpp/issues/1586) | `BACKEND-ROCM` | Adopt the ROCm 7.14 container toolchain (first TheRock production release, gfx1100 supported) and open the gfx1100 optimization campaign; baseline recorded in the issue: build 586/586 green on `rocm-dev:7.14.0`, focused gate 4/5 with the `MoeSiluMul` bf16 exactness failure at `tests/vt/test_backend_cross_device.cpp:2063` | feature | +| [#1587](https://github.com/mudler/vllm.cpp/issues/1587) | `KERNEL-QUANT-CIQ-GEMM-ROCM` | The ROCm backend has no quantized-weight GEMM provider, so every GGUF k-quant on an AMD card computes off device; upstream pins a `csrc/rocm` W4A16 family (`gptq_gemm_rdna3` gated `VLLM_ROCM_GFX1100`) we cannot reach yet because no GPTQ consumer exists. W1 registers the `kMatmulBTQuant`/`kMatmulBTQuantGrouped` keep-quant providers mirroring the CUDA sibling; W2 (owed) ports the upstream family behind a loader consumer. Spec [`kernel-quant-ciq-gemm-rocm.md`](specs/kernel-quant-ciq-gemm-rocm.md) | feature | +| [#1588](https://github.com/mudler/vllm.cpp/issues/1588) | `BACKEND-ROCM` | Characterize Qwen3.5-0.8B CPU against ROCm numerics on gfx1100: the backend matrix records the all-native run with its numerical characterization open; owns the `MoeSiluMul` bf16 exactness failure found in the 7.14 baseline | verification | diff --git a/.agents/specs/kernel-quant-ciq-gemm-rocm.md b/.agents/specs/kernel-quant-ciq-gemm-rocm.md new file mode 100644 index 000000000..8d29933cd --- /dev/null +++ b/.agents/specs/kernel-quant-ciq-gemm-rocm.md @@ -0,0 +1,160 @@ +# KERNEL-QUANT-CIQ-GEMM-ROCM — keep-quant GEMM providers on kROCM + +- Issue: [#1587](https://github.com/mudler/vllm.cpp/issues/1587) +- Base: `e2a9e035d` (upstream/main) +- State at commit: `SPIKE` accepted; W1 implementation rides this pull request +- Pull request shape: one pull request for spec and implementation + (developer decision 2026-08-21) + +## Scope + +The ROCm backend registers roughly 44 ops and has no quantized-weight GEMM +provider. A search for `MatmulBTQuant`, `kMatmulBTQuant`, and `vec_dot` over +`src/vt/rocm/` and `include/vt/` returns nothing. Every GGUF k-quant weight +on an AMD card therefore computes off device today. + +Two waves, one row: + +1. **W1 (this change).** `kROCM` providers for `OpId::kMatmulBTQuant` and + `OpId::kMatmulBTQuantGrouped`: the GGUF Q8_K-family keep-quant GEMM, + mirroring the CUDA sibling's contract. Registering the provider flips + `GgufQuantComputeAvailable()` true on the platform, so every GGUF + k-quant model reaches it with zero model-code edits. +2. **W2 (owed, see `## Owed`).** The upstream `csrc/rocm` W4A16 GPTQ/AWQ + family (`wvSplitK_int4_g`, `gptq_gemm_rdna3`, + `gptq_gemm_rdna3_wmma`, `moe_gptq_gemm_rdna3`). This project cannot + reach those kernels yet: GPTQ and AWQ checkpoints have only a host + dequant path (`awq_gptq_dequant.cpp`) and no W4A16 consumer. Porting + them before a consumer exists would land dead code. + +Out of scope: FP8 on gfx1100 (upstream refuses it on this arch; +`supports_fp8()` is gfx9 or gfx12x only), Triton-on-ROCm families, and any +loader work for AWQ/GPTQ checkpoints. + +## Upstream anchors + +Pinned vLLM `555967922`: + +- `csrc/rocm/torch_bindings.cpp` names the whole HIP quant-GEMM surface: + `LLMM1`, `wvSplitK`, `wvSplitKrc`, `wvSplitK_int4_g`, `wvSplitKQ`, and + the `VLLM_ROCM_GFX1100`-gated `gptq_gemm_rdna3`, + `gptq_gemm_rdna3_wmma`, `moe_gptq_gemm_rdna3`. +- `csrc/rocm/q_gemm_rdna3.cu:1-40` (header) records the RDNA3 hardware + facts W2 inherits: wave32 geometry, no native packed fp16/bf16 atomic + add (emulated with `global_atomic_cmpswap_b64`), `v_dot2_f32_f16` for + fp16, fp32-widened accumulate for bf16, and the WMMA forward at + `M >= 16`. +- `vllm/platforms/rocm.py` `supports_fp8` excludes gfx1100. + +Classification per `.agents/porting.md`: W1 has **no upstream +counterpart** — vLLM has no GGUF keep-quant device path anywhere. It is +derived from our own CUDA sibling plus the ggml CPU reference semantics, +and it is recorded as such in `porting-inventory.md` section 9. W2 is a +1:1 port of the pinned files. + +## Local anchors + +- `include/vt/ops.h:176` `kMatmulBTQuant`; `:184` `kMatmulBTQuantGrouped`; + `:1629` the `MatmulBTQuant` entry signature. +- `src/vt/ops.cpp:186-211` validation and dispatch through + `GetOp(OpId::kMatmulBTQuant, q.device.type)`. +- `src/vt/cpu/cpu_quant_gemm.cpp:302-310` the CPU registrar — the exact + oracle. +- `src/vt/cuda/cuda_quant_dot.cu:1-18` the oracle chain (kernel wiring, + per-block dot, activation quant); `:1814` the provider; `:1990-1993` + the registrar whose registration flips the loader default. +- `src/vllm/model_executor/model_loader/gguf_keep_quant.cpp:75-78` + `GgufQuantComputeAvailable()` — the reachability flip. +- `src/vt/rocm/rocm_ops.hip:101` the kROCM registration pattern; + `src/vt/rocm/rocm_backend.hip:328` the unified-memory bool that decides + CPU-reference fallthrough on APUs versus discrete cards. + +## Design + +W1 adds `src/vt/rocm/rocm_quant_dot.hip`, structured like +`cuda_quant_dot.cu`: + +- Quantize each activation row to `Q8_K` on the device, then run an + integer dot against the compressed weight blocks per output element. + Integer arithmetic is exact, so the provider gates **bit-exact** + against the CPU provider — the same bar the CUDA sibling meets. +- First-wave weight types: `Q4_K`, `Q5_K`, `Q6_K`, `Q2_K`, `Q3_K`. The + IQ codebook types (`IQ2_XXS`, `IQ3_XXS`) join when their tables port + cleanly; `Q8_0`/`Q4_0` activations fall back to the CPU provider over a + drained queue, exactly as `cuda_quant_dot.cu:1834` does. +- Geometry sized for wave32 on gfx1100; the CUDA warp-per-output shape + carries over with wave-size adjustments. Geometry is a performance + concern only; correctness comes from the exact integer core. +- Memory: gfx1100 is a discrete card, so weight blocks must be + device-resident. The unified-memory assumptions in the model paths do + not hold here. The provider requires device pointers and relies on the + existing weight-staging path; the implementer verifies + `needs_weight_staging()` reports true for `kROCM` so loaders stage + blocks once. If staging needs model-path edits beyond the platform + seam, that is a stop condition below. +- Registration follows `rocm_ops.hip:101`. From registration onward, + `GgufQuantComputeAvailable()` is true on `kROCM` and the GGUF loader + routes keep-quant towers to the device. A rollback env kill switch + mirrors whichever flag the CUDA side exposes. + +## Risks + +- R1: weight residency on a discrete card. The DeepSeek-V4 and Qwen3.5 + paths stage weights through `ResidentWeight` gated on + `needs_weight_staging()`; if that predicate is CUDA-only, W1 grows the + platform-seam fix and says so. +- R2: wave32 geometry differences make the first build slower than the + CPU tier on some shapes. That is a recorded measurement, not a + correctness failure; the provider stays default-on only if it wins or + ties, else ships behind the kill switch with the numbers in the spec. +- R3: hipcc `-O0` device code starts a hostcall listener that can deadlock + at exit (#132). Builds set a `CMAKE_BUILD_TYPE`; the container baseline + uses Release. +- R4: IQ codebook tables grow `.rodata`; deferring them keeps W1 small. + +## Tests + +Red-first, in the same change: + +1. Extend the quant-dot operator tests with `kROCM` arms: per-type + bit-exact equality against the CPU provider on random and boundary + inputs, the `K % 256` refusal, the grouped variant's expert-index + contract, and the unsupported-dtype CPU-fallback arm. Capture the red + before the provider exists (`OpRegistered(kMatmulBTQuant, kROCM)` is + false and keep-quant stays off). +2. Focused gate: `ctest -R 'rocm|cross_device|quant'` inside the + `rocm-dev:7.14.0` container under the host GPU mutex. +3. Model-level smoke: one small GGUF checkpoint (Qwen3.5-0.8B Q4_K_M, + fetched under the recorded authority) decodes end to end on gfx1100 + with keep-quant routed to the device, token-identical to the same + build forced onto the CPU provider. + +## Gates + +Correctness gate: bit-exact versus the CPU `kMatmulBTQuant` provider on +the declared types, plus the model smoke above. The pinned-vLLM ROCm +oracle does not cover GGUF keep-quant (`BACKEND-GATE-ROCM-VLLM` stays +`INVENTORIED`), so vLLM parity for this wave is out of reach by +construction and said so. Performance axes are measured and recorded; no +throughput floor is claimed in W1. + +## Evidence + +- Container baseline on 7.14: build 586 of 586 targets green; focused + gate 4 of 5 with the `MoeSiluMul` bf16 exactness failure recorded on + [#1586](https://github.com/mudler/vllm.cpp/issues/1586). +- This row appends its measurements to `## Outcome` when it reaches DONE. + +## Stop conditions + +- `NEEDS_DECISION`: weight staging requires edits to model forward paths + rather than the platform seam. +- Stop and report if bit-exactness cannot be reached; the integer-dot + premise would be violated, which means the port is wrong somewhere. + +## Owed + +- W2: the upstream `csrc/rocm` W4A16 family port together with the loader + consumer that makes it reachable. Stays owed unless it lands in this + pull request. +- `porting-inventory.md` section 9 entry for the W1 derivation. From 6236e9e55eb49e441e74fd130c0a8c56c86cef8c Mon Sep 17 00:00:00 2001 From: ghazni Date: Fri, 21 Aug 2026 19:52:08 +0000 Subject: [PATCH 2/2] feat(KERNEL-QUANT-CIQ-GEMM-ROCM): land the W1 keep-quant providers on kROCM MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The GGUF loader routes a block-typed weight to MatmulBTQuant whenever the running device has the provider, so registering these two ops lights up keep-quant compute on every ROCm board with no model-path change: the dense and grouped MoE towers stage once through ResidentWeight and dispatch to the new device GEMM. Coverage mirrors the CUDA sibling exactly — the ten Q8_K-family encodings plus a native Q8_0 arm. The integer dots are the portable scalar forms of the CPU reference bodies in the CPU accumulation order, because gfx1100 exposes no signed byte dot (v_dot4_i32_iu8 is unsigned-only; sdot4 needs a feature this target does not offer), and the gate is bit-exactness against the CPU tier at NMSE 1e-6 with the f64 dequant band at 5e-4. Unsupported dtypes throw naming the dtype instead of silently falling back to a host kernel that cannot follow device pointers; VT_GGUF_KEEP_QUANT=0 restores load-time expansion. Gates on gfx1100 / ROCm 7.14.0: test_rocm_quant_dot 132,094 assertions green across all ten encodings (decode through prefill shapes, broadcast and per-row grouped arms over a poisoned output buffer), focused ctest 'rocm|cross_device|quant' 20/21 with only the pre-existing MoeSiluMul bf16 exactness failure (#1588) remaining, and an end-to-end Qwen3.5-0.8B Q4_K_M decode that is deterministic on device. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- CMakeLists.txt | 4 +- src/vt/rocm/rocm_quant_dot.hip | 1044 ++++++++++++++++++++++++++++++ tests/CMakeLists.txt | 6 + tests/vt/test_rocm_quant_dot.cpp | 338 ++++++++++ 4 files changed, 1391 insertions(+), 1 deletion(-) create mode 100644 src/vt/rocm/rocm_quant_dot.hip create mode 100644 tests/vt/test_rocm_quant_dot.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index e649fa265..8277e5f8d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1588,7 +1588,8 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gdn_postconv.hip src/vt/rocm/rocm_gdn_scan.hip src/vt/rocm/rocm_gdn_fused.hip - src/vt/rocm/rocm_ops.hip) + src/vt/rocm/rocm_ops.hip + src/vt/rocm/rocm_quant_dot.hip) if(VLLM_CPP_HIP_ARCHITECTURES) set_source_files_properties( src/vt/rocm/rocm_backend.hip @@ -1609,6 +1610,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gdn_scan.hip src/vt/rocm/rocm_gdn_fused.hip src/vt/rocm/rocm_ops.hip + src/vt/rocm/rocm_quant_dot.hip PROPERTIES HIP_ARCHITECTURES "${VLLM_CPP_HIP_ARCHITECTURES}") endif() # Prefer the absolute path inside ${ROCM_PATH}/lib, fall back to the bare name, diff --git a/src/vt/rocm/rocm_quant_dot.hip b/src/vt/rocm/rocm_quant_dot.hip new file mode 100644 index 000000000..fdf11f138 --- /dev/null +++ b/src/vt/rocm/rocm_quant_dot.hip @@ -0,0 +1,1044 @@ +// ROCm keep-quant GGUF k-quant GEMM (KERNEL-QUANT-CIQ-GEMM-ROCM W1) — the +// kROCM provider for `OpId::kMatmulBTQuant` and `OpId::kMatmulBTQuantGrouped`. +// +// Port of src/vt/cuda/cuda_quant_dot.cu (the kCUDA provider), which is itself +// a port of the CPU oracle: +// src/vt/cpu/cpu_quant_gemm.cpp MatmulBTQuantKernel (the GEMM wiring) +// src/vt/cpu/cpu_quant_dot.cpp VecDot{Q2_K,Q3_K,Q4_K,Q5_K,Q6_K, +// IQ2_XXS,IQ3_XXS,IQ2_S,IQ1_S,IQ1_XXXS}Q8_K +// src/vt/cpu/cpu_quant_act.cpp QuantizeRowQ8_K (the activation quant) +// The device numeric helpers below are the SAME bit-exact ports of +// src/vt/dtype.cpp + cpu_quant_act.cpp, so the Q8_K activation bytes — and +// therefore the whole INTEGER dot — are IDENTICAL to the CPU reference. Only +// the per-super-block float scale sum is reassociated (warp reduction vs the +// CPU's sequential add), so the gate is the CUDA sibling's gate: INTEGER core +// bit-exact, final scale within the NMSE band test_ops_quant_dot uses. +// +// * Scratch uses hipMallocAsync (present in ROCm 7.14) with the same retire +// -never-free discipline as the CUDA side (graph_safe_scratch.h). +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "vt/cpu/cpu_quant_blocks.h" // vt::cpu::Block* struct mirror + // (single source; plain C++) +#include "vt/cuda/cuda_quant_iq_tables.cuh" // d_iq2xxs_grid / d_iq3xxs_grid / + // d_iq2s_grid / d_iq1s_grid / + // d_iq1xxxs_grid / d_ksigns / + // d_kmask (single source; pure + // __device__ syntax, no CUDA) +#include "vt/cuda/graph_safe_scratch.h" // RetireGraphScratch (portable) +#include "vt/ops.h" +#include "vt/quant.h" + +namespace vt::rocm { +namespace { + +// dtype.cpp F32ToF16 — round-to-nearest-even, subnormals, inf/nan. Used only +// for the Q8_0 activation scale (the CPU Q8_0 vec_dot's f16 round-trip). +__device__ inline uint16_t DF32ToF16(float f) { + uint32_t u = __float_as_uint(f); + uint16_t sign = static_cast((u >> 16) & 0x8000); + int32_t exp = static_cast((u >> 23) & 0xFF) - 127 + 15; + uint32_t mant = u & 0x7FFFFF; + if (((u >> 23) & 0xFF) == 0xFF) + return static_cast(sign | 0x7C00 | (mant ? 0x200 | (mant >> 13) : 0)); + if (exp >= 0x1F) return static_cast(sign | 0x7C00); + if (exp <= 0) { + if (exp < -10) return sign; + mant |= 0x800000; + uint32_t shift = static_cast(14 - exp); + uint32_t half = mant >> shift; + uint32_t rem = mant & ((1u << shift) - 1); + uint32_t mid = 1u << (shift - 1); + if (rem > mid || (rem == mid && (half & 1))) ++half; + return static_cast(sign | half); + } + uint32_t half = static_cast(exp << 10) | (mant >> 13); + uint32_t rem = mant & 0x1FFF; + if (rem > 0x1000 || (rem == 0x1000 && (half & 1))) ++half; + return static_cast(sign | half); +} + +// Load one activation element (dtype-decoded, exactly like cpu LoadActF32). + +using vt::cpu::BlockIQ1_S; +using vt::cpu::BlockIQ1_XXXS; +using vt::cpu::BlockIQ2_S; +using vt::cpu::BlockIQ2_XXS; +using vt::cpu::BlockIQ3_XXS; +using vt::cpu::BlockQ2_K; +using vt::cpu::BlockQ3_K; +using vt::cpu::BlockQ4_K; +using vt::cpu::BlockQ5_K; +using vt::cpu::BlockQ6_K; +using vt::cpu::BlockQ8_K; +using vt::cpu::BlockQ8_0; +using vt::cpu::kQK_K; +using vt::cpu::kQK8_0; + +void CheckHip(hipError_t err, const char* what) { + if (err != hipSuccess) { + throw std::runtime_error(std::string("vt rocm: matmul_bt_quant: ") + what + + ": " + hipGetErrorString(err)); + } +} + +// --- device numeric helpers — bit-exact ports of src/vt/dtype.cpp ------------- +__device__ inline float DF16ToF32(uint16_t h) { + uint32_t sign = static_cast(h & 0x8000) << 16; + uint32_t exp = (h >> 10) & 0x1F; + uint32_t mant = h & 0x3FF; + if (exp == 0x1F) return __int_as_float(sign | 0x7F800000 | (mant << 13)); + if (exp == 0) { + if (mant == 0) return __int_as_float(sign); + int shift = 0; + while ((mant & 0x400) == 0) { + mant <<= 1; + ++shift; + } + mant &= 0x3FF; + return __int_as_float(sign | ((113 - shift) << 23) | (mant << 13)); + } + return __int_as_float(sign | ((exp + 112) << 23) | (mant << 13)); +} + +__device__ inline float DBF16ToF32(uint16_t b) { + return __int_as_float(static_cast(b) << 16); +} + +__device__ inline uint16_t DF32ToBF16(float f) { + uint32_t u = __float_as_int(f); + if ((u & 0x7F800000) == 0x7F800000 && (u & 0x7FFFFF)) { + return static_cast((u >> 16) | 0x0040); + } + uint32_t rounding = 0x7FFF + ((u >> 16) & 1); + return static_cast((u + rounding) >> 16); +} + +// cpu_quant_act.cpp NearestInt (ggml-quants.c:563) — magic-constant round-to-even. +__device__ inline int DNearestInt(float fval) { + float val = fval + 12582912.0f; + int i = __float_as_int(val); + return (i & 0x007fffff) - 0x00400000; +} + +enum class ActDT : int { kF32 = 0, kF16 = 1, kBF16 = 2 }; + +__device__ inline float DLoadAct(const void* base, ActDT dt, int64_t idx) { + switch (dt) { + case ActDT::kF32: return static_cast(base)[idx]; + case ActDT::kF16: return DF16ToF32(static_cast(base)[idx]); + default: return DBF16ToF32(static_cast(base)[idx]); + } +} + +// --------------------------------------------------------------------------- +// GPU activation quantizer — ds4-parity grid: ONE BLOCK per (super-block, row), +// 256 threads (one thread per element). Bit-exact port of QuantizeRowQ8_K with +// the argmax tie broken by LOWEST original index (== the sequential first- +// occurrence scan). Byte-identical to both CUDA quantizers by construction. +// --------------------------------------------------------------------------- +__global__ void QuantizeQ8KKernel(BlockQ8_K* __restrict__ scratch, + const void* __restrict__ a, ActDT adt, + int64_t a_rs, int64_t m, int64_t nsb) { + const int64_t b = static_cast(blockIdx.x); // super-block within row + const int64_t i = static_cast(blockIdx.y); // activation row + if (b >= nsb || i >= m) return; + const int tid = static_cast(threadIdx.x); + const int64_t elem0 = i * a_rs + b * kQK_K; + const float v = DLoadAct(a, adt, elem0 + tid); + + __shared__ float sabs[kQK_K]; + __shared__ float sval[kQK_K]; + __shared__ int sidx[kQK_K]; + sabs[tid] = fabsf(v); + sval[tid] = v; + sidx[tid] = tid; + __syncthreads(); +#pragma unroll + for (int stride = kQK_K >> 1; stride > 0; stride >>= 1) { + if (tid < stride) { + const float oa = sabs[tid + stride]; + if (oa > sabs[tid] || (oa == sabs[tid] && sidx[tid + stride] < sidx[tid])) { + sabs[tid] = oa; + sval[tid] = sval[tid + stride]; + sidx[tid] = sidx[tid + stride]; + } + } + __syncthreads(); + } + const float mx = sval[0]; + const float amax = sabs[0]; + + BlockQ8_K& y = scratch[i * nsb + b]; + if (amax == 0.0f) { + if (tid == 0) y.d = 0.0f; + y.qs[tid] = 0; + if (tid < kQK_K / 16) y.bsums[tid] = 0; + return; + } + const float iscale = -127.0f / mx; + const int qv = DNearestInt(iscale * v); + y.qs[tid] = static_cast(qv < 127 ? qv : 127); + __syncthreads(); + if (tid < kQK_K / 16) { + int sum = 0; + for (int ii = 0; ii < 16; ++ii) sum += y.qs[tid * 16 + ii]; + y.bsums[tid] = static_cast(sum); + } + if (tid == 0) y.d = 1.0f / iscale; +} + +// --------------------------------------------------------------------------- +// Per-super-block integer dots. PORTABLE SCALAR forms of the CPU reference +// bodies (cpu_quant_dot.cpp), NOT the CUDA __dp4a forms: gfx1100 has no +// signed byte dot (see the HIP DELTAS header note). Exactness is preserved by +// keeping each dot's accumulation ORDER identical to the CPU body it mirrors: +// Q2_K sums sub-blocks in k/j/l order; Q3_K/Q6_K use the 8-wide aux32 split; +// Q4_K/Q5_K accumulate per-16 bsums then per-32 scale groups in order. +// --------------------------------------------------------------------------- +__device__ inline float DotQ2K(const BlockQ2_K* xb, const BlockQ8_K* yb) { + const uint8_t* q2 = xb->qs; + const int8_t* q8 = yb->qs; + const uint8_t* sc = xb->scales; + int summs = 0; + for (int j = 0; j < 16; ++j) summs += yb->bsums[j] * (sc[j] >> 4); + const float dall = yb->d * DF16ToF32(xb->d); + const float dmin = yb->d * DF16ToF32(xb->dmin); + int isum = 0; + int is = 0; + for (int k = 0; k < kQK_K / 128; ++k) { + int shift = 0; + for (int j = 0; j < 4; ++j) { + int d = sc[is++] & 0xF; + int isuml = 0; + for (int l = 0; l < 16; ++l) isuml += q8[l] * ((q2[l] >> shift) & 3); + isum += d * isuml; + d = sc[is++] & 0xF; + isuml = 0; + for (int l = 16; l < 32; ++l) isuml += q8[l] * ((q2[l] >> shift) & 3); + isum += d * isuml; + shift += 2; + q8 += 32; + } + q2 += 32; + } + return dall * isum - dmin * summs; +} + +__device__ inline float DotQ3K(const BlockQ3_K* xb, const BlockQ8_K* yb) { + const uint32_t kmask1 = 0x03030303; + const uint32_t kmask2 = 0x0f0f0f0f; + const uint8_t* hm = xb->hmask; + const int8_t* q8 = yb->qs; + int8_t aux8[kQK_K]; + int8_t* a = aux8; + const uint8_t* q3 = xb->qs; + uint8_t m = 1; + for (int jj = 0; jj < kQK_K; jj += 128) { + for (int l = 0; l < 32; ++l) a[l] = q3[l] & 3; + for (int l = 0; l < 32; ++l) a[l] = static_cast(a[l] - ((hm[l] & m) ? 0 : 4)); + a += 32; m = static_cast(m << 1); + for (int l = 0; l < 32; ++l) a[l] = (q3[l] >> 2) & 3; + for (int l = 0; l < 32; ++l) a[l] = static_cast(a[l] - ((hm[l] & m) ? 0 : 4)); + a += 32; m = static_cast(m << 1); + for (int l = 0; l < 32; ++l) a[l] = (q3[l] >> 4) & 3; + for (int l = 0; l < 32; ++l) a[l] = static_cast(a[l] - ((hm[l] & m) ? 0 : 4)); + a += 32; m = static_cast(m << 1); + for (int l = 0; l < 32; ++l) a[l] = (q3[l] >> 6) & 3; + for (int l = 0; l < 32; ++l) a[l] = static_cast(a[l] - ((hm[l] & m) ? 0 : 4)); + a += 32; m = static_cast(m << 1); + q3 += 32; + } + uint32_t auxs[4]; + memcpy(auxs, xb->scales, 12); + const int8_t* scales = reinterpret_cast(auxs); + uint32_t tmp = auxs[2]; + auxs[2] = ((auxs[0] >> 4) & kmask2) | (((tmp >> 4) & kmask1) << 4); + auxs[3] = ((auxs[1] >> 4) & kmask2) | (((tmp >> 6) & kmask1) << 4); + auxs[0] = (auxs[0] & kmask2) | (((tmp >> 0) & kmask1) << 4); + auxs[1] = (auxs[1] & kmask2) | (((tmp >> 2) & kmask1) << 4); + a = aux8; + const int8_t* q8p = q8; + int32_t aux32[8] = {0, 0, 0, 0, 0, 0, 0, 0}; + for (int j = 0; j < kQK_K / 16; ++j) { + for (int l = 0; l < 8; ++l) aux32[l] += (scales[j] - 32) * (q8p[l] * a[l]); + q8p += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += (scales[j] - 32) * (q8p[l] * a[l]); + q8p += 8; a += 8; + } + const float d = DF16ToF32(xb->d) * yb->d; + int isum = 0; + for (int l = 0; l < 8; ++l) isum += aux32[l]; + return d * isum; +} + +__device__ inline float DotQ4K(const BlockQ4_K* xb, const BlockQ8_K* yb) { + const uint32_t kmask1 = 0x3f3f3f3f; + const uint32_t kmask2 = 0x0f0f0f0f; + const uint32_t kmask3 = 0x03030303; + const uint8_t* q4 = xb->qs; + const int8_t* q8 = yb->qs; + uint32_t utmp[4]; + memcpy(utmp, xb->scales, 12); + utmp[3] = ((utmp[2] >> 4) & kmask2) | (((utmp[1] >> 6) & kmask3) << 4); + const uint32_t uaux = utmp[1] & kmask1; + utmp[1] = (utmp[2] & kmask2) | (((utmp[0] >> 6) & kmask3) << 4); + utmp[2] = uaux; + utmp[0] &= kmask1; + const uint8_t* scales = reinterpret_cast(&utmp[0]); + const uint8_t* mins = reinterpret_cast(&utmp[2]); + int sumi = 0; + for (int j = 0; j < kQK_K / 16; ++j) sumi += yb->bsums[j] * mins[j / 2]; + // Portable nibble walk in the CPU body's order: 64-element groups decode + // low nibble then high nibble into aux8, then the per-32 scale groups. + int8_t aux8[kQK_K]; + int8_t* a = aux8; + for (int j = 0; j < kQK_K / 64; ++j) { + for (int l = 0; l < 32; ++l) a[l] = static_cast(q4[l] & 0xF); + a += 32; + for (int l = 0; l < 32; ++l) a[l] = static_cast(q4[l] >> 4); + a += 32; + q4 += 32; + } + a = aux8; + int32_t aux32[8] = {0, 0, 0, 0, 0, 0, 0, 0}; + for (int j = 0; j < kQK_K / 32; ++j) { + const int32_t scale = scales[j]; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + } + const float d = DF16ToF32(xb->d) * yb->d; + const float dmin = DF16ToF32(xb->dmin) * yb->d; + int isum = 0; + for (int l = 0; l < 8; ++l) isum += aux32[l]; + return d * isum - dmin * sumi; +} + +__device__ inline float DotQ5K(const BlockQ5_K* xb, const BlockQ8_K* yb) { + const uint32_t kmask1 = 0x3f3f3f3f; + const uint32_t kmask2 = 0x0f0f0f0f; + const uint32_t kmask3 = 0x03030303; + const uint8_t* q4 = xb->qs; + const uint8_t* hm = xb->qh; + const int8_t* q8 = yb->qs; + uint32_t utmp[4]; + memcpy(utmp, xb->scales, 12); + utmp[3] = ((utmp[2] >> 4) & kmask2) | (((utmp[1] >> 6) & kmask3) << 4); + const uint32_t uaux = utmp[1] & kmask1; + utmp[1] = (utmp[2] & kmask2) | (((utmp[0] >> 6) & kmask3) << 4); + utmp[2] = uaux; + utmp[0] &= kmask1; + const uint8_t* scales = reinterpret_cast(&utmp[0]); + const uint8_t* mins = reinterpret_cast(&utmp[2]); + int sumi = 0; + for (int j = 0; j < kQK_K / 16; ++j) sumi += yb->bsums[j] * mins[j / 2]; + int8_t aux8[kQK_K]; + int8_t* a = aux8; + uint8_t m = 1; + for (int j = 0; j < kQK_K / 64; ++j) { + for (int l = 0; l < 32; ++l) a[l] = static_cast(q4[l] & 0xF); + for (int l = 0; l < 32; ++l) + a[l] = static_cast(a[l] + ((hm[l] & m) ? 16 : 0)); + a += 32; + m = static_cast(m << 1); + for (int l = 0; l < 32; ++l) a[l] = static_cast(q4[l] >> 4); + for (int l = 0; l < 32; ++l) + a[l] = static_cast(a[l] + ((hm[l] & m) ? 16 : 0)); + a += 32; + m = static_cast(m << 1); + q4 += 32; + } + a = aux8; + int32_t aux32[8] = {0, 0, 0, 0, 0, 0, 0, 0}; + for (int j = 0; j < kQK_K / 32; ++j) { + const int32_t scale = scales[j]; + for (int r = 0; r < 4; ++r) { + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + } + } + const float d = DF16ToF32(xb->d) * yb->d; + const float dmin = DF16ToF32(xb->dmin) * yb->d; + int isum = 0; + for (int l = 0; l < 8; ++l) isum += aux32[l]; + return d * isum - dmin * sumi; +} + +__device__ inline float DotQ6K(const BlockQ6_K* xb, const BlockQ8_K* yb) { + const uint8_t* q4 = xb->ql; + const uint8_t* qh = xb->qh; + const int8_t* q8 = yb->qs; + int8_t aux8[kQK_K]; + int8_t* a = aux8; + for (int j = 0; j < kQK_K; j += 128) { + for (int l = 0; l < 32; ++l) { + a[l + 0] = static_cast( + static_cast((q4[l + 0] & 0xF) | (((qh[l] >> 0) & 3) << 4)) - 32); + a[l + 32] = static_cast( + static_cast((q4[l + 32] & 0xF) | (((qh[l] >> 2) & 3) << 4)) - 32); + a[l + 64] = static_cast( + static_cast((q4[l + 0] >> 4) | (((qh[l] >> 4) & 3) << 4)) - 32); + a[l + 96] = static_cast( + static_cast((q4[l + 32] >> 4) | (((qh[l] >> 6) & 3) << 4)) - 32); + } + a += 128; q4 += 64; qh += 32; + } + a = aux8; + const int8_t* q8p = q8; + int32_t aux32[8] = {0, 0, 0, 0, 0, 0, 0, 0}; + for (int j = 0; j < kQK_K / 16; ++j) { + const int scale = xb->scales[j]; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8p[l] * a[l]); + q8p += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8p[l] * a[l]); + q8p += 8; a += 8; + } + const float d = DF16ToF32(xb->d) * yb->d; + int isum = 0; + for (int l = 0; l < 8; ++l) isum += aux32[l]; + return d * isum; +} + +__device__ inline float DotIQ2XXS(const BlockIQ2_XXS* xb, const BlockQ8_K* yb) { + const float d = DF16ToF32(xb->d) * yb->d; + const uint16_t* qs = xb->qs; + const int8_t* q8 = yb->qs; + int32_t bsum = 0; + for (int ib32 = 0; ib32 < kQK_K / 32; ++ib32) { + uint32_t aux32[2]; + memcpy(aux32, qs + 4 * ib32, 2 * sizeof(uint32_t)); + const uint32_t ls = 2 * (aux32[1] >> 28) + 1; + int32_t sumi = 0; + for (int l = 0; l < 4; ++l) { + const uint8_t* grid = reinterpret_cast( + &vt::cuda::d_iq2xxs_grid[(aux32[0] >> (8 * l)) & 0xff]); + const uint8_t signs = + vt::cuda::d_ksigns_iq2xs[(aux32[1] >> (7 * l)) & 127]; + for (int j = 0; j < 8; ++j) + sumi += grid[j] * q8[j] * ((signs & vt::cuda::d_kmask_iq2xs[j]) ? -1 : 1); + q8 += 8; + } + bsum += sumi * static_cast(ls); + } + return d * bsum; // final *0.125 applied after the warp reduction +} + +__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 = vt::cuda::d_iq3xxs_grid[q3[2 * l + 0]]; + const uint32_t g2 = vt::cuda::d_iq3xxs_grid[q3[2 * l + 1]]; + const uint8_t signs = vt::cuda::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 & vt::cuda::d_kmask_iq2xs[j + 0]) ? -1 : 1); + sumi += b2 * q8[j + 4] * ((signs & vt::cuda::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 +} + +__device__ inline float DotIQ2S(const BlockIQ2_S* xb, const BlockQ8_K* yb) { + const float d = DF16ToF32(xb->d) * yb->d; + const int8_t* q8 = yb->qs; + const uint8_t* qs = xb->qs; + const uint8_t* qh = xb->qh; + const uint8_t* signs = qs + kQK_K / 8; + int32_t bsum = 0; + for (int ib32 = 0; ib32 < kQK_K / 32; ++ib32) { + const int ls1 = 1 + 2 * (xb->scales[ib32] & 0xf); + const int ls2 = 1 + 2 * (xb->scales[ib32] >> 4); + int sumi1 = 0; + int sumi2 = 0; + for (int l = 0; l < 2; ++l) { + const uint8_t* grid = reinterpret_cast( + &vt::cuda::d_iq2s_grid[qs[l] | ((qh[ib32] << (8 - 2 * l)) & 0x300)]); + for (int j = 0; j < 8; ++j) + sumi1 += q8[j] * grid[j] * ((signs[l] & vt::cuda::d_kmask_iq2xs[j]) ? -1 : 1); + q8 += 8; + } + for (int l = 2; l < 4; ++l) { + const uint8_t* grid = reinterpret_cast( + &vt::cuda::d_iq2s_grid[qs[l] | ((qh[ib32] << (8 - 2 * l)) & 0x300)]); + for (int j = 0; j < 8; ++j) + sumi2 += q8[j] * grid[j] * ((signs[l] & vt::cuda::d_kmask_iq2xs[j]) ? -1 : 1); + q8 += 8; + } + bsum += ls1 * sumi1 + ls2 * sumi2; + qs += 4; + signs += 4; + } + return d * bsum; // final *0.125 applied after the warp reduction +} + +__device__ inline float DotIQ1S(const BlockIQ1_S* xb, const BlockQ8_K* yb) { + const int8_t* q8 = yb->qs; + const uint8_t* qs = xb->qs; + const uint16_t* qh = xb->qh; + int32_t sumi = 0; + int32_t sumi1 = 0; + for (int ib = 0; ib < kQK_K / 32; ++ib) { + const int ls = 2 * ((qh[ib] >> 12) & 7) + 1; + const int delta = (qh[ib] & 0x8000) ? -1 : 1; + int lsum = 0; + for (int l = 0; l < 4; ++l) { + const int8_t* grid = reinterpret_cast( + &vt::cuda::d_iq1s_grid[qs[l] | (((qh[ib] >> (3 * l)) & 7) << 8)]); + for (int j = 0; j < 8; ++j) lsum += q8[j] * grid[j]; + q8 += 8; + } + sumi += ls * lsum; + sumi1 += ls * delta * (yb->bsums[2 * ib + 0] + yb->bsums[2 * ib + 1]); + qs += 4; + } + return DF16ToF32(xb->d) * yb->d * + (static_cast(sumi) + 0.125f * static_cast(sumi1)); +} + +__device__ inline float DotIQ1XXXS(const BlockIQ1_XXXS* xb, const BlockQ8_K* yb) { + const int8_t* q8 = yb->qs; + const uint8_t* qs = xb->qs; + const uint8_t* sc = xb->sc; + int32_t sumi = 0; + int32_t sumi1 = 0; + for (int ib = 0; ib < kQK_K / 32; ++ib) { + const int nib = (sc[ib / 2] >> (4 * (ib & 1))) & 0xf; + const int ls = 2 * (nib & 7) + 1; + const int delta = (nib & 8) ? -1 : 1; + int lsum = 0; + for (int l = 0; l < 4; ++l) { + const int8_t* grid = + reinterpret_cast(&vt::cuda::d_iq1xxxs_grid[qs[l]]); + for (int j = 0; j < 8; ++j) lsum += q8[j] * grid[j]; + q8 += 8; + } + sumi += ls * lsum; + sumi1 += ls * delta * (yb->bsums[2 * ib + 0] + yb->bsums[2 * ib + 1]); + qs += 4; + } + return DF16ToF32(xb->d) * yb->d * + (static_cast(sumi) + 0.125f * static_cast(sumi1)); +} + +// --------------------------------------------------------------------------- +// WType tags + DotSuperblock dispatch — mirrors the CUDA file's table. +// --------------------------------------------------------------------------- +enum class WType : int { + kIQ2_XXS = 0, + kIQ3_XXS = 1, + kQ2_K = 2, + kQ3_K = 3, + kQ4_K = 4, + kQ5_K = 5, + kQ6_K = 6, + kIQ2_S = 7, + kIQ1_S = 8, + kIQ1_XXXS = 9, +}; + +template +__device__ inline float DotSuperblock(const void* w_sb, const BlockQ8_K* a_sb); + +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ2XXS(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ3XXS(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ2K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ3K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ4K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ5K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ6K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ2S(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ1S(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ1XXXS(static_cast(w), a); +} + +template +__device__ constexpr float FinalFactor() { + return (W == WType::kIQ2_XXS || W == WType::kIQ2_S) + ? 0.125f + : (W == WType::kIQ3_XXS ? 0.25f : 1.0f); +} + +// --------------------------------------------------------------------------- +// The MMVQ-style GEMM: one WARP per output element (i,j). Lanes split the K +// super-blocks; the warp reduction sums the partials. HIP delta: the shuffle +// mask is 64-bit on this target. Determinism note unchanged from CUDA: the +// integer core is exact; only the scale sum reassociates (within NMSE). +// --------------------------------------------------------------------------- +template +__global__ void QuantDotGemmKernel(OutT* __restrict__ out, + const uint8_t* __restrict__ weight, + const BlockQ8_K* __restrict__ act, int64_t m, + int64_t n, int64_t nsb, size_t w_row_bytes, + size_t w_block_bytes) { + const int64_t warp = static_cast(blockIdx.x) * (blockDim.x >> 5) + + (threadIdx.x >> 5); + if (warp >= m * n) return; + const int64_t i = warp / n; + const int64_t j = warp % n; + const int lane = threadIdx.x & 31; + + const uint8_t* w_row = weight + static_cast(j) * w_row_bytes; + const BlockQ8_K* a_row = act + i * nsb; + + float partial = 0.0f; + for (int64_t sb = lane; sb < nsb; sb += 32) { + const void* w_sb = w_row + static_cast(sb) * w_block_bytes; + partial += DotSuperblock(w_sb, a_row + sb); + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) + partial += __shfl_down_sync(0xffffffffffffffffull, partial, off); + + if (lane == 0) { + const float v = FinalFactor() * partial; + if constexpr (sizeof(OutT) == 4) { + out[i * n + j] = v; + } else { + out[i * n + j] = DF32ToBF16(v); + } + } +} + +// GROUPED variant: warp per (p, n); weight row selected by expert_ids[p]. +template +__global__ void QuantDotGemmGroupedKernel(OutT* __restrict__ out, + const uint8_t* __restrict__ weight, + const BlockQ8_K* __restrict__ act, + const int32_t* __restrict__ expert_ids, + int64_t P, int64_t n, int64_t nsb, + size_t w_row_bytes, + size_t w_block_bytes, bool bcast) { + const int64_t warp = static_cast(blockIdx.x) * (blockDim.x >> 5) + + (threadIdx.x >> 5); + if (warp >= P * n) return; + const int64_t p = warp / n; + const int64_t j = warp % n; + const int lane = threadIdx.x & 31; + + const int64_t e = expert_ids[p]; + const uint8_t* w_row = weight + static_cast(e * n + j) * w_row_bytes; + // Broadcast activation: the routed gate/up share ONE quantized hidden. + const BlockQ8_K* a_row = act + (bcast ? 0 : p) * nsb; + + float partial = 0.0f; + for (int64_t sb = lane; sb < nsb; sb += 32) { + const void* w_sb = w_row + static_cast(sb) * w_block_bytes; + partial += DotSuperblock(w_sb, a_row + sb); + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) + partial += __shfl_down_sync(0xffffffffffffffffull, partial, off); + + if (lane == 0) { + const float v = FinalFactor() * partial; + if constexpr (sizeof(OutT) == 4) { + out[p * n + j] = v; + } else { + out[p * n + j] = DF32ToBF16(v); + } + } +} + +// --------------------------------------------------------------------------- +// Host wiring: scratch, launches, providers, registrar. +// --------------------------------------------------------------------------- +struct StreamScratch { + void* buf = nullptr; + size_t bytes = 0; +}; + +std::mutex& ScratchMutex() { + static std::mutex mu; + return mu; +} + +StreamScratch& ScratchFor(hipStream_t s) { + static std::unordered_map map; + return map[s]; +} + +void CheckHipLaunch(const char* what) { CheckHip(hipGetLastError(), what); } + +void* EnsureScratch(size_t need, hipStream_t s) { + std::lock_guard lock(ScratchMutex()); + StreamScratch& sc = ScratchFor(s); + if (need > sc.bytes) { + // Retire (never free): a captured hipGraph may have baked this pointer. + vt::cuda::RetireGraphScratch(sc.buf); + CheckHip(hipMallocAsync(&sc.buf, need, s), "hipMallocAsync q8_K act scratch"); + sc.bytes = need; + } + return sc.buf; +} + +inline ActDT ActDtOf(DType dt) { + return dt == DType::kF32 ? ActDT::kF32 : dt == DType::kF16 ? ActDT::kF16 : ActDT::kBF16; +} + +void LaunchQuantizeQ8K(BlockQ8_K* qact, const void* data, ActDT adt, int64_t a_rs, + int64_t rows, int64_t nsb, hipStream_t s) { + dim3 qgrid(static_cast(nsb), static_cast(rows), 1); + QuantizeQ8KKernel<<>>(qact, data, adt, a_rs, rows, nsb); + CheckHipLaunch("quantize_q8_K launch"); +} + +bool IsRocmKeepQuantSupported(DType dt, WType* out) { + switch (dt) { + case DType::kIQ2_XXS: *out = WType::kIQ2_XXS; return true; + case DType::kIQ3_XXS: *out = WType::kIQ3_XXS; return true; + case DType::kQ2_K: *out = WType::kQ2_K; return true; + case DType::kQ3_K: *out = WType::kQ3_K; return true; + case DType::kQ4_K: *out = WType::kQ4_K; return true; + case DType::kQ5_K: *out = WType::kQ5_K; return true; + case DType::kQ6_K: *out = WType::kQ6_K; return true; + case DType::kIQ2_S: *out = WType::kIQ2_S; return true; + case DType::kIQ1_S: *out = WType::kIQ1_S; return true; + case DType::kIQ1_XXXS: *out = WType::kIQ1_XXXS; return true; + // Q4_0 / Q8_0 / MXFP4 dot a Q8_0 activation and have no native arm here. + default: return false; + } +} + +template +void LaunchGemm(Tensor& out, const uint8_t* weight, const BlockQ8_K* act, + int64_t m, int64_t n, int64_t nsb, size_t w_row_bytes, + size_t w_block_bytes, hipStream_t s) { + // Wave32 geometry: one warp per output; 8 warps (256 threads) per block. + constexpr int kWarpsPerBlock = 8; + dim3 block(32 * kWarpsPerBlock, 1, 1); + const int64_t warps = m * n; + const unsigned grid = + static_cast((warps + kWarpsPerBlock - 1) / kWarpsPerBlock); + if (out.dtype == DType::kF32) { + QuantDotGemmKernel<<>>( + static_cast(out.data), weight, act, m, n, nsb, w_row_bytes, + w_block_bytes); + } else { + QuantDotGemmKernel<<>>( + static_cast(out.data), weight, act, m, n, nsb, w_row_bytes, + w_block_bytes); + } + CheckHipLaunch("matmul_bt_quant launch"); +} + +template +void LaunchGroupedGemm(Tensor& out, const uint8_t* weight, const BlockQ8_K* act, + const int32_t* expert_ids, int64_t P, int64_t n, + int64_t nsb, size_t w_row_bytes, size_t w_block_bytes, + bool bcast, hipStream_t s) { + constexpr int kWarpsPerBlock = 8; + dim3 block(32 * kWarpsPerBlock, 1, 1); + const int64_t warps = P * n; + const unsigned grid = + static_cast((warps + kWarpsPerBlock - 1) / kWarpsPerBlock); + if (out.dtype == DType::kF32) { + QuantDotGemmGroupedKernel<<>>( + static_cast(out.data), weight, act, expert_ids, P, n, nsb, + w_row_bytes, w_block_bytes, bcast); + } else { + QuantDotGemmGroupedKernel<<>>( + static_cast(out.data), weight, act, expert_ids, P, n, nsb, + w_row_bytes, w_block_bytes, bcast); + } + CheckHipLaunch("matmul_bt_quant_grouped launch"); +} + + + +// Q8_0 (legacy 32-block, Q8_0-activation) arm. Self-contained: quantize the +// activation to Q8_0 on the device, then the Q8_0xQ8_0 integer dot. The dot is +// the PORTABLE SCALAR form of cpu_quant_dot.cpp VecDotQ8_0Q8_0 (quants.c:400): +// sumi0 over the first 16 elements, sumi1 over the second, ONE int32 add, then +// the f16-scale fold — accumulation order preserved for exactness. The CUDA +// file's __dp4a form is byte-identical to this on a signed-dot target; gfx1100 +// has no signed byte dot, so W1 ships the scalar body. +// --------------------------------------------------------------------------- +__global__ void QuantizeQ8_0Kernel(BlockQ8_0* __restrict__ scratch, + const void* __restrict__ a, ActDT adt, + int64_t a_rs, int64_t m, int64_t nb) { + const int64_t t = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (t >= m * nb) return; + const int64_t i = t / nb; + const int64_t b = t % nb; + const int64_t elem0 = i * a_rs + b * kQK8_0; + float amax = 0.0f; + for (int j = 0; j < kQK8_0; ++j) { + const float av = fabsf(DLoadAct(a, adt, elem0 + j)); + amax = amax > av ? amax : av; + } + BlockQ8_0& y = scratch[t]; + const float d = amax / 127.0f; + const float id = d != 0.0f ? 1.0f / d : 0.0f; + y.d = DF32ToF16(d); + for (int j = 0; j < kQK8_0; ++j) { + y.qs[j] = static_cast(roundf(DLoadAct(a, adt, elem0 + j) * id)); + } +} + +template +__global__ void QuantDotGemmQ8_0Kernel(OutT* __restrict__ out, + const uint8_t* __restrict__ weight, + const BlockQ8_0* __restrict__ act, + int64_t m, int64_t n, int64_t nb, + size_t w_row_bytes) { + const int64_t warp = static_cast(blockIdx.x) * (blockDim.x >> 5) + + (threadIdx.x >> 5); + if (warp >= m * n) return; + const int64_t i = warp / n; + const int64_t j = warp % n; + const int lane = threadIdx.x & 31; + const uint8_t* w_row = weight + static_cast(j) * w_row_bytes; + const BlockQ8_0* a_row = act + i * nb; + float partial = 0.0f; + for (int64_t b = lane; b < nb; b += 32) { + const BlockQ8_0* wb = + reinterpret_cast(w_row + static_cast(b) * + sizeof(BlockQ8_0)); + const BlockQ8_0* ab = a_row + b; + // CPU order: first half then second half, one add, then scale. + int sumi0 = 0; + for (int j2 = 0; j2 < kQK8_0 / 2; ++j2) + sumi0 += ab->qs[j2] * wb->qs[j2]; + int sumi1 = 0; + for (int j2 = kQK8_0 / 2; j2 < kQK8_0; ++j2) + sumi1 += ab->qs[j2] * wb->qs[j2]; + partial += (sumi0 + sumi1) * (DF16ToF32(wb->d) * DF16ToF32(ab->d)); + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) + partial += __shfl_down_sync(0xffffffffffffffffull, partial, off); + if (lane == 0) { + if constexpr (sizeof(OutT) == 4) + out[i * n + j] = partial; + else + out[i * n + j] = DF32ToBF16(partial); + } +} + +void MatmulQ8_0Rocm(Tensor& out, const Tensor& a, const Tensor& b, + hipStream_t s) { + const int64_t m = a.shape[0], k = a.shape[1], n = b.shape[0]; + if (m == 0 || n == 0) return; + if (k % kQK8_0 != 0) + throw std::runtime_error( + "vt rocm: matmul_bt_quant Q8_0: K must be a multiple of 32"); + const int64_t nb = k / kQK8_0; + const size_t w_row_bytes = static_cast(nb) * sizeof(BlockQ8_0); + const size_t act_bytes = + static_cast(m) * static_cast(nb) * sizeof(BlockQ8_0); + BlockQ8_0* act = static_cast(EnsureScratch(act_bytes, s)); + { + constexpr int kQBlock = 128; + const int64_t grid = (m * nb + kQBlock - 1) / kQBlock; + QuantizeQ8_0Kernel<<(grid), kQBlock, 0, s>>>( + act, a.data, ActDtOf(a.dtype), a.stride[0], m, nb); + CheckHipLaunch("quantize_q8_0 launch"); + } + constexpr int kWarpsPerBlock = 8; + dim3 block(32 * kWarpsPerBlock, 1, 1); + const unsigned grid = + static_cast((m * n + kWarpsPerBlock - 1) / kWarpsPerBlock); + if (out.dtype == DType::kF32) + QuantDotGemmQ8_0Kernel<<>>( + static_cast(out.data), static_cast(b.data), + act, m, n, nb, w_row_bytes); + else + QuantDotGemmQ8_0Kernel<<>>( + static_cast(out.data), static_cast(b.data), + act, m, n, nb, w_row_bytes); + CheckHipLaunch("matmul_bt_quant Q8_0 launch"); +} + +// The kROCM provider for OpId::kMatmulBTQuant. Validation already done by +// vt::MatmulBTQuant (ops.cpp). Contract: b is [N,K] block-quant, a [M,K] +// f32/bf16 row-packed, out [M,N]. On this DISCRETE backend an unsupported +// dtype cannot fall back to the CPU kernel (it would follow device pointers), +// so it throws naming the dtype — VT_GGUF_KEEP_QUANT=0 restores the load-time +// bf16 expansion for such files. +void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, + const Tensor& b) { + hipStream_t s = static_cast(q.handle); + const int64_t m = a.shape[0]; + const int64_t k = a.shape[1]; + const int64_t n = b.shape[0]; + if (b.dtype == DType::kQ8_0) { + MatmulQ8_0Rocm(out, a, b, s); + return; + } + + WType w{}; + if (!IsRocmKeepQuantSupported(b.dtype, &w)) { + throw std::runtime_error( + std::string("vt rocm: matmul_bt_quant: no keep-quant kernel for dtype ") + + Name(b.dtype) + + " (this discrete backend has no host fallback; set VT_GGUF_KEEP_QUANT=0 " + "to expand the file to bf16 at load)"); + } + if (k % kQK_K != 0) { + throw std::runtime_error( + "vt rocm: matmul_bt_quant: K must be a whole number of 256-element " + "Q8_K super-blocks"); + } + const int64_t nsb = k / kQK_K; + const size_t w_block_bytes = static_cast(vt::BlockBytes(b.dtype)); + const size_t w_row_bytes = static_cast(nsb) * w_block_bytes; + + const size_t act_bytes = + static_cast(m) * static_cast(nsb) * sizeof(BlockQ8_K); + BlockQ8_K* act = static_cast(EnsureScratch(act_bytes, s)); + LaunchQuantizeQ8K(act, a.data, ActDtOf(a.dtype), a.stride[0], m, nsb, s); + + const uint8_t* weight = static_cast(b.data); + switch (w) { + case WType::kIQ2_XXS: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + case WType::kIQ3_XXS: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + case WType::kQ2_K: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + case WType::kQ3_K: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + case WType::kQ4_K: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + case WType::kQ5_K: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + case WType::kQ6_K: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + case WType::kIQ2_S: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + case WType::kIQ1_S: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + case WType::kIQ1_XXXS: LaunchGemm(out, weight, act, m, n, nsb, w_row_bytes, w_block_bytes, s); break; + // IsRocmKeepQuantSupported answered yes, so a missing case must be LOUD: + // launching nothing leaves `out` stale while callers see success (#967). + default: + throw std::runtime_error( + std::string("vt rocm: matmul_bt_quant: no keep-quant kernel for dtype ") + + Name(b.dtype)); + } +} + +void MatmulBTQuantGroupedKernelRocm(Queue& q, Tensor& out, const Tensor& act, + const Tensor& weight, + const Tensor& expert_ids) { + hipStream_t s = static_cast(q.handle); + const int64_t P = out.shape[0]; + const int64_t n = out.shape[1]; + const int64_t k = act.shape[1]; + if (P == 0 || n == 0) return; + + WType w{}; + if (!IsRocmKeepQuantSupported(weight.dtype, &w)) { + throw std::runtime_error( + std::string("vt rocm: matmul_bt_quant_grouped: no keep-quant kernel for " + "dtype ") + + Name(weight.dtype) + + " (set VT_GGUF_KEEP_QUANT=0 to expand at load)"); + } + if (k % kQK_K != 0) { + throw std::runtime_error( + "vt rocm: matmul_bt_quant_grouped: K must be a whole number of " + "256-element Q8_K super-blocks"); + } + const int64_t nsb = k / kQK_K; + const size_t w_block_bytes = static_cast(vt::BlockBytes(weight.dtype)); + const size_t w_row_bytes = static_cast(nsb) * w_block_bytes; + + // Broadcast activation (preq-reuse): ONE quantized hidden feeds every routed + // expert slot; bit-identical because identical input yields identical Q8_K. + const int64_t Pa = act.shape[0]; + const bool bcast = (Pa == 1 && P > 1); + + const size_t act_bytes = + static_cast(Pa) * static_cast(nsb) * sizeof(BlockQ8_K); + BlockQ8_K* qact = static_cast(EnsureScratch(act_bytes, s)); + LaunchQuantizeQ8K(qact, act.data, ActDtOf(act.dtype), act.stride[0], Pa, nsb, + s); + + const uint8_t* wt = static_cast(weight.data); + const int32_t* eids = static_cast(expert_ids.data); + switch (w) { + case WType::kIQ2_XXS: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + case WType::kIQ3_XXS: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + case WType::kQ2_K: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + case WType::kQ3_K: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + case WType::kQ4_K: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + case WType::kQ5_K: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + case WType::kQ6_K: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + case WType::kIQ2_S: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + case WType::kIQ1_S: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + case WType::kIQ1_XXXS: LaunchGroupedGemm(out, wt, qact, eids, P, n, nsb, w_row_bytes, w_block_bytes, bcast, s); break; + default: + throw std::runtime_error( + std::string("vt rocm: matmul_bt_quant_grouped: no grouped kernel for " + "keep-quant dtype ") + + Name(weight.dtype)); + } +} + +// Registers the ROCm keep-quant GEMM during static init (table fill only, no +// HIP calls — same contract as every other registrar). This makes +// GgufQuantComputeAvailable). +struct Registrar { + Registrar() { + RegisterOp(OpId::kMatmulBTQuant, DeviceType::kROCM, + reinterpret_cast( + static_cast(&MatmulBTQuantKernelRocm))); + RegisterOp(OpId::kMatmulBTQuantGrouped, DeviceType::kROCM, + reinterpret_cast( + static_cast( + &MatmulBTQuantGroupedKernelRocm))); + } +} registrar; + +} // namespace +} // namespace vt::rocm diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 62da9ec97..5af38f101 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1734,6 +1734,12 @@ if(VLLM_CPP_HIP) # Plain C++ (no HIP header): every assertion goes through the vt:: seam. Each # case no-ops when the build has HIP but the box has no AMD GPU. vllm_cpp_add_test(test_rocm_backend vt/test_rocm_backend.cpp) +# KERNEL-QUANT-CIQ-GEMM-ROCM: the ROCm keep-quant GEMM (kROCM provider for +# kMatmulBTQuant/kMatmulBTQuantGrouped). Gates the device dequant-in-kernel dot +# against the CPU keep-quant oracle + an f64 dequant reference on the ten +# Q8_K-family encodings. Skips coherently with no AMD GPU. +vllm_cpp_add_test(test_rocm_quant_dot vt/test_rocm_quant_dot.cpp) +target_include_directories(test_rocm_quant_dot PRIVATE ${CMAKE_SOURCE_DIR}/src) # #785 P1 GPU product-seam witness. Executable only — NOT add_test. # Ordinary CTest must not see this target. Runner fail-closes on 77/nonzero. add_executable(test_ops_paged_attn_sharedk_wmma_p1_gpu diff --git a/tests/vt/test_rocm_quant_dot.cpp b/tests/vt/test_rocm_quant_dot.cpp new file mode 100644 index 000000000..c12bf04fe --- /dev/null +++ b/tests/vt/test_rocm_quant_dot.cpp @@ -0,0 +1,338 @@ +// ROCm keep-quant GEMM gate (KERNEL-QUANT-CIQ-GEMM-ROCM W1). The kROCM +// provider for `OpId::kMatmulBTQuant` / `kMatmulBTQuantGrouped` +// (src/vt/rocm/rocm_quant_dot.hip) is measured against the LANDED CPU +// keep-quant reference (src/vt/cpu/cpu_quant_gemm.cpp — the oracle) and an +// INDEPENDENT f64 dequantize-then-dot, on the ten Q8_K-family encodings the +// CUDA sibling serves (test_cuda_quant_dot.cpp's WeightCase table). +// +// THE GATE mirrors the CUDA file: the Q8_K activation quant and the whole +// INTEGER dot are bit-identical to the CPU reference by construction, so +// ROCm-vs-CPU is asserted at a TIGHT NMSE (1e-6, f32 out) — only the per- +// super-block float scale sum is reassociated (warp reduction vs the CPU's +// sequential add). ROCm-vs-f64-dequant uses the same 5e-4 band +// test_ops_quant_dot.cpp applies. A wrong codebook index / scale unpack / +// sign blows both bands (RED-first). +// +// Skips cleanly when no AMD GPU is present, so CPU-only CI stays green. +#include + +#include +#include +#include +#include +#include +#include + +#include "vt/backend.h" +#include "vt/device.h" +#include "vt/dtype.h" +#include "vt/ops.h" +#include "vt/quant.h" +#include "vt/tensor.h" + +using vt::Backend; +using vt::Device; +using vt::DeviceType; +using vt::DType; +using vt::Queue; +using vt::Tensor; + +namespace { + +constexpr double kMaxNmseErr = 5e-4; // test-backend-ops.cpp:4277 band +constexpr double kMaxNmseVsCpu = 1e-6; // integer core exact; scale sum only + +bool HasRocm() { + try { + vt::GetBackend(DeviceType::kROCM); + return true; + } catch (const std::runtime_error&) { + return false; + } +} + +Device Cpu() { return Device{DeviceType::kCPU, 0}; } +Device Gpu() { return Device{DeviceType::kROCM, 0}; } + +struct WeightCase { + DType dtype; + int64_t block_elems; + int64_t block_bytes; + int d_off; + int dmin_off; + const char* name; + // f64-dequant ceiling override (0 = kMaxNmseErr); see the CUDA table for why + // the IQ1 family needs a wider ACTIVATION-error band while the ROCm-vs-CPU + // bound below stays shared and unrelaxed. + double nmse_ref_max = 0.0; +}; + +const WeightCase kCases[] = { + {DType::kIQ2_XXS, 256, 66, 0, -1, "iq2_xxs"}, + {DType::kIQ3_XXS, 256, 98, 0, -1, "iq3_xxs"}, + {DType::kIQ2_S, 256, 82, 0, -1, "iq2_s"}, + {DType::kIQ1_S, 256, 50, 0, -1, "iq1_s", 2e-3}, + {DType::kIQ1_XXXS, 256, 38, 0, -1, "iq1_xxxs", 2e-3}, + {DType::kQ2_K, 256, 84, 80, 82, "q2_K"}, + {DType::kQ3_K, 256, 110, 108, -1, "q3_K"}, + {DType::kQ4_K, 256, 144, 0, 2, "q4_K"}, + {DType::kQ5_K, 256, 176, 0, 2, "q5_K"}, + {DType::kQ6_K, 256, 210, 208, -1, "q6_K"}, +}; + +void GenerateData(float offset, size_t n, float* dst) { + for (size_t i = 0; i < n; i++) + dst[i] = 0.1F + 2 * std::cos(static_cast(i) + offset); +} + +std::vector RandomBlocks(const WeightCase& c, int64_t nblocks, + uint32_t seed) { + std::mt19937 rng(seed); + std::vector bytes(static_cast(nblocks * c.block_bytes)); + for (uint8_t& b : bytes) b = static_cast(rng() & 0xFF); + for (int64_t i = 0; i < nblocks; ++i) { + uint8_t* blk = bytes.data() + i * c.block_bytes; + auto put_f16 = [&](int off, float v) { + const uint16_t h = vt::F32ToF16(v); + std::memcpy(blk + off, &h, sizeof(h)); + }; + const float jitter = 1.0F + 0.05F * static_cast(i % 7); + if (c.d_off >= 0) put_f16(c.d_off, 0.0125F * jitter); + if (c.dmin_off >= 0) put_f16(c.dmin_off, 0.0075F * jitter); + // IQ1 sub-block scales live INSIDE the weight (qh bits 12-14 / sc nibbles): + // narrow them to encoder-plausible values exactly as the CUDA table does. + if (c.dtype == DType::kIQ1_S) { + for (int ib = 0; ib < 8; ++ib) { + uint16_t qh = 0; + std::memcpy(&qh, blk + 34 + 2 * ib, sizeof(qh)); + const uint16_t ls = static_cast(2 + ((i + ib) % 3)); + qh = static_cast((qh & 0x8FFFU) | (ls << 12)); + std::memcpy(blk + 34 + 2 * ib, &qh, sizeof(qh)); + } + } + if (c.dtype == DType::kIQ1_XXXS) { + for (int ib = 0; ib < 8; ++ib) { + uint8_t& byte = blk[34 + ib / 2]; + const int shift = 4 * (ib & 1); + const uint8_t ls = static_cast(2 + ((i + ib) % 3)); + const uint8_t keep_sign = static_cast((byte >> shift) & 0x8); + byte = static_cast((byte & ~(0xFU << shift)) | + ((keep_sign | ls) << shift)); + } + } + } + return bytes; +} + +Tensor DevTensor(void* p, DType dt, const std::vector& shape) { + Tensor t; + t.data = p; + t.dtype = dt; + t.device = Gpu(); + t.rank = static_cast(shape.size()); + int64_t stride = 1; + for (int i = t.rank - 1; i >= 0; --i) { + t.shape[i] = shape[static_cast(i)]; + t.stride[i] = stride; + stride *= shape[static_cast(i)]; + } + return t; +} + +} // namespace + +TEST_CASE("ROCm keep-quant GEMM == CPU reference and f64 dequant (Q8_K family)") { + if (!HasRocm()) { + MESSAGE("no ROCm backend on this host; ROCm keep-quant gate skipped"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + Queue cq{Cpu(), nullptr}; + + for (const WeightCase& c : kCases) { + const int64_t k = 8 * c.block_elems; + for (int64_t m : {int64_t{1}, int64_t{4}, int64_t{32}, int64_t{512}}) { + for (int64_t n : {int64_t{1}, int64_t{7}, int64_t{16}}) { + CAPTURE(std::string(c.name)); + CAPTURE(m); + CAPTURE(k); + CAPTURE(n); + + std::vector wq = + RandomBlocks(c, n * (k / c.block_elems), 0x5EEDU); + std::vector a(static_cast(m * k)); + GenerateData(1.0F, a.size(), a.data()); + + // --- CPU oracle (the landed keep-quant kernel over host tensors) ------ + std::vector cpu_out(static_cast(m * n), 0.0F); + { + Tensor at = Tensor::Contiguous(a.data(), DType::kF32, Cpu(), {m, k}); + Tensor bt = + Tensor::Contiguous(wq.data(), DType::kF32, Cpu(), {n, k}); + bt.dtype = c.dtype; + Tensor ot = + Tensor::Contiguous(cpu_out.data(), DType::kF32, Cpu(), {m, n}); + vt::MatmulBTQuant(cq, ot, at, bt); + } + + // --- ROCm path (device tensors; discrete card, so real staging) ------ + void* d_a = gpu.Alloc(a.size() * sizeof(float)); + void* d_w = gpu.Alloc(wq.size()); + void* d_o = gpu.Alloc(static_cast(m * n) * sizeof(float)); + gpu.Copy(gq, d_a, a.data(), a.size() * sizeof(float)); + gpu.Copy(gq, d_w, wq.data(), wq.size()); + Tensor at = DevTensor(d_a, DType::kF32, {m, k}); + Tensor bt = DevTensor(d_w, c.dtype, {n, k}); + Tensor ot = DevTensor(d_o, DType::kF32, {m, n}); + vt::MatmulBTQuant(gq, ot, at, bt); + std::vector rocm_out(static_cast(m * n), 0.0F); + gpu.Copy(gq, rocm_out.data(), d_o, rocm_out.size() * sizeof(float)); + gpu.Synchronize(gq); + gpu.Free(d_a); + gpu.Free(d_w); + gpu.Free(d_o); + + // --- f64 independent reference -------------------------------------- + std::vector w(static_cast(n * k)); + vt::cpu::BlockToFloat(c.dtype)(wq.data(), w.data(), n * k); + + double num_ref = 0, den_ref = 0, num_cpu = 0, den_cpu = 0; + for (int64_t i = 0; i < m; ++i) { + for (int64_t jj = 0; jj < n; ++jj) { + double ref = 0; + for (int64_t p = 0; p < k; ++p) + ref += static_cast(a[static_cast(i * k + p)]) * + static_cast(w[static_cast(jj * k + p)]); + const double got = + rocm_out[static_cast(i * n + jj)]; + const double cpu = cpu_out[static_cast(i * n + jj)]; + num_ref += (got - ref) * (got - ref); + den_ref += ref * ref; + num_cpu += (got - cpu) * (got - cpu); + den_cpu += cpu * cpu; + REQUIRE(std::isfinite(got)); + } + } + const double nmse_ref = den_ref > 0 ? num_ref / den_ref : num_ref; + const double nmse_cpu = den_cpu > 0 ? num_cpu / den_cpu : num_cpu; + CAPTURE(nmse_ref); + CAPTURE(nmse_cpu); + const double ref_ceiling = + c.nmse_ref_max > 0 ? c.nmse_ref_max : kMaxNmseErr; + CHECK(nmse_ref <= ref_ceiling); // quantization error vs f64 dequant + CHECK(nmse_cpu <= kMaxNmseVsCpu); // matches the CPU oracle (int core exact) + } + } + } + gpu.DestroyQueue(gq); +} + +TEST_CASE("ROCm keep-quant registers the native kROCM providers") { + // The registration flips the GGUF loader's keep-quant default ON on a ROCm + // device (GgufQuantComputeAvailable -> OpRegistered(kMatmulBTQuant,kROCM)). + // Present only in a HIP build. + if (!HasRocm()) return; + CHECK(vt::OpRegistered(vt::OpId::kMatmulBTQuant, DeviceType::kROCM)); + CHECK(vt::OpRegistered(vt::OpId::kMatmulBTQuantGrouped, DeviceType::kROCM)); +} + +TEST_CASE( + "ROCm grouped keep-quant GEMM == CPU grouped golden and it WRITES the " + "output") { + if (!HasRocm()) return; + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + Queue cq{Cpu(), nullptr}; + + // All ten encodings, decode + prefill shapes, broadcast and per-row arms — + // the same matrix the CUDA grouped gate runs, over a POISONED output buffer. + struct GroupedShape { + int64_t P; + int64_t n; + int64_t E; + bool bcast; + }; + const GroupedShape kGroupedShapes[] = { + {6, 3, 4, false}, {32, 7, 8, false}, {16, 5, 2, true}}; + int64_t combos = 0; + for (const WeightCase& c : kCases) { + const int64_t k = 8 * c.block_elems; + for (const GroupedShape& g : kGroupedShapes) { + CAPTURE(std::string(c.name)); + CAPTURE(g.P); + CAPTURE(g.n); + CAPTURE(g.E); + CAPTURE(g.bcast); + const int64_t arows = g.bcast ? 1 : g.P; + std::vector wq = + RandomBlocks(c, g.E * g.n * (k / c.block_elems), 0x5EEDU); + std::vector af(static_cast(arows * k)); + GenerateData(1.0F, af.size(), af.data()); + std::vector ids(g.P); + for (int64_t p = 0; p < g.P; ++p) ids[static_cast(p)] = p % g.E; + const size_t outn = static_cast(g.P * g.n); + + // --- CPU golden (the landed grouped keep-quant kernel over host tensors) + std::vector cpu_out(outn, 1337.0F); + { + Tensor at = + Tensor::Contiguous(af.data(), DType::kF32, Cpu(), {arows, k}); + Tensor wt = + Tensor::Contiguous(wq.data(), DType::kF32, Cpu(), {g.E * g.n, k}); + wt.dtype = c.dtype; + Tensor et = + Tensor::Contiguous(ids.data(), DType::kI32, Cpu(), {g.P}); + Tensor ot = + Tensor::Contiguous(cpu_out.data(), DType::kF32, Cpu(), {g.P, g.n}); + vt::MatmulBTQuantGrouped(cq, ot, at, wt, et); + } + + // --- ROCm path over a POISONED output buffer ------------------------- + void* d_a = gpu.Alloc(af.size() * sizeof(float)); + void* d_w = gpu.Alloc(wq.size()); + void* d_e = gpu.Alloc(ids.size() * sizeof(int32_t)); + void* d_o = gpu.Alloc(outn * sizeof(float)); + std::vector poison(outn, 1337.0F); + gpu.Copy(gq, d_a, af.data(), af.size() * sizeof(float)); + gpu.Copy(gq, d_w, wq.data(), wq.size()); + gpu.Copy(gq, d_e, ids.data(), ids.size() * sizeof(int32_t)); + gpu.Copy(gq, d_o, poison.data(), poison.size() * sizeof(float)); + gpu.Synchronize(gq); + Tensor at = DevTensor(d_a, DType::kF32, {arows, k}); + Tensor wt = DevTensor(d_w, c.dtype, {g.E * g.n, k}); + Tensor et = DevTensor(d_e, DType::kI32, {g.P}); + Tensor ot = DevTensor(d_o, DType::kF32, {g.P, g.n}); + vt::MatmulBTQuantGrouped(gq, ot, at, wt, et); + std::vector got(outn, 0.0F); + gpu.Copy(gq, got.data(), d_o, got.size() * sizeof(float)); + gpu.Synchronize(gq); + gpu.Free(d_a); + gpu.Free(d_w); + gpu.Free(d_e); + gpu.Free(d_o); + + int poisoned = 0; + int nonfinite = 0; + double num = 0, den = 0; + for (size_t i = 0; i < got.size(); ++i) { + if (got[i] == 1337.0F) ++poisoned; + if (!std::isfinite(got[i])) ++nonfinite; + num += (got[i] - cpu_out[i]) * (got[i] - cpu_out[i]); + den += cpu_out[i] * cpu_out[i]; + } + const double nmse = den > 0 ? num / den : num; + CAPTURE(nmse); + CHECK(poisoned == 0); // a dispatch that launches nothing lands HERE + CHECK(nonfinite == 0); + CHECK(nmse <= kMaxNmseVsCpu); + ++combos; + } + } + // doctest prints "SUCCESS!" for a loop that never ran. Say how many it ran. + CAPTURE(combos); + CHECK(combos == + static_cast(std::size(kCases) * std::size(kGroupedShapes))); + CHECK(combos > 0); + gpu.DestroyQueue(gq); +}