Change set 13: chunked gated-delta-net on CDNA (+12% prefill) - #28
Merged
Conversation
Four defects kept upstream PR #26001's ggml_cuda_mma path from working on gfx90a: a 32-thread block against 64-lane MFMA tiles, __launch_bounds__(32) blocking the wider launch, plain load_ldmatrix having no AMD_MFMA branch, and -- the real one -- the 16x16 f32 accumulator declared I_MAJOR when CDNA needs J_MAJOR, which transposed every writeback. Measured on 2x MI210, Qwen3.6-40B IQ4_NL: pp2048 606->678 (+11.8%), pp8192 911->1021 (+12.1%). test-backend-ops 50/51 on both devices; temp-0 output byte-identical to the recurrent path. Patch 13 applies to PR #26001, not the pinned base, so it is deliberately excluded from the Dockerfile and documented separately in How to apply. Also refreshes that section, which still said 'change sets 4-10'.
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
1 participant
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.
Makes upstream PR #26001's chunked GDN kernel work on gfx90a. Four defects, the load-bearing one being the accumulator fragment layout (
DATA_LAYOUT_J_MAJORon CDNA, notI_MAJOR) which transposed every writeback.Measured on 2x MI210, Qwen3.6-40B IQ4_NL, same binary toggled by
GGML_CUDA_DISABLE_GDN_CHUNK:test-backend-ops -o GATED_DELTA_NET50/51 on both devices (from 36/51); temp-0 output byte-identical to the recurrent path. The one failure is a precision margin at n_seq_tokens=2048 (NMSE ~3e-7 vs 2e-7), disclosed rather than papered over.Patch 13 targets an unmerged upstream PR, so it is excluded from the Dockerfile and documented separately under How to apply — which this also refreshes, since it still said 'change sets 4-10'.