Skip to content

Conversation

@LeiWang1999
Copy link
Member

@LeiWang1999 LeiWang1999 commented Oct 21, 2025

This pull request adds a new configuration option to control inplace buffer detection in the StorageRewrite pass, improves documentation for pass configuration options, and introduces a test to verify the feature. The main change is the introduction of the tl.storage_rewrite_detect_inplace option, which allows users to enable or disable aggressive buffer reuse in kernel compilation.

New StorageRewrite configuration option:

  • Added the kStorageRewriteDetectInplace option to TVM's pass config registry (src/op/builtin.cc, src/op/builtin.h). [1] [2]
  • Updated the StorageRewrite pass to respect the new config and allow inplace buffer detection when enabled (src/transform/storage_rewrite.cc). [1] [2]
  • Added TL_STORAGE_REWRITE_DETECT_INPLACE to PassConfigKey enum with a detailed docstring explaining its behavior and usage (tilelang/transform/pass_config.py).

Documentation improvements:

  • Updated docstrings in tilelang/autotuner/param.py, tilelang/jit/__init__.py, and tilelang/jit/kernel.py to refer users to PassConfigKey for supported options, rather than listing them inline. [1] [2] [3]

Testing:

  • Added a new test file, test_storage_rewrite_detect_inplace.py, to verify that enabling the inplace detection option changes kernel code generation as expected.

Summary by CodeRabbit

  • New Features

    • Configurable inplace-detection for storage rewrite was added (config key "tl.storage_rewrite_detect_inplace", default: false), allowing users to toggle buffer reuse behavior during storage planning.
  • Tests

    • Added test coverage verifying behavior with the inplace-detection flag enabled and disabled.
  • Documentation

    • Updated docstrings to point users to the centralized pass configuration keys for supported options.

  - Add kStorageRewriteDetectInplace constant and register the flag with PassContext so C++ code no longer hard-codes the key.
  - Wire StorageRewrite to include TileLang builtin constants and honor the new config toggle when deciding inplace reuse.
  - Document the flag across Python surfaces (PassConfigKey, JIT/autotuner docs) with usage guidance and simplified IR examples.
@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! 🚀

@LeiWang1999
Copy link
Member Author

relevant to issue #994 .

@coderabbitai
Copy link
Contributor

coderabbitai bot commented Oct 21, 2025

Walkthrough

A new configurable flag TL_STORAGE_REWRITE_DETECT_INPLACE was added across C++ and Python layers. StorageRewrite now reads this PassContext option (default false) and forwards the detect_inplace value to StoragePlanRewriter.Rewrite; tests and docs updated accordingly.

Changes

Cohort / File(s) Summary
C++ Core: builtins & registration
src/op/builtin.h, src/op/builtin.cc
Added public constant kStorageRewriteDetectInplace and registered TVM_REGISTER_PASS_CONFIG_OPTION(kStorageRewriteDetectInplace, Bool) to expose the new pass option.
C++ Core: storage rewrite pass
src/transform/storage_rewrite.cc
Read kStorageRewriteDetectInplace from PassContext (default false) and pass detect_inplace into StoragePlanRewriter::Rewrite instead of a hard-coded true; added include of ../op/builtin.h.
Python: pass config enum
tilelang/transform/pass_config.py
Added TL_STORAGE_REWRITE_DETECT_INPLACE to PassConfigKey with docstring describing semantics and default behavior.
Python: docs / API docstrings
tilelang/autotuner/param.py, tilelang/jit/__init__.py, tilelang/jit/kernel.py
Updated pass_configs docstrings to refer to tilelang.PassConfigKey for supported options instead of enumerating specific keys.
Tests
testing/python/components/test_storage_rewrite_detect_inplace.py
New test verifying storage-rewrite inplace detection behavior with and without the flag via TL storage rewrite pass config and inspecting generated kernel source.

Sequence Diagram(s)

sequenceDiagram
    participant User
    participant Pass as StorageRewritePass
    participant PC as PassContext
    participant Rewriter as StoragePlanRewriter

    User->>Pass: run StorageRewrite pass
    Pass->>PC: Get(kStorageRewriteDetectInplace)
    alt option set
        PC-->>Pass: detect_inplace = user_value
    else default
        PC-->>Pass: detect_inplace = false
    end
    Pass->>Rewriter: Rewrite(plan, detect_inplace)
    Rewriter-->>Pass: rewritten_plan
    Pass-->>User: pass complete
Loading

Estimated code review effort

🎯 3 (Moderate) | ⏱️ ~25 minutes

Suggested reviewers

  • chengyupku

Poem

🐰 I nibble flags and hop in place,

Detecting reuse with gentle grace.
False by default, but toggle me true,
Storage rewrites choose how buffers accrue.
A little config, a big bright trace.

Pre-merge checks and finishing touches

❌ Failed checks (1 warning)
Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 0.00% 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 PR title "[PassConfig] Introduce PassConfig TL_STORAGE_REWRITE_DETECT_INPLACE" directly and accurately describes the primary objective of the changeset. The title clearly identifies the main change—introducing a new PassConfig option—and specifies the exact configuration key name. The title is concise, specific, and free of vague terms or noise. A teammate scanning commit history would immediately understand that this PR adds a new configuration option to the PassConfig system.
✨ 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 dc9784c and ec42246.

📒 Files selected for processing (1)
  • testing/python/components/test_storage_rewrite_detect_inplace.py (1 hunks)
🚧 Files skipped from review as they are similar to previous changes (1)
  • testing/python/components/test_storage_rewrite_detect_inplace.py

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.

Copy link
Contributor

@coderabbitai coderabbitai bot left a comment

Choose a reason for hiding this comment

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

Actionable comments posted: 0

🧹 Nitpick comments (1)
testing/python/components/test_storage_rewrite_detect_inplace.py (1)

7-20: Consider renaming function for clarity.

The function name buggy_kernel is somewhat misleading since the kernel isn't actually buggy—it's demonstrating storage rewrite behavior. Consider using a more descriptive name like test_kernel or storage_rewrite_test_kernel.

📜 Review details

Configuration used: CodeRabbit UI

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 60e9c7e and dc9784c.

📒 Files selected for processing (8)
  • src/op/builtin.cc (1 hunks)
  • src/op/builtin.h (1 hunks)
  • src/transform/storage_rewrite.cc (3 hunks)
  • testing/python/components/test_storage_rewrite_detect_inplace.py (1 hunks)
  • tilelang/autotuner/param.py (1 hunks)
  • tilelang/jit/__init__.py (1 hunks)
  • tilelang/jit/kernel.py (1 hunks)
  • tilelang/transform/pass_config.py (1 hunks)
🧰 Additional context used
🧬 Code graph analysis (1)
testing/python/components/test_storage_rewrite_detect_inplace.py (3)
tilelang/jit/__init__.py (1)
  • jit (237-310)
tilelang/language/symbolics.py (1)
  • symbolic (26-28)
tilelang/transform/pass_config.py (1)
  • PassConfigKey (6-144)
🔇 Additional comments (11)
tilelang/autotuner/param.py (1)

40-40: LGTM! Documentation updated to reference centralized config.

The change correctly redirects users to tilelang.PassConfigKey for supported options, reducing duplication and improving maintainability.

tilelang/jit/__init__.py (1)

62-62: LGTM! Documentation reference updated.

Correctly directs users to tilelang.transform.PassConfigKey for supported options, maintaining consistency with the broader PR changes.

tilelang/jit/kernel.py (1)

74-74: LGTM! Docstring updated consistently.

The change aligns with the PR's goal of centralizing pass config documentation in tilelang.PassConfigKey.

src/op/builtin.cc (1)

36-36: LGTM! Pass config option registered correctly.

The new kStorageRewriteDetectInplace boolean option is registered following the established pattern and matches the constant defined in src/op/builtin.h.

tilelang/transform/pass_config.py (1)

72-110: LGTM! Comprehensive config key documentation.

The new TL_STORAGE_REWRITE_DETECT_INPLACE enum member is well-documented with:

  • Clear explanation of default behavior (False)
  • Concrete IR examples showing aliasing differences
  • Proper usage example with PassContext

The string value correctly matches the C++ constant kStorageRewriteDetectInplace defined in src/op/builtin.h.

src/op/builtin.h (1)

51-52: LGTM! Config constant defined correctly.

The kStorageRewriteDetectInplace constant is properly defined with the value "tl.storage_rewrite_detect_inplace", matching the Python PassConfigKey.TL_STORAGE_REWRITE_DETECT_INPLACE enum value.

testing/python/components/test_storage_rewrite_detect_inplace.py (1)

52-57: Test validates inplace detection correctly.

The test appropriately verifies that:

  • Without the flag: "read = (read * 2);" does not appear (separate allocations)
  • With the flag: "read = (read * 2);" appears (inplace reuse)

The string-based assertion is acceptable for this code generation test, though it's somewhat fragile to code generation changes.

src/transform/storage_rewrite.cc (4)

41-41: Include added for new config constant.

The #include "../op/builtin.h" is necessary to access kStorageRewriteDetectInplace used at line 1919.


470-472: Rewrite signature updated to accept inplace flag.

The detect_inplace parameter is now passed to Rewrite() instead of being hardcoded. This enables runtime configuration via PassContext.


1918-1919: Default value conservative and correct.

The config defaults to Bool(false), which keeps distinct temporaries and avoids implicit aliasing. This is a safe default that maintains backward compatibility for users who don't explicitly enable the flag.


1945-1947: Config flag properly propagated to rewriter.

The detect_inplace flag read from PassContext is correctly passed to StoragePlanRewriter().Rewrite(), enabling user control over inplace buffer detection.

@LeiWang1999 LeiWang1999 merged commit cdc67fc into tile-ai:main Oct 21, 2025
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