Compare commits

...

2 Commits

Author SHA1 Message Date
Oliver Simons 683f0c72e5 Only index by compile times + always multiply/add (#25445)
The first one avoids relying on compile to optimize local memory away,
and the second is cheaper than issuing control flow statements
2026-07-09 13:23:57 +02:00
Adrien Gallouët 259f2e2a53 llama-bench : init params.offline (#25476)
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
2026-07-09 11:56:56 +02:00
2 changed files with 24 additions and 32 deletions
+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 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;
}
}
+1
View File
@@ -520,6 +520,7 @@ static cmd_params parse_cmd_params(int argc, char ** argv) {
params.delay = cmd_params_defaults.delay;
params.progress = cmd_params_defaults.progress;
params.no_warmup = cmd_params_defaults.no_warmup;
params.offline = cmd_params_defaults.offline;
if (const char * env = getenv("HF_TOKEN")) {
params.hf_token = env;