CUDA: fuse FFN gate/up matmuls and GLU in MMQ with interleaved gate/up rows by anujj · Pull Request #29948 · ggml-org/llama.cpp

anujj · 2026-10-04T13:27:39Z

Extends the gate/up + GLU fusion, which today only covers small batches (MMVQ/MMVF), to prompt-processing batch
sizes on the MMQ path.

Alternative to #28702 implementing @am17an's suggestion there: keep one accumulator and apply the GLU in the write-back instead of carrying a second accumulator tile for the gate as discussed here: #28702 (comment)

For prompt processing, the gate and up MUL_MATs of a dense FFN and the following GLU run as a single mul_mat_q launch. The small-batch MMVQ/MMVF fusion, other backends and the weight layout are unchanged.

Implimentation

The loaders interleave up and gate rows in chunks of 8, so each thread holds the up and gate value of the same output element (rows r and r + 8 of the MMA tile). Chunks of 8 rather than 16 also cover the one-tile-per-warp configs.
MMA loop, tile config and sum[] are unchanged; the write-back applies the GLU from registers and writes 64 output rows per tile.
Q4_K, Q5_K, Q6_K, Q8_0, IQ4_XS on NVIDIA Turing+; SWIGLU, GEGLU, SWIGLU_OAI, SWIGLU_CLAMP.
Please note: I have a plan to validate and enable it for the other GGML also.

Performance

DGX Spark (GB10), llama-bench -ngl 999 -fa on -p 16384 -n 0 -b 2048 -ub 2048 -ctk f16 -ctv f16 -r 5, tok/s:

model	type	master	PR
Qwen3.8-27B	Q4_K_M	840.1	899.9	+7.1%
	Q5_K_M	840.1	880.6	+4.8%
	Q6_K	751.2	785.8	+4.6%
	Q8_0	876.2	940.1	+7.3%
	IQ4_XS	882.4	948.0	+7.4%
Qwen3.6-27B	Q4_K_M	816.6	865.6	+6.0%
	Q5_K_M	800.0	851.1	+6.4%
	Q6_K	720.8	770.4	+6.9%
	Q8_0	876.9	940.2	+7.2%
	IQ4_XS	882.2	948.5	+7.5%
Llama-3.1-8B	Q4_K_M	2749.3	3000.6	+9.1%
	Q5_K_M	2699.4	2960.3	+9.7%
	Q6_K	2392.2	2616.0	+9.4%
	IQ4_XS	2933.9	3249.0	+10.7%

RTX 5090, llama-bench -ngl 999 -fa on -p 16384 -n 0 -b 2048 -ub 2048 -ctk f16 -ctv f16 -r 5, tok/s:

model	type	master	PR
Qwen3.8-27B	Q4_K_M	3688.4	3790.9	+2.8%
	Q5_K_M	3589.5	3659.6	+2.0%
	Q6_K	3220.5	3281.0	+1.9%
	Q8_0	3617.3	3766.9	+4.1%
	IQ4_XS	3796.2	3926.8	+3.4%
Qwen3.6-27B	Q4_K_M	3567.9	3660.3	+2.6%
	Q5_K_M	3429.8	3525.3	+2.8%
	Q6_K	3105.2	3187.8	+2.7%
	Q8_0	3628.7	3788.4	+4.4%
	IQ4_XS	3785.6	3915.4	+3.4%
Llama-3.1-8B	Q4_K_M	12071.5	12625.9	+4.6%
	Q5_K_M	11823.0	12297.3	+4.0%
	Q6_K	10513.1	10928.6	+4.0%
	IQ4_XS	12997.6	13603.3	+4.7%

Correctness

test-backend-ops: 19 new MUL_MAT_VEC_FUSION cases hit the MMQ path (129 rows with bounds checks at one and two MMA tiles per warp, 256 rows without, GEGLU / SWIGLU_OAI / SWIGLU_CLAMP, batched). Full suite passes on GB10; the MUL_MAT suites also pass on an RTX 5090.
Perplexity (wiki.test, 8 chunks) is bit-identical with GGML_CUDA_DISABLE_FUSION=1 for all five types, also with -ot ffn_gate=CPU.

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: YES - Paired with claude

anujj · 2026-10-04T13:32:49Z

Adding @am17an and @JohannesGaessler for review

anujj · 2026-10-08T10:25:36Z

Adding @am17an and @JohannesGaessler for review

gentle reminder for the review

JohannesGaessler · 2026-10-08T11:36:52Z

Sorry, I hadn't prioritized looking at this PR for review because it is still marked as a draft.

JohannesGaessler · 2026-10-08T11:41:43Z

+    if (up->op != GGML_OP_MUL_MAT || !ggml_cuda_mmq_fusion_supported(x_up->type) || y->type != GGML_TYPE_F32 ||
+            ggml_get_op_params_i32(up, 1) != GGML_HINT_NONE || ggml_get_op_params_i32(gate, 1) != GGML_HINT_NONE) {
+        return false;
+    }
+
+    if (!ggml_cuda_should_fuse_mul_mat(up, gate, glu)) {
+        return false;
+    }


I'm not sure this is logically correct. You can fuse matrices with the Hadamard hind, but whether you should is a different question. So the correct place to have this check would be in ggml_cuda_should_fuse_mul_mat_q.

Moved the hint check into ggml_cuda_should_fuse_mul_mat_q. If I understand currently it does not trigger here: the hint is only set on the KV-cache rotation matmuls (F32 src0, not followed by a GLU), while the fused path only takes quantized gate/up weights.

JohannesGaessler · 2026-10-08T11:55:00Z

+    // Fused gate/up/GLU: src0 is the up weight, dst the GLU output.
+    const ggml_tensor * gate = fusion ? fusion->gate : nullptr;
+    if (fusion) {
+        GGML_ASSERT(gate && !ids && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale);


Suggested change

GGML_ASSERT(gate && !ids && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale);

// TODO MoE support

GGML_ASSERT(gate && !ids && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale);

JohannesGaessler · 2026-10-08T11:57:15Z

+    // With fusion dst is the GLU output, not a MUL_MAT; the fused types always quantize src1 to Q8_1.
+    const ggml_prec prec_src1 = gate ? GGML_PREC_Q8 : ggml_cuda_mmq_get_prec_src1(src0, dst, cc);


Suggested change

// With fusion dst is the GLU output, not a MUL_MAT; the fused types always quantize src1 to Q8_1.

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 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);

JohannesGaessler · 2026-10-08T12:07:45Z

-    if (!config.stream_k) {
-        mul_mat_q<type, J, fallback, prec_src1><<<block_nums_xy_tiling, block_dims, nbytes_shared, stream>>>
-            (args.x, args.y, args.ids_dst, args.expert_bounds, args.dst, nullptr, args.y_scale,
+    if (has_fusion || !config.stream_k) {


Since fusion is incompatible with stream-k, this PR should make it so that the expected efficiency of tiling vs. stream-k can be determined ahead of time and used in ggml_cuda_should_fuse_mul_mat_q.

JohannesGaessler · 2026-10-08T12:13:29Z

+        const int jt = has_fusion ? blockIdx.x : blockIdx.y;
+        const int it = has_fusion ? blockIdx.y : blockIdx.x;


Why are you swapping blockIdx.x and blockIdx.y?

The fused kernel cannot use stream-k, so it runs the xy tiling with the column tiles innermost, the same order as the stream-k path (jt = tmp % ntx). With the row tiles innermost the concurrently running blocks each read a different weight tile, so the gate+up weights of a layer (100 MB for Qwen3.8-27B Q4_K vs 24 MB of L2 on spark) are streamed from VRAM once per column tile; with the column tiles innermost they share one weight tile and only y (~12 MB at 2048 tokens) has to stay in L2. Fused gate/up time per layer on spark, Qwen3.8-27B pp2048, row tiles innermost vs this PR: Q4_K 14.5 vs 10.1 ms, Q8_0 20.7 vs 9.1 ms, IQ4_XS 10.7 vs 8.7 ms.

JohannesGaessler · 2026-10-08T12:17:22Z

+// Fused gate/up/GLU write-back: accumulators l and l + 2 hold the up and gate values of the same output row and column.
+template<ggml_type type, int J, bool fallback>
+static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma_glu(
+            const float * __restrict__ sum, const ggml_glu_op glu_op, const float glu_limit,
+            const int * __restrict__ ids_dst, float * __restrict__ dst, const int stride, const int i_max, const int j_max) {
+#ifdef TURING_MMA_AVAILABLE
+    typedef tile<16, 8, int> tile_C;
+    static_assert(tile_C::I == 2*GGML_CUDA_MMQ_FUSION_ROWS, "each C tile must hold one chunk of up rows and one of gate rows");
+
+    constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
+    constexpr int ntx           = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
+
+    const int i0 = (threadIdx.y / ntx) * (ntx*tile_C::I);
+
+#pragma unroll
+    for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) {
+#pragma unroll
+        for (int n = 0; n < ntx; ++n) {
+#pragma unroll
+            for (int l = 0; l < tile_C::ne/2; ++l) {
+                const int j = j0 + (threadIdx.y % ntx) * tile_C::J + tile_C::get_j(l);
+
+                if (j > j_max) {
+                    continue;
+                }
+
+                const int i = (i0 + n*tile_C::I)/2 + tile_C::get_i(l);
+
+                if (fallback && i > i_max) {
+                    continue;
+                }
+
+                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 + 2], sum[sum_index], glu_limit);
+            }
+        }
+    }
+#else
+    GGML_UNUSED_VARS(sum, glu_op, glu_limit, ids_dst, dst, stride, i_max, j_max);
+    NO_DEVICE_CODE;
+#endif // TURING_MMA_AVAILABLE
+}
+


Avoid adding a whole new function and extend the template for the existing one instead.

JohannesGaessler · 2026-10-08T12:20:57Z

+    if (up->op != GGML_OP_MUL_MAT || !ggml_cuda_mmq_fusion_supported(x_up->type) || y->type != GGML_TYPE_F32 ||
+            ggml_get_op_params_i32(up, 1) != GGML_HINT_NONE || ggml_get_op_params_i32(gate, 1) != GGML_HINT_NONE) {
+        return false;
+    }


I think this is missing a check that up and gate have the same type.

That is covered by ggml_cuda_should_fuse_mul_mat, called right below: it returns false unless up->src[0] and gate->src[0] have the same type, shape and strides, and ggml_cuda_mul_mat_q asserts the same. Happy to make it explicit here if you prefer.

JohannesGaessler · 2026-10-08T12:35:03Z


+template <ggml_type type, bool fallback, ggml_prec prec_src1>
+static void mul_mat_q_switch_fusion(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) {
+    if constexpr (ggml_cuda_mmq_fusion_supported(type)) {


Suggested change

if constexpr (ggml_cuda_mmq_fusion_supported(type)) {

if constexpr (!fallback && ggml_cuda_mmq_fusion_supported(type)) {

In order to avoid compiling template specializations that make no sense to use. Add a check to disable fusion for cases with fallback == false and then you can also simplify the code a bit by using the assumption that nrows_x is a multiple of 128. Relevant PR: #29953 .

For prompt processing, run the gate and up matmuls of a dense FFN as one mul_mat_q launch and apply the GLU in the write-back. - The loaders take the gate weights as a second pointer and fill the weight tile with up and gate rows interleaved in chunks of 8, so each thread holds an up value and the matching gate value (rows r and r + 8 of an MMA tile). - The MMA loop is unchanged: same tile, warp layout, vec_dot and one sum[] per thread. The fused write-back computes the GLU from registers and writes half the rows. - The fused path uses the regular tiling (no stream-k) with the column tiles innermost so that neighbouring blocks share weight tiles. - Enabled for Q4_K, Q5_K, Q6_K, Q8_0 and IQ4_XS on NVIDIA Turing and newer.

…write-back

anujj · 2026-10-12T09:34:39Z

Addressed all the comments

SongXiaoXi · 2026-10-14T02:54:08Z

I ran this on Ada (RTX 4090, clocks locked to 2235 MHz, Qwen3.8-27B Q4_K, llama-bench -ngl 99 -r 5):

batch	master	PR
pp512	2984	3023 (+1.3%)
pp2048	2979	3015 (+1.2%)
pp4096	2951	2988 (+1.3%)

It helps on Ada too, though by less than on GB10 / RTX 5090. I also timed a single FFN block at the 27B shape (Q5_K, 5120 -> 17408 -> 5120, 512 tokens, median of 30 runs): 1441 us on master, 1414 us with this PR, and 1401 us with the down-projection GLU fusion in #30489. The fused launch is roughly at parity with the stock gate/up pair it replaces (~884 us vs ~880 us), so on Ada the gain comes mostly from removing the GLU kernel rather than from MMA-side savings; the down-side approach measures slightly faster per FFN.

One interaction worth knowing about: With #30489 also in the tree, the two split coverage per FFN: this PR anchors the GLU first, so on the FFNs it accepts the down-side fusion does not run. On Ada the down-side fusion alone measures slightly faster (3030 vs 3023 t/s at pp512); both together 3040.

Gating the fusion to cases where the plain dispatch would not pick stream-k might be worth checking on Ada.

anujj · 2026-10-14T07:33:42Z

Thanks for the Ada numbers and the per-FFN breakdown.

Scope. This PR is deliberately limited to the dense path and five types to keep it reviewable. The MUL_MAT_ID extension is already implemented locally on top of it (Qwen3.6-35B-A3B Q4_K_M, pp16K vs master: +10.2% on GB10, +6.1% on RTX 5090, perplexity bit-identical) and I plan to push it, and the remaining MMQ types, once the dense part is settled.

Fused matmul vs the two stock matmuls. It depends on the type and the GPU, though. Per FFN layer, pp2048, 27B shape (5120 -> 17408), median time of the two stock MMQ launches vs the fused one, and the whole gate/up block:

GPU	type	stock matmuls	fused matmul	matmul	block
GB10	Q4_K	10.32 ms	9.82 ms	-5%	12.70 -> 10.13 ms (-20%)
GB10	Q5_K	10.26 ms	10.12 ms	-1%	12.62 -> 10.35 ms (-18%)
GB10	Q8_0	9.69 ms	9.10 ms	-6%	11.90 -> 9.50 ms (-20%)
GB10	IQ4_XS	9.19 ms	8.37 ms	-9%	11.52 -> 8.68 ms (-25%)
RTX 5090	Q8_0	2.35 ms	2.26 ms	-4%	3.17 -> 2.29 ms (-28%)
RTX 5090	IQ4_XS	2.29 ms	2.26 ms	-1%	2.54 -> 2.23 ms (-12%)
RTX 5090	Q4_K	2.44 ms	2.49 ms	+2%	2.88 -> 2.48 ms (-14%)

Gating. Already in the PR since b6e0f03: ggml_cuda_should_fuse_mul_mat_q only fuses if master's tiling-vs-stream-k rule (>= 90% wave efficiency, ggml_cuda_mmq_tiling_is_efficient, shared with launch_mul_mat_q) would pick tiling for the fused grid. On 128 SMs all three of your shapes pass it (pp512: 1088 tiles, 94%), so your numbers are with fusion on. With the gate, pp4096 vs master: RTX 5090 Llama-8B ub 128/256/512 unfused (+0.6/+0.7/+0.1%), ub 1024/2048 +4.9/+6.3%; GB10 fuses at every ub, +5..+11%.

#30489. Unsloth dynamic quants are a concrete case: some layers have different gate and up types, which this PR cannot interleave, so #30489 should take those layers while this PR takes the rest. I am measuring master / this PR / #30489 / both on Qwen3.8-27B UD-Q4_K_M on spark and will post the numbers here.

anujj · 2026-10-14T09:46:25Z

Measured on spark with #30489 cherry-picked onto this PR. Qwen3.8-27B unsloth UD quants, pp16K, tok/s vs master on the same base:

model	master	#29948	both
UD-Q4_K_M	900.2	935.1 (+3.9%)	950.6 (+5.6%)
UD-Q4_K_XL	879.4	922.3 (+4.9%)	934.5 (+6.3%)

UD-Q4_K_M has 30 of 64 FFN layers with different gate and up types.

anujj · 2026-10-15T16:54:32Z

Addressed all the comments

@JohannesGaessler : pinging here to move this up in your review stack

github-actions Bot added testing Everything test related ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Oct 4, 2026

anujj marked this pull request as ready for review October 8, 2026 12:05

anujj requested review from a team and ggerganov as code owners October 8, 2026 12:05

JohannesGaessler reviewed Oct 8, 2026

View reviewed changes

anujj added 5 commits October 9, 2026 12:23

cuda: assert the fused MMQ gate argument and trim redundant comments

8b2875d

cuda: move the fused MMQ hint check to should_fuse and add TODOs

711b308

cuda: pass prec_src1 to the fused MMQ write-back

6bbeddf

cuda: fuse only with efficient tiling and no fallback, merge the GLU …

b6e0f03

…write-back

anujj force-pushed the cuda-mmq-fused-gate-up-interleaved branch from 0071207 to b6e0f03 Compare October 9, 2026 12:19

LXD-8 mentioned this pull request Oct 10, 2026

Feature Request: Metal prefill speedup (~68%) on M5 Max (NAX / tensor API kernels) #30287

Closed

4 tasks done

SongXiaoXi mentioned this pull request Oct 14, 2026

ggml-cuda: fuse GLU into the MMQ Q8_1 activation quantization #30489

Open

GenerelSchwerz mentioned this pull request Oct 16, 2026

cuda : emit routed MMQ down inputs from paired GLU GenerelSchwerz/llama.cpp#175

Open

	GGML_ASSERT(gate && !ids && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale);
	// TODO MoE support
	GGML_ASSERT(gate && !ids && !fusion->x_bias && !fusion->gate_bias && !fusion->x_scale && !fusion->gate_scale);

		// With fusion dst is the GLU output, not a MUL_MAT; the fused types always quantize src1 to Q8_1.
		const ggml_prec prec_src1 = gate ? GGML_PREC_Q8 : ggml_cuda_mmq_get_prec_src1(src0, dst, cc);

		const int jt = has_fusion ? blockIdx.x : blockIdx.y;
		const int it = has_fusion ? blockIdx.y : blockIdx.x;

	if constexpr (ggml_cuda_mmq_fusion_supported(type)) {
	if constexpr (!fallback && ggml_cuda_mmq_fusion_supported(type)) {

Conversation

anujj commented Oct 4, 2026

Implimentation

Performance

Correctness

Requirements

Uh oh!

anujj commented Oct 4, 2026

Uh oh!

anujj commented Oct 8, 2026

Uh oh!

JohannesGaessler commented Oct 8, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

anujj Oct 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

anujj commented Oct 12, 2026

Uh oh!

SongXiaoXi commented Oct 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

anujj commented Oct 14, 2026

Uh oh!

anujj commented Oct 14, 2026

Uh oh!

anujj commented Oct 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

This branch has not been deployed

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

anujj Oct 9, 2026 •

edited

Loading

SongXiaoXi commented Oct 14, 2026 •

edited

Loading

anujj commented Oct 15, 2026 •

edited

Loading