add cuda arch 3.5 compilation support for ABACUS (#2794)

deepmodeling · Aug 4, 2023 · cf1270f · cf1270f
1 parent ce14dcc
commit cf1270f
Show file tree

Hide file tree

Showing 3 changed files with 17 additions and 0 deletions.
diff --git a/source/module_hamilt_pw/hamilt_pwdft/kernels/cuda/force_op.cu b/source/module_hamilt_pw/hamilt_pwdft/kernels/cuda/force_op.cu
@@ -1,4 +1,5 @@
 #include "module_hamilt_pw/hamilt_pwdft/kernels/force_op.h"
+#include "module_psi/kernels/device.h"
 
 #include <complex>
 

diff --git a/source/module_hamilt_pw/hamilt_pwdft/kernels/cuda/stress_op.cu b/source/module_hamilt_pw/hamilt_pwdft/kernels/cuda/stress_op.cu
@@ -1,4 +1,5 @@
 #include "module_hamilt_pw/hamilt_pwdft/kernels/stress_op.h"
+#include "module_psi/kernels/device.h"
 
 #include <complex>
 #include <thrust/complex.h>

diff --git a/source/module_psi/kernels/device.h b/source/module_psi/kernels/device.h
@@ -8,6 +8,21 @@
 #include <complex>
 #include <iostream>
 
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 600
+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);
+  } while (assumed != old);
+  return __longlong_as_double(old);
+}
+#endif
+
 namespace psi {
 namespace device {