ball_query.cpp 1.71 KB
Newer Older
wuyuefeng's avatar
Credit  
wuyuefeng committed
1
2
3
// Modified from
// https://github.com/sshaoshuai/Pointnet2.PyTorch/tree/master/pointnet2/src/ball_query.cpp

wuyuefeng's avatar
wuyuefeng committed
4
5
6
7
#include <THC/THC.h>
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <torch/extension.h>
zhangwenwei's avatar
zhangwenwei committed
8
9
10
#include <torch/serialize/tensor.h>

#include <vector>
wuyuefeng's avatar
wuyuefeng committed
11
12
13

extern THCState *state;

zhangwenwei's avatar
zhangwenwei committed
14
15
16
17
18
19
20
#define CHECK_CUDA(x) \
  TORCH_CHECK(x.type().is_cuda(), #x, " must be a CUDAtensor ")
#define CHECK_CONTIGUOUS(x) \
  TORCH_CHECK(x.is_contiguous(), #x, " must be contiguous ")
#define CHECK_INPUT(x) \
  CHECK_CUDA(x);       \
  CHECK_CONTIGUOUS(x)
wuyuefeng's avatar
wuyuefeng committed
21

22
int ball_query_wrapper(int b, int n, int m, float min_radius, float max_radius, int nsample,
zhangwenwei's avatar
zhangwenwei committed
23
24
                       at::Tensor new_xyz_tensor, at::Tensor xyz_tensor,
                       at::Tensor idx_tensor);
wuyuefeng's avatar
wuyuefeng committed
25

26
27
void ball_query_kernel_launcher(int b, int n, int m, float min_radius, float max_radius,
                                int nsample, const float *xyz, const float *new_xyz,
zhangwenwei's avatar
zhangwenwei committed
28
                                int *idx, cudaStream_t stream);
wuyuefeng's avatar
wuyuefeng committed
29

30
int ball_query_wrapper(int b, int n, int m, float min_radius, float max_radius, int nsample,
zhangwenwei's avatar
zhangwenwei committed
31
32
33
34
35
36
37
38
                       at::Tensor new_xyz_tensor, at::Tensor xyz_tensor,
                       at::Tensor idx_tensor) {
  CHECK_INPUT(new_xyz_tensor);
  CHECK_INPUT(xyz_tensor);
  const float *new_xyz = new_xyz_tensor.data_ptr<float>();
  const float *xyz = xyz_tensor.data_ptr<float>();
  int *idx = idx_tensor.data_ptr<int>();

39
  cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream();
40
41
  ball_query_kernel_launcher(b, n, m, min_radius, max_radius,
                             nsample, new_xyz, xyz, idx, stream);
zhangwenwei's avatar
zhangwenwei committed
42
  return 1;
wuyuefeng's avatar
wuyuefeng committed
43
44
45
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
zhangwenwei's avatar
zhangwenwei committed
46
  m.def("ball_query_wrapper", &ball_query_wrapper, "ball_query_wrapper");
wuyuefeng's avatar
wuyuefeng committed
47
}