From 88ef64264c243d06fadcaf8338f0e3a318854698 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Mon, 17 Aug 2026 03:08:38 -0600 Subject: [PATCH] cuda/hip: optional wave64 for the quantized mat-vec kernels on RDNA Adds GGML_HIP_MMVQ_WAVE64, off by default, which builds mmvq.cu with -mwavefrontsize64. Exploratory: it is a clear win on a couple of shapes and a clear loss on more of them, so it is exposed as a flag rather than turned on, and the data is recorded here so the next person does not have to rediscover it. Not applied target-wide because fattn-mma-f16.cuh static_asserts on wave32 tiling. -mwavefrontsize64 is per translation unit, so this flips all 265 kernels in mmvq.cu, not only the q4_K one. A function attribute cannot narrow it: clang refuses to inline any wave32 function into a target("wavefrontsize64") one, down to the threadIdx accessors, and every kernel here is built from __forceinline__ helpers. Narrowing it properly needs the q4_K instantiation moved to its own translation unit. q4_K, wave32 vs wave64, rocprofv3 kernel time, median of 3 passes, gfx1151: m=21504 n=1 k=5376 289.6 -> 285.6 us -1.4% m=4096 n=4 k=14336 117.6 -> 110.0 us -6.5% m=4096 n=8 k=14336 210.0 -> 223.1 us +6.2% m=4096 n=4 k=4096 35.1 -> 40.5 us +15.4% m=4096 n=1 k=4096 32.3 -> 37.5 us +16.2% m=4096 n=1 k=14336 54.1 -> 69.6 us +28.7% Only the gemma-4-31B-it fused FFN gate/up shape and one batched shape benefit. The k=4096 and k=14336 n=1 shapes lose badly. The likely mechanism is the K-loop trip count: at nwarps=2 and wave64, blocks_per_iter doubles, so k=4096 gives two iterations instead of four and the 4x unroll has nothing left to pipeline. Across all 21 mmvq types at m=4096 n=1 k=4096, median of 3: 5 improve by more than 2% (iq2_xs -12.8, iq2_xxs -12.5, iq3_xxs -12.5, iq2_s -11.2, iq3_s -5.7), 9 regress by more than 2% (q4_K +18.7, iq4_xs +15.5, mxfp4 +9.8, nvfp4 +9.1, q2_K +7.7, q4_0 +5.4, iq1_s +4.7, iq4_nl +4.4, q5_0 +2.4), 7 are within +/-2%. In-model impact on gfx1151 is narrower than that table suggests: with GGML_CUDA_DQ_MMV on by default, q4_K/q5_K/q6_K at ne11==1 go to mul_mat_vec_dq_*, which this does not touch. The only mmvq user left for gemma-4-31B-it Q4_K_M is the fused GEGLU gate/up op, which is the 21504x5376 shape above. Models built from non-K quants route every n=1 matvec through mmvq and would take the regressions. Verification: test-backend-ops test -o MUL_MAT,MUL_MAT_ID,MUL_MAT_VEC_FUSION,MUL_MAT_ID_FUSION passes with the flag on, with and without GGML_CUDA_DQ_MMV=0. Note the flag miscompiles on ROCm 7.12.0a20260211 (clang 22): 371 of 1175 MUL_MAT cases fail, including the benchmarked shape at ERR 0.47. 7.15.0a20260728 is clean. There is no version guard; that is a reason to keep this opt-in. Assisted-by: Claude Opus 5 --- ggml/src/ggml-cuda/common.cuh | 9 ++++++++- ggml/src/ggml-cuda/mmvq.cu | 5 ++++- ggml/src/ggml-hip/CMakeLists.txt | 9 +++++++++ 3 files changed, 21 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index ceb01ad342c3..a0df80ae5003 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -381,8 +381,15 @@ struct block_q8_1_x4 { }; static_assert(sizeof(block_q8_1_x4) == 4*sizeof(block_q8_1), "block_q8_1_x4 must alias 4 q8_1 blocks"); +// Set by ggml-hip/CMakeLists.txt on the translation units built with -mwavefrontsize64. +#ifndef GGML_CUDA_FORCE_WAVE64 +#define GGML_CUDA_FORCE_WAVE64 0 +#endif + static constexpr __device__ int ggml_cuda_get_physical_warp_size() { -#if defined(GGML_USE_HIP) && (defined(__GFX9__) || defined(__GFX8__)) +#if GGML_CUDA_FORCE_WAVE64 + return 64; +#elif defined(GGML_USE_HIP) && (defined(__GFX9__) || defined(__GFX8__)) return 64; #else return 32; diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 5f425871da4e..2c7597e2a6f1 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -898,7 +898,10 @@ static void mul_mat_vec_q_switch_ncols_dst( const int device = ggml_cuda_get_device(); const int cc = ggml_cuda_info().devices[device].cc; - const int warp_size = ggml_cuda_info().devices[device].warp_size; + // This TU is compiled -mwavefrontsize64 when GGML_HIP_MMVQ_WAVE64 is set, in which case its + // kernels have 64 lanes per wave regardless of what the device properties report (32 on + // RDNA), and the launch geometry has to match. + const int warp_size = GGML_CUDA_FORCE_WAVE64 ? 64 : ggml_cuda_info().devices[device].warp_size; const mmvq_parameter_table_id table_id = get_device_table_id(cc); const bool has_ids = ids != nullptr; diff --git a/ggml/src/ggml-hip/CMakeLists.txt b/ggml/src/ggml-hip/CMakeLists.txt index efb3a7984922..8e9a46c9cfb0 100644 --- a/ggml/src/ggml-hip/CMakeLists.txt +++ b/ggml/src/ggml-hip/CMakeLists.txt @@ -61,6 +61,15 @@ file(GLOB GGML_HEADERS_ROCM "../ggml-cuda/*.cuh") list(APPEND GGML_HEADERS_ROCM "../../include/ggml-cuda.h") file(GLOB GGML_SOURCES_ROCM "../ggml-cuda/*.cu") + +# Opt-in: build the quantized mat-vec kernels for wave64. RDNA selects the wave size per kernel, +# but -mwavefrontsize64 is per translation unit, so this flips every kernel in mmvq.cu, not just +# the q4_K one. Not applied target-wide: fattn-mma-f16.cuh static_asserts on wave32 tiling. +if (GGML_HIP_MMVQ_WAVE64) + set_source_files_properties("../ggml-cuda/mmvq.cu" PROPERTIES + COMPILE_OPTIONS "-mwavefrontsize64;-DGGML_CUDA_FORCE_WAVE64=1") +endif() + file(GLOB SRCS "../ggml-cuda/template-instances/fattn-tile*.cu") list(APPEND GGML_SOURCES_ROCM ${SRCS}) file(GLOB SRCS "../ggml-cuda/template-instances/fattn-mma*.cu")