Repository navigation
feat(cuda): tensor-core decode attention composed with GQA-group grid, default ON (#199) - #243
Merged
Merged
Conversation
…lit-KV grid (#199 v2) v1 (issue/199-tensor-core-decode-attention, not merged) built a decode-only tensor-core kernel with genuine HMMA/LDSM SASS but scoped to one warp/block, grid=numHeads=24 -- measured 4-5x SLOWER than attention_f32 at every realistic depth (~4% theoretical occupancy). This v2 composes the tensor-core math with #197/#198's GQA-group + split-KV grid design instead: grid=(numKvHeads, kvSplit), with the GQA group packed into the mma.sync instruction's M=16 dimension (free -- same instruction count as v1, up to 8x more useful throughput) and PV work split 8-way across warps within each block. Real interleaved wall-clock (5 independent runs, real Bonsai-27B shape): at depth >= 512 this is a clean 2-2.6x win over attention_f32 and a 2-2.5x win over attention_f32_gqa_split_kv, reproduced every run -- the first positive wall-clock result in the #197/#198/#199 investigation family. Correctness: 5e-3 abs-OR-rel vs CPU oracle, F32 GPU baseline, and attention_f32_gqa_split_kv itself, plus a 300-step non-compounding- drift check (CudaAttentionMmaDecodeGqaSplitTests.cs, 11 tests). ptxas: 55 registers, 0 spill, 42696B static shared; cuobjdump --dump-sass: 20 HMMA + 37 LDSM, matching the hand-computed instruction count exactly. Ships opt-in/default-OFF (DOTLLM_ATTN_MMA_DECODE_GQA_SPLIT=1) per the project's #180/#183 precedent -- this session validated synthetic-fixture parity and wall-clock, not a real generation-parity pass, and is not wired into any model's forward path yet. Full writeup in .perf-runs/issue199-v2-mma-decode-gqa-split/README.md and docs/CUDA.md. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…ridDense + real generation-parity validation Model integration: new top-priority dispatch tier in CudaQwen3HybridDenseTransformerModel, above attention_f32_gqa_split_kv, requiring numKvHeads < numHeads in addition to CudaAttentionMmaDecodeGqaSplit.CanUse's own shape checks (group=1 would degenerate to v1's failed one-warp-per-block shape). Converts Q to FP16 into new scratch, reads K/V straight from the FP16 KV cache, routes dispatch through CudaAttentionMmaDecodeGqaSplit.Run() (not CudaKernels directly) so DispatchCount correctly proves the branch fires. New generation-parity test (same harness shape as #222's CudaAttentionSplitKvGenerationParityTests) caught two real bugs before landing a trustworthy result: an initial direct-dispatch wiring bypassed the wrapper's DispatchCount tracking (fixed by routing through Run()), and the test's own depth=t step-index proxy (copied from #222's test) undercounted by one at the gate boundary (fixed to depth=t+1 -- real seqKv at step t is t+1 since the KV write precedes that step's read). Result: pre-gate bit-identical (0.0 exactly), post-gate perplexity improves slightly (-0.173%), generation diverges at the same step/depth (225/257) #222's investigation already documented for the sibling split-KV kernel -- a small borderline-margin flip consistent with that already-characterized reassociation sensitivity, not a new correctness problem. Full CUDA suite (423 tests) passes, zero regressions. Still opt-in/default-OFF (DOTLLM_ATTN_MMA_DECODE_GQA_SPLIT=1) pending a maintainer decision on default-on. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
The real generation-parity validation this project's #180/#183 precedent for new precision axes calls for before a default-on flip is now done (see the prior commit): pre-gate bit-identical, post-gate perplexity improves slightly (-0.173%), generation divergence matches an already-accepted sibling-kernel pattern rather than a new correctness problem. CudaAttentionMmaDecodeGqaSplit.Enabled now defaults true; opt back out with DOTLLM_ATTN_MMA_DECODE_GQA_SPLIT=0. Full CUDA suite (423 tests) re-verified clean with the new default active -- 0 failures. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
There was a problem hiding this comment.
Pull request overview
Adds a new CUDA decode-attention fast path that composes FP16 tensor-core (mma.sync) compute with the existing GQA-group + split-KV grid design, then wires it into the CUDA Qwen3 hybrid dense model with accompanying PTX/build updates and correctness/perf validation.
Changes:
- Introduce a new cooperative-launch decode attention kernel (
attention_flash_mma_decode_gqa_split_f16) + loader/dispatcher (CudaAttentionMmaDecodeGqaSplit,CudaKernels.LaunchAttentionMmaDecodeGqaSplit). - Integrate the kernel as the top-priority eligible decode-attention tier in
CudaQwen3HybridDenseTransformerModel, including FP16 Q scratch conversion. - Add unit + integration parity tests, and update CUDA docs / perf-run notes and PTX build plumbing.
Reviewed changes
Copilot reviewed 11 out of 11 changed files in this pull request and generated 5 comments.
Show a summary per file
| File | Description |
|---|---|
| tests/DotLLM.Tests.Unit/Cuda/CudaAttentionMmaDecodeGqaSplitTests.cs | New synthetic-fixture parity suite and microbench harness for the composed mma decode kernel. |
| tests/DotLLM.Tests.Integration/Cuda/CudaAttentionMmaDecodeGqaSplitGenerationParityTests.cs | Real-model greedy generation + decode-perplexity parity harness for the new default path. |
| src/DotLLM.Cuda/DotLLM.Cuda.csproj | Ensures the new kernel is built with required C++17 flags. |
| src/DotLLM.Cuda/CudaKernels.cs | Loads the new PTX, exposes cooperative-launch API + co-residency query for safe kvSplit. |
| src/DotLLM.Cuda/CudaAttentionMmaDecodeGqaSplit.cs | New gate/dispatch wrapper and dispatch counter for engagement verification. |
| src/DotLLM.Cuda/Architectures/CudaQwen3HybridDenseTransformerModel.cs | Wires the new kernel into the decode attention tiering and allocates FP16 Q scratch. |
| native/ptx/attention_flash_mma_decode_gqa_split.ptx | Generated sm_86 PTX for the new kernel. |
| native/kernels/attention_flash_mma_decode_gqa_split.cu | New mma-based decode attention kernel implementation + combine phase. |
| native/build_ptx.bat | Builds the new kernel at compute_86 and applies required C++17 list. |
| docs/CUDA.md | Documents the investigation results and the new kernel’s behavior/validation. |
| .perf-runs/issue199-v2-mma-decode-gqa-split/README.md | Captures perf methodology and results for the new kernel. |
Comment on lines
+35
to
+36
| [Trait("Category", "GPU")] | ||
| public sealed unsafe class CudaAttentionMmaDecodeGqaSplitTests |
Comment on lines
+98
to
+100
| Skip.IfNot(CudaDevice.IsAvailable(), "No CUDA GPU available."); | ||
|
|
||
| string? ggufPath = ResolveModelPath(); |
Comment on lines
+93
to
+95
| // Ships opt-in (DOTLLM_ATTN_MMA_DECODE_GQA_SPLIT=1, default OFF), same #180/#183 precedent as | ||
| // v1 and the plain GQA-split kernel: a new precision + reassociation axis on this | ||
| // architecture, not yet backed by a real generation-level validation pass. |
Comment on lines
+2449
to
+2452
| /// <summary> | ||
| /// OPT-IN, default-OFF composed tensor-core decode kernel (issue #199 v2), decode-only | ||
| /// (<c>seqQ==1</c> implicit). Grid = <c>(numKvHeads, kvSplit)</c> using CUDA Cooperative | ||
| /// Groups <c>grid.sync()</c> -- still exactly ONE kernel launch, same shape and combine |
Comment on lines
+932
to
+937
| shallowest depth is where fixed per-launch/grid.sync overhead is least amortized. Kept | ||
| opt-in/default-OFF (`DOTLLM_ATTN_MMA_DECODE_GQA_SPLIT=1`) per the standing #180/#183 | ||
| precedent — this session validated synthetic-fixture parity and wall-clock, not real | ||
| end-to-end generation parity (the #222-style test that caught #183's real-world precision | ||
| problem despite it passing synthetic parity), so a default-on flip should wait on that. Not | ||
| wired into any model's forward pass this session (kernel-level validation was the scope, |
2 tasks
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
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.
Summary
Closes #199. Follow-up to #197/#198 (PR #201) and #230.
A 2026-07-30 re-profile found the #197/#198 GQA-split attention kernel (
attention_f32_gqa_split_kv) genuinely reduces redundant KV traffic but only nets ~6% real wall-clock improvement, still deeply occupancy-limited. Two parallel investigations followed:__syncthreads(). Real, zero-risk cleanup — but an honest negative result on wall-clock (dead tie / noise-floor). Already merged (78082b9).mma.sync) math with the perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197/perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198 grid design. v1 (a separate, non-merged attempt) tried tensor cores alone and was 4-5x slower due to one-warp-per-block scoping (~4% occupancy). v2 packs the GQA group into themma.synctile's M=16 dimension itself — one set of tensor-core instructions computes the whole group's attention, free relative to v1's per-head cost.Results
Correctness:
CudaAttentionMmaDecodeGqaSplitTests.cs(10 tests) — CPU-oracle and F32-baseline parity at 5e-3 abs-or-rel across seqKv up to 2048, three-way agreement vs the sibling kernel, 300-step drift check.ptxas -v: 55 registers, 0 spill.cuobjdump --dump-sass: real HMMA/LDSM instructions confirmed.Performance (5 independent interleaved runs, real Bonsai-27B shape): 2-2.6x faster than the plain F32 kernel, 2-2.5x faster than the #197/#198 kernel itself, reproduced at depth ≥512.
Real generation-parity (
CudaAttentionMmaDecodeGqaSplitGenerationParityTests.cs, same harness shape as #222's split-KV test — the one that previously caught a real precision problem despite synthetic-fixture parity passing): pre-gate exactly bit-identical, post-gate perplexity improves slightly (-0.173%), greedy generation diverges at the same step/depth (225/257) #222 already documented and accepted for the sibling kernel — not a new correctness concern.Two real bugs were caught and fixed by the generation-parity test before trusting the result (see commit
88fb3b9for the full story): an initial dispatch-tracking gap (routed throughCudaKernelsdirectly instead of the wrapper'sRun(), which is what incrementsDispatchCount) and an off-by-one in the test's own depth-bucketing (copied from #222's test, which has the same unasserted approximation).Default flipped to ON (
30a48ed) now that real-generation validation is done, per this project's #180/#183 precedent for new precision axes. Opt back out withDOTLLM_ATTN_MMA_DECODE_GQA_SPLIT=0.Full CUDA unit suite (423 tests) passes with the new default active — 0 failures, 0 regressions.
Test plan
CudaAttentionMmaDecodeGqaSplitTests.cs— kernel-level correctness (10/10 pass)CudaAttentionMmaDecodeGqaSplitGenerationParityTests.cs— real Bonsai-27B generation + perplexity (pass)DotLLM.Tests.Unit.Cudasuite re-verified with new default (423 tests, 0 failed)🤖 Generated with Claude Code