tvm_tirx.backend.cuda.op.nvshmem_putmem_signal_nbi

tvm_tirx.backend.cuda.op.nvshmem_putmem_signal_nbi#

tvm_tirx.backend.cuda.op.nvshmem_putmem_signal_nbi(dst, src, nelems, sig_addr, signal, sig_op, pe, *, ty=None, loc=UNKNOWN_LOC)#

TVM intrinsic to call nvshmem_putmem_signal_nbi()

Parameters:

  • dst (Expr) – The pointer to the symmetric address of the data object to be updated on the remote PE.

  • src (Expr) – The pointer to the symmetric address or host/device address of data object containing the data to be copied.

  • nelems (int) – The number of bytes to put per thread.

  • sig_addr (Expr) – The pointer to the symmetric address of the signal data object to be updated on the remote PE as a signal, must be uint64_t*.

  • signal (uint64_t) – The unsigned 64-bit value that is used for updating the remote sig_addr signal data object.

  • sig_op (str) – Signal operator that represents the type of update to be performed on the remote sig_addr signal data object.

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

  • loc (Location)

Returns:

call (Expr) – The call expression.