Merge pull request #44 from Uakh/windows

Adding Windows support

Author: nicolas-chaulet

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>>>(

cuda/src/metrics.cpp

@@ -2,11 +2,11 @@
 #include "compat.h"
 #include "utils.h"
 
-void instance_iou_kernel_wrapper(long total_gt_instances, long max_gt_instances,
-                                 const long* nInstance, int nProposal, const long* proposals_idx,
-                                 const long* proposals_offset, const long* instance_labels,
-                                 const long* offset_num_gt_instances, const long* batch,
-                                 const long* instance_pointnum, float* proposals_iou);
+void instance_iou_kernel_wrapper(int64_t total_gt_instances, int64_t max_gt_instances,
+                                 const int64_t* nInstance, int nProposal, const int64_t* proposals_idx,
+                                 const int64_t* proposals_offset, const int64_t* instance_labels,
+                                 const int64_t* offset_num_gt_instances, const int64_t* batch,
+                                 const int64_t* instance_pointnum, float* proposals_iou);
 
 at::Tensor instance_iou_cuda(at::Tensor instance_idx, at::Tensor instance_offsets,
                              at::Tensor gt_instances, at::Tensor gt_instance_sizes,
@@ -25,7 +25,7 @@ at::Tensor instance_iou_cuda(at::Tensor instance_idx, at::Tensor instance_offset
     CHECK_CUDA(gt_instance_sizes);
 
     cudaSetDevice(instance_idx.get_device());
-    long num_proposed_instances = instance_offsets.size(0) - 1;
+    int64_t num_proposed_instances = instance_offsets.size(0) - 1;
     auto total_gt_instances = (int64_t*)malloc(sizeof(int64_t));
     cudaMemcpy(total_gt_instances, num_gt_instances.sum().DATA_PTR<int64_t>(), sizeof(int64_t),
                cudaMemcpyDeviceToHost);
@@ -40,10 +40,10 @@ at::Tensor instance_iou_cuda(at::Tensor instance_idx, at::Tensor instance_offset
     at::Tensor offset_num_gt_instances =
         at::cat({at::zeros(1, num_gt_instances.options()), num_gt_instances.cumsum(0)}, 0);
     instance_iou_kernel_wrapper(
-        total_gt_instances[0], max_gt_instances[0], num_gt_instances.DATA_PTR<long>(),
-        num_proposed_instances, instance_idx.DATA_PTR<long>(), instance_offsets.DATA_PTR<long>(),
-        gt_instances.DATA_PTR<long>(), offset_num_gt_instances.DATA_PTR<long>(),
-        batch.DATA_PTR<long>(), gt_instance_sizes.DATA_PTR<long>(), output.DATA_PTR<float>());
+        total_gt_instances[0], max_gt_instances[0], num_gt_instances.DATA_PTR<int64_t>(),
+        num_proposed_instances, instance_idx.DATA_PTR<int64_t>(), instance_offsets.DATA_PTR<int64_t>(),
+        gt_instances.DATA_PTR<int64_t>(), offset_num_gt_instances.DATA_PTR<int64_t>(),
+        batch.DATA_PTR<int64_t>(), gt_instance_sizes.DATA_PTR<int64_t>(), output.DATA_PTR<float>());
 
     return output;
 }

cuda/src/metrics_gpu.cu

@@ -7,10 +7,10 @@
 #define THREADS 512
 
 __global__ void instance_iou_cuda_kernel(
-    long total_gt_instances, const long* __restrict__ nInstance, int nProposal,
-    const long* __restrict__ proposals_idx, const long* __restrict__ proposals_offset,
-    const long* __restrict__ instance_labels, const long* __restrict__ offset_num_gt_instances,
-    const long* __restrict__ batch, const long* __restrict__ instance_pointnum,
+    int64_t total_gt_instances, const int64_t* __restrict__ nInstance, int nProposal,
+    const int64_t* __restrict__ proposals_idx, const int64_t* __restrict__ proposals_offset,
+    const int64_t* __restrict__ instance_labels, const int64_t* __restrict__ offset_num_gt_instances,
+    const int64_t* __restrict__ batch, const int64_t* __restrict__ instance_pointnum,
     float* proposals_iou)
 {
     for (int proposal_id = blockIdx.x; proposal_id < nProposal; proposal_id += gridDim.x)
@@ -44,18 +44,18 @@ __global__ void instance_iou_cuda_kernel(
 
 // input: proposals_idx (sumNPoint), int
 // input: proposals_offset (nProposal + 1), int
-// input: instance_labels (N), long, 0~total_nInst-1, -100
+// input: instance_labels (N), int64_t, 0~total_nInst-1, -100
 // input: instance_pointnum (total_nInst), int
 // output: proposals_iou (nProposal, total_nInst), float
-void instance_iou_kernel_wrapper(long total_gt_instances, long max_gt_instances,
-                                 const long* nInstance, int nProposal, const long* proposals_idx,
-                                 const long* proposals_offset, const long* instance_labels,
-                                 const long* offset_num_gt_instances, const long* batch,
-                                 const long* instance_pointnum, float* proposals_iou)
+void instance_iou_kernel_wrapper(int64_t total_gt_instances, int64_t max_gt_instances,
+                                 const int64_t* nInstance, int nProposal, const int64_t* proposals_idx,
+                                 const int64_t* proposals_offset, const int64_t* instance_labels,
+                                 const int64_t* offset_num_gt_instances, const int64_t* batch,
+                                 const int64_t* instance_pointnum, float* proposals_iou)
 {
     auto stream = at::cuda::getCurrentCUDAStream();
     instance_iou_cuda_kernel<<<std::min(nProposal, THREADS * THREADS),
-                               std::min(max_gt_instances, (long)THREADS), 0, stream>>>(
+                               std::min(max_gt_instances, (int64_t)THREADS), 0, stream>>>(
         total_gt_instances, nInstance, nProposal, proposals_idx, proposals_offset, instance_labels,
         offset_num_gt_instances, batch, instance_pointnum, proposals_iou);
 }