Skip to content

Enable SM110 for Blackwell FMHA and warpgroup reg reconfig - #3360

Open
denverlloyd-debug wants to merge 1 commit into
NVIDIA:mainfrom
denverlloyd-debug:sm110-thor-enablement
Open

Enable SM110 for Blackwell FMHA and warpgroup reg reconfig#3360
denverlloyd-debug wants to merge 1 commit into
NVIDIA:mainfrom
denverlloyd-debug:sm110-thor-enablement

Conversation

@denverlloyd-debug

@denverlloyd-debug denverlloyd-debug commented Jul 1, 2026

Copy link
Copy Markdown

Adds SM110 (sm_110a / Thor) support that CUTLASS does not have:

  • Extend Blackwell FMHA kernel arch guards to accept SM110A/SM110F (same pattern as existing SM103 guards on main). This allows SM110 to work again with CUDA >= 13.0.
  • Add SM110 (1100) to warpgroup setmaxnreg activation in reg_reconfig.h so FMHA softmax warpgroups get their intended register budget instead of compiling reg alloc/dealloc to no-ops. This restores perf for SM110 with CUDA >= 13.0.

Cutlass 4.2 Change log has the following comment:

From CUDA 13.0, the Blackwell SM101 for Thor GPUs is renamed to SM110.
For CUDA toolkit version < 13.0, SM101 is still used for Thor GPUs.
For CUDA toolkit version >= 13.0, SM110 is used for Thor GPUs and SM101 is no longer valid.

@github-actions

Copy link
Copy Markdown

This PR has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this PR if it is no longer required. Otherwise, please respond with a comment indicating any updates. This PR will be labeled inactive-90d if there is no activity in the next 60 days.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant