Skip to content

Conversation

@LeiWang1999
Copy link
Member

Enhance Global Read Detection in Pipeline Planning

  • Improved the handling of global read detection by considering condition expressions within IfThenElse nodes, ensuring accurate identification of global memory accesses.
  • Added a dedicated flag to track whether the visitor is currently within a condition expression, enhancing the precision of buffer access analysis.
  • Refactored the VisitStmt_ method to more cleanly process the structure of IfThenElse nodes, improving both readability and maintainability of the code.
// Skip condition expression of an IfThenElse node
// Example 1 (not a global read):
// shared[i] = T.if_then_else(global[i] < n, register_a[i], register_b[i])
//
// Example 2 (is a global read):
// shared[i] = T.if_then_else(global[i] < n, global_a[i], global_b[i])

…ndition handling

- Introduced a new function to replace IfThenElse nodes with their then_case while preserving attributes.
- Streamlined the PipelineBodyRewriter to improve buffer access rewriting and async state management.
- Enhanced the handling of pipeline loop conditions and added support for predicate conditions in the pipeline body.
- Removed obsolete code and improved overall code clarity and maintainability.
…y std::move calls

- Updated return statements in multiple methods to return objects directly instead of using std::move, improving code clarity and potentially avoiding unnecessary moves.
- Ensured consistent handling of BufferStore and BufferLoad nodes during pipeline transformations.
- Updated the handling of global reads to account for condition expressions within IfThenElse nodes, ensuring accurate identification of global memory accesses.
- Introduced a new flag to track whether the visitor is within a condition expression, improving the correctness of buffer access analysis.
- Refactored the VisitStmt_ method to properly handle the structure of IfThenElse nodes, enhancing the clarity and maintainability of the code.
@github-actions
Copy link

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run bash format.sh in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work!

🚀

Copy link
Contributor

@gemini-code-assist gemini-code-assist bot left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Summary of Changes

Hello @LeiWang1999, I'm Gemini Code Assist1! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!

This pull request refines the global memory access detection within the pipeline planning phase, specifically addressing how IfThenElse conditional expressions are handled. Previously, global reads within these conditions might have been incorrectly flagged. The changes ensure that only actual data accesses in the then or else branches contribute to global read detection, improving the accuracy of pipeline analysis.

Highlights

  • Accurate Global Read Detection: The core improvement ensures that global memory accesses are only considered "global reads" if they occur within the then or else branches of an IfThenElse statement, not within the condition expression itself. This prevents misclassification and leads to more precise pipeline planning.
  • Enhanced BufferRegionCollector: A new internal flag, within_condition_expr_, has been introduced in the BufferRegionCollector to accurately track whether the visitor is currently traversing a conditional expression, enabling the refined global read detection.
  • Streamlined Pipeline Rewriting Logic: Significant refactoring has been performed within the PipelineRewriter and PipelineBodyRewriter classes. This includes removing redundant parameters (like fragment_info and double_buffers), simplifying constructors, and consolidating buffer access rewriting logic.
  • Improved Asynchronous Operation Handling: The internal structures and logic for managing asynchronous states (AsyncStateGlobal, AsyncStateLocal) and calculating wait counts have been revised, leading to a more robust and simplified implementation for handling async operations within the software pipeline.
Using Gemini Code Assist

The full guide for Gemini Code Assist can be found on our documentation page, here are some quick tips.

Invoking Gemini

You can request assistance from Gemini at any point in your pull request via creating an issue comment (i.e. comment on the pull request page) using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands.

Feature Command Description
Code Review /gemini review Performs a code review for the current pull request in its current state.
Pull Request Summary /gemini summary Provides a summary of the current pull request in its current state.
Comment @gemini-code-assist Responds in comments when explicitly tagged, both in issue comments and review comments.
Help /gemini help Displays a list of available commands.

Customization

To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a .gemini/ folder in the base of the repository. Detailed instructions can be found here.

Limitations & Feedback

Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with 👍 and 👎 on @gemini-code-assist comments. If you're interested in giving your feedback about your experience with Gemini Code Assist for Github and other Google products, sign up here.

You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the Gemini Code Assist IDE Extension.

Footnotes

  1. Review the Privacy Notices, Generative AI Prohibited Use Policy, Terms of Service, and learn how to configure Gemini Code Assist in GitHub here. Gemini can make mistakes, so double check it and use code with caution.

Copy link
Contributor

@gemini-code-assist gemini-code-assist bot left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request introduces significant refactoring to the software pipeline injection logic, aiming to simplify async handling and improve global read detection. The changes are extensive, especially in PipelineRewriter, where the logic for managing async states and generating pipeline stages has been completely rewritten for clarity. A key feature is the new mechanism to skip analysis of global reads within condition expressions in IfThenElse nodes, which is a good improvement for accuracy.

However, I've identified a few areas of concern. There appears to be a critical regression where handling for several important tensor core intrinsics (like mma_sync, load_matrix_sync) was removed during the refactoring of PipelineOpaqueAccessRewriter. Additionally, there's a minor issue with swapped error messages and a potential bug in the IfThenElse visitor in pipeline_planning.cc that could cause global reads in else blocks to be missed.

Overall, the refactoring is a positive step towards maintainability, but the identified issues, particularly the potential regression, should be addressed.

@LeiWang1999 LeiWang1999 merged commit c1eef51 into tile-ai:main Aug 13, 2025
4 checks passed
RubiaCx pushed a commit to RubiaCx/tilelang that referenced this pull request Nov 24, 2025
…e-ai#713)

* Update submodule 'tvm' to commit e11521e6936a827efa334588d29571fbb4620107

* Refactor inject_pipeline.cc to enhance pipeline body rewriting and condition handling

- Introduced a new function to replace IfThenElse nodes with their then_case while preserving attributes.
- Streamlined the PipelineBodyRewriter to improve buffer access rewriting and async state management.
- Enhanced the handling of pipeline loop conditions and added support for predicate conditions in the pipeline body.
- Removed obsolete code and improved overall code clarity and maintainability.

* lint fix

* Refactor return statements in inject_pipeline.cc to remove unnecessary std::move calls

- Updated return statements in multiple methods to return objects directly instead of using std::move, improving code clarity and potentially avoiding unnecessary moves.
- Ensured consistent handling of BufferStore and BufferLoad nodes during pipeline transformations.

* test fix

* Enhance global read detection in pipeline planning

- Updated the handling of global reads to account for condition expressions within IfThenElse nodes, ensuring accurate identification of global memory accesses.
- Introduced a new flag to track whether the visitor is within a condition expression, improving the correctness of buffer access analysis.
- Refactored the VisitStmt_ method to properly handle the structure of IfThenElse nodes, enhancing the clarity and maintainability of the code.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant