Point-to-point synchronization routines

Point-to-point synchronization routines#

ROCSHMEM_WAIT_UNTIL#

__device__ void rocshmem_TYPENAME_wait_until(TYPE *ivars, int cmp, TYPE val)#
Parameters:

  • ivars – Pointer to memory on the symmetric heap to wait for.

  • cmp – Operation for the comparison.

  • val – Value to compare the memory at ivars to.

Returns:

None.

Description: This routine blocks the caller until the condition (*ivars cmp val) is true.

Valid cmp values are listed in Supported comparisons.

Valid TYPENAME and TYPE values are listed in Standard AMO Data Types.

ROCSHMEM_WAIT_UNTIL_ALL#

__device__ void rocshmem_TYPENAME_wait_until_all(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE val)#
Parameters:

  • ivars – Pointer to memory on the symmetric heap to wait for.

  • nelems – Number of elements in the ivars array.

  • status – Array of length nelems to exclude elements from the wait.

  • cmp – Operation for the comparison.

  • val – Value to compare.

Returns:

None.

Description: This routine blocks the caller until the condition (ivars[i] cmp val) is true for all ivars.

Valid cmp values are listed in Supported comparisons.

Valid TYPENAME and TYPE values are listed in Standard AMO Data Types.

ROCSHMEM_WAIT_UNTIL_ANY#

__device__ size_t rocshmem_TYPENAME_wait_until_any(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE val)#
Parameters:

  • ivars – Pointer to memory on the symmetric heap to wait for.

  • nelems – Number of elements in the ivars array.

  • status – Array of length nelems to exclude elements from the wait.

  • cmp – Operation for the comparison.

  • val – Value to compare.

Returns:

The index of an element in the ivars array that satisfies the wait condition. If the wait set is empty, this routine returns SIZE_MAX.

Description: This routine blocks the caller until any of the condition (ivars[i] cmp val) is true.

Valid cmp values are listed in Supported comparisons.

Valid TYPENAME and TYPE values are listed in Standard AMO Data Types.

ROCSHMEM_WAIT_UNTIL_SOME#

__device__ size_t rocshmem_TYPENAME_wait_until_some(TYPE *ivars, size_t nelems, size_t *indices, const int *status, int cmp, TYPE val)#
Parameters:

  • ivars – Pointer to memory on the symmetric heap to wait for.

  • nelems – Number of elements in the ivars array.

  • indices – List of indices with a length of at least nelems.

  • status – Array of length nelems to exclude elements from the wait.

  • cmp – Operation for the comparison.

  • val – Value to compare.

Returns:

The number of indices returned in the indices array. If the wait set is empty, this routine returns 0.

Description: This routine blocks the caller until any of the conditions (ivars[i] cmp val) is true.

Valid cmp values are listed in Supported comparisons.

Valid TYPENAME and TYPE values are listed in Standard AMO Data Types.

ROCSHMEM_TEST#

__device__ int rocshmem_TYPENAME_test(TYPE *ivars, int cmp, TYPE val)#
Parameters:

  • ivars – Pointer to memory on the symmetric heap to wait for.

  • cmp – Operation for the comparison.

  • val – Value to compare the memory at ivars to.

Returns:

1 if the evaluation is true. 0 otherwise.

Description: This routine tests if the condition (*ivars cmp val) is true.

ROCSHMEM_SIGNAL_WAIT_UNTIL_ON_STREAM#

__host__ void rocshmem_signal_wait_until_on_stream(uint64_t *sig_addr, int cmp, uint64_t cmp_value, hipStream_t stream)#
Parameters:

  • sig_addr – Address of the signal variable on the symmetric heap.

  • cmp – Comparison operator (e.g., ROCSHMEM_CMP_EQ, ROCSHMEM_CMP_GE, etc.).

  • cmp_value – Value to compare against.

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

Returns:

None.

Description: This routine enqueues a wait operation on a HIP stream without blocking the calling thread. Work enqueued on stream after this call does not start until the signal variable at sig_addr satisfies the comparison condition (*sig_addr cmp cmp_value). The caller must synchronize the stream (e.g., using hipStreamSynchronize) to ensure the wait condition has been satisfied.

Valid cmp values are listed in Supported comparisons.

Supported comparisons#

The following table lists the point-to-point comparison constants:

Point-to-Point Comparison Constants#

Constant

Description

ROCSHMEM_CMP_EQ

Equal

ROCSHMEM_CMP_NE

Not equal

ROCSHMEM_CMP_GT

Greater than

ROCSHMEM_CMP_GE

Greater than or equal to

ROCSHMEM_CMP_LT

Less than

ROCSHMEM_CMP_LE

Less than or equal to
Contents