Atomics¶

RAJA provides portable atomic operations that can be used to update values at arbitrary memory locations while avoiding data races. They are described in this section.

A complete working example code that shows RAJA atomic usage can be found in Computing a Histogram with Atomic Operations.

Note

• All RAJA atomic operations are in the namespace RAJA::atomic.

Atomic Operations¶

RAJA atomic support includes a variety of the most common atomic operations.

Note

• Each RAJA atomic operation is templated on an atomic policy.
• Each method described in the table below returns the value of the potentially modified argument (i.e., *acc) immediately before the atomic operation is applied, in case it is needed by a user.

Arithmetic¶

• atomicAdd< atomic_policy >(T* acc, T value) - Add value to *acc.
• atomicSub< atomic_policy >(T* acc, T value) - Subtract value from *acc.

Min/max¶

• atomicMin< atomic_policy >(T* acc, T value) - Set *acc to min of *acc and value.
• atomicMax< atomic_policy >(T* acc, T value) - Set *acc to max of *acc and value.

Increment/decrement¶

• atomicInc< atomic_policy >(T* acc) - Add 1 to *acc.
• atomicDec< atomic_policy >(T* acc) - Subtract 1 from *acc.
• atomicInc< atomic_policy >(T* acc, T compare) - Add 1 to *acc if *acc < compare, else set *acc to zero.
• atomicDec< atomic_policy >(T* acc, T compare) - Subtract 1 from *acc if *acc != 0 and *acc <= compare, else set *acc to compare.

Bitwise operations¶

• atomicAnd< atomic_policy >(T* acc, T value) - Bitwise ‘and’ equivalent: Set *acc to *acc & value. Only works with integral data types.
• atomicOr< atomic_policy >(T* acc, T value) - Bitwise ‘or’ equivalent: Set *acc to *acc | value. Only works with integral data types.
• atomicXor< atomic_policy >(T* acc, T value) - Bitwise ‘xor’ equivalent: Set *acc to *acc ^ value. Only works with integral data types.

Replace¶

• atomicExchange< atomic_policy >(T* acc, T value) - Replace *acc with value.
• atomicCAS< atomic_policy >(T* acc, Tcompare, T value) - Compare and swap: Replace *acc with value if and only if *acc is equal to compare.

Here is a simple example that shows how to use an atomic operation to compute an integral sum on a CUDA GPU device:

//
// Use CUDA UM to share data pointer with host and device code.
// RAJA mechanics work the same way if device data allocation
// and host-device copies are done with traditional cudaMalloc
// and cudaMemcpy.
//
int* sum = nullptr;
cudaMallocManaged((void **)&sum, sizeof(int));
*sum = 0;

RAJA::forall< RAJA::cuda_exec >(RAJA::RangeSegment(0, N),
[=] RAJA_DEVICE (RAJA::Index_type i) {

});

After this kernel executes, ‘*sum’ will be equal to ‘N’.

AtomicRef¶

RAJA also provides an atomic interface similar to the C++20 ‘std::atomic_ref’, but which works for arbitrary memory locations. The class RAJA::atomic::AtomicRef provides an object-oriented interface to the atomic methods described above. For example, after the following operations:

double val = 2.0;
RAJA::atomic::AtomicRef<double,  RAJA::omp_atomic > sum(&val);

sum++;
++sum;
sum += 1.0;

the value of ‘val’ will be 5.