Skip to content

Commit 77b9d08

Browse files
authored
[Bugfix] Use access_ptr("r") instead of access_ptr("w") for correct pipeline analysis (#983)
* remove debug print * pipeline fix * use the correct buffer access scope
1 parent 117f2b8 commit 77b9d08

File tree

1 file changed

+2
-2
lines changed

1 file changed

+2
-2
lines changed

tilelang/intrinsics/wgmma_macro_generator.py

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -245,9 +245,9 @@ def _warp_mma(A_buf, B_buf, C_local_buf):
245245
# TODO(lei): inject warpgroup_fence_operand for C_local_buf
246246
desc_a = T.alloc_descriptor()
247247
desc_b = T.alloc_descriptor()
248-
T.initialize_descriptor(desc_a, A_buf.access_ptr("w"), a_swizzle_mode,
248+
T.initialize_descriptor(desc_a, A_buf.access_ptr("r"), a_swizzle_mode,
249249
int(a_leading_byte_offset >> 4), int(a_stride_byte_offset >> 4))
250-
T.initialize_descriptor(desc_b, B_buf.access_ptr("w"), b_swizzle_mode,
250+
T.initialize_descriptor(desc_b, B_buf.access_ptr("r"), b_swizzle_mode,
251251
int(b_leading_byte_offset >> 4), int(b_stride_byte_offset >> 4))
252252
T.warpgroup_arrive()
253253
for ki in T.serial(0, (k_dim // micro_size_k)):

0 commit comments

Comments
 (0)