16#ifndef dealii_cuda_atomic_h
17#define dealii_cuda_atomic_h
21#ifdef DEAL_II_COMPILER_CUDA_AWARE
36 inline __device__
float
39 return atomicAdd(address, val);
51 inline __device__
double
54 return atomicAdd(address, val);
64 inline __device__
float
67 int *address_as_int =
reinterpret_cast<int *
>(address);
68 int old = *address_as_int, assumed;
72 old = atomicCAS(address_as_int,
74 atomicMax(address_as_int, __float_as_int(val)));
76 while (assumed != old);
78 return __longlong_as_double(old);
88 inline __device__
double
91 unsigned long long int *address_as_ull =
92 reinterpret_cast<unsigned long long int *
>(address);
93 unsigned long long int old = *address_as_ull, assumed;
97 old = atomicCAS(address_as_ull,
99 atomicMax(address_as_ull,
100 static_cast<unsigned long long int>(
101 __double_as_longlong(val))));
103 while (assumed != old);
105 return __longlong_as_double(old);
#define DEAL_II_DEPRECATED
#define DEAL_II_NAMESPACE_OPEN
#define DEAL_II_NAMESPACE_CLOSE
float atomicMax_wrapper(float *address, float val)
float atomicAdd_wrapper(float *address, float val)