cuda, hip: GDN chunked kernel by am17an · Pull Request #29353 · ggml-org/llama.cpp

am17an · 2026-09-24T06:23:16Z

upstreaming a kernel developed in downstream QVAC fabric. Credits also to @lukdmine for initial idea for the kernel.

Performance

All rows use -b 2048 -ub 2048 -fa 1, full GPU offload and F16 KV cache. Retested on 2026-10-13 against master 296538a (includes #30087). TL;DR 3-6% better PP on dense 27B, 11-21% on 35B-A3B on RTX 4090/5090/PRO 6000, 3-6% on Spark.

Details 👇

Performance on	Model	Prompt tokens	master tok/s	pr tok/s	Change
RTX PRO 6000 Blackwell	Qwen3.5-35B-A3B-Q8_0	2048	13,653.58	16,459.56	+20.6%
RTX PRO 6000 Blackwell	Qwen3.5-35B-A3B-Q8_0	4096	14,090.26	16,595.47	+17.8%
RTX PRO 6000 Blackwell	Qwen3.8-27B-Q8_0	2048	4,470.39	4,757.20	+6.4%
RTX PRO 6000 Blackwell	Qwen3.8-27B-Q8_0	4096	4,442.45	4,699.08	+5.8%
Spark GB10	Qwen3.6-35B-A3B-UD-Q4_K_M	2048	3,271.18	3,406.68	+4.1%
Spark GB10	Qwen3.6-35B-A3B-UD-Q4_K_M	4096	3,254.79	3,366.88	+3.4%
Spark GB10	Qwen3.8-27B-UD-Q8_K_L	2048	942.84	984.86	+4.5%
Spark GB10	Qwen3.8-27B-UD-Q8_K_L	4096	928.16	973.34	+4.9%
Spark GB10	Qwen3.8-Flash-Next-UD-IQ4_XS (`-lzm on`)	2048	1,044.71	1,104.31	+5.7%
Spark GB10	Qwen3.8-Flash-Next-UD-IQ4_XS (`-lzm on`)	4096	1,068.70	1,124.01	+5.2%
RTX 5090	Qwen3.5-35B-A3B-Q4_K_S	2048	12,477.32	14,524.81	+16.4%
RTX 5090	Qwen3.5-35B-A3B-Q4_K_S	4096	12,978.57	15,298.00	+17.9%
RTX 5090	Qwen3.8-27B-Q4_K_M	2048	4,095.61	4,320.75	+5.5%
RTX 5090	Qwen3.8-27B-Q4_K_M	4096	4,114.11	4,353.09	+5.8%
RTX 4090	Qwen3.5-35B-A3B-Q4_K_S	2048	10,872.34	12,044.25	+10.8%
RTX 4090	Qwen3.5-35B-A3B-Q4_K_S	4096	11,187.34	12,428.05	+11.1%
RTX 4090	Qwen3.8-27B-Q4_K_M	2048	3,336.69	3,438.65	+3.1%
RTX 4090	Qwen3.8-27B-Q4_K_M	4096	3,340.90	3,446.88	+3.2%
RTX 3090	Qwen3.5-35B-A3B-Q4_K_S	2048	5,672.07	6,061.59	+6.9%
RTX 3090	Qwen3.5-35B-A3B-Q4_K_S	4096	5,505.49	5,683.79	+3.2%
RTX 3090	Qwen3.8-27B-Q4_K_M	2048	1,400.46	1,412.40	+0.9%
RTX 3090	Qwen3.8-27B-Q4_K_M	4096	1,337.41	1,367.46	+2.2%

Strix Halo (Radeon 8060S) was not retested this round. On the previous master it was +6.75% (pp2048) and +6.98% (pp4096) on Qwen3.8-27B-UD-Q8_K_L.

The RTX 3090 reported thermal throttling during these measurements, so its numbers should be read with that limitation.

Correctness

llama-perplexity --kl-divergence on wikitext-2

Matched scored-token perplexity: master 5.777325, pr 5.781777.
Perplexity difference: +0.004451 +/- 0.010372; the difference is smaller than the reported uncertainty in this short comparison.
Mean KL divergence: 0.000463 +/- 0.000029; median 0.000196; maximum 0.011380.
Top-token agreement: 98.824% +/- 0.338%.

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: YES, paired with Codex.

Ported-from: tetherto#257 Source-commit: 53f0027 Co-authored-by: Aman Gupta <aman.gupta@tether.io>

DanoPTT · 2026-09-24T07:26:58Z

I noticed the kernel only activates when K == 1 (ggml_get_op_params_i32(dst, 0) == 1).
With MTP/EAGLE3/DFlash speculation, n_rs_seq = draft.n_max, so K = n_max + 1 even during prefill (src/models/delta-net-base.cpp#L564), which means the kernel won't be used there.
Is that an intentional limitation, or is K > 1 planned?

am17an · 2026-09-24T07:28:55Z

it is a prefill kernel. So it won't be activated while drafting (i.e. when K > 1)

DanoPTT · 2026-09-24T07:46:42Z

I mean prefill with speculation enabled.
K is set for the context, not per ubatch (line 564), so with --spec-type draft-mtp enabled, K > 1 even for prefill on the target context.

am17an · 2026-09-24T08:01:05Z

Yes that's correct - prefill with speculation enabled would not use this kernel for now. It would require some changes from the llama.cpp side to use it for prefill, perhaps something like llama_set_rs_snapshots

IMbackK · 2026-09-24T12:27:54Z

Also, would be pretty easy to port this to rdna4's wmma regsiter layout and wave64 for mfma, might not perform better but might as well support everything in the kernel since its minimal effort, do you want to do it or shall i do a followup?

am17an · 2026-09-24T12:29:43Z

@IMbackK I don't have RDNA4 hardware. You can do it in a follow-up.

DanoPTT · 2026-09-24T13:26:09Z

When required, I can check it on gfx1201/RDNA (R9700, Windows 11, ROCm 10.0.0/HIP)

Enable SM86 and SM89 after RTX 3090 and RTX 4090 correctness and prefill benchmarks. Keep other untested architectures on the existing fallback. Assisted-by: Codex

JohannesGaessler

I currently still consider the swizzling code in mma.cuh to be unstable. If you intend to use it in this PR I would ask that we delay merging it until I've taken another look at it.

am17an · 2026-09-25T15:02:22Z

I intentionally used it for reuse purpose. I can keep it completely separate if it's easier for review.

JohannesGaessler · 2026-09-26T08:39:57Z

I've started work on a refactor of the swizzling code to include AMD GPUs as well as strides that are not multiples of 128 bytes. I should be able to open a corresponding PR by the end of the weekend at the latest.

IMbackK · 2026-09-26T09:41:08Z

Porbubly easier to just separate the prs? just remove the swizzling in this pr and add back later as as followup.

Changes made satisfy identified defects, but i have not fully re-reviewed so no approval for now.

0cc4m · 2026-09-28T05:59:27Z

I have a prototype Vulkan version of this. One thing I found is that CUDA's GDN scan kernel is pretty bad, about half as fast as Vulkan's, which blows up the relative improvement shown here.

ORippler · 2026-09-28T15:22:11Z

+    constexpr int N = 8;
+#endif // GGML_USE_HIP
+    static_assert(C >= 16 && (C & (C - 1)) == 0 && V % N == 0 && D % V == 0, "invalid GDN tile");
+    extern __shared__ __align__(128) unsigned char bytes[];


Given we only have one source of __shared__ in the kernel, we should be good. But note that extern __shared__ __align__ + __shared__ with optional align can get you to SMEM-bank-conflict-land quickly (it's not clearly documented how these are supposed to be handled/aligned iirc)

ORippler · 2026-09-28T15:22:40Z

+        __syncthreads();
+        for (int ti = warp; ti < (C / 16) * (V / N); ti += WARPS) {
+            const int r = (ti / (V / N)) * 16, c = (ti % (V / N)) * N;
+            acc       x;
+            product<C, false, true>(x, s.inverse, s.delta, r, c);
+            store(s.scratch.solved, x, r, c);
+        }
+        __syncthreads();


Feels like warp-specialization may help here, but something that can be done in a follow-up PR

ORippler · 2026-09-28T15:23:05Z

One thing I found is that CUDA's GDN scan kernel is pretty bad

It's tuned specifically for bs=1 - given the increasing popularity for spec-dec we could try to revive #22587

Porbubly easier to just separate the prs?

Swizzling may (or may not) be required for good performance of the implementation at hand.

Yes that's correct - prefill with speculation enabled would not use this kernel for now. It would require some changes from the llama.cpp side to use it for prefill, perhaps something like llama_set_rs_snapshots

Would love if we could prioritize this, as spec-dec is increasingly popular. LMK if I can assist here, though I'll focus on FP8 support for the time being

JohannesGaessler · 2026-09-28T20:19:09Z

Refactor of the swizzling code: #29612 . Sorry, I forgot that I had conflicting obligations on the weekend so it was a bit delayed. Practically speaking the way the interface changed is that you need to explicitly call _swizzled functions templated with the stride. For stride_in_bytes % 128 == 16 the code is equivalent to padding so in the kernel you only need to set a stride either with or without padding depending on whether or you intend to use swizzling or padding.

…s K == 1, we run K = 3); prefill unchanged Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_018co7fcJoEQXMxLezg7HpMU

…l-org#30087, drop ggml-org#29030)

0cc4m · 2026-10-09T06:33:15Z

Perf test of this branch compared to master and Vulkan scan/chunked in #30207

ORippler

From my side I'm generally fine with accepting & co-maintaining this, especially as it keeps accuracy high by approximating F32 via 2 BF16s and keeping states in F32 (though some more comments/variable naming could help). What I would like for us to do before merging:

Resolve roll-back-after-prefill-with-chunked-GDN gap (though this should not happen with current control flow)
Adapt swizzling to refactored #29612
Potentially recapture perf numbers as #30087 was merged in the mean-time. Given it brings ~4-8%, ~5-10% should be gained by this chunked implementation I guess.

For the interested, I looked to optimize it a bit more & to understand the code. Changes are over at am17an#59. It's about 10% operator-level speedup on a B4500 (and thus 0.5% E2E speed-up), so not really required to get this merged. Still was fun/helpful to me

am17an · 2026-10-12T11:39:35Z

@ggerganov can you guide on how to make 1) in the above happen?

ggerganov · 2026-10-12T11:43:11Z

@am17an For what batch sizes does the new kernel run in the CUDA backend?

am17an · 2026-10-12T11:58:07Z

It runs for batch sizes greater than 2048 mostly

ggerganov · 2026-10-12T12:37:41Z

I think the simplest way is to declare that ggml_gated_delta_net is not guaranteed to return extra states when n_tokens > K:

llama.cpp/ggml/include/ggml.h

Lines 2641 to 2661 in 097cf33

    
           // 
        
           // tensor shapes (S_k == S_v, H_v % H_k == 0): 
        
           //   q, k  : [S_k, H_k, n_tokens, n_seqs] 
        
           //   v     : [S_v, H_v, n_tokens, n_seqs] 
        
           //   g     : [1, H_v, n_tokens, n_seqs] (scalar gate) or [S_v, H_v, n_tokens, n_seqs] (KDA) 
        
           //   beta  : [1, H_v, n_tokens, n_seqs] 
        
           //   state : [S_v, S_v, H_v, n_seqs] -- initial recurrent state s0 
        
           // 
        
           // the output packs the attention scores [S_v, H_v, n_tokens, n_seqs] followed by K state 
        
           // snapshots, most-recent first (slot 0 = final state, slot s = state s tokens back). K == 1 
        
           // keeps only the final state; when n_tokens < K only slots 0..n_tokens-1 are written. 
        
           GGML_API struct ggml_tensor * ggml_gated_delta_net( 
        
                   struct ggml_context * ctx, 
        
                   struct ggml_tensor  * q, 
        
                   struct ggml_tensor  * k, 
        
                   struct ggml_tensor  * v, 
        
                   struct ggml_tensor  * g, 
        
                   struct ggml_tensor  * beta, 
        
                   struct ggml_tensor  * state, 
        
                   int64_t               K);

With this restriction we'll later add guardrails in llama_context to assert if we attempt rollbacks that violate this condition. I just want to merge the #30023 first and will add the necessary logic.

So for this PR, I think we can just update the ggml_gated_delta_net comment in ggml.h to state that.

ggerganov · 2026-10-12T12:39:05Z

If we agree on that, then the following tests would also need to be removed:

llama.cpp/tests/test-backend-ops.cpp

Lines 11930 to 11932 in 097cf33

    
           // overflow: n_tokens > K — only the last K snapshots kept. 
        
           test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 32,   8, 1, 1, false, false, /*K=*/3)); 
        
           test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 4, 64,  16, 2, 1, false, false, /*K=*/4));

# Conflicts: # ggml/src/ggml-cuda/gated_delta_net.cu # ggml/src/ggml-cuda/mma.cuh

am17an · 2026-10-13T12:35:35Z

@ORippler @ggerganov PTAL

ggerganov

The ggml comment as discussed is OK. I will follow-up with the necessary libllama changes.

Prefill-sized GATED_DELTA_NET ops run a chunked tensor-core kernel instead of the per-token recurrent kernel. The raw gates are activated into temporaries first. With K > 1 state snapshots (one sequence), the MMA kernel runs the first n - K tokens and the recurrent kernel the last K, which write the snapshot slots. The fused state gather is supported: the live state rows are copied to a contiguous s0 first. One WARN line when the chunked path first engages, and one when a prefill-sized op skips it. BF16 MMA on Ampere GA10x / Ada and newer. Other cards are checked before any temporaries are allocated and keep the recurrent kernel (one WARN line with the device cc). GGML_CUDA_GDN_MMA_DISABLE=1 turns it off. Tests in the 27B shape. RTX 4070: prefill +3 to +6 % (drafting on); perplexity 8.0714 -> 8.0733 (+0.024 %, -c 2048, 32 chunks); exact copies and needles unchanged. RTX 2060 SUPER (sm_75): not taken; 32k prefill and decode unchanged.

…lama.cpp#29353) (#40) * patches: 0054 chunked gated-delta-net MMA prefill (port of ggml-org/llama.cpp#29353); README: rows 0050 and 0054, patch counts 54 * patches: 0054 checks the device before the GDN MMA temporaries (RTX 20: clean skip line, no allocations) --------- Co-authored-by: Cary Palmer <24235924+professorpalmer@users.noreply.github.com>

…#29353) Prefill-sized GATED_DELTA_NET ops run a chunked tensor-core kernel instead of the per-token recurrent kernel. The raw gates are activated into temporaries first. With K > 1 state snapshots (one sequence), the MMA kernel runs the first n - K tokens and the recurrent kernel the last K, which write the snapshot slots. The fused state gather is supported: the live state rows are copied to a contiguous s0 first. One WARN line when the chunked path first engages, and one when a prefill-sized op skips it. BF16 MMA on Ampere GA10x / Ada and newer. Other cards are checked before any temporaries are allocated and keep the recurrent kernel (one WARN line with the device cc). GGML_CUDA_GDN_MMA_DISABLE=1 turns it off. Tests in the 27B shape. RTX 4070: prefill +3 to +6 % (drafting on); perplexity 8.0714 -> 8.0733 (+0.024 %, -c 2048, 32 chunks); exact copies and needles unchanged. RTX 2060 SUPER (sm_75): not taken; 32k prefill and decode unchanged.

…ml-org#29963 rewrite re-tested and rejected

153 upstream commits (merge-base 1623d8c). Conflicts (6 files) and semantic fixes: - gated_delta_net.cu: fork kernels (column_lanes, indexed state, RDNA4 chunked WY prefill) kept; upstream GDN MMA path (ggml-org#29353, HIP only on RDNA3.5) runs after ours, skipped when the state gather was deferred to the op. - ggml-cuda.cu: n_stable reset kept, warmup reset log behind GGML_CUDA_GRAPH_DEBUG (ggml-org#30327). - llama-batch: upstream state embeddings for MTP hooks (ggml-org#30257) next to the fork embd alias; R57 extended to them: the allocator points at the batch state rows and a ubatch of consecutive rows reads them in place (no copies). - llama-context: upstream sched_new (realloc callback, ggml-org#30329) also used for the second MTP scheduler. - llama-graph: draft LM head and logits top-k kept next to build_hadamard_input / build_logits; the draft head applies the Hadamard input transform too. - llama-model: load_tensors_hadamard, then build_mtp_draft_head. - common: R59 full-ubatch warmup also sets zero state embeddings.

… of ggml-org#29353) Ported from professorpalmer/llama.cpp-ada-mirai b6bf51cdd.

The 16-head cases of ggml-org/llama.cpp#29353 (b11632): 3 and 4 sequences, GQA 2-4, 257..2048 tokens. Taken now so they also cover our chunked kernels. One V100: 56/56 GATED_DELTA_NET pass with GGML_CUDA_GDN_CHUNKED on and off; under ncu the cases from 257 tokens run chunk_prep / chunk_scan<32> / chunk_out, the 64/65 token cases the sequential kernel. The n_tokens > K snapshot cases that the PR deletes stay: our rollback reads those slots. Assisted-by: Claude Opus 5.5

cuda, hip: GDN chunked kernel

7e1113f

Ported-from: tetherto#257 Source-commit: 53f0027 Co-authored-by: Aman Gupta <aman.gupta@tether.io>

am17an requested review from a team and ggerganov as code owners September 24, 2026 06:23

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 24, 2026

fix older builds

43c3f76

IMbackK previously requested changes Sep 24, 2026

View reviewed changes

Comment thread ggml/src/ggml-cuda/gated-delta-net-mma.cu Outdated

Comment thread ggml/src/ggml-cuda/gated-delta-net-mma.cu Outdated

Comment thread ggml/src/ggml-cuda/gated-delta-net-mma.cu

am17an added 2 commits September 24, 2026 23:22

review fixes

9962e80

cuda: enable GDN MMA on tested Ampere and Ada GPUs

377e111

Enable SM86 and SM89 after RTX 3090 and RTX 4090 correctness and prefill benchmarks. Keep other untested architectures on the existing fallback. Assisted-by: Codex

JohannesGaessler reviewed Sep 25, 2026

View reviewed changes

ORippler reviewed Sep 28, 2026

View reviewed changes

anskumar01 mentioned this pull request Sep 28, 2026

CUDA: Support of GDN chunked kernel for prefill #26001

Closed

SongXiaoXi mentioned this pull request Oct 7, 2026

ggml-cuda: assign four GDN state columns per warp #30087

Merged

sourenaraya added a commit to sourenaraya/llama.cpp that referenced this pull request Oct 7, 2026

docs: update INFERENCE.md to the 448147d rebase (ggml-org#29353 + ggm…

4e2623b

…l-org#30087, drop ggml-org#29030)

0cc4m mentioned this pull request Oct 9, 2026

vulkan: chunked GATED_DELTA_NET implementation #30207

Draft

ORippler reviewed Oct 9, 2026

View reviewed changes

am17an added 4 commits October 13, 2026 19:05

Merge remote-tracking branch 'upstream/master' into aman/gdn-mma-prefill

d098253

# Conflicts: # ggml/src/ggml-cuda/gated_delta_net.cu # ggml/src/ggml-cuda/mma.cuh

cuda : use refactored swizzle API in GDN MMA

6577068

ggml : GDN only guarantees the final state when n_tokens > K

6d5e4c4

cuda : add short GDN MMA phase comments

dde6a4d

ggerganov approved these changes Oct 13, 2026

View reviewed changes

edwardyoon pushed a commit to edwardyoon/focus-llama that referenced this pull request Oct 14, 2026

cuda, hip: GDN chunked kernel (ggml-org#29353)

d0e4699

ORippler approved these changes Oct 14, 2026

View reviewed changes

am17an merged commit 62df181 into ggml-org:master Oct 14, 2026
22 of 24 checks passed

gaetan-puleo mentioned this pull request Oct 14, 2026

Rebase onto upstream llama.cpp c114bd4c5 (2026-10-14) halo-box/strix-llama.cpp#178

Closed

3 of 5 tasks

ggerganov mentioned this pull request Oct 15, 2026

ggml : bump version to 0.27.0 ggml-org/ggml#1659

Merged

sourenaraya added a commit to sourenaraya/llama.cpp that referenced this pull request Oct 15, 2026

docs: INFERENCE.md - fd55cb9 base, ggml-org#29353 merged upstream, gg…

24f6716

…ml-org#29963 rewrite re-tested and rejected

BrewTestBot mentioned this pull request Oct 15, 2026

ggml 0.27.0 Homebrew/homebrew-core#318733

Merged

1 task done

benjigill mentioned this pull request Oct 18, 2026

Sync upstream master (2026-10-18, ggml-org 797de63b2) benjigill/llama.cpp#24

Merged

6 tasks

zgsndzq pushed a commit to zgsndzq/llama.cpp-qwen38-flash that referenced this pull request Oct 18, 2026

port mirai b6bf51cdd: cuda: chunked gated-delta-net MMA prefill (port…

842ce4f

… of ggml-org#29353) Ported from professorpalmer/llama.cpp-ada-mirai b6bf51cdd.

Conversation

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

Uh oh!

Performance

Correctness

Requirements

Uh oh!

DanoPTT commented Sep 24, 2026

Uh oh!

am17an commented Sep 24, 2026

Uh oh!

DanoPTT commented Sep 24, 2026

Uh oh!

am17an commented Sep 24, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

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

Uh oh!

Uh oh!

am17an commented Sep 24, 2026

Uh oh!

DanoPTT commented Sep 24, 2026

Uh oh!

JohannesGaessler left a comment

Choose a reason for hiding this comment

Uh oh!

am17an commented Sep 25, 2026

Uh oh!

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

Uh oh!

Uh oh!

IMbackK commented Sep 26, 2026

Uh oh!

0cc4m commented Sep 28, 2026

Uh oh!

Uh oh!

ORippler Sep 28, 2026

Choose a reason for hiding this comment

Uh oh!

ORippler Sep 28, 2026

Choose a reason for hiding this comment

Uh oh!

ORippler commented Sep 28, 2026

Uh oh!

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

Uh oh!

Uh oh!

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

Uh oh!

Uh oh!

ORippler left a comment • 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!

am17an commented Oct 12, 2026

Uh oh!

ggerganov commented Oct 12, 2026

Uh oh!

am17an commented Oct 12, 2026

Uh oh!

ggerganov commented Oct 12, 2026

Uh oh!

ggerganov commented Oct 12, 2026

Uh oh!

am17an commented Oct 13, 2026

Uh oh!

ggerganov left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

8 participants

am17an commented Sep 24, 2026 •

edited

Loading

IMbackK commented Sep 24, 2026 •

edited

Loading

JohannesGaessler commented Sep 26, 2026 •

edited

Loading

JohannesGaessler commented Sep 28, 2026 •

edited

Loading

0cc4m commented Oct 9, 2026 •

edited

Loading

ORippler left a comment •

edited

Loading