diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
index d0c8233e..ee10d9ed 100644
--- a/ggml/src/ggml-cuda/mmq.cuh
+++ b/ggml/src/ggml-cuda/mmq.cuh
@@ -121,12 +121,15 @@ struct tile_x_sizes {
 };
 
 static constexpr int get_mmq_x_max_host(const int cc) {
+    // On NVIDIA Volta the device-side MMQ kernels are only instantiated up to
+    // MMQ_DP4A_MAX_BATCH_SIZE when GGML_CUDA_FORCE_MMQ is defined (see
+    // get_mmq_x_max_device() below). If the host dispatch picks a larger tile,
+    // the kernel hits NO_DEVICE_CODE at runtime. Without GGML_CUDA_FORCE_MMQ
+    // the host already capped at MMQ_DP4A_MAX_BATCH_SIZE so cuBLAS handles
+    // larger batches. Either way the bound on Volta is the same, so collapse
+    // the two paths and document the constraint.
     return int8_mma_available(cc) ? 128 :
-#ifdef GGML_CUDA_FORCE_MMQ
-        cc >= CC_VOLTA && cc < CC_OFFSET_AMD ? 128                     : 64;
-#else
         cc >= CC_VOLTA && cc < CC_OFFSET_AMD ? MMQ_DP4A_MAX_BATCH_SIZE : 64;
-#endif // GGML_CUDA_FORCE_MMQ
 }
 
 static constexpr __device__ int get_mmq_x_max_device() {