atomic.cuh

Go to the documentation of this file.

#pragma once

#if defined(__CUDA_ARCH__) 

#if __COMPUTE_CAPABILITY__ < 600

static inline __device__ double atomicAdd(double *addr, double val){
  double old = *addr, assumed;
  do {
    assumed = old;
    old = __longlong_as_double( atomicCAS((unsigned long long int*)addr,
            __double_as_longlong(assumed),
            __double_as_longlong(val + assumed)));
  } while ( __double_as_longlong(assumed) != __double_as_longlong(old) );
  
  return old;
}
#endif

static inline __device__ double2 atomicAdd(double2 *addr, double2 val){
  double2 old = *addr;
  old.x = atomicAdd((double*)addr, val.x);
  old.y = atomicAdd((double*)addr + 1, val.y);
  return old;
}

static inline __device__ float2 atomicAdd(float2 *addr, float2 val){
  float2 old = *addr;
  old.x = atomicAdd((float*)addr, val.x);
  old.y = atomicAdd((float*)addr + 1, val.y);
  return old;
}

#endif