Signaling operations

Signaling operations#

ROCSHMEM_PUTMEM_SIGNAL#

__device__ void rocshmem_putmem_signal(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_putmem_signal_wave(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_putmem_signal_wg(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_putmem_signal_nbi(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_putmem_signal_nbi_wave(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_putmem_signal_nbi_wg(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_putmem_signal(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_putmem_signal_wave(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_putmem_signal_wg(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_putmem_signal_nbi(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_putmem_signal_nbi_wave(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_putmem_signal_nbi_wg(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
Parameters:

  • ctx – Context with which to perform this operation.

  • dest – Destination address. Must be an address on the symmetric heap.

  • source – Source address. Must be an address on the symmetric heap.

  • nelems – The number of bytes to transfer.

  • sig_addr – Signal address. Must be an address on the symmetric heap.

  • signal – Signal value.

  • sig_op – Atomic operation to apply the signal value.

  • pe – PE of the remote process.

Returns:

None.

Description: This function writes contiguous data of nelems bytes from source on the calling PE to dest at pe, then applies sig_op at sig_addr with the signal value. Valid sig_op values are listed in SIGNAL_OPERATORS.

ROCSHMEM_PUT_SIGNAL#

__device__ void rocshmem_TYPENAME_put_signal(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_TYPENAME_put_signal_wave(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_TYPENAME_put_signal_wg(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_TYPENAME_put_signal_nbi(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_TYPENAME_put_signal_nbi_wave(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_TYPENAME_put_signal_nbi_wg(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_TYPENAME_put_signal(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_TYPENAME_put_signal_wave(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_TYPENAME_put_signal_wg(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_TYPENAME_put_signal_nbi(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_TYPENAME_put_signal_nbi_wave(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
__device__ void rocshmem_ctx_TYPENAME_put_signal_nbi_wg(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)#
Parameters:

  • ctx – Context with which to perform this operation.

  • dest – Destination address. Must be an address on the symmetric heap.

  • source – Source address. Must be an address on the symmetric heap.

  • nelems – The number of elements of size TYPE to transfer.

  • sig_addr – Signal address. Must be an address on the symmetric heap.

  • signal – Signal value.

  • sig_op – Atomic operation to apply the signal value.

  • pe – PE of the remote process.

Returns:

None.

Description: This function writes contiguous data of nelems elements of TYPE from source on the calling PE to dest at pe, then applies sig_op at sig_addr with the signal value. Valid sig_op values are listed in SIGNAL_OPERATORS. Valid TYPENAME and TYPE values are listed in RMA Data Types.

ROCSHMEM_PUTMEM_SIGNAL_ON_STREAM#

__host__ void rocshmem_putmem_signal_on_stream(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe, hipStream_t stream)#
Parameters:

  • dest – Destination address on the remote PE. Must be an address on the symmetric heap.

  • source – Source address on the local PE. Must be an address on the symmetric heap.

  • nelems – Size of the transfer in bytes.

  • sig_addr – Address of signal variable on the remote PE. Must be an address on the symmetric heap.

  • signal – Signal value to be written.

  • sig_op – Signal operation (ROCSHMEM_SIGNAL_SET or ROCSHMEM_SIGNAL_ADD).

  • pe – PE number of the remote PE.

  • stream – HIP stream on which to enqueue the operation.

Returns:

None.

Description: This routine enqueues a put-with-signal operation on a HIP stream. The function writes contiguous data of nelems bytes from source on the calling PE to dest at pe, then applies sig_op at sig_addr with the signal value. The operation is enqueued on the specified stream and will execute asynchronously. The caller must synchronize the stream (e.g., using hipStreamSynchronize) to ensure completion.

Valid sig_op values are listed in SIGNAL_OPERATORS.

ROCSHMEM_SIGNAL_FETCH#

__device__ uint64_t rocshmem_signal_fetch(const uint64_t *sig_addr)#
__device__ uint64_t rocshmem_signal_fetch_wg(const uint64_t *sig_addr)#
__device__ uint64_t rocshmem_signal_fetch_wave(const uint64_t *sig_addr)#
Parameters:

sig_addr – Signal address. Must be an address on the symmetric heap.

Returns:

Value at sig_addr.

Description: This function atomically fetches the value stored at sig_addr.

Signal operators#

Table 5 Signal Operators#

Value

Description

ROCSHMEM_SIGNAL_SET

The signaling operation routines will atomically set the signal value at sig_addr.

ROCSHMEM_SIGNAL_ADD

The signaling operation routines will atomically add the signal value at sig_addr.
Contents