Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
20 changes: 7 additions & 13 deletions src/transform/storage_rewrite.cc
Original file line number Diff line number Diff line change
Expand Up @@ -1109,6 +1109,9 @@ class StoragePlanRewriter : public StmtExprMutator {
// when not divided, no reuse, eg, float4 vs float3
if (e->bits_offset % op_elem_bits != 0)
continue;
// must check element type to avoid type mismatch in codegen
if (e->elem_type != op->dtype.element_of())
continue;
if (reuse_require_exact_matched_dtype && e->elem_type != op->dtype) {
continue;
}
Expand Down Expand Up @@ -1962,21 +1965,12 @@ Pass StorageRewrite() {
ctx->GetConfig<Bool>(kStorageRewriteDetectInplace, Bool(false)).value();
bool enable_reuse = true;
bool reuse_require_exact_matched_dtype = false;
bool merge_static_smem =
ctx->GetConfig<Bool>("tir.merge_static_smem", Bool(false)).value();
AllocateCollector collector;
collector(f->body);
bool has_dynamic = collector.dyn_shmem_allocs_.size() > 1;
if (has_dynamic || merge_static_smem) {
// For IRModule utilizing dynamic shared memory, reuse is not enabled
// Because dynamic doesn't require maintaining the readability and
// it benefits from a more optimized allocation strategy through the
// Pass `MergeSharedMemoryAllocations`.
// When `merge_static_smem` is true, we will reuse and merge shared
// memory in a dedicated pass `MergeSharedMemoryAllocations`.
// And so we don't enable reuse in this pass.
enable_reuse = false;
}
// Always disable reuse currently, for shared memory reuse we depend on
// MergeSharedMemoryAllocations pass, for register reuse we depend on nvcc
// or other compiler its self.
enable_reuse = false;

Optional<Target> target = f->GetAttr<Target>("target");
if (target.defined() && (target.value()->kind->name == "vulkan" ||
Expand Down
26 changes: 26 additions & 0 deletions testing/python/issue/test_tilelang_issue_1678.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,26 @@
# ruff: noqa
import tilelang
import tilelang.testing
import tilelang.language as T


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()


if __name__ == "__main__":
tilelang.testing.main()
Loading