• Lei Wang's avatar
    [TMA] Automatically lower 1d tma in appropriate cases (#788) · 9d7d45be
    Lei Wang authored
    * Enhance layout inference and copy operations with 1D TMA support
    
    - Updated `CopyNode` to introduce separate handling for 1D bulk load/store operations, including new methods for checking and lowering these operations.
    - Modified `InferLayout` and `GetCopyInst` to accommodate additional parameters for layout maps and analyzers.
    - Enhanced `AtomicAddNode` and `FillNode` to utilize the updated layout inference logic.
    - Improved buffer out-of-bounds checks during layout inference to ensure safe memory access.
    
    This update improves the efficiency and correctness of memory operations in the TileLang framework.
    
    * Refactor layout inference calls for improved readability
    
    - Updated `InferLayout` calls in `AtomicAddNode`, `CopyNode`, and `FillNode` to enhance code clarity by formatting parameters across multiple lines.
    - Cleaned up whitespace and formatting in `copy.h` and `layout_inference.cc` to adhere to coding standards and improve maintainability.
    
    This refactor aims to streamline the layout inference logic and improve overall code organization.
    
    * Fix shared tensor check in CopyNode for bulk copy operations
    
    - Updated the condition in `CheckBulkCopy1D` to verify contiguity of `shared_tensor` instead of `dst`, ensuring correct handling of shared memory layouts during bulk copy operations.
    - This change enhances the accuracy of memory operations in the TileLang framework.
    
    * Update test_example_gdn_compilation.py to invoke test function directly
    
    - Commented out the call to `tilelang.testing.main()` in `test_example_gdn_compilation.py` and replaced it with a direct call to `test_example_chunk_delta_bwd_compilation()`. This change simplifies the test execution flow and focuses on the specific test case.
    
    * Enhance bulk load/store checks in CopyNode with last dimension validation
    
    - Updated `CheckBulkLoad` and `CheckBulkStore` methods in `CopyNode` to include an optional parameter for validating the last dimension during bulk copy operations.
    - Adjusted related methods `CheckBulkLoad1D` and `CheckBulkStore1D` to pass the new parameter, improving the accuracy of bulk copy checks.
    - This change enhances the robustness of memory operations in the TileLang framework by ensuring compliance with dimensional requirements.
    
    * Refactor CheckBulkLoad and CheckBulkStore methods for improved readability
    
    - Reformatted the parameter lists of `CheckBulkLoad` and `CheckBulkStore` methods in `CopyNode` to enhance code clarity by aligning parameters across multiple lines.
    - This change improves the maintainability of the code and adheres to coding standards.
    9d7d45be
operator.h 2.66 KB