Skip to content

Commit 487d4d6

Browse files
gkioxarifacebook-github-bot
gkioxari
authored and
facebook-github-bot
committed
point mesh distances
Summary: Implementation of point to mesh distances. The current diff contains two types: (a) Point to Edge (b) Point to Face ``` Benchmark Avg Time(μs) Peak Time(μs) Iterations -------------------------------------------------------------------------------- POINT_MESH_EDGE_4_100_300_5000_cuda:0 2745 3138 183 POINT_MESH_EDGE_4_100_300_10000_cuda:0 4408 4499 114 POINT_MESH_EDGE_4_100_3000_5000_cuda:0 4978 5070 101 POINT_MESH_EDGE_4_100_3000_10000_cuda:0 9076 9187 56 POINT_MESH_EDGE_4_1000_300_5000_cuda:0 1411 1487 355 POINT_MESH_EDGE_4_1000_300_10000_cuda:0 4829 5030 104 POINT_MESH_EDGE_4_1000_3000_5000_cuda:0 7539 7620 67 POINT_MESH_EDGE_4_1000_3000_10000_cuda:0 12088 12272 42 POINT_MESH_EDGE_8_100_300_5000_cuda:0 3106 3222 161 POINT_MESH_EDGE_8_100_300_10000_cuda:0 8561 8648 59 POINT_MESH_EDGE_8_100_3000_5000_cuda:0 6932 7021 73 POINT_MESH_EDGE_8_100_3000_10000_cuda:0 24032 24176 21 POINT_MESH_EDGE_8_1000_300_5000_cuda:0 5272 5399 95 POINT_MESH_EDGE_8_1000_300_10000_cuda:0 11348 11430 45 POINT_MESH_EDGE_8_1000_3000_5000_cuda:0 17478 17683 29 POINT_MESH_EDGE_8_1000_3000_10000_cuda:0 25961 26236 20 POINT_MESH_EDGE_16_100_300_5000_cuda:0 8244 8323 61 POINT_MESH_EDGE_16_100_300_10000_cuda:0 18018 18071 28 POINT_MESH_EDGE_16_100_3000_5000_cuda:0 19428 19544 26 POINT_MESH_EDGE_16_100_3000_10000_cuda:0 44967 45135 12 POINT_MESH_EDGE_16_1000_300_5000_cuda:0 7825 7937 64 POINT_MESH_EDGE_16_1000_300_10000_cuda:0 18504 18571 28 POINT_MESH_EDGE_16_1000_3000_5000_cuda:0 65805 66132 8 POINT_MESH_EDGE_16_1000_3000_10000_cuda:0 90885 91089 6 -------------------------------------------------------------------------------- Benchmark Avg Time(μs) Peak Time(μs) Iterations -------------------------------------------------------------------------------- POINT_MESH_FACE_4_100_300_5000_cuda:0 1561 1685 321 POINT_MESH_FACE_4_100_300_10000_cuda:0 2818 2954 178 POINT_MESH_FACE_4_100_3000_5000_cuda:0 15893 16018 32 POINT_MESH_FACE_4_100_3000_10000_cuda:0 16350 16439 31 POINT_MESH_FACE_4_1000_300_5000_cuda:0 3179 3278 158 POINT_MESH_FACE_4_1000_300_10000_cuda:0 2353 2436 213 POINT_MESH_FACE_4_1000_3000_5000_cuda:0 16262 16336 31 POINT_MESH_FACE_4_1000_3000_10000_cuda:0 9334 9448 54 POINT_MESH_FACE_8_100_300_5000_cuda:0 4377 4493 115 POINT_MESH_FACE_8_100_300_10000_cuda:0 9728 9822 52 POINT_MESH_FACE_8_100_3000_5000_cuda:0 26428 26544 19 POINT_MESH_FACE_8_100_3000_10000_cuda:0 42238 43031 12 POINT_MESH_FACE_8_1000_300_5000_cuda:0 3891 3982 129 POINT_MESH_FACE_8_1000_300_10000_cuda:0 5363 5429 94 POINT_MESH_FACE_8_1000_3000_5000_cuda:0 20998 21084 24 POINT_MESH_FACE_8_1000_3000_10000_cuda:0 39711 39897 13 POINT_MESH_FACE_16_100_300_5000_cuda:0 5955 6001 84 POINT_MESH_FACE_16_100_300_10000_cuda:0 12082 12144 42 POINT_MESH_FACE_16_100_3000_5000_cuda:0 44996 45176 12 POINT_MESH_FACE_16_100_3000_10000_cuda:0 73042 73197 7 POINT_MESH_FACE_16_1000_300_5000_cuda:0 8292 8374 61 POINT_MESH_FACE_16_1000_300_10000_cuda:0 19442 19506 26 POINT_MESH_FACE_16_1000_3000_5000_cuda:0 36059 36194 14 POINT_MESH_FACE_16_1000_3000_10000_cuda:0 64644 64822 8 -------------------------------------------------------------------------------- ``` Reviewed By: jcjohnson Differential Revision: D20590462 fbshipit-source-id: 42a39837b514a546ac9471bfaff60eefe7fae829
1 parent 474c8b4 commit 487d4d6

33 files changed

+3437
-84
lines changed

‎pytorch3d/csrc/compositing/alpha_composite.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
// Copyright (c) Facebook, Inc. and its affiliates. All rights reserved.
22

33
#include <torch/extension.h>
4-
#include "pytorch3d_cutils.h"
4+
#include "utils/pytorch3d_cutils.h"
55

66
#include <vector>
77

‎pytorch3d/csrc/compositing/norm_weighted_sum.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
// Copyright (c) Facebook, Inc. and its affiliates. All rights reserved.
22

33
#include <torch/extension.h>
4-
#include "pytorch3d_cutils.h"
4+
#include "utils/pytorch3d_cutils.h"
55

66
#include <vector>
77

‎pytorch3d/csrc/compositing/weighted_sum.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
// Copyright (c) Facebook, Inc. and its affiliates. All rights reserved.
22

33
#include <torch/extension.h>
4-
#include "pytorch3d_cutils.h"
4+
#include "utils/pytorch3d_cutils.h"
55

66
#include <vector>
77

‎pytorch3d/csrc/ext.cpp

+18
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,8 @@
99
#include "knn/knn.h"
1010
#include "nearest_neighbor_points/nearest_neighbor_points.h"
1111
#include "packed_to_padded_tensor/packed_to_padded_tensor.h"
12+
#include "point_mesh/point_mesh_edge.h"
13+
#include "point_mesh/point_mesh_face.h"
1214
#include "rasterize_meshes/rasterize_meshes.h"
1315
#include "rasterize_points/rasterize_points.h"
1416

@@ -39,4 +41,20 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
3941
m.def("_rasterize_meshes_naive", &RasterizeMeshesNaive);
4042
m.def("_rasterize_meshes_coarse", &RasterizeMeshesCoarse);
4143
m.def("_rasterize_meshes_fine", &RasterizeMeshesFine);
44+
45+
// PointEdge distance functions
46+
m.def("point_edge_dist_forward", &PointEdgeDistanceForward);
47+
m.def("point_edge_dist_backward", &PointEdgeDistanceBackward);
48+
m.def("edge_point_dist_forward", &EdgePointDistanceForward);
49+
m.def("edge_point_dist_backward", &EdgePointDistanceBackward);
50+
m.def("point_edge_array_dist_forward", &PointEdgeArrayDistanceForward);
51+
m.def("point_edge_array_dist_backward", &PointEdgeArrayDistanceBackward);
52+
53+
// PointFace distance functions
54+
m.def("point_face_dist_forward", &PointFaceDistanceForward);
55+
m.def("point_face_dist_backward", &PointFaceDistanceBackward);
56+
m.def("face_point_dist_forward", &FacePointDistanceForward);
57+
m.def("face_point_dist_backward", &FacePointDistanceBackward);
58+
m.def("point_face_array_dist_forward", &PointFaceArrayDistanceForward);
59+
m.def("point_face_array_dist_backward", &PointFaceArrayDistanceBackward);
4260
}

‎pytorch3d/csrc/knn/knn.cu

+2-2
Original file line numberDiff line numberDiff line change
@@ -5,8 +5,8 @@
55
#include <iostream>
66
#include <tuple>
77

8-
#include "dispatch.cuh"
9-
#include "mink.cuh"
8+
#include "utils/dispatch.cuh"
9+
#include "utils/mink.cuh"
1010

1111
// A chunk of work is blocksize-many points of P1.
1212
// The number of potential chunks to do is N*(1+(P1-1)/blocksize)

‎pytorch3d/csrc/knn/knn.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
#pragma once
44
#include <torch/extension.h>
55
#include <tuple>
6-
#include "pytorch3d_cutils.h"
6+
#include "utils/pytorch3d_cutils.h"
77

88
// Compute indices of K nearest neighbors in pointcloud p2 to points
99
// in pointcloud p1.

‎pytorch3d/csrc/nearest_neighbor_points/nearest_neighbor_points.cu

+1-37
Original file line numberDiff line numberDiff line change
@@ -2,43 +2,7 @@
22

33
#include <ATen/ATen.h>
44
#include <float.h>
5-
6-
template <typename scalar_t>
7-
__device__ void WarpReduce(
8-
volatile scalar_t* min_dists,
9-
volatile int64_t* min_idxs,
10-
const size_t tid) {
11-
// s = 32
12-
if (min_dists[tid] > min_dists[tid + 32]) {
13-
min_idxs[tid] = min_idxs[tid + 32];
14-
min_dists[tid] = min_dists[tid + 32];
15-
}
16-
// s = 16
17-
if (min_dists[tid] > min_dists[tid + 16]) {
18-
min_idxs[tid] = min_idxs[tid + 16];
19-
min_dists[tid] = min_dists[tid + 16];
20-
}
21-
// s = 8
22-
if (min_dists[tid] > min_dists[tid + 8]) {
23-
min_idxs[tid] = min_idxs[tid + 8];
24-
min_dists[tid] = min_dists[tid + 8];
25-
}
26-
// s = 4
27-
if (min_dists[tid] > min_dists[tid + 4]) {
28-
min_idxs[tid] = min_idxs[tid + 4];
29-
min_dists[tid] = min_dists[tid + 4];
30-
}
31-
// s = 2
32-
if (min_dists[tid] > min_dists[tid + 2]) {
33-
min_idxs[tid] = min_idxs[tid + 2];
34-
min_dists[tid] = min_dists[tid + 2];
35-
}
36-
// s = 1
37-
if (min_dists[tid] > min_dists[tid + 1]) {
38-
min_idxs[tid] = min_idxs[tid + 1];
39-
min_dists[tid] = min_dists[tid + 1];
40-
}
41-
}
5+
#include "utils/warp_reduce.cuh"
426

437
// CUDA kernel to compute nearest neighbors between two batches of pointclouds
448
// where each point is of dimension D.

‎pytorch3d/csrc/nearest_neighbor_points/nearest_neighbor_points.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22

33
#pragma once
44
#include <torch/extension.h>
5-
#include "pytorch3d_cutils.h"
5+
#include "utils/pytorch3d_cutils.h"
66

77
// Compute indices of nearest neighbors in pointcloud p2 to points
88
// in pointcloud p1.

0 commit comments

Comments
 (0)