Skip to content

Commit e4bd948

Browse files
committed
add name check
1 parent 551f7b8 commit e4bd948

File tree

2 files changed

+2
-1
lines changed

2 files changed

+2
-1
lines changed

src/transform/lower_shared_barrier.cc

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -69,6 +69,8 @@ class SharedBarrierRewriter : public StmtExprMutator {
6969
ICHECK(thread_var_.defined()) << "thread_var_ is not defined";
7070

7171
for (auto buffer : barrier_buffers) {
72+
ICHECK(buffer->name != "mbarrier")
73+
<< "Shared barrier's name 'mbarrier' is reserved";
7274
buffer_data_to_buffer_.Set(buffer->data, buffer);
7375
}
7476

tilelang/tileop/gemm/gemm_wgmma.py

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -35,7 +35,6 @@ def infer_shared_layout(self, continuity: int) -> Callable[[tir.Buffer], Layout]
3535
See: https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TENSOR__MEMORY.html
3636
"""
3737
vectorized_size = 128 // self.in_dtype.bits
38-
print(f"continuity: {continuity}, vectorized_size: {vectorized_size}")
3938
if continuity % (vectorized_size * 8) == 0:
4039
return make_full_bank_swizzled_layout
4140
elif continuity % (vectorized_size * 4) == 0:

0 commit comments

Comments
 (0)