From aa538005fc5909f986d827fa45aa2ff3986961bf Mon Sep 17 00:00:00 2001 From: "Theera K." Date: Thu, 1 Oct 2026 20:21:26 +0700 Subject: [PATCH] cuda : route sm70 to the Turing MMVQ nwarps table (llama/29753) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * cuda : route sm70 to the Turing MMVQ nwarps table Volta (sm_70) has no MMVQ parameter table of its own and falls through to GENERIC, which launches K-quant batch-1 decode (ncols_dst == 1) at nwarps=4. sm_70 shares TURING's tuning: the K-quant vec_dot prefers nwarps=2 there. Route sm_70 to the existing MMVQ_PARAMETERS_TURING table in both the device and the host table selector. Measured on one Tesla V100 32GB PCIe (PG500-216, driver 580.178.04, CUDA 12.0.140) with Qwen3.8-27B Q4_K_M, tg128, interleaved A/B in 6 ABBA blocks with paired per-block deltas: +1.091 t/s = +3.17 % (t = +49.0, all six per-block deltas positive); perplexity bit-identical (6.3697 +/- 0.04066 both builds, wiki.test.raw). The patched build's K-quant mul_mat_vec_q kernels launch at nwarps=2 (cubin EIATTR_MAX_THREADS) while Q4_0/Q8_0 stay at nwarps=4, and the same measurement on the September master base gave +3.84 % (t = 85). The tuning originates from the V100-focused fork anyei/llamacpp-v100 (MIT), commit b912d1b1e, which carries a dedicated MMVQ_PARAMETERS_VOLTA table; a cubin-level comparison confirmed that routing sm_70 to the existing TURING table is equivalent for the K-quant batch-1 path this change affects, so this is the minimal 2-line form. https://github.com/anyei/llamacpp-v100/commit/b912d1b1e Original-patch-by: anyei * Update ggml/src/ggml-cuda/mmvq.cu --------- Co-authored-by: tkittich Co-authored-by: Johannes Gäßler --- ggml/src/ggml-cuda/mmvq.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index dcf484be0..a35dc9370 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -112,7 +112,7 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() { return MMVQ_PARAMETERS_RDNA2; #elif defined(GCN) || defined(CDNA) return MMVQ_PARAMETERS_GCN; -#elif __CUDA_ARCH__ >= GGML_CUDA_CC_TURING && __CUDA_ARCH__ < GGML_CUDA_CC_AMPERE +#elif __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA && __CUDA_ARCH__ < GGML_CUDA_CC_AMPERE return MMVQ_PARAMETERS_TURING; #elif __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK return MMVQ_PARAMETERS_GB10; @@ -134,7 +134,7 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) { if (GGML_CUDA_CC_IS_GCN(cc) || GGML_CUDA_CC_IS_CDNA(cc)) { return MMVQ_PARAMETERS_GCN; } - if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_TURING && ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_AMPERE) { + if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA && ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_AMPERE) { return MMVQ_PARAMETERS_TURING; } if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) == GGML_CUDA_CC_DGX_SPARK) {