Merge pull request #52 from CCInc/atomicAdd

added atomicadd implementation (fix #50)

Author: nicolas-chaulet

cuda/include/cuda_utils.h

@@ -29,6 +29,28 @@ inline dim3 opt_block_config(int x, int y)
     return block_config;
 }
 
+// from https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions
+#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)));
+
+        // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
+        } while (assumed != old);
+
+        return __longlong_as_double(old);
+    }
+#endif
+
 #define CUDA_CHECK_ERRORS()                                                                        \
     do                                                                                             \
     {                                                                                              \

cuda/src/chamfer_dist.cu

@@ -3,6 +3,7 @@
 #include <torch/extension.h>
 
 #include <vector>
+#include "cuda_utils.h"
 
 template <typename scalar_t>
 __global__ void chamfer_dist_kernel(int batch_size, int n, const scalar_t* __restrict__ xyz1, int m,