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")