RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations

[liangliangchang]

Summary

RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations and a documented HIP/clang codegen workaround.

MMQ optimizations (mmq.cuh)

• Occupancy / tile sizing: reduce mmq_x when LDS limits workgroups/CU; dual-WG mmq_x=64 for narrow-N K-quant prefill (Q6_K FFN-down), gated to K-quants only.
• Q4_K: improved dequant/vec_dot for WMMA; ldmatrix-friendly qs layout; precompute scale factors before WMMA accumulate.
• Q6_K: hoist vec_dot scales; mmq_x=128 for narrow-N prefill; batched B-tile loads (load_generic); pipelined B loads in vec_dot; restore load_ldmatrix path.
• Tiny-M prefill: specialized path for gate Q8_0 on RDNA3.5.
• LDS / pipelining: tuned LDS layout; activation y-pipeline experiments were tried and the net-negative y-pipeline was dropped.

Benchmarks (gfx1151, ROCm)

Model	TTFT Before	TTFT After
Qwen3.5-9B_Q4_K_M_GGUF_128	117.21	104.07
Qwen3.5-4B_Q4_K_M_GGUF_128	70.65	64.43
llama-2-7b.Q4_K_M_128	100.6	87.85
Qwen2.5-7B-Instruct_Q4_K_M_GGUF_128	96.4	85.04
Qwen2.5-0.5B-Instruct_Q4_K_M_GGUF_128	14.4	12.04

Run with llama-bench (-p 128 -n 0 -ngl 999, flash-attn on). Rebuild MMQ after mmq.cuh changes: touch ggml/src/ggml-cuda/mmq.cuh && cmake --build build -j$(nproc).

Files changed

• ggml/src/ggml-cuda/mmq.cuh — RDNA3.5 MMQ optimizations + MMQ_CODEGEN_SPLIT_COLD
• docs/backend/HIP-MMQ-compiler-codegen.md — compiler issue write-up (new)

Test plan

• Build with GGML_HIP=ON, target gfx11 / gfx1151
• llama-bench on Q4_K_M and Q6_K models: pp128, narrow-N FFN-down shapes
• Compare tok/s vs gfx11 base; confirm no regression vs pre-strip branch with codegen split
• Optional ASM check: compile template-instances/mmq-instance-q4_k.cu and verify hot-path ds_load_b128 ≈ 48 (not 24) per mmq_x=128 kernel body — see doc

AI usage disclosure

Used Cursor/AI for profiling analysis, ASM comparison, codegen-split investigation, and documentation.

[mgehre-amd]

This approach looks brittle. Let's please investigate further what affects the unroll decision. Can you try with an explicit unroll pragma?

[liangliangchang]

I ran this change with all models. Here is the current result: http://fisweb:8080/wrk/xcohdnobkup2/lichang/regression_comparison.html, will investigate it more.

[mgehre-amd]

I suggest to split this PR for improvements across kernels types (e.g. one PR to focus on Q8_0).
Then extract all shapes of of Q8_0 GEMMs appearing in our QoR suite, and use/build a kernel benchmark that runs the Q8_0 kernel with those shapes. (That benchmark should be committed; I don't know if llama.cpp already has one). Then you can run all shapes before your changes, and after your changes, and ensure that there are no regressions and some improvements.

Important: Benchmarking can be quite hard (dynamic frequency, thermal ramp up, ...) so it's important the the kernel benchmark reproduces the same per-kernel GPU time that the trace of the full model shows in our QoR suite.

[liangliangchang]

I suggest to split this PR for improvements across kernels types (e.g. one PR to focus on Q8_0). Then extract all shapes of of Q8_0 GEMMs appearing in our QoR suite, and use/build a kernel benchmark that runs the Q8_0 kernel with those shapes. (That benchmark should be committed; I don't know if llama.cpp already has one). Then you can run all shapes before your changes, and after your changes, and ensure that there are no regressions and some improvements.

Important: Benchmarking can be quite hard (dynamic frequency, thermal ramp up, ...) so it's important the the kernel benchmark reproduces the same per-kernel GPU time that the trace of the full model shows in our QoR suite.

Thanks a lot @mgehre-amd, this sounds good to me. I'm aware we have one in vLLM for kernel benchmarking, will learn from it, and integrate with the existing one in llama.cpp.

[Annieren]

Looks like rocprofiler-compute (https://rocm.docs.amd.com/projects/rocprofiler-compute/en/latest/) can do a pretty good on profiling kernel but I haven't dig it deep. Fyr.