[BugFix][Carver] Parse lettered SM arch strings in check_sm_version

[adityasingh2400]

check_sm_version gated on str.isdigit(), so any arch string carrying the trailing feature-set letter that nvcc and CUTLASS use for Hopper and newer, sm_90a, sm_100a, sm_103a, fell through to the -1 sentinel instead of parsing to 90, 100 or 103.

Those lettered strings are the real value of target.attrs["arch"], not a hypothetical one. tilelang/contrib/nvcc.py already does .rstrip("af") on that same attribute, tilelang/contrib/nvrtc.py documents "90a" as a valid arch, tilelang/contrib/cutedsl/cpasync.py validates against ["sm_90", "sm_90a", "sm_100a"], and tilelang/carver/roller/policy/tensorcore.py compares compute_capability against "sm_90a" directly. One path handled the suffix and this one did not.

The -1 then silently corrupted capability dispatch. CUDA.sm_version is set from check_sm_version and feeds is_volta_arch, is_ampere_arch, is_ada_arch, is_hopper_arch and has_mma_support, and matmul_analysis compares the same value against 70, 80 and 90. With -1 a real Hopper sm_90a target failed every one of those checks, so it was treated as pre-sm_70 and quietly lost its arch-specific pipeline, block-reduce and MMA dispatch. No error was raised, which is why no test caught it.

Fix

Parse the leading digits instead of requiring an all-digit string, so sm_90a maps to 90 and sm_100a maps to 100. -1 is kept only for genuinely non-CUDA input such as a HIP gfx942 target, and the bare numeric form that the old code already accepted stays accepted.

tilelang/carver/matmul_analysis.py carried a second copy of the identical function body inside get_tensorized_func_and_tags, so fixing one place would have left the other wrong. It now imports the shared helper from tilelang.carver.arch.cuda. That module was already an import dependency of matmul_analysis, so no new import edge is added.

Testing

New CPU regression test at testing/python/carver/test_tilelang_carver_sm_version.py. No GPU is needed, the defect and the fix are pure string parsing.

Verified fail-before and pass-after against the base ref rather than a stash:

git checkout upstream/main -- tilelang/carver/arch/cuda.py tilelang/carver/matmul_analysis.py
python -m pytest testing/python/carver/test_tilelang_carver_sm_version.py -q
# 3 failed, 2 passed
#   test_check_sm_version_lettered_arch
#   test_lettered_arch_keeps_capability_dispatch
#   test_matmul_analysis_shares_one_parser

With the fix applied, the same file reports 5 passed, and testing/python/carver/test_tilelang_carver_hint.py still passes alongside it.

The test covers four things: the numeric arches that already worked, the lettered arches that did not, the non-CUDA strings that must stay at -1, and the downstream consequence, that a sm_90a arch now satisfies is_hopper_arch and has_mma_support and none of the older-arch predicates. The last test pins the deduplication so a future re-copy of the parser fails.

Formatting checked with the pinned ruff==0.14.14 from requirements-lint.txt: ruff format --check reports already formatted and ruff check passes on all three files.

Fixes #2852

Summary

• Updated check_sm_version to parse numeric and lettered CUDA architectures, including sm_90a, sm_100a, and sm_103a.
• Anchored the optional sm_ prefix to reject malformed inputs such as sm_sm_90.
• Preserved -1 for non-CUDA and unparseable inputs.
• Removed the duplicate parser in matmul_analysis.py.
• Added CPU-only regression tests for parsing, capability dispatch, invalid inputs, and parser deduplication.
• Formatting and lint checks pass.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: 071f0a60-bccb-440f-bfdf-fb9d82b6c168

📥 Commits

Reviewing files that changed from the base of the PR and between 1a3d4c0 and 6e51257.

📒 Files selected for processing (2)

• testing/python/carver/test_tilelang_carver_sm_version.py
• tilelang/carver/arch/cuda.py

🚧 Files skipped from review as they are similar to previous changes (2)

• tilelang/carver/arch/cuda.py
• testing/python/carver/test_tilelang_carver_sm_version.py

---

📝 Walkthrough

Walkthrough

The shared CUDA SM parser now accepts letter-suffixed architectures and returns their numeric versions. matmul_analysis uses this parser. New CPU-only tests validate parsing, invalid inputs, capability predicates, and parser reuse.

Changes

CUDA SM parsing

Layer / File(s)	Summary
Shared SM parser implementation tilelang/carver/arch/cuda.py, tilelang/carver/matmul_analysis.py	check_sm_version now parses numeric capabilities from CUDA architecture strings, including trailing letters. matmul_analysis imports the shared implementation and removes its duplicate.
Parser and capability regression coverage testing/python/carver/test_tilelang_carver_sm_version.py	Tests cover numeric and letter-suffixed architectures, invalid inputs, CPU-only capability predicates, shared function identity, and test execution.

Estimated code review effort: 3 (Moderate) | ~20 minutes

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title clearly identifies the bug fix for parsing lettered SM architecture strings in check_sm_version.
Linked Issues check	✅ Passed	The changes parse lettered SM architectures, preserve invalid-input handling, deduplicate the parser, and add regression coverage required by issue #2852.
Out of Scope Changes check	✅ Passed	All changes support issue #2852 by fixing parsing, sharing the implementation, and adding focused regression tests.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

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.}

[github-actions]

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run pre-commit run --all-files in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀

[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 `@tilelang/carver/arch/cuda.py`:
- Around line 21-23: Update the CUDA arch parsing logic in the function that
strips the sm prefix and matches _SM_VERSION_PATTERN so it removes only a single
leading "sm_" from arch instead of replacing every occurrence. Keep the existing
int(match.group(1)) success path unchanged, and ensure malformed values like
repeated-prefix inputs still fail the pattern match and return -1.

🪄 Autofix

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: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: cbc05763-8598-46a0-aa11-78dad8fd939f

📥 Commits

Reviewing files that changed from the base of the PR and between 1d155f4 and 1a3d4c0.

📒 Files selected for processing (3)

• testing/python/carver/test_tilelang_carver_sm_version.py
• tilelang/carver/arch/cuda.py
• tilelang/carver/matmul_analysis.py

[SiriusNEO]

LGTM

[adityasingh2400]

The ROCm job went red here, and I do not think it is this change.

It is failing on 4 of the 5 other open PRs right now as well (#2905, #2898, #2895, #2892), so it looks repo-wide rather than branch specific.

On the change itself, the only behavior differences are the intended ones. Comparing old and new across arch strings:

arch	before	after
sm_90a	-1	90
sm_100a	-1	100
sm_sm_90	90	-1
gfx942	-1	-1
gfx1100	-1	-1
sm_90, sm_80, 80, cuda, sm_	unchanged	unchanged

HIP targets parse identically, and the matmul_analysis.py hunk only removes a duplicated local copy of the same function in favour of importing it, behind an is_cuda_tensorcore_target guard that already requires target.kind.name == "cuda". So there is no path from this diff to a ROCm test.

The CUDA and Metal jobs are still queued, and those are the ones that actually exercise it.

[SiriusNEO]

The ROCm CI failure is not related to this PR, let me handle it

[adityasingh2400]

Update now that the queue cleared. CUDA-auto passed in 14m48s and Metal passed in 24m21s, along with Quick Lint. Those are the jobs that actually exercise check_sm_version.

ROCm was re-run and failed again, so it is not flaky. I want to be more careful than I was earlier though. I said it was failing on 4 of 5 other open PRs, and the current picture is 5 of 7 failing with 2 passing, so widespread but not universal:

PR	ROCm	completed
#2887	pass	08-06 06:49
#2885	pass	08-06 06:55
#2892	fail	08-05 19:03
#2895	fail	08-06 05:02
#2905	fail	08-07 02:00
#2898	fail	08-07 06:08

The failures straddle the two passes in time, so I cannot pin it to a single breakage point, and the self-hosted job logs are not readable from outside so I cannot see the actual error.

What I can still stand behind is that this diff has no path to a ROCm test. HIP arch strings parse identically before and after, and the matmul_analysis.py hunk only removes a duplicated local function behind a guard that already requires target.kind.name == "cuda". Happy to dig further if someone can share the failing output.

[adityasingh2400]

I was wrong that the ROCm logs are unreadable, they are available through gh api .../actions/jobs/<id>/logs. So here is the actual failure rather than more inference.

ROCm is 2 failed, 1818 passed, 1227 skipped:

FAILED python/jit/test_tilelang_jit_diagnostics.py::test_cuda_compile_callback_uses_fatbin_for_multiple_target_code - KeyError: 'target_format'
FAILED python/issue/test_tilelang_issue_2682.py::test_vectorized_select - RuntimeError: #include <hip/hip_runtime.h>

Both pass on the CUDA job of the same run, on this same commit:

test	CUDA	ROCm
test_cuda_compile_callback_uses_fatbin_for_multiple_target_code	PASSED	FAILED
test_vectorized_select	PASSED	FAILED

So they are runner dependent, not commit dependent. test_vectorized_select fails on a HIP include, which lines up with the agent/fix-hip-vector-select-calls work already in flight. The fatbin one calls tilelang_callback_cuda_compile and never reaches nvcc.compile_cuda, hence the empty captured dict, which is what you would expect where nvcc is not the active toolchain.

I did check the one thing that could plausibly have been mine. That test builds a target with arch: "sm_100f", and my regex does change check_sm_version("sm_100f") from -1 to 100 where the old isdigit check rejected it. But that function lives in tilelang/carver, the callback under test is in tilelang/engine/lower.py, and the test passes on CUDA where that path is actually exercised. If you would rather the regex not accept f suffixed arches, say so and I will narrow it.