Repository navigation
Conversation
…pt-in) SKYRL_DSA_QUERY_CHUNK=<tokens> runs the fused absorbed sparse attention in query chunks, each checkpointed. Each query attends only its own top-k keys, so chunks are independent (key grads sum). The checkpoint is what saves memory: without it autograd keeps every chunk's padded query (TileLang pads 8 local heads to 16 at TP8, FlashMLA to 64) and kernel output -- as much as one unchunked call. With it only one chunk's are alive at a time, for one extra kernel forward in backward (~5.6 GiB less at 512k, TP8). Wraps whichever dsa_kernel_backend is selected; on tilelang the kernel takes GLM-5.3-Flash's NoPE / 2051-wide layout through patch_sparse_mla_nope (#2388). The first call logs whether the fused kernel was used or declined. Test: test_glm5_next_fused_sparse_attention.py (chunked == one-shot fused kernel, fwd + bwd, ragged last chunk). Independent of the k-pool indexer PRs. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
There was a problem hiding this comment.
Code Review
This pull request introduces query-chunked fused sparse attention for GLM-5.3-Flash, which runs the fused absorbed sparse attention in checkpointed query chunks to reduce memory usage during training. It also adds a corresponding GPU test to verify that the chunked implementation matches the unchunked version. Feedback on the implementation highlights two issues with the checkpointing logic: the closure capture of the key tensor, which can prevent autograd from correctly tracking gradients, and a potential RuntimeError if checkpoint is called when neither input tensor requires gradients.
| def chunked(config, query, key, topk_indices, softmax_scale, v_channels, topk_length=None): | ||
| def run(q, idx, lengths): | ||
| return fused_fn(config, q, key, idx, softmax_scale, v_channels, topk_length=lengths) | ||
|
|
||
| # Each query attends only its own top-k keys, so query chunks are independent (key grads | ||
| # just sum). Checkpointing each chunk keeps only one chunk's padded q / kernel output | ||
| # alive instead of the whole sequence's -- at TP8 the kernel pads 8 local heads to 16, so | ||
| # that is ~13 GiB/GPU at 512k tokens -- for one extra kernel forward in backward. | ||
| sq = query.size(0) | ||
| if _DSA_QUERY_CHUNK and sq > _DSA_QUERY_CHUNK and topk_length is None: | ||
| outs = [] | ||
| for q0 in range(0, sq, _DSA_QUERY_CHUNK): | ||
| q1 = min(sq, q0 + _DSA_QUERY_CHUNK) | ||
| args = (query[q0:q1], topk_indices[:, q0:q1], None) | ||
| if torch.is_grad_enabled(): | ||
| o = torch.utils.checkpoint.checkpoint(run, *args, use_reentrant=False) | ||
| else: | ||
| o = run(*args) | ||
| if o is None: | ||
| return None | ||
| outs.append(o) | ||
| out = torch.cat(outs, dim=0) | ||
| else: | ||
| out = run(query, topk_indices, topk_length) |
There was a problem hiding this comment.
There are two issues with the current checkpointing implementation:\n1. Closure Capture of key: The key tensor is captured via closure inside run instead of being passed as an explicit argument to checkpoint. In PyTorch, capturing tensors that require gradients inside a checkpointed function can prevent autograd from correctly tracking and backpropagating gradients to them.\n2. Potential RuntimeError: If torch.is_grad_enabled() is True but neither query nor key requires gradients (e.g., during validation/evaluation passes where requires_grad=False but torch.no_grad() is not used), torch.utils.checkpoint.checkpoint will raise a RuntimeError: none of the inputs have requires_grad=True.\n\nTo fix both issues, pass key as an explicit argument to run and only invoke checkpoint if at least one of the tensors requires gradients.
def chunked(config, query, key, topk_indices, softmax_scale, v_channels, topk_length=None):\n def run(q, k, idx, lengths):\n return fused_fn(config, q, k, idx, softmax_scale, v_channels, topk_length=lengths)\n\n # Each query attends only its own top-k keys, so query chunks are independent (key grads\n # just sum). Checkpointing each chunk keeps only one chunk's padded q / kernel output\n # alive instead of the whole sequence's -- at TP8 the kernel pads 8 local heads to 16, so\n # that is ~13 GiB/GPU at 512k tokens -- for one extra kernel forward in backward.\n sq = query.size(0)\n if _DSA_QUERY_CHUNK and sq > _DSA_QUERY_CHUNK and topk_length is None:\n outs = []\n for q0 in range(0, sq, _DSA_QUERY_CHUNK):\n q1 = min(sq, q0 + _DSA_QUERY_CHUNK)\n args = (query[q0:q1], key, topk_indices[:, q0:q1], None)\n if torch.is_grad_enabled() and (query.requires_grad or key.requires_grad):\n o = torch.utils.checkpoint.checkpoint(run, *args, use_reentrant=False)\n else:\n o = run(*args)\n if o is None:\n return None\n outs.append(o)\n out = torch.cat(outs, dim=0)\n else:\n out = run(query, key, topk_indices, topk_length)
|
| if o is None: | ||
| return None |
There was a problem hiding this comment.
Chunk declines go unlogged When a fused backend declines on a chunk, this early return skips the new “declined” log. The caller then uses the much more memory-intensive dense fallback, but an opted-in run has no diagnostic explaining why the fused path was not used.
Note: If this suggestion doesn't match your team's coding style, reply to this and let me know. I'll remember it for next time!
| fused = glm_dsa._query_chunked_fused_absorbed_sparse_attention(dsa_kernels.run_fused_absorbed_sparse_attention) | ||
| grad_out = torch.randn(seqlen, 1, HEADS, LATENT, device=device) | ||
|
|
||
| def run(): | ||
| out = fused(cfg, q, k, idx, scale, LATENT) |
There was a problem hiding this comment.
Production wiring remains untested This test constructs and calls the wrapper directly. It would still pass if
Glm5NextDSAttention.forward stopped installing or invoking the wrapper, leaving the opt-in flag ineffective in training. An assertion through the attention forward path would cover that integration.
| return None | ||
|
|
||
| saved = (mcore_dsa.fused_qk_topk_naive, dsa_kernels.run_fused_dsa_attention, dsa_kernels.run_fused_qk_topk) | ||
| saved = ( |
There was a problem hiding this comment.
we ended up changing this structure in a previous PR. maybe we can go towards that one
…n query chunks (#2413) ## Summary First of three PRs split from the original GLM-5.3-Flash DSA long-context change: - **this PR**: chunked k-pool indexer scoring (on by default, exact) - **#2422**: TP-sharded indexer (opt-in), stacked on this one - **#2423**: query-chunked sparse attention (opt-in), independent, on main `mcore_ext/dsa_kpool.py`: `fused_qk_topk_kpool` scores and selects pools in query chunks. Verbatim #7522 materializes FP32 `[sq, b, heads, sq / kpool]` per-head scores (32 GiB/GPU at 32k), so GLM-5.3-Flash could not train past ~16k. Top-k is per query row, so chunking is exact. Only affects GLM-5.3-Flash's k-pool path. `SKYRL_DSA_KPOOL_SCORE_CHUNK_ELEMS` sets the chunk cap (default 2 GiB of FP32 scores, which is 128 queries per chunk at 512k). ### Aligned with upstream #7522 @ `b27efd8` #7522 has since added its own query chunking (`_kpool_topk_from_score_chunks`). The second commit here (`15fde54d`) adopts its structure past one chunk: scoring under `no_grad` with ReLU and head weighting in place (one FP32 score buffer per chunk instead of two or three), `k_pooled` cast to FP32 once instead of per chunk, and a preallocated `[batch, sq, budget]` selection instead of `torch.cat`. Selections are bitwise unchanged; the one-chunk path still returns full scores with autograd. What stays SkyRL-specific is the configurable cap: upstream's is a fixed 256 MiB. One selection at 512k (one TP8 rank's 65,536 query rows x 131,072 pools, 32 heads, packed varlen masking, 1 B200; median of 3, spread <0.1%; all runs select bitwise-identical pools): | implementation | 256 MiB (16 queries/chunk) | 2 GiB (128) | 8 GiB (512) | |---|---|---|---| | upstream #7522 @ `b27efd8` | 3.308 s, 0.61 GiB | 2.640 s, 2.45 GiB | 2.598 s, 8.88 GiB | | this PR, first commit | 3.480 s, 0.76 GiB | 2.652 s, 4.31 GiB | 2.593 s, 16.5 GiB | | this PR, with `15fde54d` | 3.282 s, 0.57 GiB | **2.626 s, 2.38 GiB** | 2.585 s, 8.57 GiB | (GiB = peak extra memory of the call.) With 11 DSA layers under full recompute (22 selections per step), upstream's 256 MiB cap costs ~15 s of a ~217 s step versus 2 GiB; 8 GiB saves under 1 s more over 2 GiB for 6 GiB more memory, so the 2 GiB default stays. ## Testing - [x] `gpu_ci/patches/megatron/mcore_ext/test_dsa_kpool.py` (chunked selection == one-shot): 5 passed with `15fde54d` (B200) - [x] CPU: first commit vs `15fde54d` vs one-shot, 72 chunked cases (packed / dense / additive mask, ReLU on and off, 1 to n-1 queries per chunk): bitwise-equal indices, identical one-shot scores, gradients still flow through them - [x] Chunk-size benchmark above - [x] Logprobs vs vLLM on the real 4-layer GLM-5.3-Flash checkpoint (TP2 / EP4, 4 B200). Re-run 2026-10-06 with every split PR of #2341 merged onto main (this PR's first commit; `15fde54d` selects bitwise-identical pools): Megatron vs vLLM 0.0658 (short) and 0.0548 (`kpool_beyond_topk`); vLLM pre/post sync 0.255 and 0.358. All pass. Originally measured on the pre-split branch: | row | Megatron vs vLLM (thr 0.1) | main | vLLM pre/post sync | main | |---|---|---|---|---| | short, 128 tokens | 0.0651 | 0.0623 | 0.2402 (thr 0.3) | 0.2467 | | `kpool_beyond_topk` (past `index_topk`) | 0.0558 | 0.0546 | 0.3446 (thr 0.5) | 0.3382 | Split out of #2341. 🤖 Generated with [Claude Code](https://claude.com/claude-code) <!-- CURSOR_SUMMARY --> --- > [!NOTE] > **High Risk** > K-pool indexer mistakes change which tokens are attended without raising, affecting Megatron vs vLLM parity past `dsa_indexer_topk`; chunking is intended to be exact per query row but this path is explicitly called the riskiest patch entry. > > **Overview** > **GLM-5.3-Flash k-pool DSA indexer** no longer materializes full-sequence FP32 index scores in `fused_qk_topk_kpool`. Long contexts now **score and top-k select pools in query chunks** via `_kpool_topk_in_query_chunks`, with chunk size capped by **`SKYRL_DSA_KPOOL_SCORE_CHUNK_ELEMS`** (default ~2 GiB of FP32 scores). When one chunk covers the whole sequence, behavior stays on the original `_compute_index_scores` path so full scores (and autograd) are still returned; otherwise only pool indices are kept and `index_scores` is `None`. > > Docs in the Megatron patches README describe this deliberate deviation from verbatim #7522 and how it relates to upstream’s smaller fixed chunk cap. > > GPU tests add **`test_kpool_query_chunking_is_exact`** (chunked vs one-shot indices) and **`test_kpool_selection_matches_reference`** (one-shot and chunked vs a brute-force definition, including packed varlen and ReLU on/off). > > <sup>Reviewed by [Cursor Bugbot](https://cursor.com/bugbot) for commit 1105a1b. Bugbot is set up for automated code reviews on this repo. Configure [here](https://www.cursor.com/dashboard/bugbot).</sup> <!-- /CURSOR_SUMMARY --> --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
|
this PR may not be that useful since sparse attention's saved tensors grow linearly |
Summary
SKYRL_DSA_QUERY_CHUNK=<tokens>(opt-in) runs GLM-5.3-Flash's fused absorbed sparse attention in query chunks, each checkpointed. Each query attends only its own top-k keys, so chunks are independent (key grads sum).The checkpoint is what saves memory, so chunking always implies the recompute. Without it, autograd keeps every chunk's padded query (TileLang pads 8 local heads to 16 at TP8; FlashMLA pads to 64) and kernel output, as much as one unchunked call. With it, only one chunk's are alive at a time, for one extra kernel forward in backward: ~5.6 GiB less at 512k, TP8. With gradients off (e.g. inference logprobs) it only chunks.
Wraps whichever
dsa_kernel_backendis selected; ontilelangthe kernel takes GLM-5.3-Flash's NoPE / 2051-wide layout throughpatch_sparse_mla_nope(#2388). The first call logs whether the fused kernel was used or declined.Independent of the k-pool indexer PRs (#2413, #2422). Both add an env flag in the same lines at the top of
glm5_next/dsa.py, so whichever merges second needs a one-line rebase.Testing
gpu_ci/patches/megatron/test_glm5_next_fused_sparse_attention.py(chunked == one-shot fused kernel, fwd + bwd, ragged last chunk): passed on this branch (B200)Split out of #2341.
🤖 Generated with Claude Code
Note
Medium Risk
Changes GLM-5.3-Flash training/inference attention execution and autograd behavior behind an env flag, but defaults off and is covered by a GPU numerical parity test.
Overview
Adds an opt-in memory path for GLM-5.3-Flash k-pool DSA: set
SKYRL_DSA_QUERY_CHUNK=<tokens>to run fused absorbed sparse attention over the query axis in chunks, with gradient checkpointing per chunk so only one chunk’s padded queries and kernel outputs stay live during backward (trades one extra kernel forward for large-sequence memory).The wrapper
_query_chunked_fused_absorbed_sparse_attentionis installed duringGlm5NextDSAttention._forward_with_kpool_topkby temporarily replacingdsa_kernels.run_fused_absorbed_sparse_attention; it is backend-agnostic and logs once whether the fused path was used or declined. Chunking is skipped whentopk_lengthis set or the sequence is shorter than the chunk size; without autograd it chunks without checkpoint.Documentation in the Megatron patches README is updated, and a new GPU CI test asserts chunked vs one-shot TileLang fused sparse attention match in forward and backward (with relaxed tolerance on key grads due to summation order).
Reviewed by Cursor Bugbot for commit 5ebaf98. Bugbot is set up for automated code reviews on this repo. Configure here.