From 45d47042676435f5fdda358252432a42ebdba446 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 29 May 2026 13:58:32 +0200 Subject: [PATCH] Add dense MMVQ fusion as well Perf numbers on B4500. Note qwen35 is FP8->Q8 + ./scripts/compare-llama-bench.py -b master -c osimons/nvfp4_fuse_mmvq --tool llama-bench -i llama-bench.sqlite | Model | Test | t/s master | t/s osimons/nvfp4_fuse_mmvq | Speedup | |:-------------------------|:-------------|-------------:|------------------------------:|----------:| | qwen35moe 35B.A3B NVFP4 | tg128@d32768 | 150.15 | 156.29 | 1.04 | | qwen35moe 35B.A3B Q4_K_M | tg128@d32768 | 157.91 | 157.64 | 1.00 | Perf numbers on DGX Spark + ./scripts/compare-llama-bench.py -b master -c osimons/nvfp4_fuse_mmvq --tool llama-bench -i llama-bench.sqlite | Model | Test | t/s master | t/s osimons/nvfp4_fuse_mmvq | Speedup | |:-------------------------|:-------------|-------------:|------------------------------:|----------:| | qwen35moe 35B.A3B NVFP4 | tg128@d32768 | 58.31 | 59.69 | 1.02 | | qwen35moe 35B.A3B Q4_K_M | tg128@d32768 | 54.94 | 54.79 | 1.00 | --- ggml/src/ggml-cuda/ggml-cuda.cu | 176 ++++++++++++++++++++++++++++++++ ggml/src/ggml-cuda/mmvq.cu | 10 +- 2 files changed, 180 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 620e71cdef..187c689458 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3950,6 +3950,61 @@ static bool ggml_cuda_parse_nvfp4_mmid_lane(const ggml_cgraph * cgraph, int i, g return true; } +static bool ggml_cuda_parse_nvfp4_mm_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mmid_lane & lane) { + if (i >= cgraph->n_nodes || cgraph->nodes[i]->op != GGML_OP_MUL_MAT) { + return false; + } + + ggml_tensor * mm = cgraph->nodes[i]; + if (mm->src[0]->type != GGML_TYPE_NVFP4 || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32) { + return false; + } + + lane = {}; + lane.mm = mm; + lane.out = mm; + lane.n_nodes = 1; + + if (i + lane.n_nodes < cgraph->n_nodes && cgraph->nodes[i + lane.n_nodes]->op == GGML_OP_ADD) { + ggml_tensor * add = cgraph->nodes[i + lane.n_nodes]; + if (add->src[0] == lane.out) { + lane.bias = add->src[1]; + } else if (add->src[1] == lane.out) { + lane.bias = add->src[0]; + } else { + return false; + } + lane.bias_node = add; + lane.out = add; + lane.n_nodes++; + } + + if (i + lane.n_nodes >= cgraph->n_nodes || cgraph->nodes[i + lane.n_nodes]->op != GGML_OP_MUL) { + return true; + } + + ggml_tensor * mul = cgraph->nodes[i + lane.n_nodes]; + const bool mul_lhs_out = mul->src[0] == lane.out; + const bool mul_rhs_out = mul->src[1] == lane.out; + if (!mul_lhs_out && !mul_rhs_out) { + return true; + } + + const ggml_tensor * scale = mul_lhs_out ? mul->src[1] : mul->src[0]; + if (scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != 1) { + return false; + } + + if (mul->type != GGML_TYPE_F32 || !ggml_are_same_shape(mul, lane.out)) { + return false; + } + + lane.scale = scale; + lane.out = mul; + lane.n_nodes++; + return true; +} + static bool ggml_cuda_can_fuse_mmid_scale_subgraph(const ggml_cgraph * cgraph, int start_idx, int count, const int * outputs, int num_outputs) { for (int j = 0; j < count; ++j) { const int idx = start_idx + j; @@ -4019,6 +4074,34 @@ static bool ggml_cuda_should_fuse_mmid_lanes(const ggml_cuda_mmid_lane & up, con return !split; } +static bool ggml_cuda_should_fuse_mm_lanes(const ggml_cuda_mmid_lane & up, const ggml_cuda_mmid_lane & gate, const ggml_tensor * glu) { + if (up.mm->src[0]->type != gate.mm->src[0]->type || !ggml_are_same_shape(up.mm->src[0], gate.mm->src[0]) || + !ggml_are_same_stride(up.mm->src[0], gate.mm->src[0])) { + return false; + } + + if (up.mm->src[1] != gate.mm->src[1]) { + return false; + } + + if (glu->op != GGML_OP_GLU || glu->src[0] != gate.out || glu->src[1] != up.out) { + return false; + } + + static constexpr std::array valid_glu_ops = { GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU, GGML_GLU_OP_SWIGLU_OAI }; + if (std::find(valid_glu_ops.begin(), valid_glu_ops.end(), ggml_get_glu_op(glu)) == valid_glu_ops.end()) { + return false; + } + + if (const bool swapped = ggml_get_op_params_i32(glu, 1); swapped) { + return false; + } + + const bool split = ggml_backend_buft_is_cuda_split(up.mm->src[0]->buffer->buft) || + ggml_backend_buft_is_cuda_split(gate.mm->src[0]->buffer->buft); + return !split; +} + static int ggml_cuda_try_fuse_nvfp4_mmid_scale_glu(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { ggml_cuda_mmid_lane lane0; if (!ggml_cuda_parse_nvfp4_mmid_lane(cgraph, i, lane0)) { @@ -4083,6 +4166,66 @@ static int ggml_cuda_try_fuse_nvfp4_mmid_scale_glu(ggml_backend_cuda_context * c return glu_idx - i; } +static int ggml_cuda_try_fuse_nvfp4_mm_scale_glu(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { + ggml_cuda_mmid_lane lane0; + if (!ggml_cuda_parse_nvfp4_mm_lane(cgraph, i, lane0)) { + return 0; + } + + ggml_cuda_mmid_lane lane1; + if (!ggml_cuda_parse_nvfp4_mm_lane(cgraph, i + lane0.n_nodes, lane1)) { + return 0; + } + + const int glu_idx = i + lane0.n_nodes + lane1.n_nodes; + if (glu_idx >= cgraph->n_nodes || cgraph->nodes[glu_idx]->op != GGML_OP_GLU) { + return 0; + } + + if (lane0.scale == nullptr && lane1.scale == nullptr) { + return 0; + } + + const ggml_tensor * glu = cgraph->nodes[glu_idx]; + ggml_cuda_mmid_lane * gate = nullptr; + ggml_cuda_mmid_lane * up = nullptr; + if (glu->src[0] == lane0.out && glu->src[1] == lane1.out) { + gate = &lane0; + up = &lane1; + } else if (glu->src[0] == lane1.out && glu->src[1] == lane0.out) { + gate = &lane1; + up = &lane0; + } else { + return 0; + } + + if (!ggml_cuda_should_fuse_mm_lanes(*up, *gate, glu)) { + return 0; + } + + const int out_nodes[] = { glu_idx }; + const int n_nodes = glu_idx - i + 1; + if (!ggml_cuda_can_fuse_mmid_scale_subgraph(cgraph, i, n_nodes, out_nodes, 1) || + !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_nodes, out_nodes, 1)) { + return 0; + } + + if (!ggml_cuda_should_fuse_mul_mat_vec_q(up->mm)) { + return 0; + } + + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.gate = gate->mm->src[0]; + fusion_data.x_bias = up->bias; + fusion_data.gate_bias = gate->bias; + fusion_data.x_scale = up->scale; + fusion_data.gate_scale = gate->scale; + fusion_data.glu_op = ggml_get_glu_op(glu); + + ggml_cuda_mul_mat_vec_q(*cuda_ctx, up->mm->src[0], up->mm->src[1], nullptr, cgraph->nodes[glu_idx], &fusion_data); + return glu_idx - i; +} + static int ggml_cuda_try_fuse_nvfp4_mmid_scale(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { ggml_cuda_mmid_lane lane; if (!ggml_cuda_parse_nvfp4_mmid_lane(cgraph, i, lane) || lane.scale == nullptr) { @@ -4113,6 +4256,31 @@ static int ggml_cuda_try_fuse_nvfp4_mmid_scale(ggml_backend_cuda_context * cuda_ return lane.n_nodes - 1; } +static int ggml_cuda_try_fuse_nvfp4_mm_scale(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { + ggml_cuda_mmid_lane lane; + if (!ggml_cuda_parse_nvfp4_mm_lane(cgraph, i, lane) || lane.scale == nullptr) { + return 0; + } + + const int out_idx = i + lane.n_nodes - 1; + const int out_nodes[] = { out_idx }; + if (!ggml_cuda_can_fuse_mmid_scale_subgraph(cgraph, i, lane.n_nodes, out_nodes, 1) || + !ggml_cuda_check_fusion_memory_ranges(cgraph, i, lane.n_nodes, out_nodes, 1)) { + return 0; + } + + if (!ggml_cuda_should_fuse_mul_mat_vec_q(lane.mm)) { + return 0; + } + + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.x_bias = lane.bias; + fusion_data.x_scale = lane.scale; + + ggml_cuda_mul_mat_vec_q(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], nullptr, lane.out, &fusion_data); + return lane.n_nodes - 1; +} + // try and fuse nodes and return the number of nodes to skip static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { @@ -4291,6 +4459,10 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph if (fused_nvfp4_mmid_nodes > 0) { return fused_nvfp4_mmid_nodes; } + fused_nvfp4_mmid_nodes = ggml_cuda_try_fuse_nvfp4_mm_scale_glu(cuda_ctx, cgraph, i); + if (fused_nvfp4_mmid_nodes > 0) { + return fused_nvfp4_mmid_nodes; + } } bool fused_mul_mat_vec = false; @@ -4428,6 +4600,10 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph if (fused_nvfp4_mmid_nodes > 0) { return fused_nvfp4_mmid_nodes; } + fused_nvfp4_mmid_nodes = ggml_cuda_try_fuse_nvfp4_mm_scale(cuda_ctx, cgraph, i); + if (fused_nvfp4_mmid_nodes > 0) { + return fused_nvfp4_mmid_nodes; + } } // gate + add + glu + up + add diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index f91afd7e73..6aad0b3059 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -572,10 +572,10 @@ static __global__ void mul_mat_vec_q( } } if (use_scale) { - x_scales = x_scale[ids ? channel_x : channel_dst]; + x_scales = x_scale[ids ? channel_x : 0]; } if (use_gate_scale) { - gate_scales = gate_scale[ids ? channel_x : channel_dst]; + gate_scales = gate_scale[ids ? channel_x : 0]; } } @@ -1194,19 +1194,17 @@ void ggml_cuda_mul_mat_vec_q( fusion_local.gate_bias = fusion->gate_bias->data; } if (fusion->x_scale) { - GGML_ASSERT(ids); GGML_ASSERT(src0->type == GGML_TYPE_NVFP4); GGML_ASSERT(fusion->x_scale->type == GGML_TYPE_F32); GGML_ASSERT(ggml_is_contiguous(fusion->x_scale)); - GGML_ASSERT(ggml_nelements(fusion->x_scale) == src0->ne[2]); + GGML_ASSERT(ggml_nelements(fusion->x_scale) == (ids ? src0->ne[2] : 1)); fusion_local.x_scale = fusion->x_scale->data; } if (fusion->gate_scale) { - GGML_ASSERT(ids); GGML_ASSERT(src0->type == GGML_TYPE_NVFP4); GGML_ASSERT(fusion->gate_scale->type == GGML_TYPE_F32); GGML_ASSERT(ggml_is_contiguous(fusion->gate_scale)); - GGML_ASSERT(ggml_nelements(fusion->gate_scale) == src0->ne[2]); + GGML_ASSERT(ggml_nelements(fusion->gate_scale) == (ids ? src0->ne[2] : 1)); fusion_local.gate_scale = fusion->gate_scale->data; } fusion_local.glu_op = fusion->glu_op;