mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-23 07:00:29 +08:00
Lei WangandClaude Opus 4.6
9246b82b24
* [Refactor] Replace create_list_of_mbarrier/get_mbarrier with buffer-based T.alloc_barrier Unify the two mbarrier representation paths into a single buffer-based approach using shared.barrier scope buffers with barrier_init block annotations. This eliminates the create_list_of_mbarrier and get_mbarrier intermediate intrinsics, making the IR cleaner and the barrier handling consistent across user-explicit (T.alloc_barrier) and compiler-generated (T.copy/T.tma_copy) paths. Key changes: - LowerTileOp now injects a shared.barrier buffer + barrier_init annotation into the root block instead of emitting create_list_of_mbarrier statements - copy.cc uses BufferLoad on the barrier buffer instead of get_mbarrier - WarpSpecializedRewriter and ProducerConsumerWS create barrier buffers with annotations instead of create_list_of_mbarrier calls - LowerSharedBarrier runs a second time after WS passes to process newly created barriers - LowerHopperIntrin no longer handles mbarrier allocation/init (fully delegated to LowerSharedBarrier) - Removed create_list_of_mbarrier and get_mbarrier op definitions from builtin.h/cc - Removed get_mbarrier codegen handling from CUDA, CuTeDSL, HIP backends - Updated WarpSpecializedDetector, inject_tma_barrier, and eliminate_storage_sync passes to detect barriers via shared.barrier scope buffers instead of the removed intrinsics Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * style: apply clang-format and ruff-format Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Refactor mbarrier handling in copy operations Updated the handling of mbarriers in the `LowerBulkCopy` and `LowerBulkCopy1D` methods to allocate a single barrier slot instead of multiple for each pipeline stage. This change simplifies the barrier management by allowing the `MultiVersionBuffer` to expand the barrier for pipelining stages. Additionally, modified the `LowerArgs` structure to use a pointer for the `mbarrier_buffer`, ensuring compatibility with lazy buffer creation. Key changes: - Simplified mbarrier allocation logic in `copy.cc`. - Updated `LowerArgs` to use a pointer for `mbarrier_buffer`. - Adjusted related comments for clarity. This refactor enhances the consistency and readability of the code while maintaining functionality. * Remove deprecated TMA barrier and warp specialization files This commit deletes the `inject_tma_barrier.cc`, `warp_specialized_rewriter.cc`, and associated test files as part of a cleanup effort. These files are no longer necessary following recent refactoring and updates to the TMA pipeline and warp specialization logic. The removal simplifies the codebase and eliminates outdated implementations that have been superseded by newer approaches. Key changes: - Deleted `inject_tma_barrier.cc` and `warp_specialized_rewriter.cc`. - Removed related test file `test_tilelang_transform_warp_specialized.py`. - Updated `phase.py` and `__init__.py` to reflect the removal of the WarpSpecialized and InjectTmaBarrier functionalities. This cleanup enhances maintainability and clarity in the transformation pipeline. --------- Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>