JosephTheOctonaut commented on code in PR #15684:
URL: https://github.com/apache/tvm/pull/15684#discussion_r1318944190
##########
include/tvm/tir/builtin.h:
##########
@@ -681,43 +680,51 @@ TVM_DLL const Op& ptx_wait_group();
/*!
* \brief tvm intrinsics for ptx async copy barrier using
cp.async.mbarrier.arrive
*
- * ptx_cp_async_barrier(Var barrier_ptr, Expr barrier_offset)
+ * ptx_cp_async_barrier(int barrier_id)
*
*/
TVM_DLL const Op& ptx_cp_async_barrier();
/*!
* \brief tvm intrinsics for ptx barrier initialization of thread count using
mbarrier.init
*
- * ptx_init_barrier_thread_count(Var barrier_ptr, Expr barrier_offset, int
thread_count)
+ * ptx_init_barrier_thread_count(int barrier_id, int thread_count)
*
*/
TVM_DLL const Op& ptx_init_barrier_thread_count();
/*!
* \brief tvm intrinsics for ptx barrier arrival using mbarrier.arrive
*
- * ptx_arrive_barrier(Var barrier_ptr, Expr barrier_offset)
+ * ptx_arrive_barrier(int barrier_id)
*
*/
TVM_DLL const Op& ptx_arrive_barrier();
/*!
* \brief tvm intrinsic for ptx barrier arrival with expect tx using
mbarrier.arrive.expect_tx
*
- * ptx_arrive_barrier_expect_tx(Var barrier_ptr, Expr barrier_offset, int
byte_count)
+ * ptx_arrive_barrier_expect_tx(int barrier_id, int byte_count)
Review Comment:
Right now this is the only way to add expected transaction counts to a
barrier, correct? Do we need a mechanism aside from `arrive` for adding to the
count? This way makes sense when a thread issues a bulk async copy and arrives
right away, but if the arrival isn't right away, the transaction count isn't
added to the barrier until later.
##########
src/target/source/codegen_cuda.h:
##########
@@ -109,6 +109,13 @@ class CodeGenCUDA final : public CodeGenC {
// Op attribute map
OpAttrMap<bool> op_need_warp_shuffle_ =
Op::GetAttrMap<bool>("cuda.need_warp_shuffle");
+ // The name of the barrier array in shared memory
+ const std::string barrier_name_ = "barrier";
Review Comment:
This is definitely fine for now, but thinking ahead a little bit, will this
cause issues if we try to have multiple Blocks in the same Thread Block
Cluster, each with their own "barrier"? Depending on scoping it might be fine.
##########
src/target/source/codegen_cuda.h:
##########
@@ -109,6 +109,13 @@ class CodeGenCUDA final : public CodeGenC {
// Op attribute map
OpAttrMap<bool> op_need_warp_shuffle_ =
Op::GetAttrMap<bool>("cuda.need_warp_shuffle");
+ // The name of the barrier array in shared memory
+ const std::string barrier_name_ = "barrier";
+ // The size of the barrier array in shared memory
+ int barrier_count_ = -1;
+ // The alignment of the barrier array in shared memory
+ const int barrier_alignment_bytes_ = 64;
Review Comment:
Why is this 64 instead of 16?
--
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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
