[Bugfix] Fix pdl related intrin handling to avoid strict annotation codegen

[LeiWang1999]

This pull request updates the handling of CUDA synchronization and trigger calls in the PDL lowering and warp specialization code, and updates the TVM submodule. The changes simplify and clarify how these synchronization points are detected and marked in the codebase.

Refactoring of CUDA Sync/Trigger Call Detection:

• Updated MarkCudaSyncCalls in lower_pdl.cc to directly check for tl::pdl_trigger() and tl::pdl_sync() instead of matching external call names, improving robustness and maintainability.
• Simplified the logic in WarpSpecializedRoleMarker in warp_specialized_rewriter.cc to directly check for pdl_sync() and pdl_trigger() instead of string-based matching, ensuring more reliable role assignment.

Testing Improvements:

• Enhanced the test_pdl_sync test in test_tilelang_language_pdl.py to assert that "__restrict__" is not present in the generated code, in addition to checking for "cudaGridDependencySynchronize".

Dependency Update:

• Updated the TVM submodule to a newer commit, ensuring compatibility with recent upstream changes.

Summary by CodeRabbit

Release Notes

• Improvements

  • Refined handling of CUDA PDL operations in the compilation pipeline.
  • Updated external dependencies to latest stable versions.

• Tests

  • Enhanced kernel generation test coverage with additional validation.

_{✏️ Tip: You can customize this high-level summary in your review settings.}

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

Caution

Review failed

The pull request is closed.

📝 Walkthrough

Walkthrough

The changes refactor PDL operator detection in the compiler from string-based matching to direct operator comparisons using tl::pdl_trigger() and tl::pdl_sync(). Additionally, the submodule pointer is updated and a test assertion is added to verify kernel source formatting.

Changes

Cohort / File(s)	Summary
Submodule Update 3rdparty/tvm	Updated git submodule pointer from commit 001022bdb2dbb337d242eed9d208f8555b8edc98 to e47e76a2a0d565e02b6474c06f9f47e1374821f3.
PDL Operator Detection Refactoring src/transform/lower_pdl.cc, src/transform/warp_specialized_rewriter.cc	Replaced string-based detection logic for CUDA PDL calls with direct operator checks. Now uses tl::pdl_trigger() and tl::pdl_sync() operators instead of parsing function names and string arguments. Sets corresponding flags (has_trigger_launch_, has_grid_sync_, Role::kBoth) when operators are detected.
Test Enhancement testing/python/language/test_tilelang_language_pdl.py	Added assertion in test_pdl_sync to verify the generated kernel source does not contain "__restrict__".

Estimated code review effort

🎯 2 (Simple) | ⏱️ ~12 minutes

Poem

🐰 The detectives trade their magnifying glass for a map,
Strings were read, now operators snap!
Direct checks bloom where comparisons grew,
Simpler pathways, cleaner and true.
PDL's now sharp as a rabbit's new tooth! ✨

✨ Finishing touches

• 📝 Generate docstrings

---

📜 Recent review details

Configuration used: defaults

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 6e43953 and 51813c7.

📒 Files selected for processing (4)

• 3rdparty/tvm
• src/transform/lower_pdl.cc
• src/transform/warp_specialized_rewriter.cc
• testing/python/language/test_tilelang_language_pdl.py

---

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 and usage tips.}

[LeiWang1999]

In future I think we can merge those into a general attribute to enable the restrct related codegen.