tvm_tirx.backend.cuda.op.nvshmem_signal_op

tvm_tirx.backend.cuda.op.nvshmem_signal_op#

tvm_tirx.backend.cuda.op.nvshmem_signal_op(sig_addr, signal, sig_op, pe, *, ty=None, loc=UNKNOWN_LOC)#

TVM intrinsic to call nvshmem_signal_op()

Parameters:

  • sig_addr (Expr) – The pointer to the symmetric address of the signal word to be updated, must be uint64_t*.

  • signal (uint64_t) – The value used to update sig_addr.

  • sig_op (str) – Operation used to update sig_addr with signal, typical sig_op values are “set” and “add”.

  • pe (int) – The PE number of the remote PE.

  • loc (Location)

Returns:

call (Expr) – The call expression.