Public/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-19 09:15:04 +02:00
Code Issues Packages Projects Releases Wiki Activity

Restrict scale-fusion to NVFP4

Browse Source 
This is necessary, as the prolog is quite heavy in GEMV for some
quants/model configs, leading to net perf regression.
We should really be looking to refactor this such that ratio of
prologue/hot-loop/epilogue is better on the hot-loop
front:

+ ./scripts/compare-llama-bench.py -b master -c c1b9381d32 --tool llama-bench -i llama-bench.sqlite
| CPU                         | Model                    | Test         |   t/s master |   t/s c1b9381d3 |   Speedup |
|:----------------------------|:-------------------------|:-------------|-------------:|----------------:|----------:|
| INTEL(R) XEON(R) GOLD 6542Y | gemma4 26B.A4B NVFP4     | tg128@d32768 |       151.70 |          154.32 |      1.02 |
| INTEL(R) XEON(R) GOLD 6542Y | gemma4 26B.A4B Q4_K_M    | tg128@d32768 |       187.95 |          185.73 |      0.99 |
| INTEL(R) XEON(R) GOLD 6542Y | gpt-oss 20B MXFP4 MoE    | tg128@d32768 |       304.62 |          300.69 |      0.99 |
| INTEL(R) XEON(R) GOLD 6542Y | qwen35moe 35B.A3B NVFP4  | tg128@d32768 |       193.72 |          211.99 |      1.09 |
| INTEL(R) XEON(R) GOLD 6542Y | qwen35moe 35B.A3B Q4_K_M | tg128@d32768 |       217.76 |          218.15 |      1.00
This commit is contained in:
Oliver Simons
2026-06-10 16:58:39 +02:00
parent 29713604f1
commit 28bde658e9
2 changed files with 47 additions and 22 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/ggml-cuda.cu
+15 -2
View File
@@ -3863,6 +3863,10 @@ static bool ggml_cuda_can_parse_mm_lane_type(ggml_type type) {
return ggml_is_quantized(type) || type == GGML_TYPE_F32 || type == GGML_TYPE_F16 || type == GGML_TYPE_BF16;
}
static bool ggml_cuda_can_fuse_mm_lane_scale(const ggml_tensor * mm) {
return mm->src[0]->type == GGML_TYPE_NVFP4;
}
static bool ggml_cuda_can_parse_mm_lane_bias(const ggml_tensor * mm, const ggml_tensor * bias) {
if (bias->type != GGML_TYPE_F32 || bias->ne[0] != mm->ne[0]) {
return false;
@@ -3904,7 +3908,7 @@ static bool ggml_cuda_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, g
lane.n_nodes++;
}
if (!ggml_is_quantized(mm->src[0]->type) || i + lane.n_nodes + 3 >= cgraph->n_nodes) {
if (!ggml_is_quantized(mm->src[0]->type) || !ggml_cuda_can_fuse_mm_lane_scale(mm) || i + lane.n_nodes + 3 >= cgraph->n_nodes) {
return true;
}
@@ -3979,7 +3983,8 @@ static bool ggml_cuda_parse_mul_mat_lane(const ggml_cgraph * cgraph, int i, ggml
lane.n_nodes++;
}
if (!ggml_is_quantized(mm->src[0]->type) || i + lane.n_nodes >= cgraph->n_nodes || cgraph->nodes[i + lane.n_nodes]->op != GGML_OP_MUL) {
if (!ggml_is_quantized(mm->src[0]->type) || !ggml_cuda_can_fuse_mm_lane_scale(mm) ||
i + lane.n_nodes >= cgraph->n_nodes || cgraph->nodes[i + lane.n_nodes]->op != GGML_OP_MUL) {
return true;
}
@@ -4094,6 +4099,10 @@ static int ggml_cuda_try_fuse_mm_glu(ggml_backend_cuda_context * cuda_ctx, ggml_
return 0;
}
if ((up->scale != nullptr || gate->scale != nullptr) && !ggml_cuda_can_fuse_mm_lane_scale(up->mm)) {
return 0;
}
const int out_nodes[] = { glu_idx };
const int n_nodes = glu_idx - i + 1;
int external_view_nodes[2];
@@ -4140,6 +4149,10 @@ static int ggml_cuda_try_fuse_mm_scale(ggml_backend_cuda_context * cuda_ctx, ggm
return 0;
}
if (!ggml_cuda_can_fuse_mm_lane_scale(lane.mm)) {
return 0;
}
if (!ggml_cuda_should_fuse_mul_mat_vec_q(lane.mm)) {
return 0;
}
ggml/src/ggml-cuda/mmvq.cu
+32 -20
View File
@@ -529,17 +529,19 @@ static __global__ void mul_mat_vec_q(
ggml_glu_op active_glu;
if constexpr (has_fusion) {
use_gate = fusion.gate != nullptr;
use_bias = fusion.x_bias != nullptr;
use_gate_bias = fusion.gate_bias != nullptr && use_gate;
use_scale = fusion.x_scale != nullptr;
use_gate_scale = fusion.gate_scale != nullptr && use_gate;
vgate = fusion.gate;
x_bias = (const float *) fusion.x_bias;
gate_bias = (const float *) fusion.gate_bias;
x_scale = (const float *) fusion.x_scale;
gate_scale = (const float *) fusion.gate_scale;
active_glu = fusion.glu_op;
use_gate = fusion.gate != nullptr;
use_bias = fusion.x_bias != nullptr;
use_gate_bias = fusion.gate_bias != nullptr && use_gate;
vgate = fusion.gate;
x_bias = (const float *) fusion.x_bias;
gate_bias = (const float *) fusion.gate_bias;
active_glu = fusion.glu_op;
if constexpr (type == GGML_TYPE_NVFP4) {
use_scale = fusion.x_scale != nullptr;
use_gate_scale = fusion.gate_scale != nullptr && use_gate;
x_scale = (const float *) fusion.x_scale;
gate_scale = (const float *) fusion.gate_scale;
}
}
@@ -569,11 +571,13 @@ static __global__ void mul_mat_vec_q(
gate_biases[j] = gate_bias[j * stride_col_dst + threadIdx.x];
}
}
if (use_scale) {
x_scales = x_scale[ids ? channel_x : 0];
}
if (use_gate_scale) {
gate_scales = gate_scale[ids ? channel_x : 0];
if constexpr (type == GGML_TYPE_NVFP4) {
if (use_scale) {
x_scales = x_scale[ids ? channel_x : 0];
}
if (use_gate_scale) {
gate_scales = gate_scale[ids ? channel_x : 0];
}
}
}
@@ -658,16 +662,20 @@ static __global__ void mul_mat_vec_q(
if (use_bias) {
result += x_biases[j];
}
if (use_scale) {
result *= x_scales;
if constexpr (type == GGML_TYPE_NVFP4) {
if (use_scale) {
result *= x_scales;
}
}
if (use_gate) {
float gate_value = tmp_gate[j][threadIdx.x];
if (use_gate_bias) {
gate_value += gate_biases[j];
}
if (use_gate_scale) {
gate_value *= gate_scales;
if constexpr (type == GGML_TYPE_NVFP4) {
if (use_gate_scale) {
gate_value *= gate_scales;
}
}
switch (active_glu) {
case GGML_GLU_OP_SWIGLU:
@@ -693,6 +701,9 @@ static __global__ void mul_mat_vec_q(
if constexpr (!has_fusion) {
GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, use_scale, use_gate_scale, active_glu, gate_bias, x_bias, x_scale, gate_scale, tmp_gate);
}
if constexpr (type != GGML_TYPE_NVFP4) {
GGML_UNUSED_VARS(use_scale, use_gate_scale, x_scale, gate_scale, x_scales, gate_scales);
}
}
// Dedicated MoE multi-token kernel.
@@ -1174,6 +1185,7 @@ void ggml_cuda_mul_mat_vec_q(
if (fusion) {
GGML_ASSERT( !ids || dst->ne[2] == 1);
GGML_ASSERT( ids || dst->ne[1] == 1);
GGML_ASSERT((fusion->x_scale == nullptr && fusion->gate_scale == nullptr) || src0->type == GGML_TYPE_NVFP4);
if (fusion->x_bias) {
GGML_ASSERT(fusion->x_bias->type == GGML_TYPE_F32);