Sync branch sync_0203: Features, Enhancements, and Bug Fixes

[LeiWang1999]

Summary

This PR synchronizes the sync_0203 branch with the main branch, including numerous features, enhancements, and bug fixes across 209 files (+12,176/-2,866 lines).

Features

• Reimplement Threadsync with ConstrVisitor ([Feature] Reimplement Threadsync with ConstrVisitor #1631)
• Support cp.reduce.async.bulk.tensor ([Feature] Support cp.reduce.async.bulk.tensor #1667)
• Add Tilelang AutoDD to reduce buggy programs ([AutoDD] Add Tilelang AutoDD to Reduce Buggy Program #1639)
• Add EagerJIT support for parameter-only kernel compilation ([EagerJIT] Add Support for Parameter Only Kernel Compilation #1664)
• Add PDL (Parallel Domain Language) support ([Feat] PDL Support #1494)
• Add hoist_broadcast_values pass ([Feature] Add hoist_broadcast_values pass #1606)
• Add T.sync_warp & T.shfl_sync intrinsics ([Feature] add T.sync_warp & T.shfl_sync; change extern pdl into intrin #1614)
• Add strong checker to detect data racing in T.Parallel ([Feat] Add strong checker to detect data racing in T.Parallel #1615)
• Add more curand operations with vectorization support ([Feature] Add more curand operations & support vectorization #1582)
• Support annotating loop layout in frontend ([Layout] Support annotating loop layout in frontend #1579)

Enhancements

• Enhance broadcast codegen for symbolic values ([CUDA] Enhance Broadcast Codegen for Symbolic Value #1669)
• Improve unroll loop functionality for dynamic extent ([Enhancement] Improve unroll loop functionality for dynamic extent and corresponding test case #1654)
• Introduce simulated load/store 256-bit access for CUDA compatibility ([CUDA] Introduce simulated load/store 256bits access for CUDA compatibility  #1656)
• Support nvidia-cuda-nvcc as nvcc ([Enhancement][CUDA] Support nvidia-cuda-nvcc as nvcc #1528)
• Allow import tilelang on CPU-only machines without CUDA libraries ([Enhancement] Allow import tilelang on CPU-only machines without CUDA libraries #1481)
• Improve symbolic shape/stride handling for subtype ([Enhancement][Subtype] Enhance symbolic shape/stride handling for subtype #1599)

Bug Fixes

• Fix thread storage synchronization logic in thread_storage_sync.cc
• Fix EagerJIT bug in handling of positional arguments ([EagerJIT] Fix bug in handling of positional arguments #1675)
• Fix variable scoping issue in InjectSoftwarePipeline ([Bugfix] Fix variable scoping issue in InjectSoftwarePipeline for transitive LetStmt dependencies #1657)
• Fix missing annotations for default CallNode Visitor ([Bugfix] Fix missing annotations for default CallNode Visitor #1659)
• Fix smem alloc for single warp reduce ([bugfix] fix smem alloc for single warp reduce #1643)
• Fix pdl related intrin handling ([Bugfix] Fix pdl related intrin handling to avoid strict annotation codegen #1650)
• Multiple bug fixes for ROCm/HIP support ([Bugfix] Open Rocm ci test and fix some bugs. #1443, [BugFix] Correct index_map selection for transposed A matrix in MFMA Layout with k_dim==4 and open rocm-ci for gemmsr #1627)

Refactoring

• Unify @jit and @lazy_jit into a single @jit decorator ([Refactor] Unify @jit and @lazy_jit into a single @jit decorator #1632)
• Move ConstrVisitor to src/transform/common/ for reuse ([Refactor] Move ConstrVisitor to src/transform/common/constr_visitor.h for reuse #1622)
• Move dtypes.py from eager to language module ([Refactor] Move dtypes.py from eager to language and add bits/bytes properties #1646)

Examples

• Add attention sink varlen examples ([Example] Add attention sink varlen examples #1645)
• Add Seesaw Sparse MLA Forward Kernel for DeepSeek-V3 ([Example] Add Seesaw Sparse MLA Forward Kernel for DeepSeek-V3.2 #1636)
• Add AMD fp8 preshuffle gemm example ([Enhancement][AMD] Add preshuffle fp8 gemm example on amd. #1605)

Test plan

• CI tests pass
• Verify existing functionality is not broken
• Run regression tests

🤖 Generated with Claude Code

Summary by CodeRabbit

• Bug Fixes
  • Fixed thread synchronization conflict detection to properly identify write-after-read memory access patterns in loop scenarios, ensuring correct ordering and preventing potential race conditions.

[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

Modifies conflict-detection logic in TileLangThreadSyncPlanner::FindConflict by removing a double-buffer-specific exception and introducing a new conflict rule for read-then-write patterns under loop-carry conditions in thread synchronization.

Changes

Cohort / File(s)	Summary
Conflict Detection Logic src/transform/thread_storage_sync.cc	Removed exception allowing reads during double-buffer swaps without loop-carry; added conflict detection for write-after-read in loop-carry scenarios. Control flow simplified by eliminating double-buffer special case.

Estimated code review effort

🎯 2 (Simple) | ⏱️ ~12 minutes

Possibly related PRs

• [ThreadSync] Use Z3 for constraint equivalence checking #1760: Modifies the same conflict-detection function FindConflict in thread synchronization logic.
• [Bugfix] Implement classic arena algorithm for shmem merge and WAW conflict detection #1146: Updates TileLangThreadSyncPlanner conflict-detection and read/write conflict handling in the same file.

Poem

🐰 A rabbit hops through sync with care,
Removes the double-buffer spare,
Loop-carry writes now caught in time,
No read-then-write conflicts to climb! ✨

✨ Finishing touches

• 📝 Generate docstrings

🧪 Generate unit tests (beta)

• Create PR with unit tests
• Post copyable unit tests in a comment
• Commit unit tests in branch sync_0203

---

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