upstream-archive/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-07-23 11:10:57 -05:00
Code Issues Packages Projects Releases Wiki Activity

Only index by compile times + always multiply/add (llama/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 Georgi Gerganov
parent 9cd788e703
commit a2c50280d8
1 changed files with 23 additions and 32 deletions
Download Patch File Download Diff File Expand all files Collapse all files

55
ggml/src/ggml-cuda/mmvq.cu
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 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<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)) {
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;
}
}