From 81139f91fe63f83ea9870429fb35f87495f8678f Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Mon, 27 Oct 2025 21:58:30 +0800 Subject: [PATCH] CUDA: add unused vars to mmvf and mmvq Add GGML_UNSUED_VARS when not taking the fusion path --- ggml/src/ggml-cuda/mmvf.cu | 4 ++++ ggml/src/ggml-cuda/mmvq.cu | 4 ++++ 2 files changed, 8 insertions(+) diff --git a/ggml/src/ggml-cuda/mmvf.cu b/ggml/src/ggml-cuda/mmvf.cu index c2c31cdaf231b..4e31783436d80 100644 --- a/ggml/src/ggml-cuda/mmvf.cu +++ b/ggml/src/ggml-cuda/mmvf.cu @@ -343,6 +343,10 @@ static __global__ void mul_mat_vec_f( } dst[tid*stride_col_dst + row] = value; + + if constexpr (!has_fusion) { + GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, glu_op, gate_x, x_bias, gate_bias, sumf_gate); + } } template diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 7a783e4fcf9b4..be04a85cc5515 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -310,6 +310,10 @@ static __global__ void mul_mat_vec_q( dst[j*stride_col_dst + threadIdx.x] = result; } } + + if constexpr (!has_fusion) { + GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, active_glu, gate_bias, x_bias, tmp_gate); + } } static std::pair calc_launch_params(