Skip to content

Document setmaxregister launch-bound requirement - #3618

Open
XFDG wants to merge 1 commit into
NVIDIA:mainfrom
XFDG:docs/setmaxregister-launch-bound-3406
Open

Document setmaxregister launch-bound requirement#3618
XFDG wants to merge 1 commit into
NVIDIA:mainfrom
XFDG:docs/setmaxregister-launch-bound-3406

Conversation

@XFDG

@XFDG XFDG commented Sep 13, 2026

Copy link
Copy Markdown

Summary

  • document the register-count and warpgroup requirements of the public setmaxregister APIs
  • warn that ptxas can ignore setmaxnreg with C7508 when the entry register count is unknown
  • show how min_blocks_per_mp establishes a launch bound, while calling out its occupancy and spilling tradeoffs

Testing

  • B200 / SM100a, CUDA 13.1 ptxas: without minnctapersm, C7508 is emitted and SASS contains 0 USETMAXREG instructions
  • B200 / SM100a, CUDA 13.1 ptxas: with minnctapersm 1, C7508 disappears and SASS contains 2 USETMAXREG instructions
  • ruff check --ignore E731 on both changed files (the ignored E731 findings are pre-existing in nvvm_wrappers.py)
  • git diff --check

Contribution disclosure

This contribution was prepared with AI-assisted tooling. I reviewed the changes and validated the documented behavior on B200.

Fixes #3406

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.

[CuTe DSL] setmaxregister_* can be silently dropped by ptxas (C7508) — nothing surfaces to Python; min_blocks_per_mp=1 prevents it

1 participant