From 4cc59ab81ea4c06c7404b7405e47abae99dd2c19 Mon Sep 17 00:00:00 2001 From: Brian <288398250+bri-prism@users.noreply.github.com> Date: Wed, 28 Oct 2026 20:12:47 +0800 Subject: [PATCH 1/2] cuda: extend the tall mmvq geometry to Q8_0/Q4_0/Q2_0/Q4_K/Q6_K rows>=4096 on Hopper (GGML_CUDA_MMVQ_TALL_EXT) Extend the 8-rows-per-CTA one-column mat-vec geometry, previously used for PQ2_0 only, to Q8_0, Q4_0, Q2_0, Q4_K and Q6_K when the matrix has at least 4096 rows (a multiple of 8) on Hopper. The K partition and reduction order per row are unchanged, so every output row is bit-identical to the default geometry. GGML_CUDA_MMVQ_TALL_EXT=0 restores the PQ2_0-only behavior (default: enabled). --- ggml/src/ggml-cuda/mmvq.cu | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 4a1712d17193..712c4030c851 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -582,6 +582,12 @@ static constexpr __host__ __device__ int calc_rows_per_block(ggml_type type, int // rows of the weight matrix per CTA for a tall one-column mat-vec (see should_use_tall below) #define MMVQ_TALL_ROWS_PER_BLOCK 8 +// Types besides PQ2_0 that also take the tall geometry (one instantiation more each): the common quantized types. +static constexpr bool mmvq_tall_ext_type(ggml_type type) { + return type == GGML_TYPE_Q8_0 || type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q2_0 || + type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q6_K; +} + template __launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id(), small_k, halve_iters)*ggml_cuda_get_physical_warp_size(), 1) static __global__ void mul_mat_vec_q( @@ -1343,8 +1349,11 @@ static void mul_mat_vec_q_switch_ncols_dst( // (Bonsai 2 9B LM head, 248320 x 4096: 249 us, about 1.1 TB/s on an H200). Eight rows per CTA keep the same K // partition and reduction order for every row (each row's sum is bit-identical) and cut that to 121 us; the // 4096-12288 row trunk projections gain 28-42% as well. Matrices under 4096 rows have too few CTAs for it. + // The same geometry helps any type whose one-column default gives one 128-thread CTA per row (tall output heads). + // GGML_CUDA_MMVQ_TALL_EXT=0 keeps it PQ2_0-only. const auto should_use_tall = [&] { - return type == GGML_TYPE_PQ2_0 && table_id == MMVQ_PARAMETERS_GENERIC && GGML_CUDA_CC_IS_NVIDIA(cc) && + static const bool ext = !getenv("GGML_CUDA_MMVQ_TALL_EXT") || atoi(getenv("GGML_CUDA_MMVQ_TALL_EXT")) != 0; + return (type == GGML_TYPE_PQ2_0 || (ext && mmvq_tall_ext_type(type))) && table_id == MMVQ_PARAMETERS_GENERIC && GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_HOPPER && cc < GGML_CUDA_CC_BLACKWELL && !has_ids && !y_soa && nrows_x >= 4096 && nrows_x % MMVQ_TALL_ROWS_PER_BLOCK == 0; // whole CTAs: rows past the end would be read, not only skipped }; @@ -1388,7 +1397,7 @@ static void mul_mat_vec_q_switch_ncols_dst( launch(std::true_type{}, std::false_type{}, std::false_type{}); } else if (should_halve_iters()) { launch(std::false_type{}, std::true_type{}, std::false_type{}); - } else if constexpr (type == GGML_TYPE_PQ2_0) { + } else if constexpr (type == GGML_TYPE_PQ2_0 || mmvq_tall_ext_type(type)) { if (should_use_tall()) { launch(std::false_type{}, std::false_type{}, std::true_type{}); } else { From 17c1fa97266aeb806c921def9df5275f44a74745 Mon Sep 17 00:00:00 2001 From: Brian <288398250+bri-prism@users.noreply.github.com> Date: Wed, 28 Oct 2026 15:15:07 -0700 Subject: [PATCH 2/2] model: allow prism.hadamard folded weights for Gemma 4 Every folded Gemma 4 weight (attention q/k/v/o, shared FFN, routed experts) goes through build_lora_mm / build_lora_mm_id, which apply the activation transform. --- src/llama-model.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 3bf8245e406c..8f6516be0227 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -1287,6 +1287,7 @@ void llama_model_base::load_hparams(llama_model_loader & ml) { case LLM_ARCH_QWEN35MOE: case LLM_ARCH_QWEN3NEXT: case LLM_ARCH_DSPARK: + case LLM_ARCH_GEMMA4: break; default: throw std::runtime_error(format(