perf(ggml-cuda mmq): Q4_K/Q5_K MMQ tuning for RDNA4 — +7.6% Qwen3.6-27B prefill

[DeanoC]

What

Tuning of the Q4_K/Q5_K MMQ path (ggml/src/ggml-cuda/mmq.cuh) for RDNA4 / gfx1201, found autonomously by the geo-evo kernel loop over 9 iterations:

• relax __launch_bounds__ minBlocks 2 → 1 (RDNA4 wave32 occupancy)
• int2-vectorized y-tile flat copy in load_tiles
• int4-vectorized x_dm LDS write in load_tiles_q4_K / load_tiles_q5_K (the biggest single win)

MMQ is the compute-bound quantized-matmul path (prefill + dflash spec-decode verify), so kernel efficiency translates here — unlike the batch-1 matvec/decode path, which is VRAM-bandwidth-bound (a parallel MMVQ experiment got +40% microbench but 0 real decode gain; not included).

Results (gfx1201 / Radeon AI PRO R9700)

• Microbench (test-backend-ops perf, MUL_MAT q4_K m=4096 n=512 k=14336): 1139 → ~1037 µs = +9.8%
• Real model (unsloth/Qwen3.6-27B-Q4_K_M, pinned clocks, same-session A/B): prefill pp512 583.6 → 627.7 tok/s = +7.6%; decode tg128 unchanged (MMQ doesn't touch batch-1 decode)

Notes

• Regressed approaches were tried and reverted (documented in the loop handoff): dsB VGPR preload, stream-K (−24% on RDNA4), y-tile half2 register prefetch (−7.9%).
• mmq.cuh is byte-identical between ac06e543 (base, perf/luce-verify-kernels) and the 30c9d7d base the change was measured on — applies cleanly.
• Complementary to feat/rdna35-mmq-tile-override.
• Staged on our fork; upstream to Luce-Org/lucebox-ggml when ready.

🤖 Generated with Claude Code

[DeanoC]

Updated to the iter-9 kernel: added the int4-vectorized x_dm LDS write (load_tiles_q4_K/q5_K), improving the result from +6.4% → +7.6% real 27B prefill (microbench +7.5% → +9.8%). Same-session pinned-clock A/B: stock 583.6 → optimized 627.7 tok/s.