Skip to content

Conversation

@LeiWang1999
Copy link
Member

@LeiWang1999 LeiWang1999 commented Dec 10, 2025

This pull request updates the TVM submodule and refactors the pipeline injection logic in inject_pipeline.cc to improve correctness and robustness of async pipeline scheduling. The main changes include more accurate wait count computation for async dependencies, enhanced symbolic analysis for loop bounds, and richer metadata for pipeline blocks.

Pipeline scheduling and async dependency handling:

  • Refactored wait count computation in PopulateWaitCounts to precisely track async producer-consumer relationships using commit groups, buffer dependencies, and symbolic analysis. This ensures correct synchronization and reduces unnecessary waits.
  • Added commit_predicate to async state tracking and propagate it through the pipeline, improving the handling of conditional execution of commit blocks. [1] [2]

Loop and block metadata improvements:

  • Extended PipelineAnnotation and RewrittenBlockInfo structs to include original_idx, start, and end fields, enabling more accurate mapping and analysis of block order and loop bounds. [1] [2] [3]
  • Enhanced symbolic analysis by binding loop variables to their iteration domains and maintaining constraint contexts for improved proof strength during transformations.

General codebase updates:

  • Updated the TVM submodule to a newer commit, ensuring compatibility with recent upstream changes.
  • Added includes for <memory> and <sstream> to support new features in the refactored code.

Summary by CodeRabbit

  • Refactor

    • Improved pipeline injection with per-block tracking, clearer start/end propagation, epilogue-aware emission, and stronger async commit/wait handling.
  • Tests

    • Updated pipeline transform test to include an additional simplify pass for more robust verification.
  • Chores

    • Updated third-party submodule.

✏️ Tip: You can customize this high-level summary in your review settings.

…lation

- Adjusted output directories for the tilelang_cython_wrapper to ensure that development builds place the extension in build/lib.
- Updated installation paths to place the extension in tilelang/lib within the wheel, improving organization and avoiding potential conflicts with other modules.
- Modified the internal library path exposure in env.py to prevent shadowing of common module names, enhancing compatibility and usability in user projects.
- Set output directories for both tilelang and tilelang_module libraries to "${CMAKE_BINARY_DIR}/lib" for consistency in development builds.
- This change enhances organization and ensures that all build artifacts are located in a unified directory structure.
- Updated the TVM subproject to commit 90581fe9e5287bbcf1844ad14255a1e1e8cdf7f0.
- Added new fields to `PipelineAnnotation` and `RewrittenBlockInfo` structures to track original statement indices and improve async state management.
- Refactored `EmitImpl` and `PopulateWaitCounts` methods to enhance clarity and functionality, including better handling of commit groups and wait counts.
- Simplified access index calculations and strengthened analyzer constraints for loop bounds.
…ne.cc

- Eliminated the Apache license block from the top of the file to streamline the code.
- Removed unused include directives for memory and stringstream to enhance code clarity and reduce unnecessary dependencies.
@github-actions
Copy link

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

Please remember to run pre-commit run --all-files 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! 🚀

@coderabbitai
Copy link
Contributor

coderabbitai bot commented Dec 10, 2025

Walkthrough

Updates a TVM submodule reference and extends the software-pipeline injection: adds per-block original index and start/end iteration bounds, makes emission and wait-count logic epilogue-aware, records per-async-stage commit predicates, and adjusts tests and phase pipeline to run extra Simplify passes.

Changes

Cohort / File(s) Summary
Submodule update
3rdparty/tvm
Subproject commit reference updated (3a32b763... -> 90581fe9...).
Pipeline injection core
src/transform/inject_pipeline.cc
Added original_idx to PipelineAnnotation; added start/end to RewrittenBlockInfo; added commit_predicate to AsyncStateLocal; extended EmitImpl(..., is_epilogue=false) and PopulateWaitCounts(..., is_epilogue=false) signatures; propagate per-block start/end and original index; make EmitImpl epilogue-aware and strengthen bound analysis; rework async wait-count computation and store per-block pending waits.
Tests
testing/python/transform/test_tilelang_transform_Inject_software_pipeline.py
Test now runs an extra Simplify pass after LowerOpaqueBlock before comparison.
Pipeline phase orchestration
tilelang/engine/phase.py
OptimizeForTarget now inserts a Simplify pass immediately after LowerOpaqueBlock in both code paths.

Sequence Diagram

sequenceDiagram
    participant Collector as BlockCollector
    participant Annotator as PipelineAnnotator
    participant Emitter as Emitter
    participant Waiter as AsyncWaitComputer

    Collector->>Annotator: collect blocks (index i)
    Annotator->>Annotator: create PipelineAnnotation(original_idx = i)
    Annotator->>Annotator: record block start/end in RewrittenBlockInfo

    rect rgb(220,240,255)
    Note over Emitter: Prologue / Body emission (is_epilogue=false)
    Collector->>Emitter: EmitImpl(start,end,..., is_epilogue=false)
    Emitter->>Emitter: bind loop var domain\napply bound guards
    end

    rect rgb(255,240,220)
    Note over Emitter,Waiter: Epilogue emission (is_epilogue=true)
    Collector->>Emitter: EmitImpl(start,end,..., is_epilogue=true)
    Emitter->>Waiter: PopulateWaitCounts(new_blocks, async_states, is_epilogue=true)
    Waiter->>Waiter: build order->access index map\nidentify dependent commit groups
    Waiter->>Waiter: compute waits (epilogue-aware)\nrecord commit_predicate in AsyncStateLocal
    Waiter-->>Emitter: store pending_waits per block
    end
Loading

Estimated code review effort

🎯 4 (Complex) | ⏱️ ~50 minutes

  • Verify all call sites updated for is_epilogue additions.
  • Inspect async wait-count logic and mapping from orders to access indices.
  • Confirm original_idx, start/end, and commit_predicate are consistently propagated and not lossy.
  • Review bound-analysis changes for correctness in symbolic and non-unit loop cases.

Possibly related PRs

Poem

🐰
I hopped through blocks with quiet cheer,
Marked every index far and near,
Start and end, prologue, epilogue too,
Async waits and commits in view,
A tiny rabbit's pipeline debut! 🥕✨

Pre-merge checks and finishing touches

❌ Failed checks (1 warning)
Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 16.67% which is insufficient. The required threshold is 80.00%. You can run @coderabbitai generate docstrings to improve docstring coverage.
✅ Passed checks (2 passed)
Check name Status Explanation
Description Check ✅ Passed Check skipped - CodeRabbit’s high-level summary is enabled.
Title check ✅ Passed The title directly reflects the main changes in the PR: refactoring pipeline injection logic to support dynamic pipeline extents with improved async dependency handling, block metadata, and loop-bound analysis.
✨ Finishing touches
  • 📝 Generate docstrings
🧪 Generate unit tests (beta)
  • Create PR with unit tests
  • Post copyable unit tests in a comment

📜 Recent review details

Configuration used: CodeRabbit UI

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between d369f52 and 0c5616e.

📒 Files selected for processing (2)
  • testing/python/transform/test_tilelang_transform_Inject_software_pipeline.py (1 hunks)
  • tilelang/engine/phase.py (1 hunks)
🧰 Additional context used
🧬 Code graph analysis (1)
tilelang/engine/phase.py (2)
src/transform/simplify.cc (2)
  • Simplify (530-538)
  • Simplify (530-530)
tilelang/transform/simplify.py (1)
  • Simplify (20-28)
⏰ Context from checks skipped due to timeout of 90000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (2)
  • GitHub Check: Test for Python 3.12 with Metal (on macos-latest)
  • GitHub Check: Test for Python 3.12 with CUDA-12.8 (on self-hosted-nvidia)
🔇 Additional comments (2)
tilelang/engine/phase.py (1)

219-220: LGTM! Simplify pass added after LowerOpaqueBlock aligns with pipeline refactoring.

The addition of Simplify() immediately after LowerOpaqueBlock() is consistent with the test changes and appears to support the enhanced symbolic analysis mentioned in the PR objectives. The refactored pipeline injection logic likely introduces expressions that benefit from early simplification before subsequent transformations.

Note that line 226 still contains tir.transform.Simplify(), so you now have two Simplify passes in sequence (tilelang's at line 220, TVM's at line 226), which is acceptable if they serve different purposes.

testing/python/transform/test_tilelang_transform_Inject_software_pipeline.py (1)

13-13: LGTM! Test updated to match the production pipeline.

The additional Simplify() pass correctly mirrors the change in tilelang/engine/phase.py, ensuring the test validates the same transformation sequence used in production. This maintains consistency between the test suite and the actual pipeline implementation.


Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands and usage tips.

- Added an additional Simplify transformation in the InjectSoftwarePipeline to improve optimization.
- Updated the test file to call `test_trival_pipeline()` directly, commenting out the previous main execution for better test isolation.
@LeiWang1999 LeiWang1999 merged commit f2858fa into tile-ai:main Dec 10, 2025
5 of 6 checks passed
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