kernel_compat.h

Go to the documentation of this file.

#ifndef KERNEL_COMPAT_H_
#define KERNEL_COMPAT_H_
 
// atomicAdd for double precision - required for CUDA architectures < 600 (pre-Pascal)
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 600 && !defined(__CUDA_ON_DCU)
static __inline__ __device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull = (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do
    {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
        // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
    } while (assumed != old);
    return __longlong_as_double(old);
}
#endif
 
#endif // KERNEL_COMPAT_H_