[GG] build: pin CUTLASS DSL 4.5.3 for SM120 W4A16

[voipmonitor]

Summary

• pin the CUDA CuTe DSL dependency to 4.5.3 on the Gilded Gnosis SM120 stack
• preserve a source-level package requirement that matches the release image
• document the concrete B12X W4A16 code-generation regression seen with 4.6.0

Why

CUTLASS DSL 4.6.0 regresses the B12X W4A16 fused-MoE prefill signature on
SM120. With the same vLLM source, B12X source, model, launcher, and benchmark:

Runtime DSL	A16 decode CC1	A16 prefill 64k median
4.6.0	87.36 tok/s	5,052 tok/s
4.5.3	87.18 tok/s	5,909 tok/s

The generated kernel changes from REG:242 STACK:0 under 4.5.3 to
REG:255 STACK:256 under 4.6.0. A rank-0 Torch trace attributes the prefill
delta to W4A16FusedMoeKernel; attention and indexer timings do not show the
same regression. Decode remains within measurement noise.

The 4.5.3 package is API-compatible with this branch and restores the previous
4.5.2 performance. The Docker build separately guarantees that the CUDA 13
wheel is installed after the overlapping base wheel.

Keeping this pin in requirements/cuda.txt, rather than overriding only the
container, is necessary because the generated vLLM wheel otherwise declares
an exact 4.6.0 requirement and fails pip check with the validated runtime.

Validation

• package availability and import on Python 3.12 / CUDA 13
• cutlass.cute.nvgpu.warp.MmaMXF8Op present
• TP8/DCP1/MTP0 Luke GLM-5.2 NVFP4 A16 original:
  • decode 87.18 tok/s
  • 64k prefill 5912 / 5909 / 5906, median 5909 tok/s
• TP8/DCP1/MTP0 A16 online MXFP8:
  • decode 92.18 tok/s
  • 64k prefill 5915 / 5909 / 5907, median 5909 tok/s

Summary by CodeRabbit

• Bug Fixes
  • Improved CUDA compatibility by updating the CUTLASS DSL dependency to a version that avoids register and stack spilling in supported prefill workloads.
  • Restored spill-free code generation for SM120 GPU kernels to improve performance consistency.

[coderabbitai]

📝 Walkthrough

Walkthrough

The CUDA requirements pin nvidia-cutlass-dsl[cu13] to version 4.5.3 and document spilling behavior observed with version 4.6.0.

Changes

CUDA dependency update

Layer / File(s)	Summary
Pin CUTLASS DSL to 4.5.3 requirements/cuda.txt	Changes the nvidia-cutlass-dsl[cu13] pin from 4.6.0 to 4.5.3 and documents the related register/stack spilling behavior.

Estimated code review effort: 1 (Trivial) | ~2 minutes

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title is concise and accurately describes the main change: pinning CUTLASS DSL 4.5.3 for the SM120 W4A16 build.

✨ Finishing Touches

📝 Generate docstrings

• Create stacked PR
• Commit on current branch

🧪 Generate unit tests (beta)

• Create PR with unit tests
• Commit unit tests in branch fix/gg-rebased-cutlass453-sm120-20260718

---

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands.}

[coderabbitai]

Actionable comments posted: 1

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@requirements/cuda.txt`:
- Around line 27-29: Update the CUDA 12 dependency rewrite in the build script
around its nvidia-cutlass-dsl matching logic so it recognizes the pinned ==4.5.3
requirement and removes the [cu13] extra consistently with the Dockerfile and
setup.py paths. Match the package independently of the version operator rather
than reverting the exact pin.

🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

• Push a commit to this branch (recommended)
• Create a new PR with the fixes

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Organization UI

Review profile: CHILL

Plan: Pro Plus

Run ID: 5e7626ea-f46c-40b7-86a2-088e924cf00f

📥 Commits

Reviewing files that changed from the base of the PR and between 533b037 and e53b92f.

📒 Files selected for processing (1)

• requirements/cuda.txt

[coderabbitai]

🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Update the CUDA 12 CI selector rewrite for the exact pin.

This change from a >= requirement to ==4.5.3 means the rewrite in .github/workflows/scripts/build.sh:11-16, which only matches nvidia-cutlass-dsl[cu13]>=, no longer strips [cu13]. CUDA 12 CI therefore resolves this differently from docker/Dockerfile:245-267 and setup.py:1080-1103.

Make the CI rewrite match the package name independently of the version operator, or retain a requirement syntax compatible with that script.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@requirements/cuda.txt` around lines 27 - 29, Update the CUDA 12 dependency
rewrite in the build script around its nvidia-cutlass-dsl matching logic so it
recognizes the pinned ==4.5.3 requirement and removes the [cu13] extra
consistently with the Dockerfile and setup.py paths. Match the package
independently of the version operator rather than reverting the exact pin.