JackWeiw commented on code in PR #16454:
URL: https://github.com/apache/tvm/pull/16454#discussion_r1467213250
##########
tests/python/tir-transform/test_tir_transform_thread_sync.py:
##########
@@ -160,8 +160,49 @@ def expected(A: T.Buffer((4, 4), "float32"), E:
T.Buffer((4, 4), "float32")):
tvm.ir.assert_structural_equal(mod["main"], expected)
+@tvm.testing.requires_cuda
+def test_sync_let_stmt():
+ @T.prim_func(private=True)
+ def func(A: T.Buffer((16 * 512), "float32")):
+ blockIdx_x = T.launch_thread("blockIdx.x", 16)
+ A_shared = T.allocate([512], "float32", "shared")
+ in_thread_A_temp = T.allocate([1], "float32", "local")
+ cross_thread_A_temp = T.allocate([1], "float32", "local")
+ threadIdx_x = T.launch_thread("threadIdx.x", 128)
+ A_shared_1 = T.Buffer((512,), data=A_shared, scope="shared")
+ for ax0 in range(512):
+ A_shared_1[ax0] = A[blockIdx_x * 512 + ax0]
+ in_thread_A_temp_1 = T.Buffer((1,), data=in_thread_A_temp,
scope="local")
+ in_thread_A_temp_1[0] = T.float32(0)
+ with T.LetStmt(in_thread_A_temp_1[0] + A_shared_1[threadIdx_x]) as
A_temp:
Review Comment:
in case here, before take change, `StorageAccessVisitor` will call
`VisitExpr_(const BufferLoadNode* op)` directly (not called by father call like
`VisitStmt_(const BufferStoreNode* op)` traverse child), so
`ICHECK(allow_append_) << op << " " << scope.to_string();` will return error.
if we take `VisitStmt_(const LetStmtNode* op)` and traverse child of
`LetStmtNode`
(case here it will traverse `BufferLoadNode` A_shared which will increase
curr_stmt_.access by 1, and then visit `BufferStoreNode` `in_thread_A_temp` ,
`ICHECK_EQ(curr_stmt_.access.size(), 0U);` will return error)
Do you have any insights on how to solve this problem? @vinx13
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscr...@tvm.apache.org
For queries about this service, please contact Infrastructure at:
us...@infra.apache.org
