Skip to content

Conversation

@LeiWang1999
Copy link
Member

@LeiWang1999 LeiWang1999 commented Jan 19, 2026

Fix for issue #1678

For shared memory reuse we depend on MergeSharedMemoryAllocations pass, for register reuse we depend on nvcc or other vendors' compilers itself.

Summary by CodeRabbit

  • Bug Fixes

    • Enforced element-type consistency for buffer reuse and disabled local reuse in this pass to rely on downstream reuse mechanisms.
  • Tests

    • Added a regression test for issue 1678 validating TileLang kernel construction, configuration, and emitted kernel source.

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

@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 Jan 19, 2026

Caution

Review failed

The pull request is closed.

📝 Walkthrough

Walkthrough

Enforces element-type checks when considering reusable StorageEntry objects and unconditionally disables storage reuse in the StorageRewrite pass (delegating reuse to downstream passes/compilers). Adds a new TileLang Python test exercising JIT kernel construction and kernel source retrieval.

Changes

Cohort / File(s) Summary
Storage Reuse Logic
src/transform/storage_rewrite.cc
Added element-type compatibility checks comparing e->elem_type with op->dtype.element_of() in both larger and smaller buffer search paths. Unconditionally set enable_reuse = false and removed prior conditional logic tied to dynamic shared memory / merge_static_smem; added comment clarifying delegation to MergeSharedMemoryAllocations and compiler.
TileLang Test
testing/python/issue/test_tilelang_issue_1678.py
New integration test test_issue_1678 that defines a JIT-wrapped TileLang kernel, constructs a primitive kernel with local allocations and kernel config, invokes it, and prints the emitted kernel source.

Sequence Diagram(s)

mermaid
sequenceDiagram
participant StorageRewrite as StorageRewrite (pass)
participant MergeSM as MergeSharedMemoryAllocations (pass)
participant Compiler as Backend Compiler
participant IR as IR/Allocator Decisions

StorageRewrite->>IR: Analyze allocations and attempt reuse
note right of StorageRewrite: Previously: reuse if size/type allowed
StorageRewrite-->>IR: Disable reuse unconditionally (enable_reuse = false)
StorageRewrite->>MergeSM: Delegate shared-memory reuse considerations
MergeSM->>IR: Merge/shared-memory-specific allocation transformations
IR->>Compiler: Emit lowered allocations and register usage
Compiler->>IR: Apply register-level reuse / backend-specific optimizations

Estimated code review effort

🎯 4 (Complex) | ⏱️ ~45 minutes

Possibly related PRs

Suggested reviewers

  • kurisu6912

Poem

🐰
I hopped through buffers, checked each type,
No mismatched bytes to start a gripe.
I handed reuse to smarter kin,
And watched the tests begin to grin! 🥕

🚥 Pre-merge checks | ✅ 2 | ❌ 1
❌ 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%. 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 '[Bugfix][Refactor] Always disable light storage reuse' clearly and specifically describes the main change—unconditionally disabling storage reuse in the StorageRewrite pass, which is the primary modification in src/transform/storage_rewrite.cc.

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

✨ Finishing touches
  • 📝 Generate docstrings

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: 1

🤖 Fix all issues with AI agents
In `@testing/python/issue/test_tilelang_issue_1678.py`:
- Around line 7-23: Replace the sole print in test_issue_1678 with an assertion
that validates the generated kernel source: call kernel = qwq(); src =
kernel.get_kernel_source() and assert that src is non-empty and contains
relevant identifiers (e.g., "amax_local" and "tmp_row" or "max") to ensure the
qwq/qwq_kernel JIT produced the expected code rather than just printing it.

Comment on lines 7 to 23
def test_issue_1678():
@tilelang.jit
def qwq():
@T.prim_func
def qwq_kernel():
with T.Kernel(4096, 1, threads=1) as (pid_y, pid_x):
i = T.alloc_var("int32")
i = 1
tmp_row = T.alloc_local((4,), "float32")
amax_local = T.alloc_var("float32")
j = 0
amax_local = T.max(amax_local, tmp_row[j])

return qwq_kernel

kernel = qwq()
print(kernel.get_kernel_source())
Copy link
Contributor

Choose a reason for hiding this comment

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

⚠️ Potential issue | 🟡 Minor

🧩 Analysis chain

🏁 Script executed:

# First, find the test file and verify it exists
find . -name "test_tilelang_issue_1678.py" -type f

Repository: tile-ai/tilelang

Length of output: 110


🏁 Script executed:

# Also list test files in the issue directory to understand the structure
fd "test_.*\.py" testing/python/issue --max-results 20

Repository: tile-ai/tilelang

Length of output: 941


🏁 Script executed:

# Search for similar tests that validate kernel source
rg "get_kernel_source" -A 5 -B 2 testing/python --type py | head -100

Repository: tile-ai/tilelang

Length of output: 8596


🏁 Script executed:

cat -n testing/python/issue/test_tilelang_issue_1678.py

Repository: tile-ai/tilelang

Length of output: 896


🏁 Script executed:

# Check a couple of similar issue tests to see their patterns
rg "def test_issue" -A 20 testing/python/issue/test_tilelang_issue_1001.py | head -60

Repository: tile-ai/tilelang

Length of output: 42


🏁 Script executed:

# Another example
rg "def test_issue" -A 20 testing/python/issue/test_tilelang_issue_1026.py | head -60

Repository: tile-ai/tilelang

Length of output: 233


Add an assertion to validate kernel source generation.

The test currently only prints the kernel source without any assertions. Similar issue tests in the codebase (e.g., test_tilelang_issue_1026.py) validate behavior by asserting on generated source content. Add an assertion to ensure the kernel source is generated correctly and prevent regressions:

     kernel = qwq()
-    print(kernel.get_kernel_source())
+    src = kernel.get_kernel_source()
+    assert src and "qwq_kernel" in src
+    print(src)
📝 Committable suggestion

‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.

Suggested change
def test_issue_1678():
@tilelang.jit
def qwq():
@T.prim_func
def qwq_kernel():
with T.Kernel(4096, 1, threads=1) as (pid_y, pid_x):
i = T.alloc_var("int32")
i = 1
tmp_row = T.alloc_local((4,), "float32")
amax_local = T.alloc_var("float32")
j = 0
amax_local = T.max(amax_local, tmp_row[j])
return qwq_kernel
kernel = qwq()
print(kernel.get_kernel_source())
def test_issue_1678():
`@tilelang.jit`
def qwq():
`@T.prim_func`
def qwq_kernel():
with T.Kernel(4096, 1, threads=1) as (pid_y, pid_x):
i = T.alloc_var("int32")
i = 1
tmp_row = T.alloc_local((4,), "float32")
amax_local = T.alloc_var("float32")
j = 0
amax_local = T.max(amax_local, tmp_row[j])
return qwq_kernel
kernel = qwq()
src = kernel.get_kernel_source()
assert src and "qwq_kernel" in src
print(src)
🤖 Prompt for AI Agents
In `@testing/python/issue/test_tilelang_issue_1678.py` around lines 7 - 23,
Replace the sole print in test_issue_1678 with an assertion that validates the
generated kernel source: call kernel = qwq(); src = kernel.get_kernel_source()
and assert that src is non-empty and contains relevant identifiers (e.g.,
"amax_local" and "tmp_row" or "max") to ensure the qwq/qwq_kernel JIT produced
the expected code rather than just printing it.

Remove print statement for kernel source in test.
@LeiWang1999 LeiWang1999 merged commit c9cec62 into tile-ai:main Jan 19, 2026
1 of 3 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