[REFACTOR][S-TIR] Lift transform passes to s_tir namespace

[tqchen]

This PR migrates the s_tir related transform passes into s_tir namespace instead. This set of changes can minimize the overall tir namespace to make it more focused.

[gemini-code-assist]

Summary of Changes

Hello @tqchen, I'm Gemini Code Assist¹! 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 undertakes a significant refactoring effort to improve the organization and clarity of the TVM codebase. It systematically moves all transformation passes specifically designed for S-TIR (Static Tensor IR) into their own distinct s_tir namespace. This separation helps to delineate responsibilities, reduce potential naming conflicts within the broader tir namespace, and makes it easier for developers to understand which passes apply to S-TIR.

Highlights

• Namespace Migration: All S-TIR related transformation passes have been migrated from the tir namespace to a dedicated s_tir namespace. This includes C++ header and source files, as well as Python bindings.
• File Renaming and Restructuring: Numerous C++ source files for S-TIR passes (e.g., canonicalize_loop.cc, lower_init_block.cc, inject_software_pipeline.cc) have been renamed and moved from src/tir/transforms/ to src/s_tir/transform/. Corresponding Python test files have also been renamed and moved.
• Build System and Import Updates: The CMakeLists.txt file has been updated to include the new src/s_tir/*.cc files. Python __init__.py files and various C++ and Python modules (e.g., s_tir/pipeline.py, meta_schedule components) have been updated to reflect the new s_tir.transform import paths and function calls.
• API Clean-up: The include/tvm/tir/transform.h and python/tvm/tir/transform/transform.py files have been cleaned up by removing the declarations and wrappers for the S-TIR specific passes, ensuring the tir namespace remains focused on general TIR transformations.

🧠 New Feature in Public Preview: You can now enable Memory to help Gemini Code Assist learn from your team's feedback. This makes future code reviews more consistent and personalized to your project's style. Click here to enable Memory in your admin console.

Changelog

• CMakeLists.txt
  • Added src/s_tir/*.cc to the COMPILER_SRCS list to ensure new S-TIR source files are compiled.
• include/tvm/s_tir/transform.h
  • New file created to declare S-TIR specific transformation passes, such as CanonicalizeLoop, LowerCrossThreadReduction, InjectSoftwarePipeline, etc.
• include/tvm/tir/transform.h
  • Removed declarations for S-TIR specific transformation passes, including LowerCrossThreadReduction, LowerInitBlock, PlanAndUpdateBufferAllocationLocation, ConvertBlocksToOpaque, LiftThreadBinding, CompactBufferAllocation, LowerMatchBuffer, InjectPermutedLayout, TransformMmaBufferLayout, LowerOpaqueBlock, UnifyThreadBinding, InjectSoftwarePipeline, LowerAutoCopy, ManifestSharedMemoryLocalStage, and AnnotateIrregularLoop.
• python/tvm/s_tir/init.py
  • Added from . import transform to expose the new S-TIR transform module.
• python/tvm/s_tir/backend/adreno/pipeline.py
  • Updated imports from tvm.tir to tvm.tir, s_tir.
  • Modified calls to S-TIR specific transform passes to use s_tir.transform. prefix instead of tir.transform..
• python/tvm/s_tir/pipeline.py
  • Updated imports from tvm.tir to tvm.tir, s_tir.
  • Modified calls to S-TIR specific transform passes to use s_tir.transform. prefix instead of tir.transform..
• python/tvm/s_tir/transform/init.py
  • New file created to define the s_tir.transform namespace.
• python/tvm/s_tir/transform/_ffi_api.py
  • New file created to initialize FFI APIs for s_tir.transform.
• python/tvm/s_tir/transform/transform.py
  • New file created, containing Python wrappers for S-TIR specific transformation passes.
• python/tvm/tir/transform/transform.py
  • Removed Python wrappers for S-TIR specific transformation passes, such as AnnotateIrregularLoop, LowerCrossThreadReduction, LowerInitBlock, PlanAndUpdateBufferAllocationLocation, ConvertBlocksToOpaque, LiftThreadBinding, CompactBufferAllocation, LowerMatchBuffer, LowerOpaqueBlock, TransformMmaBufferLayout, InjectPermutedLayout, UnifyThreadBinding, InjectSoftwarePipeline, LowerAutoCopy, and ManifestSharedMemoryLocalStage.
• src/meta_schedule/feature_extractor/per_store_feature.cc
  • Added #include <tvm/s_tir/transform.h>.
  • Updated calls to S-TIR specific transform passes to use s_tir::transform:: prefix.
• src/meta_schedule/postproc/disallow_async_strided_mem_copy.cc
  • Added #include <tvm/s_tir/transform.h>.
  • Updated calls to S-TIR specific transform passes to use s_tir::transform:: prefix.
• src/meta_schedule/postproc/verify_gpu_code.cc
  • Added #include <tvm/s_tir/transform.h>.
  • Updated calls to S-TIR specific transform passes to use s_tir::transform:: prefix.
• src/tir/analysis/calculate_allocated_memory.cc
  • Added #include <tvm/s_tir/transform.h>.
  • Updated calls to S-TIR specific transform passes to use s_tir::transform:: prefix.
• src/tir/transforms/annotate_irregular_loop.cc
  • Renamed to src/s_tir/transform/annotate_irregular_loop.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to attr:: to tir::attr:: and CreatePrimFuncPass pass name to s_tir.AnnotateIrregularLoop.
• src/tir/transforms/canonicalize_loop.cc
  • Renamed to src/s_tir/transform/canonicalize_loop.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.CanonicalizeLoop.
• src/tir/transforms/compact_buffer_region.cc
  • Renamed to src/s_tir/transform/compact_buffer_region.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to attr:: to tir::attr:: and CreatePrimFuncPass pass name to s_tir.CompactBufferAllocation.
• src/tir/transforms/convert_blocks_to_opaque.cc
  • Renamed to src/s_tir/transform/convert_blocks_to_opaque.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.ConvertBlocksToOpaque.
• src/tir/transforms/inject_permuted_layout.cc
  • Renamed to src/s_tir/transform/inject_permuted_layout.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.InjectPermutedLayout.
• src/tir/transforms/inject_software_pipeline.cc
  • Renamed to src/s_tir/transform/inject_software_pipeline.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to attr:: to tir::attr:: and CreatePrimFuncPass pass name to s_tir.InjectSoftwarePipeline.
• src/tir/transforms/lift_thread_binding.cc
  • Renamed to src/s_tir/transform/lift_thread_binding.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.LiftThreadBinding.
• src/tir/transforms/lower_cross_thread_reduction.cc
  • Renamed to src/s_tir/transform/lower_cross_thread_reduction.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.LowerCrossThreadReduction.
• src/tir/transforms/lower_init_block.cc
  • Renamed to src/s_tir/transform/lower_init_block.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.LowerInitBlock.
• src/tir/transforms/lower_match_buffer.cc
  • Renamed to src/s_tir/transform/lower_match_buffer.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.LowerMatchBuffer.
• src/tir/transforms/lower_opaque_block.cc
  • Renamed to src/s_tir/transform/lower_opaque_block.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to attr:: to tir::attr:: and CreatePrimFuncPass pass name to s_tir.LowerOpaqueBlock.
• src/tir/transforms/manifest_shared_memory_local_stage.cc
  • Renamed to src/s_tir/transform/manifest_shared_memory_local_stage.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to attr:: to tir::attr:: and CreatePrimFuncPass pass name to s_tir.ManifestSharedMemoryLocalStage.
• src/tir/transforms/memhammer_lower_auto_copy.cc
  • Renamed to src/s_tir/transform/memhammer_lower_auto_copy.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.LowerAutoCopy.
• src/tir/transforms/plan_update_buffer_allocation_location.cc
  • Renamed to src/s_tir/transform/plan_update_buffer_allocation_location.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.PlanAndUpdateBufferAllocationLocation.
• src/tir/transforms/transform_mma_buffer_layout.cc
  • Renamed to src/s_tir/transform/transform_mma_buffer_layout.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to CreatePrimFuncPass pass name to s_tir.TransformMmaBufferLayout.
• src/tir/transforms/unify_thread_binding.cc
  • Renamed to src/s_tir/transform/unify_thread_binding.cc.
  • Updated namespace from tvm::tir to tvm::s_tir.
  • Updated internal references to attr:: to tir::attr:: and CreatePrimFuncPass pass name to s_tir.UnifyThreadBinding.
• tests/python/tir-analysis/test_tir_analysis_calculate_allocated_memory.py
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_annotate_irregular_loop.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_annotate_irregular_loop.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_canonicalize_loop.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_canonicalize_loop.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_compact_buffer_region.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_compact_buffer_region.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_convert_blocks_to_opaque.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_convert_blocks_to_opaque.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_inject_permuted_layout.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_inject_permuted_layout.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py
  • Modified calls to LowerOpaqueBlock to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_inject_software_pipeline.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_inject_software_pipeline.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_lift_thread_binding.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_lift_thread_binding.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_loop_partition.py
  • Modified calls to LowerOpaqueBlock to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_lower_cross_thread_reduction.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_lower_cross_thread_reduction.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_lower_init_block.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_lower_init_block.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_lower_match_buffer.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_lower_match_buffer.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_lower_opaque_block.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_lower_opaque_block.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_manifest_shared_memory_local_stage.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_manifest_shared_memory_local_stage.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_memhammer_lower_auto_copy.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_memhammer_lower_auto_copy.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_plan_update_buffer_allocation_location.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_plan_update_buffer_allocation_location.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tir-transform/test_tir_transform_unify_thread_binding.py
  • Renamed to tests/python/s_tir/transform/test_s_tir_transform_unify_thread_binding.py.
  • Updated imports to include tvm.s_tir.
  • Modified calls to transform passes to use tvm.s_tir.transform. prefix.
• tests/python/tvmscript/test_tvmscript_roundtrip.py
  • Modified calls to LowerOpaqueBlock to use tvm.s_tir.transform. prefix.

Activity

• The author tqchen initiated a large-scale refactoring to move S-TIR specific transformation passes into a dedicated s_tir namespace.
• This involved creating new C++ header and Python files for the s_tir.transform module.
• Existing S-TIR related passes were removed from the general tir namespace in both C++ and Python.
• The build system (CMakeLists.txt) was updated to compile the newly organized source files.
• Numerous Python and C++ files that utilize these passes were updated to import and call them from the new s_tir.transform namespace.
• Associated test files were also renamed and modified to align with the new namespace structure.

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 by creating a comment using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands on the current page.

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 pull request 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. ↩

[gemini-code-assist]

Code Review

This pull request is a large-scale refactoring that moves S-TIR specific transform passes into their own s_tir namespace. This is a great improvement for code organization and modularity. The changes are extensive, touching C++ source files, Python bindings, and build configurations. The file moves, namespace updates, and adjustments in dependent code seem correct and consistent with the goal. I've found one minor pre-existing issue with some unreachable code that was moved as part of this refactoring and have left a suggestion to fix it. Overall, this is a solid refactoring effort.