1 //////////////////////////////////////////////////////////////////////
2 // This file is distributed under the University of Illinois/NCSA Open Source
3 // License. See LICENSE file in top directory for details.
4 //
5 // Copyright (c) 2020 QMCPACK developers.
6 //
7 // File developed by:
8 // Lawrence Livermore National Laboratory
9 //
10 // File created by:
11 // Miguel A. Morales, moralessilva2@llnl.gov
12 // Lawrence Livermore National Laboratory
13 ////////////////////////////////////////////////////////////////////////////////
14
15
16 #ifndef CUDA_WORKAROUND_LEGACY_HARDWARE_H
17 #define CUDA_WORKAROUND_LEGACY_HARDWARE_H
18
19 namespace kernels
20 {
21 #if __CUDA_ARCH__ < 600
atomicAdd(double * address,double val)22 inline __device__ double atomicAdd(double* address, double val)
23 {
24 unsigned long long int* address_as_ull = (unsigned long long int*)address;
25 unsigned long long int old = *address_as_ull, assumed;
26
27 do
28 {
29 assumed = old;
30 old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
31
32 // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
33 } while (assumed != old);
34
35 return __longlong_as_double(old);
36 }
37
atomicAdd(float * address,float val)38 inline __device__ float atomicAdd(float* address, float val)
39 {
40 unsigned long long int* address_as_ull = (unsigned long long int*)address;
41 unsigned long long int old = *address_as_ull, assumed;
42
43 do
44 {
45 assumed = old;
46 old = atomicCAS(address_as_ull, assumed, __float_as_int(val + __int_as_float(assumed)));
47
48 // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
49 } while (assumed != old);
50
51 return __int_as_float(old);
52 }
53
54 #endif
55 } // namespace kernels
56 #endif
57