Skip to content

fix: Use is_family_of() for SM90 arch guard in warpgroup MmaOp#3084

Open
blake-snc wants to merge 1 commit intoNVIDIA:mainfrom
blake-snc:fix/sm90-warpgroup-family-arch-guard
Open

fix: Use is_family_of() for SM90 arch guard in warpgroup MmaOp#3084
blake-snc wants to merge 1 commit intoNVIDIA:mainfrom
blake-snc:fix/sm90-warpgroup-family-arch-guard

Conversation

@blake-snc
Copy link
Copy Markdown

Summary

  • Replace hardcoded arch == Arch.sm_90a with arch.is_family_of(Arch.sm_90a) in warpgroup/mma.py's MmaOp.__post_init__

Problem

MmaOp in warpgroup/mma.py uses if not arch == Arch.sm_90a to guard its arch check. While functionally correct today (sm_90a is the only Hopper "a"-suffix arch), this is inconsistent with the is_family_of() pattern used elsewhere in CuTe DSL (see #3082 for the SM12x equivalent fix).

Fix

# Before
if not arch == Arch.sm_90a:

# After
if not arch.is_family_of(Arch.sm_90a):

Validation

$ python3 -c "from cutlass.base_dsl.arch import Arch; print(Arch.sm_90a.is_family_of(Arch.sm_90a))"
True

$ python3 -c "from cutlass.base_dsl.arch import Arch; print(Arch.sm_120a.is_family_of(Arch.sm_90a))"
False

$ python3 -c "from cutlass.base_dsl.arch import Arch; print(Arch.sm_100a.is_family_of(Arch.sm_90a))"
False

sm_90a accepted, non-Hopper arches correctly rejected.

Related

Contributed by Second Nature Computing (https://joinsecondnature.com)

Replace hardcoded `arch == Arch.sm_90a` with
`arch.is_family_of(Arch.sm_90a)` in warpgroup/mma.py's MmaOp for
consistency with the warp-level MMA fix in NVIDIA#3082.

While functionally equivalent today (sm_90a is the only Hopper
"a"-suffix arch in practice), this makes the arch guard consistent
with the is_family_of() pattern and future-proofs against potential
Hopper variants.

Validated:
- is_family_of(Arch.sm_90a) returns True for sm_90a
- is_family_of(Arch.sm_90a) returns False for sm_120a, sm_100a

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
@blake-snc
Copy link
Copy Markdown
Author

Ping for review — same is_family_of() pattern as #3082, applied to the SM90 warpgroup MmaOp guard.

@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