[FEA] Hopper CuTe DSL FMHA example: variable-length (cu_seqlens) inputs
Maintainers usually reply within 4 days
Nobody has claimed this yet.
Assessment
- Difficulty
- 5/5
- Estimated time
- Over a week
- Newbie friendliness
- 35/100
- Issue type
- Feature
- Clarity
- Mostly clear
- Activity status
- Active
- Domain
- machine-learning, performance
Research direction
Start with examples/python/CuTeDSL/cute/hopper/kernel/attention/fmha.py and compare its interface with cute/blackwell/kernel/attention/fmha/fmha.py. The requested work spans variable-length inputs, per-sequence tile scheduling, tail handling, and coverage for empty and length-1 sequences, causal and sliding-window modes, and GQA ratios 1/4/8. Done means those cases pass and the PR includes a TFLOPS comparison on GH200; the issue asks whether the scope is acceptable and notes it may depend on #3399.
Written by the indexing model from the issue text.
Description
Is your feature request related to a problem?
examples/python/CuTeDSL/cute/hopper/kernel/attention/fmha.py runs only fixed-length batches. The Blackwell DSL FMHA (cute/blackwell/kernel/attention/fmha/fmha.py) takes cum_seqlen_q / cum_seqlen_k, so the Hopper example can't serve packed/ragged batches, which most training and serving stacks use.
(GQA/MQA already works in the Hopper example's run(). Only the docstring still says Q and K head counts must match. #3770 fixes the docstring and adds GQA 8:2 and MQA cases, which pass the reference check on GH200.)
Describe the solution you'd like
Bring the Blackwell example's varlen interface to the Hopper example:
cum_seqlen_q/cum_seqlen_k;- a tile scheduler that issues work per sequence length;
- tile tails handled by TMA out-of-bounds zero fill plus the existing mask.
Tests would cover empty and length-1 sequences, causal and sliding window, and GQA ratios 1/4/8. The PR would include a TFLOPS comparison on GH200.
Additional context
Before I write this: would you take a change this size to the Hopper example, or would you prefer the Hopper example stay fixed-length? I'd build on #3399 (max-reduction fix) if it's still open.
- Dominant language
- C++
- Stars
- 10.7k
- Forks
- 2.2k
- Avg merge
- 3d 20h
- Merged PRs (30d)
- 4
Getting set up
This project ships no dev container, Dockerfile or contributing guide, so setting up is up to you: start from its README, and see our first-contribution guide for the general steps.
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
More from NVIDIA/cutlass
-
[Bug][CuTeDSL][Fuzz] `utils.gemm.sm100.epilogue` releases the accumulator before `tcgen05.ld` completes: no `tcgen05.wait::ld` before the arrive that lets the next MMA overwrite TMEMPossibly taken @0z5a claimed this 1 day ago. Open
Difficulty 2/5 1-3 hours Newbie friendliness 78/100
Maintainers usually reply within 4 days
-
? - Needs Triage question
Difficulty 2/5 1-3 hours Newbie friendliness 70/100
NVIDIA/cutlass#3796 · 1 comment · 1 reaction ·
Maintainers usually reply within 4 days
-
[DOC] IKET quick start omits the GEMM GPU requirement and references a missing examplePossibly taken @0z5a claimed this 11 days ago. Open
Difficulty 1/5 Under an hour Newbie friendliness 82/100
NVIDIA/cutlass#3761 · 1 comment ·
Maintainers usually reply within 4 days
-
[BUG] CuTe DSL from_dlpack rejects float8_e4m3fn tensors (dtype_map lists Float8E5M2 only)Possibly taken @Ammar-Alnagar claimed this 21 days ago. Open
Difficulty 2/5 1-3 hours Newbie friendliness 74/100
NVIDIA/cutlass#3716 · 1 comment ·
Maintainers usually reply within 4 days
-
EllGemm ColumnMajor-output workspace sizing uses the unswapped problem shape and undersizes split-K semaphoresPossibly taken @XFDG claimed this 49 days ago. OpenCUTLASS C++
Difficulty 2/5 1-3 hours Newbie friendliness 88/100
NVIDIA/cutlass#3539 · 2 comments ·
Maintainers usually reply within 4 days
Similar issues
-
Difficulty 1/5 Under an hour Newbie friendliness 74/100
scp-fs2open/fs2open.github.com#7984 ·
Maintainers usually reply within 1 day
-
bug
Difficulty 2/5 1-3 hours Newbie friendliness 72/100
jellyfin/jellyfin-desktop#1290 ·
Maintainers usually reply within 1 day
-
bug
Difficulty 2/5 1-3 hours Newbie friendliness 65/100
nasa/fprime#6301 · 1 comment ·
Maintainers usually reply within 1 day
-
c++ first-contributor first-contributors good first issue help wanted medium up-for-grabs
Difficulty 2/5 1-3 hours Newbie friendliness 85/100
fork-commit-merge/fork-commit-merge#8531 ·
Maintainers usually reply within 1 day
-
Difficulty 2/5 1-3 hours Newbie friendliness 66/100
godotengine/godot#125035 ·
Maintainers usually reply within 1 day