• Lei Wang's avatar
    [Refactor] Update buffer handling in copy and atomic operations (#1247) · 2c0072a8
    Lei Wang authored
    * [Refactor] Update buffer handling in copy and atomic operations
    
    * Refactored the `copy` and `atomic_add` functions to use element-wise minimum for defining copy extents, ensuring correct handling of overlapping regions.
    * Updated utility functions to create `BufferLoad` instances with explicit extents, improving memory management and clarity.
    * Removed unused imports from `atomic.py` and `copy.py` to streamline the codebase.
    * Adjusted logging in `copy.cc` to provide clearer warnings for fallback scenarios in bulk copy operations.
    
    * Remove obsolete .git_commit.txt file
    
    * Add unit test for dynamic copy extent handling in TileLang
    
    * Introduced a new test file `test_tilelang_issue_1237.py` to verify that the `T.copy` function correctly manages dynamic extents during primitive function building.
    * The test reproduces a specific issue related to dynamic slice lengths and static buffer sizes, ensuring robustness in the handling of such scenarios.
    * The test does not require execution of the kernel, as building the primitive function is sufficient to validate the fix.
    
    * lint fix
    
    * fix
    
    * Revert "fix"
    
    This reverts commit 828b4c1e4de76a7d11e4d4092927303fbbe00097.
    
    * Update TVM submodule and refactor atomic and copy functions
    
    * Updated the TVM submodule to a dirty state.
    * Refactored `atomic_add` and `copy` functions to pass extents explicitly to the `_to_region` helper, improving clarity and correctness in handling buffer regions.
    * Commented out the main execution call in the test example for `cast` and added a new function call to better demonstrate the example usage.
    
    * Enhance extent handling in atomic and copy functions
    
    * Introduced `legalize_pairwise_extents` utility to align and broadcast extent lists for `atomic_add` and `copy` functions, ensuring compatibility and correctness in buffer operations.
    * Updated both functions to utilize the new utility, improving clarity and robustness in handling dynamic and static extents.
    * Added comments to clarify the extent handling logic.
    
    * Enhance `legalize_pairwise_extents` function with early-exit rule
    
    * Added an early-exit condition to the `legalize_pairwise_extents` function to return original extents if the number of non-1 dimensions in both source and destination extents is equal, improving performance by avoiding unnecessary adjustments.
    * Updated the function's documentation to clarify the new behavior and maintain clarity in the extent handling logic.
    
    * lint fix
    2c0072a8
copy.cc 79.5 KB