diff --git a/.agents/specs/rocm-qwen35-08b-cpu-gfx1100-numerics.md b/.agents/specs/rocm-qwen35-08b-cpu-gfx1100-numerics.md new file mode 100644 index 000000000..c4b7ddff8 --- /dev/null +++ b/.agents/specs/rocm-qwen35-08b-cpu-gfx1100-numerics.md @@ -0,0 +1,496 @@ +# Qwen3.5-0.8B CPU and gfx1100 numerical characterization + +Row: `BACKEND-ROCM`. + +Issue: [#2773](https://github.com/mudler/vllm.cpp/issues/2773), which supersedes +issue #1588 for implementation traceability. + +Branch: `row/BACKEND-ROCM-NUMERICS-1588`. + +Active primary oracle pin: vLLM +`e126687a9a828d513c01a07cd69f025f27d63280`. + +## Now + +`PENDING`. This repaired spec is the only completed deliverable in this commit. +Implementation cannot start until a fresh reviewer passes this immutable spec. +Results cannot be accepted until both correctness prerequisites pass: + +1. Issue #2772 and PR #2856 own the default-on ROCm `wvSplitK` sacred-anchor + repair. The reviewed old base failed prompt 10/token 10 (`369` instead of + `488`), while `VT_ROCM_SKINNY=0` passed. The operator reran PR #2856 at + `f06619e4c213e3de28359ee10995e682e8c06932`: CPU mode 2/2, prerequisites + 77/77/77/86, `wvSplitK` 79796/79796, and sacred 137/137 with 15 strict, one + tied, maximum gap zero, and zero divergence. PR #2856 is still open and the + contributor lacks merge authority. This prerequisite remains `PENDING` + until a maintainer lands it and the unchanged default gate passes on the + implementation base. Disabling skinny GEMM is diagnostic only. +2. The configured source and runnable ROCm wrapper are still at historical + vLLM `5559679229bc961848b121ccdeaa8fa5d79bec98`. The active source object is + available for read-only inspection, but no runnable active-pin ROCm runtime + has been supplied or proved. Issue #2794 records repository pin-validation + context; #2773 itself owns the cache-matched active-pin Qwen3.5-0.8B captures. + Active-pin capture and token revalidation remain `PENDING` until that runtime + exists and runs the model under the GPU mutex. + +Do not create characterization goldens from a known-regressed local default or +from the historical oracle revision. + +## Scope + +### In scope + +- Add opt-in production-boundary dumps for the residual stream, attention + output, MLP output, stored K/V cache, and persistent GDN convolution/SSM + state. +- Compare CPU and gfx1100 on identical weights and identical token prefixes. +- Report descriptive `max_abs`, RMS, and relative-L2 deltas. +- Port applicable active-pin upstream operation tests with their parameters, + modes, fixtures, failure cases, and exact tolerances. +- Audit runtime dtype, byte width, requested/resolved cache dtype, selected + attention backend, and provider counts for `auto`, `bfloat16`, and + `fp8_e4m3`. +- Run cache-matched active-pin vLLM captures and the established end-to-end + oracle-backed token/near-tie gate after the prerequisites pass. + +### Out of scope + +- Changing a numerical tolerance after observing results. +- Treating CPU as the conformance oracle for FP8 physical-cache behavior. +- Permanently expanding the sacred 16-prompt gate merely because this + characterization measures three cache modes. A new permanent case or golden + may land only when #2773's active-pin capture proves it is necessary and the + final implementation records the evidence. Broader gate policy needs its own + issue. +- Enabling ROCm static graph mode; issue #332 owns that work. +- Performance claims or tuning. + +## Artifact and oracle pins + +Use `Qwen/Qwen3.5-0.8B@2fc06364715b967f1860aea9cf38778875588b17` +from `${CHECKPOINT_ROOT}/Qwen3.5-0.8B`. Before every model run, require: + +```sh +sha256sum \ + "${CHECKPOINT_ROOT}/Qwen3.5-0.8B/model.safetensors-00001-of-00001.safetensors" \ + "${CHECKPOINT_ROOT}/Qwen3.5-0.8B/model.safetensors.index.json" \ + "${CHECKPOINT_ROOT}/Qwen3.5-0.8B/config.json" +``` + +Expected hashes are: + +```text +04b1c301231dd422b8860db31311ab2721511346a32cb1e079c4c4e5f1fe4696 model.safetensors-00001-of-00001.safetensors +d8a08838a613b025eb7952ed9db11696213e57e76a375661ef5c12f9dd5dcf4e model.safetensors.index.json +b90b86f35c8e6925ef74ee04d0e758f0a845c83a42089ad82bbaa948de9b4204 config.json +``` + +Resolve the active oracle source with: + +```sh +git -C "${VLLM_SOURCE}" rev-parse e126687a9a828d513c01a07cd69f025f27d63280^{commit} +``` + +The configured checkout may remain detached at the historical revision. Read +active-pin files with `git -C "${VLLM_SOURCE}" show :`; do not move +that checkout. Source availability does not prove runtime gateability. + +## Upstream executing chain + +All cross-file anchors name symbols and refer to active pin `e126687a9` unless +explicitly labeled historical: + +- `vllm/model_executor/models/qwen3_5.py::Qwen3_5ForCausalLM` defines the text + model and layer composition. +- `vllm/model_executor/layers/mamba/mamba_utils.py::MambaStateShapeCalculator.gated_delta_net_state_shape` + defines the GDN convolution and SSM state shapes. It starts at line 258 at + the active pin. `MambaStateShapeCalculator.mamba2_state_shape` is a different + architecture and is not the anchor. +- `tests/kernels/mamba/cpu/test_cpu_gdn_ops.py::test_fused_sigmoid_gating_delta_rule_update_cpu` + and `test_chunk_gated_delta_rule_cpu` define CPU operation tolerances. +- `tests/kernels/mamba/test_gdn_forward_core_split.py::test_forward_core_split_matches_unified` + defines split/unified state and output comparisons. +- `tests/kernels/attention/test_cache.py::test_reshape_and_cache` defines + cache-store dtype behavior and FP8 comparison tolerances. +- `tests/kernels/attention/test_cpu_attn.py::varlen_with_paged_kv` starts at + line 415. Its `_FP8_ATOL` and `_FP8_RTOL` constants at lines 45 and 46 + set `atol` to `0.2` for E4M3 and `0.3` for E5M2, with `rtol=0.1`. + The comparisons at lines 629 to 653 select these bounds for FP8 and + `atol=1.5e-2`, `rtol=1e-2` otherwise. Preserve both split and unsplit checks. + +The earlier spec cited these surfaces from historical `555967922`; those +citations remain diagnostic history but cannot satisfy the active-pin gate. The +implementer must refresh line evidence from the active object and run the active +runtime before acceptance. + +The local production entry is `ModelRegistry::Forward`, which dispatches the +registered `ForwardQwen3_5Dense` factory. Its paged forward reaches +`DenseForwardBody` -> `DenseForwardLayers` -> `RunDenseLayerPaged` -> +`FullAttnBlockPaged` -> `dense_attn::WriteKvCache` -> `vt::ReshapeAndCache` or +`vt::ReshapeAndCacheFp8`. The full-attention caller is +`src/vllm/model_executor/models/qwen3_5.cpp:7758`. `FullAttnBlockPaged` starts +at line 5726 and writes through `dense_attn::WriteKvCache` at +`src/vllm/model_executor/models/qwen3_5.cpp:5909`, before `vt::PagedAttention` +at line 5931. Place the stored-K/V probes after that write. Place GDN probes +after the persistent writes in `GdnBlockPaged` for prefill and decode. + +This Qwen3.5 path does not call `dense_attn::AttnBlock`. Issue #2923 under +`Owed` records this existing shared-seam debt. Characterize the actual path +and cite the executing local and active-pin upstream symbols in the evidence. + +## Dtype and state contract + +The checkpoint resolves model dtype BF16 and explicitly sets +`mamba_ssm_dtype=float32`. The measured surfaces are: + +| Surface | Shape per active unit | Physical storage | +|---|---|---| +| Hidden/residual | `[T, 1024]` | BF16 | +| Full-attention K/V | `[T, 2, 256]` per layer | BF16 or FP8 E4M3 | +| GDN convolution state | Per layer and slot: oracle SD `[3, 6144]`, local DS `[6144, 3]` | BF16 | +| GDN SSM state | `[1, 16, 128, 128]` per layer | FP32 | + +The active oracle defaults to SD physical convolution storage, with axes +`(state_len, dim)`. DS reverses these axes to `(dim, state_len)`. +`vllm/model_executor/layers/mamba/mamba_utils.py::get_conv_state_layout`, lines +28 to 44, returns `SD` without an override. `_orient_conv_shape`, lines 162 +to 166, and `gated_delta_net_state_shape`, lines 268 to 272, produce +`[3, 6144]` per slot for this workload. The kernel consumes a DS view +`[6144, 3]` after the transpose in +`vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py:1296` to line 1303. +Local `MakeQwen3_5KVCacheSpec` stores DS directly in +`src/vllm/model_executor/models/qwen3_5_common.cpp:85`. Leave +`VLLM_SSM_CONV_STATE_LAYOUT` unset for the production oracle. Record physical +layout, kernel-view layout, and actual strides on both sides. A transposed +view does not change the underlying physical layout or byte count. + +Across six full-attention layers, BF16 K/V use 12,288 bytes per cached token +and FP8 K/V use 6,144 bytes. Across 18 GDN layers, convolution state uses +663,552 bytes and SSM state uses 18,874,368 bytes per active slot. Confirm +allocated bytes from runtime specs; arithmetic alone is not evidence. + +For every compared buffer and GEMM output, log the local symbol, active-pin +upstream symbol, runtime dtype, shape, active elements, bytes on each side, and +the reason beside every FP32 exception. An unjustified wider local dtype is a +`DTYPE_WIDTH_MISMATCH` even when tokens agree. + +## State dump design + +Enable dumps only when `VT_DUMP_ACT` and `VT_DUMP_ACT_SUB` name the same empty, +writable directory. Preserve existing keys and add: + +| Key | Production source | Dump dtype | Shape | +|---|---|---|---| +| `state_fa_k` | stored K after `dense_attn::WriteKvCache` | FP32 | `[active_tokens, 512]` | +| `state_fa_v` | stored V after `dense_attn::WriteKvCache` | FP32 | `[active_tokens, 512]` | +| `state_gdn_conv` | persistent convolution cache after update | FP32 | `[active_requests, 18432]` | +| `state_gdn_ssm` | persistent SSM cache after update | FP32 | `[active_requests, 262144]` | + +Widen stored BF16 to FP32. Dequantize stored FP8 with the recorded per-layer +scale. The FP32 file format never conceals separately logged physical dtype and +allocation. + +Use canonical DS order for `state_gdn_conv`, with `0 <= d < 6144` and +`0 <= s < 3`: + +```text +local_dump[r, d * 3 + s] = widen(local[slot_local(r), d, s]) +oracle_dump[r, d * 3 + s] = widen(oracle[slot_oracle(r), s, d]) +``` + +Here `r` follows scheduler request order, and each arm resolves its own slot. +Use the recorded strides and storage offsets for each lookup. Record this +mapping in the manifest. The canonical DS dump order does not prescribe +physical storage or make the oracle's transposed view contiguous. + +Gather nonnegative slot mappings in input-token order and active GDN indices in +scheduler request order. Refuse unexplained duplicate destinations, missing +rows, partial joins, or capacity-wide dumps. A 24-layer step emits 48 state +blobs, two per layer. Place each probe after its production write and preserve +queue ordering. + +## Workload and backend identity + +Use one binary and checkpoint copy. Run one request at concurrency one with +`The capital of France is`, greedy sampling, seed 0, MTP disabled, and eight +output tokens. Keep block size, block count, and scheduler budget equal. Run +each arm and cache mode in a separate process and empty directory. Take two +enabled repeats and one dump-disabled identity control. + +Log and assert requested cache dtype, resolved storage and FP8 interpretation, +normalized selector input, physical cache dtype, selected attention backend, +and actual provider execution for every arm and mode. +`LoadedEngine::ApplyResolvedCacheDType` in +`src/vllm/entrypoints/model_loader.cpp:1835` applies `ParseCacheDType` from +`include/vllm/v1/kv_cache_dtype.h:58` to the cache spec. The runner constructs +`cfg.kv_cache_dtype` from resolved storage and interpretation at +`src/vllm/v1/worker/gpu/runner.cpp:1590`. `KvCacheDTypeName` maps `kBF16` to +`auto` in `src/vllm/v1/attention/backend.cpp:91` to line 96. Therefore the +production CPU selector receives these values for this BF16 workload: + +| Requested cache dtype | Resolved physical storage | Normalized selector input | Selected CPU backend | +|---|---|---|---| +| `auto` | BF16 (`kBF16`) | `auto` | `CPU_ATTN` | +| `bfloat16` | BF16 (`kBF16`) | `auto` | `CPU_ATTN` | +| `fp8_e4m3` | E4M3 bytes (`kI8` plus `kFp8E4M3`) | `fp8_e4m3` | `CPU_ATTN` | + +CPU priority starts with `CPU_ATTN` in `src/vllm/platforms/cpu.cpp:51`. +Its accepted selector values include `auto` and `fp8_e4m3` in +`include/vllm/v1/attention/backends/cpu_attn.h:108` to line 110. +`SelectAttentionBackendName` returns the first valid candidate in +`src/vllm/v1/attention/registry.cpp:124` to line 128. +Selection alone does not prove execution. Independently record the executing +attention and cache-store operator providers, selection counts, declines, and +reference-tier hits. The runner's selected name validates cache configuration +but does not dispatch the model's attention call. If main changes this chain +before implementation, update the cited contract before code and obtain fresh +review. + +CPU is a diagnostic comparison arm for CPU-versus-ROCm deltas, including FP8 +state after dequantization. Active-pin vLLM and its upstream tests are the +conformance authority. CPU results cannot accept or reject FP8 physical-cache +semantics. + +Join only steps with identical input token prefixes. Include the first divergent +output step and exclude all later steps. Record the first divergence and number +of excluded steps. + +## Descriptive metrics and acceptance + +Decode joined blobs to FP32. Let `A` be ROCm and `B` be CPU. Report: + +```text +max_abs = max_i(abs(A_i - B_i)) +rms = sqrt(sum_i((A_i - B_i)^2) / N) +rel_l2 = sqrt(sum_i((A_i - B_i)^2)) / sqrt(sum_i(B_i^2)) +``` + +Define relative L2 as zero when both norms are zero and infinity when only `B` +has zero norm. Report the worst index, both manifests, joined/rejected rows, and +excluded steps. Refuse missing or duplicate keys, dtype/shape/byte mismatches, +nonfinite values, and partial joins. + +These layer metrics are descriptive. They have no propagated numerical +acceptance envelope, discontinuity multiplier, or post-measurement threshold. +The rejected `E(k)` construction is removed because operation counts do not +bound cancellation, conditioning, nonlinear sensitivity, correlated +reductions, FP8 scaling/saturation, or subnormals. + +Acceptance consists of both: + +1. Ported active-pin operation tests with their exact upstream per-surface + tolerances. Examples include GDN `atol=rtol=1e-2` where + `test_cpu_gdn_ops.py` specifies it, exact convolution-state equality in the + split-core test, its dtype-specific output/SSM tolerances, and cache/CPU + attention FP8 tolerances selected by the upstream fixtures. Do not collapse + them into one project-wide tolerance. +2. The established end-to-end 16-prompt strict-token and ratified 500 + milli-nat near-tie gate against deterministic, cache-matched, active-pin + vLLM captures. A new physical cache mode needs its own capture and + teacher-forced gaps before it can be accepted. + +If no upstream analogue exists for a layer boundary, report its metrics only. +A future stage-level acceptance rule requires an independently justified bound +or explicit developer ratification before values are inspected. + +## Provider and end-to-end contract + +For `auto` and `bfloat16`, require the established native operator set, +including nonzero `kReshapeAndCache`, zero `kReshapeAndCacheFp8`, zero declines, +and zero reference-tier hits. The two modes resolve to physical BF16 and must +produce identical local token streams. + +For `fp8_e4m3`, replace the store requirement: require nonzero +`kReshapeAndCacheFp8`, zero `kReshapeAndCache`, nonzero native selections for +all other applicable operators, zero declines, and zero reference-tier hits. +Provider sets are mode-specific; never require a BF16 store from a correct FP8 +run. + +The existing gate remains the permanent gate until #2773 produces reviewed +active-pin evidence for an additional case. Characterization must run all three +modes, but this spec alone does not authorize new permanent goldens. + +## Required capture-tool implementation + +The current `scripts/qwen3-oracle-capture.py` and +`scripts/qwen3-neartie-gap.py` hard-code `enforce_eager=True` and expose no +cache-dtype argument. Before active-pin capture, add reviewed options equivalent +to: + +```text +--kv-cache-dtype {auto,bfloat16,fp8_e4m3} +--execution-mode {production,eager} +--seed INT +--max-tokens INT +--repetitions INT +--model-revision REV +--vllm-revision REV +--provenance-out PATH +``` + +`production` must instantiate vLLM without `enforce_eager=True`; eager is a +diagnostic arm and never the denominator. The output must record source, wheel, +image and artifact hashes, complete arguments, a hash of the scripts' shared +16-entry `PROMPTS` list, batching, concurrency, sampling, seed, token count, +repetitions, cache mode, resolved +cache dtype, execution mode, and output hash. + +Extend `test_qwen35_paged_engine` with future environment inputs +`VT_QWEN35_GATE_DIR` and `VT_QWEN35_KV_CACHE_DTYPE`. The first selects an empty +issue-evidence directory instead of the committed golden directory. The second +sets `EngineParams::kv_cache_dtype` and prints requested, resolved, and physical +dtype. These commands describe the intended interface and data flow after +those changes. They are future commands and will fail today: + +```sh +VLLM_PIN=e126687a9a828d513c01a07cd69f025f27d63280 +MODEL_REV=2fc06364715b967f1860aea9cf38778875588b17 +MODEL="${CHECKPOINT_ROOT}/Qwen3.5-0.8B" +for mode in auto bfloat16 fp8_e4m3; do + GOLDEN_DIR="evidence/2773/sacred-${mode}" + test ! -e "${GOLDEN_DIR}" + mkdir -p "${GOLDEN_DIR}" + "${VLLM_ORACLE}" scripts/qwen3-oracle-capture.py \ + --model "${MODEL}" --model-revision "${MODEL_REV}" \ + --vllm-revision "${VLLM_PIN}" \ + --kv-cache-dtype "${mode}" --execution-mode production --seed 0 \ + --max-tokens 16 --runs 10 --per-prompt --out-dir "${GOLDEN_DIR}" \ + --provenance-out "evidence/2773/oracle-${mode}.json" + VT_QWEN35_GATE_DIR="${GOLDEN_DIR}" \ + VT_QWEN35_KV_CACHE_DTYPE="${mode}" \ + VT_DUMP_IDS=1 build-rocm/tests/test_qwen35_paged_engine + test -s "${GOLDEN_DIR}/our_ids.i32" + "${VLLM_ORACLE}" scripts/qwen3-neartie-gap.py \ + --model "${MODEL}" --model-revision "${MODEL_REV}" \ + --vllm-revision "${VLLM_PIN}" \ + --kv-cache-dtype "${mode}" --execution-mode production --seed 0 \ + --max-tokens 16 --topk 20 --golden-dir "${GOLDEN_DIR}" \ + --provenance-out "evidence/2773/neartie-${mode}.json" + VT_QWEN35_GATE_DIR="${GOLDEN_DIR}" \ + VT_QWEN35_KV_CACHE_DTYPE="${mode}" \ + build-rocm/tests/test_qwen35_paged_engine +done +``` + +The tools must refuse when their `PROMPTS` lists differ from each other or from +`tests/parity/test_qwen35_paged_engine.cpp::Prompts`. Record exact as-run +commands. Do not create a golden unless all 10 repeats are deterministic. + +The internal CPU/ROCm state characterization remains eight output tokens as +specified under `Workload and backend identity`. The permanent sacred-gate +candidate uses the existing 16-prompt, 16-output-token regime above. Never use +the eight-token characterization files as sacred-gate goldens. + +## Tests and review mutations + +The later implementation starts with focused tests that fail because the four +state rows and capture options are absent. It must test row shape/dtype, +nonmonotonic and negative slot mappings, inactive-capacity exclusion, stored +BF16 widening, post-quantization FP8 dequantization, persistent GDN prefill and +decode state, incomplete-step refusal, comparator structural refusals and +zero-denominator rules, dtype-width refusal, backend/dtype logging, and +mode-appropriate provider sets. The checkpoint-backed case must enter through +`LoadedEngine::FromModelDir`. + +A fresh reviewer mutates each guarantee in a scratch copy: remove each +production probe, substitute pre-write tensors, remove negative-slot filtering, +sort by physical slot, dump capacity, misreport FP8 as BF16, swap the mode's +cache-store operator, inject a reference-tier hit, remove the production entry +point, and force eager oracle mode. Each focused test must fail for the intended +reason and the reviewer must restore the tree byte for byte. + +This spec repair changes no runtime behavior. `IMP-TEST-FIRST` and +`IMP-MUTATE` are future implementation/review gates, not evidence claimed by +this commit. + +## Gate order + +The later implementation must satisfy, in order: + +1. #2772 lands and the unchanged default local gate passes on the chosen base. +2. A runnable active-pin ROCm vLLM runtime is identified and proves this model. +3. Artifact, source, wrapper, wheel, image, prompt, and output hashes are saved. +4. The capture tools fail first for missing options, then pass focused tests. +5. Ported active-pin operation tests pass with unchanged upstream tolerances. +6. Dump/comparator CPU tests and checkpoint-backed CPU trace pass. +7. The controlled full preflight runs; every skip remains `PENDING`. +8. Fresh immutable review detects every required mutation. +9. Two CPU and two gfx1100 trace repeats plus disabled controls pass per mode. +10. Cache-matched active-pin production captures are deterministic and the + established end-to-end gate passes per measured mode. +11. Mode-appropriate native selections are nonzero, with zero declines and + zero reference-tier hits. +12. The operator independently reruns focused, full, oracle, and hardware gates. + +No GPU result, oracle result, model execution, runtime mutation, or +implementation test is claimed by this spec-only repair. + +## Evidence required + +Store evidence under an issue-specific durable path and record commit/tree +hashes, clean status, exact commands/statuses, every skip, active oracle and +artifact hashes, compiler/build/ROCm/driver/board identity, GPU mutex evidence, +requested/resolved/physical cache dtype, selected CPU and ROCm attention +backends, provider counts, enabled repeats, disabled controls, manifests, raw +dumps, metric tables, first divergence, excluded steps, upstream tolerance +anchors, red-first output, mutations, fresh review, and operator reruns. + +Classify structural and execution failures before interpreting metrics: +`ARTIFACT_MISMATCH`, `INSTRUMENTATION_FAIL`, `NONDETERMINISTIC`, +`STRUCTURE_MISMATCH`, `DTYPE_WIDTH_MISMATCH`, `CORRECTNESS_FAIL`, or +`NONFINITE`. When none applies, report descriptive metrics and the separate +operation/end-to-end acceptance results. Do not invent `WITHIN_DTYPE_ENVELOPE` +or `ORDERING_DRIFT` labels. + +## Risks + +- Probes can synchronize queues; disabled controls detect output perturbation. +- FP8 boundaries can change codes; compare dequantized values and record bytes. +- Different valid reduction orders can produce descriptive CPU/ROCm deltas. +- Backend fallback can change the diagnostic denominator; explicit backend + logging exposes it. +- Greedy divergence makes later states incomparable; exclude them. +- The active source object can be read while the active runtime remains absent; + never convert source availability into a gate pass. + +## Stop conditions + +- Stop before implementation until this spec receives a fresh `PASS`. +- Keep correctness `PENDING` until #2772 lands and the unchanged default gate + passes on the implementation base. +- Stop active-pin capture until a runnable active-pin ROCm runtime is proved. +- Stop on a revision/hash/configuration mismatch, dump perturbation, + nondeterminism, incomplete manifest, unexplained backend, provider decline, + reference-tier hit, wrong board, missing GPU mutex, or divergent input prefix. +- Stop before changing tolerances or permanent sacred-gate scope after seeing + results. +- File and assign a new issue before fixing any unexpected defect outside + #2773. + +## Git integration + +Use one pull request for the committed spec and later implementation, following +the recorded repository default. Preserve rejected spec commit `7bc2546e9` in +history. The eventual pull request body must name row `BACKEND-ROCM`, link and +close #2773, and carry the required trailers. The spec implementer does not +push, open, or merge that pull request. + +## Owed + +- A maintainer owes the merge decision for reviewed PR #2856; #2772 remains + pending until it lands. +- #2773 owes the runnable active-pin Qwen3.5-0.8B ROCm captures, even though + #2794 supplies repository sync context. +- [#2923](https://github.com/mudler/vllm.cpp/issues/2923), owned by + `BACKEND-ROCM`, owes routing Qwen3.5 paged attention through + `dense_attn::AttnBlock`. This tracked exception records existing debt. It + does not waive or satisfy the shared-seam requirement. Wiring needs its own + reviewed spec and implementation, outside #2773's instrumentation scope. +- A fresh reviewer owes this repaired spec a verdict. +- A fresh implementer owes the capture options, state probes, tests, comparator, + and provider checks after the prerequisites pass. +- A fresh implementation reviewer owes static review and every mutation. +- The operator owes the independent controlled and gfx1100 gates. +- The final implementation adds `## Outcome` with measured results, rejected + alternatives, and reasons for defaults. +- Issue #332 owns the graph-enabled repetition. diff --git a/CMakeLists.txt b/CMakeLists.txt index 57af96778..62ea46813 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1738,7 +1738,8 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_mla_ops.hip src/vt/rocm/rocm_mla_attn.hip src/vt/rocm/rocm_skinny_gemm.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 @@ -1767,6 +1768,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_mla_attn.hip src/vt/rocm/rocm_skinny_gemm.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/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md b/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md index 7af460e65..b5f7320a8 100644 --- a/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md +++ b/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md @@ -229,84 +229,3 @@ Gate: 16/16, 839 assertions. Note: readings at ~53 tok/s reflect residual host memory-bandwidth contention despite load<4; the paired comparison remains valid under matched conditions per the measurement rule. - -## IDLE-WINDOW ACCEPTANCE GATE + T13 + COPY-STORM ATTRIBUTION (2026-08-26 ~18:55Z) - -### Acceptance gate rerun (load 1.45-2.20, idle host) - -Full 12-lever config (YT4 now default), 6 reps, 256 tokens, seed 0: -- Run 1 (warmup): 90.197 tok/s -- Runs 2-6: 100.534, 100.482, 100.462, 100.392, 100.407 -- **Median: 100.46 tok/s** (runs 2-6, warmup discarded) - -Crossing the 100 tok/s milestone. The YT4 adoption contributes more -under unconstrained memory bandwidth than the contended paired sweep -showed (+1.8% under load → +8.1% idle: 92.8 → 100.4). - -### T13 async-runner paired A/B (idle host, load 1.45) - -OFF median 89.984 vs ON 89.819 (−0.18%, WASH). All 5 pairs byte-identical. -Confirms: the CLI sync loop drains depth-1 regardless of -VT_ASYNC_RUNNER; the batch-queue pipelining only engages under -AsyncScheduler (serving mode). T13 CLOSED for the CLI path. - -### Copy-storm attribution (rocprofv3 trace, 64 tokens) - -318 memory copies total, ALL >64KB. Per-step small copies (160KB×2 + -64KB×1 + 1.4MB every 4 steps) total ~734KB/step at ~35µs/step = **0.035 -ms/tok — NEGLIGIBLE**. The large copies (33MB×76, 20MB×48, etc.) are -model-loading artifacts, not steady-state decode. The "small copy storm" -is CLOSED as a lever — it was a profiling artifact of aggregate counting. - -### Roofline analysis - -Model: 2.74 GB. At 800 GB/s effective, minimum weight read = 3.43 ms/tok. -At 200 tok/s (5.0 ms/tok), leaves 1.57 ms for all compute + attention + -dispatch. Current kernel budget: 8.89 ms/tok (2.6x minimum). The GEMV/GEMM -family accounts for 6.33 ms/tok = 63% of wall. - -| Kernel | ms/tok | % of roofline | headroom | -|---|---|---|---| -| KQuantGemvMmvqK | 2.76 | 78-88% | limited | -| wvSplitKSml<1,4,bf16> | 2.30 | ~57% | **significant** | -| KQuantGemvMmvqK | 1.27 | ~85% | limited | - -**Next attack: wvSplitKSml compute-memory balance.** The inner loop -unpacks bf16→f32 then does 3 FLOPs per pair. RDNA3's v_dot2_f32_bf16 -does this in 1 instruction. If compute is the bottleneck at 57% -bandwidth, dot2 should raise utilization toward 80-90%. - -## T17 v_dot2_f32_bf16 — CLOSED NOT-ADOPTED (2026-08-26, idle host load 0.55) - -### Hypothesis -wvSplitKSml at 57% bandwidth utilization might be compute-bound. The inner -loop does 599 v_mul_f32 + 1158 v_add_f32 = 1757 scalar f32 ops. RDNA3's -v_dot2_f32_bf16 does a.x*b.x + a.y*b.y + c in 1 instruction, replacing 5 -ops per bf16x2 pair. - -### Implementation -Env-gated VT_WVSPLIT_DOT2=1 selects the dot2 MAC path. ISA verified: 1120 -v_dot2_f32_bf16 instructions generated for the ON path. Kernel parameter -threads the flag through WvSplitKBTDispatch. - -### A/B result (idle host, load 0.55, 5 paired runs) -- OFF median: 88.784 tok/s -- ON median: 88.897 tok/s (+0.13%, WASH) -- All 5 pairs DIFFER (reduction order change) -- Both outputs coherent analytic prose (207 vs 223 words, same topic) - -### Root cause: memory-bound, not compute-bound -The kernel is memory-bound at 57% bandwidth. Compute is already fully -hidden behind memory latency. Reducing compute instructions doesn't help -when waiting for memory. Same lesson as T15. - -### Additional finding: runtime branch regression -The OFF arm regressed from 100.4 → 88.8 tok/s (−12%) because the runtime -branch in the inner loop increased code size and register pressure for -both paths. Reverted; 100.47 tok/s confirmed restored post-revert. - -### Decision -CLOSED not-adopted. The dot2 instruction is architecturally correct but -targets the wrong bottleneck. To utilize dot2, the kernel would need to -first become compute-bound (e.g., by increasing memory reuse or reducing -memory traffic), which is a different optimization. diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index ff438c7ad..a636dc31f 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -3,6 +3,7 @@ VLLM_GEMMA4_MM_DEBUG VLLM_MM_TOWER_PROFILE VT_ACT_F32 VT_ARCH_TACTIC_STATS +VT_ASYNC_DEBUG VT_ASYNC_EXECUTOR_NO_DBUF VT_ASYNC_EXECUTOR_POISON VT_ASYNC_EXECUTOR_TRACE @@ -218,6 +219,8 @@ VT_V4_RESIDENT_W VT_V4_ROPE_FLOAT VT_V4_ROUTE_WARP_TOPK VT_W4A4_TRUE +VT_WVSPLIT_PRGRP +VT_WVSPLIT_YTILE # (VT_ASYNC_EXECUTOR_TRACE). Never set in production. # Gemma-4 MoE expert-execution tuning (#154, gemma4_moe.cpp / gemma4.cpp). Each # Inherited ROCm/Gemma-4 bring-up tactics. These are default-off diagnostic, diff --git a/src/vllm/v1/worker/gpu/runner.cpp b/src/vllm/v1/worker/gpu/runner.cpp index 2bd04cdf5..ec37c6976 100644 --- a/src/vllm/v1/worker/gpu/runner.cpp +++ b/src/vllm/v1/worker/gpu/runner.cpp @@ -518,6 +518,11 @@ GPUModelRunner::GPUModelRunner( // scheduling ON for the Eagle-type family (vllm/config/vllm.py:1064-1112). async_sched_supported_ = AsyncRunnerEnvDefault() && QueueSupportsAsyncInputCombine(queue_); + if (std::getenv("VT_ASYNC_DEBUG")) + fprintf(stderr, "[async-debug] env=%d queue_combine=%d -> sched_supported=%d\n", + (int)AsyncRunnerEnvDefault(), + (int)QueueSupportsAsyncInputCombine(queue_), + (int)async_sched_supported_); // ARCH-ONE-SURFACE ROW 6 (mirror gpu/model_runner.py:368-369): a POOLING // model's runner pools instead of sampling — build the PoolingRunner over // the model-owned Pooler. Null for every text arch (byte-identical). @@ -623,6 +628,11 @@ GPUModelRunner::GPUModelRunner( // scheduling ON for the Eagle-type family (vllm/config/vllm.py:1064-1112). async_sched_supported_ = AsyncRunnerEnvDefault() && QueueSupportsAsyncInputCombine(queue_); + if (std::getenv("VT_ASYNC_DEBUG")) + fprintf(stderr, "[async-debug] env=%d queue_combine=%d -> sched_supported=%d\n", + (int)AsyncRunnerEnvDefault(), + (int)QueueSupportsAsyncInputCombine(queue_), + (int)async_sched_supported_); // ARCH-ONE-SURFACE ROW 6 (mirror gpu/model_runner.py:368-369): a POOLING // model's runner pools instead of sampling — build the PoolingRunner over // the model-owned Pooler. Null for every text arch (byte-identical). diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index 42eea3d3b..8b13bec5c 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -215,7 +215,56 @@ class RocmBackend final : public Backend { int DeviceCapabilityMajor() const override { return major_; } int DeviceCapabilityMinor() const override { return minor_; } - // Both branches satisfy the >= 64B contract Backend::Alloc owes StepArena: + // T13 (GFX1100-TG200): REAL event primitives. The base-class defaults are + // NO-OPS (only CudaBackend overrode them), which made the async sampled-id + // delivery read its pinned host buffer BEFORE the D2H copy completed — + // batch-1 greedy decode degenerated to token-0 garbage (the R9700 + // incident's root cause, now fixed at the source). Bodies mirror + // CudaBackend one-to-one via hipEvent*. + Event CreateEvent(bool blocking = false) override { + hipEvent_t ev = nullptr; + const unsigned flags = + hipEventDisableTiming | (blocking ? hipEventBlockingSync : 0u); + Check(hipEventCreateWithFlags(&ev, flags), "hipEventCreateWithFlags"); + return Event{Device{DeviceType::kROCM, device_}, reinterpret_cast(ev)}; + } + void DestroyEvent(Event& e) override { + if (e.handle == nullptr) return; + Check(hipEventDestroy(reinterpret_cast(e.handle)), + "hipEventDestroy"); + e.handle = nullptr; + } + void RecordEvent(Event& e, Queue& q) override { + Check(hipEventRecord(reinterpret_cast(e.handle), AsStream(q)), + "hipEventRecord"); + } + void SynchronizeEvent(Event& e) override { + Check(hipEventSynchronize(reinterpret_cast(e.handle)), + "hipEventSynchronize"); + } + bool QueryEvent(Event& e) override { + const hipError_t status = hipEventQuery(reinterpret_cast(e.handle)); + if (status == hipErrorNotReady) { + (void)hipGetLastError(); + return false; + } + Check(status, "hipEventQuery"); + return true; + } + void QueueWaitEvent(Queue& q, Event& e) override { + Check(hipStreamWaitEvent(AsStream(q), + reinterpret_cast(e.handle), 0), + "hipStreamWaitEvent"); + } + + // T13 (GFX1100-TG200): TRUE. Attempt 1 flipped this while the event + // primitives above were still base-class NO-OPS — the async output read + // its pinned buffer before the D2H landed (token-0 loops, ~2.5x slowdown + // from host spinning). The primitives are now real (this file), so the + // delivery contract pinned-slot -> copy queue -> ready event -> host read + // holds end to end, matching CudaBackend call for call. + bool SupportsAsyncSampledTokenReadback() const override { return true; } + // hipMalloc aligns to at least 256 bytes as cudaMalloc does, and managed // allocations are at least page-granular. void* Alloc(size_t bytes) override { diff --git a/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index bcfa1c848..e82906b9e 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -1320,7 +1320,7 @@ void* Q8KSetKernelExecutionWitnessForTest(void* device_counts) { return previous; } -void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) { +void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) { EnsureQueueDevice(q); const int64_t m = a.shape[0], k = a.shape[1], n = b.shape[0]; if (m == 0 || n == 0) return; @@ -1453,7 +1453,7 @@ void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tenso // kMatmulBTQuantGrouped for ROCm: Q8_0 / Q4_K / Q6_K natively (the formats the // target GDN-MoE GGUFs carry); anything else throws loudly (never a silent // CPU-pointer deref on a discrete card). -void MatmulBTQuantGroupedKernelRocm(Queue& q, Tensor& out, const Tensor& act, +void MatmulBTQuantGroupedKernelRocmGdn(Queue& q, Tensor& out, const Tensor& act, const Tensor& weight, const Tensor& expert_ids) { EnsureQueueDevice(q); const int64_t P = out.shape[0], n = out.shape[1], k = act.shape[1]; diff --git a/src/vt/rocm/rocm_matmul_hipblaslt.hip b/src/vt/rocm/rocm_matmul_hipblaslt.hip index 77deaea03..c9c393173 100644 --- a/src/vt/rocm/rocm_matmul_hipblaslt.hip +++ b/src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -535,6 +535,10 @@ void MatmulKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) { void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, int N, int K, int device); +// Effective YTILE after the T16 sweep knobs (VT_WVSPLIT_YTILE, default 4) — +// the gate must test the SAME tile the kernel will run. +int WvSplitKYtile(); + // Row-major trick: gemm(OP_T, OP_N, N, M, K, B, K, A, a_rs, C, N) // BLAS: C = op(A)*op(B) with opA=T => A is KxN in col form = row B[N,K] // opB=N, B is KxM col = row A[M,K] with ld=a_rs @@ -569,14 +573,16 @@ void MatmulBTKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) // * N > 8 (upstream `m > 8`, utils.py:181 — the feature-dim lower bound; // at N==1 the kernel's y=1 store writes C[1] past the end of a 1-wide // output); - // * N % YTILE == 0 (skinny_gemms.cu:1217 `M_in % _YTILE == 0`; YTILE=2 — - // the last wave's y=1 store lands at C[N] on odd N); + // * N % YTILE == 0 (skinny_gemms.cu:1217 `M_in % _YTILE == 0` — the last + // wave's stores land past C[N] otherwise; the tile is the T16-effective + // one, NOT the donor's fixed 2: with the YT=4 default an N ≡ 2 mod 4 + // shape would write two columns past the row under the old N%2 gate); // * wave32 arch only (the gfx9 ROW_BCAST reduction arm is not ported); // * K % 8 == 0 and the LDS-fit (K*M activation fits the staging buffer); // everything else stays on the BLAS path. if (bf16 && out.dtype == DType::kBF16 && M >= 1 && M <= 4 && (K % 8) == 0 && - N > 8 && (N % 2) == 0 && a.stride[0] == K && K * M <= 32768 && - SkinnyGemmEnabled() && + N > 8 && (N % vt::rocm::WvSplitKYtile()) == 0 && a.stride[0] == K && + K * M <= 32768 && SkinnyGemmEnabled() && vt::rocm::SkinnyGemmArchOk(q.device.index, vt::rocm::DeviceArchName)) { WvSplitKBT(s, out.data, a.data, b.data, static_cast(M), static_cast(N), static_cast(K), q.device.index); diff --git a/src/vt/rocm/rocm_quant_dot.hip b/src/vt/rocm/rocm_quant_dot.hip new file mode 100644 index 000000000..17d1137eb --- /dev/null +++ b/src/vt/rocm/rocm_quant_dot.hip @@ -0,0 +1,1076 @@ +// ROCm keep-quant GGUF k-quant GEMM (KERNEL-QUANT-CIQ-GEMM-ROCM W1) — the +// kROCM provider for `OpId::kMatmulBTQuant` and `OpId::kMatmulBTQuantGrouped`. +// +// Port of src/vt/cuda/cuda_quant_dot.cu (the kCUDA provider), which is itself +// a port of the CPU oracle: +// src/vt/cpu/cpu_quant_gemm.cpp MatmulBTQuantKernel (the GEMM wiring) +// src/vt/cpu/cpu_quant_dot.cpp VecDot{Q2_K,Q3_K,Q4_K,Q5_K,Q6_K, +// IQ2_XXS,IQ3_XXS,IQ2_S,IQ1_S,IQ1_XXXS}Q8_K +// src/vt/cpu/cpu_quant_act.cpp QuantizeRowQ8_K (the activation quant) +// The device numeric helpers below are the SAME bit-exact ports of +// src/vt/dtype.cpp + cpu_quant_act.cpp, so the Q8_K activation bytes — and +// therefore the whole INTEGER dot — are IDENTICAL to the CPU reference. Only +// the per-super-block float scale sum is reassociated (warp reduction vs the +// CPU's sequential add), so the gate is the CUDA sibling's gate: INTEGER core +// bit-exact, final scale within the NMSE band test_ops_quant_dot uses. +// +// * Scratch uses hipMallocAsync (present in ROCm 7.14) with the same retire +// -never-free discipline as the CUDA side (graph_safe_scratch.h). +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "vt/cpu/cpu_quant_blocks.h" // vt::cpu::Block* struct mirror + // (single source; plain C++) +#include "vt/cuda/cuda_quant_iq_tables.cuh" // d_iq2xxs_grid / d_iq3xxs_grid / + // d_iq2s_grid / d_iq1s_grid / + // d_iq1xxxs_grid / d_ksigns / + // d_kmask (single source; pure + // __device__ syntax, no CUDA) +#include "vt/cuda/graph_safe_scratch.h" // RetireGraphScratch (portable) +#include "vt/ops.h" +#include "vt/quant.h" + +// Forward declarations: the optimized K-quant GEMM kernels in +// rocm_grouped_gemm.hip (renamed *Gdn). F1 delegates Q4_K/Q5_K/Q6_K to these +// to preserve the baseline-optimized path; the new kernels in this file +// handle IQ types and Q2_K/Q3_K that the old code did not support. +namespace vt::rocm { +void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, + const Tensor& b); +void MatmulBTQuantGroupedKernelRocmGdn(Queue& q, Tensor& out, const Tensor& act, + const Tensor& weight, + const Tensor& expert_ids); +} + +namespace vt::rocm { +namespace { + +// dtype.cpp F32ToF16 — round-to-nearest-even, subnormals, inf/nan. Used only +// for the Q8_0 activation scale (the CPU Q8_0 vec_dot's f16 round-trip). +__device__ inline uint16_t DF32ToF16(float f) { + uint32_t u = __float_as_uint(f); + uint16_t sign = static_cast((u >> 16) & 0x8000); + int32_t exp = static_cast((u >> 23) & 0xFF) - 127 + 15; + uint32_t mant = u & 0x7FFFFF; + if (((u >> 23) & 0xFF) == 0xFF) + return static_cast(sign | 0x7C00 | (mant ? 0x200 | (mant >> 13) : 0)); + if (exp >= 0x1F) return static_cast(sign | 0x7C00); + if (exp <= 0) { + if (exp < -10) return sign; + mant |= 0x800000; + uint32_t shift = static_cast(14 - exp); + uint32_t half = mant >> shift; + uint32_t rem = mant & ((1u << shift) - 1); + uint32_t mid = 1u << (shift - 1); + if (rem > mid || (rem == mid && (half & 1))) ++half; + return static_cast(sign | half); + } + uint32_t half = static_cast(exp << 10) | (mant >> 13); + uint32_t rem = mant & 0x1FFF; + if (rem > 0x1000 || (rem == 0x1000 && (half & 1))) ++half; + return static_cast(sign | half); +} + +// Load one activation element (dtype-decoded, exactly like cpu LoadActF32). + +using vt::cpu::BlockIQ1_S; +using vt::cpu::BlockIQ1_XXXS; +using vt::cpu::BlockIQ2_S; +using vt::cpu::BlockIQ2_XXS; +using vt::cpu::BlockIQ3_XXS; +using vt::cpu::BlockQ2_K; +using vt::cpu::BlockQ3_K; +using vt::cpu::BlockQ4_K; +using vt::cpu::BlockQ5_K; +using vt::cpu::BlockQ6_K; +using vt::cpu::BlockQ8_K; +using vt::cpu::BlockQ8_0; +using vt::cpu::kQK_K; +using vt::cpu::kQK8_0; + +void CheckHip(hipError_t err, const char* what) { + if (err != hipSuccess) { + throw std::runtime_error(std::string("vt rocm: matmul_bt_quant: ") + what + + ": " + hipGetErrorString(err)); + } +} + +// --- device numeric helpers — bit-exact ports of src/vt/dtype.cpp ------------- +__device__ inline float DF16ToF32(uint16_t h) { + uint32_t sign = static_cast(h & 0x8000) << 16; + uint32_t exp = (h >> 10) & 0x1F; + uint32_t mant = h & 0x3FF; + if (exp == 0x1F) return __int_as_float(sign | 0x7F800000 | (mant << 13)); + if (exp == 0) { + if (mant == 0) return __int_as_float(sign); + int shift = 0; + while ((mant & 0x400) == 0) { + mant <<= 1; + ++shift; + } + mant &= 0x3FF; + return __int_as_float(sign | ((113 - shift) << 23) | (mant << 13)); + } + return __int_as_float(sign | ((exp + 112) << 23) | (mant << 13)); +} + +__device__ inline float DBF16ToF32(uint16_t b) { + return __int_as_float(static_cast(b) << 16); +} + +__device__ inline uint16_t DF32ToBF16(float f) { + uint32_t u = __float_as_int(f); + if ((u & 0x7F800000) == 0x7F800000 && (u & 0x7FFFFF)) { + return static_cast((u >> 16) | 0x0040); + } + uint32_t rounding = 0x7FFF + ((u >> 16) & 1); + return static_cast((u + rounding) >> 16); +} + +// cpu_quant_act.cpp NearestInt (ggml-quants.c:563) — magic-constant round-to-even. +__device__ inline int DNearestInt(float fval) { + float val = fval + 12582912.0f; + int i = __float_as_int(val); + return (i & 0x007fffff) - 0x00400000; +} + +enum class ActDT : int { kF32 = 0, kF16 = 1, kBF16 = 2 }; + +__device__ inline float DLoadAct(const void* base, ActDT dt, int64_t idx) { + switch (dt) { + case ActDT::kF32: return static_cast(base)[idx]; + case ActDT::kF16: return DF16ToF32(static_cast(base)[idx]); + default: return DBF16ToF32(static_cast(base)[idx]); + } +} + +// --------------------------------------------------------------------------- +// GPU activation quantizer — ds4-parity grid: ONE BLOCK per (super-block, row), +// 256 threads (one thread per element). Bit-exact port of QuantizeRowQ8_K with +// the argmax tie broken by LOWEST original index (== the sequential first- +// occurrence scan). Byte-identical to both CUDA quantizers by construction. +// --------------------------------------------------------------------------- +__global__ void QuantizeQ8KKernel(BlockQ8_K* __restrict__ scratch, + const void* __restrict__ a, ActDT adt, + int64_t a_rs, int64_t m, int64_t nsb) { + const int64_t b = static_cast(blockIdx.x); // super-block within row + const int64_t i = static_cast(blockIdx.y); // activation row + if (b >= nsb || i >= m) return; + const int tid = static_cast(threadIdx.x); + const int64_t elem0 = i * a_rs + b * kQK_K; + const float v = DLoadAct(a, adt, elem0 + tid); + + __shared__ float sabs[kQK_K]; + __shared__ float sval[kQK_K]; + __shared__ int sidx[kQK_K]; + sabs[tid] = fabsf(v); + sval[tid] = v; + sidx[tid] = tid; + __syncthreads(); +#pragma unroll + for (int stride = kQK_K >> 1; stride > 0; stride >>= 1) { + if (tid < stride) { + const float oa = sabs[tid + stride]; + if (oa > sabs[tid] || (oa == sabs[tid] && sidx[tid + stride] < sidx[tid])) { + sabs[tid] = oa; + sval[tid] = sval[tid + stride]; + sidx[tid] = sidx[tid + stride]; + } + } + __syncthreads(); + } + const float mx = sval[0]; + const float amax = sabs[0]; + + BlockQ8_K& y = scratch[i * nsb + b]; + if (amax == 0.0f) { + if (tid == 0) y.d = 0.0f; + y.qs[tid] = 0; + if (tid < kQK_K / 16) y.bsums[tid] = 0; + return; + } + const float iscale = -127.0f / mx; + const int qv = DNearestInt(iscale * v); + y.qs[tid] = static_cast(qv < 127 ? qv : 127); + __syncthreads(); + if (tid < kQK_K / 16) { + int sum = 0; + for (int ii = 0; ii < 16; ++ii) sum += y.qs[tid * 16 + ii]; + y.bsums[tid] = static_cast(sum); + } + if (tid == 0) y.d = 1.0f / iscale; +} + +// --------------------------------------------------------------------------- +// Per-super-block integer dots. PORTABLE SCALAR forms of the CPU reference +// bodies (cpu_quant_dot.cpp), NOT the CUDA __dp4a forms: gfx1100 has no +// signed byte dot (see the HIP DELTAS header note). Exactness is preserved by +// keeping each dot's accumulation ORDER identical to the CPU body it mirrors: +// Q2_K sums sub-blocks in k/j/l order; Q3_K/Q6_K use the 8-wide aux32 split; +// Q4_K/Q5_K accumulate per-16 bsums then per-32 scale groups in order. +// --------------------------------------------------------------------------- +__device__ inline float DotQ2K(const BlockQ2_K* xb, const BlockQ8_K* yb) { + const uint8_t* q2 = xb->qs; + const int8_t* q8 = yb->qs; + const uint8_t* sc = xb->scales; + int summs = 0; + for (int j = 0; j < 16; ++j) summs += yb->bsums[j] * (sc[j] >> 4); + const float dall = yb->d * DF16ToF32(xb->d); + const float dmin = yb->d * DF16ToF32(xb->dmin); + int isum = 0; + int is = 0; + for (int k = 0; k < kQK_K / 128; ++k) { + int shift = 0; + for (int j = 0; j < 4; ++j) { + int d = sc[is++] & 0xF; + int isuml = 0; + for (int l = 0; l < 16; ++l) isuml += q8[l] * ((q2[l] >> shift) & 3); + isum += d * isuml; + d = sc[is++] & 0xF; + isuml = 0; + for (int l = 16; l < 32; ++l) isuml += q8[l] * ((q2[l] >> shift) & 3); + isum += d * isuml; + shift += 2; + q8 += 32; + } + q2 += 32; + } + return dall * isum - dmin * summs; +} + +__device__ inline float DotQ3K(const BlockQ3_K* xb, const BlockQ8_K* yb) { + const uint32_t kmask1 = 0x03030303; + const uint32_t kmask2 = 0x0f0f0f0f; + const uint8_t* hm = xb->hmask; + const int8_t* q8 = yb->qs; + int8_t aux8[kQK_K]; + int8_t* a = aux8; + const uint8_t* q3 = xb->qs; + uint8_t m = 1; + for (int jj = 0; jj < kQK_K; jj += 128) { + for (int l = 0; l < 32; ++l) a[l] = q3[l] & 3; + for (int l = 0; l < 32; ++l) a[l] = static_cast(a[l] - ((hm[l] & m) ? 0 : 4)); + a += 32; m = static_cast(m << 1); + for (int l = 0; l < 32; ++l) a[l] = (q3[l] >> 2) & 3; + for (int l = 0; l < 32; ++l) a[l] = static_cast(a[l] - ((hm[l] & m) ? 0 : 4)); + a += 32; m = static_cast(m << 1); + for (int l = 0; l < 32; ++l) a[l] = (q3[l] >> 4) & 3; + for (int l = 0; l < 32; ++l) a[l] = static_cast(a[l] - ((hm[l] & m) ? 0 : 4)); + a += 32; m = static_cast(m << 1); + for (int l = 0; l < 32; ++l) a[l] = (q3[l] >> 6) & 3; + for (int l = 0; l < 32; ++l) a[l] = static_cast(a[l] - ((hm[l] & m) ? 0 : 4)); + a += 32; m = static_cast(m << 1); + q3 += 32; + } + uint32_t auxs[4]; + memcpy(auxs, xb->scales, 12); + const int8_t* scales = reinterpret_cast(auxs); + uint32_t tmp = auxs[2]; + auxs[2] = ((auxs[0] >> 4) & kmask2) | (((tmp >> 4) & kmask1) << 4); + auxs[3] = ((auxs[1] >> 4) & kmask2) | (((tmp >> 6) & kmask1) << 4); + auxs[0] = (auxs[0] & kmask2) | (((tmp >> 0) & kmask1) << 4); + auxs[1] = (auxs[1] & kmask2) | (((tmp >> 2) & kmask1) << 4); + a = aux8; + const int8_t* q8p = q8; + int32_t aux32[8] = {0, 0, 0, 0, 0, 0, 0, 0}; + for (int j = 0; j < kQK_K / 16; ++j) { + for (int l = 0; l < 8; ++l) aux32[l] += (scales[j] - 32) * (q8p[l] * a[l]); + q8p += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += (scales[j] - 32) * (q8p[l] * a[l]); + q8p += 8; a += 8; + } + const float d = DF16ToF32(xb->d) * yb->d; + int isum = 0; + for (int l = 0; l < 8; ++l) isum += aux32[l]; + return d * isum; +} + +__device__ inline float DotQ4K(const BlockQ4_K* xb, const BlockQ8_K* yb) { + const uint32_t kmask1 = 0x3f3f3f3f; + const uint32_t kmask2 = 0x0f0f0f0f; + const uint32_t kmask3 = 0x03030303; + const uint8_t* q4 = xb->qs; + const int8_t* q8 = yb->qs; + uint32_t utmp[4]; + memcpy(utmp, xb->scales, 12); + utmp[3] = ((utmp[2] >> 4) & kmask2) | (((utmp[1] >> 6) & kmask3) << 4); + const uint32_t uaux = utmp[1] & kmask1; + utmp[1] = (utmp[2] & kmask2) | (((utmp[0] >> 6) & kmask3) << 4); + utmp[2] = uaux; + utmp[0] &= kmask1; + const uint8_t* scales = reinterpret_cast(&utmp[0]); + const uint8_t* mins = reinterpret_cast(&utmp[2]); + int sumi = 0; + for (int j = 0; j < kQK_K / 16; ++j) sumi += yb->bsums[j] * mins[j / 2]; + // Portable nibble walk in the CPU body's order: 64-element groups decode + // low nibble then high nibble into aux8, then the per-32 scale groups. + int8_t aux8[kQK_K]; + int8_t* a = aux8; + for (int j = 0; j < kQK_K / 64; ++j) { + for (int l = 0; l < 32; ++l) a[l] = static_cast(q4[l] & 0xF); + a += 32; + for (int l = 0; l < 32; ++l) a[l] = static_cast(q4[l] >> 4); + a += 32; + q4 += 32; + } + a = aux8; + int32_t aux32[8] = {0, 0, 0, 0, 0, 0, 0, 0}; + for (int j = 0; j < kQK_K / 32; ++j) { + const int32_t scale = scales[j]; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + } + const float d = DF16ToF32(xb->d) * yb->d; + const float dmin = DF16ToF32(xb->dmin) * yb->d; + int isum = 0; + for (int l = 0; l < 8; ++l) isum += aux32[l]; + return d * isum - dmin * sumi; +} + +__device__ inline float DotQ5K(const BlockQ5_K* xb, const BlockQ8_K* yb) { + const uint32_t kmask1 = 0x3f3f3f3f; + const uint32_t kmask2 = 0x0f0f0f0f; + const uint32_t kmask3 = 0x03030303; + const uint8_t* q4 = xb->qs; + const uint8_t* hm = xb->qh; + const int8_t* q8 = yb->qs; + uint32_t utmp[4]; + memcpy(utmp, xb->scales, 12); + utmp[3] = ((utmp[2] >> 4) & kmask2) | (((utmp[1] >> 6) & kmask3) << 4); + const uint32_t uaux = utmp[1] & kmask1; + utmp[1] = (utmp[2] & kmask2) | (((utmp[0] >> 6) & kmask3) << 4); + utmp[2] = uaux; + utmp[0] &= kmask1; + const uint8_t* scales = reinterpret_cast(&utmp[0]); + const uint8_t* mins = reinterpret_cast(&utmp[2]); + int sumi = 0; + for (int j = 0; j < kQK_K / 16; ++j) sumi += yb->bsums[j] * mins[j / 2]; + int8_t aux8[kQK_K]; + int8_t* a = aux8; + uint8_t m = 1; + for (int j = 0; j < kQK_K / 64; ++j) { + for (int l = 0; l < 32; ++l) a[l] = static_cast(q4[l] & 0xF); + for (int l = 0; l < 32; ++l) + a[l] = static_cast(a[l] + ((hm[l] & m) ? 16 : 0)); + a += 32; + m = static_cast(m << 1); + for (int l = 0; l < 32; ++l) a[l] = static_cast(q4[l] >> 4); + for (int l = 0; l < 32; ++l) + a[l] = static_cast(a[l] + ((hm[l] & m) ? 16 : 0)); + a += 32; + m = static_cast(m << 1); + q4 += 32; + } + a = aux8; + int32_t aux32[8] = {0, 0, 0, 0, 0, 0, 0, 0}; + for (int j = 0; j < kQK_K / 32; ++j) { + const int32_t scale = scales[j]; + for (int r = 0; r < 4; ++r) { + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8[l] * a[l]); + q8 += 8; a += 8; + } + } + const float d = DF16ToF32(xb->d) * yb->d; + const float dmin = DF16ToF32(xb->dmin) * yb->d; + int isum = 0; + for (int l = 0; l < 8; ++l) isum += aux32[l]; + return d * isum - dmin * sumi; +} + +__device__ inline float DotQ6K(const BlockQ6_K* xb, const BlockQ8_K* yb) { + const uint8_t* q4 = xb->ql; + const uint8_t* qh = xb->qh; + const int8_t* q8 = yb->qs; + int8_t aux8[kQK_K]; + int8_t* a = aux8; + for (int j = 0; j < kQK_K; j += 128) { + for (int l = 0; l < 32; ++l) { + a[l + 0] = static_cast( + static_cast((q4[l + 0] & 0xF) | (((qh[l] >> 0) & 3) << 4)) - 32); + a[l + 32] = static_cast( + static_cast((q4[l + 32] & 0xF) | (((qh[l] >> 2) & 3) << 4)) - 32); + a[l + 64] = static_cast( + static_cast((q4[l + 0] >> 4) | (((qh[l] >> 4) & 3) << 4)) - 32); + a[l + 96] = static_cast( + static_cast((q4[l + 32] >> 4) | (((qh[l] >> 6) & 3) << 4)) - 32); + } + a += 128; q4 += 64; qh += 32; + } + a = aux8; + const int8_t* q8p = q8; + int32_t aux32[8] = {0, 0, 0, 0, 0, 0, 0, 0}; + for (int j = 0; j < kQK_K / 16; ++j) { + const int scale = xb->scales[j]; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8p[l] * a[l]); + q8p += 8; a += 8; + for (int l = 0; l < 8; ++l) aux32[l] += scale * (q8p[l] * a[l]); + q8p += 8; a += 8; + } + const float d = DF16ToF32(xb->d) * yb->d; + int isum = 0; + for (int l = 0; l < 8; ++l) isum += aux32[l]; + return d * isum; +} + +__device__ inline float DotIQ2XXS(const BlockIQ2_XXS* xb, const BlockQ8_K* yb) { + const float d = DF16ToF32(xb->d) * yb->d; + const uint16_t* qs = xb->qs; + const int8_t* q8 = yb->qs; + int32_t bsum = 0; + for (int ib32 = 0; ib32 < kQK_K / 32; ++ib32) { + uint32_t aux32[2]; + memcpy(aux32, qs + 4 * ib32, 2 * sizeof(uint32_t)); + const uint32_t ls = 2 * (aux32[1] >> 28) + 1; + int32_t sumi = 0; + for (int l = 0; l < 4; ++l) { + const uint8_t* grid = reinterpret_cast( + &vt::cuda::d_iq2xxs_grid[(aux32[0] >> (8 * l)) & 0xff]); + const uint8_t signs = + vt::cuda::d_ksigns_iq2xs[(aux32[1] >> (7 * l)) & 127]; + for (int j = 0; j < 8; ++j) + sumi += grid[j] * q8[j] * ((signs & vt::cuda::d_kmask_iq2xs[j]) ? -1 : 1); + q8 += 8; + } + bsum += sumi * static_cast(ls); + } + return d * bsum; // final *0.125 applied after the warp reduction +} + +__device__ inline float DotIQ3XXS(const BlockIQ3_XXS* xb, const BlockQ8_K* yb) { + const float d = DF16ToF32(xb->d) * yb->d; + const uint8_t* q3 = xb->qs; + const uint8_t* gas = xb->qs + kQK_K / 4; + const int8_t* q8 = yb->qs; + int32_t bsum = 0; + for (int ib32 = 0; ib32 < kQK_K / 32; ++ib32) { + uint32_t a32; + memcpy(&a32, gas, sizeof(uint32_t)); + gas += sizeof(uint32_t); + const uint32_t ls = 2 * (a32 >> 28) + 1; + int32_t sumi = 0; + for (int l = 0; l < 4; ++l) { + const uint32_t g1 = vt::cuda::d_iq3xxs_grid[q3[2 * l + 0]]; + const uint32_t g2 = vt::cuda::d_iq3xxs_grid[q3[2 * l + 1]]; + const uint8_t signs = vt::cuda::d_ksigns_iq2xs[(a32 >> (7 * l)) & 127]; + for (int j = 0; j < 4; ++j) { + const int b1 = static_cast((g1 >> (8 * j)) & 0xff); + const int b2 = static_cast((g2 >> (8 * j)) & 0xff); + sumi += b1 * q8[j + 0] * ((signs & vt::cuda::d_kmask_iq2xs[j + 0]) ? -1 : 1); + sumi += b2 * q8[j + 4] * ((signs & vt::cuda::d_kmask_iq2xs[j + 4]) ? -1 : 1); + } + q8 += 8; + } + q3 += 8; + bsum += sumi * static_cast(ls); + } + return d * bsum; // final *0.25 applied after the warp reduction +} + +__device__ inline float DotIQ2S(const BlockIQ2_S* xb, const BlockQ8_K* yb) { + const float d = DF16ToF32(xb->d) * yb->d; + const int8_t* q8 = yb->qs; + const uint8_t* qs = xb->qs; + const uint8_t* qh = xb->qh; + const uint8_t* signs = qs + kQK_K / 8; + int32_t bsum = 0; + for (int ib32 = 0; ib32 < kQK_K / 32; ++ib32) { + const int ls1 = 1 + 2 * (xb->scales[ib32] & 0xf); + const int ls2 = 1 + 2 * (xb->scales[ib32] >> 4); + int sumi1 = 0; + int sumi2 = 0; + for (int l = 0; l < 2; ++l) { + const uint8_t* grid = reinterpret_cast( + &vt::cuda::d_iq2s_grid[qs[l] | ((qh[ib32] << (8 - 2 * l)) & 0x300)]); + for (int j = 0; j < 8; ++j) + sumi1 += q8[j] * grid[j] * ((signs[l] & vt::cuda::d_kmask_iq2xs[j]) ? -1 : 1); + q8 += 8; + } + for (int l = 2; l < 4; ++l) { + const uint8_t* grid = reinterpret_cast( + &vt::cuda::d_iq2s_grid[qs[l] | ((qh[ib32] << (8 - 2 * l)) & 0x300)]); + for (int j = 0; j < 8; ++j) + sumi2 += q8[j] * grid[j] * ((signs[l] & vt::cuda::d_kmask_iq2xs[j]) ? -1 : 1); + q8 += 8; + } + bsum += ls1 * sumi1 + ls2 * sumi2; + qs += 4; + signs += 4; + } + return d * bsum; // final *0.125 applied after the warp reduction +} + +__device__ inline float DotIQ1S(const BlockIQ1_S* xb, const BlockQ8_K* yb) { + const int8_t* q8 = yb->qs; + const uint8_t* qs = xb->qs; + const uint16_t* qh = xb->qh; + int32_t sumi = 0; + int32_t sumi1 = 0; + for (int ib = 0; ib < kQK_K / 32; ++ib) { + const int ls = 2 * ((qh[ib] >> 12) & 7) + 1; + const int delta = (qh[ib] & 0x8000) ? -1 : 1; + int lsum = 0; + for (int l = 0; l < 4; ++l) { + const int8_t* grid = reinterpret_cast( + &vt::cuda::d_iq1s_grid[qs[l] | (((qh[ib] >> (3 * l)) & 7) << 8)]); + for (int j = 0; j < 8; ++j) lsum += q8[j] * grid[j]; + q8 += 8; + } + sumi += ls * lsum; + sumi1 += ls * delta * (yb->bsums[2 * ib + 0] + yb->bsums[2 * ib + 1]); + qs += 4; + } + return DF16ToF32(xb->d) * yb->d * + (static_cast(sumi) + 0.125f * static_cast(sumi1)); +} + +__device__ inline float DotIQ1XXXS(const BlockIQ1_XXXS* xb, const BlockQ8_K* yb) { + const int8_t* q8 = yb->qs; + const uint8_t* qs = xb->qs; + const uint8_t* sc = xb->sc; + int32_t sumi = 0; + int32_t sumi1 = 0; + for (int ib = 0; ib < kQK_K / 32; ++ib) { + const int nib = (sc[ib / 2] >> (4 * (ib & 1))) & 0xf; + const int ls = 2 * (nib & 7) + 1; + const int delta = (nib & 8) ? -1 : 1; + int lsum = 0; + for (int l = 0; l < 4; ++l) { + const int8_t* grid = + reinterpret_cast(&vt::cuda::d_iq1xxxs_grid[qs[l]]); + for (int j = 0; j < 8; ++j) lsum += q8[j] * grid[j]; + q8 += 8; + } + sumi += ls * lsum; + sumi1 += ls * delta * (yb->bsums[2 * ib + 0] + yb->bsums[2 * ib + 1]); + qs += 4; + } + return DF16ToF32(xb->d) * yb->d * + (static_cast(sumi) + 0.125f * static_cast(sumi1)); +} + +// --------------------------------------------------------------------------- +// WType tags + DotSuperblock dispatch — mirrors the CUDA file's table. +// --------------------------------------------------------------------------- +enum class WType : int { + kIQ2_XXS = 0, + kIQ3_XXS = 1, + kQ2_K = 2, + kQ3_K = 3, + kQ4_K = 4, + kQ5_K = 5, + kQ6_K = 6, + kIQ2_S = 7, + kIQ1_S = 8, + kIQ1_XXXS = 9, +}; + +template +__device__ inline float DotSuperblock(const void* w_sb, const BlockQ8_K* a_sb); + +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ2XXS(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ3XXS(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ2K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ3K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ4K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ5K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotQ6K(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ2S(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ1S(static_cast(w), a); +} +template <> +__device__ inline float DotSuperblock(const void* w, const BlockQ8_K* a) { + return DotIQ1XXXS(static_cast(w), a); +} + +template +__device__ constexpr float FinalFactor() { + return (W == WType::kIQ2_XXS || W == WType::kIQ2_S) + ? 0.125f + : (W == WType::kIQ3_XXS ? 0.25f : 1.0f); +} + +// --------------------------------------------------------------------------- +// The MMVQ-style GEMM: one WARP per output element (i,j). Lanes split the K +// super-blocks; the warp reduction sums the partials. HIP delta: the shuffle +// mask is 64-bit on this target. Determinism note unchanged from CUDA: the +// integer core is exact; only the scale sum reassociates (within NMSE). +// --------------------------------------------------------------------------- +template +__global__ void QuantDotGemmKernel(OutT* __restrict__ out, + const uint8_t* __restrict__ weight, + const BlockQ8_K* __restrict__ act, int64_t m, + int64_t n, int64_t nsb, size_t w_row_bytes, + size_t w_block_bytes) { + const int64_t warp = static_cast(blockIdx.x) * (blockDim.x >> 5) + + (threadIdx.x >> 5); + if (warp >= m * n) return; + const int64_t i = warp / n; + const int64_t j = warp % n; + const int lane = threadIdx.x & 31; + + const uint8_t* w_row = weight + static_cast(j) * w_row_bytes; + const BlockQ8_K* a_row = act + i * nsb; + + float partial = 0.0f; + for (int64_t sb = lane; sb < nsb; sb += 32) { + const void* w_sb = w_row + static_cast(sb) * w_block_bytes; + partial += DotSuperblock(w_sb, a_row + sb); + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) + partial += __shfl_down_sync(0xffffffffffffffffull, partial, off); + + if (lane == 0) { + const float v = FinalFactor() * partial; + if constexpr (sizeof(OutT) == 4) { + out[i * n + j] = v; + } else { + out[i * n + j] = DF32ToBF16(v); + } + } +} + +// GROUPED variant: warp per (p, n); weight row selected by expert_ids[p]. +template +__global__ void QuantDotGemmGroupedKernel(OutT* __restrict__ out, + const uint8_t* __restrict__ weight, + const BlockQ8_K* __restrict__ act, + const int32_t* __restrict__ expert_ids, + int64_t P, int64_t n, int64_t nsb, + size_t w_row_bytes, + size_t w_block_bytes, bool bcast) { + const int64_t warp = static_cast(blockIdx.x) * (blockDim.x >> 5) + + (threadIdx.x >> 5); + if (warp >= P * n) return; + const int64_t p = warp / n; + const int64_t j = warp % n; + const int lane = threadIdx.x & 31; + + const int64_t e = expert_ids[p]; + const uint8_t* w_row = weight + static_cast(e * n + j) * w_row_bytes; + // Broadcast activation: the routed gate/up share ONE quantized hidden. + const BlockQ8_K* a_row = act + (bcast ? 0 : p) * nsb; + + float partial = 0.0f; + for (int64_t sb = lane; sb < nsb; sb += 32) { + const void* w_sb = w_row + static_cast(sb) * w_block_bytes; + partial += DotSuperblock(w_sb, a_row + sb); + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) + partial += __shfl_down_sync(0xffffffffffffffffull, partial, off); + + if (lane == 0) { + const float v = FinalFactor() * partial; + if constexpr (sizeof(OutT) == 4) { + out[p * n + j] = v; + } else { + out[p * n + j] = DF32ToBF16(v); + } + } +} + +// --------------------------------------------------------------------------- +// Host wiring: scratch, launches, providers, registrar. +// --------------------------------------------------------------------------- +struct StreamScratch { + void* buf = nullptr; + size_t bytes = 0; +}; + +std::mutex& ScratchMutex() { + static std::mutex mu; + return mu; +} + +StreamScratch& ScratchFor(hipStream_t s) { + static std::unordered_map map; + return map[s]; +} + +void CheckHipLaunch(const char* what) { CheckHip(hipGetLastError(), what); } + +void* EnsureScratch(size_t need, hipStream_t s) { + std::lock_guard lock(ScratchMutex()); + StreamScratch& sc = ScratchFor(s); + if (need > sc.bytes) { + // Retire (never free): a captured hipGraph may have baked this pointer. + vt::cuda::RetireGraphScratch(sc.buf); + CheckHip(hipMallocAsync(&sc.buf, need, s), "hipMallocAsync q8_K act scratch"); + sc.bytes = need; + } + return sc.buf; +} + +inline ActDT ActDtOf(DType dt) { + return dt == DType::kF32 ? ActDT::kF32 : dt == DType::kF16 ? ActDT::kF16 : ActDT::kBF16; +} + +void LaunchQuantizeQ8K(BlockQ8_K* qact, const void* data, ActDT adt, int64_t a_rs, + int64_t rows, int64_t nsb, hipStream_t s) { + dim3 qgrid(static_cast(nsb), static_cast(rows), 1); + QuantizeQ8KKernel<<>>(qact, data, adt, a_rs, rows, nsb); + CheckHipLaunch("quantize_q8_K launch"); +} + +bool IsRocmKeepQuantSupported(DType dt, WType* out) { + switch (dt) { + case DType::kIQ2_XXS: *out = WType::kIQ2_XXS; return true; + case DType::kIQ3_XXS: *out = WType::kIQ3_XXS; return true; + case DType::kQ2_K: *out = WType::kQ2_K; return true; + case DType::kQ3_K: *out = WType::kQ3_K; return true; + case DType::kQ4_K: *out = WType::kQ4_K; return true; + case DType::kQ5_K: *out = WType::kQ5_K; return true; + case DType::kQ6_K: *out = WType::kQ6_K; return true; + case DType::kIQ2_S: *out = WType::kIQ2_S; return true; + case DType::kIQ1_S: *out = WType::kIQ1_S; return true; + case DType::kIQ1_XXXS: *out = WType::kIQ1_XXXS; return true; + // Q4_0 / Q8_0 / MXFP4 dot a Q8_0 activation and have no native arm here. + default: return false; + } +} + +template +void LaunchGemm(Tensor& out, const uint8_t* weight, const BlockQ8_K* act, + int64_t m, int64_t n, int64_t nsb, size_t w_row_bytes, + size_t w_block_bytes, hipStream_t s) { + // Wave32 geometry: one warp per output; 8 warps (256 threads) per block. + constexpr int kWarpsPerBlock = 8; + dim3 block(32 * kWarpsPerBlock, 1, 1); + const int64_t warps = m * n; + const unsigned grid = + static_cast((warps + kWarpsPerBlock - 1) / kWarpsPerBlock); + if (out.dtype == DType::kF32) { + QuantDotGemmKernel<<>>( + static_cast(out.data), weight, act, m, n, nsb, w_row_bytes, + w_block_bytes); + } else { + QuantDotGemmKernel<<>>( + static_cast(out.data), weight, act, m, n, nsb, w_row_bytes, + w_block_bytes); + } + CheckHipLaunch("matmul_bt_quant launch"); +} + +template +void LaunchGroupedGemm(Tensor& out, const uint8_t* weight, const BlockQ8_K* act, + const int32_t* expert_ids, int64_t P, int64_t n, + int64_t nsb, size_t w_row_bytes, size_t w_block_bytes, + bool bcast, hipStream_t s) { + constexpr int kWarpsPerBlock = 8; + dim3 block(32 * kWarpsPerBlock, 1, 1); + const int64_t warps = P * n; + const unsigned grid = + static_cast((warps + kWarpsPerBlock - 1) / kWarpsPerBlock); + if (out.dtype == DType::kF32) { + QuantDotGemmGroupedKernel<<>>( + static_cast(out.data), weight, act, expert_ids, P, n, nsb, + w_row_bytes, w_block_bytes, bcast); + } else { + QuantDotGemmGroupedKernel<<>>( + static_cast(out.data), weight, act, expert_ids, P, n, nsb, + w_row_bytes, w_block_bytes, bcast); + } + CheckHipLaunch("matmul_bt_quant_grouped launch"); +} + + + +// Q8_0 (legacy 32-block, Q8_0-activation) arm. Self-contained: quantize the +// activation to Q8_0 on the device, then the Q8_0xQ8_0 integer dot. The dot is +// the PORTABLE SCALAR form of cpu_quant_dot.cpp VecDotQ8_0Q8_0 (quants.c:400): +// sumi0 over the first 16 elements, sumi1 over the second, ONE int32 add, then +// the f16-scale fold — accumulation order preserved for exactness. The CUDA +// file's __dp4a form is byte-identical to this on a signed-dot target; gfx1100 +// has no signed byte dot, so W1 ships the scalar body. +// --------------------------------------------------------------------------- +__global__ void QuantizeQ8_0Kernel(BlockQ8_0* __restrict__ scratch, + const void* __restrict__ a, ActDT adt, + int64_t a_rs, int64_t m, int64_t nb) { + const int64_t t = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (t >= m * nb) return; + const int64_t i = t / nb; + const int64_t b = t % nb; + const int64_t elem0 = i * a_rs + b * kQK8_0; + float amax = 0.0f; + for (int j = 0; j < kQK8_0; ++j) { + const float av = fabsf(DLoadAct(a, adt, elem0 + j)); + amax = amax > av ? amax : av; + } + BlockQ8_0& y = scratch[t]; + const float d = amax / 127.0f; + const float id = d != 0.0f ? 1.0f / d : 0.0f; + y.d = DF32ToF16(d); + for (int j = 0; j < kQK8_0; ++j) { + y.qs[j] = static_cast(roundf(DLoadAct(a, adt, elem0 + j) * id)); + } +} + +template +__global__ void QuantDotGemmQ8_0Kernel(OutT* __restrict__ out, + const uint8_t* __restrict__ weight, + const BlockQ8_0* __restrict__ act, + int64_t m, int64_t n, int64_t nb, + size_t w_row_bytes) { + const int64_t warp = static_cast(blockIdx.x) * (blockDim.x >> 5) + + (threadIdx.x >> 5); + if (warp >= m * n) return; + const int64_t i = warp / n; + const int64_t j = warp % n; + const int lane = threadIdx.x & 31; + const uint8_t* w_row = weight + static_cast(j) * w_row_bytes; + const BlockQ8_0* a_row = act + i * nb; + float partial = 0.0f; + for (int64_t b = lane; b < nb; b += 32) { + const BlockQ8_0* wb = + reinterpret_cast(w_row + static_cast(b) * + sizeof(BlockQ8_0)); + const BlockQ8_0* ab = a_row + b; + // CPU order: first half then second half, one add, then scale. + int sumi0 = 0; + for (int j2 = 0; j2 < kQK8_0 / 2; ++j2) + sumi0 += ab->qs[j2] * wb->qs[j2]; + int sumi1 = 0; + for (int j2 = kQK8_0 / 2; j2 < kQK8_0; ++j2) + sumi1 += ab->qs[j2] * wb->qs[j2]; + partial += (sumi0 + sumi1) * (DF16ToF32(wb->d) * DF16ToF32(ab->d)); + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) + partial += __shfl_down_sync(0xffffffffffffffffull, partial, off); + if (lane == 0) { + if constexpr (sizeof(OutT) == 4) + out[i * n + j] = partial; + else + out[i * n + j] = DF32ToBF16(partial); + } +} + +void MatmulQ8_0Rocm(Tensor& out, const Tensor& a, const Tensor& b, + hipStream_t s) { + const int64_t m = a.shape[0], k = a.shape[1], n = b.shape[0]; + if (m == 0 || n == 0) return; + if (k % kQK8_0 != 0) + throw std::runtime_error( + "vt rocm: matmul_bt_quant Q8_0: K must be a multiple of 32"); + const int64_t nb = k / kQK8_0; + const size_t w_row_bytes = static_cast(nb) * sizeof(BlockQ8_0); + const size_t act_bytes = + static_cast(m) * static_cast(nb) * sizeof(BlockQ8_0); + BlockQ8_0* act = static_cast(EnsureScratch(act_bytes, s)); + { + constexpr int kQBlock = 128; + const int64_t grid = (m * nb + kQBlock - 1) / kQBlock; + QuantizeQ8_0Kernel<<(grid), kQBlock, 0, s>>>( + act, a.data, ActDtOf(a.dtype), a.stride[0], m, nb); + CheckHipLaunch("quantize_q8_0 launch"); + } + constexpr int kWarpsPerBlock = 8; + dim3 block(32 * kWarpsPerBlock, 1, 1); + const unsigned grid = + static_cast((m * n + kWarpsPerBlock - 1) / kWarpsPerBlock); + if (out.dtype == DType::kF32) + QuantDotGemmQ8_0Kernel<<>>( + static_cast(out.data), static_cast(b.data), + act, m, n, nb, w_row_bytes); + else + QuantDotGemmQ8_0Kernel<<>>( + static_cast(out.data), static_cast(b.data), + act, m, n, nb, w_row_bytes); + CheckHipLaunch("matmul_bt_quant Q8_0 launch"); +} + +// The kROCM provider for OpId::kMatmulBTQuant. Validation already done by +// vt::MatmulBTQuant (ops.cpp). Contract: b is [N,K] block-quant, a [M,K] +// f32/bf16 row-packed, out [M,N]. On this DISCRETE backend an unsupported +// dtype cannot fall back to the CPU kernel (it would follow device pointers), +// so it throws naming the dtype — VT_GGUF_KEEP_QUANT=0 restores the load-time +// bf16 expansion for such files. +void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, + const Tensor& b) { + hipStream_t s = static_cast(q.handle); + const int64_t m = a.shape[0]; + const int64_t k = a.shape[1]; + const int64_t n = b.shape[0]; + if (b.dtype == DType::kQ8_0) { + MatmulQ8_0Rocm(out, a, b, s); + return; + } + + // Delegate Q4_K/Q5_K/Q6_K to the optimized kernels in rocm_grouped_gemm.hip + // (renamed *Gdn). These have cooperative-warp dispatch and format-specific + // templates tuned for gfx1100; the new CUDA-ported kernels in this file are + // slower for these types. The new kernels handle IQ types and Q2_K/Q3_K. + if (b.dtype == DType::kQ4_K || b.dtype == DType::kQ5_K || + b.dtype == DType::kQ6_K) { + MatmulBTQuantKernelRocmGdn(q, out, a, b); + 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; + + // Delegate Q8_0/Q4_K/Q5_K/Q6_K to the optimized grouped kernels. Q8_0 has no + // Q8_K-superblock arm in this file (it dots a Q8_0 activation), so leaving it + // out of this list turns a format the *Gdn kernel serves today into a throw. + if (weight.dtype == DType::kQ8_0 || weight.dtype == DType::kQ4_K || + weight.dtype == DType::kQ5_K || weight.dtype == DType::kQ6_K) { + MatmulBTQuantGroupedKernelRocmGdn(q, out, act, weight, expert_ids); + 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). The registration is what +// makes GgufQuantComputeAvailable() answer yes for kROCM, so the GGUF loader +// keeps the file quantized instead of expanding it to bf16 at load. +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/src/vt/rocm/rocm_skinny_gemm.hip b/src/vt/rocm/rocm_skinny_gemm.hip index b4d7dd755..c77ee1250 100644 --- a/src/vt/rocm/rocm_skinny_gemm.hip +++ b/src/vt/rocm/rocm_skinny_gemm.hip @@ -52,11 +52,28 @@ inline int mindiv(int N, int div1, int div2) { } // A (activation) fits LDS. N = decode batch (our M). bf16, f32 accum. -template +// Output-store polymorphism: the adopted bf16 decode arm rounds the f32 +// accumulator to bf16; lever B2 (evidence 15.1) serves the Qwen3.5 GDN BA +// projections, which must EMIT f32 — same geometry/reduction tree, only the +// store type differs. No rounding on the f32 arm. +template +__device__ __forceinline__ OutT WvStoreCast(float v); +template <> +__device__ __forceinline__ __hip_bfloat16 WvStoreCast<__hip_bfloat16>(float v) { + return __float2bfloat16(v); +} +template <> +__device__ __forceinline__ float WvStoreCast(float v) { + return v; +} + +// A (activation) fits LDS. N = decode batch (our M). bf16 in, OutT out, +// f32 accum. +template __global__ void __launch_bounds__(kWvPrGrp * kThrds) wvSplitKSml(const int K, const int Kbp, const int Kap, const int M, const __hip_bfloat16* __restrict__ B, const __hip_bfloat16* __restrict__ A, - __hip_bfloat16* C, const int _WvPrGrp, const int CuCount) { + OutT* C, const int _WvPrGrp, const int CuCount) { constexpr int max_lds_len = kLdsSize / 2; // bf16 elements __shared__ __hip_bfloat16 s[max_lds_len]; @@ -69,19 +86,19 @@ __global__ void __launch_bounds__(kWvPrGrp * kThrds) if (threadIdx.y >= static_cast(_WvPrGrp)) return; - uint32_t m = (blockIdx.x * _WvPrGrp + (threadIdx.y % _WvPrGrp)) * kYtile; + uint32_t m = (blockIdx.x * _WvPrGrp + (threadIdx.y % _WvPrGrp)) * YT; while (m < static_cast(M)) { - float sum[N][kYtile] = {}; + float sum[N][YT] = {}; for (uint32_t k1 = 0; k1 < static_cast(K); k1 += kThrds * kAChunk * kUnrl) { bigType bigA[N][kUnrl] = {}; - bigType bigB[kYtile][kUnrl]; + bigType bigB[YT][kUnrl]; #pragma unroll for (uint32_t k2 = 0; k2 < kUnrl; k2++) { uint32_t k = k1 + k2 * kThrds * kAChunk; uint32_t k_ = k + threadIdx.x * kAChunk; const __hip_bfloat16* B_ = &B[min__(k_, K - kAChunk)]; - for (int y = 0; y < kYtile; y++) + for (int y = 0; y < YT; y++) bigB[y][k2].h8 = __builtin_nontemporal_load( (const scalar8*)(&B_[min__(y + m, M - 1) * Kbp])); } @@ -95,7 +112,7 @@ __global__ void __launch_bounds__(kWvPrGrp * kThrds) // Interleaved MAC; bf16 pairs unpacked to f32 (donor DOT2C bf16 branch). for (uint32_t k2 = 0; k2 < kUnrl; k2++) { for (int n = 0; n < N; n++) { - for (int y = 0; y < kYtile; y++) { + for (int y = 0; y < YT; y++) { #pragma unroll for (uint32_t b = 0; b < kAChunk / 2; b++) { float2 a2 = __bfloat1622float2(*((__hip_bfloat162*)(&(bigA[n][k2].h[b * 2])))); @@ -109,7 +126,7 @@ __global__ void __launch_bounds__(kWvPrGrp * kThrds) __builtin_amdgcn_sched_barrier(0); // Wave32 reduction: DPP row_shr 8/4/2/1 then shfl_xor(16). for (int n = 0; n < N; n++) { - for (int y = 0; y < kYtile; y++) { + for (int y = 0; y < YT; y++) { sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x118, 0xf, 0xf, 1); sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x114, 0xf, 0xf, 1); sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x112, 0xf, 0xf, 1); @@ -119,9 +136,10 @@ __global__ void __launch_bounds__(kWvPrGrp * kThrds) } if (threadIdx.x == (kThrds - 1)) { for (int n = 0; n < N; n++) - for (int y = 0; y < kYtile; y++) C[m + y + n * M] = __float2bfloat16(sum[n][y]); + for (int y = 0; y < YT; y++) + C[m + y + n * M] = WvStoreCast(sum[n][y]); } - m += CuCount * _WvPrGrp * kYtile; + m += CuCount * _WvPrGrp * YT; } } @@ -139,31 +157,69 @@ int DeviceCuCount(int device) { return cache; } + +// Shared launch math; the store dtype is the only difference between the +// adopted bf16 arm (WvSplitKBT) and lever B2's f32-out arm (WvSplitKBTToF32). +// T16 (GFX1100-TG200): VT_WVSPLIT_YTILE (1|2|4, default 2 = donor proven +// decode config) selects precompiled kYtile variants; VT_WVSPLIT_PRGRP +// (0 = mindiv default, else explicit) overrides the work-groups-per-grouping +// so the three hot GDN shapes can be swept on this part. Read once. +namespace { +struct WvCfg { + int yt = 4; // T16 sweep: YT=4 wins 5/5 paired (+1.8% median), bit-identical + int prgrp = 0; +}; +inline WvCfg WvCfgFromEnv() { + WvCfg c; + if (const char* e = std::getenv("VT_WVSPLIT_YTILE")) { + const int v = atoi(e); + if (v == 1 || v == 2 || v == 4) c.yt = v; + } + if (const char* e = std::getenv("VT_WVSPLIT_PRGRP")) { + const int v = atoi(e); + if (v > 0 && v <= kWvPrGrp) c.prgrp = v; + } + return c; +} } // namespace -// out[M,N] = a[M,K] @ b[N,K]^T, bf16 in/out. Only called for the decode-skinny -// gate below (M in 1..4, K%8==0); the caller keeps every other shape on the -// BLAS path. -void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, int N, - int K, int device) { +template +void WvSplitKBTDispatch(hipStream_t s, void* out, const void* a, const void* b, + int M, int N, int K, int device, int yt, int prgrp_ovr) { const int cu = DeviceCuCount(device); dim3 grid(cu), block(kThrds, kWvPrGrp); - const int wvPrGrp = mindiv(N, cu * kYtile, kWvPrGrp); - auto* C = static_cast<__hip_bfloat16*>(out); + const int wvPrGrp = prgrp_ovr > 0 ? prgrp_ovr : mindiv(N, cu * yt, kWvPrGrp); + auto* C = static_cast(out); auto* A = static_cast(a); auto* B = static_cast(b); switch (M) { case 1: - wvSplitKSml<1><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); + switch (yt) { + case 1: wvSplitKSml<1, 1, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<1, 2, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<1, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } break; case 2: - wvSplitKSml<2><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); + switch (yt) { + case 1: wvSplitKSml<2, 1, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<2, 2, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<2, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } break; case 3: - wvSplitKSml<3><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); + switch (yt) { + case 1: wvSplitKSml<3, 1, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<3, 2, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<3, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } break; case 4: - wvSplitKSml<4><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); + switch (yt) { + case 1: wvSplitKSml<4, 1, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<4, 2, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<4, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } break; default: throw std::runtime_error("vt rocm: wvSplitK unsupported M=" + std::to_string(M)); @@ -172,4 +228,58 @@ void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, i throw std::runtime_error("vt rocm: wvSplitK launch failed"); } +template +void WvSplitKBTLaunch(hipStream_t s, void* out, const void* a, const void* b, + int M, int N, int K, int device) { + static const WvCfg cfg = WvCfgFromEnv(); + // f32-out arm (lever B2) instantiates its own store type at YT=2 only. + if constexpr (std::is_same_v) { + if (cfg.yt != 2) { + // keep donor geometry for the f32 arm regardless of sweep knob + const int cu = DeviceCuCount(device); + dim3 grid(cu), block(kThrds, kWvPrGrp); + const int wvPrGrp = mindiv(N, cu * kYtile, kWvPrGrp); + auto* C = static_cast(out); + auto* A = static_cast(a); + auto* B = static_cast(b); + switch (M) { + case 1: wvSplitKSml<1, kYtile, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<2, kYtile, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 3: wvSplitKSml<3, kYtile, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<4, kYtile, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } + return; + } + } + WvSplitKBTDispatch(s, out, a, b, M, N, K, device, cfg.yt, cfg.prgrp); +} + +} // namespace + +// Effective YTILE after the T16 sweep knobs, cached identically to the launch +// path's static cfg. The dispatch gate (rocm_matmul_hipblaslt.hip) tests +// N % YTILE == 0 against THIS value: the kernel's unguarded y-tile stores are +// memory-safe only for N a multiple of the tile actually launched, and the +// T16 default flip to 4 made the old N%2 gate too weak. +int WvSplitKYtile() { + static const WvCfg cfg = WvCfgFromEnv(); + return cfg.yt; +} + +// out[M,N] = a[M,K] @ b[N,K]^T, bf16 in/out. Only called for the decode-skinny +// gate below (M in 1..4, K%8==0); the caller keeps every other shape on the +// BLAS path. +void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, int N, + int K, int device) { + WvSplitKBTLaunch<__hip_bfloat16>(s, out, a, b, M, N, K, device); +} + +// Lever B2 (evidence 15.1): same geometry, f32 OUTPUT — the Qwen3.5 GDN BA +// projections consume f32 (g/beta chain) and today starve on rocBLAS's +// large-M tile. Dispatched by MatmulBTKernelRocm under VT_SKINNY_BF16=1. +void WvSplitKBTToF32(hipStream_t s, void* out, const void* a, const void* b, + int M, int N, int K, int device) { + WvSplitKBTLaunch(s, out, a, b, M, N, K, device); +} + } // namespace vt::rocm diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index f14de2923..b6198f2d3 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -2659,6 +2659,12 @@ vllm_cpp_add_test(test_rocm_fp8_kv_cache vt/test_rocm_fp8_kv_cache.cpp) # the multiply, matching the CPU oracle RoundThrough and upstream vLLM silu_kernel. # Self-skipping without a ROCm device, mirroring test_rocm_backend.cpp guard. vllm_cpp_add_test(test_ops_rocm_silu_rounding vt/test_ops_rocm_silu_rounding.cpp) +# KERNEL-QUANT-CIQ-GEMM-ROCM: the ROCm keep-quant GEMM (kROCM provider for +# kMatmulBTQuant/kMatmulBTQuantGrouped). Gates the device dequant-in-kernel dot +# against the CPU keep-quant oracle + an f64 dequant reference on the ten +# Q8_K-family encodings. Skips coherently with no AMD GPU. +vllm_cpp_add_test(test_rocm_quant_dot vt/test_rocm_quant_dot.cpp) +target_include_directories(test_rocm_quant_dot PRIVATE ${CMAKE_SOURCE_DIR}/src) # #785 P1 GPU product-seam witness. Executable only — NOT add_test. # Ordinary CTest must not see this target. Runner fail-closes on 77/nonzero. add_executable(test_ops_paged_attn_sharedk_wmma_p1_gpu diff --git a/tests/vt/test_rocm_quant_dot.cpp b/tests/vt/test_rocm_quant_dot.cpp new file mode 100644 index 000000000..593475c64 --- /dev/null +++ b/tests/vt/test_rocm_quant_dot.cpp @@ -0,0 +1,344 @@ +// 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"}, + // Q8_0 is NOT a Q8_K-superblock encoding: it dots a Q8_0 activation and is + // served by the *Gdn kernels this file delegates to. It sits in the same + // table because both arms of the provider must keep serving it -- dropping + // it from the grouped delegation list turns a working MoE format into a + // throw, and only a case here catches that. + {DType::kQ8_0, 32, 34, 0, -1, "q8_0"}, +}; + +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") { + 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); +}