diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 9a66f7d73..70289a628 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -527,3 +527,7 @@ 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 | +| [#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/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. 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. 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/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 && diff --git a/src/vt/rocm/rocm_paged_attn.hip b/src/vt/rocm/rocm_paged_attn.hip index 001ca79f6..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. @@ -1751,8 +1909,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 +2092,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(), @@ -1979,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); diff --git a/src/vt/rocm/rocm_quant_dot.hip b/src/vt/rocm/rocm_quant_dot.hip new file mode 100644 index 000000000..c240c3e15 --- /dev/null +++ b/src/vt/rocm/rocm_quant_dot.hip @@ -0,0 +1,1404 @@ +// 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 +#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) { + // 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; + const uint8_t* q3 = xb->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); + 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) { + 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; + 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) { + // 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; + 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]; + 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]; +#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; + 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) { + // 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 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]; + 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]; +#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; + 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) { + // 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}; +#pragma unroll + 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); + } + } + 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); +} + +// --------------------------------------------------------------------------- +// 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 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 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); +} + + +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) + ? 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; + + // 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 < 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); + + 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); + } + } +} + +// --------------------------------------------------------------------------- +// 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, + 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) { + // 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 = + 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, + 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/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); +} 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); +}