Replace atomicadd with PyTorch implementation (#75)

Author: clee-ai

torchsparse/src/common/gpu.cuh

@@ -104,11 +104,6 @@ template <typename Dtype1, typename Dtype2>
 void print(const thrust::device_vector<Dtype1> &v1,
            const thrust::device_vector<Dtype2> &v2);
 
-// atomicadd for half types (from aten/src/THC/THCAtomics.cuh)
-static inline  __device__ at::Half atomicAdd(at::Half *address, at::Half val) {
-  return atomicAdd(reinterpret_cast<__half*>(address), val);
-}
-
 // AtomicAddition for double with cuda arch <= 600
 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600
 #else

torchsparse/src/interpolation/devox_gpu.cu

@@ -2,7 +2,7 @@
 #include <stdlib.h>
 #include <thrust/device_vector.h>
 #include <torch/extension.h>
-#include "../common/gpu.cuh"
+#include <THC/THCAtomics.cuh>
 
 
 //input features (n, c), indices (N, 8), weight (N, 8) -> output features (N, c)

torchsparse/src/others/insertion_gpu.cu

@@ -2,7 +2,7 @@
 #include <stdlib.h>
 #include <cmath>
 #include <torch/torch.h>
-#include "../common/gpu.cuh"
+#include <THC/THCAtomics.cuh>
 
 //hashing
 //input N*F float tensor, pointer to output N'*F int64 tensor, N*1 count tensor, N*1 index tensor