fix: restore ROCm Q8 prequantized decode kernels

[Scorp1o117]

Summary

Restore the ROCm Q8 prequantized decode kernels that were dropped by ef8d923 ("Add ROCm GLM 5.2 support"). On DeepSeek-V4-Flash IQ2XXS this regression halves decode throughput: ~15 t/s → 8.4 t/s on Strix Halo (gfx1151, ROCm 7.14). All four restored kernels already exist in rocm/ds4_rocm_q8.cuh; this PR only reconnects their call sites (2 files, +155 lines, pure additions).

Root cause

ef8d923 refactored the ROCm matmul/attention code and removed the prequantized integer (Q8xQ8 + dp4a) single-token decode paths, leaving only fp32 dequant matmuls:

// old (15 t/s):
quantize_q8_0_f32_kernel<<<...>>>(xq, xscale, x, ...);       // quantize activation once
matmul_q8_0_preq_rows_w32_kernel<<<...>>>(out, w, xq, xscale, ..., use_dp4a);  // int dp4a dot

// new after ef8d923 (8.4 t/s):
matmul_q8_0_f32_sharedx_warp_rows_w32_kernel<<<...>>>(out, w, (float*)x->ptr, ...);  // fp32 dequant

Q8_0 weights cover the every-token decode path of DeepSeek (attention projections, shared experts, output projection — AProjQ8/SExpQ8/OutQ8 in the IQ2XXS layout), while prefill (n_tok>1, batch kernels) is unaffected — matching the observed profile: prefill ~150 t/s unchanged, decode halved.

Bisect evidence (same machine, same ROCm 7.14, same model file, only code version changed; ds4-bench --ctx-start 2048 --ctx-max 2048 --gen-tokens 64-128):

commit	date	decode t/s
80ebbc3 (baseline)	6/17	14.97
519c4d8 (#51)	7/12	15.08
ef8d923 ("Add ROCm GLM 5.2 support")	7/18	8.43 ← regression
646db66 (#39)	7/20	8.42
fa8b0b9 (#26)	7/22	8.43
54b36ed (main)	7/28	8.44

(Commits #45–#50 in between fail to compile on the ROCm path — no CI coverage for ROCm — and were skipped per bisect rules; ef8d923 is the first compilable bad commit.)

Changes

Restored four preq decode call sites, all guarded by !g_quality_mode to preserve the historical q8_prequant_decode = !g_quality_mode semantics (quality mode keeps the fp32 path for numeric precision):

1. cuda_matmul_q8_0_tensor_labeled — matmul_q8_0_preq_rows_w32_kernel (rows_per_block: quality 8 / normal 1, matching old q8_decode_rpb)
2. ds4_gpu_matmul_q8_0_pair_tensor — matmul_q8_0_pair_preq_warp8_kernel
3. cuda_matmul_q8_0_hc_expand_tensor_labeled — matmul_q8_0_hc_expand_preq_rows_w32_kernel (rows_per_block: quality 8 / normal 16, matching old q8_hc_decode_rpb)
4. ds4_gpu_attention_output_low_q8_tensor — grouped_q8_0_a_preq_warp8_kernel (rows_per_block: quality 8 / normal 32, matching old attn_out_low_decode_rpb)

Each preq block keeps the existing fp32 path as fallback (tmp alloc or launch failure → falls through to fp32; strictly more robust than the old hard-fail behavior). Prefill/batch paths and GLM paths are untouched.

Verification

Benchmark: DeepSeek-V4-Flash IQ2XXS (...-imatrix-0731.gguf, 80.8GB), ROCm 7.14, gfx1151, ctx=2048, 128 gen tokens, full residency:

state	decode t/s
main (before)	8.44
+ matmul preq	9.42
+ attention preq	10.78
+ pair + hc_expand preq (this PR)	15.03
old baseline (80ebbc3)	14.97

Real chat smoke test: ds4 --rocm -m <model> -p "..." -n 32 → 16.63 t/s, first-token latency 120 ms → 68 ms, output correct, no errors.

Other: make rocm -j32 clean; ds4-eval --self-test-extractors passes. Note: the full make test suite requires CUDA tooling ($(NVCC), Makefile:247) and cannot run on a ROCm-only host — pre-existing environment limitation, unrelated to this change.

Notes for reviewers

• The four kernels themselves are untouched upstream code (rocm/ds4_rocm_q8.cuh); only call sites are restored.
• Launch parameters (use_dp4a=1, grid/block shape, 16-byte-aligned tmp layout) match the pre-ef8d923 code.
• Optionally consider re-introducing q8_decode_rpb / attn_out_low_decode_rpb / q8_hc_decode_rpb as runtime config (this PR uses equivalent hardcoded defaults) and adding a ROCm compile check to CI — six commits between 7/12 and 7/18 currently fail to compile on the ROCm path.

[kyuz0]

I think I have these already in this PR: #623. I'm just checking if anything is missing.

[Scorp1o117]

Thanks for the heads-up, kyuz0! We independently bisected this to the same root cause — ef8d923 ("Add ROCm GLM 5.2 support") — across the 51 commits between 80ebbc3 (6/17) and 54b36ed (7/28), and restored the same four preq decode paths (single matmul, paired, HC-expand, attention output). Our bench matches yours:

• 8.44 → 15.03 t/s (ROCm 7.14, gfx1151, IQ2XXS ...-imatrix-0731.gguf, ctx 2048, 128 gen tokens, full residency)
• Real chat smoke test: 16.63 t/s, first-token latency 120 ms → 68 ms

Your #623 is more complete: model-family guard so GLM is untouched, DS4_ROCM_DSV4_PREQUANT_DECODE=0 diagnostic rollback, and the Q4 SSD streaming fix. Nothing is missing on our side — I'll close #640 in favor of #623 and post the independent bisect data there as cross-validation.

[Scorp1o117]

Closing in favor of #623 (kyuz0): same root cause and fix, independently bisected to ef8d923; my bench data (8.44 → 15.03 t/s) and independent bisect table are posted on #623 as cross-validation. Thanks kyuz0!