[1/2] cuda: native TQ4_1S MMQ prefill on CDNA

[jasstrong]

Series: TQ4_1S MMQ prefill on AMD — part 1/2

• [1/2] this PR — native TQ4_1S MMQ prefill on CDNA
• [2/2] [2/2] cuda: TQ4_1S MMQ prefill on RDNA3 and RDNA2 #337 — extend the same path to RDNA3 (WMMA) and RDNA2 (dp4a), branched on this

Independent of #334 / #335, though those add the MoE MUL_MAT_ID coverage this type otherwise lacks.

Overview

TQ4_1S has no MMQ path, so prefill dequantises the weights to f16 and calls hipBLAS. That throws away the point of the format: the 5bpw footprint is paid for at load time and then not used at compute time.

This adds a TQ4_1S tile loader, MMQ type traits, and dispatch. The activation is pre-rotated so the block-local turbo WHT cancels against the weights, which lets the weights stay in their rotated-domain int8 centroid form and be consumed directly by the stock MFMA-i8 MMQ kernel.

The path is env-gated behind GGML_TQ_MMQ so it can be A/B'd against the hipBLAS fallback.

Additional information

MI210 (gfx90a), ROCm 7.2.3, Qwen3.6-35B-A3B-TQ4max, measured against this same branch with the path disabled:

	prefill (pp512)	decode (tg128)	perplexity
MMQ on	2124.09 ± 17.38	92.58 ± 0.13	3.2653
MMQ off	970.99 ± 20.24	95.08 ± 0.15	3.2653

2.19x prefill at identical perplexity, and roughly 13GB less VRAM than converting the experts to Q8_0. The ~2.6% decode cost is real; this is a prefill-only path.

test-backend-ops -o MUL_MAT with GGML_TQ_MMQ=1: 147/147 tq4_1s cases pass.

Note for reviewers: the tile geometry is not free to choose

TQ4_1S reuses GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, because its per-16 scales have the same shape. It therefore has to take the same tile geometry Q3_K and IQ2_XS use on CDNA — occupancy 1, I=128. Other geometries do not match what that layout expects and write out of bounds; the symptom is non-deterministic NaN plus test-harness sentinel mismatches rather than a clean failure.

And a warning about how to validate it

Because the path is env-gated, test-backend-ops run without GGML_TQ_MMQ=1 silently exercises the hipBLAS fallback and reports a clean pass for a kernel that was never invoked. llama-bench does not check output at all, so a broken kernel still reports a plausible tokens/s. During development we spent a while trusting exactly that combination without realising the path under test was switched off.

This is the failure mode AGENTS.md already describes — "a green run means the cases that ran passed, not that your change was exercised" — so the reliable check is llama-perplexity with the flag on versus off. Any corpus works, since only the delta matters.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES — Claude Code (Opus 5) was used to develop and verify this: kernel and dispatch work, the commit message, and the benchmark/perplexity runs on the MI210. I reviewed the change and am responsible for every line of it.

[TheTom]

Pushed one commit to this branch (03ee3ed): the CUDA and MUSA jobs were failing to compile because tq4_1s_qs_to_int8x4() calls __builtin_amdgcn_perm unconditionally. It now keeps v_perm_b32 under GGML_USE_HIP and uses __byte_perm(v.x, v.y, 0x5140) elsewhere, which is the same byte order. The MMQ path is still AMD-only at dispatch, this just makes the header compile on nvcc and MUSA.

Same builtin shows up in #337 (inherited) and twice more in #338's v_perm_b32 decode patch, so those need the same guard. I'll carry the fix over once CI is green here. Landing order: #336, then #337, then #338.

The windows job failure on this PR is test_completion_unified in the server unit tests (context size exceeded), which has nothing to do with these files. I'll rerun it.

[jasstrong]

Thanks for the guard, and for carrying it into #337 and #338. I checked the two selectors against each other: __builtin_amdgcn_perm(v.y, v.x, 0x05010400) and __byte_perm(v.x, v.y, 0x5140) both produce [x0, y0, x1, y1], so the nvcc/MUSA form is exact, not just close.

One thing on the ubuntu-22-hip-quality-check job, since it is red on #337: it fails with the same four kernels on the already-merged #334 and #335, so it is the branch baseline rather than anything in these PRs. Two of the four (rwkv_wkv_f32 and the mul_mat_q<Q2_K, 64, fallback> instance) are on upstream's current ignore list in scripts/hip/gcn-cdna-vgpr-check.py; the fork's copy of the list predates that. The other two are the flash_attn_ext_vec<128, 2, turbo3, q8_0> instances at 259 VGPRs (255 + 4 spilled), so just over the line. I can send a small PR that syncs the ignore list with upstream and, if the 4-VGPR spill in the turbo3 vec kernel turns out to be easy to trim, includes that too.