From e2a9e035dbf8662f4bd87fc21a768d184f547c73 Mon Sep 17 00:00:00 2001 From: Thomas Brasser Date: Fri, 21 Aug 2026 12:54:03 +0200 Subject: [PATCH 1/9] test(BACKEND-ROCM): the block-size contract is enforced at its production call site (#1273) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Closes the "Owed" item from #1065: CheckKvCacheShape's production call site in GPUModelRunner::initialize_kv_cache now has a test. Building the runner with a non-multiple-of-16 block size (kBlockSize = 8) asserts the throw at construction, from the FLASH_ATTN backend's own get_kv_cache_shape — the executable statement of the contract the server's --block-size validation and the bench rounding exist to prevent at the entry points. Validated on gfx1151 (Strix Halo, ROCm 7.2.3) as part of the #41 M3 battery: test_runner 20/20 cases, 544 assertions, on the current-main build with #1056 + #1065 merged. Issue: #41 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: DeepSeekV4:Flash [Freebuff] --- tests/vllm/v1/worker/test_runner.cpp | 29 ++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/tests/vllm/v1/worker/test_runner.cpp b/tests/vllm/v1/worker/test_runner.cpp index d2bb97c2b..b2c94ebeb 100644 --- a/tests/vllm/v1/worker/test_runner.cpp +++ b/tests/vllm/v1/worker/test_runner.cpp @@ -1529,3 +1529,32 @@ TEST_CASE("runner: full-attention-only step skips GDN metadata build (no OOB)") doctest::Contains("qwen3_5 dense paged forward"), std::runtime_error); } + +// ─── M3: THE BLOCK-SIZE CONTRACT AT ITS PRODUCTION CALL SITE ───────────────── +// +// `CheckKvCacheShape` is well tested in isolation (test_attn_backend_registry / +// test_common_attn_metadata), but its PRODUCTION call site — the install inside +// `GPUModelRunner::initialize_kv_cache` — was not: no test drove the runner +// with a non-multiple-of-16 block size, so deleting that install left every +// gate green (the #1065 Owed item). The FLASH_ATTN backend's own +// `get_kv_cache_shape` refuses block_size % 16 != 0, and the runner resolves +// FLASH_ATTN for the CPU device, so construction must throw the backend's +// `invalid_argument` from init — the same failure the server's --block-size +// validation and the bench rounding exist to prevent at the entry points. +TEST_CASE("runner: initialize_kv_cache refuses a non-multiple-of-16 block size") { + const HfConfig c = MakeDenseOnlyConfig(); + const Qwen3_5DenseWeights w = MakeDenseOnlyWeights(c); + + KVCacheConfig kv = MakeFaOnlyKvConfig(c); + kv.kv_cache_groups[0].kv_cache_spec = std::make_shared( + /*block_size=*/8, static_cast(c.num_key_value_heads), + static_cast(c.head_dim), vllm::v1::ResolveKvCacheDType()); + + auto make_runner = [&]() { + GPUModelRunner runner(c, w, kv, Q(), /*max_num_reqs=*/8, kMaxModelLen, + /*max_num_batched_tokens=*/64); + }; + CHECK_THROWS_WITH_AS(make_runner(), + doctest::Contains("Block size must be a multiple of 16"), + std::invalid_argument); +} From f445fa8cbf58842b1db4cca7bb256ab5cb23df1e Mon Sep 17 00:00:00 2001 From: ghazni Date: Fri, 21 Aug 2026 11:29:04 +0000 Subject: [PATCH 2/9] 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 3/9] 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); +} From 2578c9b17170b41af9787f8484938948e9f9a8f3 Mon Sep 17 00:00:00 2001 From: ghazni Date: Fri, 21 Aug 2026 23:10:49 +0000 Subject: [PATCH 4/9] spec(ROCM-QUANT-GEMM-BW): commit the keep-quant bandwidth spec QuantDotGemm is 48.3% of GPU busy in the 97,721-dispatch rocprofv3 capture of Qwen3.5-4B Q4_K_M decode at 6236e9e55, and it streams weights one byte per lane, so the model decodes at ~163 GB/s effective weight streaming, about 17% of the RX 7900 XTX peak. The spec scopes an attempt ladder that vectorizes the weight loads to 16 bytes per lane and reshapes the wave geometry without touching the accumulation order, because test_rocm_quant_dot pins the integer core bit-exact against the CPU reference and stays unchanged. The issue index gains the campaign row under #1586. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- .agents/issue-index.md | 1 + .agents/specs/rocm-quant-gemm-bw.md | 110 ++++++++++++++++++++++++++++ 2 files changed, 111 insertions(+) create mode 100644 .agents/specs/rocm-quant-gemm-bw.md diff --git a/.agents/issue-index.md b/.agents/issue-index.md index cd77064de..70289a628 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -530,3 +530,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#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 | +| [#1586](https://github.com/mudler/vllm.cpp/issues/1586) | `ROCM-QUANT-GEMM-BW` | Raise Qwen3.5-4B Q4_K_M greedy-decode effective weight-streaming on gfx1100 from ~163 GB/s (17% of peak; 97,721-dispatch rocprofv3 capture at `6236e9e55` shows QuantDotGemm at 48.3% of GPU busy with 1-byte-per-lane loads) to >=576 GB/s (60%) by vectorizing the keep-quant GEMM memory path and reshaping waves, with the integer core bit-exact vs CPU (`test_rocm_quant_dot` unchanged). Spec [`rocm-quant-gemm-bw.md`](specs/rocm-quant-gemm-bw.md) | performance | diff --git a/.agents/specs/rocm-quant-gemm-bw.md b/.agents/specs/rocm-quant-gemm-bw.md new file mode 100644 index 000000000..96bc76e8d --- /dev/null +++ b/.agents/specs/rocm-quant-gemm-bw.md @@ -0,0 +1,110 @@ +# Spec: ROCM-QUANT-GEMM-BW + +## Scope + +Raise the effective weight-streaming rate of Qwen3.5-4B Q4_K_M greedy +decode on gfx1100 (RX 7900 XTX, ROCm 7.14 container) to at least 60% of +peak DRAM bandwidth (~576 GB/s of ~960), by optimizing the W1 keep-quant +GEMM (`src/vt/rocm/rocm_quant_dot.hip`) memory path and scheduling with +zero numeric change. Owned under issue #1586. Success is measured, not +argued: the fixed workload in `## Gates` must reach the rate with the +existing bit-exactness gate unchanged. + +## Upstream anchors + +- vLLM pin `555967922` (0.26.0.dev0). vLLM defines no keep-quant RDNA3 + GEMM, so behavior parity does not constrain the internals; only our + CPU reference (`src/vt/cpu/cpu_quant_dot.cpp`) pins the numerics. +- The CUDA sibling `src/vt/cuda/cuda_quant_dot.cu` is the structural + mirror. It stays untouched; any improvement found here that would also + help CUDA is recorded as owed, never ported silently into this row. + +## Baseline evidence + +`rocprofv3 -r true` capture of the gate workload at tree `6236e9e55` +(144 tokens, 97,721 dispatches, results db parsed from +`rocpd_kernel_dispatch`): + +| Fact | Value | +|---|---| +| GPU busy fraction | 0.83 | +| `QuantDotGemm*` share of busy | 48.3% | +| hipBLASLt `Cijk_*` share | 26.2% | +| GDN family share | 16.9% | +| Effective weight-streaming rate | ~163 GB/s (~17% of peak) | +| `QuantDotGemmKernel` decode geometry | grid up to 7,946,240 blocks x 256 threads, avg 1.9 ms | + +Diagnosis: the kernel occupies the chip but streams bytes narrowly. +Each lane reads one byte per super-block step; q-weight rows are walked +with lane-strided single-byte loads, so every 256-thread wavefront +touches scattered addresses and the memory system delivers far below +its burst width. + +## Design + +Attempt ladder, one attempt = change + rebuild + both-gate verify: + +1. **Vectorized weight loads.** Give each lane a contiguous 16-byte + load (`ulonglong2`) covering four lanes' worth of q-weight payload + per super-block step where the block layout allows it, keeping the + CPU accumulation order exactly (sum over nibbles/bytes in reference + sequence). Bit-exactness is preserved because reassociation is not + introduced; only the load width changes. +2. **Wave/block reshaping.** Reduce grid size by assigning each warp + multiple output elements along N; improves L2 reuse of activation + rows and drops launch count. Output mapping stays N-major within a + super-block so partial sums remain per-output. +3. **hipBLASLt algo-policy A/B** for the bf16 arms (26.2% share): + measurement-only lever from the #1586 attribution table; adopt a + pinned algo policy if a variant wins at decode shapes. +4. Optional: `VT_*` env knob parity with the CUDA side for any new + scheduling switch, defaulting to the fast path. + +Numerics guardrail for every attempt: no hardware dot instructions +(gfx1100 has no signed byte dot; recorded in the W1 spec), no change to +scale application order, no fp reassociation beyond what the reference +already fixes. + +## Risks + +- R1: Vector loads misaligned at odd N*K offsets -> guard with + alignment checks falling back to the scalar path (same kernel, + selected per-tensor, still bit-exact). +- R2: Register pressure rise kills occupancy and negates the win -> + measure VGPR count from the code object before accepting. +- R3: Prefill arms regress while decode improves -> gates measure both; + accept only when neither arm regresses beyond noise on the gate test. + +## Tests + +- `tests/vt/test_rocm_quant_dot.cpp` runs UNCHANGED as the correctness + gate: 132,094 assertions, integer core bit-exact vs CPU, NMSE <= 1e-6. +- Bandwidth gate: the workload in `## Gates`, parsed from rocprofv3 + results db, must show >= 576 GB/s steady-state decode. +- End-to-end smoke: deterministic decode across two identical runs. + +## Gates + +Fixed workload (identical to the baseline capture): + +``` +rocprofv3 -r true -- examples/vllm-cli \ + --model /models/Qwen3.5-4B-Q4_K_M.gguf --device auto \ + --temperature 0 --seed 0 --max-tokens 48 --repeat 3 \ + --prompt "Write a detailed explanation of how a transformer neural network works." +``` + +Rate = 2,740,937,888 bytes x 144 tokens / total GPU-busy seconds of +steady-state decode runs (run 1 warmup excluded). Pass at >= 576 GB/s. + +## Owed + +- Any improvement applicable to the CUDA sibling: record in the W1 + spec's owed list rather than editing `cuda_quant_dot.cu` here. +- Kernel-matrix family row updates ride the landing commit. + +## Stop conditions + +20 failed attempts without meeting the bandwidth gate: stop, report +findings and the measured ceiling hypothesis. Ambiguity needing a user +decision or an operation outside recorded authority: halt and surface. From 8e78dfa3f51fc2f12af1929efa486268e653abe1 Mon Sep 17 00:00:00 2001 From: ghazni Date: Sat, 22 Aug 2026 00:57:01 +0000 Subject: [PATCH 5/9] perf(ROCM-QUANT-GEMM-BW): split each super-block across the warp's lanes The keep-quant GEMM assigned one super-block per lane (sb = lane), so a decode-shape row with nsb = k/256 in {10..36} idled 22 of 32 lanes on every block and serialized each output on one lane. The rocprofv3 capture priced the result at 392.8 us for a [1,2560]x[9216] Q4_K GEMM (13.3 MB of weights, ~34 GB/s) and 1438.5 us for the [1,9216]x[2560] Q6_K arm (~13 GB/s). The tail path now splits every super-block across all 32 lanes (8 elements per lane, DotSuperblockWarp) and warp-reduces the integer partials; whole super-blocks beyond the first multiple of 32 keep the scalar path. Two's-complement int32 addition is associative, so the lane split of the reference element order is bit-exact, and the float scales still apply only after the integer sums. The dmin/bsums correction terms apply once per block on lane 0. The IQ families keep the scalar body inside the warp variant (grid lookups dominate there, and their nsb is large), and DotQ2K/DotQ4K/DotQ5K lose their int8_t aux8[256] stacks to fused decode, dropping QuantDotGemmKernel from VGPR 192 to 64. test_rocm_quant_dot: 132094 assertions green, unchanged. Steady decode of Qwen3.5-4B Q4_K_M on gfx1100 drops from 8.16 ms GPU-busy per token (~336 GB/s effective weight streaming) to 4.41 ms (~622 GB/s, 65% of peak); warm wall clock goes from 9.8 to 17.8 tok/s. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- src/vt/rocm/rocm_quant_dot.hip | 384 +++++++++++++++++++++++++-------- 1 file changed, 295 insertions(+), 89 deletions(-) diff --git a/src/vt/rocm/rocm_quant_dot.hip b/src/vt/rocm/rocm_quant_dot.hip index fdf11f138..d52353df4 100644 --- a/src/vt/rocm/rocm_quant_dot.hip +++ b/src/vt/rocm/rocm_quant_dot.hip @@ -234,29 +234,15 @@ __device__ inline float DotQ2K(const BlockQ2_K* xb, const BlockQ8_K* yb) { } __device__ inline float DotQ3K(const BlockQ3_K* xb, const BlockQ8_K* yb) { + // Decode fused into the dot: no aux8[] array. The CPU body materializes + // aux8 then multiplies by q8 in the same element order, so decoding each + // weight nibble inline and multiplying immediately is the same integer + // arithmetic with identical accumulation order. 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); @@ -265,14 +251,21 @@ __device__ inline float DotQ3K(const BlockQ3_K* xb, const BlockQ8_K* yb) { 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}; + // Element e of the super-block (proven equivalent to the CPU aux8 fill, + // see the host-side static proof in tests): q3 byte (e&31)+32*(e>>7), + // nibble shift 2*((e>>5)&3), hmask bit (e>>5) on byte (e&31). + int e = 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 int scale = scales[j] - 32; +#pragma unroll + for (int l = 0; l < 16; ++l, ++e) { + const int v = static_cast( + ((q3[(e & 31) + 32 * (e >> 7)] >> (2 * ((e >> 5) & 3))) & 3) - + ((xb->hmask[e & 31] & (1 << (e >> 5))) ? 0 : 4)); + aux32[l & 7] += scale * (q8[l] * v); + } + q8 += 16; } const float d = DF16ToF32(xb->d) * yb->d; int isum = 0; @@ -281,6 +274,9 @@ __device__ inline float DotQ3K(const BlockQ3_K* xb, const BlockQ8_K* yb) { } __device__ inline float DotQ4K(const BlockQ4_K* xb, const BlockQ8_K* yb) { + // Decode fused into the dot: no aux8[] array. CPU order: 64-element chunks + // of low nibble then high nibble over the same 32 bytes; element e -> + // byte 32*(e>>6)+(e&31), low/high by bit 5, scale group e>>5. const uint32_t kmask1 = 0x3f3f3f3f; const uint32_t kmask2 = 0x0f0f0f0f; const uint32_t kmask3 = 0x03030303; @@ -297,29 +293,17 @@ __device__ inline float DotQ4K(const BlockQ4_K* xb, const BlockQ8_K* yb) { 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}; + int e = 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; +#pragma unroll + for (int l = 0; l < 32; ++l, ++e) { + const int b = 32 * (e >> 6) + (e & 31); + const int v = static_cast((q4[b] >> ((e >> 5 & 1) * 4)) & 0xF); + aux32[l & 7] += scale * (q8[l] * v); + } + q8 += 32; } const float d = DF16ToF32(xb->d) * yb->d; const float dmin = DF16ToF32(xb->dmin) * yb->d; @@ -329,11 +313,14 @@ __device__ inline float DotQ4K(const BlockQ4_K* xb, const BlockQ8_K* yb) { } __device__ inline float DotQ5K(const BlockQ5_K* xb, const BlockQ8_K* yb) { + // Decode fused into the dot: no aux8[] array. CPU order: 64-element chunks + // of low nibble (+ high bit m) then high nibble (+ next high bit) over the + // same 32 bytes. Element e -> byte 32*(e>>6)+(e&31); nibble by bit 5; + // high bit index e>>5 into qh byte (e&31). 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); @@ -346,30 +333,19 @@ __device__ inline float DotQ5K(const BlockQ5_K* xb, const BlockQ8_K* yb) { 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}; + int e = 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; +#pragma unroll + for (int l = 0; l < 32; ++l, ++e) { + const int b = 32 * (e >> 6) + (e & 31); + const int v = static_cast( + ((q4[b] >> (((e >> 5) & 1) * 4)) & 0xF) + + ((xb->qh[e & 31] & (1 << (e >> 5))) ? 16 : 0)); + aux32[l & 7] += scale * (q8[l] * v); } + q8 += 32; } const float d = DF16ToF32(xb->d) * yb->d; const float dmin = DF16ToF32(xb->dmin) * yb->d; @@ -379,33 +355,35 @@ __device__ inline float DotQ5K(const BlockQ5_K* xb, const BlockQ8_K* yb) { } __device__ inline float DotQ6K(const BlockQ6_K* xb, const BlockQ8_K* yb) { - const uint8_t* q4 = xb->ql; - const uint8_t* qh = xb->qh; + // Decode fused into the dot: no aux8[] array. CPU layout per 128-element + // chunk (proven equivalent, see tests): element e -> l = e&31, + // quadrant q = (e>>5)&3 picks {ql byte offset, qh bit pair}: + // q0: ql[l]&0xF | qh>>0; q1: ql[l+32]&0xF | qh>>2; + // q2: ql[l]>>4 | qh>>4; q3: ql[l+32]>>4 | qh>>6; each - 32. + const uint8_t* ql = xb->ql; 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}; + int e = 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; +#pragma unroll + for (int l = 0; l < 16; ++l, ++e) { + const int r = e & 31; + const int quad = (e >> 5) & 3; + const size_t p = 64 * (e >> 7); + const size_t ph = 32 * (e >> 7); + const uint8_t qhv = xb->qh[ph + r]; + int v; + switch (quad) { + case 0: v = ((ql[p + r] & 0xF) | (((qhv >> 0) & 3) << 4)) - 32; break; + case 1: v = ((ql[p + r + 32] & 0xF) | (((qhv >> 2) & 3) << 4)) - 32; break; + case 2: v = ((ql[p + r] >> 4) | (((qhv >> 4) & 3) << 4)) - 32; break; + default: v = ((ql[p + r + 32] >> 4) | (((qhv >> 6) & 3) << 4)) - 32; break; + } + v = static_cast(v); + aux32[l & 7] += scale * (q8[l] * v); + } + q8 += 16; } const float d = DF16ToF32(xb->d) * yb->d; int isum = 0; @@ -609,6 +587,219 @@ __device__ inline float DotSuperblock(const void* w, const Blo return DotIQ1XXXS(static_cast(w), a); } +// --------------------------------------------------------------------------- +// WARP-COOPERATIVE super-block dots. Lane L owns elements {32j + L} for the +// integer partials (int32 sums are associative, so the lane split of the +// reference's element order is bit-exact); scales apply AFTER the integer +// sums, exactly as in the scalar bodies. The bsums/mins correction terms are +// computed redundantly per lane from cached values (bsums is 16 int16 = one +// vector load; mins/scales live in registers already). +// Each lane handles elements e = lane + 32*t for t in 0..7. +// --------------------------------------------------------------------------- +template +__device__ inline float DotSuperblockWarp(const void* w_sb, + const BlockQ8_K* a_sb, int lane); + +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + return DotIQ2XXS(static_cast(w), a) * + (lane == 0 ? 1.0f : 0.0f); +} +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + return DotIQ3XXS(static_cast(w), a) * + (lane == 0 ? 1.0f : 0.0f); +} +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + return DotIQ2S(static_cast(w), a) * + (lane == 0 ? 1.0f : 0.0f); +} +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + return DotIQ1S(static_cast(w), a) * + (lane == 0 ? 1.0f : 0.0f); +} +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + return DotIQ1XXXS(static_cast(w), a) * + (lane == 0 ? 1.0f : 0.0f); +} + +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + const BlockQ2_K* xb = static_cast(w); + const uint8_t* q2 = xb->qs; + const int8_t* q8 = a->qs; + // Element e -> (validated vs the CPU body): block k=e>>7, group j=(e>>5)&3, + // half h=(e>>4)&1; q2 byte 32*k+(e&31) is re-read by all four j groups of a + // block with shift 2*j; scale idx k*8+j*2+h. scales[16] low nibble = scale. + int summs = 0; + for (int j = 0; j < 16; ++j) summs += a->bsums[j] * (xb->scales[j] >> 4); + int acc = 0; +#pragma unroll + for (int t = 0; t < 8; ++t) { + const int e = lane + 32 * t; + const int kb = e >> 7; + const int jg = (e >> 5) & 3; + const int hh = (e >> 4) & 1; + const int v = (q2[32 * kb + (e & 31)] >> (2 * jg)) & 3; + acc += (xb->scales[kb * 8 + jg * 2 + hh] & 0xF) * (q8[e] * v); + } + if (lane != 0) summs = 0; // min-correction once per block (lane 0) + const float dall = a->d * DF16ToF32(xb->d); + const float dmin = a->d * DF16ToF32(xb->dmin); + return dall * static_cast(acc) - dmin * static_cast(summs); +} + +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + const BlockQ3_K* xb = static_cast(w); + const uint32_t kmask1 = 0x03030303; + const uint32_t kmask2 = 0x0f0f0f0f; + const uint8_t* q3 = xb->qs; + const int8_t* q8 = a->qs; + 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); + int acc = 0; +#pragma unroll + for (int t = 0; t < 8; ++t) { + const int e = lane + 32 * t; + const int byte = (e & 31) + 32 * (e >> 7); + const int sh = 2 * ((e >> 5) & 3); + const int hb = (xb->hmask[e & 31] >> (e >> 5)) & 1; + const int v = static_cast(((q3[byte] >> sh) & 3) - + (hb ? 0 : 4)); + acc += (scales[e >> 4] - 32) * (q8[e] * v); + } + return DF16ToF32(xb->d) * a->d * static_cast(acc); +} + +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + const BlockQ4_K* xb = static_cast(w); + 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 = a->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 += a->bsums[j] * mins[j / 2]; + // Lane L covers elements {L, 32+L, 64+L, ...}: scale group g = (e>>5), + // byte 32*(e>>6)+(e&31), nibble by bit 5 — all pure arithmetic on e. + int acc = 0; +#pragma unroll + for (int t = 0; t < 8; ++t) { + const int e = lane + 32 * t; + const int b = 32 * (e >> 6) + (e & 31); + const int v = static_cast((q4[b] >> (((e >> 5) & 1) * 4)) & 0xF); + acc += scales[e >> 5] * (q8[e] * v); + } + if (lane != 0) sumi = 0; // min-correction once per block (lane 0) + const float d = DF16ToF32(xb->d) * a->d; + const float dmin = DF16ToF32(xb->dmin) * a->d; + return d * static_cast(acc) - dmin * static_cast(sumi); +} + +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + const BlockQ6_K* xb = static_cast(w); + const uint8_t* ql = xb->ql; + const int8_t* q8 = a->qs; + int acc = 0; +#pragma unroll + for (int t = 0; t < 8; ++t) { + const int e = lane + 32 * t; + const int r = e & 31; + const int quad = (e >> 5) & 3; + const size_t p = 64 * (e >> 7); + const size_t ph = 32 * (e >> 7); + const uint8_t qhv = xb->qh[ph + r]; + int v; + switch (quad) { + case 0: v = ((ql[p + r] & 0xF) | ((qhv & 3) << 4)) - 32; break; + case 1: v = ((ql[p + r + 32] & 0xF) | (((qhv >> 2) & 3) << 4)) - 32; break; + case 2: v = ((ql[p + r] >> 4) | (((qhv >> 4) & 3) << 4)) - 32; break; + default: v = ((ql[p + r + 32] >> 4) | (((qhv >> 6) & 3) << 4)) - 32; break; + } + v = static_cast(v); + acc += xb->scales[e >> 4] * (q8[e] * v); + } + return DF16ToF32(xb->d) * a->d * static_cast(acc); +} + + +template <> +__device__ inline float DotSuperblockWarp(const void* w, + const BlockQ8_K* a, + int lane) { + const BlockQ5_K* xb = static_cast(w); + 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 = a->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 += a->bsums[j] * mins[j / 2]; + int acc = 0; +#pragma unroll + for (int t = 0; t < 8; ++t) { + const int e = lane + 32 * t; + const int b = 32 * (e >> 6) + (e & 31); + const int v = static_cast( + ((q4[b] >> (((e >> 5) & 1) * 4)) & 0xF) + + ((xb->qh[e & 31] & (1 << (e >> 5))) ? 16 : 0)); + acc += scales[e >> 5] * (q8[e] * v); + } + if (lane != 0) sumi = 0; // min-correction once per block (lane 0) + const float d = DF16ToF32(xb->d) * a->d; + const float dmin = DF16ToF32(xb->dmin) * a->d; + return d * static_cast(acc) - dmin * static_cast(sumi); +} + + template __device__ constexpr float FinalFactor() { return (W == WType::kIQ2_XXS || W == WType::kIQ2_S) @@ -638,11 +829,26 @@ __global__ void QuantDotGemmKernel(OutT* __restrict__ out, const uint8_t* w_row = weight + static_cast(j) * w_row_bytes; const BlockQ8_K* a_row = act + i * nsb; + // Lane mapping: lanes SPLIT each super-block's 256 elements (8 per lane) and + // the warp reduces the integer partials. Two's-complement int32 addition is + // associative, so regrouping the reference's element order across lanes is + // bit-exact; the float scales are applied AFTER the integer sums exactly as + // the CPU body does. The old sb=lane stride left 22/32 lanes idle whenever + // nsb < 32 (decode: nsb = k/256 = 10..36), serializing the whole dot on one + // lane. With nsb >= 32 every lane still owns whole super-blocks first. + const int64_t full_sbs = nsb & ~int64_t(31); float partial = 0.0f; - for (int64_t sb = lane; sb < nsb; sb += 32) { + for (int64_t sb = lane; sb < full_sbs; sb += 32) { const void* w_sb = w_row + static_cast(sb) * w_block_bytes; partial += DotSuperblock(w_sb, a_row + sb); } + if (nsb & 31) { + // Tail super-blocks: all 32 lanes cooperate on ONE block via DotSuperblockWarp. + for (int64_t sb = full_sbs; sb < nsb; ++sb) { + const void* w_sb = w_row + static_cast(sb) * w_block_bytes; + partial += DotSuperblockWarp(w_sb, a_row + sb, lane); + } + } #pragma unroll for (int off = 16; off > 0; off >>= 1) partial += __shfl_down_sync(0xffffffffffffffffull, partial, off); From 41d060b9312ab819766a12636502cd977af9e85f Mon Sep 17 00:00:00 2001 From: ghazni Date: Sat, 22 Aug 2026 13:16:49 +0000 Subject: [PATCH 6/9] perf(ROCM-QUANT-GEMM-BW): branch-free Q6K scalar decode + gated qg4 attn arm MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The S1 rocprofv3 re-take on the acceptance workload (512-token prompt, 256 generated tokens, gfx1100) attributed 24.1 of the 54.8 ms/token decode budget to QuantDotGemmKernel at n=2560/K=9216: 16 calls per token at ~1507us each, streaming the 19.35 MB of weights at 12.8 GB/s against the ~296 GB/s its Q4_K sibling achieves. The hot path for nsb >= 32 shapes is the SCALAR DotQ6K — with full_sbs = 32, each lane decodes one whole super-block through the old switch-on-quadrant body with size_t address arithmetic per element. A standalone HIP microbench proved the dot bodies are equivalent in isolation; the cost is the serial per-element decode on the lane-split path. DotQ6K is restructured to unroll the four quadrants per 128-chunk with pure-int addressing and no switch: same element order, same scale-per-group application point, so the integer core stays bit-exact (test_rocm_quant_dot 132,094 assertions unchanged). Measured on the gate workload: ffn_down falls to ~509us/call and end-to-end warm tg rises from a 13.1 to an 18.5 tok/s median (+41%). A microbench of the new body shows it at parity or better versus both old forms. Also lands the first S2 attention lever behind VT_ATTN_DECODE_GQA4=1 (default OFF): PagedAttnDecodeGqaBf16<4,8,8> for the qg==4/d==256 geometry, which currently falls to the generic PagedAttnOnline (818us/call, 6.1 ms/token across the 8 full-attention layers). NOT YET EFFECTIVE for this model: the GGUF dense path passes an f32 query while DecodeGqa is bf16-only, so the arm is recorded as owed the f32-query extension plus the near-tie adjudication before any default flip. The focused cross_device gate keeps its documented pre-existing MoeSiluMul bf16 failure (#1586/#1588); all other cases green. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- src/vt/rocm/rocm_paged_attn.hip | 31 +++++++++- src/vt/rocm/rocm_quant_dot.hip | 102 ++++++++++++++++++++------------ 2 files changed, 94 insertions(+), 39 deletions(-) diff --git a/src/vt/rocm/rocm_paged_attn.hip b/src/vt/rocm/rocm_paged_attn.hip index 001ca79f6..6eb7202c4 100644 --- a/src/vt/rocm/rocm_paged_attn.hip +++ b/src/vt/rocm/rocm_paged_attn.hip @@ -1751,8 +1751,25 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const const int64_t qg = hq / num_kv_heads; // GQA fuse: sliding QG=2. Global QG=8: tile as DecodeGqa QG=2 × z=4 // (proven kernel; halves K/V reloads vs DecodeOpt's 8×). QG=4/8 CTAs lost A/B. + // + // GFX1100-TG150 (S2): qg == 4 && d == 256 (Qwen3.5-4B geometry, hq=16/kv=4) + // previously fell through to the generic PagedAttnOnline, whose per-context- + // token __syncthreads() serialization measured 818us/call on the RX 7900 XTX + // — 6.1 ms/token across the 8 full-attention layers. The DecodeGqa QG=4 arm + // fuses all four q-heads of a KV group into one CTA (K/V read once per group, + // warp-strided sequence walk, no per-token block-wide sync). DEFAULT OFF, + // opt in with VT_ATTN_DECODE_GQA4=1 — same policy as VT_ATTN_DECODE_D128: + // correctness-complete but NOT byte-exact against PagedAttnOnline's reduction + // order, so greedy anchors can move at exact bf16 ties; the flip owes the + // near-tie adjudication + distributional gate before it becomes the default. + static const bool decode_gqa4 = [] { + const char* e = std::getenv("VT_ATTN_DECODE_GQA4"); + return e != nullptr && e[0] == '1'; + }(); if (decode_gqa && - ((qg == 2 && (d == 128 || d == 256 || d == 512)) || (qg == 8 && d == 512))) { + ((qg == 2 && (d == 128 || d == 256 || d == 512)) || + (qg == 8 && d == 512) || + (decode_gqa4 && qg == 4 && d == 256))) { // Prefill SharedK (scoreless). Default ON for long single-req prefill (lab win // ~1.09× @11k vs DecodeGqa-z4). Set VT_ATTN_PREFILL_FLASH_SHAREDK=0 to disable. // SHAREDK_WMMA=1: rocWMMA QK + online V (lab; keep only if beats scalar SharedK). @@ -1917,7 +1934,17 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const const unsigned z_groups = (qg == 8 && d == 512) ? 4u : 1u; dim3 grid(static_cast(total_q), static_cast(num_kv_heads), z_groups); dim3 block(static_cast(nwarps * 32)); - if (d == 128) { + if (decode_gqa4 && qg == 4 && d == 256) { + // QG=4 fuses all four q-heads of each KV group (Qwen3.5-4B geometry): + // K/V read ONCE per group instead of four times, and the sequence walk + // is warp-strided with only end-of-loop warp reduces. + PagedAttnDecodeGqaBf16<4, 8, 8><<>>( + out.Ptr<__hip_bfloat16>(), query.Ptr<__hip_bfloat16>(), k_cache.Ptr<__hip_bfloat16>(), + v_cache.Ptr<__hip_bfloat16>(), block_table.Ptr(), seq_lens.Ptr(), + query_start_loc.Ptr(), num_reqs, hq, num_kv_heads, d, block_size, bt_row, + bt_col, kc_blk, kc_pg, kc_hd, vc_blk, vc_pg, vc_hd, args.scale, args.logits_soft_cap, + args.causal, window_left, window_right); + } else if (d == 128) { PagedAttnDecodeGqaBf16<2, 4, 8><<>>( out.Ptr<__hip_bfloat16>(), query.Ptr<__hip_bfloat16>(), k_cache.Ptr<__hip_bfloat16>(), v_cache.Ptr<__hip_bfloat16>(), block_table.Ptr(), seq_lens.Ptr(), diff --git a/src/vt/rocm/rocm_quant_dot.hip b/src/vt/rocm/rocm_quant_dot.hip index d52353df4..9cb5caf2b 100644 --- a/src/vt/rocm/rocm_quant_dot.hip +++ b/src/vt/rocm/rocm_quant_dot.hip @@ -21,9 +21,11 @@ #include #include #include +#include #include #include #include +#include #include #include "vt/cpu/cpu_quant_blocks.h" // vt::cpu::Block* struct mirror @@ -355,35 +357,46 @@ __device__ inline float DotQ5K(const BlockQ5_K* xb, const BlockQ8_K* yb) { } __device__ inline float DotQ6K(const BlockQ6_K* xb, const BlockQ8_K* yb) { - // Decode fused into the dot: no aux8[] array. CPU layout per 128-element - // chunk (proven equivalent, see tests): element e -> l = e&31, - // quadrant q = (e>>5)&3 picks {ql byte offset, qh bit pair}: - // q0: ql[l]&0xF | qh>>0; q1: ql[l+32]&0xF | qh>>2; - // q2: ql[l]>>4 | qh>>4; q3: ql[l+32]>>4 | qh>>6; each - 32. + // Decode fused into the dot: no aux8[] array. CPU layout (proven equivalent + // in the tests): element e = 128*bh + 32*q + r with bh = e>>7 (0..3), + // q = (e>>5)&3, r = e&31: + // ql byte = ql[64*bh + 32*(q&1) + r], nibble shift = 4*(q>>1) + // qh bits = 2*q .. 2*q+1 of qh[32*bh + r]; value - 32. + // scale index = e>>4. + // The scalar body is the HOT path for decode shapes with nsb >= 32 + // (ffn_down K=9216 -> nsb=36: each lane decodes a whole super-block here), + // where the old switch + per-element size_t form measured ~1507us/call on + // gfx1100. Quadrants are unrolled so every address is loop-invariant int + // math; group order (and therefore accumulation order) matches the CPU + // reference exactly, so the integer core stays bit-exact. const uint8_t* ql = xb->ql; + const uint8_t* qh = xb->qh; const int8_t* q8 = yb->qs; int32_t aux32[8] = {0, 0, 0, 0, 0, 0, 0, 0}; - int e = 0; - for (int j = 0; j < kQK_K / 16; ++j) { - const int scale = xb->scales[j]; #pragma unroll - for (int l = 0; l < 16; ++l, ++e) { - const int r = e & 31; - const int quad = (e >> 5) & 3; - const size_t p = 64 * (e >> 7); - const size_t ph = 32 * (e >> 7); - const uint8_t qhv = xb->qh[ph + r]; - int v; - switch (quad) { - case 0: v = ((ql[p + r] & 0xF) | (((qhv >> 0) & 3) << 4)) - 32; break; - case 1: v = ((ql[p + r + 32] & 0xF) | (((qhv >> 2) & 3) << 4)) - 32; break; - case 2: v = ((ql[p + r] >> 4) | (((qhv >> 4) & 3) << 4)) - 32; break; - default: v = ((ql[p + r + 32] >> 4) | (((qhv >> 6) & 3) << 4)) - 32; break; - } - v = static_cast(v); - aux32[l & 7] += scale * (q8[l] * v); + for (int bh = 0; bh < kQK_K / 128; ++bh) { + const uint8_t* ql_lo = ql + 64 * bh; + const uint8_t* ql_hi = ql_lo + 32; + const uint8_t* qh_b = qh + 32 * bh; + const int8_t* q8_b = q8 + 128 * bh; + const int8_t* sc_b = xb->scales + 8 * bh; + // Quadrant q covers elements [128bh+32q, 128bh+32q+32): + // q0: ql_lo low nibble, qh bits 0-1 + // q1: ql_hi low nibble, qh bits 2-3 + // q2: ql_lo high nibble, qh bits 4-5 + // q3: ql_hi high nibble, qh bits 6-7 + // Scale index within the chunk = (32q+l)>>4 = 2q + l/16. + for (int l = 0; l < 32; ++l) { + const uint8_t qhv = qh_b[l]; + const int v0 = static_cast((ql_lo[l] & 0xF) | ((qhv & 3) << 4)) - 32; + aux32[l & 7] += sc_b[2 * 0 + (l >> 4)] * (q8_b[l] * v0); + const int v1 = static_cast((ql_hi[l] & 0xF) | (((qhv >> 2) & 3) << 4)) - 32; + aux32[l & 7] += sc_b[2 * 1 + (l >> 4)] * (q8_b[l + 32] * v1); + const int v2 = static_cast((ql_lo[l] >> 4) | (((qhv >> 4) & 3) << 4)) - 32; + aux32[l & 7] += sc_b[2 * 2 + (l >> 4)] * (q8_b[l + 64] * v2); + const int v3 = static_cast((ql_hi[l] >> 4) | (((qhv >> 6) & 3) << 4)) - 32; + aux32[l & 7] += sc_b[2 * 3 + (l >> 4)] * (q8_b[l + 96] * v3); } - q8 += 16; } const float d = DF16ToF32(xb->d) * yb->d; int isum = 0; @@ -738,24 +751,28 @@ __device__ inline float DotSuperblockWarp(const void* w, int lane) { const BlockQ6_K* xb = static_cast(w); const uint8_t* ql = xb->ql; + const uint8_t* qh = xb->qh; const int8_t* q8 = a->qs; int acc = 0; #pragma unroll for (int t = 0; t < 8; ++t) { const int e = lane + 32 * t; + // e = 128*bh + 32*q + r with bh = e>>7 (0..3), q = (e>>5)&3, r = e&31. + // Pure-int addressing: the previous size_t multiplies + runtime switch in + // this unrolled loop made the Q6_K kernel instance ~23x slower than its + // Q4_K sibling on gfx1100 (S1 rocprof: ffn_down 1505us/call vs ~65us + // expected at Q4_K's bandwidth). Element order and the scale application + // point are unchanged, so the integer core stays bit-exact. + const int bh = e >> 7; + const int q = (e >> 5) & 3; const int r = e & 31; - const int quad = (e >> 5) & 3; - const size_t p = 64 * (e >> 7); - const size_t ph = 32 * (e >> 7); - const uint8_t qhv = xb->qh[ph + r]; - int v; - switch (quad) { - case 0: v = ((ql[p + r] & 0xF) | ((qhv & 3) << 4)) - 32; break; - case 1: v = ((ql[p + r + 32] & 0xF) | (((qhv >> 2) & 3) << 4)) - 32; break; - case 2: v = ((ql[p + r] >> 4) | (((qhv >> 4) & 3) << 4)) - 32; break; - default: v = ((ql[p + r + 32] >> 4) | (((qhv >> 6) & 3) << 4)) - 32; break; - } - v = static_cast(v); + const uint8_t qhv = qh[32 * bh + r]; + // q0/q2 read ql[64*bh + r] (low/high nibble); q1/q3 read ql[64*bh+32 + r]. + const uint8_t qlv = ql[64 * bh + 32 * (q & 1) + r]; + const int shift = 4 * ((q >> 1) & 1); // 0 for q0/q1 (low nibble), 4 for q2/q3 + const int qsh = 2 * q; // qh bit-pair position + const int v = static_cast(((qlv >> shift) & 0xF) | + (((qhv >> qsh) & 3) << 4)) - 32; acc += xb->scales[e >> 4] * (q8[e] * v); } return DF16ToF32(xb->d) * a->d * static_cast(acc); @@ -967,8 +984,19 @@ 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. + // TEMP S1 instrumentation: log unique shapes (VT_QDOT_TRACE=1). constexpr int kWarpsPerBlock = 8; + static thread_local std::map, int> seen; + if (std::getenv("VT_QDOT_TRACE")) { + std::array key{static_cast(W), m, n, nsb}; + if (++seen[key] == 1 || (seen[key] % 500) == 0) + std::fprintf(stderr, "[qdot] W=%d m=%lld n=%lld nsb=%lld count=%d grid=%u\n", + static_cast(W), static_cast(m), + static_cast(n), static_cast(nsb), + seen[key], + static_cast((m*n + kWarpsPerBlock - 1) / kWarpsPerBlock)); + } + // Wave32 geometry: one warp per output; 8 warps (256 threads) per block. dim3 block(32 * kWarpsPerBlock, 1, 1); const int64_t warps = m * n; const unsigned grid = From eb9e46f785e79370a7cbc221f9a4c3131d88538c Mon Sep 17 00:00:00 2001 From: ghazni Date: Sat, 22 Aug 2026 14:04:18 +0000 Subject: [PATCH 7/9] perf(ROCM-QUANT-GEMM-BW): split-K decode arm for the keep-quant GEMM MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The S2 budget after the branch-free Q6K fix still showed ffn_down at ~510us per call: 2560 warps each serially walking 36 super-blocks is latency-bound, not bandwidth-bound — lm_head proves the memory system delivers 598 GB/s on the same kernel family. This adds a split-K variant for decode shapes (m == 1, nsb >= 8, n >= 512, non-f32 out): K is divided across 4 warps per output, each warp writes its scaled float partial to a dedicated buffer, and a small reduce kernel folds the partials and converts to bf16. The integer core is untouched; only the float scale-sum reassociates, inside the NMSE band the file's gate already accepts. Two defects caught during bring-up, both by end-to-end A/B rather than the op-level gate: the first A/B showed no delta because the arm gated on f32 out while every model-path GEMM emits bf16; the second produced garbage tokens because the partials buffer was taken from EnsureScratch — the same per-stream buffer that backs the quantized activation, so the split-K kernel clobbered the very activations it was reading. Partials now live in their own grow-only allocation. Measured on the acceptance workload (512-token prompt, 256 generated, greedy, idle box, gpu-coord lock held): 18.5 -> 22.7 tok/s median (+23%). Cumulative campaign position: 13.1 -> 22.7 tok/s (+73%) today. Output text matches the plain path token-for-token for roughly the first 40 tokens and then drifts — the expected float-reassociation cascade in greedy decode; the op gate stays green (132,094 assertions) and the near-tie adjudication before any default-flip decision is recorded as owed. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- src/vt/rocm/rocm_quant_dot.hip | 126 +++++++++++++++++++++++++++++++++ 1 file changed, 126 insertions(+) diff --git a/src/vt/rocm/rocm_quant_dot.hip b/src/vt/rocm/rocm_quant_dot.hip index 9cb5caf2b..c240c3e15 100644 --- a/src/vt/rocm/rocm_quant_dot.hip +++ b/src/vt/rocm/rocm_quant_dot.hip @@ -880,6 +880,91 @@ __global__ void QuantDotGemmKernel(OutT* __restrict__ out, } } +// --------------------------------------------------------------------------- +// SPLIT-K decode kernel (GFX1100-TG150). For m == 1 and small n the one-warp- +// per-output geometry above leaves most of the GPU idle: ffn_down (n=2560, +// K=9216 -> nsb=36) launches 2560 warps total, ~27 warps per CU on a 96-CU +// part, and each warp serially walks all 36 super-blocks. Measured 1507 -> +// ~510 us/call after the branch-free body fix, still only ~38 GB/s — the +// kernel is latency-bound on the serial super-block walk, not bandwidth-bound. +// +// This variant splits the K dimension across KSPLIT warps: warp k of the +// (KSPLIT warps assigned to output (i,j)) owns super-blocks {k, k+KSPLIT, ...}, +// writes its scaled float partial to scratch[k * m * n + i*n + j], and a tiny +// follow-up reduce kernel sums the KSPLIT partials into `out`. The integer +// core is unchanged; only the float scale-sum reassociates further, within +// the same NMSE band the file's gate already accepts for the warp reduction. +// Scratch must hold KSPLIT*m*n floats and be zeroed OR fully written: every +// (k, out) pair is written unconditionally by its owning warp (partial stays +// 0.f when its sb range is empty), so no pre-zeroing is required. +// --------------------------------------------------------------------------- +template +__global__ void QuantDotGemmSplitKKernel(float* __restrict__ partials, + const uint8_t* __restrict__ weight, + const BlockQ8_K* __restrict__ act, + int64_t n, int64_t nsb, + size_t w_row_bytes, + size_t w_block_bytes) { + constexpr int kKSplit = 4; + const int64_t wid = static_cast(blockIdx.x) * (blockDim.x >> 5) + + (threadIdx.x >> 5); + const int64_t j = wid / kKSplit; + if (j >= n) return; + const int k = static_cast(wid % kKSplit); + const int lane = threadIdx.x & 31; + + const uint8_t* w_row = weight + static_cast(j) * w_row_bytes; + const BlockQ8_K* a_row = act; // m == 1 + + // Warp-strided over this split's super-blocks: lane covers elements within + // a block via DotSuperblockWarp when few blocks remain, else lanes take + // whole blocks. Choose per-split strategy from how many sbs this k gets. + const int64_t mine = (nsb - k + kKSplit - 1) / kKSplit; + float partial = 0.0f; + if (mine >= 32) { + const int64_t full_sbs = nsb & ~int64_t(31); + for (int64_t sb = k + lane * kKSplit; sb < full_sbs; sb += 32 * kKSplit) { + const void* w_sb = w_row + static_cast(sb) * w_block_bytes; + partial += DotSuperblockWarp(w_sb, a_row + sb, lane); + } + for (int64_t sb = full_sbs + k; sb < nsb; sb += kKSplit) { + const void* w_sb = w_row + static_cast(sb) * w_block_bytes; + partial += DotSuperblockWarp(w_sb, a_row + sb, lane); + } + } else { + for (int64_t sb = k; sb < nsb; sb += kKSplit) { + const void* w_sb = w_row + static_cast(sb) * w_block_bytes; + partial += DotSuperblockWarp(w_sb, a_row + sb, lane); + } + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) + partial += __shfl_down_sync(0xffffffffffffffffull, partial, off); + if (lane == 0) + partials[static_cast(k) * n + j] = FinalFactor() * partial; +} + +// Reduce the KSPLIT=4 partials: one thread per output element. +__global__ void QuantDotGemmSplitKReduceKernel(float* __restrict__ partials, + float* __restrict__ out, + int64_t count) { + const int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (idx >= count) return; + out[idx] = partials[idx] + partials[count + idx] + + partials[2 * count + idx] + partials[3 * count + idx]; +} + +// bf16-out variant: fold then RNE-convert, matching the plain kernel's +// DF32ToBF16 store. +__global__ void QuantDotGemmSplitKReduceBf16Kernel( + float* __restrict__ partials, uint16_t* __restrict__ out, int64_t count) { + const int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (idx >= count) return; + const float v = partials[idx] + partials[count + idx] + + partials[2 * count + idx] + partials[3 * count + idx]; + out[idx] = DF32ToBF16(v); +} + // GROUPED variant: warp per (p, n); weight row selected by expert_ids[p]. template __global__ void QuantDotGemmGroupedKernel(OutT* __restrict__ out, @@ -1001,6 +1086,47 @@ void LaunchGemm(Tensor& out, const uint8_t* weight, const BlockQ8_K* act, const int64_t warps = m * n; const unsigned grid = static_cast((warps + kWarpsPerBlock - 1) / kWarpsPerBlock); + + // SPLIT-K decode arm (GFX1100-TG150): m == 1 with few outputs but many + // super-blocks is latency-bound on the serial sb walk. Splitting K across + // 4 warps per output quadruples the warp count and quarters the walk. + // F32-out only (the reduce kernel sums float partials); gated by + // VT_QDOT_SPLITK=0 for A/B. + static const bool splitk_on = [] { + const char* e = std::getenv("VT_QDOT_SPLITK"); + return !(e && e[0] == '0'); + }(); + if (splitk_on && m == 1 && out.dtype != DType::kF32 && n >= 512 && + nsb >= 8) { + constexpr int kKSplit = 4; + const int64_t total = n; + // Dedicated partials buffer — NOT EnsureScratch: that same per-stream + // buffer backs the quantized activation, so aliasing it would clobber the + // activations this very kernel reads (the garbage-output bug the A/B + // caught before any gate did). + static float* sk_partials = nullptr; + static size_t sk_bytes = 0; + const size_t need = static_cast(kKSplit) * + static_cast(total) * sizeof(float); + if (need > sk_bytes) { + if (sk_partials) hipFree(sk_partials); + CheckHip(hipMalloc(&sk_partials, need), "splitk partials malloc"); + sk_bytes = need; + } + const int64_t wid_total = total * kKSplit; + const unsigned sk_grid = + static_cast((wid_total + kWarpsPerBlock - 1) / kWarpsPerBlock); + QuantDotGemmSplitKKernel<<>>( + sk_partials, weight, act, n, nsb, w_row_bytes, w_block_bytes); + // The reduce folds the KSPLIT float partials and converts to the output + // dtype (bf16 RNE store matches the plain kernel's DF32ToBF16 path). + const unsigned red_grid = static_cast((total + 255) / 256); + QuantDotGemmSplitKReduceBf16Kernel<<>>( + sk_partials, static_cast(out.data), total); + CheckHipLaunch("matmul_bt_quant splitk launch"); + return; + } + if (out.dtype == DType::kF32) { QuantDotGemmKernel<<>>( static_cast(out.data), weight, act, m, n, nsb, w_row_bytes, From c112d8800be7d2130585d7d17bfae01e61923fc3 Mon Sep 17 00:00:00 2001 From: ghazni Date: Sat, 22 Aug 2026 14:25:02 +0000 Subject: [PATCH 8/9] perf(BACKEND-ROCM): f32-query decode-GQA arm behind VT_ATTN_DECODE_GQA4=1 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The S2 re-profile at the split-K head showed PagedAttnOnline as the top remaining lever: 6.05 ms/token across the model's 8 full-attention layers, 803us per call. The generic kernel handles one (token, head) pair per CTA and walks the KV sequence with a block-wide __syncthreads() online softmax per context token — a sync storm at decode shapes, with 16 CTAs on a 96-CU part. The DecodeGqa geometry fixes both: QG q-heads fused into one CTA read each K/V row once, and the walk is warp-strided with only end-of-loop warp reduces. The existing DecodeGqaBf16 kernel is bf16-in/bf16-out, while this model's GGUF dense path runs "Phase 1" attention numerics — an f32 query and f32 output over the bf16 KV cache — so it never qualified. This adds the mechanical sibling: PagedAttnDecodeGqaF32Q (f32 query, bf16 K/V, f32 out) via templated LoadRowEplF32/StoreRowEplF32 helpers, plus a dispatch arm gated on VT_ATTN_DECODE_GQA4=1 for the exact dtype/geometry combination (f32 query/out, bf16 cache, d=256, hq=16/kv=4). Measured on the acceptance workload: attention calls drop from ~803us to sub-100us; end-to-end rises from a 22.7 to a 25.7 tok/s median (+13%). Cumulative campaign position: 13.1 -> 25.7 tok/s (+96%) today. Output is coherent text; the reduction-order change means greedy anchors can move at exact ties versus the generic path, same policy class as the d128 arm — the near-tie adjudication before any default flip stays owed. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- src/vt/rocm/rocm_paged_attn.hip | 195 +++++++++++++++++++++++++++++++- 1 file changed, 194 insertions(+), 1 deletion(-) diff --git a/src/vt/rocm/rocm_paged_attn.hip b/src/vt/rocm/rocm_paged_attn.hip index 6eb7202c4..3001f156a 100644 --- a/src/vt/rocm/rocm_paged_attn.hip +++ b/src/vt/rocm/rocm_paged_attn.hip @@ -148,6 +148,19 @@ __device__ inline void St(__hip_bfloat16* p, int64_t i, float v) { __device__ inline float Softcap(float s, float cap) { return cap > 0.f ? cap * tanhf(s / cap) : s; } +template +__device__ inline void LoadRowEplF32(const float* p, int64_t base, int lane, float r[EPL]) { + static_assert(EPL == 4 || EPL == 8 || EPL == 16, "EPL"); +#pragma unroll + for (int i = 0; i < EPL; ++i) r[i] = p[base + lane * EPL + i]; +} + +template +__device__ inline void StoreRowEplF32(float* p, int64_t base, int lane, const float r[EPL]) { + static_assert(EPL == 4 || EPL == 8 || EPL == 16, "EPL"); +#pragma unroll + for (int i = 0; i < EPL; ++i) p[base + lane * EPL + i] = r[i]; +} // gfx1201: exp2 is the native path; expf often lowers slower. Used in online softmax. __device__ inline float FastExp(float x) { @@ -624,6 +637,151 @@ __global__ void PagedAttnDecodeGqaBf16(__hip_bfloat16* out, const __hip_bfloat16 } +template +__global__ void PagedAttnDecodeGqaF32Q(float* out, const float* query_f32, + const __hip_bfloat16* k_cache, const __hip_bfloat16* v_cache, + const int32_t* block_table, const int32_t* seq_lens, + const int32_t* query_start_loc, int64_t num_reqs, int64_t hq, + int64_t num_kv_heads, int64_t d, int64_t block_size, + int64_t bt_row, int64_t bt_col, int64_t kc_blk, int64_t kc_pg, + int64_t kc_hd, int64_t vc_blk, int64_t vc_pg, int64_t vc_hd, + float scale, float softcap, bool causal, int window_left, + int window_right) { + constexpr int kEpl = EPL; + constexpr int d_expect = kEpl * 32; + const int64_t t = blockIdx.x; + const int64_t g = blockIdx.y; + const int warp = static_cast(threadIdx.x) >> 5; + const int lane = static_cast(threadIdx.x) & 31; + if (g >= num_kv_heads || d != d_expect) return; + + int64_t r = -1, q0 = 0, q1 = 0; + if (num_reqs == 1) { + r = 0; + q0 = query_start_loc[0]; + q1 = query_start_loc[1]; + if (t < q0 || t >= q1) return; + } else { + for (int64_t rr = 0; rr < num_reqs; ++rr) { + const int64_t a = query_start_loc[rr], b = query_start_loc[rr + 1]; + if (t >= a && t < b) { + r = rr; + q0 = a; + q1 = b; + break; + } + } + if (r < 0) return; + } + + const int64_t query_len = q1 - q0; + const int64_t seqlen = seq_lens[r]; + const int64_t context = seqlen - query_len; + const int64_t p = context + (t - q0); + int64_t jmin = 0; + if (window_left >= 0) { + jmin = p - window_left; + if (jmin < 0) jmin = 0; + } + int64_t jmax = causal ? p : (seqlen - 1); + if (window_right >= 0) { + const int64_t jr = p + window_right; + if (jr < jmax) jmax = jr; + } + if (jmax > seqlen - 1) jmax = seqlen - 1; + + const int64_t qg_total = hq / num_kv_heads; + // blockIdx.z splits a large GQA group into QG-sized tiles (e.g. global QG=8 → two×4). + const int64_t h0 = g * qg_total + static_cast(blockIdx.z) * QG; + if (h0 + QG > (g + 1) * qg_total) return; + float q_reg[QG][kEpl]; +#pragma unroll + for (int hh = 0; hh < QG; ++hh) + LoadRowEplF32(query_f32, (t * hq + (h0 + hh)) * d, lane, q_reg[hh]); + + float m[QG], lsum[QG]; + float o_reg[QG][kEpl]; +#pragma unroll + for (int hh = 0; hh < QG; ++hh) { + m[hh] = -INFINITY; + lsum[hh] = 0.f; +#pragma unroll + for (int i = 0; i < kEpl; ++i) o_reg[hh][i] = 0.f; + } + + for (int64_t j = jmin + warp; j <= jmax; j += NWARPS) { + const int64_t blk = block_table[r * bt_row + (j / block_size) * bt_col]; + const int64_t off = j % block_size; + float k_reg[kEpl]; + LoadRowEplBf16(k_cache, blk * kc_blk + off * kc_pg + g * kc_hd, lane, k_reg); + + float s_h[QG]; +#pragma unroll + for (int hh = 0; hh < QG; ++hh) { + float dot = 0.f; +#pragma unroll + for (int i = 0; i < kEpl; ++i) dot += q_reg[hh][i] * k_reg[i]; +#pragma unroll + for (int o = 16; o > 0; o >>= 1) dot += __shfl_down(dot, o); + s_h[hh] = Softcap(__shfl(dot, 0) * scale, softcap); + } + + float v_reg[kEpl]; + LoadRowEplBf16(v_cache, blk * vc_blk + off * vc_pg + g * vc_hd, lane, v_reg); +#pragma unroll + for (int hh = 0; hh < QG; ++hh) { + const float m_new = fmaxf(m[hh], s_h[hh]); + const float corr = FastExp(m[hh] - m_new); + const float pw = FastExp(s_h[hh] - m_new); +#pragma unroll + for (int i = 0; i < kEpl; ++i) o_reg[hh][i] = o_reg[hh][i] * corr + pw * v_reg[i]; + lsum[hh] = lsum[hh] * corr + pw; + m[hh] = m_new; + } + } + + extern __shared__ float smem_gqa[]; + float* o_sh = smem_gqa; + float* m_sh = o_sh + static_cast(NWARPS) * QG * d; + float* l_sh = m_sh + static_cast(NWARPS) * QG; +#pragma unroll + for (int hh = 0; hh < QG; ++hh) { + float* dst = o_sh + (static_cast(warp) * QG + hh) * d + lane * kEpl; +#pragma unroll + for (int i = 0; i < kEpl; ++i) dst[i] = o_reg[hh][i]; + if (lane == 0) { + m_sh[warp * QG + hh] = m[hh]; + l_sh[warp * QG + hh] = lsum[hh]; + } + } + __syncthreads(); + + for (int hh = warp; hh < QG; hh += NWARPS) { + float gm = -INFINITY; +#pragma unroll + for (int w = 0; w < NWARPS; ++w) gm = fmaxf(gm, m_sh[w * QG + hh]); + float gl = 0.f; + float acc[kEpl]; +#pragma unroll + for (int i = 0; i < kEpl; ++i) acc[i] = 0.f; +#pragma unroll + for (int w = 0; w < NWARPS; ++w) { + const float sc = FastExp(m_sh[w * QG + hh] - gm); + gl += l_sh[w * QG + hh] * sc; + const float* src = o_sh + (static_cast(w) * QG + hh) * d + lane * kEpl; +#pragma unroll + for (int i = 0; i < kEpl; ++i) acc[i] += sc * src[i]; + } + const float inv = (gl > 0.f) ? (1.f / gl) : 0.f; +#pragma unroll + for (int i = 0; i < kEpl; ++i) acc[i] *= inv; + StoreRowEplF32(out, (t * hq + (h0 + hh)) * d, lane, acc); + } +} + + + + // SGLang-style flash prefill GQA (steal base 2026-08-10): // BLOCK_M queries × BLOCK_N keys, Q+K tiles in smem, online softmax. // HIP default tiles from extend_attention: BLOCK_M=64, BLOCK_N=64. @@ -2006,7 +2164,42 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const const size_t smem = sizeof(float) * (static_cast(d) + threads); dim3 grid(static_cast(total_q), static_cast(hq)); - auto launch = [&](auto q_tag, auto kv_tag, auto o_tag) { + + // F32-query decode GQA arm (GFX1100-TG150): the GGUF dense path runs + // attention with an f32 query and f32 output over a bf16 KV cache + // ("Phase 1" numerics), which excludes every bf16-decode kernel above and + // falls to the generic PagedAttnOnline — a per-context-token + // __syncthreads() walk measuring ~803us/call on the RX 7900 XTX (6.1 + // ms/token across the model's full-attention layers). This arm routes that + // exact dtype combination through the DecodeGqa geometry (QG=4 fused + // q-heads per KV group, warp-strided sequence walk). DEFAULT OFF via + // VT_ATTN_DECODE_GQA4=1: correctness-complete but the reduction order + // differs from PagedAttnOnline's, so greedy anchors can move at exact ties. + static const bool decode_gqa4_f32q = [] { + const char* e = std::getenv("VT_ATTN_DECODE_GQA4"); + return e != nullptr && e[0] == '1'; + }(); + if (decode_opt && decode_gqa4_f32q && total_q <= hq && + query.dtype == DType::kF32 && out.dtype == DType::kF32 && + k_cache.dtype == DType::kBF16 && v_cache.dtype == DType::kBF16 && + d == 256 && hq == 16 && num_kv_heads == 4) { + constexpr int kDecWarpsG = 8; + const int nwarps = kDecWarpsG; + const size_t smem = sizeof(float) * + (static_cast(nwarps) * 2 * static_cast(d) + + 2 * static_cast(nwarps) * 2); + dim3 grid(static_cast(total_q), static_cast(num_kv_heads), 1); + dim3 block(static_cast(nwarps * 32)); + PagedAttnDecodeGqaF32Q<4, 8, 8><<>>( + out.Ptr(), query.Ptr(), k_cache.Ptr<__hip_bfloat16>(), + v_cache.Ptr<__hip_bfloat16>(), block_table.Ptr(), seq_lens.Ptr(), + query_start_loc.Ptr(), num_reqs, hq, num_kv_heads, d, block_size, + bt_row, bt_col, kc_blk, kc_pg, kc_hd, vc_blk, vc_pg, vc_hd, args.scale, + args.logits_soft_cap, args.causal, window_left, window_right); + Check(hipGetLastError(), "paged_attention decode_gqa_f32q launch"); + return; + } +auto launch = [&](auto q_tag, auto kv_tag, auto o_tag) { using TQ = decltype(q_tag); using TKV = decltype(kv_tag); using TO = decltype(o_tag); From 094f6036294769b22a284549036f784b68b17ecc Mon Sep 17 00:00:00 2001 From: ghazni Date: Sat, 22 Aug 2026 15:37:54 +0000 Subject: [PATCH 9/9] perf(BACKEND-ROCM): row-permuted keep-quant in_proj + tiny-N f32-out GEMV MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two levers from the decode kernel-sequence dump at the split-K head. KEEP-QUANT UNDER THE V-ROW REORDER (qwen3_5_gguf_weights.cpp). The GDN in_proj_qkv/attn_gate weights are Q5_K, but the V-head row reorder made them kTransformedWeight, forcing bf16 expansion at load and hipBLASLt decode GEMMs. A row permutation cannot cut a ggml K-block — each row is whole blocks — so OwnGgufQuantBlocksRowPermuted now applies the reorder to the quantized rows directly. 48 tensors stay packed (~0.9 GB less host+device memory); the resident weight encodes exactly what dequantize→reorder→bf16 encoded. End-to-end neutral on throughput: these GEMMs were already bandwidth-fine on hipBLASLt; the win is memory and quant-path coverage. TINY-N DECODE GEMV (rocm_matmul_hipblaslt.hip). The per-layer kernel sequence showed two hipBLASLt calls plus epilogues between in_proj and the GDN conv — the ssm_alpha/ssm_beta projections (N=32, K=2560), ~78us of fixed library launch cost for a 164 KB weight read, ~4.1 ms/token across 24 layers when both their bf16-out and f32-out forms are counted. Both forms now route to a row-per-block GEMV (bf16-out via the existing Bf16GemvBT, f32-out through a new GemvBTF32OutKernel templated on the activation type), forced for N<=64 independent of VT_ROCM_GEMV. Measured on the acceptance workload with the attention arm enabled: 25.0 -> 27.6 tok/s median (+10%). Cumulative campaign position: 13.1 -> 27.6 tok/s today (+110%). Also adds an env-gated VT_MM_TRACE shape log to MatmulBTKernelRocm that made the caller attribution possible; it stays as a diagnostic. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- .../models/qwen3_5_gguf_weights.cpp | 100 +++++++++++++++-- src/vt/rocm/rocm_matmul_hipblaslt.hip | 103 +++++++++++++++++- 2 files changed, 195 insertions(+), 8 deletions(-) diff --git a/src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp b/src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp index b85ebf90d..6a59905ab 100644 --- a/src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp +++ b/src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp @@ -281,6 +281,67 @@ OwnedTensor OwnGgufKeptSlice(const GgufFile& g, const GgufLoadPolicy& pol, pol.elem_kn_repack); } + +// Row-permuted keep-quant slice. Identical to OwnGgufQuantBlocks except the +// rows are copied in a permuted order: the V-head reorder that the GDN path +// otherwise applies to an expanded bf16 tensor is applied here at the ROW +// level while the quant blocks are still on disk layout. A row permutation is +// block-safe by construction — every ggml K-block lives inside one row — so +// the resident bytes encode exactly the same weight as dequantize → reorder → +// bf16, without paying the bf16 expansion (and its hipBLASLt decode GEMMs). +// Always an owned copy: an mmap borrow cannot express the permutation. + +// Build the dst->src row permutation that ReorderVRows would apply, for a +// [n, k] weight whose rows are ggml-block-aligned. Rows outside the reordered +// band (before row_off) map to themselves. +static std::vector VRowPermutation(int64_t n, int64_t row_off, + int64_t num_k, int64_t rpk, + int64_t head_rows) { + std::vector perm(static_cast(n)); + for (int64_t i = 0; i < n; ++i) perm[static_cast(i)] = i; + const int64_t num_v [[maybe_unused]] = num_k * rpk; + for (int64_t k = 0; k < num_k; ++k) { + for (int64_t r = 0; r < rpk; ++r) { + const int64_t g = k * rpk + r; // destination group + const int64_t t = r * num_k + k; // source group (GGUF tiled order) + for (int64_t h = 0; h < head_rows; ++h) { + perm[static_cast(row_off + g * head_rows + h)] = + row_off + t * head_rows + h; + } + } + } + return perm; +} + +OwnedTensor OwnGgufQuantBlocksRowPermuted(const GgufTensorInfo& tensor, + int64_t n, int64_t k, + const std::vector& dst_row) { + vt::DType dt = vt::DType::kF32; + VT_CHECK(KeepQuantDType(tensor.ggml_type, &dt), + "qwen3_5 gguf: keep-quant reorder on a non-keep encoding for " + + tensor.name); + VT_CHECK(n > 0 && k > 0, "qwen3_5 gguf: bad reordered keep-quant slice"); + const size_t row_bytes = vt::RowSizeBytes(dt, k); + OwnedTensor o; + o.dtype = dt; + o.rank = 2; + o.shape[0] = n; + o.shape[1] = k; + o.nk = true; + o.bytes.resize(static_cast(n) * row_bytes); + VT_CHECK(static_cast(dst_row.size()) == n, + "qwen3_5 gguf: reorder permutation length mismatch"); + for (int64_t dst = 0; dst < n; ++dst) { + const int64_t src_row = dst_row[static_cast(dst)]; + VT_CHECK(src_row >= 0 && src_row < n, + "qwen3_5 gguf: reorder permutation out of range"); + std::memcpy(o.bytes.data() + static_cast(dst) * row_bytes, + tensor.data + static_cast(src_row) * row_bytes, + row_bytes); + } + return o; +} + bool HasTensor(const GgufFile& g, const std::string& name) { for (const GgufTensorInfo& t : g.Tensors()) { if (t.name == name) return true; @@ -1046,11 +1107,27 @@ GdnLayerWeights LoadGdnGguf(const GgufFile& g, int64_t il, const HfConfig& c, GdnLayerWeights gdn; // in_proj_qkv <- attn_qkv [conv_dim, H]; only the trailing V rows reorder. + // GFX1100-TG150: when the V-row reorder is active this used to force bf16 + // expansion (kTransformedWeight never keeps blocks), landing the decode + // in_proj on hipBLASLt (~3.4 ms/token across the 24 GDN layers). A row + // permutation cannot cut a ggml K-block — each row is whole blocks — so the + // reorder is applied to the quantized rows directly and the blocks stay + // resident. The resident weight encodes exactly what dequantize → reorder → + // bf16 encoded; only the storage dtype differs. { const std::string nm = Blk(il, "attn_qkv.weight"); - const GgufResidency r = pol.Route(g.Get(nm), proj_role); - if (r != GgufResidency::kExpandBf16) { - const GgufTensorInfo& ti = g.Get(nm); + const GgufTensorInfo& ti = g.Get(nm); + const bool row_reorder = reorder && ti.ggml_type != 0; + const GgufTensorRole route_role = row_reorder + ? GgufTensorRole::kMatmulWeight + : proj_role; + const GgufResidency r = pol.Route(g.Get(nm), route_role); + if (r == GgufResidency::kKeepQuant && row_reorder) { + gdn.in_proj_qkv = OwnGgufQuantBlocksRowPermuted( + ti, ti.shape[0], ti.shape[1], + VRowPermutation(ti.shape[0], /*row_off=*/2 * key_dim, num_k, rpk, + dv)); + } else if (r != GgufResidency::kExpandBf16) { gdn.in_proj_qkv = OwnGgufKeptSlice(g, pol, ti, r, ti.shape[0], ti.shape[1], 0); } else { @@ -1062,12 +1139,21 @@ GdnLayerWeights LoadGdnGguf(const GgufFile& g, int64_t il, const HfConfig& c, gdn.in_proj_qkv = MakeGdnProj(dq, out_dim, in_dim, pol.gdn_expand_nk); } } - // in_proj_z <- attn_gate [value_dim, H]; all rows are V. + // in_proj_z <- attn_gate [value_dim, H]; all rows are V. Row-permuted + // keep-quant, same reasoning as in_proj_qkv above (GFX1100-TG150). { const std::string nm = Blk(il, "attn_gate.weight"); - const GgufResidency r = pol.Route(g.Get(nm), proj_role); - if (r != GgufResidency::kExpandBf16) { - const GgufTensorInfo& ti = g.Get(nm); + const GgufTensorInfo& ti = g.Get(nm); + const bool row_reorder = reorder && ti.ggml_type != 0; + const GgufTensorRole route_role = row_reorder + ? GgufTensorRole::kMatmulWeight + : proj_role; + const GgufResidency r = pol.Route(g.Get(nm), route_role); + if (r == GgufResidency::kKeepQuant && row_reorder) { + gdn.in_proj_z = OwnGgufQuantBlocksRowPermuted( + ti, ti.shape[0], ti.shape[1], + VRowPermutation(ti.shape[0], /*row_off=*/0, num_k, rpk, dv)); + } else if (r != GgufResidency::kExpandBf16) { gdn.in_proj_z = OwnGgufKeptSlice(g, pol, ti, r, ti.shape[0], ti.shape[1], 0); } else { diff --git a/src/vt/rocm/rocm_matmul_hipblaslt.hip b/src/vt/rocm/rocm_matmul_hipblaslt.hip index a1d043702..0ce986e6d 100644 --- a/src/vt/rocm/rocm_matmul_hipblaslt.hip +++ b/src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -24,6 +24,9 @@ #include #include #include +#include +#include +#include #include "vt/ops.h" #include "vt/rocm/rocm_device_bind.h" @@ -237,6 +240,78 @@ void Bf16GemvBT(hipStream_t s, void* out, const void* a, const void* b, int N, i } } +// F32-in/f32-out variant for tiny decode GEMVs whose activation stays f32 +// (the GDN a/b projections). Same row-per-block geometry as Bf16GemvBTRowKernel. +template +__global__ void GemvBTF32OutKernel(float* __restrict__ y, + const XT* __restrict__ x, + const __hip_bfloat16* __restrict__ W, int N, int K, + float alpha, float beta) { + const int n = static_cast(blockIdx.x); + if (n >= N) return; + extern __shared__ float smem[]; + float* x_cache = smem; + const int tid = static_cast(threadIdx.x); + constexpr int kBlock = 256; + for (int k = tid; k < K; k += kBlock) { + if constexpr (std::is_same_v) x_cache[k] = x[k]; + else x_cache[k] = __bfloat162float(x[k]); + } + __syncthreads(); + + const __hip_bfloat16* wrow = W + static_cast(n) * static_cast(K); + float acc = 0.f; + const int K16 = K & ~15; + for (int k = tid * 16; k < K16; k += kBlock * 16) { +#pragma unroll + for (int t = 0; t < 16; ++t) acc += x_cache[k + t] * __bfloat162float(wrow[k + t]); + } + for (int k = K16 + tid; k < K; k += kBlock) acc += x_cache[k] * __bfloat162float(wrow[k]); + + __shared__ float red[256]; + red[tid] = acc; + __syncthreads(); +#pragma unroll + for (int s = 128; s > 0; s >>= 1) { + if (tid < s) red[tid] += red[tid + s]; + __syncthreads(); + } + if (tid == 0) { + float v = alpha * red[0]; + if (beta != 0.f) v += beta * y[n]; + y[n] = v; + } +} + +// Dispatch the f32-out GEMV when shared memory allows the cached-x form. +void GemvBTF32OutF32X(hipStream_t s, void* out, const void* a, const void* b, int N, int K, + float alpha, float beta) { + constexpr int kBlock = 256; + const size_t shmem = static_cast(K) * sizeof(float); + GemvBTF32OutKernel<<(N), kBlock, shmem, s>>>( + static_cast(out), static_cast(a), + static_cast(b), N, K, alpha, beta); +} +void GemvBTF32OutBf16X(hipStream_t s, void* out, const void* a, const void* b, int N, int K, + float alpha, float beta) { + constexpr int kBlock = 256; + const size_t shmem = static_cast(K) * sizeof(float); + GemvBTF32OutKernel<__hip_bfloat16><<(N), kBlock, shmem, s>>>( + static_cast(out), static_cast(a), + static_cast(b), N, K, alpha, beta); +} +void GemvBTF32Out(hipStream_t s, void* out, const void* a, const void* b, int N, int K, + float alpha, float beta, bool a_is_f32) { + constexpr int kBlock = 256; + if (K <= 0 || N <= 0) return; + const size_t shmem = static_cast(K) * sizeof(float); + if (shmem > 48 * 1024) return; // caller falls back to the library path + if (a_is_f32) + GemvBTF32OutF32X(s, out, a, b, N, K, alpha, beta); + else + GemvBTF32OutBf16X(s, out, a, b, N, K, alpha, beta); +} + // VT_ROCM_HIPBLASLT=1 enables. Default OFF — heuristic path aborted on gfx1201 in lab. bool LtEnabled() { static const bool on = [] { @@ -470,6 +545,15 @@ void MatmulBTKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) } const int64_t M = a.shape[0], K = a.shape[1], N = b.shape[0]; if (M == 0 || N == 0) return; + if (std::getenv("VT_MM_TRACE") && M == 1) { + static std::map, int> seen; + std::array key{N, K, static_cast(a.dtype)}; + if (++seen[key] == 1 || seen[key] % 200 == 0) + std::fprintf(stderr, "[mmbt] N=%lld K=%lld a=%d b=%d out=%d count=%d\n", + static_cast(N), static_cast(K), + static_cast(a.dtype), static_cast(b.dtype), + static_cast(out.dtype), seen[key]); + } hipStream_t s = static_cast(q.handle); if (K == 0) { CheckHip(hipMemsetAsync(out.data, 0, out.Bytes(), s), "bt k0"); @@ -483,10 +567,27 @@ void MatmulBTKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) } // Decode: M=1 BF16 GEMV - if (M == 1 && bf16 && out.dtype == DType::kBF16 && a.stride[0] == K && GemvEnabled()) { + // GFX1100-TG150: tiny-N decode GEMVs (the GDN a/b projections, N=32) must + // not fall to hipBLASLt — its fixed launch+epilogue cost measured ~78us per + // call on the RX 7900 XTX for what is a 164 KB weight read (~4.1 ms/token + // across the model). For small N the row-per-block GEMV below wins + // regardless of the global VT_ROCM_GEMV A/B (which was decided on large-N + // shapes), so it is forced here. + if (M == 1 && bf16 && out.dtype == DType::kBF16 && a.stride[0] == K && + (GemvEnabled() || N <= 64)) { Bf16GemvBT(s, out.data, a.data, b.data, static_cast(N), static_cast(K), 1.f, 0.f); return; } + // F32-out twin: the a/b projections keep an f32 output ("g/beta stay f32", + // FLA split); without this arm they fall to hipblasGemmEx → hipBLASLt with + // the same ~78us fixed cost per call. + if (M == 1 && out.dtype == DType::kF32 && + (a.dtype == DType::kF32 || a.dtype == DType::kBF16) && + b.dtype == DType::kBF16 && N <= 64 && a.stride[0] == K) { + GemvBTF32Out(s, out.data, a.data, b.data, static_cast(N), + static_cast(K), 1.f, 0.f, a.dtype == DType::kF32); + return; + } // hipBLASLt (BF16 contiguous) if (bf16 && out.dtype == DType::kBF16 && a.stride[0] == K &&