From 28bde658e9a3ab12fad983efe2956f5a38703799 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Wed, 10 Jun 2026 16:58:39 +0200 Subject: [PATCH] Restrict scale-fusion to NVFP4 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 c1b9381d327e063cc846b46b59708444b66dc4d8 --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 --- ggml/src/ggml-cuda/ggml-cuda.cu | 17 +++++++++-- ggml/src/ggml-cuda/mmvq.cu | 52 ++++++++++++++++++++------------- 2 files changed, 47 insertions(+), 22 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 6c14c462f3..6d77f64d99 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -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; } diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 9cb4c313f6..9c07ada06d 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -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);