feat(KERNEL-QUANT-CIQ-GEMM-ROCM): keep-quant GEMM providers on kROCM

[ghazni101]

What changed

W1 of KERNEL-QUANT-CIQ-GEMM-ROCM (#1587): the kROCM providers for
OpId::kMatmulBTQuant and OpId::kMatmulBTQuantGrouped, in the new
src/vt/rocm/rocm_quant_dot.hip, plus the committed spec
(.agents/specs/kernel-quant-ciq-gemm-rocm.md) and a red-first gate
(tests/vt/test_rocm_quant_dot.cpp).

Registering the providers flips GgufQuantComputeAvailable() true on
ROCm, so every GGUF k-quant model now computes on the AMD GPU through
the existing seams (loader policy -> vt::MatmulBT[Quant] ->
GetOp(kROCM)) with zero model-path edits. Towers stage once through
the existing ResidentWeight path.

Coverage mirrors the CUDA sibling exactly: ten Q8_K-family encodings
(Q2_K, Q3_K, Q4_K, Q5_K, Q6_K, IQ2_XXS, IQ3_XXS, IQ2_S, IQ1_S,
IQ1_XXXS) plus a native Q8_0 arm. The IQ codebooks reuse
cuda_quant_iq_tables.cuh as the single source — pure __device__
syntax, no CUDA — so no transcription drift is possible.

Why

The ROCm backend had no quantized-weight GEMM provider: every GGUF
k-quant weight on an AMD card computed off device (bf16 expansion at
load). AGENTS.md makes quantized arms a standing requirement. Upstream
pins an RDNA3 W4A16 family in csrc/rocm (gptq_gemm_rdna3 gated
VLLM_ROCM_GFX1100) but nothing in this tree can reach it yet — no
AWQ/GPTQ consumer exists — so that family stays under ## Owed in the
spec (W2), per the reachability rule.

HIP delta that shaped the implementation (measured on gfx1100 /

ROCm 7.14.0)

gfx1100 exposes NO signed byte dot: __dp4a does not exist;
__builtin_amdgcn_sdot4 needs target feature dot1-insts, which clang
rejects for this target ("not a recognized feature"); the one hardware
dot that assembles (v_dot4_i32_iu8) is UNSIGNEDxUNSIGNED and computes
wrong over signed Q8 activations (verified by device experiment:
byte 0xFD contributed +253x, not -3x). The integer dots therefore ship
as portable scalar loops in the CPU reference's exact accumulation
order — which is what makes the bit-exactness gate achievable. The ISA
dots are recorded as W2 levers with their own exactness arguments.
__shfl_down_sync also requires a 64-bit mask on this target.

Unsupported dtypes (Q4_0/MXFP4, Q8_0-activation) THROW naming the dtype
instead of falling back to a host kernel that cannot follow device
pointers on a discrete card; VT_GGUF_KEEP_QUANT=0 restores load-time
bf16 expansion.

How to verify

On gfx1100 / ROCm 7.14.0 (container rocm-dev:7.14.0):

cmake -B build-hip -S . -GNinja -DCMAKE_BUILD_TYPE=Release \
  -DVLLM_CPP_HIP=ON \
  -DCMAKE_EXE_LINKER_FLAGS="-Wl,-rpath,/opt/rocm/lib -Wl,-rpath-link,/opt/rocm/lib" \
  -DCMAKE_SHARED_LINKER_FLAGS="-Wl,-rpath,/opt/rocm/lib -Wl,-rpath-link,/opt/rocm/lib"
cmake --build build-hip -j 8
ctest --test-dir build-hip -R 'test_rocm_quant_dot|test_gguf_keep_quant|test_ops_quant_dot'
VT_GGUF_KEEP_QUANT=1 ./build-hip/examples/vllm-cli \
  --model <qwen3.5 gguf> --device auto --temperature 0 \
  --max-tokens 64 --prompt "The capital of France is"

Measured on this tree:

• test_rocm_quant_dot: 132,094 assertions green — all ten encodings x
  decode/prefill shapes x broadcast/per-row grouped arms, output buffer
  poisoned before every launch so a silent no-op fails (ENG-EXPERT-STREAM: the CUDA arm for IQ1_S and IQ1_XXXS, verified on GB10 #967 class).
• Focused sweep ctest -R 'rocm|cross_device|quant': 20 of 21 pass;
  the one failure is the PRE-EXISTING MoeSiluMul bf16 exactness case
  tracked in Characterize Qwen3.5-0.8B CPU against ROCm numerics on gfx1100 #1588 (present on main before this change).
• End-to-end Qwen3.5-0.8B and Qwen3.5-4B Q4_K_M decode on gfx1100:
  deterministic across runs, coherent output, warm ~27.6 tok/s at 4B.

What remains unverified or out of scope

• No fresh independent review yet: this was implemented in-session under
  a recorded deviation (subagent runtime had no model configured). A
  fresh review pass against the spec is still owed before merge.
• Performance vs the CPU tier is parity-at-these-shapes, not a claimed
  win; the unsigned-dot/dot2 ISA levers are W2 with their own exactness
  arguments.
• W2 upstream RDNA3 family port + AWQ/GPTQ loader consumer stays owed in
  the spec's ## Owed.
• docs/ENVIRONMENT.md keep-quant row should gain the ROCm throw caveat;
  riding the follow-up rather than widening this PR.

Closes #1587

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true

AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]

[ghazni101]

Closing at the author's direction: work continues on the author's fork (ghazni101/vllm.cpp) rather than upstream. The branch row/KERNEL-QUANT-CIQ-GEMM-ROCM remains available on the fork.