Commit 5ce06fc71 for llama.cpp
commit 5ce06fc71f5e938301f7dc51697a59ae308bd6b7
Author: anujj <ajalota@nvidia.com>
Date: Thu Nov 5 08:17:33 2026 +0530
Title: CUDA: fuse NVFP4 gate/up matmuls and GLU in MMQ (#31296)
* cuda: fuse NVFP4 gate/up and GLU in MMQ
* cuda: honor the requested src1 precision in fused gate/up MMQ
* cuda: share NVFP4 scale checks, skip unused fused FP4 kernels
* cuda: test NVFP4 fused MMQ at W4A8 without scales and with batches, clean up naming
* cuda: check gate_scale before loading the fused NVFP4 gate scale
diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh
index 00476f0fd..bc83ea375 100644
--- a/ggml/src/ggml-cuda/common.cuh
+++ b/ggml/src/ggml-cuda/common.cuh
@@ -1585,6 +1585,7 @@ struct ggml_cuda_mm_fusion_args_host {
const ggml_tensor * gate_scale = nullptr;
ggml_glu_op glu_op;
float glu_limit = 0.0f;
+ ggml_prec prec_src1 = GGML_PREC_UNDEFINED; // src1 precision requested by the fused MUL_MATs (MMQ)
const ggml_tensor * shared_up = nullptr;
const ggml_tensor * shared_gate = nullptr;
ggml_tensor * shared_dst = nullptr;
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index ee28955f4..9c467ce95 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -2151,8 +2151,50 @@ static bool ggml_cuda_match_shared_expert(const ggml_cgraph * graph, int routed_
return true;
}
+// The NVFP4 weight scale that scale_node (MUL) applies to the output of mm_node (MUL_MAT), or nullptr.
+// Shared by ggml_cuda_try_fuse and graph_optimize so that both accept the same patterns.
+static const ggml_tensor * ggml_cuda_get_mul_mat_scale(const ggml_tensor * scale_node, const ggml_tensor * mm_node) {
+ 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;
+}
+
+// The same for MUL_MAT_ID with one scale per expert: RESHAPE, REPEAT and GET_ROWS (by the ids) of the scales, then the MUL.
+static const ggml_tensor * ggml_cuda_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) {
+ 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;
+}
+
// Buffer- and batch-size independent, so graph_optimize produces the same graph topology for every ubatch size (see #29986).
-static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggml_tensor * gate, const ggml_tensor * glu, const int device) {
+// up_scale/gate_scale: the MUL nodes that apply the NVFP4 weight scales (per tensor or per expert) between the MUL_MATs and the GLU.
+static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggml_tensor * gate, const ggml_tensor * glu, const int device,
+ const ggml_tensor * up_scale = nullptr, const ggml_tensor * gate_scale = nullptr) {
const ggml_tensor * x_up = up->src[0];
const ggml_tensor * y = up->src[1];
@@ -2165,7 +2207,12 @@ static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggm
return false;
}
- if (!ggml_cuda_should_fuse_mul_mat(up, gate, glu)) {
+ if (!ggml_cuda_should_fuse_mul_mat(up, gate, glu, nullptr, nullptr, up_scale, gate_scale)) {
+ return false;
+ }
+
+ // The fused kernel runs both MUL_MATs at one src1 precision, see ggml_cuda_mul_mat_q.
+ if (ggml_prec_get_src(up, 1) != ggml_prec_get_src(gate, 1)) {
return false;
}
@@ -2178,7 +2225,8 @@ static bool ggml_cuda_mul_mat_q_fusion_matches(const ggml_tensor * up, const ggm
return turing_mma_available(ggml_cuda_info().devices[device].cc);
}
-static bool ggml_cuda_should_fuse_mul_mat_q(const ggml_tensor * up, const ggml_tensor * gate, const ggml_tensor * glu) {
+static bool ggml_cuda_should_fuse_mul_mat_q(const ggml_tensor * up, const ggml_tensor * gate, const ggml_tensor * glu,
+ const ggml_tensor * up_scale = nullptr, const ggml_tensor * gate_scale = nullptr) {
// A hinted matmul can be fused, but its unfused path may be faster (e.g. FWHT for GGML_HINT_SRC0_IS_HADAMARD).
if (ggml_get_op_params_i32(up, 1) != GGML_HINT_NONE || ggml_get_op_params_i32(gate, 1) != GGML_HINT_NONE) {
return false;
@@ -2190,7 +2238,7 @@ static bool ggml_cuda_should_fuse_mul_mat_q(const ggml_tensor * up, const ggml_t
ggml_nbytes(x) != ggml_backend_buffer_get_alloc_size(x->buffer, x) && x->view_src;
};
const int device = ggml_cuda_get_device();
- if (!ggml_cuda_mul_mat_q_fusion_matches(up, gate, glu, device) || bad_padding_clear(up->src[0]) || bad_padding_clear(gate->src[0])) {
+ if (!ggml_cuda_mul_mat_q_fusion_matches(up, gate, glu, device, up_scale, gate_scale) || bad_padding_clear(up->src[0]) || bad_padding_clear(gate->src[0])) {
return false;
}
@@ -2209,7 +2257,7 @@ static bool ggml_cuda_should_fuse_mul_mat_q(const ggml_tensor * up, const ggml_t
}
// The fused kernel needs one block per tile (no stream-k split), so only fuse if the tiling is efficient.
- return ggml_cuda_mmq_fusion_is_efficient(x_up->type, x_up->ne[1], ncols, nchannels, device);
+ return ggml_cuda_mmq_fusion_is_efficient(x_up->type, ggml_prec_get_src(up, 1), x_up->ne[1], ncols, nchannels, device);
}
static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
@@ -4284,43 +4332,6 @@ 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;
- 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) {
@@ -4387,8 +4398,8 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
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);
+ const ggml_tensor * gate_scale = ggml_cuda_get_mul_mat_scale(gate_scale_n, gate_n);
+ const ggml_tensor * up_scale = ggml_cuda_get_mul_mat_scale(up_scale_n, up_n);
if (!gate_scale || !up_scale) {
continue;
}
@@ -4419,6 +4430,14 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
fused_node_count = n_ops;
break;
}
+
+ if (!with_bias && ggml_cuda_should_fuse_mul_mat_q(up_n, gate_n, glu, up_scale_n, gate_scale_n)) {
+ fusion_data.prec_src1 = ggml_prec_get_src(up_n, 1);
+ ggml_cuda_mul_mat_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) {
@@ -4483,9 +4502,9 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
continue;
}
- const ggml_tensor * gate_scale = get_mul_mat_id_scale(cgraph->nodes[gate_idx + 1], cgraph->nodes[gate_idx + 2],
+ const ggml_tensor * gate_scale = ggml_cuda_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],
+ const ggml_tensor * up_scale = ggml_cuda_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;
@@ -4513,6 +4532,14 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
fused_node_count = n_ops;
break;
}
+
+ if (!with_bias && ggml_cuda_should_fuse_mul_mat_q(up_n, gate_n, glu, up_scale_n, gate_scale_n)) {
+ fusion_data.prec_src1 = ggml_prec_get_src(up_n, 1);
+ ggml_cuda_mul_mat_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) {
@@ -4628,6 +4655,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
fusion_data.gate = gate->src[0];
fusion_data.glu_op = ggml_get_glu_op(glu);
fusion_data.glu_limit = ggml_get_op_params_f32(glu, 3);
+ fusion_data.prec_src1 = ggml_prec_get_src(up, 1);
ggml_cuda_mul_mat_q(*cuda_ctx, src0, src1, ids, glu, &fusion_data);
fused_mul_mat_vec = true;
@@ -4689,9 +4717,9 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
const ggml_tensor * scale = nullptr;
if (op == GGML_OP_MUL_MAT) {
- scale = get_mul_mat_scale(scale_node, mm_node);
+ scale = ggml_cuda_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);
+ scale = ggml_cuda_get_mul_mat_id_scale(cgraph->nodes[i + 1], cgraph->nodes[i + 2], cgraph->nodes[i + 3], scale_node, mm_node);
}
if (!scale) {
continue;
@@ -5242,6 +5270,33 @@ static void ggml_backend_cuda_graph_optimize(ggml_backend_t backend, ggml_cgraph
continue;
}
}
+ // Same with the NVFP4 scales: MUL_MAT(gate), MUL(gate scale), MUL_MAT(up), MUL(up scale), GLU.
+ if (i + 4 < cgraph->n_nodes && cgraph->nodes[i + 4]->op == GGML_OP_GLU) {
+ ggml_tensor * glu = cgraph->nodes[i + 4];
+ ggml_tensor * gate_scale = cgraph->nodes[i + 1];
+ ggml_tensor * up_scale = cgraph->nodes[i + 3];
+ if (glu->src[0] == gate_scale && glu->src[1] == up_scale &&
+ ggml_cuda_mul_mat_q_fusion_matches(cgraph->nodes[i + 2], cgraph->nodes[i], glu, cuda_ctx->device, up_scale, gate_scale) &&
+ ggml_cuda_get_mul_mat_scale(gate_scale, cgraph->nodes[i]) && ggml_cuda_get_mul_mat_scale(up_scale, cgraph->nodes[i + 2])) {
+ add_alloc_deps(i, i + 4);
+ i += 4;
+ continue;
+ }
+ }
+ // Same with per-expert NVFP4 scales: MUL_MAT_ID, RESHAPE, REPEAT, GET_ROWS, MUL for gate and for up, GLU.
+ if (i + 10 < cgraph->n_nodes && cgraph->nodes[i + 10]->op == GGML_OP_GLU) {
+ ggml_tensor * glu = cgraph->nodes[i + 10];
+ ggml_tensor * gate_scale = cgraph->nodes[i + 4];
+ ggml_tensor * up_scale = cgraph->nodes[i + 9];
+ if (glu->src[0] == gate_scale && glu->src[1] == up_scale &&
+ ggml_cuda_mul_mat_q_fusion_matches(cgraph->nodes[i + 5], cgraph->nodes[i], glu, cuda_ctx->device, up_scale, gate_scale) &&
+ ggml_cuda_get_mul_mat_id_scale(cgraph->nodes[i + 1], cgraph->nodes[i + 2], cgraph->nodes[i + 3], gate_scale, cgraph->nodes[i]) &&
+ ggml_cuda_get_mul_mat_id_scale(cgraph->nodes[i + 6], cgraph->nodes[i + 7], cgraph->nodes[i + 8], up_scale, cgraph->nodes[i + 5])) {
+ add_alloc_deps(i, i + 10);
+ i += 10;
+ continue;
+ }
+ }
ggml_cuda_moe_weighted_reduction_match match;
if (ggml_cuda_match_moe_weighted_reduction(cgraph, i, match)) {
diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh
index 4951db72e..8b4c1e8b3 100644
--- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh
+++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh
@@ -1724,10 +1724,8 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
}
}
-template <ggml_type type, int J, bool fallback, ggml_prec prec_src1> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4(
+template <ggml_type type, int J, bool fallback, ggml_prec prec_src1, bool has_fusion> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4(
const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kb0, const int i_max, const int stride) {
- GGML_UNUSED(x_gate);
-
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, prec_src1);
@@ -1755,7 +1753,7 @@ template <ggml_type type, int J, bool fallback, ggml_prec prec_src1> static __de
i = min(i, i_max);
}
- const block_nvfp4 * bxi = (const block_nvfp4 *) x + kb0 + i * stride + kbx;
+ const block_nvfp4 * bxi = ggml_cuda_mmq_x_row<block_nvfp4, has_fusion>(x, x_gate, kb0, i, stride) + kbx;
const uint32_t * __restrict__ src_qs = reinterpret_cast<const uint32_t *>(bxi->qs);
const int kqs = 16 * kbx;
const int ksc = 4 * kbx;
@@ -1782,10 +1780,8 @@ template <ggml_type type, int J, bool fallback, ggml_prec prec_src1> static __de
}
}
-template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4_nvfp4(
+template <ggml_type type, int J, bool fallback, bool has_fusion> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4_nvfp4(
const char * __restrict__ x, const char * __restrict__ x_gate, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
- GGML_UNUSED(x_gate);
-
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q4) / warp_size;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q4);
@@ -1800,7 +1796,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
const int kbx = txi % threads_per_row;
const int row_in_warp = txi / threads_per_row;
- const block_nvfp4 * bxi_base = (const block_nvfp4 *) x + kbx0 + kbx;
uint32_t * x_u32_scale = x_u32 + 64 + kbx;
#pragma unroll
@@ -1811,7 +1806,7 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
i = min(i, i_max);
}
- const block_nvfp4 * bxi = bxi_base + i * stride;
+ const block_nvfp4 * bxi = ggml_cuda_mmq_x_row<block_nvfp4, has_fusion>(x, x_gate, kbx0, i, stride) + kbx;
const uint32_t * src_qs = reinterpret_cast<const uint32_t *>(bxi->qs);
diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu
index 49e0ce905..3d4b01d93 100644
--- a/ggml/src/ggml-cuda/mmq.cu
+++ b/ggml/src/ggml-cuda/mmq.cu
@@ -114,17 +114,17 @@ static ggml_prec ggml_cuda_mmq_get_prec_env() {
// src1 is quantized to Q8_1 unless the FP4 types can use 4-bit activations, in which case they
// default to the native W4A4 instructions on Blackwell.
-static ggml_prec ggml_cuda_mmq_get_prec_src1(const ggml_tensor * src0, const ggml_tensor * dst, const int cc) {
+static ggml_prec ggml_cuda_mmq_get_prec_src1(const ggml_type type, const ggml_prec prec_requested, const int cc) {
static const ggml_prec prec_env = ggml_cuda_mmq_get_prec_env();
ggml_prec prec = prec_env;
if (prec == GGML_PREC_UNDEFINED) {
- prec = (ggml_prec) ggml_get_op_params_i32(dst, 3);
+ prec = prec_requested;
}
// Q4 only for the FP4 types on Blackwell
GGML_ASSERT(prec == GGML_PREC_UNDEFINED || prec == GGML_PREC_Q8 || prec == GGML_PREC_Q4);
- const bool can_use_q4 = (src0->type == GGML_TYPE_NVFP4 || src0->type == GGML_TYPE_MXFP4) && blackwell_mma_available(cc);
+ const bool can_use_q4 = (type == GGML_TYPE_NVFP4 || type == GGML_TYPE_MXFP4) && blackwell_mma_available(cc);
if (prec == GGML_PREC_Q8 || !can_use_q4) {
return GGML_PREC_Q8;
}
@@ -172,11 +172,18 @@ void ggml_cuda_mul_mat_q(
const ggml_tensor * gate = fusion ? fusion->gate : nullptr;
ggml_cuda_mm_fusion_args_device fusion_device = {};
if (fusion) {
- GGML_ASSERT(gate && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale);
+ GGML_ASSERT(gate && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale == !fusion->gate_scale);
GGML_ASSERT(gate->type == src0->type && ggml_are_same_shape(gate, src0) && ggml_are_same_stride(gate, src0));
- fusion_device.gate = gate->data;
- fusion_device.glu_op = fusion->glu_op;
- fusion_device.glu_limit = fusion->glu_limit;
+ // NVFP4 weight scales (only applied for NVFP4): one value per tensor, or one per expert with ids.
+ for (const ggml_tensor * scale : { fusion->x_scale, fusion->gate_scale }) {
+ GGML_ASSERT(!scale || (src0->type == GGML_TYPE_NVFP4 && scale->type == GGML_TYPE_F32 && ggml_is_contiguous(scale) &&
+ ggml_nelements(scale) == (ids ? src0->ne[2] : 1)));
+ }
+ fusion_device.gate = gate->data;
+ fusion_device.x_scale = fusion->x_scale ? fusion->x_scale->data : nullptr;
+ fusion_device.gate_scale = fusion->gate_scale ? fusion->gate_scale->data : nullptr;
+ fusion_device.glu_op = fusion->glu_op;
+ fusion_device.glu_limit = fusion->glu_limit;
}
GGML_TENSOR_BINARY_OP_LOCALS;
@@ -224,9 +231,8 @@ void ggml_cuda_mul_mat_q(
const bool fallback = ggml_cuda_mmq_needs_fallback(ne01);
- // With fusion dst is the GLU output, not a MUL_MAT; the fused types always quantize src1 to Q8_1.
- // TODO support for Q4 precision if both tensors allow it
- const ggml_prec prec_src1 = gate ? GGML_PREC_Q8 : ggml_cuda_mmq_get_prec_src1(src0, dst, cc);
+ // With fusion dst is the GLU output, not a MUL_MAT; the src1 precision requested by both fused MUL_MATs is in the fusion args.
+ const ggml_prec prec_src1 = ggml_cuda_mmq_get_prec_src1(src0->type, gate ? fusion->prec_src1 : ggml_prec_get_src(dst, 1), cc);
const bool use_native_fp4 = prec_src1 == GGML_PREC_Q4;
const size_t y_block_size = use_native_fp4 ? sizeof(block_fp4_mmq) : sizeof(block_q8_1_mmq);
@@ -513,18 +519,20 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
return (!GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
}
-bool ggml_cuda_mmq_fusion_is_efficient(const ggml_type type, const int64_t nrows_x, const int64_t ncols, const int64_t nchannels, const int device) {
+bool ggml_cuda_mmq_fusion_is_efficient(
+ const ggml_type type, const ggml_prec prec_requested, const int64_t nrows_x, const int64_t ncols, const int64_t nchannels, const int device) {
const int cc = ggml_cuda_info().devices[device].cc;
const int nsm = ggml_cuda_info().devices[device].nsm;
const size_t smpbo = ggml_cuda_info().devices[device].smpbo;
- // Same tile width as ggml_cuda_mul_mat_q; the fused kernel has no fallback, quantizes src1 to Q8_1 and writes I/2 output rows per tile.
- const bool fallback = false;
- const int J_best = ggml_cuda_mmq_get_J_best(type, ncols, fallback, cc, smpbo, GGML_PREC_Q8);
+ // Same tile width as ggml_cuda_mul_mat_q; the fused kernel has no fallback and writes I/2 output rows per tile.
+ const bool fallback = false;
+ const ggml_prec prec_src1 = ggml_cuda_mmq_get_prec_src1(type, prec_requested, cc);
+ const int J_best = ggml_cuda_mmq_get_J_best(type, ncols, fallback, cc, smpbo, prec_src1);
if (J_best == 0) {
return false;
}
- const int64_t nrows_tile = ggml_cuda_mmq_get_config(type, J_best, fallback, cc, GGML_PREC_Q8).I / 2;
+ const int64_t nrows_tile = ggml_cuda_mmq_get_config(type, J_best, fallback, cc, prec_src1).I / 2;
const int64_t ntiles = (ncols + J_best - 1)/J_best * ((nrows_x + nrows_tile - 1)/nrows_tile) * nchannels;
// Same condition as launch_mul_mat_q uses to run the stream-k kernel with one block per tile.
const int64_t tiles_nwaves = (ntiles + nsm - 1) / nsm;
diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
index 308968be9..474b19dfe 100644
--- a/ggml/src/ggml-cuda/mmq.cuh
+++ b/ggml/src/ggml-cuda/mmq.cuh
@@ -16,7 +16,7 @@ typedef void (*ggml_cuda_mmq_load_tiles_t)(const char * __restrict__ x, const ch
typedef void (*ggml_cuda_mmq_vec_dot_t)(const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00);
typedef void (*ggml_cuda_mmq_write_back_t)(const float * __restrict__ sum, const int32_t * __restrict__ get_rows_to_sorted,
float * __restrict__ dst, const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max,
- const ggml_glu_op glu_op, const float glu_limit);
+ const ggml_glu_op glu_op, const float glu_limit, const float * __restrict__ x_scale, const float * __restrict__ gate_scale);
enum mmq_q8_1_ds_layout {
MMQ_Q8_1_DS_LAYOUT_D4,
@@ -436,7 +436,8 @@ static __host__ int ggml_cuda_mmq_get_nbytes_shared_x(const ggml_cuda_mmq_config
// Weight types with a fused gate/up/GLU path (see ggml_cuda_mmq_get_util_funcs).
static constexpr __host__ __device__ bool ggml_cuda_mmq_fusion_supported(const ggml_type type) {
- return type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K || type == GGML_TYPE_Q6_K || type == GGML_TYPE_Q8_0 || type == GGML_TYPE_IQ4_XS;
+ return type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K || type == GGML_TYPE_Q6_K || type == GGML_TYPE_Q8_0 || type == GGML_TYPE_IQ4_XS ||
+ type == GGML_TYPE_NVFP4;
}
// Fused gate/up/GLU: the weight tile alternates chunks of this many up and gate rows, so rows r and r + 8 of an MMA tile (same thread) pair up.
@@ -462,8 +463,8 @@ static __device__ __forceinline__ const block_t * ggml_cuda_mmq_x_row(
template <ggml_type type, int J, bool fallback, ggml_prec prec_src1> static __device__ __forceinline__ void ggml_cuda_mmq_write_back_dp4a(
const float * __restrict__ sum, const int32_t * __restrict__ ids_dst, float * __restrict__ dst,
const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max,
- const ggml_glu_op glu_op, const float glu_limit) {
- GGML_UNUSED_VARS(glu_op, glu_limit);
+ const ggml_glu_op glu_op, const float glu_limit, const float * __restrict__ x_scale, const float * __restrict__ gate_scale) {
+ GGML_UNUSED_VARS(glu_op, glu_limit, x_scale, gate_scale);
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size;
@@ -503,11 +504,12 @@ template <ggml_type type, int J, bool fallback, ggml_prec prec_src1> static __de
}
// has_fusion: accumulators l and l + ne/2 hold the up and gate values of one output element, the GLU of the two is written.
+// x_scale/gate_scale: optional NVFP4 scales of the up and gate weights (per tensor, or of the expert of this tile).
template<ggml_type type, int J, bool fallback, ggml_prec prec_src1, bool has_fusion>
static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma(
const float * __restrict__ sum, const int * __restrict__ ids_dst, float * __restrict__ dst,
const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max,
- const ggml_glu_op glu_op, const float glu_limit) {
+ const ggml_glu_op glu_op, const float glu_limit, const float * __restrict__ x_scale, const float * __restrict__ gate_scale) {
#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C;
@@ -523,6 +525,11 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma(
const bool y_scale_used = y_scale != nullptr;
+ // Fused NVFP4: load the weight scales once, 1 without scales.
+ constexpr bool fused_nvfp4 = has_fusion && type == GGML_TYPE_NVFP4;
+ const float scale_up = fused_nvfp4 && x_scale ? *x_scale : 1.0f;
+ const float scale_gate = fused_nvfp4 && gate_scale ? *gate_scale : 1.0f;
+
#pragma unroll
for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) {
#pragma unroll
@@ -543,7 +550,17 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma(
if constexpr (has_fusion) {
const int sum_index = (j0/tile_C::J + n)*tile_C::ne + l;
- dst[ids_dst[j]*stride + i] = ggml_cuda_op_glu_single(glu_op, sum[sum_index + tile_C::ne/2], sum[sum_index], glu_limit);
+ float up = sum[sum_index];
+ float gate = sum[sum_index + tile_C::ne/2];
+ if constexpr (type == GGML_TYPE_NVFP4) {
+ if (y_scale_used) {
+ up *= y_scale[j];
+ gate *= y_scale[j];
+ }
+ up *= scale_up;
+ gate *= scale_gate;
+ }
+ dst[ids_dst[j]*stride + i] = ggml_cuda_op_glu_single(glu_op, gate, up, glu_limit);
} else if constexpr (type == GGML_TYPE_NVFP4) {
if (y_scale_used) {
dst[ids_dst[j]*stride + i] = y_scale[j] * sum[(j0/tile_C::J + n)*tile_C::ne + l];
@@ -558,7 +575,7 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma(
}
}
- GGML_UNUSED_VARS(y_scale_used, glu_op, glu_limit);
+ GGML_UNUSED_VARS(y_scale_used, glu_op, glu_limit, x_scale, gate_scale, scale_up, scale_gate);
}
// -------------------------------------------------------------------------------------------------------------------------------------
@@ -712,7 +729,7 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func
case GGML_TYPE_NVFP4:
return ggml_cuda_mmq_util_funcs(
VDR_NVFP4_Q8_1_MMQ,
- ggml_cuda_mmq_load_tiles_nvfp4<type, J, fallback, prec_src1>,
+ ggml_cuda_mmq_load_tiles_nvfp4<type, J, fallback, prec_src1, false>,
ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_dp4a<type, J, fallback>,
ggml_cuda_mmq_write_back_dp4a<type, J, fallback, prec_src1>);
default:
@@ -737,9 +754,9 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func
if (prec_src1 == GGML_PREC_Q4) {
return ggml_cuda_mmq_util_funcs(
-1,
- ggml_cuda_mmq_load_tiles_nvfp4_nvfp4<type, J, fallback>,
+ ggml_cuda_mmq_load_tiles_nvfp4_nvfp4<type, J, fallback, has_fusion>,
ggml_cuda_mmq_vec_dot_fp4_fp4_mma<type, J, fallback>,
- ggml_cuda_mmq_write_back_mma<type, J, fallback, prec_src1, false>);
+ ggml_cuda_mmq_write_back_mma<type, J, fallback, prec_src1, has_fusion>);
}
break;
default:
@@ -882,9 +899,9 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func
case GGML_TYPE_NVFP4:
return ggml_cuda_mmq_util_funcs(
-1,
- ggml_cuda_mmq_load_tiles_nvfp4<type, J, fallback, prec_src1>,
+ ggml_cuda_mmq_load_tiles_nvfp4<type, J, fallback, prec_src1, has_fusion>,
ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma<type, J, fallback, prec_src1>,
- ggml_cuda_mmq_write_back_mma<type, J, fallback, prec_src1, false>);
+ ggml_cuda_mmq_write_back_mma<type, J, fallback, prec_src1, has_fusion>);
default:
return ggml_cuda_mmq_util_funcs(1, nullptr, nullptr, nullptr);
}
@@ -918,6 +935,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile(
const char * __restrict__ x, const char * __restrict__ x_gate, const int offset_x, const int * __restrict__ y,
const int * __restrict__ ids_dst, float * __restrict__ dst, float * __restrict__ tmp_fixup,
const float * __restrict__ y_scale, const ggml_glu_op glu_op, const float glu_limit,
+ const float * __restrict__ x_scale, const float * __restrict__ gate_scale,
const int stride_row_x, const int ncols_y, const int stride_col_dst,
const int tile_x_max_i, const int tile_y_max_j, const int kb0_start, const int kb0_stop) {
static_assert(!(fixup && has_fusion), "the fused gate/up path does not split K");
@@ -987,9 +1005,9 @@ static __device__ __forceinline__ void mul_mat_q_process_tile(
}
if (fixup) {
- write_back(sum, ids_dst, tmp_fixup + blockIdx.x*(J*I), y_scale, I, I, J, glu_op, glu_limit);
+ write_back(sum, ids_dst, tmp_fixup + blockIdx.x*(J*I), y_scale, I, I, J, glu_op, glu_limit, x_scale, gate_scale);
} else {
- write_back(sum, ids_dst, dst, y_scale, stride_col_dst, tile_x_max_i, tile_y_max_j, glu_op, glu_limit);
+ write_back(sum, ids_dst, dst, y_scale, stride_col_dst, tile_x_max_i, tile_y_max_j, glu_op, glu_limit, x_scale, gate_scale);
}
}
@@ -1002,6 +1020,7 @@ static __global__ void mul_mat_q(
const char * __restrict__ x, const char * __restrict__ x_gate, const int * __restrict__ y, const int32_t * __restrict__ ids_dst,
const int32_t * __restrict__ expert_bounds, float * __restrict__ dst, float * __restrict__ tmp_fixup,
const float * __restrict__ y_scale, const ggml_glu_op glu_op, const float glu_limit,
+ const float * __restrict__ x_scale, const float * __restrict__ gate_scale,
const uint3 blocks_per_ne00, const int nrows_x, const int ncols_dst, const int stride_row_x, const int ncols_y, const int stride_col_dst,
const uint3 channel_ratio, const uint3 nchannels_y, const int stride_channel_x, const int stride_channel_y, const int stride_channel_dst,
const uint3 sample_ratio, const uint3 nsamples_y, const int stride_sample_x, const int stride_sample_y, const int stride_sample_dst,
@@ -1012,9 +1031,15 @@ static __global__ void mul_mat_q(
#else
constexpr bool fusion_available = false;
#endif // TURING_MMA_AVAILABLE
+#ifdef BLACKWELL_MMA_AVAILABLE
+ constexpr bool native_fp4_available = true;
+#else
+ constexpr bool native_fp4_available = false;
+#endif // BLACKWELL_MMA_AVAILABLE
- // Skip unused template specializations for faster compilation:
- if (ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).type == GGML_TYPE_COUNT || (has_fusion && !fusion_available)) {
+ // Skip unused template specializations for faster compilation (the host only uses 4-bit src1 with native FP4):
+ if (ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).type == GGML_TYPE_COUNT ||
+ (has_fusion && (!fusion_available || (prec_src1 == GGML_PREC_Q4 && !native_fp4_available)))) {
NO_DEVICE_CODE;
return;
}
@@ -1108,10 +1133,14 @@ static __global__ void mul_mat_q(
const int offset_x = fastdiv(wt, sample_ratio)*stride_sample_x + fastdiv(zt, channel_ratio)*stride_channel_x + it*I_dst*stride_row_x;
+ // With ids the NVFP4 weight scales are per expert.
+ const float * x_scale_tile = x_scale && ids_dst ? x_scale + zt : x_scale;
+ const float * gate_scale_tile = gate_scale && ids_dst ? gate_scale + zt : gate_scale;
+
constexpr bool fixup = false;
mul_mat_q_process_tile<type, J, fallback, fixup, prec_src1, has_fusion>
(x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile,
- glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst,
+ glu_op, glu_limit, x_scale_tile, gate_scale_tile, stride_row_x, ncols_y, stride_col_dst,
tile_x_max_i, tile_y_max_j, 0, blocks_per_ne00.z);
return;
}
@@ -1205,7 +1234,7 @@ static __global__ void mul_mat_q(
constexpr bool fixup = false; // All but (potentially) the last iterations write their data to dst rather than the fixup buffer.
mul_mat_q_process_tile<type, J, fallback, fixup, prec_src1, false>
(x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile,
- glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst,
+ glu_op, glu_limit, x_scale, gate_scale, stride_row_x, ncols_y, stride_col_dst,
tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop);
kbc += blocks_per_ne00.z;
@@ -1289,7 +1318,7 @@ static __global__ void mul_mat_q(
constexpr bool fixup = true; // Last index writes its data to fixup buffer to avoid data races with other blocks.
mul_mat_q_process_tile<type, J, fallback, fixup, prec_src1, false>
(x, x_gate, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, y_scale_tile,
- glu_op, glu_limit, stride_row_x, ncols_y, stride_col_dst,
+ glu_op, glu_limit, x_scale, gate_scale, stride_row_x, ncols_y, stride_col_dst,
tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop);
}
@@ -1491,6 +1520,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a
GGML_ASSERT(nty <= 65535); // gridDim.y limit
mul_mat_q<type, J, fallback, prec_src1, has_fusion><<<block_nums_xy_tiling, block_dims, nbytes_shared, stream>>>
(args.x, (const char *) args.fusion.gate, args.y, args.ids_dst, args.expert_bounds, args.dst, nullptr, args.y_scale, args.fusion.glu_op, args.fusion.glu_limit,
+ (const float *) args.fusion.x_scale, (const float *) args.fusion.gate_scale,
blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst,
channel_ratio_fd, nchannels_y_fd, args.stride_channel_x, args.stride_channel_y, args.stride_channel_dst,
sample_ratio_fd, nsamples_y_fd, args.stride_sample_x, args.stride_sample_y, args.stride_sample_dst,
@@ -1522,6 +1552,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a
mul_mat_q<type, J, fallback, prec_src1, has_fusion><<<block_nums_stream_k, block_dims, nbytes_shared, stream>>>
(args.x, (const char *) args.fusion.gate, args.y, args.ids_dst, args.expert_bounds, args.dst, tmp_fixup.ptr, args.y_scale, args.fusion.glu_op, args.fusion.glu_limit,
+ (const float *) args.fusion.x_scale, (const float *) args.fusion.gate_scale,
blocks_per_ne00_fd, args.nrows_x, args.ncols_dst, args.stride_row_x, args.ncols_y, args.nrows_dst,
channel_ratio_fd, nchannels_y_fd, args.stride_channel_x, args.stride_channel_y, args.stride_channel_dst,
sample_ratio_fd, nsamples_y_fd, args.stride_sample_x, args.stride_sample_y, args.stride_sample_dst,
@@ -1661,4 +1692,5 @@ void ggml_cuda_mul_mat_q(
bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t n_experts);
-bool ggml_cuda_mmq_fusion_is_efficient(ggml_type type, int64_t nrows_x, int64_t ncols, int64_t nchannels, int device);
+// prec_requested: the src1 precision requested by the fused MUL_MATs.
+bool ggml_cuda_mmq_fusion_is_efficient(ggml_type type, ggml_prec prec_requested, int64_t nrows_x, int64_t ncols, int64_t nchannels, int device);
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index f9cb12c3b..87067b64d 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -7993,6 +7993,36 @@ 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 {
+ // 4-bit activations on Blackwell, see test_mul_mat
+ if (type == GGML_TYPE_NVFP4 && !graph_mul_mat_hi_prec_act(gf, use_id ? GGML_OP_MUL_MAT_ID : GGML_OP_MUL_MAT) &&
+ backend_has_feature(backend, "BLACKWELL_NATIVE_FP4")) {
+ return 2e-2;
+ }
+ return max_nmse_err();
+ }
+};
+
+// Fused gate/up/GLU with GGML_PREC_Q8 on src1, e.g. a W4A16 NVFP4 checkpoint (W4A8, see test_mul_mat_w4a8)
+struct test_mul_mat_vec_fusion_w4a8 : public test_mul_mat_vec_fusion {
+ using test_mul_mat_vec_fusion::test_mul_mat_vec_fusion;
+ using test_mul_mat_vec_fusion::build_graph;
+
+ ggml_tensor * build_graph(ggml_context * ctx, ggml_context * ctx_weights) override {
+ ggml_tensor * out = test_mul_mat_vec_fusion::build_graph(ctx, ctx_weights);
+ for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
+ if (t->op == GGML_OP_MUL_MAT || t->op == GGML_OP_MUL_MAT_ID) {
+ ggml_prec_set_src(t, GGML_PREC_Q8, 1);
+ }
+ }
+ return out;
+ }
+
+ std::string op_desc(ggml_tensor * t) override {
+ GGML_UNUSED(t);
+ return "MUL_MAT_VEC_FUSION_W4A8";
+ }
};
// GGML_OP_SUM
@@ -12630,9 +12660,17 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q6_K, GGML_GLU_OP_SWIGLU_OAI, 150, 40960, 256, false, 1, 1, false, false, true, false, { 1, 1 }));
test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q8_0, GGML_GLU_OP_SWIGLU_CLAMP, 150, 6912, 256, false, 1, 1, false, false, true, false, { 2, 3 }));
test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_IQ4_XS, GGML_GLU_OP_SWIGLU, 150, 40960, 512, false, 1, 1, false, false, true, false, { 1, 1 }));
+ test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 6912, 256, false, 1, 1, false, false, true, true, { 2, 3 }));
// The same with experts (~1280 tiles each): 17 tokens with broadcast input (one MMA tile per warp), 2048 tokens without (two, several column tiles per expert).
test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 17, 5120, 256, true, 16, 8, true));
test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q8_0, GGML_GLU_OP_SWIGLU, 2048, 128, 256, true, 40, 8, false));
+ // NVFP4 with experts, without and with per-expert weight scales.
+ test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 8, false));
+ test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 17, 5120, 256, true, 16, 8, true, false, true, true));
+ // NVFP4 with 8-bit activations requested (W4A16 checkpoints), without and with weight scales.
+ test_cases.emplace_back(new test_mul_mat_vec_fusion_w4a8(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 8, false));
+ test_cases.emplace_back(new test_mul_mat_vec_fusion_w4a8(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 40960, 256, false, 1, 1, false, false, true, true, { 1, 1 }));
+ test_cases.emplace_back(new test_mul_mat_vec_fusion_w4a8(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 8, false, false, true, true));
for (auto gate : {GATING_FUNC_SOFTMAX, GATING_FUNC_SIGMOID, GATING_FUNC_SOFTMAX_WEIGHT, GATING_FUNC_SQRT_SOFTPLUS}) {
for (bool with_norm : {false, true}) {