adstraw commented on code in PR #15684: URL: https://github.com/apache/tvm/pull/15684#discussion_r1319144179
########## 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: Good point. I am tempted to leave this as-is for now and solve it when the need arises. -- 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]
