Atomic memory operations

You can call these functions from divergent control paths at the per-thread level.

ROSHMEM_ATOMIC_FETCH

__device__ TYPE rocshmem_TYPENAME_atomic_fetch(TYPE *source, int pe)

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch(rocshmem_ctx_t ctx, TYPE *source, int pe)

Parameters:

• ctx – Context with which to perform this operation.

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

• pe – PE of the remote process.

Returns:

The value of dest.

Description: This function atomically returns the value of dest to the calling PE.

Valid TYPENAME and TYPE values are listed in EXTENDED_AMO_TYPES.

SHMEM_ATOMIC_SET

__device__ void rocshmem_TYPENAME_atomic_set(TYPE *dest, TYPE value, int pe);

__device__ void rocshmem_ctx_TYPENAME_atomic_set(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, int pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically set.

• pe – PE of the remote process.

Returns:

None.

Description: This function atomically sets the value value to dest on pe.

Valid TYPENAME and TYPE values are listed in EXTENDED_AMO_TYPES.

SHMEM_ATOMIC_COMPARE_SWAP

__device__ TYPE rocshmem_TYPENAME_atomic_compare_swap(TYPE *dest, TYPE cond, TYPE value, TYPE pe);

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_compare_swap(rocshmem_ctx_t ctx, TYPE *dest, TYPE cond, TYPE value, TYPE pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• cond – The value to be compare with.

• value – The value to be atomically swapped.

• pe – PE of the remote process.

Returns:

The old value of dest.

Description: This function atomically compares the value in dest with cond. If they are equal, it stores value in dest. The operation returns the older value of dest to the calling PE. The operation is blocking.

Valid TYPENAME and TYPE values are listed in STANDARD_AMO_TYPES.

SHMEM_ATOMIC_SWAP

__device__ TYPE rocshmem_TYPENAME_atomic_swap(TYPE *dest, TYPE value, TYPE pe);

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_swap(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically swapped.

• pe – PE of the remote process.

Returns:

The old value of dest.

Description: This function atomically swaps the value val with dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in EXTENDED_AMO_TYPES.

SHMEM_ATOMIC_FETCH_INC

__device__ TYPE rocshmem_TYPENAME_atomic_fetch_inc(TYPE *dest, TYPE pe);

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_inc(rocshmem_ctx_t ctx, TYPE *dest, TYPE pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• pe – PE of the remote process.

Returns:

The old value of dest.

Description: This function atomically adds 1 to dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in STANDARD_AMO_TYPES.

SHMEM_ATOMIC_INC

__device__ void rocshmem_TYPENAME_atomic_inc(TYPE *dest, TYPE pe);

__device__ void rocshmem_ctx_TYPENAME_atomic_inc(rocshmem_ctx_t ctx, TYPE *dest, TYPE pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• pe – PE of the remote process.

Returns:

None.

Description: This function atomically adds 1 to dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in STANDARD_AMO_TYPES.

SHMEM_ATOMIC_FETCH_ADD

__device__ TYPE rocshmem_TYPENAME_atomic_fetch_add(TYPE *dest, TYPE value, TYPE pe);

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_add(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically added.

• pe – PE of the remote process.

Returns:

The old value of dest.

Description: This function atomically adds value to dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in STANDARD_AMO_TYPES.

SHMEM_ATOMIC_ADD

__device__ void rocshmem_TYPENAME_atomic_add(TYPE *dest, TYPE value, TYPE pe);

__device__ void rocshmem_ctx_TYPENAME_atomic_add(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically added.

• pe – PE of the remote process.

Returns:

None.

Description: This function atomically adds value to dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values can be seen in STANDARD_AMO_TYPES.

SHMEM_ATOMIC_FETCH_AND

__device__ TYPE rocshmem_TYPENAME_atomic_fetch_and(TYPE *dest, TYPE value, TYPE pe);

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_and(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically AND.

• pe – PE of the remote process.

Returns:

The old value of dest.

Description: This function atomically bitwise-and value to the value at dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in BITWISE_AMO_TYPES.

SHMEM_ATOMIC_AND

__device__ TYPE rocshmem_TYPENAME_atomic_and(TYPE *dest, TYPE value, TYPE pe);

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_and(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically AND.

• pe – PE of the remote process.

Returns:

None

Description: This function atomically bitwise-and value to the value at dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in BITWISE_AMO_TYPES.

SHMEM_ATOMIC_FETCH_OR

__device__ TYPE rocshmem_TYPENAME_atomic_fetch_or(TYPE *dest, TYPE value, TYPE pe)

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_or(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe)

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically OR.

• pe – PE of the remote process.

Returns:

The old value of dest.

Description: This function atomically bitwise-or value to the value at dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in BITWISE_AMO_TYPES.

SHMEM_ATOMIC_OR

__device__ TYPE rocshmem_TYPENAME_atomic_or(TYPE *dest, TYPE value, TYPE pe)

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_or(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe)

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically OR.

• pe – PE of the remote process.

Returns:

None.

Description: This function atomically bitwise-or value to the value at dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in BITWISE_AMO_TYPES.

SHMEM_ATOMIC_FETCH_XOR

__device__ TYPE rocshmem_TYPENAME_atomic_fetch_xor(TYPE *dest, TYPE value, TYPE pe);

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_xor(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe);

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically XOR.

• pe – PE of the remote process.

Returns:

The old value of dest.

Description: This function atomically bitwise-xor value to the value at dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in BITWISE_AMO_TYPES.

SHMEM_ATOMIC_XOR

__device__ TYPE rocshmem_TYPENAME_atomic_xor(TYPE *dest, TYPE value, TYPE pe)

__device__ TYPE rocshmem_ctx_TYPENAME_atomic_xor(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe)

Parameters:

• ctx – Context with which to perform this operation.

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

• value – The value to be atomically XOR.

• pe – PE of the remote process.

Returns:

None.

Description: This function atomically bitwise-xor value to the value at dest on pe. The operation is blocking.

Valid TYPENAME and TYPE values are listed in BITWISE_AMO_TYPES.

Supported AMO data types

Table 3 Standard AMO Data Types

TYPE	TYPENAME	Supported
int	int	Yes
long	long	Yes
long long	longlong	Yes
unsigned int	uint	Yes
unsigned long	ulong	Yes
unsigned long long	ulonglong	Yes
int32_t	int32	Yes
int64_t	int64	Yes
uint32_t	uint32	Yes
uint64_t	uint64	Yes
size_t	size	Yes
ptrdiff_t	ptrdiff	Yes

Table 4 Extended AMO Data Types

TYPE	TYPENAME	Supported
float	float	Yes
double	double	Yes
int	int	Yes
long	long	Yes
long long	longlong	Yes
unsigned int	uint	Yes
unsigned long	ulong	Yes
unsigned long long	ulonglong	Yes
int32_t	int32	Yes
int64_t	int64	Yes
uint32_t	uint32	Yes
uint64_t	uint64	Yes
size_t	size	Yes
ptrdiff_t	ptrdiff	Yes

Table 5 Bitwise AMO Data Types

TYPE	TYPENAME	Supported
unsigned int	uint	Yes
unsigned long	ulong	Yes
unsigned long long	ulonglong	Yes
int32_t	int32	Yes
int64_t	int64	Yes
uint32_t	uint32	Yes
uint64_t	uint64	Yes