Commit 0d69255c5 for llama.cpp
commit 0d69255c5205c3af9eb15d8a7fc027594848041f
Author: Konrad Moren <kmoren@nvidia.com>
Date: Thu Nov 12 10:28:36 2026 +0100
CUDA: Optimize prefil via fuse of w_s scale in epilogue MMQ for nvfp4 checkpoints (#26048)
* ggml-cuda: Apply NVFP4 weight scale in MMQ write-back
* ggml-cuda: Fuse mul_mat(_id) + scale into MMQ
* ggml-cuda: Tests for the fused scale
* ggml-cuda: reuse existing graph to test fusion in mmq epiloque
* ggml-cuda: rename test_mul_mat_vec_fusion to test_mul_mat_fusion for consistency
* ggml-cuda: refine comments for clarity on MMQ dispatch and fusion checks in mul_mat functions
* ggml-cuda: refactor fusion checks for mul_mat operations and remove unused function
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 92ae9848a..28073a17f 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -5009,6 +5009,23 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
return 1;
}
+ if (!with_bias && src0->type == GGML_TYPE_NVFP4 && src1->type == GGML_TYPE_F32 && mm_node->type == GGML_TYPE_F32) {
+ const int cc = ggml_cuda_info().devices[cuda_ctx->device].cc;
+ const bool is_id = op == GGML_OP_MUL_MAT_ID;
+ const int64_t ncols = is_id ? src1->ne[2] : src1->ne[1];
+ const bool use_mmvq = is_id ? ncols <= get_mmvq_mmid_max_batch(src0->type, cc) :
+ ggml_cuda_should_use_mmvq(src0->type, cc, ncols);
+ const bool bad_padding_clear = ggml_backend_buffer_get_usage(src0->buffer) == GGML_BACKEND_BUFFER_USAGE_COMPUTE &&
+ ggml_nbytes(src0) != ggml_backend_buffer_get_alloc_size(src0->buffer, src0) &&
+ src0->view_src;
+ if (!use_mmvq && !bad_padding_clear &&
+ ggml_cuda_should_use_mmq(src0->type, cc, ncols, /*n_experts =*/ is_id ? src0->ne[2] : 0)) {
+ fusion_data.prec_src1 = ggml_prec_get_src(mm_node, 1);
+ ggml_cuda_mul_mat_q(*cuda_ctx, src0, src1, ids, out_node, &fusion_data);
+ return n_ops - 1;
+ }
+ }
+
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;
diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu
index 367657508..2e6199766 100644
--- a/ggml/src/ggml-cuda/mmq.cu
+++ b/ggml/src/ggml-cuda/mmq.cu
@@ -169,21 +169,27 @@ void ggml_cuda_mul_mat_q(
GGML_ASSERT(!ids || ids->type == GGML_TYPE_I32); // Optional, used for batched GGML_MUL_MAT_ID.
// Fused gate/up/GLU: src0 is the up weight, dst the GLU output.
+ // Fused scale (no gate): dst is the MUL that applies the weight scale to the MUL_MAT.
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->type == src0->type && ggml_are_same_shape(gate, src0) && ggml_are_same_stride(gate, src0));
+ GGML_ASSERT(!fusion->x_bias && !fusion->gate_bias);
// 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;
+ if (gate) {
+ GGML_ASSERT(!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.gate_scale = fusion->gate_scale ? fusion->gate_scale->data : nullptr;
+ fusion_device.glu_op = fusion->glu_op;
+ fusion_device.glu_limit = fusion->glu_limit;
+ } else {
+ GGML_ASSERT(!fusion->gate_scale && fusion->x_scale);
+ }
+ fusion_device.x_scale = fusion->x_scale ? fusion->x_scale->data : nullptr;
}
GGML_TENSOR_BINARY_OP_LOCALS;
@@ -231,8 +237,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 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);
+ // With fusion dst is the GLU output or the scale node, not a MUL_MAT; the src1 precision requested by the fused MUL_MATs is in the fusion args.
+ const ggml_prec prec_src1 = ggml_cuda_mmq_get_prec_src1(src0->type, fusion ? 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);
diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
index 42a3afd24..d645a8aac 100644
--- a/ggml/src/ggml-cuda/mmq.cuh
+++ b/ggml/src/ggml-cuda/mmq.cuh
@@ -465,7 +465,11 @@ template <ggml_type type, int J, bool fallback, ggml_prec prec_src1> static __de
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_sup,
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);
+ GGML_UNUSED_VARS(glu_op, glu_limit, gate_scale);
+
+ // NVFP4 weight scale of a fused MUL (per tensor, or of the expert of this tile), 1 without.
+ const float scale = type == GGML_TYPE_NVFP4 && x_scale ? *x_scale : 1.0f;
+
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);
@@ -489,11 +493,15 @@ template <ggml_type type, int J, bool fallback, ggml_prec prec_src1> static __de
continue;
}
- if (y_scale_used) {
- dst[ids_dst[j]*stride + i] = y_scale[j] * sum[(j0/nwarps) * (I/warp_size) + i0/warp_size];
+ float value = sum[(j0/nwarps) * (I/warp_size) + i0/warp_size];
+ if constexpr (type == GGML_TYPE_NVFP4) {
+ if (y_scale_used) {
+ value *= y_scale[j];
+ }
} else {
- dst[ids_dst[j]*stride + i] = sum[(j0/nwarps) * (I/warp_size) + i0/warp_size];
+ GGML_UNUSED(y_scale_used);
}
+ dst[ids_dst[j]*stride + i] = value * scale;
}
}
}
@@ -520,9 +528,9 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma(
const bool y_scale_used = type == GGML_TYPE_NVFP4 && y_scale != nullptr;
- // Fused NVFP4: load the weight scales once, 1 without scales.
+ // NVFP4: load the weight scales once, 1 without scales. Without fusion x_scale is the scale of a fused MUL.
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_up = type == GGML_TYPE_NVFP4 && x_scale ? *x_scale : 1.0f;
const float scale_gate = fused_nvfp4 && gate_scale ? *gate_scale : 1.0f;
#pragma unroll
@@ -556,10 +564,16 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma(
gate *= scale_gate;
}
dst[ids_dst[j]*stride + i] = ggml_cuda_op_glu_single(glu_op, gate, up, glu_limit);
- } else if (y_scale_used) {
- dst[ids_dst[j]*stride + i] = y_scale[j] * sum[(j0/tile_C::J + n)*tile_C::ne + l];
} else {
- dst[ids_dst[j]*stride + i] = sum[(j0/tile_C::J + n)*tile_C::ne + l];
+ float value = sum[(j0/tile_C::J + n)*tile_C::ne + l];
+ if constexpr (type == GGML_TYPE_NVFP4) {
+ if (y_scale_used) {
+ value *= y_scale[j];
+ }
+ } else {
+ GGML_UNUSED(y_scale_used);
+ }
+ dst[ids_dst[j]*stride + i] = value * scale_up;
}
}
}
@@ -1201,10 +1215,12 @@ 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*stride_row_x;
+ const float * x_scale_tile = x_scale && ids_dst ? x_scale + zt : x_scale;
+
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, x_scale, gate_scale, stride_row_x, ncols_y, stride_col_dst,
+ glu_op, glu_limit, x_scale_tile, gate_scale, stride_row_x, ncols_y, stride_col_dst,
tile_x_i_max, col_high - col_low, kb0_start, kb0_stop);
kbc += blocks_per_ne00.z;
@@ -1275,10 +1291,12 @@ 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*stride_row_x;
+ const float * x_scale_tile = x_scale && ids_dst ? x_scale + zt : x_scale;
+
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, x_scale, gate_scale, stride_row_x, ncols_y, stride_col_dst,
+ glu_op, glu_limit, x_scale_tile, gate_scale, stride_row_x, ncols_y, stride_col_dst,
tile_x_i_max, col_high - col_low, kb0_start, kb0_stop);
}
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index a63d5e24d..b783cccb5 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -8118,7 +8118,7 @@ struct test_mul_mat_add : public test_case {
}
};
-struct test_mul_mat_vec_fusion : public test_case {
+struct test_mul_mat_fusion : public test_case {
const ggml_type type;
const ggml_glu_op glu_op;
const int64_t m;
@@ -8134,7 +8134,7 @@ struct test_mul_mat_vec_fusion : public test_case {
std::array<int64_t, 2> batch_dims;
const bool skip; // set some ids to a negative value (only for use_id)
- test_mul_mat_vec_fusion(ggml_type type, ggml_glu_op op, int64_t m, int64_t n, int64_t k,
+ test_mul_mat_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<int64_t, 2> batch_dims = {4, 2}, bool skip = 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),
@@ -8150,7 +8150,7 @@ struct test_mul_mat_vec_fusion : public test_case {
std::string op_desc(ggml_tensor * t) override {
GGML_UNUSED(t);
- return "MUL_MAT_VEC_FUSION";
+ return "MUL_MAT_FUSION";
}
bool run_whole_graph() override { return true; }
@@ -8315,12 +8315,12 @@ struct test_mul_mat_vec_fusion : public test_case {
};
// 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;
+struct test_mul_mat_fusion_w4a8 : public test_mul_mat_fusion {
+ using test_mul_mat_fusion::test_mul_mat_fusion;
+ using test_mul_mat_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);
+ ggml_tensor * out = test_mul_mat_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);
@@ -8331,7 +8331,7 @@ struct test_mul_mat_vec_fusion_w4a8 : public test_mul_mat_vec_fusion {
std::string op_desc(ggml_tensor * t) override {
GGML_UNUSED(t);
- return "MUL_MAT_VEC_FUSION_W4A8";
+ return "MUL_MAT_FUSION_W4A8";
}
};
@@ -11800,6 +11800,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
}
}
}
+
#if 0
// > 2^31-element dst (8.9 GB), about 45 GB peak memory. Too large to be enabled by default.
// kernel_mul_mm dst offset of the last batch is past INT32_MAX: rows 0-31 use full tiles, rows 32-39 use edge tiles.
@@ -13137,12 +13138,12 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_opt_step_sgd(GGML_TYPE_F32, {10, 5, 4, 3}));
for (bool with_bias : {false, true}) {
- test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_F8_E4M3, GGML_GLU_OP_SWIGLU, 16, 32, 256,
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_F8_E4M3, GGML_GLU_OP_SWIGLU, 16, 32, 256,
false, 1, 1, false, with_bias, false, true, {1, 1}));
}
- test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_F8_E4M3, GGML_GLU_OP_SWIGLU, 16, 32, 256,
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_F8_E4M3, GGML_GLU_OP_SWIGLU, 16, 32, 256,
false, 1, 1, false, false, false, true, {2, 1}));
- test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_F8_E4M3, GGML_GLU_OP_SWIGLU, 16, 31, 257,
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_F8_E4M3, GGML_GLU_OP_SWIGLU, 16, 31, 257,
false, 1, 1, false, false, false, true, {1, 1}));
for (ggml_type type : base_types) {
@@ -13167,18 +13168,18 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
if (with_lane_scale && type != GGML_TYPE_NVFP4 && type != GGML_TYPE_F8_E4M3) {
continue;
}
- test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256,
+ test_cases.emplace_back(new test_mul_mat_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,
+ test_cases.emplace_back(new test_mul_mat_fusion(type, glu_op, 1, 32, 256,
use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, {1, 1}));
// multi-token batches (spec decoding)
for (int64_t m_batch : { 2, 4, 8 }) {
- test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, m_batch, 32, 256,
+ test_cases.emplace_back(new test_mul_mat_fusion(type, glu_op, m_batch, 32, 256,
use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, {1, 1}));
}
// the lane scale does a GET_ROWS on the ids, which does not accept -1
if (use_id && !with_lane_scale) {
- test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 8, 32, 256,
+ test_cases.emplace_back(new test_mul_mat_fusion(type, glu_op, 8, 32, 256,
use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, {1, 1}, true));
}
}
@@ -13190,7 +13191,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
}
for (bool b : {false, true}) {
- test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_IQ2_S, GGML_GLU_OP_SWIGLU_CLAMP, 1, 32, 256,
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_IQ2_S, GGML_GLU_OP_SWIGLU_CLAMP, 1, 32, 256,
true, 16, 8, b, false, true, false));
}
@@ -13200,7 +13201,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
GGML_TYPE_Q8_0}) {
for (ggml_glu_op glu_op : {GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU}) {
for (int64_t m_batch : {1, 2, 3, 5, 8, 9, 16, 17}) {
- test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, m_batch, 1000, 1024,
+ test_cases.emplace_back(new test_mul_mat_fusion(type, glu_op, m_batch, 1000, 1024,
false, 16, 8, false, false, true, false, { 1, 1 }));
}
}
@@ -13208,7 +13209,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
for (int64_t k : {96, 896, 1024, 2880}) {
for (ggml_glu_op glu_op : {GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU}) {
- test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q8_0, glu_op, 1, 32, k,
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_Q8_0, glu_op, 1, 32, k,
false, 16, 8, false, false, true, false, {1, 1}));
}
}
@@ -13218,7 +13219,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
//for (ggml_glu_op glu_op : { GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU }) {
// for (int64_t m_batch : { 2, 3, 4 }) {
// for (int64_t rows : { 1, 2, 3 }) {
- // test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, glu_op, m_batch, rows, 256,
+ // test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_Q4_K, glu_op, m_batch, rows, 256,
// false, 16, 8, false, false, true, false, { 1, 1 }));
// }
// }
@@ -13226,27 +13227,40 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
// Both sides of the same row-count boundary as above, on the fused path.
for (int64_t rows : {6271, 6272, 6273}) {
- test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 2, rows, 256,
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 2, rows, 256,
false, 16, 8, false, false, true, false, { 1, 1 }));
}
// Fused gate/up + GLU on the MMQ path, only used if its tiling is efficient: ~1280 tiles per case, one case per type and GLU op.
- test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 17, 81920, 256, false, 1, 1, false, false, true, false, { 1, 1 }));
- test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q5_K, GGML_GLU_OP_GEGLU, 150, 40960, 256, false, 1, 1, false, false, true, false, { 1, 1 }));
- 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 }));
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 17, 81920, 256, false, 1, 1, false, false, true, false, { 1, 1 }));
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_Q5_K, GGML_GLU_OP_GEGLU, 150, 40960, 256, false, 1, 1, false, false, true, false, { 1, 1 }));
+ test_cases.emplace_back(new test_mul_mat_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_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_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_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 6912, 256, false, 1, 1, false, false, true, true, { 2, 3 }));
+ // NVFP4 weight scale without gate/up: MMQ applies it in the write-back, the trailing bias add stays a separate node
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, false, 1, 1, false, false, false, true, { 1, 1 }));
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 8, false, false, false, true));
+ // Same pattern with broadcast input, 4 of 16 experts used, 8-bit activations (W4A8), and Q4_K, which keeps the scale as a separate node
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 8, true, false, false, true));
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 4, false, false, false, true));
+ test_cases.emplace_back(new test_mul_mat_fusion_w4a8(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, false, 1, 1, false, false, false, true, { 1, 1 }));
+ test_cases.emplace_back(new test_mul_mat_fusion_w4a8(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 8, false, false, false, true));
+ // Few columns and tiles with a long K: split-K through the stream-k fixup, the scale is applied to the partial sums
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 11, 64, 2048, false, 1, 1, false, false, false, true, { 1, 1 }));
+ test_cases.emplace_back(new test_mul_mat_fusion_w4a8(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 11, 64, 2048, false, 1, 1, false, false, false, true, { 1, 1 }));
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 150, 2560, 256, false, 1, 1, false, false, false, true, { 1, 1 }));
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 8, false, false, false, true));
// 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));
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_Q4_K, GGML_GLU_OP_SWIGLU, 17, 5120, 256, true, 16, 8, true));
+ test_cases.emplace_back(new test_mul_mat_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));
+ test_cases.emplace_back(new test_mul_mat_fusion(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 8, false));
+ test_cases.emplace_back(new test_mul_mat_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));
+ test_cases.emplace_back(new test_mul_mat_fusion_w4a8(GGML_TYPE_NVFP4, GGML_GLU_OP_SWIGLU, 150, 2560, 256, true, 16, 8, false));
+ test_cases.emplace_back(new test_mul_mat_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_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}) {