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

anujj · 2026-09-10T13:36:11Z

Overview

Extends the gate/up + GLU fusion, which today only covers small batches (MMVQ/MMVF), to prompt-processing batch
sizes on the MMQ path. For dense FFN layers where ffn_gate and ffn_up have the same type and shape, the two
matmuls and the GLU now run as one MMQ kernel: the activations are quantized once, the GLU is applied in the
write-back, and the standalone GLU kernel and the second activation quantization disappear.

Enabled for Q4_K, Q5_K, Q6_K, Q8_0 and IQ4_XS on GPUs with the MMA data layout; other cases take the existing path, I have a plan to enable it other types also.

Implementation:

mul_mat_q gets a fused variant (has_fusion): every thread keeps separate up and gate accumulators, so the
block uses half the threads and rows of the stock tile (128 threads, 64 rows, 2 blocks/SM) and J is capped at 64.
Per K step both Y halves are loaded once (cp.async where available, plain copy otherwise), then the up and the gate
tiles are loaded and multiplied in turn. No stream-K: the GLU needs the full K reduction in one block. Column tiles
are the innermost grid dimension, so neighbouring blocks reuse the same weight tiles from L2.
Only the q8_0 / q4_K / q5_K / q6_K / iq4_xs loaders take the has_fusion parameter; the fused kernel is only
compiled for the MMA layout.
ggml_cuda_op_glu_single (unary.cuh) is now shared by the fused MMQ write-back, MMVQ and MMVF; the padding clear
in ggml_cuda_mul_mat_q / MMVQ moved to ggml_cuda_clear_padding.
Graph side: the MUL_MAT + MUL_MAT + GLU pattern is matched in ggml_cuda_try_fuse; graph_optimize keeps the
activation alive until the GLU node so the allocator cannot alias it with the fused output.

Additional information

Numbers for different Q4_K_M models

DGX Spark (GB10), CUDA 13.1, pp16384, -b 2048 -ub 2048 -fa on, f16 KV, t/s, master and PR in one session:

model (Q4_K_M)	master	PR	gain
Llama-3.1-8B-Instruct	2704.50	3081.35	+13.9%
Qwen3-8B	2610.06	2940.20	+12.6%
Qwen3.6-27B	800.85	875.14	+9.3%
Qwen3.8-27B	834.54	916.31	+9.8%
Qwen3.8-27B UD-Q4_K_M (mixed types)	844.44	873.27	+3.4%

Numbers with supported GGML types in this PR

DGX Spark (GB10), CUDA 13.1, pp16384, -b 2048 -ub 2048 -fa on, f16 KV, t/s:

model (gate/up type)	master	PR	gain
Qwen3.8-27B Q4_K_M (Q4_K)	838.42	919.68	+9.7%
Qwen3.8-27B Q5_K_M (Q5_K)	833.29	872.99	+4.8%
Qwen3.8-27B Q6_K (Q6_K)	714.58	747.19	+4.6%
Qwen3.8-27B Q8_0 (Q8_0)	876.22	901.18	+2.8%
Qwen3.8-27B IQ4_XS (IQ4_XS)	876.91	914.39	+4.3%
Qwen3.8-27B UD-Q4_K_XL (mixed types)	828.69	861.88	+4.0%
Gemma 3 27B-it Q4_K_M (Q4_K, GeGLU)	904.59	1025.26	+13.3%

WikiText-2 perplexity is bit-identical to master (all 36 chunks) for every model above.
test-backend-ops: MUL_MAT_FUSION and the new MUL_MAT_FUSION_DOWN (gate/up + GLU + down projection) pass; the
full suite passes.
UD-Q4_K_XL gains less because 25 of its 64 FFN layers have different gate and up types and stay unfused.
Not covered: MoE (mul_mat_id), mixed gate/up types, have a plan to extend this fusion to the MOE also in follow up PR Only tested on NVIDIA; AMD MFMA/WMMA testing would be appreciated, I will also try I arrange the AMD also and test.

Requirements

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

anujj · 2026-09-13T08:19:23Z

moving to review state, adding @ORippler and @gaugarg-nv for review

ORippler

What's your proposal of extending this to other data-types? Ideally, we want to apply this fusion to all supported ggml_type's in the long term, and I fear simply adding a new kernel will fragment things.
We should look to make this code potentially compatible with HIP parts as well. I guess current blockers on this front are the async copies.

ORippler · 2026-09-18T16:06:54Z

    return use_mul_mat_vec_f;
 }

+static bool ggml_cuda_mul_mat_bad_padding_clear(const ggml_tensor * src0) {


Suggested change

static bool ggml_cuda_mul_mat_bad_padding_clear(const ggml_tensor * src0) {

static bool ggml_cuda_mul_mat_bad_padding_clear(const ggml_tensor * src) {

The helper only access a single tensor, so suffix is unneeded

ORippler · 2026-09-18T16:10:38Z

+            if (i + 2 < cgraph->n_nodes && cgraph->nodes[i + 2]->op == GGML_OP_GLU) {
+                ggml_tensor * glu  = cgraph->nodes[i + 2];
+                ggml_tensor * gate = cgraph->nodes[i];
+                ggml_tensor * up   = cgraph->nodes[i + 1];
+                if (glu->src[0] == gate && glu->src[1] == up &&
+                        ggml_cuda_mul_mat_q_gate_up_swiglu_matches(up, gate, glu, cuda_ctx->device)) {
+                    params->add_alloc_dep(params->user_data, up->src[1], glu);
+                    i += 2;
+                    continue;
+                }


Please do add_alloc_dep for both orderings of gate/up being passed into GLU

Done, graph_optimize now accepts gate/up in either order before adding the alloc dep

ORippler · 2026-09-18T16:19:30Z

+        mmq_gate_up_swiglu_copy_y<J, nwarps, warp_size>(y + ncols_y * (kb0 * qk / QK8_1_MMQ) * sz, ncols_y, tile_y0);
+
+        mmq_gate_up_swiglu_load_x<type, J, fallback>(x_up, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x);
+        cp_async_wait_all();


Is the register pressure/L1 utilization so high at this point to warrant an async copy? What's the perf with regular copies? Note LDGSTS requires 128-byte alignment to be SMEM-conflict-free iirc, at least if done via cuda::memcpy_async.

Can use NCU to analyze this

ORippler · 2026-09-18T16:20:35Z

+static constexpr int GGML_CUDA_MMQ_GATE_UP_SWIGLU_J_MAX = 64;
+static constexpr int GGML_CUDA_MMQ_GATE_UP_SWIGLU_I     = 64;


How were these magic constants determined?

They are no longer hard-coded: the fused config is derived from the stock per-arch config (ggml_cuda_mmq_get_config_fusion halves threads and rows, e.g. 128 threads / I = 64 at two blocks per SM on
Ampere+). GGML_CUDA_MMQ_FUSION_J_MAX = 64 is the register limit: with gate and up accumulators per thread, J = 64 already uses 255 registers without spills (J = 8/16/32: 213/253/254); J = 128 would double
the accumulators. The half-size config beat the stock config with two accumulators (Spark pp16K +8.1% vs +5.8%), and forcing the current kernel to one block per SM makes it ~45% slower.

ORippler · 2026-09-18T16:28:25Z

Effectively we are now testing mul_mat_fusion, and no longer mul_mat_vec. Please rename or justify why current naming is still accurate

The rename is done, and the latest commit also widened what the test covers.

ORippler · 2026-09-18T16:34:12Z

+    for (int kb0 = 0; kb0 < kb0_stop; kb0 += blocks_per_iter) {
+        mmq_gate_up_swiglu_copy_y<J, nwarps, warp_size>(y + ncols_y * (kb0 * qk / QK8_1_MMQ) * sz, ncols_y, tile_y0);
+
+        mmq_gate_up_swiglu_load_x<type, J, fallback>(x_up, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x);
+        cp_async_wait_all();
+        __syncthreads();
+
+        ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma<type, J, fallback>(tile_x, tile_y0, sum_up, 0);
+        ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma<type, J, fallback>(tile_x, tile_y1, sum_up, MMQ_TILE_NE_K);
+        __syncthreads();
+
+        mmq_gate_up_swiglu_load_x<type, J, fallback>(x_gate, tile_x, offset_x + kb0, tile_x_max_i, stride_row_x);
+        __syncthreads();
+
+        ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma<type, J, fallback>(tile_x, tile_y0, sum_gate, 0);
+        ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma<type, J, fallback>(tile_x, tile_y1, sum_gate, MMQ_TILE_NE_K);
+        __syncthreads();
+    }


Have you tried pipe-lining/warp-specialization (though this would make cross-IHV use worse)? How high is tensor-core utilization in this variant

Tensor-pipe utilization (Nsight Compute, GB10): Q4_K is 37% fused versus 32% stock; Q8_0 and IQ4_XS are 28% versus 34%. Tensor cores do not appear to be the limiter. For Q4_K, the bottleneck is FP32 scaling around each MMA; for Q8_0 and IQ4_XS, it is weight loads and decoding.

Pipelining: The Y tile is copied with cp.async while the weights load. This gives about a 1% gain at pp16K. Running two blocks per SM hides most of the latency; one block per SM is 42–47% slower. Register-staged weight prefetch gave a 0.7% gain in an earlier version, but this version has no register budget left (255 registers).

Warp specialization: I tried splitting gate and up across two warp groups with a 128-token tile. It helps Q8_0 and IQ4_XS but hurts Q4_K and Q5_K, so it may be worth a per-type follow-up. I have not tried a producer/consumer design; that would also be NVIDIA-specific.

anujj · 2026-09-22T14:35:34Z

What's your proposal of extending this to other data-types? Ideally, we want to apply this fusion to all supported ggml_type's in the long term, and I fear simply adding a new kernel will fragment things.

I’m refactoring how the decode phase handles this fusion. This will help avoid kernel fragmentation and extend the fusion support to all GGML types.

…tream ggml-org#28702) Assisted-by: Claude Opus 5.5

…ill (upstream ggml-org#28702)" KLD vs master under -sm tensor: mean 0.0098, max 19.1 (master vs master: 0). Assisted-by: Claude Opus 5.5

…tream ggml-org#28702, re-apply) Verified on 2x RTX 5070/5070 Ti, -sm tensor, q8_0/q4_0 KV, wikitext-2 8x4096: - vs master: mean KLD 0.0098, same top p 97.46% - master -sm layer vs master -sm tensor: mean KLD 0.0079, same top p 97.39% - fused vs unfused per call: max abs err ~2.5e-5, no out-of-bounds writes pp4096 +2-3.5%, tg unchanged. Assisted-by: Claude Opus 5.5

Prefill-only fusion of the two dense Q4_K FFN projections and the following SwiGLU into one MMQ kernel. The activation is quantized to Q8_1 once and both gate and up tiles are computed against the same shared copy of it; up * silu(gate) is applied in the F32 write-back, so the separate gate matmul and the SwiGLU kernel disappear. The kernel uses the stock int8 MMA dot product and produces results identical to the unfused path. It runs with 128 threads and 64 rows per block (two blocks per SM) and a vectorized Q4_K tile loader with bank-conflict-free shared memory stores. Dispatched only for Q4_K weights, F32 activations, split SwiGLU and dense MUL_MAT where MMQ would have been selected, on devices with int8 MMA and cp.async. test-backend-ops gains dense Q4_K fusion cases and a whole-FFN case (MUL_MAT_Q_GATE_UP_SWIGLU_DOWN) that feeds the fused output into a quantized down projection.

Replace the Q4_K-only fused kernel with a has_fusion template flag on mul_mat_q, mirroring the mmvq decode fusion: both Y half tiles stay resident, a second accumulator holds the gate product, and the GLU is applied during write-back. Fused instances run with half the threads and rows at two blocks per SM (J <= 64), copy Y with cp.async where available (with a synchronous fallback otherwise), and are compiled for Q4_K, Q5_K, Q6_K, Q8_0, and IQ4_XS. SwiGLU, GeGLU, SwiGLU-OAI, and SwiGLU-clamp are handled by ggml_cuda_op_glu_single, which is now shared with the mmvq and mmvf fusions.

am17an · 2026-09-23T16:14:06Z

How much more effective is this vs fusing the gate-up weights? we can possibly extend some allocation primitive to ask for gate and up to contigiously allocated

The helper only inspects a single tensor, so the src0 name is misleading.

The fused MMQ kernel reads the up and gate weights while it writes the GLU output. Under op offload the weights are copies in the CUDA compute buffer, which the allocator could reuse for the GLU output once the matmul nodes are done, corrupting the weights the kernel still reads (NaN logits with -ot ffn_gate=CPU at prompt batch sizes). Add alloc deps for both weights next to the existing one for the activation; for weights in weight buffers they have no effect.

The fused write-back applies SWIGLU_OAI with the default alpha and limit of ggml_cuda_op_swiglu_oai_single and ignores the values of the GLU node. Decline the fusion when they differ so those nodes take the unfused path, whose GLU kernel reads them. The in-tree users (gpt-oss, MiniMax-M3) use the defaults.

Restore the early-return layout of master instead of wrapping the stream-k part in an else: only its two mul_mat_q_process_tile calls change, and qk is declared at the top of the kernel again. The generated code is unchanged.

mmq.cuh now includes cp-async.cuh, which puts it into many translation units. Guard it against double inclusion like the other CUDA headers.

The MMQ-size fusion cases only ran the fused kernel with bounds-checked rows and small tiles, so the variant used for prompt processing (J = 48 and 64, row count a multiple of 128) was not tested. Add 48- and 128-token cases with 256 rows for each fused type.

anujj · 2026-09-28T10:41:15Z

How much more effective is this vs fusing the gate-up weights? we can possibly extend some allocation primitive to ask for gate and up to contigiously allocated

Thanks, I prototyped it to answer this: the loader allocates ffn_gate/ffn_up as the two halves of one tensor and
build_ffn does one mul_mat + two views + GLU. DGX Spark (GB10), Qwen3.8-27B

Time per FFN layer for gate/up (+ GLU), ms:

type	master	contiguous gate/up	this PR
Q4_K	12.58	12.43	8.88
Q5_K	12.76	12.30	10.02

With contiguous allocation alone, pp16K improves by only +0.6% for Q4_K, +1.2% for Q5_K, compared with +9.7%, +4.8% from this PR's fusion.

anujj · 2026-09-28T10:45:12Z

What's your proposal of extending this to other data-types? Ideally, we want to apply this fusion to all supported ggml_type's in the long term, and I fear simply adding a new kernel will fragment things.

We should look to make this code potentially compatible with HIP parts as well. I guess current blockers on this front are the async copies.

@ORippler : Done with the refactoring, i am now following how this fusion is implemented for the decode.
The fusion no longer adds a kernel: it is a has_fusion variant of the stock mul_mat_q that reuses the per-type loaders, vec_dot and write-back. Enabling another type means giving its loader the
has_fusion parameter and adding it to ggml_cuda_mmq_fusion_supported. I enabled Q4_K, Q5_K, Q6_K, Q8_0 and IQ4_XS for now, the gate/up types of the common quants.

The async Y copy falls back to regular copies where CP_ASYNC_AVAILABLE is not defined (HIP, pre-Ampere). The fused variant uses the per-arch MMQ configs, so it is enabled for AMD

am17an · 2026-09-28T10:51:18Z

With contiguous allocation alone, pp16K improves by only +0.6% for Q4_K, +1.2% for Q5_K, compared with +9.7%, +4.8% from this PR's fusion.

So what's the real speed-up here? fusing the SWIGLU?

anujj · 2026-09-28T10:56:21Z

With contiguous allocation alone, pp16K improves by only +0.6% for Q4_K, +1.2% for Q5_K, compared with +9.7%, +4.8% from this PR's fusion.

So what's the real speed-up here? fusing the SWIGLU?

The actual speedup is 0.6% for Q4_K and 1.2% for Q5_K. I haven’t fused SwiGLU yet; I’ll share the numbers after adding the fusion.

am17an · 2026-09-28T12:57:49Z

We should also try interleaving them so that we can keep one accumulator. i.e. warp N calculates up and warp N + 1 calculates gate and we can do silu(g) * u at the writeback. I think it should be faster than this PR and also less changes.

anujj · 2026-09-28T16:35:47Z

@am17an Agreed on keeping one accumulator. I looked at interleaved vs contiguous gate/up, and prototyped the one-accumulator fusion on contiguous weights.

Interleaved vs contiguous

For the fused kernel both should perform the same. A block reads the same bytes in the same pieces either way: for Q4_K, one 144-byte block per row per K step, with rows 2880 bytes apart. Let me know if you think otherwise

Interleaving needs extra work outside the kernel:

the loader has to shuffle gate/up rows into chunks at load time;
gate/up are no longer plain matrices, so the existing decode fusion (MMVQ gate + up + GLU) stops matching
the tensor split needs a new rule; it already handles [gate; up] for ffn_up with 2*n_ff in llama-model.cpp;
the chunk size becomes tied to the MMA accumulator layout, which differs by arch (NVIDIA: a thread holds rows r and r+8; CDNA/RDNA4: consecutive rows), so one layout baked into the weights can't fit every GPU.

Contiguous [gate; up] is already used: Phi-3, ChatGLM, GLM4 and PLaMo2 store ffn_up as 2*n_ff and apply ggml_swiglu to the two halves, and MoE ffn_gate_up_exps runs one mul_mat_id plus two views. So I'd keep the gate/up pairing in the kernel and the weights contiguous, or even separate, since the kernel only needs two pointers.

Prototype: contiguous + one accumulator

ffn_gate/ffn_up are loaded as the two halves of one tensor. The kernel uses master's 128 x 128 tile with 8 warps:

warps 0-3 compute 64 gate rows and warps 4-7 the matching 64 up rows, with one accumulator per thread;
after the K loop, the gate sums go through shared memory and the up warps apply silu(g) * u in the write-back.

DGX Spark (GB10), Qwen3.8-27B. Gate/up (+ GLU) time per FFN layer, and end-to-end pp16K gain vs master:

type	master	this PR	contiguous + one acc.	pp16K end-to-end, this PR	pp16K end-to-end, contiguous + one acc.
Q4_K	12.58 ms	8.88 ms	10.97 ms	+9.7%	+4.1%
Q5_K	12.76 ms	10.02 ms	10.81 ms	+4.8%	+4.0%
Q8_0	11.92 ms	10.63 ms	9.91 ms	+2.8%	+5.2%
IQ4_XS	11.57 ms	10.38 ms	9.00 ms	+4.3%	+8.1%

Why each wins where it does

The two kernels split the work differently:

this PR: half-size blocks (4 warps, 64 rows x 64 tokens, gate and up accumulators per thread), 2 blocks per SM;
the prototype: master's tile (8 warps, 64 gate + 64 up rows x 128 tokens), 1 block per SM.

Q4_K/Q5_K: Q4_K is limited by the FP32 scale/min math around each MMA, and Q5_K has the same scale structure. The second block per SM hides that latency: forcing this PR's kernel to one block per SM makes it 42-47% slower.

Q8_0/IQ4_XS: these are limited by loading and decoding weights. Q8_0 has the largest weights (~1.06 bytes per weight) and IQ4_XS decodes through a lookup table. The 128-token tile loads and decodes each weight once per 128 tokens instead of once per 64: 16 times instead of 32 per 2048-token ubatch.

Tensor-pipe utilization matches this. This PR vs master is 37% vs 32% for Q4_K, but 28% vs 34% for Q8_0 and IQ4_XS.

So the one-accumulator design fits Q8_0/IQ4_XS better, and this PR's kernel fits Q4_K/Q5_K better. How would you like to proceed: keep this PR's kernel and add the one-accumulator path for Q8_0/IQ4_XS as a follow-up, or pick the variant per type in this PR?

am17an · 2026-09-28T16:38:58Z

warps 0-3 compute 64 gate rows and warps 4-7 the matching 64 up rows, with one accumulator per thread;
after the K loop, the gate sums go through shared memory and the up warps apply silu(g) * u in the write-back.

I think this can be avoided with the idea before that the interleaving is across 16 rows (i.e. the MMA tile width), that way 1 accumulator keeps the gate tile and 1 keeps the up tile. No shared memory interlude needed

anujj · 2026-09-29T09:09:42Z

warps 0-3 compute 64 gate rows and warps 4-7 the matching 64 up rows, with one accumulator per thread;
after the K loop, the gate sums go through shared memory and the up warps apply silu(g) * u in the write-back.

I think this can be avoided with the idea before that the interleaving is across 16 rows (i.e. the MMA tile width), that way 1 accumulator keeps the gate tile and 1 keeps the up tile. No shared memory interlude needed

okey, so each warp covers 32 rows as two 16-row MMA tiles, so with gate rows in the first and up rows in the second, every thread holds g and the matching u at the same positions and can apply silu(g) * u straight from its registers.

I'd do the interleave in the loader rather than in the weights:

Weights and graph unchanged: ffn_gate/ffn_up stay separate (or contiguous), and the graph stays gate mul_mat + up mul_mat + GLU, so the decode fusion, the other backends and the tensor split are untouched.
Loader: takes the gate and up pointers and fills the shared-memory tile with gate and up rows interleaved in chunks of 8 (the C fragment gives each thread rows r and r+8 of a 16-row MMA tile, so this pairs them for every J, including J <= 40 where a warp has a single tile). Only the source row address changes, through one helper; same bytes, same decoding. Each loader gets the gate pointer, only the fused types use it.
MMA loop: master's kernel as is: same tile, same warp layout, same vec_dot, one sum[] per thread.
Write-back: each thread combines its up and gate values with silu(g) * u and writes half the rows. The fused path uses the regular tiling (no stream-k), since the GLU needs the full K sum.
AMD: the accumulator holds consecutive rows there, so it needs a 2-row pattern, a per-arch constant in the loader. I'd start with NVIDIA and add AMD once I can test it.

Does this match what you had in mind?

anujj · 2026-09-30T14:48:01Z

below are the perf numbers i am getting with the approach discussed in the #28702 (comment)

DGX Spark (GB10), Qwen3.8-27B, pp16K end-to-end (-b/-ub 2048, fa on, f16 KV):

type	PR tok/s	PR vs master	interleaved tok/s	interleaved vs master	interleaved vs PR
Q4_K_M	923.4	+9.7%	899.9	+7.1%	-2.5%
Q5_K_M	879.5	+4.9%	880.6	+4.8%	+0.1%
Q6_K	747.2	+4.6%	785.8	+4.6%	tie*
Q8_0	899.9	+2.3%	940.1	+7.3%	+4.5%
IQ4_XS	910.8	+3.9%	948.0	+7.4%	+4.1%

am17an · 2026-09-30T15:00:39Z

can you create a separate PR or branch with that approach?

anujj · 2026-09-30T16:07:45Z

sure, I will share the changes soon

anujj · 2026-10-01T13:20:33Z

can you create a separate PR or branch with that approach?

Changes are in the branch: https://github.com/anujj/llama.cpp/tree/cuda-mmq-fused-gate-up-interleaved

am17an · 2026-10-02T09:20:05Z

Please create a draft PR with that change after a self review

matcher Upstream ggml-org#29633 added a warp_size parameter; the fused Q4_K gate/up SwiGLU matcher still used the old signature and failed to compile.

anujj requested review from a team and ggerganov as code owners September 10, 2026 13:36

anujj marked this pull request as draft September 10, 2026 13:36

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 Sep 10, 2026

anujj marked this pull request as ready for review September 13, 2026 08:18

ORippler reviewed Sep 18, 2026

View reviewed changes

ORippler self-assigned this Sep 18, 2026

benjigill added a commit to benjigill/llama.cpp that referenced this pull request Sep 22, 2026

CUDA: fuse dense Q4_K gate/up projections and SwiGLU for prefill (ups…

5ad4d7b

…tream ggml-org#28702) Assisted-by: Claude Opus 5.5

anujj added 4 commits September 23, 2026 14:00

cuda : preserve input lifetime for Q4_K gate/up fusion

fbd4c4c

remove extra comments

e8497d9

anujj force-pushed the q4k-gate-up-swiglu-fp32-output branch from 13ecfb0 to 27ac308 Compare September 23, 2026 08:50

cuda: pass fused gate weights as __restrict__ to restore read-only loads

2163399

benjigill mentioned this pull request Sep 23, 2026

Sync upstream master (2026-09-23, ggml-org 6e60f3560) benjigill/llama.cpp#8

Merged

4 tasks

cuda: limit fused MMQ changes to supported types and MMA layout

ed85722

anujj changed the title ~~CUDA: fuse dense Q4_K Gate/Up projections and SwiGLU for prefill~~ CUDA: fuse FFN gate/up matmuls and GLU in MMQ Sep 24, 2026

anujj added 6 commits September 28, 2026 14:03

cuda: rename the tensor parameter of ggml_cuda_mul_mat_bad_padding_clear

0b6d253

The helper only inspects a single tensor, so the src0 name is misleading.

cuda: add #pragma once to cp-async.cuh

4d661fd

mmq.cuh now includes cp-async.cuh, which puts it into many translation units. Guard it against double inclusion like the other CUDA headers.

benjigill mentioned this pull request Sep 28, 2026

Sync upstream master (2026-09-28, ggml-org 6c7a87f7e) benjigill/llama.cpp#12

Merged

IMbackK self-assigned this Sep 28, 2026

anujj mentioned this pull request Sep 29, 2026

CUDA: fuse ffn_gate + glu into the mul_mat_q epilogue #27341

Closed

GenerelSchwerz mentioned this pull request Oct 2, 2026

cuda : pair routed MMQ gate/up projections GenerelSchwerz/llama.cpp#122

Draft

anujj mentioned this pull request Oct 4, 2026

CUDA: fuse FFN gate/up matmuls and GLU in MMQ with interleaved gate/up rows #29948

Open

benjigill mentioned this pull request Oct 6, 2026

Sync upstream master (2026-10-05, ggml-org 50569eb87) benjigill/llama.cpp#17

Merged

5 tasks

benjigill mentioned this pull request Oct 8, 2026

Sync upstream master (2026-10-08, ggml-org a11f57ba9) benjigill/llama.cpp#18

Open

2 of 6 tasks

	static bool ggml_cuda_mul_mat_bad_padding_clear(const ggml_tensor * src0) {
	static bool ggml_cuda_mul_mat_bad_padding_clear(const ggml_tensor * src) {

		static constexpr int GGML_CUDA_MMQ_GATE_UP_SWIGLU_J_MAX = 64;
		static constexpr int GGML_CUDA_MMQ_GATE_UP_SWIGLU_I = 64;

Conversation

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

Uh oh!

Overview

Additional information

Numbers for different Q4_K_M models

Numbers with supported GGML types in this PR

Requirements

Uh oh!

anujj commented Sep 13, 2026

Uh oh!

ORippler left a comment

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!

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 Sep 22, 2026

Uh oh!

am17an commented Sep 23, 2026

Uh oh!

anujj commented Sep 28, 2026

Uh oh!

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

Uh oh!

Uh oh!

am17an commented Sep 28, 2026

Uh oh!

anujj commented Sep 28, 2026

Uh oh!

am17an commented Sep 28, 2026

Uh oh!

anujj commented Sep 28, 2026

Uh oh!

am17an commented Sep 28, 2026

Uh oh!

anujj commented Sep 29, 2026

Uh oh!

anujj commented Sep 30, 2026

Uh oh!

am17an commented Sep 30, 2026

Uh oh!

anujj commented Sep 30, 2026

Uh oh!

anujj commented Oct 1, 2026

Uh oh!

am17an commented Oct 2, 2026

Uh oh!

This branch has not been deployed

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

anujj commented Sep 10, 2026 •

edited

Loading

anujj commented Sep 28, 2026 •

edited

Loading