Atomic memory operations

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