From f3d3398b975170a0674851966200a06af652b27a Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Tue, 2 Jun 2026 11:48:35 +0200 Subject: [PATCH 01/28] CUDA: Fuse MMVQ for NVFP4 and BS 1 TODO: 1. Add tests to test-backend-ops (did verify correctness manually for one model) 2. Reorder bias/scale once PRs for NVFP4 are merged/landed --- ggml/src/ggml-cuda/common.cuh | 4 + ggml/src/ggml-cuda/ggml-cuda.cu | 268 +++++++++++++++++++++++++++++++- ggml/src/ggml-cuda/mmvq.cu | 60 +++++-- 3 files changed, 316 insertions(+), 16 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index e6e50e041195..290dc4aff259 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1505,12 +1505,16 @@ struct ggml_cuda_mm_fusion_args_host { const ggml_tensor * x_bias = nullptr; const ggml_tensor * gate = nullptr; const ggml_tensor * gate_bias = nullptr; + const ggml_tensor * x_scale = nullptr; + const ggml_tensor * gate_scale = nullptr; ggml_glu_op glu_op; }; struct ggml_cuda_mm_fusion_args_device { const void * x_bias = nullptr; const void * gate = nullptr; const void * gate_bias = nullptr; + const void * x_scale = nullptr; + const void * gate_scale = nullptr; ggml_glu_op glu_op; }; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index e779a9be9e95..620e71cdef92 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3875,10 +3875,251 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, return false; } + +struct ggml_cuda_mmid_lane { + ggml_tensor * mm = nullptr; + ggml_tensor * bias_node = nullptr; + ggml_tensor * out = nullptr; + const ggml_tensor * bias = nullptr; + const ggml_tensor * scale = nullptr; + int n_nodes = 0; +}; + +static bool ggml_cuda_parse_nvfp4_mmid_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_ID) { + 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 || mm->src[2] == nullptr) { + 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_ID) { + ggml_tensor * add = cgraph->nodes[i + lane.n_nodes]; + if (add->src[0] != lane.out || add->src[2] != mm->src[2]) { + return false; + } + lane.bias_node = add; + lane.bias = add->src[1]; + lane.out = add; + lane.n_nodes++; + } + + if (i + lane.n_nodes + 3 >= cgraph->n_nodes) { + return true; + } + + ggml_tensor * reshape = cgraph->nodes[i + lane.n_nodes + 0]; + ggml_tensor * repeat = cgraph->nodes[i + lane.n_nodes + 1]; + ggml_tensor * getrows = cgraph->nodes[i + lane.n_nodes + 2]; + ggml_tensor * mul = cgraph->nodes[i + lane.n_nodes + 3]; + + if (reshape->op != GGML_OP_RESHAPE || repeat->op != GGML_OP_REPEAT || + getrows->op != GGML_OP_GET_ROWS || mul->op != GGML_OP_MUL) { + return true; + } + + if (repeat->src[0] != reshape || getrows->src[0] != repeat || getrows->src[1] != mm->src[2]) { + return true; + } + + const bool mul_has_out = mul->src[0] == lane.out || mul->src[1] == lane.out; + const bool mul_has_scale = mul->src[0] == getrows || mul->src[1] == getrows; + if (!mul_has_out || !mul_has_scale) { + return true; + } + + const ggml_tensor * scale = reshape->src[0]; + if (scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != mm->src[0]->ne[2]) { + 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 += 4; + 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; + if (idx >= cgraph->n_nodes) { + return false; + } + + const ggml_tensor * node = cgraph->nodes[idx]; + bool is_output = false; + for (int k = 0; k < num_outputs; ++k) { + if (outputs[k] < cgraph->n_nodes && cgraph->nodes[outputs[k]] == node) { + is_output = true; + break; + } + } + if (is_output) { + continue; + } + + if (node->flags & GGML_TENSOR_FLAG_OUTPUT) { + return false; + } + + int subgraph_uses = 0; + for (int k = j + 1; k < count; ++k) { + const ggml_tensor * other_node = cgraph->nodes[start_idx + k]; + for (int src_idx = 0; src_idx < GGML_MAX_SRC; ++src_idx) { + if (other_node->src[src_idx] == node) { + subgraph_uses++; + } + } + } + + if (subgraph_uses != ggml_node_get_use_count(cgraph, idx)) { + return false; + } + } + + return true; +} + +static bool ggml_cuda_should_fuse_mmid_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] || up.mm->src[2] != gate.mm->src[2]) { + 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)) { + return 0; + } + + ggml_cuda_mmid_lane lane1; + if (!ggml_cuda_parse_nvfp4_mmid_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_mmid_lanes(*up, *gate, glu)) { + return 0; + } + + std::vector ops; + for (int j = i; j <= glu_idx; ++j) { + ops.push_back(cgraph->nodes[j]->op); + } + + int out_nodes[] = { glu_idx }; + if (!ggml_cuda_can_fuse_mmid_scale_subgraph(cgraph, i, (int) ops.size(), out_nodes, 1) || + !ggml_cuda_check_fusion_memory_ranges(cgraph, i, (int) ops.size(), 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], up->mm->src[2], 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) { + return 0; + } + + std::vector ops; + for (int j = 0; j < lane.n_nodes; ++j) { + ops.push_back(cgraph->nodes[i + j]->op); + } + + const int out_idx = i + lane.n_nodes - 1; + 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], lane.mm->src[2], 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) { static bool disable_fusion = getenv("GGML_CUDA_DISABLE_FUSION") != nullptr && std::atoi(getenv("GGML_CUDA_DISABLE_FUSION")); + const char * disable_nvfp4_mmid_scale_fusion_env = getenv("GGML_CUDA_DISABLE_NVFP4_MMID_SCALE_FUSION"); + const bool disable_nvfp4_mmid_scale_fusion = + disable_nvfp4_mmid_scale_fusion_env != nullptr && std::atoi(disable_nvfp4_mmid_scale_fusion_env); if (disable_fusion) { return 0; } @@ -4044,6 +4285,14 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph } } + int fused_nvfp4_mmid_nodes = 0; + if (!disable_nvfp4_mmid_scale_fusion) { + fused_nvfp4_mmid_nodes = ggml_cuda_try_fuse_nvfp4_mmid_scale_glu(cuda_ctx, cgraph, i); + if (fused_nvfp4_mmid_nodes > 0) { + return fused_nvfp4_mmid_nodes; + } + } + bool fused_mul_mat_vec = false; int fused_node_count = 0; @@ -4174,6 +4423,13 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph fused_mul_mat_vec = false; fused_node_count = 0; + if (!disable_nvfp4_mmid_scale_fusion) { + fused_nvfp4_mmid_nodes = ggml_cuda_try_fuse_nvfp4_mmid_scale(cuda_ctx, cgraph, i); + if (fused_nvfp4_mmid_nodes > 0) { + return fused_nvfp4_mmid_nodes; + } + } + // gate + add + glu + up + add for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) { const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID; @@ -4405,12 +4661,6 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud } } -#ifdef GGML_CUDA_DEBUG - const int nodes_fused = i - prev_i - 1; - if (nodes_fused > 0) { - GGML_LOG_INFO("nodes_fused: %d\n", nodes_fused); - } -#endif prev_i = i; if (ggml_is_empty(node) || node->op == GGML_OP_RESHAPE || node->op == GGML_OP_TRANSPOSE || node->op == GGML_OP_VIEW || node->op == GGML_OP_PERMUTE || node->op == GGML_OP_NONE) { @@ -4424,6 +4674,12 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud int nodes_to_skip = ggml_cuda_try_fuse(cuda_ctx, cgraph, i); if (nodes_to_skip != 0) { +#ifdef GGML_CUDA_DEBUG + const int last_fused = i + nodes_to_skip; + GGML_LOG_INFO("nodes_fused: %d, first: %s (%s), last: %s (%s)\n", + nodes_to_skip, ggml_op_name(node->op), node->name, + ggml_op_name(cgraph->nodes[last_fused]->op), cgraph->nodes[last_fused]->name); +#endif i += nodes_to_skip; continue; } diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index bdfbfd2d387f..f91afd7e7300 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -519,24 +519,34 @@ static __global__ void mul_mat_vec_q( bool use_gate = false; bool use_bias = false; bool use_gate_bias = false; + bool use_scale = false; + bool use_gate_scale = false; [[maybe_unused]] const void * vgate = nullptr; const float * x_bias = nullptr; const float * gate_bias = nullptr; + const float * x_scale = nullptr; + const float * gate_scale = nullptr; 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; - vgate = fusion.gate; - x_bias = (const float *) fusion.x_bias; - gate_bias = (const float *) fusion.gate_bias; - 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; + 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; } [[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; if constexpr (has_fusion) { const uint32_t channel_bias = ids ? channel_x : channel_dst; if (use_bias) { @@ -561,6 +571,12 @@ static __global__ void mul_mat_vec_q( } } } + if (use_scale) { + x_scales = x_scale[ids ? channel_x : channel_dst]; + } + if (use_gate_scale) { + gate_scales = gate_scale[ids ? channel_x : channel_dst]; + } } // partial sum for each thread @@ -644,11 +660,17 @@ static __global__ void mul_mat_vec_q( if (use_bias) { result += x_biases[j]; } + 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; + } switch (active_glu) { case GGML_GLU_OP_SWIGLU: result *= ggml_cuda_op_silu_single(gate_value); @@ -671,7 +693,7 @@ static __global__ void mul_mat_vec_q( } if constexpr (!has_fusion) { - GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, active_glu, gate_bias, x_bias, tmp_gate); + 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); } } @@ -767,7 +789,8 @@ static void mul_mat_vec_q_switch_fusion( const dim3 & block_nums, const dim3 & block_dims, const int nbytes_shared, const uint32_t ids_stride, cudaStream_t stream) { - const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr; + const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr || + fusion.x_scale != nullptr || fusion.gate_scale != nullptr; if constexpr (c_ncols_dst == 1) { if (has_fusion) { const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream); @@ -832,7 +855,8 @@ static void mul_mat_vec_q_switch_ncols_dst( const int warp_size = ggml_cuda_info().devices[device].warp_size; const mmvq_parameter_table_id table_id = get_device_table_id(cc); - const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr; + const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr || + fusion.x_scale != nullptr || fusion.gate_scale != nullptr; const bool has_ids = ids != nullptr; const auto should_use_small_k = [&](int c_ncols_dst) { @@ -1169,6 +1193,22 @@ void ggml_cuda_mul_mat_vec_q( GGML_ASSERT(!ids || fusion->gate_bias->ne[1] == src0->ne[2]); 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]); + 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]); + fusion_local.gate_scale = fusion->gate_scale->data; + } fusion_local.glu_op = fusion->glu_op; } From 45d47042676435f5fdda358252432a42ebdba446 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 29 May 2026 13:58:32 +0200 Subject: [PATCH 02/28] 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 620e71cdef92..187c689458a9 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 f91afd7e7300..6aad0b3059c5 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; From 2173d97dd41850e28d836aeae5378bff76bd951a Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 29 May 2026 14:09:01 +0200 Subject: [PATCH 03/28] Add tests for the added fusion ops --- tests/test-backend-ops.cpp | 60 ++++++++++++++++++++++++++++++++++---- 1 file changed, 54 insertions(+), 6 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index f561d09b5b71..2de84bcae91c 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5788,19 +5788,21 @@ struct test_mul_mat_vec_fusion : public test_case { const bool b; // broadcast b matrix (only for use_id) const bool with_bias; const bool with_gate; + const bool with_scale; std::array batch_dims; test_mul_mat_vec_fusion(ggml_type type, ggml_glu_op op, int64_t m, int64_t n, int64_t k, bool use_id = false, int n_mats = 1, int n_used = 1, bool b = false, bool with_bias = false, bool with_gate = true, - std::array batch_dims = {4, 2}) - : type(type), glu_op(op), m(m), n(n), k(k), use_id(use_id), n_mats(n_mats), n_used(n_used), b(b), with_bias(with_bias), with_gate(with_gate), batch_dims(batch_dims) { + std::array batch_dims = {4, 2}, bool with_scale = false) + : type(type), glu_op(op), m(m), n(n), k(k), use_id(use_id), n_mats(n_mats), n_used(n_used), b(b), with_bias(with_bias), + with_gate(with_gate), with_scale(with_scale), batch_dims(batch_dims) { if (use_id) { GGML_ASSERT(n_used <= n_mats); } } std::string vars() override { - return VARS_TO_STR12(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, batch_dims); + return VARS_TO_STR13(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_scale, batch_dims); } std::string op_desc(ggml_tensor * t) override { @@ -5824,6 +5826,19 @@ struct test_mul_mat_vec_fusion : public test_case { return out; } + ggml_tensor * build_dense_scale(ggml_context * ctx, ggml_tensor * out) { + ggml_tensor * scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 1); + return ggml_mul(ctx, out, scale); + } + + ggml_tensor * build_id_scale(ggml_context * ctx, ggml_tensor * out, ggml_tensor * ids) { + ggml_tensor * scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_mats); + ggml_tensor * s = ggml_reshape_3d(ctx, scale, 1, n_mats, 1); + s = ggml_repeat_4d(ctx, s, 1, n_mats, m, 1); + s = ggml_get_rows(ctx, s, ids); + return ggml_mul(ctx, out, s); + } + ggml_tensor * build_graph(ggml_context * ctx) override { if (!use_id) { const int channels = batch_dims[0]; @@ -5841,6 +5856,9 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * up_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); ffn_up = ggml_add(ctx, ffn_up, up_bias); } + if (with_scale) { + ffn_up = build_dense_scale(ctx, ffn_up); + } ggml_tensor * ffn_gate = with_gate ? ggml_mul_mat(ctx, gate, cur) : nullptr; if (with_bias && with_gate) { @@ -5848,6 +5866,9 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * gate_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); ffn_gate = ggml_add(ctx, ffn_gate, gate_bias); } + if (with_scale && with_gate) { + ffn_gate = build_dense_scale(ctx, ffn_gate); + } ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; @@ -5874,18 +5895,26 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * up_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_up->ne[0], n_mats); ffn_up = ggml_add_id(ctx, ffn_up, up_bias_param, ids); } + if (with_scale) { + ffn_up = build_id_scale(ctx, ffn_up, ids); + } ggml_tensor * ffn_gate = with_gate? ggml_mul_mat_id(ctx, gates, cur, ids) : nullptr; if (with_bias && with_gate) { ggml_tensor * gate_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_gate->ne[0], n_mats); ffn_gate = ggml_add_id(ctx, ffn_gate, gate_bias_param, ids); } + if (with_scale && with_gate) { + ffn_gate = build_id_scale(ctx, ffn_gate, ids); + } ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; - std::array scale_ne { 1, out->ne[1], out->ne[2], out->ne[3] }; - ggml_tensor * scale = ggml_new_tensor(ctx, out->type, 4, scale_ne.data()); - out = ggml_mul(ctx, out, scale); + if (!with_scale) { + std::array scale_ne { 1, out->ne[1], out->ne[2], out->ne[3] }; + ggml_tensor * scale = ggml_new_tensor(ctx, out->type, 4, scale_ne.data()); + out = ggml_mul(ctx, out, scale); + } ggml_set_name(out, "out"); return out; @@ -5905,6 +5934,13 @@ struct test_mul_mat_vec_fusion : public test_case { double max_nmse_err() override { return 5e-3; } + + double max_nmse_err(ggml_backend_t backend) override { + if (type == GGML_TYPE_NVFP4 && backend_has_feature(backend, "BLACKWELL_NATIVE_FP4")) { + return 2e-2; + } + return max_nmse_err(); + } }; // GGML_OP_SUM @@ -9094,6 +9130,18 @@ static std::vector> make_test_cases_eval() { } } + // NVFP4 scale2 fusion patterns: dense/MMID, GLU/two-lane, single-lane, with and without bias. + for (bool with_bias : {false, true}) { + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 1, 32, 256, + false, 1, 1, false, with_bias, true, {1, 1}, true)); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 1, 32, 256, + true, 16, 8, false, with_bias, true, {1, 1}, true)); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 1, 32, 256, + false, 1, 1, false, with_bias, false, {1, 1}, true)); + test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 1, 32, 256, + true, 16, 8, false, with_bias, false, {1, 1}, true)); + } + for (auto gate : {GATING_FUNC_SOFTMAX, GATING_FUNC_SIGMOID, GATING_FUNC_SOFTMAX_WEIGHT}) { for (bool with_norm : {false, true}) { for (bool bias_probs : {false, true}) { From 4ed8ef226fee5b80793a8e0ccc8ce10213a7a118 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Tue, 2 Jun 2026 12:16:01 +0200 Subject: [PATCH 04/28] Cleanup test-backend-ops --- tests/test-backend-ops.cpp | 65 ++++++++++++++------------------------ 1 file changed, 23 insertions(+), 42 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 2de84bcae91c..84e6b8e9fca5 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5788,21 +5788,21 @@ struct test_mul_mat_vec_fusion : public test_case { const bool b; // broadcast b matrix (only for use_id) const bool with_bias; const bool with_gate; - const bool with_scale; + const bool with_lane_scale; std::array batch_dims; test_mul_mat_vec_fusion(ggml_type type, ggml_glu_op op, int64_t m, int64_t n, int64_t k, bool use_id = false, int n_mats = 1, int n_used = 1, bool b = false, bool with_bias = false, bool with_gate = true, - std::array batch_dims = {4, 2}, bool with_scale = false) + bool with_lane_scale = false, std::array batch_dims = {4, 2}) : type(type), glu_op(op), m(m), n(n), k(k), use_id(use_id), n_mats(n_mats), n_used(n_used), b(b), with_bias(with_bias), - with_gate(with_gate), with_scale(with_scale), batch_dims(batch_dims) { + with_gate(with_gate), with_lane_scale(with_lane_scale), batch_dims(batch_dims) { if (use_id) { GGML_ASSERT(n_used <= n_mats); } } std::string vars() override { - return VARS_TO_STR13(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_scale, batch_dims); + return VARS_TO_STR13(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_lane_scale, batch_dims); } std::string op_desc(ggml_tensor * t) override { @@ -5826,12 +5826,12 @@ struct test_mul_mat_vec_fusion : public test_case { return out; } - ggml_tensor * build_dense_scale(ggml_context * ctx, ggml_tensor * out) { + ggml_tensor * build_dense_lane_scale(ggml_context * ctx, ggml_tensor * out) { ggml_tensor * scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 1); return ggml_mul(ctx, out, scale); } - ggml_tensor * build_id_scale(ggml_context * ctx, ggml_tensor * out, ggml_tensor * ids) { + ggml_tensor * build_id_lane_scale(ggml_context * ctx, ggml_tensor * out, ggml_tensor * ids) { ggml_tensor * scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_mats); ggml_tensor * s = ggml_reshape_3d(ctx, scale, 1, n_mats, 1); s = ggml_repeat_4d(ctx, s, 1, n_mats, m, 1); @@ -5856,8 +5856,8 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * up_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); ffn_up = ggml_add(ctx, ffn_up, up_bias); } - if (with_scale) { - ffn_up = build_dense_scale(ctx, ffn_up); + if (with_lane_scale) { + ffn_up = build_dense_lane_scale(ctx, ffn_up); } ggml_tensor * ffn_gate = with_gate ? ggml_mul_mat(ctx, gate, cur) : nullptr; @@ -5866,8 +5866,8 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * gate_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); ffn_gate = ggml_add(ctx, ffn_gate, gate_bias); } - if (with_scale && with_gate) { - ffn_gate = build_dense_scale(ctx, ffn_gate); + if (with_lane_scale && with_gate) { + ffn_gate = build_dense_lane_scale(ctx, ffn_gate); } ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; @@ -5895,8 +5895,8 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * up_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_up->ne[0], n_mats); ffn_up = ggml_add_id(ctx, ffn_up, up_bias_param, ids); } - if (with_scale) { - ffn_up = build_id_scale(ctx, ffn_up, ids); + if (with_lane_scale) { + ffn_up = build_id_lane_scale(ctx, ffn_up, ids); } ggml_tensor * ffn_gate = with_gate? ggml_mul_mat_id(ctx, gates, cur, ids) : nullptr; @@ -5904,17 +5904,15 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * gate_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_gate->ne[0], n_mats); ffn_gate = ggml_add_id(ctx, ffn_gate, gate_bias_param, ids); } - if (with_scale && with_gate) { - ffn_gate = build_id_scale(ctx, ffn_gate, ids); + if (with_lane_scale && with_gate) { + ffn_gate = build_id_lane_scale(ctx, ffn_gate, ids); } ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; - if (!with_scale) { - std::array scale_ne { 1, out->ne[1], out->ne[2], out->ne[3] }; - ggml_tensor * scale = ggml_new_tensor(ctx, out->type, 4, scale_ne.data()); - out = ggml_mul(ctx, out, scale); - } + std::array scale_ne { 1, out->ne[1], out->ne[2], out->ne[3] }; + ggml_tensor * scale = ggml_new_tensor(ctx, out->type, 4, scale_ne.data()); + out = ggml_mul(ctx, out, scale); ggml_set_name(out, "out"); return out; @@ -5934,13 +5932,6 @@ struct test_mul_mat_vec_fusion : public test_case { double max_nmse_err() override { return 5e-3; } - - double max_nmse_err(ggml_backend_t backend) override { - if (type == GGML_TYPE_NVFP4 && backend_has_feature(backend, "BLACKWELL_NATIVE_FP4")) { - return 2e-2; - } - return max_nmse_err(); - } }; // GGML_OP_SUM @@ -9119,10 +9110,12 @@ static std::vector> make_test_cases_eval() { if (!with_gate && glu_op != GGML_GLU_OP_SWIGLU) { continue; } - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, - use_id, 16, 8, b, with_bias, with_gate)); - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, - use_id, 16, 8, b, with_bias, with_gate, {1, 1})); + for (bool with_lane_scale : {false, true}) { + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, + use_id, 16, 8, b, with_bias, with_gate, with_lane_scale)); + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, + use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, {1, 1})); + } } } } @@ -9130,18 +9123,6 @@ static std::vector> make_test_cases_eval() { } } - // NVFP4 scale2 fusion patterns: dense/MMID, GLU/two-lane, single-lane, with and without bias. - for (bool with_bias : {false, true}) { - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 1, 32, 256, - false, 1, 1, false, with_bias, true, {1, 1}, true)); - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 1, 32, 256, - true, 16, 8, false, with_bias, true, {1, 1}, true)); - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 1, 32, 256, - false, 1, 1, false, with_bias, false, {1, 1}, true)); - test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 1, 32, 256, - true, 16, 8, false, with_bias, false, {1, 1}, true)); - } - for (auto gate : {GATING_FUNC_SOFTMAX, GATING_FUNC_SIGMOID, GATING_FUNC_SOFTMAX_WEIGHT}) { for (bool with_norm : {false, true}) { for (bool bias_probs : {false, true}) { From d6190b97a560e427c19c6d72bbc329383aa2cf76 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Tue, 2 Jun 2026 15:08:52 +0200 Subject: [PATCH 05/28] Cleanup ggml-cuda/mmvq 1. Unrestrict post-scale fusion 2. Rename names accordingly 3. Remove env variable to disable fusion --- ggml/src/ggml-cuda/ggml-cuda.cu | 64 +++++++++++++++------------------ ggml/src/ggml-cuda/mmvq.cu | 2 -- 2 files changed, 28 insertions(+), 38 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 187c689458a9..12f07dceef6a 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3885,13 +3885,13 @@ struct ggml_cuda_mmid_lane { int n_nodes = 0; }; -static bool ggml_cuda_parse_nvfp4_mmid_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mmid_lane & lane) { +static bool ggml_cuda_parse_mmid_scale_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_ID) { 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 || mm->src[2] == nullptr) { + if (!ggml_is_quantized(mm->src[0]->type) || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32 || mm->src[2] == nullptr) { return false; } @@ -3950,13 +3950,13 @@ 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) { +static bool ggml_cuda_parse_mm_scale_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) { + if (!ggml_is_quantized(mm->src[0]->type) || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32) { return false; } @@ -4102,14 +4102,14 @@ static bool ggml_cuda_should_fuse_mm_lanes(const ggml_cuda_mmid_lane & up, const return !split; } -static int ggml_cuda_try_fuse_nvfp4_mmid_scale_glu(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { +static int ggml_cuda_try_fuse_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)) { + if (!ggml_cuda_parse_mmid_scale_lane(cgraph, i, lane0)) { return 0; } ggml_cuda_mmid_lane lane1; - if (!ggml_cuda_parse_nvfp4_mmid_lane(cgraph, i + lane0.n_nodes, lane1)) { + if (!ggml_cuda_parse_mmid_scale_lane(cgraph, i + lane0.n_nodes, lane1)) { return 0; } @@ -4166,14 +4166,14 @@ 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) { +static int ggml_cuda_try_fuse_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)) { + if (!ggml_cuda_parse_mm_scale_lane(cgraph, i, lane0)) { return 0; } ggml_cuda_mmid_lane lane1; - if (!ggml_cuda_parse_nvfp4_mm_lane(cgraph, i + lane0.n_nodes, lane1)) { + if (!ggml_cuda_parse_mm_scale_lane(cgraph, i + lane0.n_nodes, lane1)) { return 0; } @@ -4226,9 +4226,9 @@ static int ggml_cuda_try_fuse_nvfp4_mm_scale_glu(ggml_backend_cuda_context * cud return glu_idx - i; } -static int ggml_cuda_try_fuse_nvfp4_mmid_scale(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { +static int ggml_cuda_try_fuse_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) { + if (!ggml_cuda_parse_mmid_scale_lane(cgraph, i, lane) || lane.scale == nullptr) { return 0; } @@ -4256,9 +4256,9 @@ 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) { +static int ggml_cuda_try_fuse_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) { + if (!ggml_cuda_parse_mm_scale_lane(cgraph, i, lane) || lane.scale == nullptr) { return 0; } @@ -4285,9 +4285,6 @@ static int ggml_cuda_try_fuse_nvfp4_mm_scale(ggml_backend_cuda_context * cuda_ct static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { static bool disable_fusion = getenv("GGML_CUDA_DISABLE_FUSION") != nullptr && std::atoi(getenv("GGML_CUDA_DISABLE_FUSION")); - const char * disable_nvfp4_mmid_scale_fusion_env = getenv("GGML_CUDA_DISABLE_NVFP4_MMID_SCALE_FUSION"); - const bool disable_nvfp4_mmid_scale_fusion = - disable_nvfp4_mmid_scale_fusion_env != nullptr && std::atoi(disable_nvfp4_mmid_scale_fusion_env); if (disable_fusion) { return 0; } @@ -4453,16 +4450,13 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph } } - int fused_nvfp4_mmid_nodes = 0; - if (!disable_nvfp4_mmid_scale_fusion) { - fused_nvfp4_mmid_nodes = ggml_cuda_try_fuse_nvfp4_mmid_scale_glu(cuda_ctx, cgraph, i); - 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; - } + int fused_scale_nodes = ggml_cuda_try_fuse_mmid_scale_glu(cuda_ctx, cgraph, i); + if (fused_scale_nodes > 0) { + return fused_scale_nodes; + } + fused_scale_nodes = ggml_cuda_try_fuse_mm_scale_glu(cuda_ctx, cgraph, i); + if (fused_scale_nodes > 0) { + return fused_scale_nodes; } bool fused_mul_mat_vec = false; @@ -4595,15 +4589,13 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph fused_mul_mat_vec = false; fused_node_count = 0; - if (!disable_nvfp4_mmid_scale_fusion) { - fused_nvfp4_mmid_nodes = ggml_cuda_try_fuse_nvfp4_mmid_scale(cuda_ctx, cgraph, i); - 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; - } + fused_scale_nodes = ggml_cuda_try_fuse_mmid_scale(cuda_ctx, cgraph, i); + if (fused_scale_nodes > 0) { + return fused_scale_nodes; + } + fused_scale_nodes = ggml_cuda_try_fuse_mm_scale(cuda_ctx, cgraph, i); + if (fused_scale_nodes > 0) { + return fused_scale_nodes; } // gate + add + glu + up + add diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 6aad0b3059c5..5a4b2fdf4233 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -1194,14 +1194,12 @@ void ggml_cuda_mul_mat_vec_q( fusion_local.gate_bias = fusion->gate_bias->data; } if (fusion->x_scale) { - 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) == (ids ? src0->ne[2] : 1)); fusion_local.x_scale = fusion->x_scale->data; } if (fusion->gate_scale) { - 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) == (ids ? src0->ne[2] : 1)); From b07d0f03d3996717c70a5c7ab1faffae13971a66 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Tue, 2 Jun 2026 15:57:11 +0200 Subject: [PATCH 06/28] Merge old mul_mat patterns into the lane-based approach --- ggml/src/ggml-cuda/ggml-cuda.cu | 543 ++++++-------------------------- tests/test-backend-ops.cpp | 116 ++++--- 2 files changed, 171 insertions(+), 488 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 12f07dceef6a..fd0b12cbe95f 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2418,91 +2418,6 @@ static void ggml_cuda_mul_mat_batched_cublas(ggml_backend_cuda_context & ctx, co } } -static bool ggml_cuda_should_fuse_mul_mat(const ggml_tensor * ffn_up, - const ggml_tensor * ffn_gate, - const ggml_tensor * glu, - const ggml_tensor * ffn_up_bias = nullptr, - const ggml_tensor * ffn_gate_bias = nullptr) { - const bool has_bias = ffn_up_bias != nullptr || ffn_gate_bias != nullptr; - - if (has_bias && (!ffn_up_bias || !ffn_gate_bias)) { - return false; - } - - const bool is_mul_mat = ffn_up->op == GGML_OP_MUL_MAT && ffn_gate->op == GGML_OP_MUL_MAT && glu->op == GGML_OP_GLU; - const bool is_mul_mat_id = ffn_up->op == GGML_OP_MUL_MAT_ID && ffn_gate->op == GGML_OP_MUL_MAT_ID && glu->op == GGML_OP_GLU; - - GGML_ASSERT(ffn_up && ffn_gate && glu); - - if (!is_mul_mat && !is_mul_mat_id) { - return false; - } - - const ggml_op expected_bias_op = is_mul_mat ? GGML_OP_ADD : GGML_OP_ADD_ID; - - if (has_bias) { - if (ffn_up_bias->op != expected_bias_op || ffn_gate_bias->op != expected_bias_op) { - return false; - } - - if (glu->src[0] != ffn_gate_bias || glu->src[1] != ffn_up_bias) { - return false; - } - - if (expected_bias_op == GGML_OP_ADD) { - const bool up_has_mul = ffn_up_bias->src[0] == ffn_up || ffn_up_bias->src[1] == ffn_up; - const bool gate_has_mul = ffn_gate_bias->src[0] == ffn_gate || ffn_gate_bias->src[1] == ffn_gate; - if (!up_has_mul || !gate_has_mul) { - return false; - } - } else { // GGML_OP_ADD_ID - if (ffn_up_bias->src[0] != ffn_up || ffn_gate_bias->src[0] != ffn_gate) { - return false; - } - if (ffn_up_bias->src[2] != ffn_up->src[2] || ffn_gate_bias->src[2] != ffn_gate->src[2]) { - return false; - } - } - } else { - if (glu->src[0] != ffn_gate && glu->src[1] != ffn_up) { - return false; - } - } - - if (ffn_up->src[0]->type != ffn_gate->src[0]->type || !ggml_are_same_shape(ffn_up->src[0], ffn_gate->src[0]) || - !ggml_are_same_stride(ffn_up->src[0], ffn_gate->src[0])) { - return false; - } - - if (ffn_up->src[1] != ffn_gate->src[1]) { - return false; - } - - if (ffn_up->src[2] && (ffn_up->src[2] != ffn_gate->src[2])) { - 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(ffn_up->src[0]->buffer->buft) || - ggml_backend_buft_is_cuda_split(ffn_gate->src[0]->buffer->buft); - - //TODO: add support for fusion for split buffers - if (split) { - return false; - } - - return true; -} - static bool ggml_cuda_should_fuse_mul_mat_vec_f(const ggml_tensor * tensor) { ggml_tensor * src0 = tensor->src[0]; ggml_tensor * src1 = tensor->src[1]; @@ -3655,6 +3570,35 @@ static bool ggml_cuda_check_fusion_memory_ranges(const ggml_cgraph * cgraph, } +static bool ggml_cuda_can_fuse_subgraph(const struct ggml_cgraph * cgraph, + int node_idx, + int count, + const enum ggml_op * ops, + const int * out_nodes, + int out_count, + bool is_topk_moe = false) { + return ggml_can_fuse_subgraph(cgraph, node_idx, count, ops, out_nodes, out_count) && + ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, count, out_nodes, out_count, is_topk_moe); +} + +static bool ggml_cuda_can_fuse_parsed_subgraph(const struct ggml_cgraph * cgraph, + int node_idx, + int count, + const int * out_nodes, + int out_count, + bool is_topk_moe = false) { + if (node_idx + count > cgraph->n_nodes) { + return false; + } + + std::vector ops; + ops.reserve(count); + for (int j = 0; j < count; ++j) { + ops.push_back(cgraph->nodes[node_idx + j]->op); + } + return ggml_cuda_can_fuse_subgraph(cgraph, node_idx, count, ops.data(), out_nodes, out_count, is_topk_moe); +} + static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, int node_idx, std::initializer_list ops, @@ -3669,41 +3613,11 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, return std::equal(list1.begin(), list1.end(), list2.begin(), list2.end()); }; - std::initializer_list mul_mat_bias_glu_ops = { GGML_OP_MUL_MAT, GGML_OP_ADD, GGML_OP_MUL_MAT, GGML_OP_ADD, GGML_OP_GLU }; - std::initializer_list mul_mat_id_bias_glu_ops = { GGML_OP_MUL_MAT_ID, GGML_OP_ADD_ID, GGML_OP_MUL_MAT_ID, GGML_OP_ADD_ID, GGML_OP_GLU }; - - std::initializer_list mul_mat_id_glu_ops = { GGML_OP_MUL_MAT_ID, GGML_OP_MUL_MAT_ID, GGML_OP_GLU }; - std::initializer_list mul_mat_glu_ops = { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT, GGML_OP_GLU }; - - if ((is_equal(mul_mat_bias_glu_ops, ops) || is_equal(mul_mat_id_bias_glu_ops, ops)) && - ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 4 })) { - const ggml_tensor * ffn_gate = cgraph->nodes[node_idx]; - const ggml_tensor * ffn_gate_bias = cgraph->nodes[node_idx + 1]; - const ggml_tensor * ffn_up = cgraph->nodes[node_idx + 2]; - const ggml_tensor * ffn_up_bias = cgraph->nodes[node_idx + 3]; - const ggml_tensor * glu = cgraph->nodes[node_idx + 4]; - - if (ggml_cuda_should_fuse_mul_mat(ffn_up, ffn_gate, glu, ffn_up_bias, ffn_gate_bias)) { - int out_nodes[] = { node_idx + 4 }; - return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1); - } - } - - if ((is_equal(mul_mat_id_glu_ops, ops) || is_equal(mul_mat_glu_ops, ops)) && - ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 2 })) { - const ggml_tensor * ffn_gate = cgraph->nodes[node_idx]; - const ggml_tensor * ffn_up = cgraph->nodes[node_idx + 1]; - const ggml_tensor * glu = cgraph->nodes[node_idx + 2]; - - if (ggml_cuda_should_fuse_mul_mat(ffn_up, ffn_gate, glu)) { - int out_nodes[] = { node_idx + 2 }; - return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1); - } - } - std::initializer_list rope_set_rows_ops = { GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS }; - if (is_equal(rope_set_rows_ops, ops) && ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 2 })) { + const int rope_set_rows_out_nodes[] = { node_idx + 2 }; + if (is_equal(rope_set_rows_ops, ops) && + ggml_cuda_can_fuse_subgraph(cgraph, node_idx, (int) ops.size(), ops.begin(), rope_set_rows_out_nodes, 1)) { const ggml_tensor * rope = cgraph->nodes[node_idx]; const ggml_tensor * view = cgraph->nodes[node_idx + 1]; const ggml_tensor * set_rows = cgraph->nodes[node_idx + 2]; @@ -3876,7 +3790,7 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, } -struct ggml_cuda_mmid_lane { +struct ggml_cuda_mm_lane { ggml_tensor * mm = nullptr; ggml_tensor * bias_node = nullptr; ggml_tensor * out = nullptr; @@ -3885,13 +3799,17 @@ struct ggml_cuda_mmid_lane { int n_nodes = 0; }; -static bool ggml_cuda_parse_mmid_scale_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mmid_lane & lane) { +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_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mm_lane & lane) { if (i >= cgraph->n_nodes || cgraph->nodes[i]->op != GGML_OP_MUL_MAT_ID) { return false; } ggml_tensor * mm = cgraph->nodes[i]; - if (!ggml_is_quantized(mm->src[0]->type) || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32 || mm->src[2] == nullptr) { + if (!ggml_cuda_can_parse_mm_lane_type(mm->src[0]->type) || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32 || mm->src[2] == nullptr) { return false; } @@ -3911,7 +3829,7 @@ static bool ggml_cuda_parse_mmid_scale_lane(const ggml_cgraph * cgraph, int i, g lane.n_nodes++; } - if (i + lane.n_nodes + 3 >= cgraph->n_nodes) { + if (!ggml_is_quantized(mm->src[0]->type) || i + lane.n_nodes + 3 >= cgraph->n_nodes) { return true; } @@ -3950,13 +3868,13 @@ static bool ggml_cuda_parse_mmid_scale_lane(const ggml_cgraph * cgraph, int i, g return true; } -static bool ggml_cuda_parse_mm_scale_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mmid_lane & lane) { +static bool ggml_cuda_parse_mul_mat_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mm_lane & lane) { if (i >= cgraph->n_nodes || cgraph->nodes[i]->op != GGML_OP_MUL_MAT) { return false; } ggml_tensor * mm = cgraph->nodes[i]; - if (!ggml_is_quantized(mm->src[0]->type) || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32) { + if (!ggml_cuda_can_parse_mm_lane_type(mm->src[0]->type) || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32) { return false; } @@ -3974,12 +3892,15 @@ static bool ggml_cuda_parse_mm_scale_lane(const ggml_cgraph * cgraph, int i, ggm } else { return false; } + if (!ggml_are_same_shape(add->src[0], add->src[1])) { + 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) { + 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) { return true; } @@ -4005,76 +3926,24 @@ static bool ggml_cuda_parse_mm_scale_lane(const ggml_cgraph * cgraph, int i, ggm 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; - if (idx >= cgraph->n_nodes) { - return false; - } - - const ggml_tensor * node = cgraph->nodes[idx]; - bool is_output = false; - for (int k = 0; k < num_outputs; ++k) { - if (outputs[k] < cgraph->n_nodes && cgraph->nodes[outputs[k]] == node) { - is_output = true; - break; - } - } - if (is_output) { - continue; - } - - if (node->flags & GGML_TENSOR_FLAG_OUTPUT) { - return false; - } - - int subgraph_uses = 0; - for (int k = j + 1; k < count; ++k) { - const ggml_tensor * other_node = cgraph->nodes[start_idx + k]; - for (int src_idx = 0; src_idx < GGML_MAX_SRC; ++src_idx) { - if (other_node->src[src_idx] == node) { - subgraph_uses++; - } - } - } - - if (subgraph_uses != ggml_node_get_use_count(cgraph, idx)) { - return false; - } - } - - return true; -} - -static bool ggml_cuda_should_fuse_mmid_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] || up.mm->src[2] != gate.mm->src[2]) { +static bool ggml_cuda_parse_mm_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mm_lane & lane) { + if (i >= cgraph->n_nodes) { return false; } - - if (glu->op != GGML_OP_GLU || glu->src[0] != gate.out || glu->src[1] != up.out) { - return false; + if (cgraph->nodes[i]->op == GGML_OP_MUL_MAT_ID) { + return ggml_cuda_parse_mul_mat_id_lane(cgraph, i, lane); } - - 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 (cgraph->nodes[i]->op == GGML_OP_MUL_MAT) { + return ggml_cuda_parse_mul_mat_lane(cgraph, i, lane); } + return false; +} - if (const bool swapped = ggml_get_op_params_i32(glu, 1); swapped) { +static bool ggml_cuda_should_fuse_mm_lanes(const ggml_cuda_mm_lane & up, const ggml_cuda_mm_lane & gate, const ggml_tensor * glu) { + if (up.mm->op != gate.mm->op) { 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 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; @@ -4084,6 +3953,10 @@ static bool ggml_cuda_should_fuse_mm_lanes(const ggml_cuda_mmid_lane & up, const return false; } + if (up.mm->op == GGML_OP_MUL_MAT_ID && up.mm->src[2] != gate.mm->src[2]) { + return false; + } + if (glu->op != GGML_OP_GLU || glu->src[0] != gate.out || glu->src[1] != up.out) { return false; } @@ -4102,78 +3975,14 @@ static bool ggml_cuda_should_fuse_mm_lanes(const ggml_cuda_mmid_lane & up, const return !split; } -static int ggml_cuda_try_fuse_mmid_scale_glu(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { - ggml_cuda_mmid_lane lane0; - if (!ggml_cuda_parse_mmid_scale_lane(cgraph, i, lane0)) { - return 0; - } - - ggml_cuda_mmid_lane lane1; - if (!ggml_cuda_parse_mmid_scale_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_mmid_lanes(*up, *gate, glu)) { - return 0; - } - - std::vector ops; - for (int j = i; j <= glu_idx; ++j) { - ops.push_back(cgraph->nodes[j]->op); - } - - int out_nodes[] = { glu_idx }; - if (!ggml_cuda_can_fuse_mmid_scale_subgraph(cgraph, i, (int) ops.size(), out_nodes, 1) || - !ggml_cuda_check_fusion_memory_ranges(cgraph, i, (int) ops.size(), 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], up->mm->src[2], cgraph->nodes[glu_idx], &fusion_data); - return glu_idx - i; -} - -static int ggml_cuda_try_fuse_mm_scale_glu(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { - ggml_cuda_mmid_lane lane0; - if (!ggml_cuda_parse_mm_scale_lane(cgraph, i, lane0)) { +static int ggml_cuda_try_fuse_mm_glu(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { + ggml_cuda_mm_lane lane0; + if (!ggml_cuda_parse_mm_lane(cgraph, i, lane0)) { return 0; } - ggml_cuda_mmid_lane lane1; - if (!ggml_cuda_parse_mm_scale_lane(cgraph, i + lane0.n_nodes, lane1)) { + ggml_cuda_mm_lane lane1; + if (!ggml_cuda_parse_mm_lane(cgraph, i + lane0.n_nodes, lane1)) { return 0; } @@ -4182,13 +3991,9 @@ static int ggml_cuda_try_fuse_mm_scale_glu(ggml_backend_cuda_context * cuda_ctx, 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; + ggml_cuda_mm_lane * gate = nullptr; + ggml_cuda_mm_lane * up = nullptr; if (glu->src[0] == lane0.out && glu->src[1] == lane1.out) { gate = &lane0; up = &lane1; @@ -4205,12 +4010,7 @@ static int ggml_cuda_try_fuse_mm_scale_glu(ggml_backend_cuda_context * cuda_ctx, 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)) { + if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, n_nodes, out_nodes, 1)) { return 0; } @@ -4222,50 +4022,29 @@ static int ggml_cuda_try_fuse_mm_scale_glu(ggml_backend_cuda_context * cuda_ctx, 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_mmid_scale(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { - ggml_cuda_mmid_lane lane; - if (!ggml_cuda_parse_mmid_scale_lane(cgraph, i, lane) || lane.scale == nullptr) { - return 0; - } - - std::vector ops; - for (int j = 0; j < lane.n_nodes; ++j) { - ops.push_back(cgraph->nodes[i + j]->op); - } - - const int out_idx = i + lane.n_nodes - 1; - 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; + const ggml_tensor * ids = up->mm->op == GGML_OP_MUL_MAT_ID ? up->mm->src[2] : nullptr; + if (up->scale == nullptr && gate->scale == nullptr && ggml_cuda_should_fuse_mul_mat_vec_f(up->mm)) { + ggml_cuda_mul_mat_vec_f(*cuda_ctx, up->mm->src[0], up->mm->src[1], ids, cgraph->nodes[glu_idx], &fusion_data); + return glu_idx - i; } - if (!ggml_cuda_should_fuse_mul_mat_vec_q(lane.mm)) { - return 0; + if (ggml_cuda_should_fuse_mul_mat_vec_q(up->mm)) { + ggml_cuda_mul_mat_vec_q(*cuda_ctx, up->mm->src[0], up->mm->src[1], ids, cgraph->nodes[glu_idx], &fusion_data); + return glu_idx - i; } - 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], lane.mm->src[2], lane.out, &fusion_data); - return lane.n_nodes - 1; + return 0; } static int ggml_cuda_try_fuse_mm_scale(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { - ggml_cuda_mmid_lane lane; - if (!ggml_cuda_parse_mm_scale_lane(cgraph, i, lane) || lane.scale == nullptr) { + ggml_cuda_mm_lane lane; + if (!ggml_cuda_parse_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)) { + if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, lane.n_nodes, out_nodes, 1)) { return 0; } @@ -4277,7 +4056,8 @@ static int ggml_cuda_try_fuse_mm_scale(ggml_backend_cuda_context * cuda_ctx, ggm 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); + const ggml_tensor * ids = lane.mm->op == GGML_OP_MUL_MAT_ID ? lane.mm->src[2] : nullptr; + ggml_cuda_mul_mat_vec_q(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.out, &fusion_data); return lane.n_nodes - 1; } @@ -4336,9 +4116,8 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph weights = cgraph->nodes[i + ops.size() - 1]; out_nodes[1] = i + ops.size() - 1; - if (ggml_can_fuse_subgraph(cgraph, i, ops.size(), ops.data(), out_nodes, 2) && - ggml_cuda_should_use_topk_moe(node, logits, weights, ids) && - ggml_cuda_check_fusion_memory_ranges(cgraph, i, ops.size(), out_nodes, 2, /*is_topk_moe=*/true)) { + if (ggml_cuda_can_fuse_subgraph(cgraph, i, (int) ops.size(), ops.data(), out_nodes, 2, /*is_topk_moe=*/true) && + ggml_cuda_should_use_topk_moe(node, logits, weights, ids)) { ggml_cuda_op_topk_moe(*cuda_ctx, logits, weights, ids, clamp, scale, bias, args); return ops.size() - 1; } @@ -4351,9 +4130,8 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph const ggml_tensor * softmax = cgraph->nodes[i + 4]; int out_nodes[2] = { i + 1, i + 5 }; - if (ggml_can_fuse_subgraph(cgraph, i, ops.size(), ops.data(), out_nodes, 2) && - ggml_cuda_should_use_topk_moe(softmax, logits, weights, ids) && - ggml_cuda_check_fusion_memory_ranges(cgraph, i, ops.size(), out_nodes, 2, /*is_topk_moe=*/true)) { + if (ggml_cuda_can_fuse_subgraph(cgraph, i, (int) ops.size(), ops.data(), out_nodes, 2, /*is_topk_moe=*/true) && + ggml_cuda_should_use_topk_moe(softmax, logits, weights, ids)) { ggml_cuda_op_topk_moe(*cuda_ctx, logits, weights, ids, clamp, scale, bias, args); return ops.size() - 1; } @@ -4450,11 +4228,14 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph } } - int fused_scale_nodes = ggml_cuda_try_fuse_mmid_scale_glu(cuda_ctx, cgraph, i); - if (fused_scale_nodes > 0) { - return fused_scale_nodes; + // Two-lane MM GLU fusion. Each lane is MUL_MAT[_ID] plus optional bias and optional quantized post-scale. + int fused_mm_glu_nodes = ggml_cuda_try_fuse_mm_glu(cuda_ctx, cgraph, i); + if (fused_mm_glu_nodes > 0) { + return fused_mm_glu_nodes; } - fused_scale_nodes = ggml_cuda_try_fuse_mm_scale_glu(cuda_ctx, cgraph, i); + + // Single-lane quantized MM post-scale fusion; + int fused_scale_nodes = ggml_cuda_try_fuse_mm_scale(cuda_ctx, cgraph, i); if (fused_scale_nodes > 0) { return fused_scale_nodes; } @@ -4462,143 +4243,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph bool fused_mul_mat_vec = false; int fused_node_count = 0; - // gate + glu + up - for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) { - const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID; - - if (ggml_cuda_can_fuse(cgraph, i, { op, bias_op, op, bias_op, GGML_OP_GLU }, {})) { - ggml_tensor * glu = cgraph->nodes[i + 4]; - ggml_tensor * gate_bias_n = glu->src[0]; - ggml_tensor * up_bias_n = glu->src[1]; - - //we don't assume the order for {gate, up}. Instead infer it from the bias tensor - ggml_tensor * gate_n = nullptr; - ggml_tensor * up_n = nullptr; - - if (gate_bias_n->src[0] == cgraph->nodes[i] || gate_bias_n->src[1] == cgraph->nodes[i]) { - gate_n = cgraph->nodes[i]; - up_n = cgraph->nodes[i + 2]; - } else if (gate_bias_n->src[0] == cgraph->nodes[i + 2] || gate_bias_n->src[1] == cgraph->nodes[i + 2]) { - gate_n = cgraph->nodes[i + 2]; - up_n = cgraph->nodes[i]; - } else { - continue; - } - - auto get_bias_tensor = [](const ggml_tensor * bias_node, const ggml_tensor * mul_node, ggml_op op_bias) { - if (op_bias == GGML_OP_ADD) { - if (bias_node->src[0] == mul_node) { - return bias_node->src[1]; - } - if (bias_node->src[1] == mul_node) { - return bias_node->src[0]; - } - return (ggml_tensor *) nullptr; - } - GGML_ASSERT(op_bias == GGML_OP_ADD_ID); - GGML_ASSERT(bias_node->src[0] == mul_node); - return bias_node->src[1]; - }; - - ggml_tensor * up_bias_tensor = get_bias_tensor(up_bias_n, up_n, bias_op); - ggml_tensor * gate_bias_tensor = get_bias_tensor(gate_bias_n, gate_n, bias_op); - - if (!up_bias_tensor || !gate_bias_tensor) { - continue; - } - - // we don't support repeating adds - if (bias_op == GGML_OP_ADD && (!ggml_are_same_shape(gate_bias_n->src[0], gate_bias_n->src[1]) || - !ggml_are_same_shape(up_bias_n->src[0], up_bias_n->src[1]))) { - continue; - } - - const ggml_tensor * src0 = up_n->src[0]; - const ggml_tensor * src1 = up_n->src[1]; - const ggml_tensor * ids = up_n->src[2]; - - if (ggml_cuda_should_fuse_mul_mat_vec_f(up_n)) { - ggml_cuda_mm_fusion_args_host fusion_data{}; - fusion_data.gate = gate_n->src[0]; - fusion_data.x_bias = up_bias_tensor; - fusion_data.gate_bias = gate_bias_tensor; - fusion_data.glu_op = ggml_get_glu_op(glu); - - ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, glu, &fusion_data); - fused_mul_mat_vec = true; - fused_node_count = 5; - break; - } - - if (ggml_cuda_should_fuse_mul_mat_vec_q(up_n)) { - ggml_cuda_mm_fusion_args_host fusion_data{}; - fusion_data.gate = gate_n->src[0]; - fusion_data.x_bias = up_bias_tensor; - fusion_data.gate_bias = gate_bias_tensor; - fusion_data.glu_op = ggml_get_glu_op(glu); - - ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, glu, &fusion_data); - fused_mul_mat_vec = true; - fused_node_count = 5; - break; - } - } else if (ggml_cuda_can_fuse(cgraph, i, { op, op, GGML_OP_GLU }, {})) { - ggml_tensor * glu = cgraph->nodes[i + 2]; - ggml_tensor * gate = glu->src[0]; - ggml_tensor * up = glu->src[1]; - - bool ok = (gate == cgraph->nodes[i] && up == cgraph->nodes[i + 1]) || - (gate == cgraph->nodes[i + 1] && up == cgraph->nodes[i]); - - if (!ok) { - continue; - } - - const ggml_tensor * src0 = up->src[0]; - const ggml_tensor * src1 = up->src[1]; - const ggml_tensor * ids = up->src[2]; - - if (ggml_cuda_should_fuse_mul_mat_vec_f(up)) { - ggml_cuda_mm_fusion_args_host fusion_data{}; - fusion_data.gate = gate->src[0]; - fusion_data.glu_op = ggml_get_glu_op(glu); - - ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, glu, &fusion_data); - fused_mul_mat_vec = true; - fused_node_count = 3; - break; - } - - if (ggml_cuda_should_fuse_mul_mat_vec_q(up)) { - ggml_cuda_mm_fusion_args_host fusion_data{}; - fusion_data.gate = gate->src[0]; - fusion_data.glu_op = ggml_get_glu_op(glu); - - ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, glu, &fusion_data); - fused_mul_mat_vec = true; - fused_node_count = 3; - break; - } - } - } - - if (fused_mul_mat_vec) { - return fused_node_count - 1; - } - - fused_mul_mat_vec = false; - fused_node_count = 0; - - fused_scale_nodes = ggml_cuda_try_fuse_mmid_scale(cuda_ctx, cgraph, i); - if (fused_scale_nodes > 0) { - return fused_scale_nodes; - } - fused_scale_nodes = ggml_cuda_try_fuse_mm_scale(cuda_ctx, cgraph, i); - if (fused_scale_nodes > 0) { - return fused_scale_nodes; - } - - // gate + add + glu + up + add + // mul_mat + add for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) { const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID; @@ -4845,7 +4490,7 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud #ifdef GGML_CUDA_DEBUG const int last_fused = i + nodes_to_skip; GGML_LOG_INFO("nodes_fused: %d, first: %s (%s), last: %s (%s)\n", - nodes_to_skip, ggml_op_name(node->op), node->name, + nodes_to_skip + 1, ggml_op_name(node->op), node->name, ggml_op_name(cgraph->nodes[last_fused]->op), cgraph->nodes[last_fused]->name); #endif i += nodes_to_skip; diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 84e6b8e9fca5..445575b62a5d 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5789,20 +5789,21 @@ struct test_mul_mat_vec_fusion : public test_case { const bool with_bias; const bool with_gate; const bool with_lane_scale; + const bool gate_first; std::array batch_dims; test_mul_mat_vec_fusion(ggml_type type, ggml_glu_op op, int64_t m, int64_t n, int64_t k, bool use_id = false, int n_mats = 1, int n_used = 1, bool b = false, bool with_bias = false, bool with_gate = true, - bool with_lane_scale = false, std::array batch_dims = {4, 2}) + bool with_lane_scale = false, bool gate_first = false, std::array batch_dims = {4, 2}) : type(type), glu_op(op), m(m), n(n), k(k), use_id(use_id), n_mats(n_mats), n_used(n_used), b(b), with_bias(with_bias), - with_gate(with_gate), with_lane_scale(with_lane_scale), batch_dims(batch_dims) { + with_gate(with_gate), with_lane_scale(with_lane_scale), gate_first(gate_first), batch_dims(batch_dims) { if (use_id) { GGML_ASSERT(n_used <= n_mats); } } std::string vars() override { - return VARS_TO_STR13(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_lane_scale, batch_dims); + return VARS_TO_STR14(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_lane_scale, gate_first, batch_dims); } std::string op_desc(ggml_tensor * t) override { @@ -5850,24 +5851,40 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * gate = with_gate ? ggml_new_tensor(ctx, type, 4, ne0.data()) : nullptr; ggml_tensor * up = ggml_new_tensor(ctx, type, 4, ne0.data()); - ggml_tensor * ffn_up = ggml_mul_mat(ctx, up, cur); - if (with_bias) { - std::array bias_ne = { ffn_up->ne[0], 1, channels, samples }; - ggml_tensor * up_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); - ffn_up = ggml_add(ctx, ffn_up, up_bias); - } - if (with_lane_scale) { - ffn_up = build_dense_lane_scale(ctx, ffn_up); - } + auto build_up_lane = [&]() { + ggml_tensor * ffn_up = ggml_mul_mat(ctx, up, cur); + if (with_bias) { + std::array bias_ne = { ffn_up->ne[0], 1, channels, samples }; + ggml_tensor * up_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); + ffn_up = ggml_add(ctx, ffn_up, up_bias); + } + if (with_lane_scale) { + ffn_up = build_dense_lane_scale(ctx, ffn_up); + } + return ffn_up; + }; - ggml_tensor * ffn_gate = with_gate ? ggml_mul_mat(ctx, gate, cur) : nullptr; - if (with_bias && with_gate) { - std::array bias_ne = { ffn_gate->ne[0], 1, channels, samples }; - ggml_tensor * gate_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); - ffn_gate = ggml_add(ctx, ffn_gate, gate_bias); - } - if (with_lane_scale && with_gate) { - ffn_gate = build_dense_lane_scale(ctx, ffn_gate); + auto build_gate_lane = [&]() { + ggml_tensor * ffn_gate = ggml_mul_mat(ctx, gate, cur); + if (with_bias) { + std::array bias_ne = { ffn_gate->ne[0], 1, channels, samples }; + ggml_tensor * gate_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); + ffn_gate = ggml_add(ctx, ffn_gate, gate_bias); + } + if (with_lane_scale) { + ffn_gate = build_dense_lane_scale(ctx, ffn_gate); + } + return ffn_gate; + }; + + ggml_tensor * ffn_up = nullptr; + ggml_tensor * ffn_gate = nullptr; + if (with_gate && gate_first) { + ffn_gate = build_gate_lane(); + ffn_up = build_up_lane(); + } else { + ffn_up = build_up_lane(); + ffn_gate = with_gate ? build_gate_lane() : nullptr; } ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; @@ -5890,22 +5907,38 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * cur = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, k, this->b ? 1 : n_used, m); ggml_set_name(cur, "cur"); - ggml_tensor * ffn_up = ggml_mul_mat_id(ctx, ups, cur, ids); - if (with_bias) { - ggml_tensor * up_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_up->ne[0], n_mats); - ffn_up = ggml_add_id(ctx, ffn_up, up_bias_param, ids); - } - if (with_lane_scale) { - ffn_up = build_id_lane_scale(ctx, ffn_up, ids); - } + auto build_up_lane = [&]() { + ggml_tensor * ffn_up = ggml_mul_mat_id(ctx, ups, cur, ids); + if (with_bias) { + ggml_tensor * up_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_up->ne[0], n_mats); + ffn_up = ggml_add_id(ctx, ffn_up, up_bias_param, ids); + } + if (with_lane_scale) { + ffn_up = build_id_lane_scale(ctx, ffn_up, ids); + } + return ffn_up; + }; - ggml_tensor * ffn_gate = with_gate? ggml_mul_mat_id(ctx, gates, cur, ids) : nullptr; - if (with_bias && with_gate) { - ggml_tensor * gate_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_gate->ne[0], n_mats); - ffn_gate = ggml_add_id(ctx, ffn_gate, gate_bias_param, ids); - } - if (with_lane_scale && with_gate) { - ffn_gate = build_id_lane_scale(ctx, ffn_gate, ids); + auto build_gate_lane = [&]() { + ggml_tensor * ffn_gate = ggml_mul_mat_id(ctx, gates, cur, ids); + if (with_bias) { + ggml_tensor * gate_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_gate->ne[0], n_mats); + ffn_gate = ggml_add_id(ctx, ffn_gate, gate_bias_param, ids); + } + if (with_lane_scale) { + ffn_gate = build_id_lane_scale(ctx, ffn_gate, ids); + } + return ffn_gate; + }; + + ggml_tensor * ffn_up = nullptr; + ggml_tensor * ffn_gate = nullptr; + if (with_gate && gate_first) { + ffn_gate = build_gate_lane(); + ffn_up = build_up_lane(); + } else { + ffn_up = build_up_lane(); + ffn_gate = with_gate ? build_gate_lane() : nullptr; } ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; @@ -9111,10 +9144,15 @@ static std::vector> make_test_cases_eval() { continue; } for (bool with_lane_scale : {false, true}) { - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, - use_id, 16, 8, b, with_bias, with_gate, with_lane_scale)); - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, - use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, {1, 1})); + for (bool gate_first : {false, true}) { + if (!with_gate && gate_first) { + continue; + } + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, + use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, gate_first)); + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, + use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, gate_first, {1, 1})); + } } } } From cf97e37d816f665d83cf41d86d506fe19400b94e Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Wed, 3 Jun 2026 09:34:00 +0200 Subject: [PATCH 07/28] Enable fusion for MoE in shared MMVQ --- ggml/src/ggml-cuda/ggml-cuda.cu | 42 +++++++++++++++++++++++++++++---- 1 file changed, 37 insertions(+), 5 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index fd0b12cbe95f..0dde6118f7b4 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3591,12 +3591,44 @@ static bool ggml_cuda_can_fuse_parsed_subgraph(const struct ggml_cgraph * cgraph return false; } - std::vector ops; - ops.reserve(count); - for (int j = 0; j < count; ++j) { - ops.push_back(cgraph->nodes[node_idx + j]->op); + const auto is_output = [&](int idx) { + for (int j = 0; j < out_count; ++j) { + if (out_nodes[j] == idx) { + return true; + } + } + return false; + }; + + // Parsed MM lanes may contain RESHAPE views of external scale tensors. The + // parser validates those scale tensors, so only require closure by use count. + for (int j = node_idx; j < node_idx + count; ++j) { + const ggml_tensor * node = cgraph->nodes[j]; + if ((node->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) { + return false; + } + if (is_output(j)) { + continue; + } + if (node->flags & GGML_TENSOR_FLAG_OUTPUT) { + return false; + } + + int subgraph_uses = 0; + for (int k = j + 1; k < node_idx + count; ++k) { + const ggml_tensor * other = cgraph->nodes[k]; + for (int src_idx = 0; src_idx < GGML_MAX_SRC; ++src_idx) { + if (other->src[src_idx] == node) { + subgraph_uses++; + } + } + } + if (subgraph_uses != ggml_node_get_use_count(cgraph, j)) { + return false; + } } - return ggml_cuda_can_fuse_subgraph(cgraph, node_idx, count, ops.data(), out_nodes, out_count, is_topk_moe); + + return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, count, out_nodes, out_count, is_topk_moe); } static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, From c1b9381d327e063cc846b46b59708444b66dc4d8 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Wed, 3 Jun 2026 10:47:40 +0200 Subject: [PATCH 08/28] Restrict scale_view_nodes, enroll MM + ADD into lane-matcher --- ggml/src/ggml-cuda/ggml-cuda.cu | 172 +++++++++++++++++++------------- 1 file changed, 103 insertions(+), 69 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 0dde6118f7b4..6c14c462f38d 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3586,6 +3586,8 @@ static bool ggml_cuda_can_fuse_parsed_subgraph(const struct ggml_cgraph * cgraph int count, const int * out_nodes, int out_count, + const int * external_view_nodes = nullptr, + int external_view_count = 0, bool is_topk_moe = false) { if (node_idx + count > cgraph->n_nodes) { return false; @@ -3600,8 +3602,24 @@ static bool ggml_cuda_can_fuse_parsed_subgraph(const struct ggml_cgraph * cgraph return false; }; - // Parsed MM lanes may contain RESHAPE views of external scale tensors. The - // parser validates those scale tensors, so only require closure by use count. + const auto is_in_subgraph = [&](const ggml_tensor * tensor) { + for (int j = node_idx; j < node_idx + count; ++j) { + if (cgraph->nodes[j] == tensor) { + return true; + } + } + return false; + }; + + const auto is_allowed_external_view = [&](int idx) { + for (int j = 0; j < external_view_count; ++j) { + if (external_view_nodes[j] == idx) { + return true; + } + } + return false; + }; + for (int j = node_idx; j < node_idx + count; ++j) { const ggml_tensor * node = cgraph->nodes[j]; if ((node->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) { @@ -3626,6 +3644,12 @@ static bool ggml_cuda_can_fuse_parsed_subgraph(const struct ggml_cgraph * cgraph if (subgraph_uses != ggml_node_get_use_count(cgraph, j)) { return false; } + + for (const ggml_tensor * view_src = node->view_src; view_src != nullptr; view_src = view_src->view_src) { + if (!is_in_subgraph(view_src) && !is_allowed_external_view(j)) { + return false; + } + } } return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, count, out_nodes, out_count, is_topk_moe); @@ -3822,12 +3846,16 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, } +// Matched MM lane forms: +// MUL_MAT [ADD] [MUL scalar_scale] +// MUL_MAT_ID [ADD_ID] [RESHAPE -> REPEAT -> GET_ROWS -> MUL expert_scale] struct ggml_cuda_mm_lane { ggml_tensor * mm = nullptr; ggml_tensor * bias_node = nullptr; ggml_tensor * out = nullptr; const ggml_tensor * bias = nullptr; const ggml_tensor * scale = nullptr; + int scale_view_node = -1; int n_nodes = 0; }; @@ -3835,6 +3863,18 @@ 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_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; + } + + if (mm->op == GGML_OP_MUL_MAT_ID && bias->ne[1] != mm->src[0]->ne[2]) { + return false; + } + + return true; +} + static bool ggml_cuda_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mm_lane & lane) { if (i >= cgraph->n_nodes || cgraph->nodes[i]->op != GGML_OP_MUL_MAT_ID) { return false; @@ -3852,7 +3892,10 @@ static bool ggml_cuda_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, g if (i + lane.n_nodes < cgraph->n_nodes && cgraph->nodes[i + lane.n_nodes]->op == GGML_OP_ADD_ID) { ggml_tensor * add = cgraph->nodes[i + lane.n_nodes]; - if (add->src[0] != lane.out || add->src[2] != mm->src[2]) { + if (add->src[0] != lane.out || add->src[2] != mm->src[2] || add->type != GGML_TYPE_F32) { + return false; + } + if (!ggml_cuda_can_parse_mm_lane_bias(mm, add->src[1])) { return false; } lane.bias_node = add; @@ -3895,6 +3938,7 @@ static bool ggml_cuda_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, g } lane.scale = scale; + lane.scale_view_node = i + lane.n_nodes; lane.out = mul; lane.n_nodes += 4; return true; @@ -3924,7 +3968,10 @@ static bool ggml_cuda_parse_mul_mat_lane(const ggml_cgraph * cgraph, int i, ggml } else { return false; } - if (!ggml_are_same_shape(add->src[0], add->src[1])) { + if (add->type != GGML_TYPE_F32 || !ggml_are_same_shape(add->src[0], add->src[1])) { + return false; + } + if (!ggml_cuda_can_parse_mm_lane_bias(mm, lane.bias)) { return false; } lane.bias_node = add; @@ -3971,6 +4018,13 @@ static bool ggml_cuda_parse_mm_lane(const ggml_cgraph * cgraph, int i, ggml_cuda return false; } +static int ggml_cuda_add_mm_lane_external_view_node(const ggml_cuda_mm_lane & lane, int * external_view_nodes, int n_external_view_nodes) { + if (lane.scale_view_node != -1) { + external_view_nodes[n_external_view_nodes++] = lane.scale_view_node; + } + return n_external_view_nodes; +} + static bool ggml_cuda_should_fuse_mm_lanes(const ggml_cuda_mm_lane & up, const ggml_cuda_mm_lane & gate, const ggml_tensor * glu) { if (up.mm->op != gate.mm->op) { return false; @@ -4042,7 +4096,11 @@ static int ggml_cuda_try_fuse_mm_glu(ggml_backend_cuda_context * cuda_ctx, ggml_ const int out_nodes[] = { glu_idx }; const int n_nodes = glu_idx - i + 1; - if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, n_nodes, out_nodes, 1)) { + int external_view_nodes[2]; + int n_external_view_nodes = 0; + n_external_view_nodes = ggml_cuda_add_mm_lane_external_view_node(*up, external_view_nodes, n_external_view_nodes); + n_external_view_nodes = ggml_cuda_add_mm_lane_external_view_node(*gate, external_view_nodes, n_external_view_nodes); + if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, n_nodes, out_nodes, 1, external_view_nodes, n_external_view_nodes)) { return 0; } @@ -4076,7 +4134,9 @@ static int ggml_cuda_try_fuse_mm_scale(ggml_backend_cuda_context * cuda_ctx, ggm const int out_idx = i + lane.n_nodes - 1; const int out_nodes[] = { out_idx }; - if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, lane.n_nodes, out_nodes, 1)) { + int external_view_nodes[1]; + int n_external_view_nodes = ggml_cuda_add_mm_lane_external_view_node(lane, external_view_nodes, 0); + if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, lane.n_nodes, out_nodes, 1, external_view_nodes, n_external_view_nodes)) { return 0; } @@ -4093,6 +4153,39 @@ static int ggml_cuda_try_fuse_mm_scale(ggml_backend_cuda_context * cuda_ctx, ggm return lane.n_nodes - 1; } +static int ggml_cuda_try_fuse_mm_bias(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { + ggml_cuda_mm_lane lane; + if (!ggml_cuda_parse_mm_lane(cgraph, i, lane) || lane.bias == nullptr) { + return 0; + } + + const int bias_idx = i + 1; + if (bias_idx >= cgraph->n_nodes || cgraph->nodes[bias_idx] != lane.bias_node) { + return 0; + } + + const int out_nodes[] = { bias_idx }; + if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, 2, out_nodes, 1)) { + return 0; + } + + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.x_bias = lane.bias; + + const ggml_tensor * ids = lane.mm->op == GGML_OP_MUL_MAT_ID ? lane.mm->src[2] : nullptr; + if (ggml_cuda_should_fuse_mul_mat_vec_f(lane.mm)) { + ggml_cuda_mul_mat_vec_f(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.bias_node, &fusion_data); + return 1; + } + + if (ggml_cuda_should_fuse_mul_mat_vec_q(lane.mm)) { + ggml_cuda_mul_mat_vec_q(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.bias_node, &fusion_data); + return 1; + } + + return 0; +} + // 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) { @@ -4266,74 +4359,15 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph return fused_mm_glu_nodes; } - // Single-lane quantized MM post-scale fusion; + // Single-lane quantized MM post-scale fusion. int fused_scale_nodes = ggml_cuda_try_fuse_mm_scale(cuda_ctx, cgraph, i); if (fused_scale_nodes > 0) { return fused_scale_nodes; } - bool fused_mul_mat_vec = false; - int fused_node_count = 0; - - // mul_mat + add - for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) { - const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID; - - if (!ggml_can_fuse(cgraph, i, { op, bias_op })) { - continue; - } - - ggml_tensor * mm_node = cgraph->nodes[i]; - ggml_tensor * bias_node = cgraph->nodes[i + 1]; - - ggml_tensor * bias_tensor = nullptr; - if (bias_op == GGML_OP_ADD) { - if (bias_node->src[0] == mm_node) { - bias_tensor = bias_node->src[1]; - } else if (bias_node->src[1] == mm_node) { - bias_tensor = bias_node->src[0]; - } else { - continue; - } - } else { - if (bias_node->src[0] != mm_node) { - continue; - } - bias_tensor = bias_node->src[1]; - } - - const ggml_tensor * src0 = mm_node->src[0]; - const ggml_tensor * src1 = mm_node->src[1]; - const ggml_tensor * ids = mm_node->src[2]; - - if (bias_op == GGML_OP_ADD_ID && bias_node->src[2] != ids) { - continue; - } - - if (bias_op == GGML_OP_ADD && !ggml_are_same_shape(bias_node->src[0], bias_node->src[1])) { - continue; - } - - ggml_cuda_mm_fusion_args_host fusion_data{}; - fusion_data.x_bias = bias_tensor; - - if (ggml_cuda_should_fuse_mul_mat_vec_f(mm_node)) { - ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, bias_node, &fusion_data); - fused_mul_mat_vec = true; - fused_node_count = 2; - break; - } - - if (ggml_cuda_should_fuse_mul_mat_vec_q(mm_node)) { - ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, bias_node, &fusion_data); - fused_mul_mat_vec = true; - fused_node_count = 2; - break; - } - } - - if (fused_mul_mat_vec) { - return fused_node_count - 1; + int fused_bias_nodes = ggml_cuda_try_fuse_mm_bias(cuda_ctx, cgraph, i); + if (fused_bias_nodes > 0) { + return fused_bias_nodes; } if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ADD }, {})) { From 29713604f1603b445239aa2cb488e984d77796f3 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Wed, 3 Jun 2026 15:53:05 +0200 Subject: [PATCH 09/28] Refactor mmvq loads, still does not help non-nvfp4 kernels --- ggml/src/ggml-cuda/mmvq.cu | 22 ++++++++++------------ 1 file changed, 10 insertions(+), 12 deletions(-) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 5a4b2fdf4233..9cb4c313f6e1 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -548,13 +548,14 @@ static __global__ void mul_mat_vec_q( [[maybe_unused]] float x_scales; [[maybe_unused]] float gate_scales; 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 const uint32_t channel_bias = ids ? channel_x : channel_dst; - if (use_bias) { - x_bias = x_bias + sample_dst*stride_sample_dst + channel_bias*stride_channel_dst + row0; - // 1. Hide latency by prefetching bias and gate here - // 2. load only on threads that won't die after partial sum calculation - if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 && - (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) { + if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 && + (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) { + const uint32_t channel_bias = ids ? channel_x : channel_dst; + if (use_bias) { + x_bias = x_bias + sample_dst * stride_sample_dst + channel_bias * stride_channel_dst + row0; #pragma unroll for (int j = 0; j < ncols_dst; ++j) { x_biases[j] = x_bias[j * stride_col_dst + threadIdx.x]; @@ -562,13 +563,10 @@ static __global__ void mul_mat_vec_q( } } if (use_gate_bias) { - gate_bias = gate_bias + sample_dst*stride_sample_dst + channel_bias*stride_channel_dst + row0; - if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 && - (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) { + gate_bias = gate_bias + sample_dst * stride_sample_dst + channel_bias * stride_channel_dst + row0; #pragma unroll - for (int j = 0; j < ncols_dst; ++j) { - gate_biases[j] = gate_bias[j * stride_col_dst + threadIdx.x]; - } + for (int j = 0; j < ncols_dst; ++j) { + gate_biases[j] = gate_bias[j * stride_col_dst + threadIdx.x]; } } if (use_scale) { From 28bde658e9a3ab12fad983efe2956f5a38703799 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Wed, 10 Jun 2026 16:58:39 +0200 Subject: [PATCH 10/28] 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 6c14c462f38d..6d77f64d9917 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 9cb4c313f6e1..9c07ada06dc6 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); From 9d9f1a0086ab84e9c7c2f4c20356d9c9901f47a2 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Wed, 10 Jun 2026 17:26:03 +0200 Subject: [PATCH 11/28] Reorder scale & bias-add to adhere to #24331 --- ggml/src/ggml-cuda/ggml-cuda.cu | 122 +++++++++++++++----------------- ggml/src/ggml-cuda/mmvq.cu | 12 ++-- tests/test-backend-ops.cpp | 24 +++---- 3 files changed, 77 insertions(+), 81 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 6d77f64d9917..36ef22a29cb5 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3847,8 +3847,8 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, // Matched MM lane forms: -// MUL_MAT [ADD] [MUL scalar_scale] -// MUL_MAT_ID [ADD_ID] [RESHAPE -> REPEAT -> GET_ROWS -> MUL expert_scale] +// MUL_MAT [MUL scalar_scale] [ADD] +// MUL_MAT_ID [RESHAPE -> REPEAT -> GET_ROWS -> MUL expert_scale] [ADD_ID] struct ggml_cuda_mm_lane { ggml_tensor * mm = nullptr; ggml_tensor * bias_node = nullptr; @@ -3894,6 +3894,40 @@ static bool ggml_cuda_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, g lane.out = mm; lane.n_nodes = 1; + if (ggml_is_quantized(mm->src[0]->type) && ggml_cuda_can_fuse_mm_lane_scale(mm) && i + lane.n_nodes + 3 < cgraph->n_nodes) { + ggml_tensor * reshape = cgraph->nodes[i + lane.n_nodes + 0]; + ggml_tensor * repeat = cgraph->nodes[i + lane.n_nodes + 1]; + ggml_tensor * getrows = cgraph->nodes[i + lane.n_nodes + 2]; + ggml_tensor * mul = cgraph->nodes[i + lane.n_nodes + 3]; + + if (reshape->op == GGML_OP_RESHAPE && repeat->op == GGML_OP_REPEAT && + getrows->op == GGML_OP_GET_ROWS && mul->op == GGML_OP_MUL) { + if (repeat->src[0] != reshape || getrows->src[0] != repeat || getrows->src[1] != mm->src[2]) { + return false; + } + + const bool mul_has_out = mul->src[0] == lane.out || mul->src[1] == lane.out; + const bool mul_has_scale = mul->src[0] == getrows || mul->src[1] == getrows; + if (!mul_has_out || !mul_has_scale) { + return false; + } + + const ggml_tensor * scale = reshape->src[0]; + if (scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != mm->src[0]->ne[2]) { + return false; + } + + if (mul->type != GGML_TYPE_F32 || !ggml_are_same_shape(mul, lane.out)) { + return false; + } + + lane.scale = scale; + lane.scale_view_node = i + lane.n_nodes; + lane.out = mul; + lane.n_nodes += 4; + } + } + if (i + lane.n_nodes < cgraph->n_nodes && cgraph->nodes[i + lane.n_nodes]->op == GGML_OP_ADD_ID) { ggml_tensor * add = cgraph->nodes[i + lane.n_nodes]; if (add->src[0] != lane.out || add->src[2] != mm->src[2] || add->type != GGML_TYPE_F32) { @@ -3908,43 +3942,6 @@ 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) || !ggml_cuda_can_fuse_mm_lane_scale(mm) || i + lane.n_nodes + 3 >= cgraph->n_nodes) { - return true; - } - - ggml_tensor * reshape = cgraph->nodes[i + lane.n_nodes + 0]; - ggml_tensor * repeat = cgraph->nodes[i + lane.n_nodes + 1]; - ggml_tensor * getrows = cgraph->nodes[i + lane.n_nodes + 2]; - ggml_tensor * mul = cgraph->nodes[i + lane.n_nodes + 3]; - - if (reshape->op != GGML_OP_RESHAPE || repeat->op != GGML_OP_REPEAT || - getrows->op != GGML_OP_GET_ROWS || mul->op != GGML_OP_MUL) { - return true; - } - - if (repeat->src[0] != reshape || getrows->src[0] != repeat || getrows->src[1] != mm->src[2]) { - return true; - } - - const bool mul_has_out = mul->src[0] == lane.out || mul->src[1] == lane.out; - const bool mul_has_scale = mul->src[0] == getrows || mul->src[1] == getrows; - if (!mul_has_out || !mul_has_scale) { - return true; - } - - const ggml_tensor * scale = reshape->src[0]; - if (scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != mm->src[0]->ne[2]) { - return false; - } - - if (mul->type != GGML_TYPE_F32 || !ggml_are_same_shape(mul, lane.out)) { - return false; - } - - lane.scale = scale; - lane.scale_view_node = i + lane.n_nodes; - lane.out = mul; - lane.n_nodes += 4; return true; } @@ -3963,6 +3960,29 @@ static bool ggml_cuda_parse_mul_mat_lane(const ggml_cgraph * cgraph, int i, ggml lane.out = mm; lane.n_nodes = 1; + 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) { + 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 false; + } + + 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++; + } + 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) { @@ -3983,30 +4003,6 @@ 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) || !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; - } - - 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; } diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 9c07ada06dc6..c2d06cbf48b5 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -659,24 +659,24 @@ static __global__ void mul_mat_vec_q( 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 (use_bias) { - result += x_biases[j]; - } if constexpr (type == GGML_TYPE_NVFP4) { if (use_scale) { result *= x_scales; } } + if (use_bias) { + result += x_biases[j]; + } if (use_gate) { float gate_value = tmp_gate[j][threadIdx.x]; - if (use_gate_bias) { - gate_value += gate_biases[j]; - } if constexpr (type == GGML_TYPE_NVFP4) { if (use_gate_scale) { 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); diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 445575b62a5d..c770822776b9 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5853,27 +5853,27 @@ struct test_mul_mat_vec_fusion : public test_case { auto build_up_lane = [&]() { ggml_tensor * ffn_up = ggml_mul_mat(ctx, up, cur); + if (with_lane_scale) { + ffn_up = build_dense_lane_scale(ctx, ffn_up); + } if (with_bias) { std::array bias_ne = { ffn_up->ne[0], 1, channels, samples }; ggml_tensor * up_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); ffn_up = ggml_add(ctx, ffn_up, up_bias); } - if (with_lane_scale) { - ffn_up = build_dense_lane_scale(ctx, ffn_up); - } return ffn_up; }; auto build_gate_lane = [&]() { ggml_tensor * ffn_gate = ggml_mul_mat(ctx, gate, cur); + if (with_lane_scale) { + ffn_gate = build_dense_lane_scale(ctx, ffn_gate); + } if (with_bias) { std::array bias_ne = { ffn_gate->ne[0], 1, channels, samples }; ggml_tensor * gate_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data()); ffn_gate = ggml_add(ctx, ffn_gate, gate_bias); } - if (with_lane_scale) { - ffn_gate = build_dense_lane_scale(ctx, ffn_gate); - } return ffn_gate; }; @@ -5909,25 +5909,25 @@ struct test_mul_mat_vec_fusion : public test_case { auto build_up_lane = [&]() { ggml_tensor * ffn_up = ggml_mul_mat_id(ctx, ups, cur, ids); + if (with_lane_scale) { + ffn_up = build_id_lane_scale(ctx, ffn_up, ids); + } if (with_bias) { ggml_tensor * up_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_up->ne[0], n_mats); ffn_up = ggml_add_id(ctx, ffn_up, up_bias_param, ids); } - if (with_lane_scale) { - ffn_up = build_id_lane_scale(ctx, ffn_up, ids); - } return ffn_up; }; auto build_gate_lane = [&]() { ggml_tensor * ffn_gate = ggml_mul_mat_id(ctx, gates, cur, ids); + if (with_lane_scale) { + ffn_gate = build_id_lane_scale(ctx, ffn_gate, ids); + } if (with_bias) { ggml_tensor * gate_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_gate->ne[0], n_mats); ffn_gate = ggml_add_id(ctx, ffn_gate, gate_bias_param, ids); } - if (with_lane_scale) { - ffn_gate = build_id_lane_scale(ctx, ffn_gate, ids); - } return ffn_gate; }; From 096597884b2cead9bce316bbba820c9cbacdc03e Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Wed, 10 Jun 2026 17:28:28 +0200 Subject: [PATCH 12/28] Restrict lane scale to NVFP4 Don't need to test unfused combinations --- tests/test-backend-ops.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index c770822776b9..9c1f16cfb68b 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9144,6 +9144,9 @@ static std::vector> make_test_cases_eval() { continue; } for (bool with_lane_scale : {false, true}) { + if (with_lane_scale && type != GGML_TYPE_NVFP4) { + continue; + } for (bool gate_first : {false, true}) { if (!with_gate && gate_first) { continue; From 9088a3c6d23d3047aef771c821ce547b63c4c6ce Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Wed, 10 Jun 2026 17:36:56 +0200 Subject: [PATCH 13/28] Cleanup --- ggml/src/ggml-cuda/mmvq.cu | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index c2d06cbf48b5..83f20b9aeb21 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -864,8 +864,6 @@ static void mul_mat_vec_q_switch_ncols_dst( const int warp_size = ggml_cuda_info().devices[device].warp_size; const mmvq_parameter_table_id table_id = get_device_table_id(cc); - const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr || - fusion.x_scale != nullptr || fusion.gate_scale != nullptr; const bool has_ids = ids != nullptr; const auto should_use_small_k = [&](int c_ncols_dst) { @@ -1004,8 +1002,6 @@ static void mul_mat_vec_q_switch_ncols_dst( GGML_ABORT("fatal error"); break; } - - GGML_UNUSED(has_fusion); } static void mul_mat_vec_q_switch_type( const void * vx, const ggml_type type_x, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst, @@ -1185,6 +1181,8 @@ void ggml_cuda_mul_mat_vec_q( if (fusion) { GGML_ASSERT( !ids || dst->ne[2] == 1); GGML_ASSERT( ids || dst->ne[1] == 1); + // Scale fusion is only allowed for NVFP4 currently as the cost of checking this at run-time in the prologue is + // non-negligible for some models such as gpt-oss-20b GGML_ASSERT((fusion->x_scale == nullptr && fusion->gate_scale == nullptr) || src0->type == GGML_TYPE_NVFP4); if (fusion->x_bias) { From c91327f44bbda6201601ec55a8165048dc1b5e1c Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Thu, 11 Jun 2026 11:30:34 +0200 Subject: [PATCH 14/28] Merge single-lane mm-fusion helpers --- ggml/src/ggml-cuda/ggml-cuda.cu | 66 ++++++++++++--------------------- 1 file changed, 23 insertions(+), 43 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 36ef22a29cb5..c1773dbcb833 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4131,9 +4131,9 @@ static int ggml_cuda_try_fuse_mm_glu(ggml_backend_cuda_context * cuda_ctx, ggml_ return 0; } -static int ggml_cuda_try_fuse_mm_scale(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { +static int ggml_cuda_try_fuse_mm_lane(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { ggml_cuda_mm_lane lane; - if (!ggml_cuda_parse_mm_lane(cgraph, i, lane) || lane.scale == nullptr) { + if (!ggml_cuda_parse_mm_lane(cgraph, i, lane) || (lane.scale == nullptr && lane.bias == nullptr)) { return 0; } @@ -4145,51 +4145,36 @@ 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_cuda_mm_fusion_args_host fusion_data{}; fusion_data.x_bias = lane.bias; fusion_data.x_scale = lane.scale; const ggml_tensor * ids = lane.mm->op == GGML_OP_MUL_MAT_ID ? lane.mm->src[2] : nullptr; - ggml_cuda_mul_mat_vec_q(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.out, &fusion_data); - return lane.n_nodes - 1; -} -static int ggml_cuda_try_fuse_mm_bias(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { - ggml_cuda_mm_lane lane; - if (!ggml_cuda_parse_mm_lane(cgraph, i, lane) || lane.bias == nullptr) { - return 0; - } + // Lane scale fusion is implemented by MMVQ only because it is limited to NVFP4. + // This path owns scale lanes, including scale followed by bias. + if (lane.scale != nullptr) { + if (!ggml_cuda_can_fuse_mm_lane_scale(lane.mm)) { + return 0; + } - const int bias_idx = i + 1; - if (bias_idx >= cgraph->n_nodes || cgraph->nodes[bias_idx] != lane.bias_node) { - return 0; - } + if (!ggml_cuda_should_fuse_mul_mat_vec_q(lane.mm)) { + return 0; + } - const int out_nodes[] = { bias_idx }; - if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, 2, out_nodes, 1)) { - return 0; + ggml_cuda_mul_mat_vec_q(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.out, &fusion_data); + return lane.n_nodes - 1; } - ggml_cuda_mm_fusion_args_host fusion_data{}; - fusion_data.x_bias = lane.bias; - - const ggml_tensor * ids = lane.mm->op == GGML_OP_MUL_MAT_ID ? lane.mm->src[2] : nullptr; + // Bias-only lanes can use either MMVF or MMVQ. if (ggml_cuda_should_fuse_mul_mat_vec_f(lane.mm)) { - ggml_cuda_mul_mat_vec_f(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.bias_node, &fusion_data); - return 1; + ggml_cuda_mul_mat_vec_f(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.out, &fusion_data); + return lane.n_nodes - 1; } if (ggml_cuda_should_fuse_mul_mat_vec_q(lane.mm)) { - ggml_cuda_mul_mat_vec_q(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.bias_node, &fusion_data); - return 1; + ggml_cuda_mul_mat_vec_q(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.out, &fusion_data); + return lane.n_nodes - 1; } return 0; @@ -4362,21 +4347,16 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph } } - // Two-lane MM GLU fusion. Each lane is MUL_MAT[_ID] plus optional bias and optional quantized post-scale. + // Two-lane MM GLU fusion. Each lane is MUL_MAT[_ID] + optional scale + optional bias int fused_mm_glu_nodes = ggml_cuda_try_fuse_mm_glu(cuda_ctx, cgraph, i); if (fused_mm_glu_nodes > 0) { return fused_mm_glu_nodes; } - // Single-lane quantized MM post-scale fusion. - int fused_scale_nodes = ggml_cuda_try_fuse_mm_scale(cuda_ctx, cgraph, i); - if (fused_scale_nodes > 0) { - return fused_scale_nodes; - } - - int fused_bias_nodes = ggml_cuda_try_fuse_mm_bias(cuda_ctx, cgraph, i); - if (fused_bias_nodes > 0) { - return fused_bias_nodes; + // Single-lane MM fusion. The lane is MUL_MAT[_ID] + optional scale + optional bias. + int fused_mm_lane_nodes = ggml_cuda_try_fuse_mm_lane(cuda_ctx, cgraph, i); + if (fused_mm_lane_nodes > 0) { + return fused_mm_lane_nodes; } if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ADD }, {})) { From 6d815ca4d18150e24c2af078d24e5cb5237799c2 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Thu, 11 Jun 2026 13:46:48 +0200 Subject: [PATCH 15/28] Refactor and clean-up host-side fusion logic --- ggml/src/ggml-cuda/ggml-cuda.cu | 62 +++++++-------------------------- 1 file changed, 13 insertions(+), 49 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index c1773dbcb833..6eb1ee0d2092 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3581,6 +3581,9 @@ static bool ggml_cuda_can_fuse_subgraph(const struct ggml_cgraph * cgraph, ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, count, out_nodes, out_count, is_topk_moe); } +// This mirrors ggml_can_fuse_subgraph for parser-validated patterns. It avoids +// rebuilding variable lane patterns as fixed op lists, and allows parser-approved +// external view sources such as the NVFP4 MUL_MAT_ID scale reshape. static bool ggml_cuda_can_fuse_parsed_subgraph(const struct ggml_cgraph * cgraph, int node_idx, int count, @@ -3867,23 +3870,7 @@ 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; - } - - if (mm->op == GGML_OP_MUL_MAT_ID && bias->ne[1] != mm->src[0]->ne[2]) { - return false; - } - - return true; -} - static bool ggml_cuda_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mm_lane & lane) { - if (i >= cgraph->n_nodes || cgraph->nodes[i]->op != GGML_OP_MUL_MAT_ID) { - return false; - } - ggml_tensor * mm = cgraph->nodes[i]; if (!ggml_cuda_can_parse_mm_lane_type(mm->src[0]->type) || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32 || mm->src[2] == nullptr) { return false; @@ -3933,11 +3920,12 @@ static bool ggml_cuda_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, g if (add->src[0] != lane.out || add->src[2] != mm->src[2] || add->type != GGML_TYPE_F32) { return false; } - if (!ggml_cuda_can_parse_mm_lane_bias(mm, add->src[1])) { + const ggml_tensor * bias = add->src[1]; + if (bias->type != GGML_TYPE_F32 || bias->ne[0] != mm->ne[0] || bias->ne[1] != mm->src[0]->ne[2]) { return false; } lane.bias_node = add; - lane.bias = add->src[1]; + lane.bias = bias; lane.out = add; lane.n_nodes++; } @@ -3946,10 +3934,6 @@ static bool ggml_cuda_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, g } static bool ggml_cuda_parse_mul_mat_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mm_lane & lane) { - if (i >= cgraph->n_nodes || cgraph->nodes[i]->op != GGML_OP_MUL_MAT) { - return false; - } - ggml_tensor * mm = cgraph->nodes[i]; if (!ggml_cuda_can_parse_mm_lane_type(mm->src[0]->type) || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32) { return false; @@ -3985,17 +3969,17 @@ static bool ggml_cuda_parse_mul_mat_lane(const ggml_cgraph * cgraph, int i, ggml 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 { + const bool add_lhs_out = add->src[0] == lane.out; + const bool add_rhs_out = add->src[1] == lane.out; + if (!add_lhs_out && !add_rhs_out) { return false; } + + lane.bias = add_lhs_out ? add->src[1] : add->src[0]; if (add->type != GGML_TYPE_F32 || !ggml_are_same_shape(add->src[0], add->src[1])) { return false; } - if (!ggml_cuda_can_parse_mm_lane_bias(mm, lane.bias)) { + if (lane.bias->type != GGML_TYPE_F32 || lane.bias->ne[0] != mm->ne[0]) { return false; } lane.bias_node = add; @@ -4095,10 +4079,6 @@ 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]; @@ -4118,7 +4098,7 @@ static int ggml_cuda_try_fuse_mm_glu(ggml_backend_cuda_context * cuda_ctx, ggml_ fusion_data.glu_op = ggml_get_glu_op(glu); const ggml_tensor * ids = up->mm->op == GGML_OP_MUL_MAT_ID ? up->mm->src[2] : nullptr; - if (up->scale == nullptr && gate->scale == nullptr && ggml_cuda_should_fuse_mul_mat_vec_f(up->mm)) { + if (ggml_cuda_should_fuse_mul_mat_vec_f(up->mm)) { ggml_cuda_mul_mat_vec_f(*cuda_ctx, up->mm->src[0], up->mm->src[1], ids, cgraph->nodes[glu_idx], &fusion_data); return glu_idx - i; } @@ -4151,22 +4131,6 @@ static int ggml_cuda_try_fuse_mm_lane(ggml_backend_cuda_context * cuda_ctx, ggml const ggml_tensor * ids = lane.mm->op == GGML_OP_MUL_MAT_ID ? lane.mm->src[2] : nullptr; - // Lane scale fusion is implemented by MMVQ only because it is limited to NVFP4. - // This path owns scale lanes, including scale followed by bias. - if (lane.scale != nullptr) { - 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_cuda_mul_mat_vec_q(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.out, &fusion_data); - return lane.n_nodes - 1; - } - - // Bias-only lanes can use either MMVF or MMVQ. if (ggml_cuda_should_fuse_mul_mat_vec_f(lane.mm)) { ggml_cuda_mul_mat_vec_f(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.out, &fusion_data); return lane.n_nodes - 1; From 6b9024e842e2b51bcb4ad499c123958a3d41f290 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Thu, 11 Jun 2026 14:00:40 +0200 Subject: [PATCH 16/28] Move gate_bias and scale into the same active-thread guard Latest perf numbers: B6000 build: 5b7d9f272 (9578) + ./scripts/compare-llama-bench.py -b master -c osimons/nvfp4_fuse_mmvq --tool llama-bench -i llama-bench.sqlite | CPU | Model | Test | t/s master | t/s osimons/nvfp4_fuse_mmvq | Speedup | |:----------------------------|:-------------------------|:-------------|-------------:|------------------------------:|----------:| | INTEL(R) XEON(R) GOLD 6542Y | gemma4 26B.A4B NVFP4 | tg128@d32768 | 151.79 | 154.10 | 1.02 | | INTEL(R) XEON(R) GOLD 6542Y | gemma4 26B.A4B Q4_K_M | tg128@d32768 | 187.90 | 187.27 | 1.00 | | INTEL(R) XEON(R) GOLD 6542Y | gpt-oss 20B MXFP4 MoE | tg128@d32768 | 303.77 | 306.56 | 1.01 | | INTEL(R) XEON(R) GOLD 6542Y | qwen35moe 35B.A3B NVFP4 | tg128@d32768 | 193.41 | 207.99 | 1.08 | | INTEL(R) XEON(R) GOLD 6542Y | qwen35moe 35B.A3B Q4_K_M | tg128@d32768 | 217.60 | 218.58 | 1.00 | DGX Spark build: 5b7d9f272 (9578) + ./scripts/compare-llama-bench.py -b master -c osimons/nvfp4_fuse_mmvq --tool llama-bench -i llama-bench.sqlite | CPU | Model | Test | t/s master | t/s osimons/nvfp4_fuse_mmvq | Speedup | |:------|:-------------------------|:-------------|-------------:|------------------------------:|----------:| | CPU | gemma4 26B.A4B NVFP4 | tg128@d32768 | 34.61 | 34.84 | 1.01 | | CPU | gemma4 26B.A4B Q4_K_M | tg128@d32768 | 46.95 | 46.90 | 1.00 | | CPU | gpt-oss 20B MXFP4 MoE | tg128@d32768 | 64.84 | 64.62 | 1.00 | | CPU | qwen35moe 35B.A3B NVFP4 | tg128@d32768 | 59.63 | 60.72 | 1.02 | | CPU | qwen35moe 35B.A3B Q4_K_M | tg128@d32768 | 56.53 | 56.55 | 1.00 | PPL values for 5 chunks: this PR model mode ppl uncertainty log /mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf fusion_enabled 5.2892 0.35389 ppl-value-checks/Qwen3.6-35B-A3B-UD-Q4_K_M.fusion_enabled.log /mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf fusion_disabled 5.2742 0.35215 ppl-value-checks/Qwen3.6-35B-A3B-UD-Q4_K_M.fusion_disabled.log /mnt/share/gguf/nvidia/Qwen3.6-35B-A3B-2.06GB-per-token-CT/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.gguf fusion_enabled 5.4487 0.36866 ppl-value-checks/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.fusion_enabled.log /mnt/share/gguf/nvidia/Qwen3.6-35B-A3B-2.06GB-per-token-CT/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.gguf fusion_disabled 5.4403 0.36782 ppl-value-checks/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.fusion_disabled.log /mnt/share/gguf/nvidia/Gemma-4-26B-A4B-NVFP4/Gemma-4-26B-A4B-NVFP4_fp8_q8.gguf fusion_enabled 17342.4348 3703.13932 ppl-value-checks/Gemma-4-26B-A4B-NVFP4_fp8_q8.fusion_enabled.log /mnt/share/gguf/nvidia/Gemma-4-26B-A4B-NVFP4/Gemma-4-26B-A4B-NVFP4_fp8_q8.gguf fusion_disabled 18627.0624 3998.42475 ppl-value-checks/Gemma-4-26B-A4B-NVFP4_fp8_q8.fusion_disabled.log /mnt/share/gguf/ggml-org/gpt-oss-20b-GGUF/gpt-oss-20b-mxfp4.gguf fusion_enabled 363.8913 33.14007 ppl-value-checks/gpt-oss-20b-mxfp4.fusion_enabled.log /mnt/share/gguf/ggml-org/gpt-oss-20b-GGUF/gpt-oss-20b-mxfp4.gguf fusion_disabled 363.8913 33.14007 ppl-value-checks/gpt-oss-20b-mxfp4.fusion_disabled.log /mnt/share/gguf/unsloth/gemma-4-26B-A4B-it-GGUF/gemma-4-26B-A4B-it-UD-Q4_K_XL.gguf fusion_enabled 17330.3926 3716.70472 ppl-value-checks/gemma-4-26B-A4B-it-UD-Q4_K_XL.fusion_enabled.log /mnt/share/gguf/unsloth/gemma-4-26B-A4B-it-GGUF/gemma-4-26B-A4B-it-UD-Q4_K_XL.gguf fusion_disabled 17933.9524 3883.17066 ppl-value-checks/gemma-4-26B-A4B-it-UD-Q4_K_XL.fusion_disabled.log master: summary: ppl-value-checks/summary.tsv model mode ppl uncertainty log /mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf fusion_enabled 5.2892 0.35389 ppl-value-checks/Qwen3.6-35B-A3B-UD-Q4_K_M.fusion_enabled.log /mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf fusion_disabled 5.2742 0.35215 ppl-value-checks/Qwen3.6-35B-A3B-UD-Q4_K_M.fusion_disabled.log /mnt/share/gguf/nvidia/Qwen3.6-35B-A3B-2.06GB-per-token-CT/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.gguf fusion_enabled 5.4487 0.36866 ppl-value-checks/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.fusion_enabled.log /mnt/share/gguf/nvidia/Qwen3.6-35B-A3B-2.06GB-per-token-CT/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.gguf fusion_disabled 5.4403 0.36782 ppl-value-checks/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.fusion_disabled.log /mnt/share/gguf/nvidia/Gemma-4-26B-A4B-NVFP4/Gemma-4-26B-A4B-NVFP4_fp8_q8.gguf fusion_enabled 17342.4348 3703.13932 ppl-value-checks/Gemma-4-26B-A4B-NVFP4_fp8_q8.fusion_enabled.log /mnt/share/gguf/nvidia/Gemma-4-26B-A4B-NVFP4/Gemma-4-26B-A4B-NVFP4_fp8_q8.gguf fusion_disabled 18627.0624 3998.42475 ppl-value-checks/Gemma-4-26B-A4B-NVFP4_fp8_q8.fusion_disabled.log /mnt/share/gguf/ggml-org/gpt-oss-20b-GGUF/gpt-oss-20b-mxfp4.gguf fusion_enabled 363.8913 33.14007 ppl-value-checks/gpt-oss-20b-mxfp4.fusion_enabled.log /mnt/share/gguf/ggml-org/gpt-oss-20b-GGUF/gpt-oss-20b-mxfp4.gguf fusion_disabled 363.8913 33.14007 ppl-value-checks/gpt-oss-20b-mxfp4.fusion_disabled.log /mnt/share/gguf/unsloth/gemma-4-26B-A4B-it-GGUF/gemma-4-26B-A4B-it-UD-Q4_K_XL.gguf fusion_enabled 17330.3926 3716.70472 ppl-value-checks/gemma-4-26B-A4B-it-UD-Q4_K_XL.fusion_enabled.log /mnt/share/gguf/unsloth/gemma-4-26B-A4B-it-GGUF/gemma-4-26B-A4B-it-UD-Q4_K_XL.gguf fusion_disabled 17933.9524 3883.17066 ppl-value-checks/gemma-4-26B-A4B-it-UD-Q4_K_XL.fusion_disabled.log --- ggml/src/ggml-cuda/mmvq.cu | 25 ++++++++++++------------- 1 file changed, 12 insertions(+), 13 deletions(-) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 83f20b9aeb21..3335f6663ba6 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -555,7 +555,6 @@ static __global__ void mul_mat_vec_q( const uint32_t channel_bias = ids ? channel_x : channel_dst; if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 && (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) { - const uint32_t channel_bias = ids ? channel_x : channel_dst; if (use_bias) { x_bias = x_bias + sample_dst * stride_sample_dst + channel_bias * stride_channel_dst + row0; #pragma unroll @@ -563,20 +562,20 @@ static __global__ void mul_mat_vec_q( x_biases[j] = x_bias[j * stride_col_dst + threadIdx.x]; } } - } - if (use_gate_bias) { - gate_bias = gate_bias + sample_dst * stride_sample_dst + channel_bias * stride_channel_dst + row0; + if (use_gate_bias) { + gate_bias = gate_bias + sample_dst * stride_sample_dst + channel_bias * stride_channel_dst + row0; #pragma unroll - for (int j = 0; j < ncols_dst; ++j) { - gate_biases[j] = gate_bias[j * stride_col_dst + threadIdx.x]; - } - } - if constexpr (type == GGML_TYPE_NVFP4) { - if (use_scale) { - x_scales = x_scale[ids ? channel_x : 0]; + for (int j = 0; j < ncols_dst; ++j) { + gate_biases[j] = gate_bias[j * stride_col_dst + threadIdx.x]; + } } - 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]; + } } } } From 8dd49e004f2cc38c3e9bad72fa42215c52ef0c7b Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Tue, 16 Jun 2026 15:20:54 +0200 Subject: [PATCH 17/28] Allow views to weights in ggml_can_fuse_subgraph --- ggml/src/ggml-cuda/ggml-cuda.cu | 106 ++++---------------------------- ggml/src/ggml.c | 9 ++- 2 files changed, 19 insertions(+), 96 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 6eb1ee0d2092..f2acd5473411 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3581,83 +3581,6 @@ static bool ggml_cuda_can_fuse_subgraph(const struct ggml_cgraph * cgraph, ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, count, out_nodes, out_count, is_topk_moe); } -// This mirrors ggml_can_fuse_subgraph for parser-validated patterns. It avoids -// rebuilding variable lane patterns as fixed op lists, and allows parser-approved -// external view sources such as the NVFP4 MUL_MAT_ID scale reshape. -static bool ggml_cuda_can_fuse_parsed_subgraph(const struct ggml_cgraph * cgraph, - int node_idx, - int count, - const int * out_nodes, - int out_count, - const int * external_view_nodes = nullptr, - int external_view_count = 0, - bool is_topk_moe = false) { - if (node_idx + count > cgraph->n_nodes) { - return false; - } - - const auto is_output = [&](int idx) { - for (int j = 0; j < out_count; ++j) { - if (out_nodes[j] == idx) { - return true; - } - } - return false; - }; - - const auto is_in_subgraph = [&](const ggml_tensor * tensor) { - for (int j = node_idx; j < node_idx + count; ++j) { - if (cgraph->nodes[j] == tensor) { - return true; - } - } - return false; - }; - - const auto is_allowed_external_view = [&](int idx) { - for (int j = 0; j < external_view_count; ++j) { - if (external_view_nodes[j] == idx) { - return true; - } - } - return false; - }; - - for (int j = node_idx; j < node_idx + count; ++j) { - const ggml_tensor * node = cgraph->nodes[j]; - if ((node->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) { - return false; - } - if (is_output(j)) { - continue; - } - if (node->flags & GGML_TENSOR_FLAG_OUTPUT) { - return false; - } - - int subgraph_uses = 0; - for (int k = j + 1; k < node_idx + count; ++k) { - const ggml_tensor * other = cgraph->nodes[k]; - for (int src_idx = 0; src_idx < GGML_MAX_SRC; ++src_idx) { - if (other->src[src_idx] == node) { - subgraph_uses++; - } - } - } - if (subgraph_uses != ggml_node_get_use_count(cgraph, j)) { - return false; - } - - for (const ggml_tensor * view_src = node->view_src; view_src != nullptr; view_src = view_src->view_src) { - if (!is_in_subgraph(view_src) && !is_allowed_external_view(j)) { - return false; - } - } - } - - return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, count, out_nodes, out_count, is_topk_moe); -} - static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, int node_idx, std::initializer_list ops, @@ -3858,7 +3781,6 @@ struct ggml_cuda_mm_lane { ggml_tensor * out = nullptr; const ggml_tensor * bias = nullptr; const ggml_tensor * scale = nullptr; - int scale_view_node = -1; int n_nodes = 0; }; @@ -3909,7 +3831,6 @@ static bool ggml_cuda_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, g } lane.scale = scale; - lane.scale_view_node = i + lane.n_nodes; lane.out = mul; lane.n_nodes += 4; } @@ -4003,13 +3924,6 @@ static bool ggml_cuda_parse_mm_lane(const ggml_cgraph * cgraph, int i, ggml_cuda return false; } -static int ggml_cuda_add_mm_lane_external_view_node(const ggml_cuda_mm_lane & lane, int * external_view_nodes, int n_external_view_nodes) { - if (lane.scale_view_node != -1) { - external_view_nodes[n_external_view_nodes++] = lane.scale_view_node; - } - return n_external_view_nodes; -} - static bool ggml_cuda_should_fuse_mm_lanes(const ggml_cuda_mm_lane & up, const ggml_cuda_mm_lane & gate, const ggml_tensor * glu) { if (up.mm->op != gate.mm->op) { return false; @@ -4081,11 +3995,12 @@ static int ggml_cuda_try_fuse_mm_glu(ggml_backend_cuda_context * cuda_ctx, ggml_ const int out_nodes[] = { glu_idx }; const int n_nodes = glu_idx - i + 1; - int external_view_nodes[2]; - int n_external_view_nodes = 0; - n_external_view_nodes = ggml_cuda_add_mm_lane_external_view_node(*up, external_view_nodes, n_external_view_nodes); - n_external_view_nodes = ggml_cuda_add_mm_lane_external_view_node(*gate, external_view_nodes, n_external_view_nodes); - if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, n_nodes, out_nodes, 1, external_view_nodes, n_external_view_nodes)) { + std::vector ops; + ops.reserve(n_nodes); + for (int j = 0; j < n_nodes; ++j) { + ops.push_back(cgraph->nodes[i + j]->op); + } + if (!ggml_cuda_can_fuse_subgraph(cgraph, i, n_nodes, ops.data(), out_nodes, 1)) { return 0; } @@ -4119,9 +4034,12 @@ static int ggml_cuda_try_fuse_mm_lane(ggml_backend_cuda_context * cuda_ctx, ggml const int out_idx = i + lane.n_nodes - 1; const int out_nodes[] = { out_idx }; - int external_view_nodes[1]; - int n_external_view_nodes = ggml_cuda_add_mm_lane_external_view_node(lane, external_view_nodes, 0); - if (!ggml_cuda_can_fuse_parsed_subgraph(cgraph, i, lane.n_nodes, out_nodes, 1, external_view_nodes, n_external_view_nodes)) { + std::vector ops; + ops.reserve(lane.n_nodes); + for (int j = 0; j < lane.n_nodes; ++j) { + ops.push_back(cgraph->nodes[i + j]->op); + } + if (!ggml_cuda_can_fuse_subgraph(cgraph, i, lane.n_nodes, ops.data(), out_nodes, 1)) { return 0; } diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 8815c67d8bcc..f96445956933 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -7371,6 +7371,10 @@ static int ggml_node_list_find_tensor(const struct ggml_cgraph * cgraph, return -1; } +static bool ggml_is_constant_view_src(const struct ggml_tensor * tensor) { + return tensor->buffer != NULL && ggml_backend_buffer_get_usage(tensor->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS; +} + bool ggml_can_fuse_subgraph_ext(const struct ggml_cgraph * cgraph, const int * node_idxs, int count, @@ -7416,10 +7420,11 @@ bool ggml_can_fuse_subgraph_ext(const struct ggml_cgraph * cgraph, return false; } - // if node is a view, check if the view_src and all it's parent view_srcs are within the subgraph + // if node is a view, check if the view_src and all its parent view_srcs are within the subgraph. + // external view sources are allowed only for weight tensors, which are constant for this graph execution. struct ggml_tensor * view_src = node->view_src; while (view_src) { - if (ggml_node_list_find_tensor(cgraph, node_idxs, count, view_src) == -1) { + if (ggml_node_list_find_tensor(cgraph, node_idxs, count, view_src) == -1 && !ggml_is_constant_view_src(view_src)) { return false; } view_src = view_src->view_src; From 89fcfc21d472485f50c9dbad8bff9e66d9c1a5b3 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Wed, 17 Jun 2026 09:28:59 +0200 Subject: [PATCH 18/28] Remove gate_first from test_mul_mat_vec_fusion --- tests/test-backend-ops.cpp | 42 ++++++++++---------------------------- 1 file changed, 11 insertions(+), 31 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 9c1f16cfb68b..fee80c47487e 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5789,21 +5789,20 @@ struct test_mul_mat_vec_fusion : public test_case { const bool with_bias; const bool with_gate; const bool with_lane_scale; - const bool gate_first; std::array batch_dims; test_mul_mat_vec_fusion(ggml_type type, ggml_glu_op op, int64_t m, int64_t n, int64_t k, bool use_id = false, int n_mats = 1, int n_used = 1, bool b = false, bool with_bias = false, bool with_gate = true, - bool with_lane_scale = false, bool gate_first = false, std::array batch_dims = {4, 2}) + bool with_lane_scale = false, std::array batch_dims = {4, 2}) : type(type), glu_op(op), m(m), n(n), k(k), use_id(use_id), n_mats(n_mats), n_used(n_used), b(b), with_bias(with_bias), - with_gate(with_gate), with_lane_scale(with_lane_scale), gate_first(gate_first), batch_dims(batch_dims) { + with_gate(with_gate), with_lane_scale(with_lane_scale), batch_dims(batch_dims) { if (use_id) { GGML_ASSERT(n_used <= n_mats); } } std::string vars() override { - return VARS_TO_STR14(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_lane_scale, gate_first, batch_dims); + return VARS_TO_STR13(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_lane_scale, batch_dims); } std::string op_desc(ggml_tensor * t) override { @@ -5877,15 +5876,8 @@ struct test_mul_mat_vec_fusion : public test_case { return ffn_gate; }; - ggml_tensor * ffn_up = nullptr; - ggml_tensor * ffn_gate = nullptr; - if (with_gate && gate_first) { - ffn_gate = build_gate_lane(); - ffn_up = build_up_lane(); - } else { - ffn_up = build_up_lane(); - ffn_gate = with_gate ? build_gate_lane() : nullptr; - } + ggml_tensor * ffn_up = build_up_lane(); + ggml_tensor * ffn_gate = with_gate ? build_gate_lane() : nullptr; ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; @@ -5931,15 +5923,8 @@ struct test_mul_mat_vec_fusion : public test_case { return ffn_gate; }; - ggml_tensor * ffn_up = nullptr; - ggml_tensor * ffn_gate = nullptr; - if (with_gate && gate_first) { - ffn_gate = build_gate_lane(); - ffn_up = build_up_lane(); - } else { - ffn_up = build_up_lane(); - ffn_gate = with_gate ? build_gate_lane() : nullptr; - } + ggml_tensor * ffn_up = build_up_lane(); + ggml_tensor * ffn_gate = with_gate ? build_gate_lane() : nullptr; ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; @@ -9147,15 +9132,10 @@ static std::vector> make_test_cases_eval() { if (with_lane_scale && type != GGML_TYPE_NVFP4) { continue; } - for (bool gate_first : {false, true}) { - if (!with_gate && gate_first) { - continue; - } - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, - use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, gate_first)); - test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, - use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, gate_first, {1, 1})); - } + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, + use_id, 16, 8, b, with_bias, with_gate, with_lane_scale)); + test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256, + use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, {1, 1})); } } } From d8955f68e1c6cb47f41f7bb46c0be577c9681a78 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Tue, 16 Jun 2026 16:07:56 +0200 Subject: [PATCH 19/28] Ditch lane-parsing approach in favor of hard-coded patterns --- ggml/src/ggml-cuda/ggml-cuda.cu | 960 +++++++++++++++++++++----------- 1 file changed, 637 insertions(+), 323 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index f2acd5473411..571e09c5d9a9 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2418,6 +2418,108 @@ static void ggml_cuda_mul_mat_batched_cublas(ggml_backend_cuda_context & ctx, co } } +static bool ggml_cuda_should_fuse_mul_mat(const ggml_tensor * ffn_up, + const ggml_tensor * ffn_gate, + const ggml_tensor * glu, + const ggml_tensor * ffn_up_bias = nullptr, + const ggml_tensor * ffn_gate_bias = nullptr, + const ggml_tensor * ffn_up_scale = nullptr, + const ggml_tensor * ffn_gate_scale = nullptr) { + const bool has_bias = ffn_up_bias != nullptr || ffn_gate_bias != nullptr; + const bool has_scale = ffn_up_scale != nullptr || ffn_gate_scale != nullptr; + + if (has_bias && (!ffn_up_bias || !ffn_gate_bias)) { + return false; + } + if (has_scale && (!ffn_up_scale || !ffn_gate_scale)) { + return false; + } + + const bool is_mul_mat = ffn_up->op == GGML_OP_MUL_MAT && ffn_gate->op == GGML_OP_MUL_MAT && glu->op == GGML_OP_GLU; + const bool is_mul_mat_id = ffn_up->op == GGML_OP_MUL_MAT_ID && ffn_gate->op == GGML_OP_MUL_MAT_ID && glu->op == GGML_OP_GLU; + + GGML_ASSERT(ffn_up && ffn_gate && glu); + + if (!is_mul_mat && !is_mul_mat_id) { + return false; + } + + const ggml_op expected_bias_op = is_mul_mat ? GGML_OP_ADD : GGML_OP_ADD_ID; + const ggml_tensor * ffn_up_bias_src = has_scale ? ffn_up_scale : ffn_up; + const ggml_tensor * ffn_gate_bias_src = has_scale ? ffn_gate_scale : ffn_gate; + const ggml_tensor * ffn_up_out = has_bias ? ffn_up_bias : ffn_up_bias_src; + const ggml_tensor * ffn_gate_out = has_bias ? ffn_gate_bias : ffn_gate_bias_src; + + if (glu->src[0] != ffn_gate_out || glu->src[1] != ffn_up_out) { + return false; + } + + if (has_scale) { + if (ffn_up_scale->op != GGML_OP_MUL || ffn_gate_scale->op != GGML_OP_MUL) { + return false; + } + const bool up_has_mm = ffn_up_scale->src[0] == ffn_up || ffn_up_scale->src[1] == ffn_up; + const bool gate_has_mm = ffn_gate_scale->src[0] == ffn_gate || ffn_gate_scale->src[1] == ffn_gate; + if (!up_has_mm || !gate_has_mm) { + return false; + } + } + + if (has_bias) { + if (ffn_up_bias->op != expected_bias_op || ffn_gate_bias->op != expected_bias_op) { + return false; + } + + if (expected_bias_op == GGML_OP_ADD) { + const bool up_has_mul = ffn_up_bias->src[0] == ffn_up_bias_src || ffn_up_bias->src[1] == ffn_up_bias_src; + const bool gate_has_mul = ffn_gate_bias->src[0] == ffn_gate_bias_src || ffn_gate_bias->src[1] == ffn_gate_bias_src; + if (!up_has_mul || !gate_has_mul) { + return false; + } + } else { // GGML_OP_ADD_ID + if (ffn_up_bias->src[0] != ffn_up_bias_src || ffn_gate_bias->src[0] != ffn_gate_bias_src) { + return false; + } + if (ffn_up_bias->src[2] != ffn_up->src[2] || ffn_gate_bias->src[2] != ffn_gate->src[2]) { + return false; + } + } + } + + if (ffn_up->src[0]->type != ffn_gate->src[0]->type || !ggml_are_same_shape(ffn_up->src[0], ffn_gate->src[0]) || + !ggml_are_same_stride(ffn_up->src[0], ffn_gate->src[0])) { + return false; + } + + if (ffn_up->src[1] != ffn_gate->src[1]) { + return false; + } + + if (is_mul_mat_id && ffn_up->src[2] != ffn_gate->src[2]) { + 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(ffn_up->src[0]->buffer->buft) || + ggml_backend_buft_is_cuda_split(ffn_gate->src[0]->buffer->buft); + + //TODO: add support for fusion for split buffers + if (split) { + return false; + } + + return true; +} + static bool ggml_cuda_should_fuse_mul_mat_vec_f(const ggml_tensor * tensor) { ggml_tensor * src0 = tensor->src[0]; ggml_tensor * src1 = tensor->src[1]; @@ -3570,17 +3672,6 @@ static bool ggml_cuda_check_fusion_memory_ranges(const ggml_cgraph * cgraph, } -static bool ggml_cuda_can_fuse_subgraph(const struct ggml_cgraph * cgraph, - int node_idx, - int count, - const enum ggml_op * ops, - const int * out_nodes, - int out_count, - bool is_topk_moe = false) { - return ggml_can_fuse_subgraph(cgraph, node_idx, count, ops, out_nodes, out_count) && - ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, count, out_nodes, out_count, is_topk_moe); -} - static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, int node_idx, std::initializer_list ops, @@ -3595,11 +3686,41 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, return std::equal(list1.begin(), list1.end(), list2.begin(), list2.end()); }; + std::initializer_list mul_mat_bias_glu_ops = { GGML_OP_MUL_MAT, GGML_OP_ADD, GGML_OP_MUL_MAT, GGML_OP_ADD, GGML_OP_GLU }; + std::initializer_list mul_mat_id_bias_glu_ops = { GGML_OP_MUL_MAT_ID, GGML_OP_ADD_ID, GGML_OP_MUL_MAT_ID, GGML_OP_ADD_ID, GGML_OP_GLU }; + + std::initializer_list mul_mat_id_glu_ops = { GGML_OP_MUL_MAT_ID, GGML_OP_MUL_MAT_ID, GGML_OP_GLU }; + std::initializer_list mul_mat_glu_ops = { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT, GGML_OP_GLU }; + + if ((is_equal(mul_mat_bias_glu_ops, ops) || is_equal(mul_mat_id_bias_glu_ops, ops)) && + ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 4 })) { + const ggml_tensor * ffn_gate = cgraph->nodes[node_idx]; + const ggml_tensor * ffn_gate_bias = cgraph->nodes[node_idx + 1]; + const ggml_tensor * ffn_up = cgraph->nodes[node_idx + 2]; + const ggml_tensor * ffn_up_bias = cgraph->nodes[node_idx + 3]; + const ggml_tensor * glu = cgraph->nodes[node_idx + 4]; + + if (ggml_cuda_should_fuse_mul_mat(ffn_up, ffn_gate, glu, ffn_up_bias, ffn_gate_bias)) { + int out_nodes[] = { node_idx + 4 }; + return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1); + } + } + + if ((is_equal(mul_mat_id_glu_ops, ops) || is_equal(mul_mat_glu_ops, ops)) && + ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 2 })) { + const ggml_tensor * ffn_gate = cgraph->nodes[node_idx]; + const ggml_tensor * ffn_up = cgraph->nodes[node_idx + 1]; + const ggml_tensor * glu = cgraph->nodes[node_idx + 2]; + + if (ggml_cuda_should_fuse_mul_mat(ffn_up, ffn_gate, glu)) { + int out_nodes[] = { node_idx + 2 }; + return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1); + } + } + std::initializer_list rope_set_rows_ops = { GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS }; - const int rope_set_rows_out_nodes[] = { node_idx + 2 }; - if (is_equal(rope_set_rows_ops, ops) && - ggml_cuda_can_fuse_subgraph(cgraph, node_idx, (int) ops.size(), ops.begin(), rope_set_rows_out_nodes, 1)) { + if (is_equal(rope_set_rows_ops, ops) && ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 2 })) { const ggml_tensor * rope = cgraph->nodes[node_idx]; const ggml_tensor * view = cgraph->nodes[node_idx + 1]; const ggml_tensor * set_rows = cgraph->nodes[node_idx + 2]; @@ -3771,297 +3892,6 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, return false; } - -// Matched MM lane forms: -// MUL_MAT [MUL scalar_scale] [ADD] -// MUL_MAT_ID [RESHAPE -> REPEAT -> GET_ROWS -> MUL expert_scale] [ADD_ID] -struct ggml_cuda_mm_lane { - ggml_tensor * mm = nullptr; - ggml_tensor * bias_node = nullptr; - ggml_tensor * out = nullptr; - const ggml_tensor * bias = nullptr; - const ggml_tensor * scale = nullptr; - int n_nodes = 0; -}; - -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_parse_mul_mat_id_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mm_lane & lane) { - ggml_tensor * mm = cgraph->nodes[i]; - if (!ggml_cuda_can_parse_mm_lane_type(mm->src[0]->type) || mm->src[1]->type != GGML_TYPE_F32 || mm->type != GGML_TYPE_F32 || mm->src[2] == nullptr) { - return false; - } - - lane = {}; - lane.mm = mm; - lane.out = mm; - lane.n_nodes = 1; - - if (ggml_is_quantized(mm->src[0]->type) && ggml_cuda_can_fuse_mm_lane_scale(mm) && i + lane.n_nodes + 3 < cgraph->n_nodes) { - ggml_tensor * reshape = cgraph->nodes[i + lane.n_nodes + 0]; - ggml_tensor * repeat = cgraph->nodes[i + lane.n_nodes + 1]; - ggml_tensor * getrows = cgraph->nodes[i + lane.n_nodes + 2]; - ggml_tensor * mul = cgraph->nodes[i + lane.n_nodes + 3]; - - if (reshape->op == GGML_OP_RESHAPE && repeat->op == GGML_OP_REPEAT && - getrows->op == GGML_OP_GET_ROWS && mul->op == GGML_OP_MUL) { - if (repeat->src[0] != reshape || getrows->src[0] != repeat || getrows->src[1] != mm->src[2]) { - return false; - } - - const bool mul_has_out = mul->src[0] == lane.out || mul->src[1] == lane.out; - const bool mul_has_scale = mul->src[0] == getrows || mul->src[1] == getrows; - if (!mul_has_out || !mul_has_scale) { - return false; - } - - const ggml_tensor * scale = reshape->src[0]; - if (scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != mm->src[0]->ne[2]) { - 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 += 4; - } - } - - if (i + lane.n_nodes < cgraph->n_nodes && cgraph->nodes[i + lane.n_nodes]->op == GGML_OP_ADD_ID) { - ggml_tensor * add = cgraph->nodes[i + lane.n_nodes]; - if (add->src[0] != lane.out || add->src[2] != mm->src[2] || add->type != GGML_TYPE_F32) { - return false; - } - const ggml_tensor * bias = add->src[1]; - if (bias->type != GGML_TYPE_F32 || bias->ne[0] != mm->ne[0] || bias->ne[1] != mm->src[0]->ne[2]) { - return false; - } - lane.bias_node = add; - lane.bias = bias; - lane.out = add; - lane.n_nodes++; - } - - return true; -} - -static bool ggml_cuda_parse_mul_mat_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mm_lane & lane) { - ggml_tensor * mm = cgraph->nodes[i]; - if (!ggml_cuda_can_parse_mm_lane_type(mm->src[0]->type) || 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 (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) { - 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 false; - } - - 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++; - } - - 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]; - const bool add_lhs_out = add->src[0] == lane.out; - const bool add_rhs_out = add->src[1] == lane.out; - if (!add_lhs_out && !add_rhs_out) { - return false; - } - - lane.bias = add_lhs_out ? add->src[1] : add->src[0]; - if (add->type != GGML_TYPE_F32 || !ggml_are_same_shape(add->src[0], add->src[1])) { - return false; - } - if (lane.bias->type != GGML_TYPE_F32 || lane.bias->ne[0] != mm->ne[0]) { - return false; - } - lane.bias_node = add; - lane.out = add; - lane.n_nodes++; - } - - return true; -} - -static bool ggml_cuda_parse_mm_lane(const ggml_cgraph * cgraph, int i, ggml_cuda_mm_lane & lane) { - if (i >= cgraph->n_nodes) { - return false; - } - if (cgraph->nodes[i]->op == GGML_OP_MUL_MAT_ID) { - return ggml_cuda_parse_mul_mat_id_lane(cgraph, i, lane); - } - if (cgraph->nodes[i]->op == GGML_OP_MUL_MAT) { - return ggml_cuda_parse_mul_mat_lane(cgraph, i, lane); - } - return false; -} - -static bool ggml_cuda_should_fuse_mm_lanes(const ggml_cuda_mm_lane & up, const ggml_cuda_mm_lane & gate, const ggml_tensor * glu) { - if (up.mm->op != gate.mm->op) { - return false; - } - - 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 (up.mm->op == GGML_OP_MUL_MAT_ID && up.mm->src[2] != gate.mm->src[2]) { - 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_mm_glu(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { - ggml_cuda_mm_lane lane0; - if (!ggml_cuda_parse_mm_lane(cgraph, i, lane0)) { - return 0; - } - - ggml_cuda_mm_lane lane1; - if (!ggml_cuda_parse_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; - } - - const ggml_tensor * glu = cgraph->nodes[glu_idx]; - ggml_cuda_mm_lane * gate = nullptr; - ggml_cuda_mm_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; - std::vector ops; - ops.reserve(n_nodes); - for (int j = 0; j < n_nodes; ++j) { - ops.push_back(cgraph->nodes[i + j]->op); - } - if (!ggml_cuda_can_fuse_subgraph(cgraph, i, n_nodes, ops.data(), out_nodes, 1)) { - 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); - - const ggml_tensor * ids = up->mm->op == GGML_OP_MUL_MAT_ID ? up->mm->src[2] : nullptr; - if (ggml_cuda_should_fuse_mul_mat_vec_f(up->mm)) { - ggml_cuda_mul_mat_vec_f(*cuda_ctx, up->mm->src[0], up->mm->src[1], ids, cgraph->nodes[glu_idx], &fusion_data); - return glu_idx - i; - } - - if (ggml_cuda_should_fuse_mul_mat_vec_q(up->mm)) { - ggml_cuda_mul_mat_vec_q(*cuda_ctx, up->mm->src[0], up->mm->src[1], ids, cgraph->nodes[glu_idx], &fusion_data); - return glu_idx - i; - } - - return 0; -} - -static int ggml_cuda_try_fuse_mm_lane(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph * cgraph, int i) { - ggml_cuda_mm_lane lane; - if (!ggml_cuda_parse_mm_lane(cgraph, i, lane) || (lane.scale == nullptr && lane.bias == nullptr)) { - return 0; - } - - const int out_idx = i + lane.n_nodes - 1; - const int out_nodes[] = { out_idx }; - std::vector ops; - ops.reserve(lane.n_nodes); - for (int j = 0; j < lane.n_nodes; ++j) { - ops.push_back(cgraph->nodes[i + j]->op); - } - if (!ggml_cuda_can_fuse_subgraph(cgraph, i, lane.n_nodes, ops.data(), out_nodes, 1)) { - return 0; - } - - ggml_cuda_mm_fusion_args_host fusion_data{}; - fusion_data.x_bias = lane.bias; - fusion_data.x_scale = lane.scale; - - const ggml_tensor * ids = lane.mm->op == GGML_OP_MUL_MAT_ID ? lane.mm->src[2] : nullptr; - - if (ggml_cuda_should_fuse_mul_mat_vec_f(lane.mm)) { - ggml_cuda_mul_mat_vec_f(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.out, &fusion_data); - return lane.n_nodes - 1; - } - - if (ggml_cuda_should_fuse_mul_mat_vec_q(lane.mm)) { - ggml_cuda_mul_mat_vec_q(*cuda_ctx, lane.mm->src[0], lane.mm->src[1], ids, lane.out, &fusion_data); - return lane.n_nodes - 1; - } - - return 0; -} - // 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) { @@ -4117,8 +3947,9 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph weights = cgraph->nodes[i + ops.size() - 1]; out_nodes[1] = i + ops.size() - 1; - if (ggml_cuda_can_fuse_subgraph(cgraph, i, (int) ops.size(), ops.data(), out_nodes, 2, /*is_topk_moe=*/true) && - ggml_cuda_should_use_topk_moe(node, logits, weights, ids)) { + if (ggml_can_fuse_subgraph(cgraph, i, ops.size(), ops.data(), out_nodes, 2) && + ggml_cuda_should_use_topk_moe(node, logits, weights, ids) && + ggml_cuda_check_fusion_memory_ranges(cgraph, i, ops.size(), out_nodes, 2, /*is_topk_moe=*/true)) { ggml_cuda_op_topk_moe(*cuda_ctx, logits, weights, ids, clamp, scale, bias, args); return ops.size() - 1; } @@ -4131,8 +3962,9 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph const ggml_tensor * softmax = cgraph->nodes[i + 4]; int out_nodes[2] = { i + 1, i + 5 }; - if (ggml_cuda_can_fuse_subgraph(cgraph, i, (int) ops.size(), ops.data(), out_nodes, 2, /*is_topk_moe=*/true) && - ggml_cuda_should_use_topk_moe(softmax, logits, weights, ids)) { + if (ggml_can_fuse_subgraph(cgraph, i, ops.size(), ops.data(), out_nodes, 2) && + ggml_cuda_should_use_topk_moe(softmax, logits, weights, ids) && + ggml_cuda_check_fusion_memory_ranges(cgraph, i, ops.size(), out_nodes, 2, /*is_topk_moe=*/true)) { ggml_cuda_op_topk_moe(*cuda_ctx, logits, weights, ids, clamp, scale, bias, args); return ops.size() - 1; } @@ -4229,16 +4061,498 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph } } - // Two-lane MM GLU fusion. Each lane is MUL_MAT[_ID] + optional scale + optional bias - int fused_mm_glu_nodes = ggml_cuda_try_fuse_mm_glu(cuda_ctx, cgraph, i); - if (fused_mm_glu_nodes > 0) { - return fused_mm_glu_nodes; + bool fused_mul_mat_vec = false; + int fused_node_count = 0; + + auto get_mul_mat_scale = [](const ggml_tensor * scale_node, const ggml_tensor * mm_node) -> const ggml_tensor * { + const bool scale_lhs_mm = scale_node->src[0] == mm_node; + const bool scale_rhs_mm = scale_node->src[1] == mm_node; + if (!scale_lhs_mm && !scale_rhs_mm) { + return nullptr; + } + + const ggml_tensor * scale = scale_lhs_mm ? scale_node->src[1] : scale_node->src[0]; + if (mm_node->src[0]->type != GGML_TYPE_NVFP4 || scale_node->type != GGML_TYPE_F32 || + scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != 1 || + !ggml_are_same_shape(scale_node, mm_node)) { + return nullptr; + } + + return scale; + }; + + auto get_mul_mat_id_scale = [](const ggml_tensor * reshape, const ggml_tensor * repeat, const ggml_tensor * getrows, + const ggml_tensor * scale_node, const ggml_tensor * mm_node) -> const ggml_tensor * { + if (repeat->src[0] != reshape || getrows->src[0] != repeat || getrows->src[1] != mm_node->src[2]) { + return nullptr; + } + if (!((scale_node->src[0] == mm_node && scale_node->src[1] == getrows) || + (scale_node->src[0] == getrows && scale_node->src[1] == mm_node))) { + return nullptr; + } + + const ggml_tensor * scale = reshape->src[0]; + if (mm_node->src[0]->type != GGML_TYPE_NVFP4 || scale_node->type != GGML_TYPE_F32 || + scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != mm_node->src[0]->ne[2] || + !ggml_are_same_shape(scale_node, mm_node)) { + return nullptr; + } + + return scale; + }; + + auto get_bias_tensor = [](const ggml_tensor * bias_node, const ggml_tensor * mul_node, ggml_op op_bias) -> const ggml_tensor * { + if (op_bias == GGML_OP_ADD) { + if (bias_node->src[0] == mul_node) { + return bias_node->src[1]; + } + if (bias_node->src[1] == mul_node) { + return bias_node->src[0]; + } + return nullptr; + } + GGML_ASSERT(op_bias == GGML_OP_ADD_ID); + GGML_ASSERT(bias_node->src[0] == mul_node); + return bias_node->src[1]; + }; + + // gate + glu + up, with optional scale/bias on both lanes. + for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) { + const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID; + + if (op == GGML_OP_MUL_MAT) { + for (const bool with_bias : { false, true }) { + const int gate_idx = i; + const int gate_scale_idx = i + 1; + const int gate_bias_idx = with_bias ? i + 2 : -1; + const int up_idx = with_bias ? i + 3 : i + 2; + const int up_scale_idx = up_idx + 1; + const int up_bias_idx = with_bias ? up_idx + 2 : -1; + const int glu_idx = with_bias ? up_idx + 3 : up_idx + 2; + + const int out_nodes[] = { glu_idx }; + ggml_op ops[7]; + if (with_bias) { + ops[0] = op; + ops[1] = GGML_OP_MUL; + ops[2] = bias_op; + ops[3] = op; + ops[4] = GGML_OP_MUL; + ops[5] = bias_op; + ops[6] = GGML_OP_GLU; + } else { + ops[0] = op; + ops[1] = GGML_OP_MUL; + ops[2] = op; + ops[3] = GGML_OP_MUL; + ops[4] = GGML_OP_GLU; + } + const int n_ops = with_bias ? 7 : 5; + + if (!ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1) || + !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1)) { + continue; + } + + ggml_tensor * gate_n = cgraph->nodes[gate_idx]; + ggml_tensor * gate_scale_n = cgraph->nodes[gate_scale_idx]; + ggml_tensor * gate_out_n = with_bias ? cgraph->nodes[gate_bias_idx] : gate_scale_n; + ggml_tensor * up_n = cgraph->nodes[up_idx]; + ggml_tensor * up_scale_n = cgraph->nodes[up_scale_idx]; + ggml_tensor * up_out_n = with_bias ? cgraph->nodes[up_bias_idx] : up_scale_n; + const ggml_tensor * glu = cgraph->nodes[glu_idx]; + + if (!ggml_cuda_should_fuse_mul_mat(up_n, gate_n, glu, + with_bias ? up_out_n : nullptr, with_bias ? gate_out_n : nullptr, up_scale_n, gate_scale_n)) { + continue; + } + + const ggml_tensor * gate_scale = get_mul_mat_scale(gate_scale_n, gate_n); + const ggml_tensor * up_scale = get_mul_mat_scale(up_scale_n, up_n); + if (!gate_scale || !up_scale) { + continue; + } + + const ggml_tensor * up_bias = with_bias ? get_bias_tensor(up_out_n, up_scale_n, bias_op) : nullptr; + const ggml_tensor * gate_bias = with_bias ? get_bias_tensor(gate_out_n, gate_scale_n, bias_op) : nullptr; + if (with_bias && (!ggml_are_same_shape(gate_out_n->src[0], gate_out_n->src[1]) || + !ggml_are_same_shape(up_out_n->src[0], up_out_n->src[1]))) { + continue; + } + + const ggml_tensor * src0 = up_n->src[0]; + const ggml_tensor * src1 = up_n->src[1]; + const ggml_tensor * ids = up_n->src[2]; + + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.gate = gate_n->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); + + if (ggml_cuda_should_fuse_mul_mat_vec_q(up_n)) { + ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, cgraph->nodes[glu_idx], &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = n_ops; + break; + } + } + + if (fused_mul_mat_vec) { + break; + } + } else { + for (const bool with_bias : { false, true }) { + const int gate_idx = i; + const int gate_scale_idx = i + 4; + const int gate_bias_idx = with_bias ? i + 5 : -1; + const int up_idx = with_bias ? i + 6 : i + 5; + const int up_scale_idx = up_idx + 4; + const int up_bias_idx = with_bias ? up_idx + 5 : -1; + const int glu_idx = with_bias ? up_idx + 6 : up_idx + 5; + + const int out_nodes[] = { glu_idx }; + ggml_op ops[13]; + if (with_bias) { + ops[0] = op; + ops[1] = GGML_OP_RESHAPE; + ops[2] = GGML_OP_REPEAT; + ops[3] = GGML_OP_GET_ROWS; + ops[4] = GGML_OP_MUL; + ops[5] = bias_op; + ops[6] = op; + ops[7] = GGML_OP_RESHAPE; + ops[8] = GGML_OP_REPEAT; + ops[9] = GGML_OP_GET_ROWS; + ops[10] = GGML_OP_MUL; + ops[11] = bias_op; + ops[12] = GGML_OP_GLU; + } else { + ops[0] = op; + ops[1] = GGML_OP_RESHAPE; + ops[2] = GGML_OP_REPEAT; + ops[3] = GGML_OP_GET_ROWS; + ops[4] = GGML_OP_MUL; + ops[5] = op; + ops[6] = GGML_OP_RESHAPE; + ops[7] = GGML_OP_REPEAT; + ops[8] = GGML_OP_GET_ROWS; + ops[9] = GGML_OP_MUL; + ops[10] = GGML_OP_GLU; + } + const int n_ops = with_bias ? 13 : 11; + + if (!ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1) || + !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1)) { + continue; + } + + ggml_tensor * gate_n = cgraph->nodes[gate_idx]; + ggml_tensor * gate_scale_n = cgraph->nodes[gate_scale_idx]; + ggml_tensor * gate_out_n = with_bias ? cgraph->nodes[gate_bias_idx] : gate_scale_n; + ggml_tensor * up_n = cgraph->nodes[up_idx]; + ggml_tensor * up_scale_n = cgraph->nodes[up_scale_idx]; + ggml_tensor * up_out_n = with_bias ? cgraph->nodes[up_bias_idx] : up_scale_n; + const ggml_tensor * glu = cgraph->nodes[glu_idx]; + + if (!ggml_cuda_should_fuse_mul_mat(up_n, gate_n, glu, + with_bias ? up_out_n : nullptr, with_bias ? gate_out_n : nullptr, up_scale_n, gate_scale_n)) { + continue; + } + + const ggml_tensor * gate_scale = get_mul_mat_id_scale(cgraph->nodes[gate_idx + 1], cgraph->nodes[gate_idx + 2], + cgraph->nodes[gate_idx + 3], gate_scale_n, gate_n); + const ggml_tensor * up_scale = get_mul_mat_id_scale(cgraph->nodes[up_idx + 1], cgraph->nodes[up_idx + 2], + cgraph->nodes[up_idx + 3], up_scale_n, up_n); + if (!gate_scale || !up_scale) { + continue; + } + + const ggml_tensor * up_bias = with_bias ? get_bias_tensor(up_out_n, up_scale_n, bias_op) : nullptr; + const ggml_tensor * gate_bias = with_bias ? get_bias_tensor(gate_out_n, gate_scale_n, bias_op) : nullptr; + + const ggml_tensor * src0 = up_n->src[0]; + const ggml_tensor * src1 = up_n->src[1]; + const ggml_tensor * ids = up_n->src[2]; + + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.gate = gate_n->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); + + if (ggml_cuda_should_fuse_mul_mat_vec_q(up_n)) { + ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, cgraph->nodes[glu_idx], &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = n_ops; + break; + } + } + + if (fused_mul_mat_vec) { + break; + } + } + + if (ggml_cuda_can_fuse(cgraph, i, { op, bias_op, op, bias_op, GGML_OP_GLU }, {})) { + ggml_tensor * glu = cgraph->nodes[i + 4]; + ggml_tensor * gate_bias_n = glu->src[0]; + ggml_tensor * up_bias_n = glu->src[1]; + + //we don't assume the order for {gate, up}. Instead infer it from the bias tensor + ggml_tensor * gate_n = nullptr; + ggml_tensor * up_n = nullptr; + + if (gate_bias_n->src[0] == cgraph->nodes[i] || gate_bias_n->src[1] == cgraph->nodes[i]) { + gate_n = cgraph->nodes[i]; + up_n = cgraph->nodes[i + 2]; + } else if (gate_bias_n->src[0] == cgraph->nodes[i + 2] || gate_bias_n->src[1] == cgraph->nodes[i + 2]) { + gate_n = cgraph->nodes[i + 2]; + up_n = cgraph->nodes[i]; + } else { + continue; + } + + const ggml_tensor * up_bias_tensor = get_bias_tensor(up_bias_n, up_n, bias_op); + const ggml_tensor * gate_bias_tensor = get_bias_tensor(gate_bias_n, gate_n, bias_op); + + if (!up_bias_tensor || !gate_bias_tensor) { + continue; + } + + // we don't support repeating adds + if (bias_op == GGML_OP_ADD && (!ggml_are_same_shape(gate_bias_n->src[0], gate_bias_n->src[1]) || + !ggml_are_same_shape(up_bias_n->src[0], up_bias_n->src[1]))) { + continue; + } + + const ggml_tensor * src0 = up_n->src[0]; + const ggml_tensor * src1 = up_n->src[1]; + const ggml_tensor * ids = up_n->src[2]; + + if (ggml_cuda_should_fuse_mul_mat_vec_f(up_n)) { + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.gate = gate_n->src[0]; + fusion_data.x_bias = up_bias_tensor; + fusion_data.gate_bias = gate_bias_tensor; + fusion_data.glu_op = ggml_get_glu_op(glu); + + ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, glu, &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = 5; + break; + } + + if (ggml_cuda_should_fuse_mul_mat_vec_q(up_n)) { + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.gate = gate_n->src[0]; + fusion_data.x_bias = up_bias_tensor; + fusion_data.gate_bias = gate_bias_tensor; + fusion_data.glu_op = ggml_get_glu_op(glu); + + ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, glu, &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = 5; + break; + } + } else if (ggml_cuda_can_fuse(cgraph, i, { op, op, GGML_OP_GLU }, {})) { + ggml_tensor * glu = cgraph->nodes[i + 2]; + ggml_tensor * gate = glu->src[0]; + ggml_tensor * up = glu->src[1]; + + bool ok = (gate == cgraph->nodes[i] && up == cgraph->nodes[i + 1]) || + (gate == cgraph->nodes[i + 1] && up == cgraph->nodes[i]); + + if (!ok) { + continue; + } + + const ggml_tensor * src0 = up->src[0]; + const ggml_tensor * src1 = up->src[1]; + const ggml_tensor * ids = up->src[2]; + + if (ggml_cuda_should_fuse_mul_mat_vec_f(up)) { + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.gate = gate->src[0]; + fusion_data.glu_op = ggml_get_glu_op(glu); + + ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, glu, &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = 3; + break; + } + + if (ggml_cuda_should_fuse_mul_mat_vec_q(up)) { + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.gate = gate->src[0]; + fusion_data.glu_op = ggml_get_glu_op(glu); + + ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, glu, &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = 3; + break; + } + } + } + + if (fused_mul_mat_vec) { + return fused_node_count - 1; + } + + fused_mul_mat_vec = false; + fused_node_count = 0; + + // mul_mat + scale + optional bias + for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) { + const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID; + + for (const bool with_bias : { false, true }) { + const int n_ops = op == GGML_OP_MUL_MAT ? (with_bias ? 3 : 2) : (with_bias ? 6 : 5); + const int out_nodes[] = { i + n_ops - 1 }; + ggml_op ops[6]; + if (op == GGML_OP_MUL_MAT) { + if (with_bias) { + ops[0] = op; + ops[1] = GGML_OP_MUL; + ops[2] = bias_op; + } else { + ops[0] = op; + ops[1] = GGML_OP_MUL; + } + } else { + if (with_bias) { + ops[0] = op; + ops[1] = GGML_OP_RESHAPE; + ops[2] = GGML_OP_REPEAT; + ops[3] = GGML_OP_GET_ROWS; + ops[4] = GGML_OP_MUL; + ops[5] = bias_op; + } else { + ops[0] = op; + ops[1] = GGML_OP_RESHAPE; + ops[2] = GGML_OP_REPEAT; + ops[3] = GGML_OP_GET_ROWS; + ops[4] = GGML_OP_MUL; + } + } + + if (!ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1) || + !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1)) { + continue; + } + + ggml_tensor * mm_node = cgraph->nodes[i]; + ggml_tensor * scale_node = op == GGML_OP_MUL_MAT ? cgraph->nodes[i + 1] : cgraph->nodes[i + 4]; + ggml_tensor * out_node = with_bias ? cgraph->nodes[i + n_ops - 1] : scale_node; + + const ggml_tensor * scale = nullptr; + if (op == GGML_OP_MUL_MAT) { + scale = get_mul_mat_scale(scale_node, mm_node); + } else { + scale = get_mul_mat_id_scale(cgraph->nodes[i + 1], cgraph->nodes[i + 2], cgraph->nodes[i + 3], scale_node, mm_node); + } + if (!scale) { + continue; + } + + const ggml_tensor * bias = with_bias ? get_bias_tensor(out_node, scale_node, bias_op) : nullptr; + if (with_bias && !bias) { + continue; + } + if (with_bias && bias_op == GGML_OP_ADD && !ggml_are_same_shape(out_node->src[0], out_node->src[1])) { + continue; + } + if (with_bias && bias_op == GGML_OP_ADD_ID && out_node->src[2] != mm_node->src[2]) { + continue; + } + + const ggml_tensor * src0 = mm_node->src[0]; + const ggml_tensor * src1 = mm_node->src[1]; + const ggml_tensor * ids = mm_node->src[2]; + + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.x_bias = bias; + fusion_data.x_scale = scale; + + if (ggml_cuda_should_fuse_mul_mat_vec_q(mm_node)) { + ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, out_node, &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = n_ops; + break; + } + } + if (fused_mul_mat_vec) { + break; + } + } + + if (fused_mul_mat_vec) { + return fused_node_count - 1; + } + + // mul_mat + add + for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) { + const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID; + + if (!ggml_can_fuse(cgraph, i, { op, bias_op })) { + continue; + } + + ggml_tensor * mm_node = cgraph->nodes[i]; + ggml_tensor * bias_node = cgraph->nodes[i + 1]; + + ggml_tensor * bias_tensor = nullptr; + if (bias_op == GGML_OP_ADD) { + if (bias_node->src[0] == mm_node) { + bias_tensor = bias_node->src[1]; + } else if (bias_node->src[1] == mm_node) { + bias_tensor = bias_node->src[0]; + } else { + continue; + } + } else { + if (bias_node->src[0] != mm_node) { + continue; + } + bias_tensor = bias_node->src[1]; + } + + const ggml_tensor * src0 = mm_node->src[0]; + const ggml_tensor * src1 = mm_node->src[1]; + const ggml_tensor * ids = mm_node->src[2]; + + if (bias_op == GGML_OP_ADD_ID && bias_node->src[2] != ids) { + continue; + } + + if (bias_op == GGML_OP_ADD && !ggml_are_same_shape(bias_node->src[0], bias_node->src[1])) { + continue; + } + + ggml_cuda_mm_fusion_args_host fusion_data{}; + fusion_data.x_bias = bias_tensor; + + if (ggml_cuda_should_fuse_mul_mat_vec_f(mm_node)) { + ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, bias_node, &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = 2; + break; + } + + if (ggml_cuda_should_fuse_mul_mat_vec_q(mm_node)) { + ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, bias_node, &fusion_data); + fused_mul_mat_vec = true; + fused_node_count = 2; + break; + } } - // Single-lane MM fusion. The lane is MUL_MAT[_ID] + optional scale + optional bias. - int fused_mm_lane_nodes = ggml_cuda_try_fuse_mm_lane(cuda_ctx, cgraph, i); - if (fused_mm_lane_nodes > 0) { - return fused_mm_lane_nodes; + if (fused_mul_mat_vec) { + return fused_node_count - 1; } if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ADD }, {})) { @@ -4411,6 +4725,12 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud } } +#ifdef GGML_CUDA_DEBUG + const int nodes_fused = i - prev_i - 1; + if (nodes_fused > 0) { + GGML_LOG_INFO("nodes_fused: %d\n", nodes_fused); + } +#endif prev_i = i; if (ggml_is_empty(node) || node->op == GGML_OP_RESHAPE || node->op == GGML_OP_TRANSPOSE || node->op == GGML_OP_VIEW || node->op == GGML_OP_PERMUTE || node->op == GGML_OP_NONE) { @@ -4424,12 +4744,6 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud int nodes_to_skip = ggml_cuda_try_fuse(cuda_ctx, cgraph, i); if (nodes_to_skip != 0) { -#ifdef GGML_CUDA_DEBUG - const int last_fused = i + nodes_to_skip; - GGML_LOG_INFO("nodes_fused: %d, first: %s (%s), last: %s (%s)\n", - nodes_to_skip + 1, ggml_op_name(node->op), node->name, - ggml_op_name(cgraph->nodes[last_fused]->op), cgraph->nodes[last_fused]->name); -#endif i += nodes_to_skip; continue; } From 0905129e9d12e2bc6f16d6d3cc4e6b40606fc893 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 26 Jun 2026 18:03:58 +0200 Subject: [PATCH 20/28] Apply suggestions from code review Co-authored-by: Georgi Gerganov --- tests/test-backend-ops.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index fee80c47487e..43e3a3dc6c1a 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5826,12 +5826,12 @@ struct test_mul_mat_vec_fusion : public test_case { return out; } - ggml_tensor * build_dense_lane_scale(ggml_context * ctx, ggml_tensor * out) { + ggml_tensor * build_lane_scale_dense(ggml_context * ctx, ggml_tensor * out) { ggml_tensor * scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 1); return ggml_mul(ctx, out, scale); } - ggml_tensor * build_id_lane_scale(ggml_context * ctx, ggml_tensor * out, ggml_tensor * ids) { + ggml_tensor * build_lane_scale_id(ggml_context * ctx, ggml_tensor * out, ggml_tensor * ids) { ggml_tensor * scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_mats); ggml_tensor * s = ggml_reshape_3d(ctx, scale, 1, n_mats, 1); s = ggml_repeat_4d(ctx, s, 1, n_mats, m, 1); @@ -5850,7 +5850,7 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * gate = with_gate ? ggml_new_tensor(ctx, type, 4, ne0.data()) : nullptr; ggml_tensor * up = ggml_new_tensor(ctx, type, 4, ne0.data()); - auto build_up_lane = [&]() { + auto build_lane_up = [&]() { ggml_tensor * ffn_up = ggml_mul_mat(ctx, up, cur); if (with_lane_scale) { ffn_up = build_dense_lane_scale(ctx, ffn_up); @@ -5863,7 +5863,7 @@ struct test_mul_mat_vec_fusion : public test_case { return ffn_up; }; - auto build_gate_lane = [&]() { + auto build_lane_gate = [&]() { ggml_tensor * ffn_gate = ggml_mul_mat(ctx, gate, cur); if (with_lane_scale) { ffn_gate = build_dense_lane_scale(ctx, ffn_gate); From 637bafd8304266898defb27e94e4a49409a4ba03 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 26 Jun 2026 18:03:09 +0200 Subject: [PATCH 21/28] Rename ggml_is_constant_view_src to ggml_is_constant --- ggml/src/ggml.c | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index f96445956933..88fb931a745e 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -7371,7 +7371,7 @@ static int ggml_node_list_find_tensor(const struct ggml_cgraph * cgraph, return -1; } -static bool ggml_is_constant_view_src(const struct ggml_tensor * tensor) { +static bool ggml_is_constant(const struct ggml_tensor * tensor) { return tensor->buffer != NULL && ggml_backend_buffer_get_usage(tensor->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS; } @@ -7424,7 +7424,7 @@ bool ggml_can_fuse_subgraph_ext(const struct ggml_cgraph * cgraph, // external view sources are allowed only for weight tensors, which are constant for this graph execution. struct ggml_tensor * view_src = node->view_src; while (view_src) { - if (ggml_node_list_find_tensor(cgraph, node_idxs, count, view_src) == -1 && !ggml_is_constant_view_src(view_src)) { + if (ggml_node_list_find_tensor(cgraph, node_idxs, count, view_src) == -1 && !ggml_is_constant(view_src)) { return false; } view_src = view_src->view_src; From b46e632804345874057367324c96254cedf57a9f Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 26 Jun 2026 18:15:57 +0200 Subject: [PATCH 22/28] Finish renaming of 0905129e9d12e2bc6f16d6d3cc4e6b40606fc893 --- tests/test-backend-ops.cpp | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 43e3a3dc6c1a..cefe9a028c09 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5853,7 +5853,7 @@ struct test_mul_mat_vec_fusion : public test_case { auto build_lane_up = [&]() { ggml_tensor * ffn_up = ggml_mul_mat(ctx, up, cur); if (with_lane_scale) { - ffn_up = build_dense_lane_scale(ctx, ffn_up); + ffn_up = build_lane_scale_dense(ctx, ffn_up); } if (with_bias) { std::array bias_ne = { ffn_up->ne[0], 1, channels, samples }; @@ -5866,7 +5866,7 @@ struct test_mul_mat_vec_fusion : public test_case { auto build_lane_gate = [&]() { ggml_tensor * ffn_gate = ggml_mul_mat(ctx, gate, cur); if (with_lane_scale) { - ffn_gate = build_dense_lane_scale(ctx, ffn_gate); + ffn_gate = build_lane_scale_dense(ctx, ffn_gate); } if (with_bias) { std::array bias_ne = { ffn_gate->ne[0], 1, channels, samples }; @@ -5876,8 +5876,8 @@ struct test_mul_mat_vec_fusion : public test_case { return ffn_gate; }; - ggml_tensor * ffn_up = build_up_lane(); - ggml_tensor * ffn_gate = with_gate ? build_gate_lane() : nullptr; + ggml_tensor * ffn_up = build_lane_up(); + ggml_tensor * ffn_gate = with_gate ? build_lane_gate() : nullptr; ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; @@ -5899,10 +5899,10 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * cur = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, k, this->b ? 1 : n_used, m); ggml_set_name(cur, "cur"); - auto build_up_lane = [&]() { + auto build_lane_up = [&]() { ggml_tensor * ffn_up = ggml_mul_mat_id(ctx, ups, cur, ids); if (with_lane_scale) { - ffn_up = build_id_lane_scale(ctx, ffn_up, ids); + ffn_up = build_lane_scale_id(ctx, ffn_up, ids); } if (with_bias) { ggml_tensor * up_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_up->ne[0], n_mats); @@ -5911,10 +5911,10 @@ struct test_mul_mat_vec_fusion : public test_case { return ffn_up; }; - auto build_gate_lane = [&]() { + auto build_lane_gate = [&]() { ggml_tensor * ffn_gate = ggml_mul_mat_id(ctx, gates, cur, ids); if (with_lane_scale) { - ffn_gate = build_id_lane_scale(ctx, ffn_gate, ids); + ffn_gate = build_lane_scale_id(ctx, ffn_gate, ids); } if (with_bias) { ggml_tensor * gate_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_gate->ne[0], n_mats); @@ -5923,8 +5923,8 @@ struct test_mul_mat_vec_fusion : public test_case { return ffn_gate; }; - ggml_tensor * ffn_up = build_up_lane(); - ggml_tensor * ffn_gate = with_gate ? build_gate_lane() : nullptr; + ggml_tensor * ffn_up = build_lane_up(); + ggml_tensor * ffn_gate = with_gate ? build_lane_gate() : nullptr; ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up; From 30422a8fec1fd9e9931eab6f2b6c198b3681460f Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 26 Jun 2026 18:21:50 +0200 Subject: [PATCH 23/28] Readd descriptive prints for fusion debugging --- ggml/src/ggml-cuda/ggml-cuda.cu | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 571e09c5d9a9..45eaa06dc845 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4725,12 +4725,6 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud } } -#ifdef GGML_CUDA_DEBUG - const int nodes_fused = i - prev_i - 1; - if (nodes_fused > 0) { - GGML_LOG_INFO("nodes_fused: %d\n", nodes_fused); - } -#endif prev_i = i; if (ggml_is_empty(node) || node->op == GGML_OP_RESHAPE || node->op == GGML_OP_TRANSPOSE || node->op == GGML_OP_VIEW || node->op == GGML_OP_PERMUTE || node->op == GGML_OP_NONE) { @@ -4744,6 +4738,12 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud int nodes_to_skip = ggml_cuda_try_fuse(cuda_ctx, cgraph, i); if (nodes_to_skip != 0) { +#ifdef GGML_CUDA_DEBUG + const int last_fused = i + nodes_to_skip; + GGML_LOG_INFO("nodes_fused: %d, first: %s (%s), last: %s (%s)\n", + nodes_to_skip + 1, ggml_op_name(node->op), node->name, + ggml_op_name(cgraph->nodes[last_fused]->op), cgraph->nodes[last_fused]->name); +#endif i += nodes_to_skip; continue; } From 764f9daa4974b553629e37642c1d6ddc2d614ee3 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 26 Jun 2026 18:51:45 +0200 Subject: [PATCH 24/28] Add weight-buffer pre-allocation to `test_case` This is required so we correctly test fusion of NVFP4. --- tests/test-backend-ops.cpp | 60 ++++++++++++++++++++++++++++++++++++++ 1 file changed, 60 insertions(+) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index cefe9a028c09..8a8f978b5e2f 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -1213,6 +1213,39 @@ struct test_case { virtual bool run_whole_graph() { return false; } virtual std::vector fusion_test_nodes() { return {}; } + virtual std::vector weight_tensors() { return {}; } + + void free_weight_buffers(std::vector & buffers) { + for (ggml_backend_buffer_t buffer : buffers) { + ggml_backend_buffer_free(buffer); + } + buffers.clear(); + } + + bool allocate_weight_tensors(ggml_backend_t backend, std::vector & buffers) { + for (ggml_tensor * tensor : weight_tensors()) { + ggml_backend_buffer_type_t buft = ggml_backend_get_default_buffer_type(backend); + const size_t size = ggml_backend_buft_get_alloc_size(buft, tensor); + + ggml_backend_buffer_t buffer = ggml_backend_buft_alloc_buffer(buft, size); + if (buffer == nullptr) { + free_weight_buffers(buffers); + return false; + } + ggml_backend_buffer_set_usage(buffer, GGML_BACKEND_BUFFER_USAGE_WEIGHTS); + + enum ggml_status status = ggml_backend_tensor_alloc(buffer, tensor, ggml_backend_buffer_get_base(buffer)); + if (status != GGML_STATUS_SUCCESS) { + ggml_backend_buffer_free(buffer); + free_weight_buffers(buffers); + return false; + } + + buffers.push_back(buffer); + } + + return true; + } ggml_cgraph * gf = nullptr; ggml_cgraph * gb = nullptr; @@ -1362,11 +1395,19 @@ struct test_case { // post-graph sentinel add_sentinel(ctx); + std::vector weight_buffers; + if (!allocate_weight_tensors(backend1, weight_buffers)) { + printf("failed to allocate weight tensors [%s] ", ggml_backend_name(backend1)); + ggml_free(ctx); + return test_status_t::FAIL; + } + // allocate ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backend1); if (buf == NULL) { printf("failed to allocate tensors [%s] ", ggml_backend_name(backend1)); + free_weight_buffers(weight_buffers); ggml_free(ctx); return test_status_t::FAIL; } @@ -1466,6 +1507,7 @@ struct test_case { fused_nodes_to_verify.size()); ggml_backend_buffer_free(buf); + free_weight_buffers(weight_buffers); ggml_free(ctx); @@ -1510,11 +1552,18 @@ struct test_case { return true; } + std::vector weight_buffers; + if (!allocate_weight_tensors(backend, weight_buffers)) { + printf("failed to allocate weight tensors\n"); + return false; + } + // allocate ggml_backend_buffer_ptr buf(ggml_backend_alloc_ctx_tensors(ctx.get(), backend)); // smart ptr if (buf == NULL) { printf("failed to allocate tensors\n"); + free_weight_buffers(weight_buffers); return false; } @@ -1584,6 +1633,7 @@ struct test_case { ggml_status status = ggml_backend_graph_compute(backend, gf); if (status != GGML_STATUS_SUCCESS) { fprintf(stderr, "%s: ggml_backend_graph_compute failed. status=%s \n", __func__, ggml_status_to_string(status)); + free_weight_buffers(weight_buffers); return false; } int64_t end_time = ggml_time_us(); @@ -1593,6 +1643,8 @@ struct test_case { total_runs += n_runs; } while (total_time_us < 1000*1000); // run for at least 1 second + free_weight_buffers(weight_buffers); + // Create test result double avg_time_us = (double) total_time_us / total_runs; double calculated_flops = (op_flops(out) > 0) ? (op_flops(out) * total_runs) / (total_time_us / 1e6) : 0.0; @@ -5790,6 +5842,7 @@ struct test_mul_mat_vec_fusion : public test_case { const bool with_gate; const bool with_lane_scale; std::array batch_dims; + std::vector lane_scale_weights; test_mul_mat_vec_fusion(ggml_type type, ggml_glu_op op, int64_t m, int64_t n, int64_t k, bool use_id = false, int n_mats = 1, int n_used = 1, bool b = false, bool with_bias = false, bool with_gate = true, @@ -5833,6 +5886,7 @@ struct test_mul_mat_vec_fusion : public test_case { ggml_tensor * build_lane_scale_id(ggml_context * ctx, ggml_tensor * out, ggml_tensor * ids) { ggml_tensor * scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_mats); + lane_scale_weights.push_back(scale); ggml_tensor * s = ggml_reshape_3d(ctx, scale, 1, n_mats, 1); s = ggml_repeat_4d(ctx, s, 1, n_mats, m, 1); s = ggml_get_rows(ctx, s, ids); @@ -5840,6 +5894,8 @@ struct test_mul_mat_vec_fusion : public test_case { } ggml_tensor * build_graph(ggml_context * ctx) override { + lane_scale_weights.clear(); + if (!use_id) { const int channels = batch_dims[0]; const int samples = batch_dims[1]; @@ -5937,6 +5993,10 @@ struct test_mul_mat_vec_fusion : public test_case { } } + std::vector weight_tensors() override { + return lane_scale_weights; + } + void initialize_tensors(ggml_context * ctx) override { if (!use_id) { for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) { From d5a82727c8f3b9e617fe0cb94ad6406f250449e2 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Thu, 2 Jul 2026 17:38:28 +0200 Subject: [PATCH 25/28] Update ggml/src/ggml.c MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Johannes Gäßler --- ggml/src/ggml.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 88fb931a745e..de61d2504b58 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -7372,7 +7372,7 @@ static int ggml_node_list_find_tensor(const struct ggml_cgraph * cgraph, } static bool ggml_is_constant(const struct ggml_tensor * tensor) { - return tensor->buffer != NULL && ggml_backend_buffer_get_usage(tensor->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS; + return tensor->buffer != NULL && ggml_backend_buffer_get_usage(tensor->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && (tensor->flags & GGML_TENSOR_FLAG_PARAM) == 0; } bool ggml_can_fuse_subgraph_ext(const struct ggml_cgraph * cgraph, From e9f8b33a9c33fd0a3ee61ab40b2cb74024ac5973 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 3 Jul 2026 14:14:43 +0200 Subject: [PATCH 26/28] Add 2nd context for weights as suggested by @JohannesGaessler This reflects more natural use of ggml compared to artifically pre-allocating weights into the same context --- tests/test-backend-ops.cpp | 131 +++++++++++++++++++------------------ 1 file changed, 69 insertions(+), 62 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 8a8f978b5e2f..d012f249f637 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -1137,6 +1137,10 @@ struct test_case { } virtual ggml_tensor * build_graph(ggml_context * ctx) = 0; + virtual ggml_tensor * build_graph(ggml_context * ctx, ggml_context * ctx_weights) { + GGML_UNUSED(ctx_weights); + return build_graph(ctx); + } virtual double max_nmse_err() { return 1e-7; @@ -1213,39 +1217,7 @@ struct test_case { virtual bool run_whole_graph() { return false; } virtual std::vector fusion_test_nodes() { return {}; } - virtual std::vector weight_tensors() { return {}; } - - void free_weight_buffers(std::vector & buffers) { - for (ggml_backend_buffer_t buffer : buffers) { - ggml_backend_buffer_free(buffer); - } - buffers.clear(); - } - - bool allocate_weight_tensors(ggml_backend_t backend, std::vector & buffers) { - for (ggml_tensor * tensor : weight_tensors()) { - ggml_backend_buffer_type_t buft = ggml_backend_get_default_buffer_type(backend); - const size_t size = ggml_backend_buft_get_alloc_size(buft, tensor); - - ggml_backend_buffer_t buffer = ggml_backend_buft_alloc_buffer(buft, size); - if (buffer == nullptr) { - free_weight_buffers(buffers); - return false; - } - ggml_backend_buffer_set_usage(buffer, GGML_BACKEND_BUFFER_USAGE_WEIGHTS); - - enum ggml_status status = ggml_backend_tensor_alloc(buffer, tensor, ggml_backend_buffer_get_base(buffer)); - if (status != GGML_STATUS_SUCCESS) { - ggml_backend_buffer_free(buffer); - free_weight_buffers(buffers); - return false; - } - - buffers.push_back(buffer); - } - - return true; - } + virtual bool use_weight_context() { return false; } ggml_cgraph * gf = nullptr; ggml_cgraph * gb = nullptr; @@ -1352,20 +1324,30 @@ struct test_case { /* .mem_base = */ NULL, /* .no_alloc = */ true, }; + const bool use_weights = use_weight_context(); + ggml_context * ctx = ggml_init(params); GGML_ASSERT(ctx); + ggml_context * ctx_weights = use_weights ? ggml_init(params) : nullptr; + GGML_ASSERT(!use_weights || ctx_weights); gf = ggml_new_graph(ctx); // pre-graph sentinel add_sentinel(ctx); + if (ctx_weights) { + add_sentinel(ctx_weights); + } - ggml_tensor * out = build_graph(ctx); + ggml_tensor * out = build_graph(ctx, ctx_weights); current_op_name = op_desc(out); check_for_f16_tensor(ctx); if (!matches_filter(out, op_names_filter)) { //printf(" %s: skipping\n", op_desc(out).c_str()); + if (ctx_weights) { + ggml_free(ctx_weights); + } ggml_free(ctx); return test_status_t::SKIPPED; } @@ -1388,18 +1370,29 @@ struct test_case { print_test_result_locked(output_printer, result); + if (ctx_weights) { + ggml_free(ctx_weights); + } ggml_free(ctx); return test_status_t::NOT_SUPPORTED; } // post-graph sentinel add_sentinel(ctx); + if (ctx_weights) { + add_sentinel(ctx_weights); + } - std::vector weight_buffers; - if (!allocate_weight_tensors(backend1, weight_buffers)) { - printf("failed to allocate weight tensors [%s] ", ggml_backend_name(backend1)); - ggml_free(ctx); - return test_status_t::FAIL; + ggml_backend_buffer_t buf_weights = nullptr; + if (ctx_weights) { + buf_weights = ggml_backend_alloc_ctx_tensors(ctx_weights, backend1); + if (buf_weights == NULL) { + printf("failed to allocate weight tensors [%s] ", ggml_backend_name(backend1)); + ggml_free(ctx_weights); + ggml_free(ctx); + return test_status_t::FAIL; + } + ggml_backend_buffer_set_usage(buf_weights, GGML_BACKEND_BUFFER_USAGE_WEIGHTS); } // allocate @@ -1407,7 +1400,10 @@ struct test_case { if (buf == NULL) { printf("failed to allocate tensors [%s] ", ggml_backend_name(backend1)); - free_weight_buffers(weight_buffers); + if (ctx_weights) { + ggml_backend_buffer_free(buf_weights); + ggml_free(ctx_weights); + } ggml_free(ctx); return test_status_t::FAIL; } @@ -1422,6 +1418,9 @@ struct test_case { // randomize tensors initialize_tensors(ctx); + if (ctx_weights) { + initialize_tensors(ctx_weights); + } // compare struct callback_userdata { @@ -1507,8 +1506,10 @@ struct test_case { fused_nodes_to_verify.size()); ggml_backend_buffer_free(buf); - free_weight_buffers(weight_buffers); - + if (ctx_weights) { + ggml_backend_buffer_free(buf_weights); + ggml_free(ctx_weights); + } ggml_free(ctx); // Create test result @@ -1532,10 +1533,14 @@ struct test_case { /* .mem_base = */ NULL, /* .no_alloc = */ true, }; + const bool use_weights = use_weight_context(); + ggml_context_ptr ctx(ggml_init(params)); // smart ptr GGML_ASSERT(ctx); + ggml_context_ptr ctx_weights(use_weights ? ggml_init(params) : nullptr); + GGML_ASSERT(!use_weights || ctx_weights); - ggml_tensor * out = build_graph(ctx.get()); + ggml_tensor * out = build_graph(ctx.get(), ctx_weights.get()); current_op_name = op_desc(out); if (!matches_filter(out, op_names_filter)) { //printf(" %s: skipping\n", op_desc(out).c_str()); @@ -1552,10 +1557,14 @@ struct test_case { return true; } - std::vector weight_buffers; - if (!allocate_weight_tensors(backend, weight_buffers)) { - printf("failed to allocate weight tensors\n"); - return false; + ggml_backend_buffer_ptr buf_weights(nullptr); + if (ctx_weights) { + buf_weights.reset(ggml_backend_alloc_ctx_tensors(ctx_weights.get(), backend)); + if (buf_weights == NULL) { + printf("failed to allocate weight tensors\n"); + return false; + } + ggml_backend_buffer_set_usage(buf_weights.get(), GGML_BACKEND_BUFFER_USAGE_WEIGHTS); } // allocate @@ -1563,12 +1572,14 @@ struct test_case { if (buf == NULL) { printf("failed to allocate tensors\n"); - free_weight_buffers(weight_buffers); return false; } // randomize tensors initialize_tensors(ctx.get()); + if (ctx_weights) { + initialize_tensors(ctx_weights.get()); + } // build graph ggml_cgraph * gf = ggml_new_graph_custom(ctx.get(), graph_nodes, false); @@ -1633,7 +1644,6 @@ struct test_case { ggml_status status = ggml_backend_graph_compute(backend, gf); if (status != GGML_STATUS_SUCCESS) { fprintf(stderr, "%s: ggml_backend_graph_compute failed. status=%s \n", __func__, ggml_status_to_string(status)); - free_weight_buffers(weight_buffers); return false; } int64_t end_time = ggml_time_us(); @@ -1643,8 +1653,6 @@ struct test_case { total_runs += n_runs; } while (total_time_us < 1000*1000); // run for at least 1 second - free_weight_buffers(weight_buffers); - // Create test result double avg_time_us = (double) total_time_us / total_runs; double calculated_flops = (op_flops(out) > 0) ? (op_flops(out) * total_runs) / (total_time_us / 1e6) : 0.0; @@ -5842,7 +5850,6 @@ struct test_mul_mat_vec_fusion : public test_case { const bool with_gate; const bool with_lane_scale; std::array batch_dims; - std::vector lane_scale_weights; test_mul_mat_vec_fusion(ggml_type type, ggml_glu_op op, int64_t m, int64_t n, int64_t k, bool use_id = false, int n_mats = 1, int n_used = 1, bool b = false, bool with_bias = false, bool with_gate = true, @@ -5864,6 +5871,7 @@ struct test_mul_mat_vec_fusion : public test_case { } bool run_whole_graph() override { return true; } + bool use_weight_context() override { return use_id && with_lane_scale; } ggml_tensor * build_gate(ggml_context * ctx, ggml_tensor * ffn_gate, ggml_tensor * ffn_up) { ggml_tensor * out = nullptr; @@ -5884,9 +5892,9 @@ struct test_mul_mat_vec_fusion : public test_case { return ggml_mul(ctx, out, scale); } - ggml_tensor * build_lane_scale_id(ggml_context * ctx, ggml_tensor * out, ggml_tensor * ids) { - ggml_tensor * scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_mats); - lane_scale_weights.push_back(scale); + ggml_tensor * build_lane_scale_id(ggml_context * ctx, ggml_context * ctx_weights, ggml_tensor * out, ggml_tensor * ids) { + GGML_ASSERT(ctx_weights); + ggml_tensor * scale = ggml_new_tensor_1d(ctx_weights, GGML_TYPE_F32, n_mats); ggml_tensor * s = ggml_reshape_3d(ctx, scale, 1, n_mats, 1); s = ggml_repeat_4d(ctx, s, 1, n_mats, m, 1); s = ggml_get_rows(ctx, s, ids); @@ -5894,8 +5902,11 @@ struct test_mul_mat_vec_fusion : public test_case { } ggml_tensor * build_graph(ggml_context * ctx) override { - lane_scale_weights.clear(); + GGML_ASSERT(!use_weight_context()); + return build_graph(ctx, nullptr); + } + ggml_tensor * build_graph(ggml_context * ctx, ggml_context * ctx_weights) override { if (!use_id) { const int channels = batch_dims[0]; const int samples = batch_dims[1]; @@ -5958,7 +5969,7 @@ struct test_mul_mat_vec_fusion : public test_case { auto build_lane_up = [&]() { ggml_tensor * ffn_up = ggml_mul_mat_id(ctx, ups, cur, ids); if (with_lane_scale) { - ffn_up = build_lane_scale_id(ctx, ffn_up, ids); + ffn_up = build_lane_scale_id(ctx, ctx_weights, ffn_up, ids); } if (with_bias) { ggml_tensor * up_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_up->ne[0], n_mats); @@ -5970,7 +5981,7 @@ struct test_mul_mat_vec_fusion : public test_case { auto build_lane_gate = [&]() { ggml_tensor * ffn_gate = ggml_mul_mat_id(ctx, gates, cur, ids); if (with_lane_scale) { - ffn_gate = build_lane_scale_id(ctx, ffn_gate, ids); + ffn_gate = build_lane_scale_id(ctx, ctx_weights, ffn_gate, ids); } if (with_bias) { ggml_tensor * gate_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_gate->ne[0], n_mats); @@ -5993,10 +6004,6 @@ struct test_mul_mat_vec_fusion : public test_case { } } - std::vector weight_tensors() override { - return lane_scale_weights; - } - void initialize_tensors(ggml_context * ctx) override { if (!use_id) { for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) { From 3b991ef6bd23e22779c0a79cd7b2b14d1a1f3eec Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Fri, 3 Jul 2026 14:16:37 +0200 Subject: [PATCH 27/28] Exclude fused tests from gradient mode I'm unsure of the current state, but naively every fusion pattern should require its own backpropagation implementation. I don't see these implemented for the CUDA backend, so we can disable tests to avoid triggering GGML_ASSERT for ggml_tensor * build_graph(ggml_context * ctx) override { GGML_ASSERT(!use_weight_context()); return build_graph(ctx, nullptr); } --- tests/test-backend-ops.cpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index d012f249f637..de4ff1424841 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9786,6 +9786,13 @@ static bool test_backend(ggml_backend_t backend, ggml_backend_dev_t dev, test_mo } if (mode == MODE_GRAD) { + test_cases.erase( + std::remove_if(test_cases.begin(), test_cases.end(), [](const std::unique_ptr & tc) { + return tc->run_whole_graph(); + }), + test_cases.end() + ); + size_t n_ok = 0; for (auto & test : test_cases) { if (test->eval_grad(backend, op_names_filter, output_printer)) { From 216097f509cda8b4ce34d7ad3d479358a741a578 Mon Sep 17 00:00:00 2001 From: Oliver Simons Date: Mon, 6 Jul 2026 09:41:12 +0200 Subject: [PATCH 28/28] Apply suggestions from code review MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Johannes Gäßler --- tests/test-backend-ops.cpp | 20 ++++++-------------- 1 file changed, 6 insertions(+), 14 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index de4ff1424841..8fae79715108 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -1345,9 +1345,7 @@ struct test_case { if (!matches_filter(out, op_names_filter)) { //printf(" %s: skipping\n", op_desc(out).c_str()); - if (ctx_weights) { - ggml_free(ctx_weights); - } + ggml_free(ctx_weights); ggml_free(ctx); return test_status_t::SKIPPED; } @@ -1370,9 +1368,7 @@ struct test_case { print_test_result_locked(output_printer, result); - if (ctx_weights) { - ggml_free(ctx_weights); - } + ggml_free(ctx_weights); ggml_free(ctx); return test_status_t::NOT_SUPPORTED; } @@ -1400,10 +1396,8 @@ struct test_case { if (buf == NULL) { printf("failed to allocate tensors [%s] ", ggml_backend_name(backend1)); - if (ctx_weights) { - ggml_backend_buffer_free(buf_weights); - ggml_free(ctx_weights); - } + ggml_backend_buffer_free(buf_weights); + ggml_free(ctx_weights); ggml_free(ctx); return test_status_t::FAIL; } @@ -1506,10 +1500,8 @@ struct test_case { fused_nodes_to_verify.size()); ggml_backend_buffer_free(buf); - if (ctx_weights) { - ggml_backend_buffer_free(buf_weights); - ggml_free(ctx_weights); - } + ggml_backend_buffer_free(buf_weights); + ggml_free(ctx_weights); ggml_free(ctx); // Create test result