58 Commits

Author SHA1 Message Date
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
Jeremy Reizenstein
2361845548 squared distance in comments
Summary: Comments were describing squared distance as absolute distance in a few places.

Reviewed By: nikhilaravi

Differential Revision: D20426020

fbshipit-source-id: 009946867c4a98f61f5ce7158542d41e22bf8346
2020-03-13 04:35:25 -07:00
Nikhila Ravi
d01e722849 Fix coordinate system conventions in point cloud renderer
Summary:
Applying the changes added for mesh rasterization to ensure that +Y is up and +X is left so that the coordinate system is right handed.

Also updated the diagram in the docs to indicate that (0,0) is in the top left hand corner.

Reviewed By: gkioxari

Differential Revision: D20394849

fbshipit-source-id: cfb7c79090eb1f55ad38b92327a74a70a8dc541e
2020-03-12 07:48:29 -07:00
Nikhila Ravi
32ad869dea Update point cloud rasterizer to support heterogeneous point clouds
Summary:
Update the point cloud rasterizer to:
- use the pointcloud datastructure (rebased on top of D19791851.)
- support rasterization of heterogeneous point clouds in the same way as with Meshes.

The main changes to the API will be as follows:
- The input to `rasterize_points` will be a `Pointclouds` object instead of a tensor. This will be easy to update e.g.
```
points = torch.randn(N, P, 3)
idx2, zbuf2, dists2 = rasterize_points(points, image_size, radius, points_per_pixel)

points = torch.randn(N, P, 3)
pointclouds = Pointclouds(points=points)
idx2, zbuf2, dists2 = rasterize_points(pointclouds, image_size, radius, points_per_pixel)
```

- The indices output from rasterization will now refer to points in `poinclouds.points_packed()`.
This may require some changes to the functions which consume the outputs of rasterization if they were previously
assuming that the indices ranged from 0 to P where P is the number of points in each pointcloud.

Making this change now so that Olivia can update her PR accordingly.

Reviewed By: gkioxari

Differential Revision: D20088651

fbshipit-source-id: 833ed659909712bcbbb6a50e2ec0189839f0413a
2020-03-12 07:48:29 -07:00
Nikhila Ravi
15c72be444 Fix coordinate system conventions in renderer
Summary:
## Updates

- Defined the world and camera coordinates according to this figure. The world coordinates are defined as having +Y up, +X left and +Z in.

{F230888499}

- Removed all flipping from blending functions.
- Updated the rasterizer to return images with +Y up and +X left.
- Updated all the mesh rasterizer tests
    - The expected values are now defined in terms of the default +Y up, +X left
    - Added tests where the triangles in the meshes are non symmetrical so that it is clear which direction +X and +Y are

## Questions:
- Should we have **scene settings** instead of raster settings?
    - To be more correct we should be [z clipping in the rasterizer based on the far/near clipping planes](https://github.com/ShichenLiu/SoftRas/blob/master/soft_renderer/cuda/soft_rasterize_cuda_kernel.cu#L400) - these values are also required in the blending functions so should we make these scene level parameters and have a scene settings tuple which is available to the rasterizer and shader?

Reviewed By: gkioxari

Differential Revision: D20208604

fbshipit-source-id: 55787301b1bffa0afa9618f0a0886cc681da51f3
2020-03-06 06:51:05 -08:00
takiyu
f358b9b14d Fix squared distance for CPU impl. (#83)
Summary:
`PointLineDistanceForward()` should return squared distance. However, it seems that it returned non-squared distance when `v0` was near by `v1` in CPU implementation.
Pull Request resolved: https://github.com/facebookresearch/pytorch3d/pull/83

Reviewed By: bottler

Differential Revision: D20097181

Pulled By: nikhilaravi

fbshipit-source-id: 7ea851c0837ab89364e42d283c999df21ff5ff02
2020-02-25 14:00:00 -08:00
merayxu
9e21659fc5 Fixed windows MSVC build compatibility (#9)
Summary:
Fixed a few MSVC compiler (visual studio 2019, MSVC 19.16.27034) compatibility issues
1. Replaced long with int64_t. aten::data_ptr\<long\> is not supported in MSVC
2. pytorch3d/csrc/rasterize_points/rasterize_points_cpu.cpp, inline function is not correctly recognized by MSVC.
3. pytorch3d/csrc/rasterize_meshes/geometry_utils.cuh
const auto kEpsilon = 1e-30;
MSVC does not compile this const into both host and device, change to a MACRO.
4. pytorch3d/csrc/rasterize_meshes/geometry_utils.cuh,
const float area2 = pow(area, 2.0);
2.0 is considered as double by MSVC and raised an error
5. pytorch3d/csrc/rasterize_points/rasterize_points_cpu.cpp
std::tuple<torch::Tensor, torch::Tensor> RasterizePointsCoarseCpu() return type does not match the declaration in rasterize_points_cpu.h.
Pull Request resolved: https://github.com/facebookresearch/pytorch3d/pull/9

Reviewed By: nikhilaravi

Differential Revision: D19986567

Pulled By: yuanluxu

fbshipit-source-id: f4d98525d088c99c513b85193db6f0fc69c7f017
2020-02-20 18:43:19 -08:00