tvm_tirx.backend.cuda.op.nvshmem_wait_until

tvm_tirx.backend.cuda.op.nvshmem_wait_until#

tvm_tirx.backend.cuda.op.nvshmem_wait_until(ivar, cmp, cmp_value, type='uint64_t', *, ty=None, loc=UNKNOWN_LOC)#

TVM intrinsic to call nvshmem_wait_until()

Parameters:

  • ivar (Expr) – The pointer to the symmetric address of a remotely accessible data object, must be TYPE*.

  • cmp (str) – The compare operator that compares ivar with cmp_value.

  • cmp_value (TYPE) – The value to be compared with ivar.

  • type (str) – The TYPE of ivar and cmp_value.

  • loc (Location)

Returns:

call (Expr) – The call expression.