Added Windows support

long type is ambiguous depending on operating system, whereas int64_t ensures consistent use of 64bit types across all platforms

Author: Uakh

cuda/src/ball_query.cpp

@@ -3,13 +3,13 @@
 #include "utils.h"
 
 void query_ball_point_kernel_dense_wrapper(int b, int n, int m, float radius, int nsample,
-                                           const float* new_xyz, const float* xyz, long* idx,
+                                           const float* new_xyz, const float* xyz, int64_t* idx,
                                            float* dist_out);
 
-void query_ball_point_kernel_partial_wrapper(long batch_size, int size_x, int size_y, float radius,
+void query_ball_point_kernel_partial_wrapper(int64_t batch_size, int size_x, int size_y, float radius,
                                              int nsample, const float* x, const float* y,
-                                             const long* batch_x, const long* batch_y,
-                                             long* idx_out, float* dist_out);
+                                             const int64_t* batch_x, const int64_t* batch_y,
+                                             int64_t* idx_out, float* dist_out);
 
 std::pair<at::Tensor, at::Tensor> ball_query_dense(at::Tensor new_xyz, at::Tensor xyz,
                                                    const float radius, const int nsample)
@@ -29,7 +29,7 @@ std::pair<at::Tensor, at::Tensor> ball_query_dense(at::Tensor new_xyz, at::Tenso
 
     query_ball_point_kernel_dense_wrapper(xyz.size(0), xyz.size(1), new_xyz.size(1), radius,
                                           nsample, new_xyz.DATA_PTR<float>(), xyz.DATA_PTR<float>(),
-                                          idx.DATA_PTR<long>(), dist.DATA_PTR<float>());
+                                          idx.DATA_PTR<int64_t>(), dist.DATA_PTR<float>());
 
     return std::make_pair(idx, dist);
 }
@@ -73,8 +73,8 @@ std::pair<at::Tensor, at::Tensor> ball_query_partial_dense(at::Tensor x, at::Ten
 
     query_ball_point_kernel_partial_wrapper(batch_size, x.size(0), y.size(0), radius, nsample,
                                             x.DATA_PTR<float>(), y.DATA_PTR<float>(),
-                                            batch_x.DATA_PTR<long>(), batch_y.DATA_PTR<long>(),
-                                            idx.DATA_PTR<long>(), dist.DATA_PTR<float>());
+                                            batch_x.DATA_PTR<int64_t>(), batch_y.DATA_PTR<int64_t>(),
+                                            idx.DATA_PTR<int64_t>(), dist.DATA_PTR<float>());
 
     return std::make_pair(idx, dist);
 }

cuda/src/ball_query_gpu.cu

@@ -9,7 +9,7 @@
 __global__ void query_ball_point_kernel_dense(int b, int n, int m, float radius, int nsample,
                                               const float* __restrict__ new_xyz,
                                               const float* __restrict__ xyz,
-                                              long* __restrict__ idx_out, 
+                                              int64_t* __restrict__ idx_out, 
                                               float* __restrict__ dist_out)
 {
     int batch_index = blockIdx.x;
@@ -53,7 +53,7 @@ __global__ void query_ball_point_kernel_dense(int b, int n, int m, float radius,
 
 __global__ void query_ball_point_kernel_partial_dense(
     int size_x, int size_y, float radius, int nsample, const float* __restrict__ x,
-    const float* __restrict__ y, const long* __restrict__ batch_x, const long* __restrict__ batch_y,
+    const float* __restrict__ y, const int64_t* __restrict__ batch_x, const int64_t* __restrict__ batch_y,
     int64_t* __restrict__ idx_out, float* __restrict__ dist_out)
 {
     // taken from
@@ -93,7 +93,7 @@ __global__ void query_ball_point_kernel_partial_dense(
 }
 
 void query_ball_point_kernel_dense_wrapper(int b, int n, int m, float radius, int nsample,
-                                           const float* new_xyz, const float* xyz, long* idx,float* dist_out)
+                                           const float* new_xyz, const float* xyz, int64_t* idx,float* dist_out)
 {
     cudaStream_t stream = at::cuda::getCurrentCUDAStream();
     query_ball_point_kernel_dense<<<b, opt_n_threads(m), 0, stream>>>(b, n, m, radius, nsample,
@@ -102,9 +102,9 @@ void query_ball_point_kernel_dense_wrapper(int b, int n, int m, float radius, in
     CUDA_CHECK_ERRORS();
 }
 
-void query_ball_point_kernel_partial_wrapper(long batch_size, int size_x, int size_y, float radius,
+void query_ball_point_kernel_partial_wrapper(int64_t batch_size, int size_x, int size_y, float radius,
                                              int nsample, const float* x, const float* y,
-                                             const long* batch_x, const long* batch_y,
+                                             const int64_t* batch_x, const int64_t* batch_y,
                                              int64_t* idx_out, float* dist_out)
 {
     query_ball_point_kernel_partial_dense<<<batch_size, TOTAL_THREADS_SPARSE>>>(