5 #ifndef FML_GPU_INTERNALS_ATOMICS_H
6 #define FML_GPU_INTERNALS_ATOMICS_H
10 #include <cuda_runtime.h>
12 #include "../../_internals/types.hh"
14 #include "../../_internals/arraytools/src/arraytools.cuh"
22 static __device__
float atomicMaxf(
float *address,
float val)
24 int *address_int = (
int*) address;
25 int old = *address_int;
28 while (val > __int_as_float(old))
31 old = atomicCAS(address_int, assumed, __float_as_int(val));
34 return __int_as_float(old);
37 static __device__
double atomicMaxf(
double *address,
double val)
39 unsigned long long *address_ull = (
unsigned long long*) address;
40 unsigned long long old = *address_ull;
41 unsigned long long assumed;
43 while (val > __longlong_as_double(old))
46 old = atomicCAS(address_ull, assumed, __double_as_longlong(val));
49 return __longlong_as_double(old);
54 static __device__
float atomicMinf(
float *address,
float val)
56 int *address_int = (
int*) address;
57 int old = *address_int;
63 old = atomicCAS(address_int, assumed,
64 __float_as_int(fmin(val, __int_as_float(assumed))));
65 }
while (old != assumed);
67 return __int_as_float(old);
70 static __device__
double atomicMinf(
double *address,
double val)
72 unsigned long long * address_as_ull = (
unsigned long long*) address;
73 unsigned long long old = *address_as_ull;
74 unsigned long long assumed;
78 old = atomicCAS(address_as_ull, assumed,
79 __double_as_longlong(fmin(val, __longlong_as_double(assumed))));
80 }
while (assumed != old);
82 return __longlong_as_double(old);
87 static __device__
float atomicMul(
float *address,
float val)
89 int *address_int = (
int*) address;
90 int old = *address_int;
96 old = atomicCAS(address_int, assumed,
97 __float_as_int(val * __float_as_int(assumed)));
98 }
while (old != assumed);
100 return __int_as_float(old);
103 static __device__
double atomicMul(
double *address,
double val)
105 unsigned long long *address_int = (
unsigned long long*) address;
106 unsigned long long old = *address_int;
107 unsigned long long assumed;
112 old = atomicCAS(address_int, assumed,
113 __double_as_longlong(val * __double_as_longlong(assumed)));
114 }
while (old != assumed);
116 return __int_as_float(old);