Skip to content

[Perf] Add CuTe DSL DSA indexer - #1991

Open
jayhenry wants to merge 2 commits into
InternLM:mainfrom
jayhenry:feat/cute-dsl-indexer
Open

[Perf] Add CuTe DSL DSA indexer#1991
jayhenry wants to merge 2 commits into
InternLM:mainfrom
jayhenry:feat/cute-dsl-indexer

Conversation

@jayhenry

@jayhenry jayhenry commented Jul 31, 2026

Copy link
Copy Markdown
Collaborator

Summary

  • Add an SM90 CuTe DSL fused DSA indexer that computes BF16 scores and exact radix top-k IDs without materializing global FP32 logits.
  • Expose it through the optional cute dependency extra and indexer_backend="cute_dsl".
  • Decouple indexer_backend (default: tilelang) from sparse_mla_backend, while preserving the upstream int32 top-k contract.
  • Cover eager/compiled execution, packed causal boundaries, short ranges, production top-k shapes, nonzero packed offsets, repeated K-stage reuse, online compaction, and non-default CUDA streams.

Integration fixes

  • Launch the compiled kernel on the current PyTorch CUDA stream.
  • Add the missing consumer-to-next-writer barrier for the shared partial-score tile; Compute Sanitizer racecheck now reports zero hazards.
  • Initialize the CuTe CUDA context on the input tensor device instead of defaulting every rank to GPU 0. This fixes the 32K SP2+EP2+micro2 DeepEP regression that previously failed at the post-backward NCCL all-reduce.

Kernel benchmark

H200 development measurements with BF16, Q=32768, 32 index heads, dim=128, and topk=2048:

K Existing TileLang CuTe DSL Speedup
32,768 28.606 ms 16.079 ms 1.78x
65,536 70.505 ms 45.442 ms 1.55x
262,144 320.990 ms 195.013 ms 1.65x

Peak incremental memory is about 0.504 GiB for all three shapes, versus 5.317/9.317/33.317 GiB for the existing TileLang path.

SFT regression

All runs use 8xH200, GLM-5.2-30B-MTP-new, FP8, model/torch compile, cuDNN DSA sparse MLA, 28 steps, and all-rank statistics from steps 13-28.

Case TileLang mean step CuTe DSL mean step Effective / sequence TGS change
16K micro1 Muon 1.9390 s 1.9153 s +1.24% / +1.24%
16K micro1 AdamW 1.4789 s 1.4642 s +1.00% / +1.00%
32K SP2+EP2+micro2 Muon 17.8147 s 17.2909 s +3.03% / +3.03%

All three runs completed 28/28 steps without OOM, NaN, CUDA, NCCL, DeepEP, or traceback failures. The TileLang comparison uses the 2026-07-30 7fb0eed4 snapshot, so it includes cross-snapshot noise and is not a causal attribution of the full delta to this PR.

Tests

  • Targeted pre-commit: merge-conflict, codespell, docformatter, pyupgrade, Ruff, Ruff format, mypy, and pydantic-extra-check passed.
  • pytest -q tests/ops/test_cute_dsl_indexer_topk.py (19 passed)
  • pytest -q tests/module/attention/test_dsa_mla.py (10 passed)
  • Compute Sanitizer on the production topk=2048 specialization: memcheck 0 errors, synccheck 0 errors, racecheck 0 hazards.

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.

1 participant