diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index a48cc48b2..e18ada537 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -549,8 +549,8 @@ static __global__ void mul_mat_vec_q( [[maybe_unused]] float x_biases[ncols_dst] = { 0.0f }; [[maybe_unused]] float gate_biases[ncols_dst] = { 0.0f }; - [[maybe_unused]] float x_scales; - [[maybe_unused]] float gate_scales; + [[maybe_unused]] float x_scales = 1.0f; + [[maybe_unused]] float gate_scales = 1.0f; if constexpr (has_fusion) { // 1. Hide latency by prefetching bias, gates and scales here // 2. load only on threads that won't die after partial sum calculation @@ -655,47 +655,38 @@ static __global__ void mul_mat_vec_q( tmp_gate[j][i] = warp_reduce_sum(tmp_gate[j][i]); } } - } - if (threadIdx.x < rows_per_cuda_block && (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) { - float result = tmp[j][threadIdx.x]; - if constexpr (has_fusion) { - if constexpr (type == GGML_TYPE_NVFP4) { - if (use_scale) { + if (threadIdx.x == i && (rows_per_cuda_block == 1 || uint32_t(row0 + i) < stride_col_dst)) { + float result = tmp[j][i]; + if constexpr (has_fusion) { + if constexpr (type == GGML_TYPE_NVFP4) { result *= x_scales; } - } - if (use_bias) { result += x_biases[j]; - } - if (use_gate) { - float gate_value = tmp_gate[j][threadIdx.x]; - if constexpr (type == GGML_TYPE_NVFP4) { - if (use_gate_scale) { + if (use_gate) { + float gate_value = tmp_gate[j][i]; + if constexpr (type == GGML_TYPE_NVFP4) { gate_value *= gate_scales; } - } - if (use_gate_bias) { gate_value += gate_biases[j]; - } - switch (active_glu) { - case GGML_GLU_OP_SWIGLU: - result *= ggml_cuda_op_silu_single(gate_value); - break; - case GGML_GLU_OP_GEGLU: - result *= ggml_cuda_op_gelu_single(gate_value); - break; - case GGML_GLU_OP_SWIGLU_OAI: { - result = ggml_cuda_op_swiglu_oai_single(gate_value, result); - break; + switch (active_glu) { + case GGML_GLU_OP_SWIGLU: + result *= ggml_cuda_op_silu_single(gate_value); + break; + case GGML_GLU_OP_GEGLU: + result *= ggml_cuda_op_gelu_single(gate_value); + break; + case GGML_GLU_OP_SWIGLU_OAI: + result = ggml_cuda_op_swiglu_oai_single(gate_value, result); + break; + default: + result = result * gate_value; + break; } - default: - result = result * gate_value; - break; } } + dst[j*stride_col_dst + i] = result; } - dst[j*stride_col_dst + threadIdx.x] = result; } }