Lumpiasty/llama.cpp
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Only index by compile times + always multiply/add (#25445)

Browse Source 
The first one avoids relying on compile to optimize local memory away,
and the second is cheaper than issuing control flow statements
This commit is contained in:
Oliver Simons
2026-07-09 13:23:57 +02:00
committed by GitHub
parent 259f2e2a53
commit 683f0c72e5
1 changed files with 23 additions and 32 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/mmvq.cu
+23 -32
View File
@@ -549,8 +549,8 @@ static __global__ void mul_mat_vec_q(
[[maybe_unused]] float x_biases[ncols_dst] = { 0.0f }; [[maybe_unused]] float x_biases[ncols_dst] = { 0.0f };
[[maybe_unused]] float gate_biases[ncols_dst] = { 0.0f }; [[maybe_unused]] float gate_biases[ncols_dst] = { 0.0f };
[[maybe_unused]] float x_scales; [[maybe_unused]] float x_scales = 1.0f;
[[maybe_unused]] float gate_scales; [[maybe_unused]] float gate_scales = 1.0f;
if constexpr (has_fusion) { if constexpr (has_fusion) {
// 1. Hide latency by prefetching bias, gates and scales here // 1. Hide latency by prefetching bias, gates and scales here
// 2. load only on threads that won't die after partial sum calculation // 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<warp_size>(tmp_gate[j][i]); tmp_gate[j][i] = warp_reduce_sum<warp_size>(tmp_gate[j][i]);
} }
} }
}
if (threadIdx.x < rows_per_cuda_block && (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) { if (threadIdx.x == i && (rows_per_cuda_block == 1 || uint32_t(row0 + i) < stride_col_dst)) {
float result = tmp[j][threadIdx.x]; float result = tmp[j][i];
if constexpr (has_fusion) { if constexpr (has_fusion) {
if constexpr (type == GGML_TYPE_NVFP4) { if constexpr (type == GGML_TYPE_NVFP4) {
if (use_scale) {
result *= x_scales; result *= x_scales;
} }
}
if (use_bias) {
result += x_biases[j]; result += x_biases[j];
} if (use_gate) {
if (use_gate) { float gate_value = tmp_gate[j][i];
float gate_value = tmp_gate[j][threadIdx.x]; if constexpr (type == GGML_TYPE_NVFP4) {
if constexpr (type == GGML_TYPE_NVFP4) {
if (use_gate_scale) {
gate_value *= gate_scales; gate_value *= gate_scales;
} }
}
if (use_gate_bias) {
gate_value += gate_biases[j]; gate_value += gate_biases[j];
} switch (active_glu) {
switch (active_glu) { case GGML_GLU_OP_SWIGLU:
case GGML_GLU_OP_SWIGLU: result *= ggml_cuda_op_silu_single(gate_value);
result *= ggml_cuda_op_silu_single(gate_value); break;
break; case GGML_GLU_OP_GEGLU:
case GGML_GLU_OP_GEGLU: result *= ggml_cuda_op_gelu_single(gate_value);
result *= ggml_cuda_op_gelu_single(gate_value); break;
break; case GGML_GLU_OP_SWIGLU_OAI:
case GGML_GLU_OP_SWIGLU_OAI: { result = ggml_cuda_op_swiglu_oai_single(gate_value, result);
result = ggml_cuda_op_swiglu_oai_single(gate_value, result); break;
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;
} }
} }