[Bugfix] Fix threadIdx variable lookup by thread_tag instead of position in ThreadSync#1829
Conversation
…s and updating module state visualization
… memory This commit introduces a new test case to verify the correct behavior of sync hoisting when a non-uniform if statement is present inside a loop that utilizes shared memory. The test ensures that the synchronization occurs before the if statement, confirming the expected transformation in the module's intermediate representation.
This commit removes the optional thread variable lookup by tag and replaces it with direct indexing based on the expected position of threadIdx variables. This change addresses a bug where the last three elements of env_threads_ were incorrectly assumed to always correspond to threadIdx.x/y/z, improving the accuracy of thread variable access in the synchronization logic.
|
👋 Hi! Thank you for contributing to the TileLang project. Please remember to run We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀 |
|
Caution Review failedThe pull request is closed. 📝 WalkthroughWalkthroughThe visitor for if-then-else expressions now visits the condition expression unguarded before applying branch guards, ensuring condition buffer accesses are recorded; a test was added to validate thread-sync hoisting for non-uniform ifs inside loops using shared (thread) memory. Changes
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~20 minutes Possibly related PRs
Suggested reviewers
Poem
🚥 Pre-merge checks | ✅ 2 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (2 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing touches
🧪 Generate unit tests (beta)
Warning Review ran into problems🔥 ProblemsGit: Failed to clone repository. Please run the 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. Comment |
There was a problem hiding this comment.
Pull request overview
Fixes ThreadSync-related guard/sync placement by ensuring buffer accesses in if-conditions are visited/recorded prior to applying guards, and adds a regression test for hoisting sync before a non-uniform if inside a loop.
Changes:
- Update
ConstrVisitorto visitIfThenElsecondition expressions before creating/applying guards. - Add a CUDA test covering sync hoisting with shared memory reads in a loop/conditional.
Reviewed changes
Copilot reviewed 2 out of 2 changed files in this pull request and generated 2 comments.
| File | Description |
|---|---|
| testing/python/transform/test_tilelang_transform_thread_sync.py | Adds a regression test ensuring shared-memory sync is hoisted before a non-uniform if inside a loop. |
| src/transform/common/constr_visitor.h | Ensures condition expressions are visited unguarded so their buffer accesses are recorded before guard application. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
testing/python/transform/test_tilelang_transform_thread_sync.py
Outdated
Show resolved
Hide resolved
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
Added condition expression visiting in
constr_visitor.hto ensure buffer accesses in if-conditions are properly recorded before applying guardsSummary by CodeRabbit
Bug Fixes
Tests