Skip to content

[megatron] GLM-5.3-Flash DSA: run the fused TileLang SparseMLA kernel - #2388

Merged
erictang000 merged 3 commits into
mainfrom
glm5p3-dsa-fused-sparse-mla
Oct 2, 2026
Merged

erictang000 merged 3 commits into
mainfrom
glm5p3-dsa-fused-sparse-mla

Conversation

@avigyabb

@avigyabb avigyabb commented Oct 1, 2026 •

Copy link
Copy Markdown
Collaborator

Summary

Lets GLM-5.3-Flash's DSA layers use megatron-core's fused TileLang SparseMLA kernel instead of
the dense fallback, so DSA attention memory is linear in sequence length. Split out of #2341 so it
can be compared against the cuDNN DSA backend on its own.

Why

GLM-5.3-Flash's DSA layers are NoPE MLA (q/k width 512), and its k-pool selection appends the
query's incomplete tail pool, so the top-k width is 2048 + 3 = 2051. The fused absorbed
sparse-attention kernels are specialized for the DeepSeek-V3.2 layout (q/k width 576, top-k width
a multiple of 64). They decline, and megatron-core falls back to a dense [b, heads, sq, sk] FP32
masked softmax: ~128 GiB/GPU at 32k tokens.

Change

glm5_next/dsa.py: _pad_for_fused_absorbed_sparse_attention wraps
dsa_kernels.run_fused_absorbed_sparse_attention inside _forward_with_kpool_topk:

  • q/k zero-padded 512 → 576 channels in the RoPE slot. Zero channels add nothing to q·k,
    softmax_scale is passed explicitly, and the value is still key[..., :512].
  • top-k indices padded to a multiple of 64 with -1, which the kernels mask out.

Exact. It needs dsa_kernel_backend="tilelang". If the backend still declines, the dense fallback
runs on the unpadded tensors as before. The first call logs whether the fused kernel was used or
declined.

Comparing against cuDNN

  • k-pool path: _forward_with_kpool_topk already forces megatron-core's full fused DSA hook
    (run_fused_dsa_attention, which does its own token-level top-k) to decline, so only the
    sparse-attention step is swapped between backends.
  • The padding wrapper is backend-agnostic: it wraps whichever backend's
    run_fused_absorbed_sparse_attention is selected.
  • What I measured and didn't: I only measured TileLang. An earlier look suggested the cuDNN /
    FlashMLA route needs flash_mla installed and pads TP-local heads to 64. That's untested here.
  • Expected impact: in [sft] Enable GLM-5.3-Flash full fine-tuning in SFT (Megatron) #2341's GLM-5.3-Flash profiles, SparseMLA forward+backward was a few
    percent of a long-context step; the k-pool indexer and inter-node communication dominated.

Testing

gpu_ci/patches/megatron/test_glm5_next_fused_sparse_attention.py: padded fused attention vs the
dense reference, forward and backward.

🤖 Generated with Claude Code


Note

Medium Risk
Changes the fused DSA attention path for GLM-5.3-Flash long-context training; incorrect padding would alter attention or reintroduce dense OOM, though behavior is covered by CPU/GPU parity tests and auto-disables when upstream #7617 is present.

Overview
Adds a Megatron runtime backport (NVIDIA/Megatron-LM#7617) so GLM-5.3-Flash DSA can use the TileLang SparseMLA fused path instead of megatron-core’s dense (O(L^2)) fallback.

patch_sparse_mla_nope replaces tilelang_dsa.fused_sparse_mla_absorbed at model build time (MegatronWorker.make_megatron_module). For NoPE absorbed MLA it zero-pads q/k from 512→576 and pads top-k indices to a multiple of 64 with -1 before calling the kernel, then trims head/output padding. It no-ops when upstream already contains the fix or TileLang is unavailable.

Documentation in patches/megatron/README.md describes removal once the megatron-core pin lands. CPU tests use a fake SparseMLA to assert padding/gradients; H100 GPU tests compare the real kernel and dsa_kernels.run_fused_absorbed_sparse_attention against references (including 2051-wide k-pool shapes). H100 CI runs the GPU test file.

Reviewed by Cursor Bugbot for commit 236ca6e. Bugbot is set up for automated code reviews on this repo. Configure here.

avigyabb and others added 2 commits October 1, 2026 11:16
GLM-5.3-Flash's DSA layers are NoPE MLA (q/k width 512) with a k-pool selection that appends the
query's incomplete tail pool (top-k width 2048 + 3 = 2051). The fused absorbed sparse-attention
kernels (TileLang SparseMLA) are specialized for the DeepSeek-V3.2 layout (q/k width 576, top-k
width a multiple of 64), so they decline and megatron-core falls back to a dense
[b, heads, sq, sk] FP32 masked softmax: O(L^2) memory, ~128 GiB/GPU at 32k tokens.

_pad_for_fused_absorbed_sparse_attention wraps dsa_kernels.run_fused_absorbed_sparse_attention
for the k-pool path: q/k are zero-padded from 512 to 576 channels in the RoPE slot (zero channels
add nothing to q.k; softmax_scale is passed explicitly; the value stays key[..., :512]) and the
indices are padded to a multiple of 64 with -1, which the kernels mask. Exact. Needs
dsa_kernel_backend="tilelang"; if the backend still declines, the dense fallback runs on the
unpadded tensors as before.

Test: padded fused attention vs the dense reference, forward and backward.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The TileLang SparseMLA kernel needs ~166 KB of shared memory per block,
more than an L4 (sm89) allows, so on the L4 megatron GPU CI both cases fail
with 'Failed to set the allowed dynamic shared memory size to 169984'.
Mark the test h100 and add it to ci/gpu_ci_run_h100.sh.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@erictang000

Copy link
Copy Markdown
Collaborator

relevant megatron-lm pr: NVIDIA/Megatron-LM#7617

….3 SparseMLA wrapper

Replace the GLM-specific _pad_for_fused_absorbed_sparse_attention wrapper in glm5_next/dsa.py
with a backport of NVIDIA/Megatron-LM#7617 (approved upstream), which teaches megatron-core's
tilelang_dsa.fused_sparse_mla_absorbed the same two exact paddings: 512-wide NoPE q/k zero-padded
to 576, and top-k widths padded to a multiple of 64 with -1.

patch_sparse_mla_nope.py rebinds tilelang_dsa.fused_sparse_mla_absorbed (its only caller resolves
it as a module global) to the upstream function from the PR head, and is a no-op that warns once
the pinned megatron-core contains the fix. Applied from MegatronWorker.make_megatron_module next to
the other megatron-core patches.

Tests:
- CPU: upstream's fake-SparseMLA adapter test (padding, -1 slots, scale, unpadding, gradients),
  plus the GLM k-pool width 2051, through tilelang_dsa.run_fused_absorbed_sparse_attention.
  25 passed.
- GPU (h100 marker): upstream's real-kernel vs FP32 reference sweep, plus the GLM-5.3-Flash-shaped
  case through dsa_kernels vs megatron-core's dense fallback. 20 passed on 1x B200.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@erictang000
erictang000 marked this pull request as ready for review October 2, 2026 08:54
@erictang000
erictang000 merged commit 99824f0 into main Oct 2, 2026
8 checks passed
@erictang000
erictang000 deleted the glm5p3-dsa-fused-sparse-mla branch October 2, 2026 08:55

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request introduces a backport patch (patch_sparse_mla_nope.py) for NVIDIA/Megatron-LM#7617 to support NoPE shapes and unaligned top-k in the TileLang SparseMLA kernel, preventing a fallback to dense softmax for GLM-5.3-Flash. It also adds corresponding CPU and GPU tests and integrates the patch application into the Megatron worker. The review feedback suggests wrapping the inspect.getsource call in a try-except block to avoid potential runtime crashes in environments where the source code is unavailable.

Comment on lines +113 to +118
if _UPSTREAM_SENTINEL in inspect.getsource(tilelang_dsa.fused_sparse_mla_absorbed):
logger.warning(
"megatron-core already supports NoPE SparseMLA (NVIDIA/Megatron-LM#7617); "
"delete skyrl/backends/skyrl_train/patches/megatron/patch_sparse_mla_nope.py"
)
return False

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

Using inspect.getsource can raise an OSError or TypeError if the source code is not available (e.g., in compiled, zipped, or packaged environments where .py files are stripped). To prevent potential startup crashes, wrap the inspect.getsource call in a try-except block and default to applying the patch if inspection fails.

Suggested change
if _UPSTREAM_SENTINEL in inspect.getsource(tilelang_dsa.fused_sparse_mla_absorbed):
logger.warning(
"megatron-core already supports NoPE SparseMLA (NVIDIA/Megatron-LM#7617); "
"delete skyrl/backends/skyrl_train/patches/megatron/patch_sparse_mla_nope.py"
)
return False
try:
upstream_source = inspect.getsource(tilelang_dsa.fused_sparse_mla_absorbed)
has_upstream_fix = _UPSTREAM_SENTINEL in upstream_source
except (OSError, TypeError) as e:
logger.warning(f"Could not inspect fused_sparse_mla_absorbed source: {e}. Proceeding with patch.")
has_upstream_fix = False
if has_upstream_fix:
logger.warning(
"megatron-core already supports NoPE SparseMLA (NVIDIA/Megatron-LM#7617); "
"delete skyrl/backends/skyrl_train/patches/megatron/patch_sparse_mla_nope.py"
)
return False

@greptile-apps

greptile-apps Bot commented Oct 2, 2026 •

Copy link
Copy Markdown
Contributor

RetriggerConfidence Score: 4/5

[High risk] Patches megatron-core sparse attention kernel for GLM model.

The PR should not merge until the GLM-5.3-Flash training configuration selects the patched TileLang backend.

Findings

  1. P1 Flash launcher misses TileLang ▶
  2. P2 Upstream fix fails fixtures ▶
Diagram
%%{init: {'theme': 'neutral'}}%%
flowchart LR
  A[GLM-5.3-Flash DSA] --> B{DSA backend}
  B -->|tilelang| C[Patched SparseMLA shape adapter]
  C --> D[Fused sparse attention]
  B -->|other backend| E[Existing backend or dense fallback]
Loading

Reviews (1) · Last reviewed commit: "[megatron] Port NVIDIA/Megatron-LM#7617 ..."

# are not a multiple of 64 (GLM-5.3-Flash k-pool: 2051); otherwise DSA falls back to a
# dense O(L^2) softmax. Delete along with the patch module once the megatron-core pin
# includes NVIDIA/Megatron-LM#7617.
patch_sparse_mla_nope()

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 Flash launcher misses TileLang

This installs the patch, but the GLM-5.3-Flash training launcher never sets dsa_kernel_backend="tilelang". The new GPU test sets it explicitly. Running the provided launcher therefore does not select the patched kernel, leaving long sequences exposed to the dense attention fallback and its O(L²) memory use. Set the backend in the Flash configuration and test that configuration.

@pytest.fixture(scope="module")
def tilelang_dsa():
pytest.importorskip("tilelang")
from megatron.core.transformer.experimental_attention_variant.ops import (

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P2 Upstream fix fails fixtures

patch_sparse_mla_nope() intentionally returns False when megatron-core already includes the fix. This assertion, and the same one in the CPU fixture, will fail after that dependency update even if the kernel works correctly. Allow the no-op case and check the resulting behavior instead, so a routine upgrade does not break both test suites.

avigyabb added a commit that referenced this pull request Oct 5, 2026
…hunked sparse attention

Lets GLM-5.3-Flash's DSA layers train at long context (tested to 1M tokens per sequence).
All changes are exact; the opt-ins default off.

- 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] 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.
  SKYRL_DSA_KPOOL_SCORE_CHUNK_ELEMS sets the chunk cap (default 2 GiB of FP32 scores); 8 GiB
  was ~5% faster per step at 512k tokens (TP8 x DP8, 64 B200) with no change in peak memory.
- SKYRL_DSA_INDEXER_TP_SHARD=1: every TP rank scores the same gathered sequence with the same
  (frozen) indexer weights, so each rank selects pools for 1/TP of the query rows and the
  selections are all-gathered. The O(sq^2) scoring was ~1/3 of a 512k step at TP8.
- SKYRL_DSA_QUERY_CHUNK=<tokens>: run the fused absorbed sparse attention in checkpointed query
  chunks (each query attends only its own top-k keys, so chunks are independent). Keeps one
  chunk's padded q / kernel output alive instead of the whole sequence's, for one extra kernel
  forward in backward. 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).

Tests (GPU): test_dsa_kpool.py (chunked == one-shot selection), test_dsa_kpool_tp_shard.py (two
ranks, bitwise-identical selections), test_glm5_next_fused_sparse_attention.py (chunked == one-shot
fused kernel, fwd + bwd).

Split out of #2341.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants