Merge pull request brucefan1983#844 from brucefan1983/no_special_atomic

no SM less than 60

Author: brucefan1983

src/force/nep_small_box.cuh

@@ -19,21 +19,6 @@
 #include "utilities/nep_utilities.cuh"
 #include "utilities/gpu_macro.cuh"
 
-#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
-
 static __device__ void apply_mic_small_box(
   const Box& box, const NEP::ExpandedBox& ebox, double& x12, double& y12, double& z12)
 {