[Bugfix] Fix safe memory legalization for fragment store #446
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
…i#441) * Added logic to use non-replicated buffers as source buffers for more accurate layout inference. * Enhanced comments to clarify the rationale behind buffer selection in layout inference process.
…g logic * Introduced TILELANG_CHECK macro for improved error handling in CUDA and HIP code, providing detailed error messages for kernel launches. * Enhanced loop partitioning logic to handle fragment buffers more effectively, ensuring correct replication based on thread extent. * Added logging for thread range in PlanLoopPartition to aid in debugging and performance analysis. * Updated pass configuration management to streamline vectorization control in the optimization process.
…cess handling * Replaced Apache License header with MIT License. * Added logic to handle local buffer conditions in memory access. * Introduced IsLocalBuffer function to check buffer scope. * Enhanced comments for clarity on memory access operations.
LeiWang1999
added a commit
to LeiWang1999/tilelang
that referenced
this pull request
Jul 18, 2025
* [Enhancement] Improve layout inference accuracy in ParallelOp (tile-ai#441) * Added logic to use non-replicated buffers as source buffers for more accurate layout inference. * Enhanced comments to clarify the rationale behind buffer selection in layout inference process. * [Enhancement] Add error handling macros and refactor loop partitioning logic * Introduced TILELANG_CHECK macro for improved error handling in CUDA and HIP code, providing detailed error messages for kernel launches. * Enhanced loop partitioning logic to handle fragment buffers more effectively, ensuring correct replication based on thread extent. * Added logging for thread range in PlanLoopPartition to aid in debugging and performance analysis. * Updated pass configuration management to streamline vectorization control in the optimization process. * lint fix * remove debug print * [Refactor] Update legalize_safe_memory_access.cc to improve memory access handling * Replaced Apache License header with MIT License. * Added logic to handle local buffer conditions in memory access. * Introduced IsLocalBuffer function to check buffer scope. * Enhanced comments for clarity on memory access operations.
LeiWang1999
added a commit
to LeiWang1999/tilelang
that referenced
this pull request
Jul 20, 2025
* [Enhancement] Improve layout inference accuracy in ParallelOp (tile-ai#441) * Added logic to use non-replicated buffers as source buffers for more accurate layout inference. * Enhanced comments to clarify the rationale behind buffer selection in layout inference process. * [Enhancement] Add error handling macros and refactor loop partitioning logic * Introduced TILELANG_CHECK macro for improved error handling in CUDA and HIP code, providing detailed error messages for kernel launches. * Enhanced loop partitioning logic to handle fragment buffers more effectively, ensuring correct replication based on thread extent. * Added logging for thread range in PlanLoopPartition to aid in debugging and performance analysis. * Updated pass configuration management to streamline vectorization control in the optimization process. * lint fix * remove debug print * [Refactor] Update legalize_safe_memory_access.cc to improve memory access handling * Replaced Apache License header with MIT License. * Added logic to handle local buffer conditions in memory access. * Introduced IsLocalBuffer function to check buffer scope. * Enhanced comments for clarity on memory access operations.
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 updates the
legalize_safe_memory_accesstransformation in thesrc/transform/legalize_safe_memory_access.ccfile. The changes include updating the licensing information, introducing support for local buffers in safe memory access legalization, and adding a helper method to identify local buffers.Licensing update:
Enhancements to safe memory access legalization:
SafeMemorysRewriterclass. This includes adding conditions to ensure proper handling of local buffers and throwing a fatal error for unsupported buffer types.IsLocalBuffer, to check if a buffer belongs to the "local" or "local.fragment" scope.Minor code cleanup:
SafeMemorysRewriterclass for improved readability.