Merge pull request #846 from brucefan1983/k80_back

keep K80 to be usable

src/force/nep_small_box.cuh

@@ -19,6 +19,23 @@
 #include "utilities/nep_utilities.cuh"
 #include "utilities/gpu_macro.cuh"
 
+#ifdef USE_KEPLER
+#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 600)
+static __device__ __inline__ double atomicAdd(double* address, double val)
+{
+  unsigned long long* address_as_ull = (unsigned long long*)address;
+  unsigned long long old = *address_as_ull, assumed;
+  do {
+    assumed = old;
+    old =
+      atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
+
+  } while (assumed != old);
+  return __longlong_as_double(old);
+}
+#endif
+#endif
+
 static __device__ void apply_mic_small_box(
   const Box& box, const NEP::ExpandedBox& ebox, double& x12, double& y12, double& z12)
 {