CUDA: fused chunked gated_delta_net kernel (RDNA3.5) - #54
Merged
Merged
Conversation
roberteg16
force-pushed
the
rogarcia.gdn-shared-kq
branch
from
July 16, 2026 13:51
e332179 to
5bb1972
Compare
This comment was marked as outdated.
This comment was marked as outdated.
roberteg16
force-pushed
the
rogarcia.gdn-shared-kq
branch
2 times, most recently
from
July 28, 2026 18:07
2619b60 to
fa2898d
Compare
3 tasks
roberteg16
force-pushed
the
rogarcia.gdn-shared-kq
branch
from
July 29, 2026 10:13
fa2898d to
f166cb5
Compare
roberteg16
force-pushed
the
rogarcia.gdn-shared-kq
branch
2 times, most recently
from
July 29, 2026 16:38
b7499a0 to
05ea898
Compare
roberteg16
requested review from
Annieren,
jimw567,
liangliangchang and
mgehre-amd
July 29, 2026 16:44
roberteg16
marked this pull request as ready for review
July 29, 2026 16:49
mgehre-amd
reviewed
Jul 29, 2026
GATED_DELTA_NET is a large share of prefill time on the Qwen3.5/3.6 delta-net models. The existing kernel is a token-by-token scan with one warp per state column, so k/q are re-fetched by every warp of a head on every token and the kernel ends up latency bound rather than limited by bandwidth or arithmetic. Add a chunked delta rule (WY representation / UT transform) as a single fused kernel: one block per (head, sequence), the state resident in registers and the per-chunk intermediates kept in LDS. Two things make a whole chunk fit in 64 KB of LDS: - A chunk of 32 tokens rather than 64. At 64 the chunk's U, W and qk alone are 80 KB. - U and W are never materialised. v_new = U - W S, with U = Tinv(V beta) and W = Tinv(K beta exp(g_cs)), is rewritten as the identical v_new = Tinv (V beta - (K beta exp(g_cs)) S), so one [chunk][head_dim] buffer holds V beta, becomes Y, then becomes v_new in place. The in-place triangular product needs no second buffer because the upper half of the rows is accumulated into registers before any row is overwritten. The two triangular products (Gram and qk) are packed across the block with a triangle-to-rectangle fold; the natural (tid/chunk, tid%chunk) mapping would make a wave issue a full dot product with most of its lanes masked off. The block stages its tiles through LDS and reduces every dot product across sixteen lanes, so the reduction competes with the staging for the same pipe when it is written as a shuffle butterfly, which lowers to ds_bpermute_b32. On RDNA the same reduction is a chain of DPP row_shl adds, where the cross-lane move is a modifier on the add and never leaves the VALU. The decay cumsum is likewise converted to base 2 once per chunk so each of the four decay tables is a single v_exp_f32 rather than an expf call with its denormal rescue; an exponent far enough negative to flush means a fully decayed state, which is the intended result. GDN op at 2048 tokens per ubatch: 220.9 -> 78.5 ms. End to end on Qwen3.5-35B-A3B at the default ubatch: +5.3% at pp128 and +7.9% at pp4096. The op only ever sees n_ubatch tokens, so a larger ubatch unlocks more. Scalar gate and K == 1 only, S_v == 128, more than one chunk, RDNA3.5 only; everything else falls back to the existing kernel, which is untouched. On by default; GGML_CUDA_GDN_CHUNKED=0 disables it. The occupancy request in __launch_bounds__ is guarded at compile time, since the kernel is still built for every target in the multiarch build and an unmet request is an error under -Wpass-failed. Also add five test-backend-ops cases: no existing GDN case combined head_size=128 with more than 64 tokens, and none covered GQA (v_repeat > 1), which these models use. Assisted-by: Claude Opus 4.7
roberteg16
force-pushed
the
rogarcia.gdn-shared-kq
branch
from
July 30, 2026 07:59
05ea898 to
6417ef6
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What
GATED_DELTA_NETis 18% of prefill GPU time on Qwen3.5-35B-A3B / gfx1151. The existing kernel is a token-by-token scan with one warp per state column, sok_tandq_tare re-fetched by all 128 warps of a head on every token and the kernel ends up latency bound rather than limited by bandwidth or arithmetic.This replaces it, for the scalar-gate case, with a chunked delta rule (WY representation / UT transform) as a single fused kernel: one block per (head, sequence), the state resident in registers, per-chunk intermediates kept in LDS, no scratch buffer.
The kernel stages its tiles through LDS and reduces every dot product across sixteen lanes, so on RDNA the reduction is a chain of DPP
row_shladds rather than a shuffle butterfly: the cross-lane move is a modifier on the add and never leaves the VALU, where__shfl_xorlowers tods_bpermute_b32and contends with the staging for the LDS pipe. The decay cumsum is converted to base 2 once per chunk so each decay table is a singlev_exp_f32instead of anexpfcall with its denormal rescue.The existing kernel is untouched apart from a dispatch branch.
Results
GDN op at 2048 tokens per ubatch: 220.9 -> 78.5 ms (2.81x).
End to end against
0f0db6292, interleaved, 3 iterations, medians, at the default-ub 512:Decode is unaffected: the op's
n_tokensis tokens per sequence, so it is 1 during generation and the multi-chunk condition never holds.The gain is bounded by
-ub, not by prompt lengthThe op only ever sees
n_ubatchtokens, so a 4096-token prompt at the default ubatch is eight separate 512-token GDN calls. Raising the ubatch unlocks the rest, at pp4096 on Qwen3.5-35B-A3B:With the sequential kernel a large ubatch is a pessimisation (1468 vs 1656) because the scan's serial length grows with
n_tokens; with this kernel it becomes the best configuration. Best-config to best-config that is +14.2%, at the cost of larger activation buffers.Scope
On by default for RDNA3.5 with the scalar gate,
K == 1,S_v == 128and more than one chunk of tokens.GGML_CUDA_GDN_CHUNKED=0disables it. Everything else falls back to the existing kernel.Testing
test-backend-ops -o GATED_DELTA_NET: 42/42 pass with the path enabled and disabled. Five cases added, because no existing GDN case combinedhead_size=128with more than 64 tokens, and none covered GQA (v_repeat > 1), which these models use.