From a98df2d100d20c3ec2fb641071f0bf4633b8109c Mon Sep 17 00:00:00 2001
From: greg <go2432@wayne.edu>
Date: Sat, 14 Oct 2023 16:56:30 -0500
Subject: [PATCH] add prepascal atomicAdd double

---
 cpp/cuda_common.cuh | 15 +++++++++++++++
 1 file changed, 15 insertions(+)

diff --git a/cpp/cuda_common.cuh b/cpp/cuda_common.cuh
index c6965a2..cca6e19 100644
--- a/cpp/cuda_common.cuh
+++ b/cpp/cuda_common.cuh
@@ -14,6 +14,21 @@ const int threads_per_block = 512;
 #define tensor_acc_3(T, N, type) (T).packed_accessor64<type, N, torch::RestrictPtrTraits>()
 
 /* Helpers */
+#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600
+#else
+__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)));
+    } while (assumed != old);
+    return __longlong_as_double(old);
+}
+#endif
+
 
 /** Error handling */
 inline void _cuda_check_err(const cudaError_t err, const char* file, const int line, const char* function) {