masahi commented on code in PR #15656: URL: https://github.com/apache/tvm/pull/15656#discussion_r1315592599
########## include/tvm/tir/builtin.h: ########## @@ -645,14 +645,29 @@ TVM_DLL const Op& ptx_mma_sp(); TVM_DLL const Op& ptx_ldmatrix(); /*! - * \brief tvm intrinsics for ptx async copy from global to shared memory - * - * void ptx_cp_async(Var shared_ptr, Expr shared_offset, Var global_ptr, Expr global_offset, size_t - * bytes); + * \brief tvm intrinsics for ptx async copy from global to shared memory using cp.async * + * void ptx_cp_async(Var shared_ptr, + * Expr shared_offset, + * Var global_ptr, + * Expr global_offset, + * size_t bytes); */ TVM_DLL const Op& ptx_cp_async(); +/*! + * \brief tvm intrinsics for ptx async copy from global to shared memory using cp.async.bulk + * + * void ptx_cp_async(Var shared_ptr, + * Expr shared_offset, + * Var global_ptr, + * Expr global_offset, + * size_t bytes, + * string barrier_arr, + * int barrier_id); + */ +TVM_DLL const Op& ptx_cp_async_bulk(); Review Comment: I just realized that a barrier is specified by a string / name in the intrin signature rather than a pointer variable. Is there a reason a barrier cannot be a `Var`, like src / dst pointers of the `cp.async` intrin? (Sorry I should have noticed this in the previous PR). -- 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]
