From 1163af96cf6bb4a4b819f998f84c153a49768b99 Mon Sep 17 00:00:00 2001 From: Horacio Vico Date: Wed, 15 Apr 2026 09:08:41 -0300 Subject: [PATCH] cuda: cap host MMQ tile size on Volta to match device kernels (#1638) get_mmq_x_max_host() returned 128 for NVIDIA Volta when built with GGML_CUDA_FORCE_MMQ, but get_mmq_x_max_device() caps at MMQ_DP4A_MAX_BATCH_SIZE (= 64) for the same configuration because the wider mmq_x device kernels are not instantiated on sm_70. The host dispatcher selects an mmq_x in [8, mmq_x_max] and instantiates the matching template; any pick > 64 hits NO_DEVICE_CODE at runtime and the kernel returns without writing output, producing zero-valued activations and incoherent generation. Repro on Tesla V100 (sm_70) with a Q6_K Qwen3.5-MoE model: cmake -B build -DGGML_CUDA=ON -DGGML_CUDA_FORCE_MMQ=ON \ -DGGML_CUDA_NO_PEER_COPY=ON -DCMAKE_CUDA_ARCHITECTURES=70 ./llama-cli -m model.gguf -ngl 99 -p "<~200-token prompt>" # stderr: mmq.cuh:3941: ERROR: CUDA kernel mul_mat_q has no # device code compatible with CUDA arch 700. Cap the host bound at MMQ_DP4A_MAX_BATCH_SIZE on Volta regardless of GGML_CUDA_FORCE_MMQ; this matches the device template availability and yields correct output. Measured on a single V100-SXM2: before patch: PP unusable (NO_DEVICE_CODE), TG ~100 t/s after patch: PP ~630 t/s, TG ~102 t/s The change is a no-op on architectures that have int8 MMA (Turing+), on AMD/HIP, and on builds without GGML_CUDA_FORCE_MMQ. --- ggml/src/ggml-cuda/mmq.cuh | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) 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() {