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_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 6 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`.