From 5d4d37182570a5c5ab24e7fc53f7de3303792951 Mon Sep 17 00:00:00 2001 From: lichang Date: Fri, 18 Sep 2026 16:32:19 -0600 Subject: [PATCH] HIP: use upstream rows_per_warp() on RDNA3.5 The RDNA3.5 override returns 32 for J>=64 && J%32==0, with carve-outs holding q4_0/q8_0/q4_K/q5_K at 16. rows_per_warp/16 is ntx, so 32 selects the ntx=2 schedule, and for every type without a carve-out that schedule is several times slower on gfx1151. test-backend-ops perf on gfx1151, m in {1024,4096}, k=4096, percent faster with upstream's unconditional 16: type n=48 n=64 n=96 n=128 iq4_nl +0.0 +553.1 +441.8 +2.0 iq4_xs -0.0 +543.1 +484.2 +2.4 mxfp4 -0.0 +539.6 +437.1 +2.6 q5_0 +0.1 +448.6 +418.1 +6.4 iq2_xxs -0.0 +414.9 +314.5 +1.6 iq3_s +0.1 +409.5 +291.6 +1.9 q5_1 -0.0 +15.5 +18.2 +8.3 q4_1 -0.0 +3.1 +0.6 +22.7 q6_K -0.2 +3.5 +0.0 +2.0 q3_K +0.1 +2.1 +0.1 -0.3 q4_K +0.0 +0.4 -0.2 +1.3 q5_K -0.0 +0.8 +0.8 +0.2 q2_K -0.2 -0.1 -0.1 +0.1 n=48 is a control, since both variants return 16 there. q4_0 and q8_0 are a second control, already 16 at every width. All four control columns stay within +-0.2%. Every effect above 5% reproduces on a second board with the same sign and within 2x, at a per-shape noise floor of 0.06-0.74%. The stock perf list jumps from n=32 to n=128, so J=64 and J=96 are not covered by it; these numbers come from an added n in {48,64,96,128} grid. Q4_K_M models are unaffected end to end (+-0.5%), because q4_K/q5_K/q6_K are exactly the types whose deltas are under 3.5%; the large wins are in the iq*, mxfp4, q5_0 and q5_1 families. One shape prefers 32, q4_K m=4096 n=96 at -1.1%, against 35 shapes where 16 wins by more than 5%, so keeping a per-type rule is not worth re-creating the hazard. Correctness: 1510/1510 MUL_MAT on both arms. Assisted-by: Claude Opus 5 Co-authored-by: Cursor --- ggml/src/ggml-cuda/mmq.cuh | 14 -------------- 1 file changed, 14 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 94365349df21..177008241194 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -183,21 +183,7 @@ struct ggml_cuda_mmq_config { constexpr __device__ int rows_per_warp() const { #if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) -#if defined(RDNA3_5) - if ((type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q8_0 || - type == GGML_TYPE_Q5_K) && J == 128) { - return 16; - } - // Block quants only have a batched WMMA vec_dot at J=128; at the other two widths - // that would take ntx=2 (J=64 and J=96, the only J>=64 multiples of 32) they fall - // back to the generic schedule, which is several times slower there. Keep ntx=1. - if ((type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q8_0) && (J == 64 || J == 96)) { - return 16; - } - return J >= 64 && J % 32 == 0 ? 32 : 16; -#else return 16; -#endif // defined(RDNA3_5) #else return J >= 48 && J % 16 == 0 ? 32 : 16; #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)