From 202f4cd11c0a0c4fc636640a1e1cc5f762dd695c Mon Sep 17 00:00:00 2001
From: brucefan1983 <brucenju@gmail.com>
Date: Fri, 3 Jan 2025 18:58:26 +0800
Subject: [PATCH] keep K80 to be usable

---
 src/force/nep_small_box.cuh | 17 +++++++++++++++++
 1 file changed, 17 insertions(+)

diff --git a/src/force/nep_small_box.cuh b/src/force/nep_small_box.cuh
index 359dc4005..6cb80c767 100644
--- a/src/force/nep_small_box.cuh
+++ b/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)
 {