[Feature] Add DumpIR PassConfig in TileLang side

[SiriusNEO]

Feat: #1896

Previously TVM has a pass instrument to dump all IRs between passes, but it's not exposed in TileLang side and it's hard for new users to find this option. This PR explictly adds this feature as a pass config in TileLang.

Summary by CodeRabbit

• New Features
  • Added optional IR dumping during kernel compilation with configurable output directory.
  • New pass configuration options to enable IR dumps and specify dump location.
  • Added configuration to toggle out-of-bound warnings in TileLang.
  • PassContext now supports optional IR instrumentation during compilation.

[github-actions]

👋 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]

Caution

Review failed

Pull request was closed or merged during review

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Pro

Run ID: 6cfd2d5a-5328-4cbb-a887-5b332f7dc077

📥 Commits

Reviewing files that changed from the base of the PR and between 3574b6f and a779386.

📒 Files selected for processing (1)

• tilelang/transform/pass_config.py

---

📝 Walkthrough

Walkthrough

Adds TileLang IR dumping controls: two new pass config keys and constants (enable flag and dump path) in C++ and Python, and optionally attaches a DumpIR instrument to the PassContext during kernel JIT compilation when enabled.

Changes

Cohort / File(s)	Summary
C++ headers & registration src/op/builtin.h, src/op/builtin.cc	Added public constants kDisableOutOfBoundWarning, kEnableDumpIR, kDumpIRDir and registered kEnableDumpIR (Bool) and kDumpIRDir (ffi::String) as pass config options.
Python pass-config keys tilelang/transform/pass_config.py	Added TL_ENABLE_DUMP_IR and TL_DUMP_IR_DIR keys to PassConfigKey under output debugging options.
JIT kernel instrumentation tilelang/jit/kernel.py	When TL_ENABLE_DUMP_IR is set, creates/attaches a DumpIR instrument (using TL_DUMP_IR_DIR or default) to the PassContext passed to TVM compilation.

Sequence Diagram(s)

sequenceDiagram
  participant Caller
  participant KernelJIT
  participant PassContext
  participant DumpIR
  participant TVMCompiler

  Caller->>KernelJIT: request kernel compile (pass_configs)
  KernelJIT->>PassContext: create with opt_level=3, config=pass_configs
  alt TL_ENABLE_DUMP_IR set
    KernelJIT->>DumpIR: instantiate DumpIR(dump_dir)
    KernelJIT->>PassContext: attach DumpIR instrument
  end
  KernelJIT->>TVMCompiler: run compilation with PassContext
  TVMCompiler->>DumpIR: Dump intermediate IR files (if instrument attached)
  TVMCompiler-->>KernelJIT: return compiled kernel
  KernelJIT-->>Caller: return result

Loading

Estimated code review effort

🎯 3 (Moderate) | ⏱️ ~20 minutes

Possibly related PRs

• [Bugfix] Convey compile_flags to ffi compilation path with pass_configs #1434: Modifies the same pass-config plumbing and PassContext instrumentation patterns (high overlap with these changes).

Suggested reviewers

• chengyupku
• LeiWang1999

Poem

🐰 I hopped through headers, keys in paw,

Flags and paths I tucked in raw,
When compile runs and passes fly,
My little dumps will wink and sigh,
Debugging carrots — oh my, oh my! 🥕

🚥 Pre-merge checks | ✅ 2 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 50.00% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ 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 accurately describes the main change: adding DumpIR PassConfig support to the TileLang side, which is reflected in changes across header files, build configurations, and kernel compilation logic.

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests
• Post copyable unit tests in a comment

---

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

• X
• Mastodon
• Reddit
• LinkedIn

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

[coderabbitai]

Actionable comments posted: 2

🧹 Nitpick comments (1)

src/op/builtin.h (1)

136-138: Consider documenting kDumpIRDir explicitly.

kEnableDumpIR has dedicated API docs, but kDumpIRDir is introduced without a matching brief. Adding a short comment keeps this header consistent and easier to scan.

🤖 Prompt for AI Agents

Verify each finding against the current code and only fix it if needed.

In `@src/op/builtin.h` around lines 136 - 138, Add a short API doc comment for
kDumpIRDir to match kEnableDumpIR’s style: immediately above the static
constexpr const char *kDumpIRDir declaration, add a brief one-line comment
indicating that this option specifies the directory/path where the IR dump files
are written (e.g., "Directory/path to write dumped IR files") so the header is
consistent and easier to scan; ensure formatting and phrasing follow nearby
kEnableDumpIR comments.

🤖 Prompt for all review comments with AI agents

Verify each finding against the current code and only fix it if needed.

Inline comments:
In `@tilelang/jit/kernel.py`:
- Around line 241-247: The wrapper path in parse_source_information constructs a
tvm.transform.PassContext without passing instruments, so DumpIR (activated via
PassConfigKey.TL_ENABLE_DUMP_IR) is not applied there; update
parse_source_information in tilelang/jit/adapter/wrapper.py to mirror kernel.py
by building a pass_instruments list (respecting PassConfigKey.TL_ENABLE_DUMP_IR
and TL_DUMP_IR_DIR default) and pass it into tvm.transform.PassContext via
instruments=pass_instruments when creating the context around compilation so
DumpIR is consistently injected across both paths.

In `@tilelang/transform/pass_config.py`:
- Around line 236-237: The docstring for TL_DUMP_IR_DIR ("TL_DUMP_IR_DIR") is
inconsistent with the runtime default used in the JIT (tilelang.jit.kernel);
update the TL_DUMP_IR_DIR docstring to state "Default: ./dump_ir" to match
runtime behavior, and while here verify the JIT code reads/uses TL_DUMP_IR_DIR
(not a hardcoded "./dump_ir") so both the constant and runtime default remain
consistent.

---

Nitpick comments:
In `@src/op/builtin.h`:
- Around line 136-138: Add a short API doc comment for kDumpIRDir to match
kEnableDumpIR’s style: immediately above the static constexpr const char
*kDumpIRDir declaration, add a brief one-line comment indicating that this
option specifies the directory/path where the IR dump files are written (e.g.,
"Directory/path to write dumped IR files") so the header is consistent and
easier to scan; ensure formatting and phrasing follow nearby kEnableDumpIR
comments.

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Pro

Run ID: 189e2ece-e810-4cea-90fa-126f544d98a0

📥 Commits

Reviewing files that changed from the base of the PR and between 00267e4 and 3574b6f.

📒 Files selected for processing (4)

• src/op/builtin.cc
• src/op/builtin.h
• tilelang/jit/kernel.py
• tilelang/transform/pass_config.py