fix: refine TCGEN05 architecture guards

[Rachmanino]

Summary

Follow-up to #2781.

• Replace the CuTe TCGEN05 config dependency with TL_CUDA_ARCH_TCGEN05_ENABLED, derived from CUDA architecture feature and family macros for SM100/101/103/110.
• Move dependent architecture checks into the actual TCGEN05 MMA and TMEM load/store entry points.
• Remove the codegen-emitted require_tcgen05* calls and the redundant tcgen_05.h requirement for MMA/load/store paths.

Guarded entry points

• MMA: tcgen05mma_ss, tcgen05mma_ts, tcgen05mma_ws_ss, tcgen05mma_blockscaled_ss
• TMEM load: tcgen05_ld_32dp32bNx, tcgen05_ld_32dp64bNx, tcgen05_ld_32dp128bNx, tcgen05_ld_32dp256bNx
• TMEM store: tcgen05_st_32dp32bNx, tcgen05_st_32dp64bNx, tcgen05_st_32dp128bNx, tcgen05_st_32dp256bNx
• Existing common helpers switched to the native CUDA capability: tmem_allocate, tmem_deallocate, tcgen05_before_thread_sync, tcgen05_after_thread_sync, tcgen05_mma_arrive, tcgen05_cp

All diagnostics name the concrete device function rather than a require_* helper or wildcard name.

Validation

• ninja -C build
• pytest -q testing/python/issue/test_tilelang_issue_2602.py (17 passed)
• sm_100f TCGEN05 compile check

Summary

• Replaced CuTe TCGEN05 configuration checks with native CUDA architecture feature/family detection for SM100/101/103/110 via a new TL_CUDA_ARCH_TCGEN05_ENABLED guard.
• Moved TCGEN05 architecture guards to the TCGEN05 MMA and TMEM load/store entry points by introducing tl::tcgen05mma_* public wrappers that compile on supported targets and otherwise fail with dependent static_assert diagnostics.
• Removed codegen-emitted require_tcgen05* calls and eliminated redundant TCGEN05 header requirements for MMA/load/store paths.
• Updated TCGEN05 helper gating in common headers and device helpers to use TL_CUDA_ARCH_TCGEN05_ENABLED instead of prior CuTe/CUTE_ARCH_* style configuration.
• Updated SM100A TMEM 32dp load/store template entry points to be conditionally compiled on TL_CUDA_ARCH_TCGEN05_ENABLED, with unsupported-target fallbacks guarded by dependent static_assert.
• Refreshed/expanded CUDA intrinsic architecture-guard testing:
  • Added testing/python/cuda/test_cuda_intrinsic_arch_guards.py to validate rejection diagnostics for unsupported architectures and ensure generated kernel sources do not contain "require_tcgen05".
  • Updated testing/python/issue/test_tilelang_issue_2602.py expectations for the new tcgen05mma/tmeme load/store diagnostic strings.
  • Removed the older testing/python/issue/test_tilelang_issue_2504.py module.

Validation included a successful build, 17 passing issue tests, and an sm_100f TCGEN05 compile check.

C++ style / lint notes

• This PR changes C++ headers and device-function APIs but does not appear to modify rules documented in docs/developer_guide/cpp_style.md.
• The C++ API Style Audit (warning only) is relevant since new C++ templates/wrappers and header guards were added; any TLCPP003/TLCPP004 findings should be treated as advisory unless they indicate a concrete API/FFI/maintainability risk.
• No correctness, build, or test issues are indicated by the provided validation.

[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]

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: 39bd373e-0ef8-4294-898f-94498324939b

📥 Commits

Reviewing files that changed from the base of the PR and between 5a16c6c and 4a46462.

📒 Files selected for processing (3)

• testing/python/cuda/test_cuda_intrinsic_arch_guards.py
• testing/python/issue/test_tilelang_issue_2504.py
• testing/python/language/test_tilelang_cast_rounding.py

💤 Files with no reviewable changes (1)

• testing/python/issue/test_tilelang_issue_2504.py

---

📝 Walkthrough

Walkthrough

TCGEN05 CUDA support now uses architecture-based compile-time guards and dependent assertions instead of emitted require_tcgen05* helpers. MMA wrappers, memory operations, code generation, and architecture tests were updated accordingly.

Changes

TCGEN05 architecture gating

Layer / File(s)	Summary
Architecture guards and guarded implementations src/tl_templates/cuda/common.h, src/tl_templates/cuda/tcgen_05.h, src/tl_templates/cuda/copy_sm100.h	Defines TL_CUDA_ARCH_TCGEN05_ENABLED and applies it to TCGEN05 memory, synchronization, copy, and fence implementations with unsupported-target assertions.
Architecture-gated MMA wrappers src/tl_templates/cuda/instruction/tcgen05mma.h	Adds public MMA wrappers that dispatch to internal implementations on supported architectures and assert otherwise.
Code generation without require helpers src/cuda/codegen/codegen_cuda.cc	Removes emitted require_tcgen05* calls, selects the MMA instruction header, and continues emitting TCGEN05 operations.
Architecture guard validation testing/python/cuda/test_cuda_intrinsic_arch_guards.py, testing/python/issue/test_tilelang_issue_2602.py, testing/python/language/test_tilelang_cast_rounding.py	Adds broad supported and unsupported architecture coverage, updates helper expectations, and removes obsolete device-compilation checks.

Estimated code review effort: 4 (Complex) | ~45 minutes

Possibly related PRs

• tile-ai/tilelang#2743: Overlaps in TCGEN05 TMEM architecture guards and dependent static assertions.
• tile-ai/tilelang#2774: Relies on the TCGEN05 MMA and TMEM code-generation paths updated here.
• tile-ai/tilelang#2781: Modifies the same TCGEN05 code-generation paths and require_tcgen05* handling.

Suggested reviewers: leiwang1999

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title is concise and accurately summarizes the main change: refining TCGEN05 architecture guards.
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.

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