RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations by liangliangchang · Pull Request #32 · AMD-Ecosystem/llama.cpp

liangliangchang · 2026-06-30T23:10:46Z

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 · 2026-07-01T08:33:42Z

+#if defined(GGML_USE_HIP)
+// ROCm/clang WMMA codegen workaround: an unprovable-false branch keeps full unroll/inline
+// on the hot K-loop. blockDim.z is always 1 for mul_mat_q. See docs/backend/HIP-MMQ-compiler-codegen.md.
+#define MMQ_CODEGEN_SPLIT_COLD (threadIdx.z > 0)


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

This is resolved by adding syncthread to let the compiler split the BB, instead of using duplicated dead if-else condition.

- Add kernel dimension logging for profiling analysis - Reduce mmq_x when LDS limits WGs/CU - Improve LDS layout and activation load pipelining for prefill Co-authored-by: Cursor <cursoragent@cursor.com>

- Add opt-in HIP MMQ phase profiling to split dequant vs MMA time - Refactor Q4_K qs load for WMMA ldmatrix layout - Precompute Q4_K vec_dot scale factors before WMMA accumulate - Gate dimension logs behind GGML_CUDA_MM_LOG - Fix MMQ phase profiler clock64 deltas on HIP Co-authored-by: Cursor <cursoragent@cursor.com>

Co-authored-by: Cursor <cursoragent@cursor.com>

- Hoist Q6_K vec_dot scales on RDNA3.5 WMMA - Use mmq_x=128 for Q6_K narrow-N prefill; dual-WG mmq_x=64 where needed - Pipeline Q6_K B-tile loads; batch B loads with load_generic - Remove always-on MMQ profiling overhead on hot path - Limit dual-WG mmq_x downgrade to K-quants Co-authored-by: Cursor <cursoragent@cursor.com>

Add MMQ_HIP_TILE_BARRIER() tile-scope __syncthreads() at K-loop phase boundaries on HIP+RDNA3.5 to restore pp128 throughput after removing dead threadIdx.z cold/hot loop split. Co-authored-by: Cursor <cursoragent@cursor.com>

The pipelined load_ldmatrix for B0 omitted the k01 term, reading the wrong activation slice and corrupting Q6_K matmul results. Restores sanity-check accuracy on Qwen2.5 and Llama-2 while keeping the prefetch. Co-authored-by: Cursor <cursoragent@cursor.com>

liangliangchang · 2026-07-10T05:02:47Z

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 · 2026-07-10T09:50:32Z

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.

- Add kernel dimension logging for profiling analysis - Reduce mmq_x when LDS limits WGs/CU - Improve LDS layout and activation load pipelining for prefill Co-authored-by: Cursor <cursoragent@cursor.com>

- Add opt-in HIP MMQ phase profiling to split dequant vs MMA time - Refactor Q4_K qs load for WMMA ldmatrix layout - Precompute Q4_K vec_dot scale factors before WMMA accumulate - Gate dimension logs behind GGML_CUDA_MM_LOG - Fix MMQ phase profiler clock64 deltas on HIP Co-authored-by: Cursor <cursoragent@cursor.com>

Co-authored-by: Cursor <cursoragent@cursor.com>

- Hoist Q6_K vec_dot scales on RDNA3.5 WMMA - Use mmq_x=128 for Q6_K narrow-N prefill; dual-WG mmq_x=64 where needed - Pipeline Q6_K B-tile loads; batch B loads with load_generic - Remove always-on MMQ profiling overhead on hot path - Limit dual-WG mmq_x downgrade to K-quants Co-authored-by: Cursor <cursoragent@cursor.com>

Add MMQ_HIP_TILE_BARRIER() tile-scope __syncthreads() at K-loop phase boundaries on HIP+RDNA3.5 to restore pp128 throughput after removing dead threadIdx.z cold/hot loop split. Co-authored-by: Cursor <cursoragent@cursor.com>

The pipelined load_ldmatrix for B0 omitted the k01 term, reading the wrong activation slice and corrupting Q6_K matmul results. Restores sanity-check accuracy on Qwen2.5 and Llama-2 while keeping the prefetch. Co-authored-by: Cursor <cursoragent@cursor.com>

liangliangchang · 2026-07-10T19:31:03Z

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 · 2026-07-10T22:01:26Z

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.

…llama.cpp into lichang.mmq_opt

Pull latest gfx11 CI fixes (hip-quality-check UI build off) and roofline/multistream/MoE updates.

Register Q4_K, Q6_K, Q8_0, and Q4_0 MUL_MAT shapes from prefill profiling in test-backend-ops eval and perf for gfx1151 regression against baseline.

Dedicated vec_dot_q4_0_q8_1_mma regressed ~35% on small-batch prefill (mmq_x=32, MTP-style prompts) while pp128 (mmq_x=128) stayed fine. Fall back to the baseline q8_0 alias for mmq_x<=32 and add kernel cases for n=22/32/128 to guard the regression. Co-authored-by: Cursor <cursoragent@cursor.com>

Hoisted scale/dA arrays in Q6_K, Q8_0, and Q4_0 WMMA vec_dot spilled to local memory and broke prefill accuracy; compute them inline instead. Route Q4_K through vec_dot_q8_1_q8_1_mma with RDNA3.5 loaders and gfx11 LDS layout so MTP quality is preserved while keeping dual-WG and narrow-N tile sizing. Co-authored-by: Cursor <cursoragent@cursor.com>

Extract add_rdna35_mmq_prefill_cases() so eval and perf suites share one shape list instead of maintaining duplicate emplace_back blocks. Co-authored-by: Cursor <cursoragent@cursor.com>

Remove redundant q4k_legacy_tile guards; RDNA3.5 heuristics already gate on GGML_CUDA_CC_IS_RDNA3_5. Inline the Q4_K off-RDNA3.5 LDS sizing check with a clear comment instead of the bisection-era legacy naming. Co-authored-by: Cursor <cursoragent@cursor.com>

Register-staged prefetch of activation tiles across vec_dot produced wrong tile_y on gfx1151; use global→LDS loads with MMQ_HIP_TILE_BARRIER instead.

Prefetch the next by0 tile only after both vec_dot WMMA phases complete. Keeping y0_next_cache live during the k=MMQ_TILE_NE_K vec_dot corrupted activations on gfx1151 and dropped Q4_0 27B MTP accept to ~52%.

Exclude Q4_K from dual-WG downshift so pp128 prefill keeps mmq_x=128, and skip HIP tile barriers when mmq_x<=32 where sync cost scales with K without scheduler benefit on gfx1151. Co-authored-by: Cursor <cursoragent@cursor.com>

liangliangchang · 2026-07-22T01:07:19Z

label	type	m	n	k	baseline us	opt us	speedup
Q4_K pp128 wide	Q4_K	12288	128	4096	532.3	549.1	0.97x
Q6_K pp128 narrow-N	Q6_K	4096	128	12288	1221.3	1069.2	1.14x
Q4_K pp128	Q4_K	4096	128	4096	234.9	238.9	0.98x
Q4_K pp128 alt-K	Q4_K	8192	128	4096	365.4	369.6	0.99x
Q6_K pp128	Q6_K	1024	128	4096	138.5	111.9	1.24x
Q4_K pp128 small-K	Q4_K	1024	128	4096	79.8	79.9	1.00x
Q8_0 tiny-M gate	Q8_0	32	128	4096	80.7	71.0	1.14x
Q4_0 tiny-M gate	Q4_0	32	128	4096	95.6	96.3	0.99x
Q8_0 tiny-M n=8	Q8_0	8	128	4096	161.9	153.5	1.05x
Q4_0 tiny-M n=8	Q4_0	8	128	4096	156.1	160.0	0.98x
Q8_0 tiny-M n=16	Q8_0	16	128	4096	141.6	132.5	1.07x
Q4_0 tiny-M n=16	Q4_0	16	128	4096	140.8	148.3	0.95x
Q8_0 tiny-M n=64	Q8_0	64	128	4096	65.2	66.2	0.98x
Q4_0 tiny-M n=64	Q4_0	64	128	4096	84.1	80.8	1.04x
Q4_K narrow-M n=1	Q4_K	4096	1	4096	32.3	32.9	0.98x
Q4_K narrow-M n=4	Q4_K	4096	4	4096	71.8	72.3	0.99x
Q4_K narrow-M n=128	Q4_K	4096	128	4096	234.9	238.9	0.98x
Q4_0 MTP n=22 wide	Q4_0	12288	22	4096	287.3	282.4	1.02x
Q4_0 MTP n=22	Q4_0	4096	22	4096	127.1	119.3	1.06x
Q4_0 MTP n=22 alt-K	Q4_0	4096	22	12288	363.7	348.3	1.04x
Q4_0 small-n n=32 wide	Q4_0	12288	32	4096	307.0	287.7	1.07x
Q4_0 small-n n=32	Q4_0	4096	32	4096	133.7	129.3	1.03x
Q4_0 small-n n=32 alt-K	Q4_0	4096	32	12288	370.5	364.0	1.02x
Q4_0 pp128 wide	Q4_0	12288	128	4096	563.7	550.9	1.02x
Q4_0 pp128	Q4_0	4096	128	4096	223.4	221.5	1.01x
Q4_0 pp128 alt-K	Q4_0	4096	128	12288	709.2	703.6	1.01x

liangliangchang · 2026-07-22T01:12:08Z

test_name	baseline_ttft_ms	opt_ttft_ms	ttft_delta_pct	baseline_decode_tps	opt_decode_tps	decode_delta_pct	baseline_status	opt_status
GLM-4.7-Flash_Q4_K_M_GGUF_128	None	None		None	None		FAIL	FAIL
Gemma-2B_Q4_K_M_GGUF_128	34.37330483213474	30.887919083304475	10.14	114.217	110.3415	-3.39	PASS	PASS
Gemma-2B_Q4_K_M_GGUF_8000	3748.257645859933	3433.1301349434716	8.41	104.483	102.6105	-1.79	PASS	PASS
Gemma-3-12B-IT_Q4_K_M_GGUF_4096	5264.125166913636	4919.978282908361	6.54	24.036099999999998	24.15535	0.5	PASS	PASS
Gemma-3-12B-IT_Q4_K_M_GGUF_VLM	1276.149119992624	1188.8642749981955	6.84	26.05266372707937	26.171224109283408	0.46	PASS	PASS
Gemma-3-4B-IT_VLM	965.2103344997158	904.0737810137216	6.33	68.50875053321738	69.11345327668894	0.88	PASS	PASS
Gemma-4-12B-it_Q4_K_M_GGUF_128	152.4273156894746	146.04003322410756	4.19	26.59685	26.82245	0.85	PASS	PASS
Gemma-4-26B-A4B-IT_VLM	673.6354265158297	648.4306405181997	3.74	48.50977721651645	48.532465610686714	0.05	PASS	PASS
Gemma-4-31B-IT_VLM	1754.980024008546	1643.60460799071	6.35	10.630898936570691	10.687247623982191	0.53	PASS	PASS
Gemma-4-E2B-IT_Q4_K_M_GGUF_128	44.657420022224045	40.78667295461208	8.67	92.95895	93.17615	0.23	PASS	PASS
Gemma-4-E2B-IT_Q4_K_M_GGUF_3968	1655.0090299761841	1600.0032258129552	3.32	88.03105	88.0478	0.02	PASS	PASS
Janus-Pro-1B_VLM	166.9617005100008	156.46417549578473	6.29	152.50403633566822	152.86600070979125	0.24	PASS	PASS
Janus-Pro-7B_VLM	497.0807415083982	454.38982799532823	8.59	44.27183738245855	44.340938867115696	0.16	PASS	PASS
Llama-2-7B_Q4_K_M_GGUF_128	99.92700644451125	88.93706313141841	11.0	45.71805	45.8912	0.38	PASS	PASS
Llama-2-7B_Q4_K_M_GGUF_1920	1690.2896381723742	1598.4481732318206	5.43	34.593050000000005	34.5792	-0.04	PASS	PASS
Llama-3.1-8B-Instruct_Q4_K_M_GGUF_128	104.05658076579141	94.79620667054246	8.9	41.875550000000004	42.0974	0.53	PASS	PASS
MiniCPM-V-2_6_VLM	1425.291773499339	1315.1637054979801	7.73	45.16911362157315	45.47620036284567	0.68	PASS	PASS
MiniCPM-o-2_6_VLM	1428.7269335036399	1339.6003385132644	6.24	45.1603646547412	45.410775680079716	0.55	PASS	PASS
Qwen2.5-0.5B-Instruct_Q4_K_M_GGUF_128	14.521477605726002	12.391813658102114	14.67	291.6845	291.592	-0.03	PASS	PASS
Qwen2.5-0.5B-Instruct_Q4_K_M_GGUF_3968	412.92342081246437	367.3291798968739	11.04	285.027	285.4145	0.14	PASS	PASS
Qwen2.5-3B-Instruct_Q4_K_M_GGUF_128	48.92237013595067	44.596117698910014	8.84	85.8288	85.9476	0.14	PASS	PASS
Qwen2.5-3B-Instruct_Q4_K_M_GGUF_3968	1557.9051515306182	1443.788489352606	7.33	83.43299999999999	83.7156	0.34	PASS	PASS
Qwen2.5-7B-Instruct_Q4_K_M_GGUF_128	91.60720547925598	83.65001617451486	8.69	44.90915	45.262	0.79	PASS	PASS
Qwen2.5-7B-Instruct_Q4_K_M_GGUF_3968	3057.034337706762	2840.9314646762964	7.07	42.059250000000006	42.5624	1.2	PASS	PASS
Qwen2.5-VL-3B-Instruct_VLM	2087.7546925039496	2051.2509319814853	1.75	85.48288229317032	86.13712551690409	0.77	PASS	PASS
Qwen2.5-VL-7B-Instruct_VLM	2484.2752645054134	2400.7497775019147	3.36	44.49583912886015	44.86812817593534	0.84	PASS	PASS
Qwen3-1.7B_Q4_K_M_GGUF_128	28.765790283904867	26.420324660380142	8.15	136.23149999999998	136.24650000000003	0.01	PASS	PASS
Qwen3-1.7B_Q4_K_M_GGUF_3968	1147.9704386039778	1073.4130097223951	6.49	104.265	108.744	4.3	PASS	PASS
Qwen3-30B-A3B-Instruct-2507_Q4_K_M_GGUF_128	112.07228662487304	107.96676647969298	3.66	71.0249	75.3551	6.1	PASS	PASS
Qwen3-4B_Q4_K_M_GGUF_128	62.50442536214722	56.57083508275696	9.49	69.60409999999999	69.6219	0.03	PASS	PASS
Qwen3-4B_Q4_K_M_GGUF_3968	2377.4003175459093	2192.9860921131535	7.76	58.13965	58.343450000000004	0.35	PASS	PASS
Qwen3-8B_Q4_K_M_GGUF_128	103.09029259924453	93.2784836416503	9.52	40.881	41.18445	0.74	PASS	PASS
Qwen3-8B_Q4_K_M_GGUF_3968	3725.8740733437558	3515.4242783989225	5.65	36.751000000000005	36.89335	0.39	PASS	PASS
Qwen3-VL-4B-Instruct_VLM	882.2759429895086	842.8677835036069	4.47	67.62743839624623	67.6738017245874	0.07	PASS	PASS
Qwen3.5-35B-A3B_Q4_K_M_GGUF_128	112.14691245531645	109.32275407931878	2.52	53.7791	58.4237	8.64	PASS	PASS
Qwen3.5-4B_Q4_0_GGUF_baseline_code	70.638	68.41	3.15	65.76539576135448	65.78367433397318	0.03	PASS	PASS
Qwen3.5-4B_Q4_0_GGUF_baseline_list	70.091	68.081	2.87	66.33887458050418	66.42969068512457	0.14	PASS	PASS
Qwen3.5-4B_Q4_0_GGUF_baseline_prose	69.668	66.851	4.04	65.85863809214844	65.92843102543128	0.11	PASS	PASS
Qwen3.5-4B_Q4_0_GGUF_mtp_code	81.462	78.318	3.86	116.24156158913839	118.121490315219	1.62	PASS	PASS
Qwen3.5-4B_Q4_0_GGUF_mtp_list	81.596	77.315	5.25	133.0991484264288	135.4891157077048	1.8	PASS	PASS
Qwen3.5-4B_Q4_0_GGUF_mtp_prose	80.179	77.043	3.91	83.31788732402869	84.75081112053469	1.72	PASS	PASS
Qwen3.5-4B_Q4_K_M_GGUF_128	70.92768423881573	62.771560559843856	11.5	58.714	61.2632	4.34	PASS	PASS
Qwen3.5-4B_Q4_K_M_GGUF_baseline_code	65.054	57.877	11.03	60.844811960629755	61.324298220062246	0.79	PASS	PASS
Qwen3.5-4B_Q4_K_M_GGUF_baseline_list	63.869	57.217	10.42	61.14697334471546	61.867225655246706	1.18	PASS	PASS
Qwen3.5-4B_Q4_K_M_GGUF_baseline_prose	63.716	56.579	11.2	60.86095427440429	61.37558019103149	0.85	PASS	PASS
Qwen3.5-4B_Q4_K_M_GGUF_mtp_code	74.129	68.467	7.64	95.44261512765449	98.6627254196619	3.37	PASS	PASS
Qwen3.5-4B_Q4_K_M_GGUF_mtp_list	74.2	67.553	8.96	110.13453689506987	115.37626235204692	4.76	PASS	PASS
Qwen3.5-4B_Q4_K_M_GGUF_mtp_prose	72.134	67.917	5.85	68.3816168331835	70.70785634991903	3.4	PASS	PASS
Qwen3.5-9B_Q4_0_GGUF_baseline_code	106.429	103.191	3.04	39.81374334587931	39.84587615104775	0.08	PASS	PASS
Qwen3.5-9B_Q4_0_GGUF_baseline_list	105.933	102.45	3.29	40.18899861623762	40.29180749048402	0.26	PASS	PASS
Qwen3.5-9B_Q4_0_GGUF_baseline_prose	104.686	102.334	2.25	39.88250862225562	39.91815530719671	0.09	PASS	PASS
Qwen3.5-9B_Q4_0_GGUF_mtp_code	118.373	114.58	3.2	79.32551104468921	79.84567428075017	0.66	PASS	PASS
Qwen3.5-9B_Q4_0_GGUF_mtp_list	117.051	113.871	2.72	86.38782033366026	87.51038113344818	1.3	PASS	PASS
Qwen3.5-9B_Q4_0_GGUF_mtp_prose	115.888	112.918	2.56	60.87427357081623	61.38837424199205	0.84	PASS	PASS
Qwen3.5-9B_Q4_K_M_GGUF_128	115.46458710134272	101.70999936431251	11.91	35.8283	36.0184	0.53	PASS	PASS
Qwen3.5-9B_Q4_K_M_GGUF_baseline_code	98.591	88.812	9.92	36.672451745929585	36.915148684990115	0.66	PASS	PASS
Qwen3.5-9B_Q4_K_M_GGUF_baseline_list	98.295	87.544	10.94	37.02609323170276	37.35382569093591	0.89	PASS	PASS
Qwen3.5-9B_Q4_K_M_GGUF_baseline_prose	96.081	87.412	9.02	36.77929519255896	37.02402209489712	0.67	PASS	PASS
Qwen3.5-9B_Q4_K_M_GGUF_mtp_code	111.275	100.476	9.7	62.33959243621258	64.65954808148655	3.72	PASS	PASS
Qwen3.5-9B_Q4_K_M_GGUF_mtp_list	109.166	100.045	8.36	67.64778055590558	70.7318527723419	4.56	PASS	PASS
Qwen3.5-9B_Q4_K_M_GGUF_mtp_prose	108.437	98.758	8.93	44.401251043525086	46.074675928206126	3.77	PASS	PASS
Qwen3.6-27B_Q4_0_GGUF_baseline_code	331.308	310.395	6.31	13.431219336288835	13.453987972678348	0.17	PASS	PASS
Qwen3.6-27B_Q4_0_GGUF_baseline_list	328.992	307.37	6.57	13.528888554056268	13.562711584949115	0.25	PASS	PASS
Qwen3.6-27B_Q4_0_GGUF_baseline_prose	328.321	305.015	7.1	13.3581802498148	13.429288241378663	0.53	PASS	PASS
Qwen3.6-27B_Q4_0_GGUF_mtp_code	364.06	336.597	7.54	35.99040692062806	36.39487780695495	1.12	PASS	PASS
Qwen3.6-27B_Q4_0_GGUF_mtp_list	358.937	334.26	6.88	36.737295918918626	37.08500246142614	0.95	PASS	PASS
Qwen3.6-27B_Q4_0_GGUF_mtp_prose	355.893	331.873	6.75	26.664493040784098	26.823771095096813	0.6	PASS	PASS
Qwen3.6-27B_Q4_K_M_GGUF_4096	11969.083763131366	11399.420844854048	4.76	12.02995	12.25985	1.91	PASS	PASS
Qwen3.6-27B_Q4_K_M_GGUF_baseline_code	297.195	263.427	11.36	12.434409863458075	12.514722796121323	0.65	PASS	PASS
Qwen3.6-27B_Q4_K_M_GGUF_baseline_list	296.993	264.469	10.95	12.54183994694064	12.623718630678013	0.65	PASS	PASS
Qwen3.6-27B_Q4_K_M_GGUF_baseline_prose	292.223	257.531	11.87	12.44692553160041	12.522593347203712	0.61	PASS	PASS
Qwen3.6-27B_Q4_K_M_GGUF_mtp_code	328.7	293.781	10.62	25.555255777579717	26.620261034039306	4.17	PASS	PASS
Qwen3.6-27B_Q4_K_M_GGUF_mtp_list	330.032	291.248	11.75	26.222239132424367	27.43997908105124	4.64	PASS	PASS
Qwen3.6-27B_Q4_K_M_GGUF_mtp_prose	324.936	288.404	11.24	19.67074843888712	20.45949345353573	4.01	PASS	PASS
Qwen3.6-35B-A3B_Q4_K_M_GGUF_128	216.8658980916648	113.20120630035464	47.8	51.7237	54.6784	5.71	PASS	PASS
Qwen3.6-35B-A3B_Q4_K_M_GGUF_4096	3079.6413601248096	2854.2161706397596	7.32	53.49225	57.63005	7.74	PASS	PASS
SmolLM2-1.7B-Instruct_Q4_K_M_GGUF_128	31.82247039329342	29.09781982311708	8.56	137.6415	137.51	-0.1	PASS	PASS
SmolLM2-1.7B-Instruct_Q4_K_M_GGUF_8000	6870.19622997982	6390.414378432351	6.98	36.5303	35.82125	-1.94	PASS	PASS

liangliangchang changed the title ~~Lichang.mmq opt~~ RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations Jun 30, 2026

liangliangchang marked this pull request as ready for review June 30, 2026 23:22

liangliangchang requested review from Annieren, jimw567 and mgehre-amd June 30, 2026 23:22

mgehre-amd reviewed Jul 1, 2026

View reviewed changes

mgehre-amd requested a review from roberteg16 July 1, 2026 08:57

liangliangchang force-pushed the lichang.mmq_opt branch from d4a89a2 to b94c48c Compare July 9, 2026 18:58

liangliangchang and others added 5 commits July 9, 2026 13:08

Improve RDNA3.5 MMQ occupancy, LDS layout, and activation pipelining

377f11e

- Add kernel dimension logging for profiling analysis - Reduce mmq_x when LDS limits WGs/CU - Improve LDS layout and activation load pipelining for prefill Co-authored-by: Cursor <cursoragent@cursor.com>

cuda: optimize RDNA3.5 tiny-M MMQ for gate Q8_0 prefill

f2c161e

Co-authored-by: Cursor <cursoragent@cursor.com>

liangliangchang force-pushed the lichang.mmq_opt branch from b94c48c to 7483e49 Compare July 9, 2026 19:09

liangliangchang marked this pull request as draft July 10, 2026 05:03

liangliangchang and others added 6 commits July 10, 2026 09:36

Improve RDNA3.5 MMQ occupancy, LDS layout, and activation pipelining

501ba43

- Add kernel dimension logging for profiling analysis - Reduce mmq_x when LDS limits WGs/CU - Improve LDS layout and activation load pipelining for prefill Co-authored-by: Cursor <cursoragent@cursor.com>

cuda: optimize RDNA3.5 tiny-M MMQ for gate Q8_0 prefill

d590354

Co-authored-by: Cursor <cursoragent@cursor.com>

jimw567 force-pushed the lichang.mmq_opt branch from a907526 to 74503ef Compare July 10, 2026 15:50

liangliangchang added 2 commits July 14, 2026 14:17

Merge branch 'lichang.mmq_opt' of https://github.com/liangliangchang/…

c9bbb21

…llama.cpp into lichang.mmq_opt

Merge latest rocm/gfx11 into lichang.mmq_opt

01637a8

liangliangchang marked this pull request as ready for review July 15, 2026 23:16

liangliangchang and others added 9 commits July 20, 2026 09:57

Merge branch 'gfx11' of rocm/llama.cpp into lichang.mmq_opt

cb01a99

Pull latest gfx11 CI fixes (hip-quality-check UI build off) and roofline/multistream/MoE updates.

test: add RDNA3.5 MMQ prefill kernel cases for PR ROCm#32.

c12b327

Register Q4_K, Q6_K, Q8_0, and Q4_0 MUL_MAT shapes from prefill profiling in test-backend-ops eval and perf for gfx1151 regression against baseline.

test: dedupe RDNA3.5 MMQ prefill cases in test-backend-ops

a2d70ae

Extract add_rdna35_mmq_prefill_cases() so eval and perf suites share one shape list instead of maintaining duplicate emplace_back blocks. Co-authored-by: Cursor <cursoragent@cursor.com>

fix(hip): drop RDNA3.5 activation B-pipeline that regressed MTP accept.

68bbe93

Register-staged prefetch of activation tiles across vec_dot produced wrong tile_y on gfx1151; use global→LDS loads with MMQ_HIP_TILE_BARRIER instead.

fix(hip): restore RDNA3.5 activation B-pipeline without MTP regression.

cf31ddd

Prefetch the next by0 tile only after both vec_dot WMMA phases complete. Keeping y0_next_cache live during the k=MMQ_TILE_NE_K vec_dot corrupted activations on gfx1151 and dropped Q4_0 27B MTP accept to ~52%.

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations#32

RDNA3.5 (gfx11 / gfx1151) MMQ prefill optimizations#32
liangliangchang wants to merge 23 commits into
AMD-Ecosystem:gfx11from
liangliangchang:lichang.mmq_opt

liangliangchang commented Jun 30, 2026 •

edited

Loading

Uh oh!

mgehre-amd Jul 1, 2026

Uh oh!

liangliangchang Jul 22, 2026

Uh oh!

liangliangchang commented Jul 10, 2026

Uh oh!

mgehre-amd commented Jul 10, 2026

Uh oh!

liangliangchang commented Jul 10, 2026

Uh oh!

Annieren commented Jul 10, 2026

Uh oh!

liangliangchang commented Jul 22, 2026

Uh oh!

liangliangchang commented Jul 22, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Uh oh!

Conversation

liangliangchang commented Jun 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

MMQ optimizations (mmq.cuh)

Benchmarks (gfx1151, ROCm)

Files changed

Test plan

AI usage disclosure

Uh oh!

mgehre-amd Jul 1, 2026

Choose a reason for hiding this comment

Uh oh!

liangliangchang Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

liangliangchang commented Jul 10, 2026

Uh oh!

mgehre-amd commented Jul 10, 2026

Uh oh!

liangliangchang commented Jul 10, 2026

Uh oh!

Annieren commented Jul 10, 2026

Uh oh!

liangliangchang commented Jul 22, 2026

Uh oh!

liangliangchang commented Jul 22, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

liangliangchang commented Jun 30, 2026 •

edited

Loading

MMQ optimizations (`mmq.cuh`)