64 Commits

Author SHA1 Message Date
Jeremy Reizenstein
124bb5e391 spelling
Summary: Collection of spelling things, mostly in docs / tutorials.

Reviewed By: gkioxari

Differential Revision: D26101323

fbshipit-source-id: 652f62bc9d71a4ff872efa21141225e43191353a
2021-04-09 09:58:54 -07:00
Nikhila Ravi
13429640d3 Bug fix for case where aspect ratio is a float
Summary:
- Fix the calculation of the non square NDC range when the H and W are not integer multiples.
- Add test for this case

Reviewed By: gkioxari

Differential Revision: D26613213

fbshipit-source-id: df6763cac602e9f1d516b41b432c4d2cfbaa356d
2021-02-24 10:07:17 -08:00
Jeremy Reizenstein
4bfe7158b1 mesh_normal_consistency speedup
Summary: One step in finding all the pairs of vertices which share faces is a simple calculation but annoying to parallelize. It was implemented in pure Python. We move it to C++. We still pull the data to the CPU and put the answer back on the device.

Reviewed By: nikhilaravi, gkioxari

Differential Revision: D26073475

fbshipit-source-id: ffbf4e2c347a511ab5084bceff600465812b6a52
2021-02-11 13:56:17 -08:00
Jeremy Reizenstein
5ac2f42184 test & compilation fixes
Summary:
Fixes mostly related to the "main" build on circleci.
-Avoid error to do with tuple copy from initializer_list which is `explicit` on old compiler.
-Add better reporting to copyright test.
-Move to PackedTensorAccessor64 from the deprecated PackedTensorAccessor
-Avoid some warnings about mismatched comparisons.

The "main" build is the only one that runs the test_build stuff. In that area
-Fix my bad copyright fix D26275931 (3463f418b8) / 965c9c
-Add test that all tutorials are valid json.

Reviewed By: nikhilaravi

Differential Revision: D26366466

fbshipit-source-id: c4ab8b7e6647987069f7cb7144aa6ab7c24bcdac
2021-02-11 11:06:08 -08:00
Nikhila Ravi
340662e98e CUDA/C++ Rasterizer updates to handle clipped faces
Summary:
- Updated the C++/CUDA mesh rasterization kernels to handle the clipped faces. In particular this required careful handling of the distance calculation for faces which are cut into a quadrilateral by the image plane and then split into two sub triangles i.e. both sub triangles can't be part of the top K faces.
- Updated `rasterize_meshes.py` to use the utils functions to clip the meshes and convert the fragments back to in terms of the unclipped mesh
- Added end to end tests

Reviewed By: jcjohnson

Differential Revision: D26169685

fbshipit-source-id: d64cd0d656109b965f44a35c301b7c81f451cfa0
2021-02-08 14:32:39 -08:00
Penn
e58a730e6a Fix dimension check (#524)
Summary:
Fixes the assertion that `p1` and `p2` have the same last dimension. The issue was that `D` is set to equal `p2.size(2)`, and then `D` is compared to `p2.size(2)`. The fix instead compares `D` to `p1.size(2).

Pull Request resolved: https://github.com/facebookresearch/pytorch3d/pull/524

Reviewed By: bottler

Differential Revision: D26008688

Pulled By: nikhilaravi

fbshipit-source-id: e32afe9da127d81b1a411d3c223b539a7400597b
2021-01-22 11:05:18 -08:00
Jeremy Reizenstein
4711665edb lint
Summary: Fix recent lint.

Reviewed By: nikhilaravi

Differential Revision: D25900168

fbshipit-source-id: 6b6e8d35b68c8415ef305dc4719f43eda9316c8f
2021-01-20 13:08:35 -08:00
Nikhila Ravi
3d769a66cb Non Square image rasterization for pointclouds
Summary:
Similar to non square image rasterization for meshes, apply the same updates to the pointcloud rasterizer.

Main API Change:
- PointRasterizationSettings now accepts a tuple/list of (H, W) for the image size.

Reviewed By: jcjohnson

Differential Revision: D25465206

fbshipit-source-id: 7370d83c431af1b972158cecae19d82364623380
2020-12-15 14:15:32 -08:00
Nikhila Ravi
d07307a451 Non square image rasterization for meshes
Summary:
There are a couple of options for supporting non square images:
1) NDC stays at [-1, 1] in both directions with the distance calculations all modified by (W/H). There are a lot of distance based calculations (e.g. triangle areas for barycentric coordinates etc) so this requires changes in many places.
2) NDC is scaled by (W/H) so the smallest side has [-1, 1]. In this case none of the distance calculations need to be updated and only the pixel to NDC calculation needs to be modified.

I decided to go with option 2 after trying option 1!

API Changes:
- Image size can now be specified optionally as a tuple

TODO:
- add a benchmark test for the non square case.

Reviewed By: jcjohnson

Differential Revision: D24404975

fbshipit-source-id: 545efb67c822d748ec35999b35762bce58db2cf4
2020-12-09 09:18:11 -08:00
Christoph Lassner
faed5405c8 Fix #442.
Summary: This fixed #442 by declaring two math functions to be device-only.

Reviewed By: bottler

Differential Revision: D24896992

fbshipit-source-id: a15918d06d2a3e6ee5cf250fec7af5f2f50a6164
2020-11-11 14:06:21 -08:00
Jeremy Reizenstein
d220ee2f66 pulsar build and CI changes
Summary:
Changes to CI and some minor fixes now that pulsar is part of pytorch3d. Most significantly, add CUB to CI builds.

Make CUB_HOME override the CUB already in cudatoolkit (important for cuda11.0 which uses cub 1.9.9 which pulsar doesn't work well with.
Make imageio available for testing.
Lint fixes.
Fix some test verbosity.
Avoid use of atomicAdd_block on older GPUs.

Reviewed By: nikhilaravi, classner

Differential Revision: D24773716

fbshipit-source-id: 2428356bb2e62735f2bc0c15cbe4cff35b1b24b8
2020-11-10 09:38:05 -08:00
Dave Schnizlein
804235b05a Remove point mesh edge kernels
Summary:
Removes the now-unnecessary kernels from point mesh edge file

Migrates all point mesh functionality into one file.

Reviewed By: gkioxari

Differential Revision: D24550086

fbshipit-source-id: f924996cd38a7c2c1cf189d8a01611de4506cfa3
2020-11-10 09:34:16 -08:00
Dave Schnizlein
8dcfe30f66 Consolidate mesh backward kernels
Summary: This diff creates the generic MeshBackwardKernel which can handle distance calculations between point, edge and faces in either direction. Replaces only point_mesh_face code for now.

Reviewed By: gkioxari

Differential Revision: D24549374

fbshipit-source-id: 2853c1da1c2a6b6de8d0e40007ba0735b8959044
2020-11-10 09:34:16 -08:00
Dave Schnizlein
c41aff23f0 Consolidate point mesh forward kernels
Summary: This diff creates the generic MeshForwardKernel which can handle distance calculations between point, edge and faces in either direction. Replaces only point_mesh_face code for now.

Reviewed By: gkioxari

Differential Revision: D24543316

fbshipit-source-id: 302707d7cec2d77a899738adf40481035c240da8
2020-11-10 09:34:16 -08:00
Christoph Lassner
194b29fb6c Fix #431.
Summary: Added missing include for cstdint for Windows and removed problematic inline assembly.

Reviewed By: bottler

Differential Revision: D24838053

fbshipit-source-id: 95496be841c2c22a82068073d4740e98ee8a02ac
2020-11-09 13:25:09 -08:00
Christoph Lassner
039e02601d examples and docs.
Summary: This diff updates the documentation and tutorials with information about the new pulsar backend. For more information about the pulsar backend, see the release notes and the paper (https://arxiv.org/abs/2004.07484). For information on how to use the backend, see the point cloud rendering notebook and the examples in the folder docs/examples.

Reviewed By: nikhilaravi

Differential Revision: D24498129

fbshipit-source-id: e312b0169a72b13590df6e4db36bfe6190d742f9
2020-11-03 13:06:35 -08:00
Christoph Lassner
b19fe1de2f pulsar integration.
Summary:
This diff integrates the pulsar renderer source code into PyTorch3D as an alternative backend for the PyTorch3D point renderer. This diff is the first of a series of three diffs to complete that migration and focuses on the packaging and integration of the source code.

For more information about the pulsar backend, see the release notes and the paper (https://arxiv.org/abs/2004.07484). For information on how to use the backend, see the point cloud rendering notebook and the examples in the folder `docs/examples`.

Tasks addressed in the following diffs:
* Add the PyTorch3D interface,
* Add notebook examples and documentation (or adapt the existing ones to feature both interfaces).

Reviewed By: nikhilaravi

Differential Revision: D23947736

fbshipit-source-id: a5e77b53e6750334db22aefa89b4c079cda1b443
2020-11-03 13:06:35 -08:00
Nikhila Ravi
ebe2693b11 Support variable size radius for points in rasterizer
Summary:
Support variable size pointclouds in the renderer API to allow compatibility with Pulsar rasterizer.

If radius is provided as a float, it is converted to a tensor of shape (P). Otherwise radius is expected to be an (N, P_padded) dimensional tensor where P_padded is the max number of points in the batch (following the convention from pulsar: https://our.intern.facebook.com/intern/diffusion/FBS/browse/master/fbcode/frl/gemini/pulsar/pulsar/renderer.py?commit=ee0342850210e5df441e14fd97162675c70d147c&lines=50)

Reviewed By: jcjohnson, gkioxari

Differential Revision: D21429400

fbshipit-source-id: 65de7d9cd2472b27fc29f96160c33687e88098a2
2020-09-18 18:48:18 -07:00
z003yctd
e40c2167ae fix incorrect variable naming (#362)
Summary: Pull Request resolved: https://github.com/facebookresearch/pytorch3d/pull/362

Reviewed By: bottler

Differential Revision: D23712242

Pulled By: nikhilaravi

fbshipit-source-id: 1c4184c8482049991356be7dbc9755b0c2018a1d
2020-09-17 16:50:34 -07:00
Jeremy Reizenstein
9a50cf800e Fix batching bug from TexturesUV packed ambiguity, other textures tidyup
Summary:
faces_uvs_packed and verts_uvs_packed were only used in one place and the definition of the former was ambiguous. This meant that the wrong coordinates could be used for meshes other than the first in the batch. I have therefore removed both functions and build their common result inline. Added a test that a simple batch of two meshes is rendered consistently with the rendering of each alone. This test would have failed before.

I hope this fixes https://github.com/facebookresearch/pytorch3d/issues/283.

Some other small improvements to the textures code.

Reviewed By: nikhilaravi

Differential Revision: D23161936

fbshipit-source-id: f99b560a46f6b30262e07028b049812bc04350a7
2020-08-21 05:53:29 -07:00
Steve Branson
9aaba0483c Temporary fix for mesh rasterization bug for traingles partially behind the camera
Summary: A triangle is culled if any vertex in a triangle is behind the camera.  This fixes incorrect rendering of triangles that are partially behind the camera, where screen coordinate calculations are strange.  It doesn't work for triangles that are partially behind the camera but still intersect with the view frustum.

Reviewed By: nikhilaravi

Differential Revision: D22856181

fbshipit-source-id: a9cbaa1327d89601b83d0dfd3e4a04f934a4a213
2020-08-20 22:24:19 -07:00
Jeremy Reizenstein
7944d24d48 gather_scatter on CPU
Summary: CPU implementation of the graph convolution op.

Reviewed By: nikhilaravi, gkioxari

Differential Revision: D21384361

fbshipit-source-id: bc96730e9727bb9aa1b0a232dcb82f0c0d12fe6b
2020-08-05 07:00:20 -07:00
Nikhila Ravi
cc70950f40 barycentric clipping in cuda/c++
Summary:
Added support for barycentric clipping in the C++/CUDA rasterization kernels which can be switched on/off via a rasterization setting.

Added tests and a benchmark to compare with the current implementation in PyTorch - for some cases of large image size/faces per pixel the cuda version is 10x faster.

Reviewed By: gkioxari

Differential Revision: D21705503

fbshipit-source-id: e835c0f927f1e5088ca89020aef5ff27ac3a8769
2020-07-16 10:17:28 -07:00
Nikhila Ravi
bce396df93 C++/CUDA implementation of sigmoid alpha blend
Summary:
C++/CUDA implementation of forward and backward passes for the sigmoid alpha blending function.

This is slightly faster than the vectorized implementation in Python, but more importantly uses less memory due to fewer tensors being created.

Reviewed By: gkioxari

Differential Revision: D19980671

fbshipit-source-id: 0779055d2c68b1f20fb0870e60046077ef4613ff
2020-07-16 10:17:28 -07:00
Justin Johnson
26d2cc24c1 CUDA kernel for interpolate_face_attributes
Summary: When rendering meshes with Phong shading, interpolate_face_attributes was taking up a nontrivial fraction of the overall shading time. This diff replaces our Python implementation of this function with a CUDA implementation.

Reviewed By: nikhilaravi

Differential Revision: D21610763

fbshipit-source-id: 2bb362a28f698541812aeab539047264b125ebb8
2020-07-13 12:59:37 -07:00
Jeremy Reizenstein
2f6387f239 Restore C++14 compatibility
Summary: Fix the new CPU implementation of point_mesh functionality to be compatible with older C++.

Reviewed By: nikhilaravi

Differential Revision: D22066785

fbshipit-source-id: a245849342019a93ff68e186a10985458b007436
2020-06-16 14:19:21 -07:00
Jeremy Reizenstein
74659aef26 CPU implementation for point_mesh functions
Summary:
point_mesh functions were missing CPU implementations.
The indices returned are not always matching, possibly due to numerical instability.

Reviewed By: gkioxari

Differential Revision: D21594264

fbshipit-source-id: 3016930e2a9a0f3cd8b3ac4c94a92c9411c0989d
2020-06-15 10:11:26 -07:00
Georgia Gkioxari
d689baac5e fix alpha compositing
Summary:
Fix division by zero when alpha is 1.0
In this case, the nominator is already 0 and we need to make sure division with 0 does not occur which would produce nans

Reviewed By: nikhilaravi

Differential Revision: D21650478

fbshipit-source-id: bc457105b3050fef1c8bd4e58e7d6d15c0c81ffd
2020-05-20 09:27:42 -07:00
Nikhila Ravi
3fef506895 Make cuda tensors contiguous in host function and remove contiguous check
Summary:
Update the cuda kernels to:
- remove contiguous checks for the grad tensors and for cpu functions which use accessors
- for cuda implementations call `.contiguous()` on all tensors in the host function before invoking the kernel

Reviewed By: gkioxari

Differential Revision: D21598008

fbshipit-source-id: 9b97bda4582fd4269c8a00999874d4552a1aea2d
2020-05-15 15:00:25 -07:00
Jeremy Reizenstein
728179e848 avoid converting a TensorOptions from float to integer
Summary: pytorch is adding checks that mean integer tensors with requires_grad=True need to be avoided. Fix accidentally creating them.

Reviewed By: jcjohnson, gkioxari

Differential Revision: D21576712

fbshipit-source-id: 008218997986800a36d93caa1a032ee91f2bffcd
2020-05-14 13:16:05 -07:00
Jeremy Reizenstein
8fc28baa27 Looser gradient check in test_rasterize_meshes
Summary: This has been failing intermittently

Reviewed By: nikhilaravi

Differential Revision: D21403157

fbshipit-source-id: 51b74d6c813b52effe72d14b565e250fcabbb463
2020-05-05 09:26:47 -07:00
Michele Sanna
f8acecb6b3 a formula for bin size for images over 64x64 (#90)
Summary:
Signed-off-by: Michele Sanna <sanna@arrival.com>

fixes the bin_size calculation with a formula for any image_size > 64. Matches the values chosen so far.

simple test:

```
import numpy as np
import matplotlib.pyplot as plt

image_size = np.arange(64, 2048)
bin_size = np.where(image_size <= 64, 8, (2 ** np.maximum(np.ceil(np.log2(image_size)) - 4, 4)).astype(int))

print(image_size)
print(bin_size)

for ims, bins in zip(image_size, bin_size):
    if ims <= 64:
        assert bins == 8
    elif ims <= 256:
        assert bins == 16
    elif ims <= 512:
        assert bins == 32
    elif ims <= 1024:
        assert bins == 64
    elif ims <= 2048:
        assert bins == 128

    assert (ims + bins - 1) // bins < 22

plt.plot(image_size, bin_size)
plt.grid()
plt.show()
```

![img](https://user-images.githubusercontent.com/54891577/75464693-795bcf00-597f-11ea-9061-26440211691c.png)
Pull Request resolved: https://github.com/facebookresearch/pytorch3d/pull/90

Reviewed By: jcjohnson

Differential Revision: D21160372

Pulled By: nikhilaravi

fbshipit-source-id: 660cf5832f4ca5be243c435a6bed969596fc0188
2020-04-24 14:56:41 -07:00
Nikhila Ravi
c3d636dc8c Cuda updates
Summary:
Updates to:
- enable cuda kernel launches on any GPU (not just the default)
- cuda and contiguous checks for all kernels
- checks to ensure all tensors are on the same device
- error reporting in the cuda kernels
- cuda tests now run on a random device not just the default

Reviewed By: jcjohnson, gkioxari

Differential Revision: D21215280

fbshipit-source-id: 1bedc9fe6c35e9e920bdc4d78ed12865b1005519
2020-04-24 09:11:04 -07:00
Jeremy Reizenstein
85c396f822 avoid using torch/extension.h in cuda
Summary:
Use aten instead of torch interface in all cuda code. This allows the cuda build to work with pytorch 1.5 with GCC 5 (e.g. the compiler of ubuntu 16.04LTS). This wasn't working. It has been failing with errors like the below, perhaps due to a bug in nvcc.

```
torch/include/torch/csrc/api/include/torch/nn/cloneable.h:68:61: error: invalid static_cast from type ‘const torch::OrderedDict<std::basic_string<char>, std::shared_ptr<torch::nn::Module> >’ to type ‘torch::OrderedDict<std::basic_string<char>, std::shared_ptr<torch::nn::Module> >
```

Reviewed By: nikhilaravi

Differential Revision: D21204029

fbshipit-source-id: ca6bdbcecf42493365e1c23a33fe35e1759fe8b6
2020-04-23 10:26:17 -07:00
Justin Johnson
9f31a4fd46 Expose knn_check_version in python
Summary:
We have multiple KNN CUDA implementations. From python, users can currently request a particular implementation via the `version` flag, but they have no way of knowing which implementations can be used for a given problem.

This diff exposes a function `pytorch3d._C.knn_check_version(version, D, K)` that returns whether a particular version can be used.

Reviewed By: nikhilaravi

Differential Revision: D21162573

fbshipit-source-id: 6061960bdcecba454fd920b00036f4e9ff3fdbc0
2020-04-22 14:30:52 -07:00
Nikhila Ravi
4bf30593ff back face culling in rasterization
Summary:
Added backface culling as an option to the `raster_settings`. This is needed for the full forward rendering of shapenet meshes with texture (some meshes contain
multiple overlapping segments which have different textures).

For a triangle (v0, v1, v2) define the vectors A = (v1 - v0) and B = (v2 − v0) and use this to calculate the area of the triangle as:
```
area = 0.5 * A  x B
area = 0.5 * ((x1 − x0)(y2 − y0) − (x2 − x0)(y1 − y0))
```
The area will be positive if (v0, v1, v2) are oriented counterclockwise (a front face), and negative if (v0, v1, v2) are oriented clockwise (a back face).

We can reuse the `edge_function` as it already calculates the triangle area.

Reviewed By: jcjohnson

Differential Revision: D20960115

fbshipit-source-id: 2d8a4b9ccfb653df18e79aed8d05c7ec0f057ab1
2020-04-22 08:22:46 -07:00
Nikhila Ravi
9ef1ee8455 coarse rasterization bug fix
Summary:
Fix a bug which resulted in a rendering artifacts if the image size was not a multiple of 16.
Fix: Revert coarse rasterization to original implementation and only update fine rasterization to reverse the ordering of Y and X axis. This is much simpler than the previous approach!

Additional changes:
- updated mesh rendering end-end tests to check outputs from both naive and coarse to fine rasterization.
- added pointcloud rendering end-end tests

Reviewed By: gkioxari

Differential Revision: D21102725

fbshipit-source-id: 2e7e1b013dd6dd12b3a00b79eb8167deddb2e89a
2020-04-20 14:54:16 -07:00
Jeremy Reizenstein
6207c359b1 spelling and flake
Summary: mostly recent lintish things

Reviewed By: nikhilaravi

Differential Revision: D21089003

fbshipit-source-id: 028733c1d875268f1879e4481da475b7100ba0b6
2020-04-17 10:50:22 -07:00
Jeremy Reizenstein
9397cd872d torch C API warnings
Summary: This is mostly replacing the old PackedTensorAccessor with the new PackedTensorAccessor64.

Reviewed By: gkioxari

Differential Revision: D21088773

fbshipit-source-id: 5973e5a29d934eafb7c70ec5ec154ca076b64d27
2020-04-17 10:46:31 -07:00
Jeremy Reizenstein
e19df58766 remove final nearest_neighbor files
Summary: A couple of files for the removed nearest_neighbor functionality are left behind.

Reviewed By: nikhilaravi

Differential Revision: D21088624

fbshipit-source-id: 4bb29016b4e5f63102765b384c363733b60032fa
2020-04-17 09:27:17 -07:00
Nikhila Ravi
3794f6753f remove nearest_neighbors
Summary: knn is more general and faster than the nearest_neighbor code, so remove the latter.

Reviewed By: gkioxari

Differential Revision: D20816424

fbshipit-source-id: 75d6c44d17180752d0c9859814bbdf7892558158
2020-04-15 20:51:41 -07:00
Georgia Gkioxari
b2b0c5a442 knn autograd
Summary:
Adds knn backward to return `grad_pts1` and `grad_pts2`. Adds `knn_gather` to return the nearest neighbors in pts2.

The BM tests include backward pass and are ran on an M40.
```
Benchmark                               Avg Time(μs)      Peak Time(μs) Iterations
--------------------------------------------------------------------------------
KNN_SQUARE_32_256_128_3_24_cpu              39558           43485             13
KNN_SQUARE_32_256_128_3_24_cuda:0            1080            1404            463
KNN_SQUARE_32_256_512_3_24_cpu              81950           85781              7
KNN_SQUARE_32_256_512_3_24_cuda:0            1519            1641            330
--------------------------------------------------------------------------------

Benchmark                               Avg Time(μs)      Peak Time(μs) Iterations
--------------------------------------------------------------------------------
KNN_RAGGED_32_256_128_3_24_cpu              13798           14650             37
KNN_RAGGED_32_256_128_3_24_cuda:0            1576            1713            318
KNN_RAGGED_32_256_512_3_24_cpu              31255           32210             16
KNN_RAGGED_32_256_512_3_24_cuda:0            2024            2162            248
--------------------------------------------------------------------------------
```

Reviewed By: jcjohnson

Differential Revision: D20945556

fbshipit-source-id: a16f616029c6b5f8c2afceb5f2bc12c5c20d2f3c
2020-04-14 17:22:56 -07:00
Georgia Gkioxari
487d4d6607 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
2020-04-11 00:21:24 -07:00
Jeremy Reizenstein
01b5f7b228 heterogenous KNN
Summary: Interface and working implementation of ragged KNN. Benchmarks (which aren't ragged) haven't slowed. New benchmark shows that ragged is faster than non-ragged of the same shape.

Reviewed By: jcjohnson

Differential Revision: D20696507

fbshipit-source-id: 21b80f71343a3475c8d3ee0ce2680f92f0fae4de
2020-04-07 01:47:37 -07:00
Jeremy Reizenstein
37c5c8e0b6 Linter, deprecated type()
Summary: Run linter after recent changes. Fix long comment in knn.h which clang-format has reflowed badly. Add crude test that code doesn't call deprecated `.type()` or `.data()`.

Reviewed By: nikhilaravi

Differential Revision: D20692935

fbshipit-source-id: 28ce0308adae79a870cb41a810b7cf8744f41ab8
2020-03-29 14:02:58 -07:00
Justin Johnson
870290df34 Implement K-Nearest Neighbors
Summary:
Implements K-Nearest Neighbors with C++ and CUDA versions.

KNN in CUDA is highly nontrivial. I've implemented a few different versions of the kernel, and we heuristically dispatch to different kernels based on the problem size. Some of the kernels rely on template specialization on either D or K, so we use template metaprogramming to compile specialized versions for ranges of D and K.

These kernels are up to 3x faster than our existing 1-nearest-neighbor kernels, so we should also consider swapping out `nn_points_idx` to use these kernels in the backend.

I've been working mostly on the CUDA kernels, and haven't converged on the correct Python API.

I still want to benchmark against FAISS to see how far away we are from their performance.

Reviewed By: bottler

Differential Revision: D19729286

fbshipit-source-id: 608ffbb7030c21fe4008f330522f4890f0c3c21a
2020-03-26 13:40:26 -07:00
Jeremy Reizenstein
81a4aa18ad type() deprecated
Summary:
Replace `tensor.type().is_cuda()` with the preferred `tensor.is_cuda()`.
Replace `AT_DISPATCH_FLOATING_TYPES(tensor.type(), ...` with `AT_DISPATCH_FLOATING_TYPES(tensor.scalar_type(), ...`.
These avoid deprecation warnings in future pytorch.

Reviewed By: nikhilaravi

Differential Revision: D20646565

fbshipit-source-id: 1a0c15978c871af816b1dd7d4a7ea78242abd95e
2020-03-26 04:01:41 -07:00
Jeremy Reizenstein
e22d431e5b data() deprecated
Summary: replace `data()` with preferred `data_ptr()`, avoiding some deprecation warnings in future pytorch.

Reviewed By: nikhilaravi

Differential Revision: D20645738

fbshipit-source-id: 8f6e02d292729b804fa2a66f94dd0517bbaf7887
2020-03-26 03:21:48 -07:00
Jeremy Reizenstein
8fa7678614 fix CPU-only hiding of cuda calls
Summary: CPU-only builds should be fixed by this change

Reviewed By: nikhilaravi

Differential Revision: D20598014

fbshipit-source-id: df098ec4c6c93d38515172805fe57cac7463c506
2020-03-24 05:04:32 -07:00
Olivia
53599770dd Accumulate points (#4)
Summary:
Code for accumulating points in the z-buffer in three ways:
1. weighted sum
2. normalised weighted sum
3. alpha compositing

Pull Request resolved: https://github.com/fairinternal/pytorch3d/pull/4

Reviewed By: nikhilaravi

Differential Revision: D20522422

Pulled By: gkioxari

fbshipit-source-id: 5023baa05f15e338f3821ef08f5552c2dcbfc06c
2020-03-19 11:23:12 -07:00