[Language] Support accumulative T.reduce_sum
#436
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
… and abs sum (tile-ai#436) * Modified reduce_sum and reduce_absmax functions to include a clear parameter, allowing for accumulation on existing values. * Updated ReduceOp::Lower method to handle initialization and buffer duplication based on the clear flag for sum and abs sum operations. * Added new tests for reduce_sum and reduce_max with clear functionality to ensure correctness in various scenarios. * Enhanced documentation for reduce functions to clarify the behavior of the clear parameter.
LeiWang1999
added a commit
to LeiWang1999/tilelang
that referenced
this pull request
Apr 26, 2025
… add atomic add test (tile-ai#436) * Removed redundant simplification step in vectorization logic to streamline performance. * Introduced a new test for atomic addition in TileLang, validating functionality with a reference implementation using PyTorch.
…vice_regions.py from Buffer to Tensor
LeiWang1999
added a commit
to LeiWang1999/tilelang
that referenced
this pull request
Jul 18, 2025
* [Enhancement] Update reduce operations to support clear option in sum and abs sum (tile-ai#436) * Modified reduce_sum and reduce_absmax functions to include a clear parameter, allowing for accumulation on existing values. * Updated ReduceOp::Lower method to handle initialization and buffer duplication based on the clear flag for sum and abs sum operations. * Added new tests for reduce_sum and reduce_max with clear functionality to ensure correctness in various scenarios. * Enhanced documentation for reduce functions to clarify the behavior of the clear parameter. * lint fix * Update tensor type annotations in test_tilelang_transform_annotate_device_regions.py from Buffer to Tensor * Update tensor type in reduce sum tests from float16 to float32 for improved precision
LeiWang1999
added a commit
to LeiWang1999/tilelang
that referenced
this pull request
Jul 18, 2025
… add atomic add test (tile-ai#436) (tile-ai#439) * Removed redundant simplification step in vectorization logic to streamline performance. * Introduced a new test for atomic addition in TileLang, validating functionality with a reference implementation using PyTorch.
LeiWang1999
added a commit
to LeiWang1999/tilelang
that referenced
this pull request
Jul 20, 2025
* [Enhancement] Update reduce operations to support clear option in sum and abs sum (tile-ai#436) * Modified reduce_sum and reduce_absmax functions to include a clear parameter, allowing for accumulation on existing values. * Updated ReduceOp::Lower method to handle initialization and buffer duplication based on the clear flag for sum and abs sum operations. * Added new tests for reduce_sum and reduce_max with clear functionality to ensure correctness in various scenarios. * Enhanced documentation for reduce functions to clarify the behavior of the clear parameter. * lint fix * Update tensor type annotations in test_tilelang_transform_annotate_device_regions.py from Buffer to Tensor * Update tensor type in reduce sum tests from float16 to float32 for improved precision
LeiWang1999
added a commit
to LeiWang1999/tilelang
that referenced
this pull request
Jul 20, 2025
… add atomic add test (tile-ai#436) (tile-ai#439) * Removed redundant simplification step in vectorization logic to streamline performance. * Introduced a new test for atomic addition in TileLang, validating functionality with a reference implementation using PyTorch.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This pull request introduces enhancements to the reduction operations in the TileLang framework, focusing on enabling accumulation without clearing output buffers, improving flexibility for specific use cases. It also adds comprehensive tests for the new functionality and updates the implementation of reduction operations to support this feature.
Enhancements to Reduction Operations:
Support for Accumulation Without Clearing Buffers:
clearparameter toreduce_sumandreduce_absmaxfunctions intilelang/language/reduce.py, allowing accumulation on existing buffer values when set toFalse. Updated the implementation to handle this behavior by using a temporary buffer during reduction. [1] [2] [3]Changes in
ReduceOp::LowerImplementation:src/op/reduce.ccto support the newclearbehavior. Introduced conditional buffer duplication and accumulation logic forkSumandkAbsSumreduction types. [1] [2] [3] [4]Testing Improvements:
clearParameter:reduce_sumandreduce_maxoperations with theclearparameter set toFalse. These tests validate the accumulation behavior and ensure correctness. [1] [2]Codebase Maintenance:
tir/transforms/ir_utils.hinsrc/op/reduce.ccto support the updated implementation.