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.


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

Atomic Operations

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


  • 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.
  • Please set the nvcc option “-arch=sm_35” (or greater), when CUDA is enabled. If a CUDA architecture does not support an sm_xy atomic operation, RAJA will generically implement it with atomicCAS.


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


  • 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.


  • 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.


  • 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) {

  RAJA::atomic::atomicAdd< RAJA::cuda_atomic >(sum, 1);


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


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 += 1.0;

the value of ‘val’ will be 5.

Atomic Policies

For more information about available RAJA atomic policies, please see Atomic Policies.