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. The function blocks the calling thread until the signal variable at sig_addr satisfies the comparison condition (*sig_addr cmp cmp_value). The wait operation is executed asynchronously on the specified stream. 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:

Table 7 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