diff --git a/.gitattributes b/.gitattributes index a6344aac8c09253b3b630fb776ae94478aa0275b..b08bbdd9522589d0deb34d17c2affc739aa2d4ed 100644 --- a/.gitattributes +++ b/.gitattributes @@ -33,3 +33,44 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text *.zip filter=lfs diff=lfs merge=lfs -text *.zst filter=lfs diff=lfs merge=lfs -text *tfevents* filter=lfs diff=lfs merge=lfs -text +packages/open3d/debugpy-1.8.1-py2.py3-none-any.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/fonttools-4.53.0-py3-none-any.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/jedi-0.19.1-py2.py3-none-any.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/kiwisolver-1.4.5-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/matplotlib-3.5.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/nbclassic-1.1.0-py3-none-any.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/numpy-1.21.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/open3d-0.10.0.0-cp38-cp38-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/Pillow-8.3.2-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/pygments-2.18.0-py3-none-any.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/pyzmq-20.0.0-cp38-cp38-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/open3d/widgetsnbextension-4.0.11-py3-none-any.whl filter=lfs diff=lfs merge=lfs -text +packages/scikit-learn/numpy-1.21.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/scikit-learn/scikit_learn-1.0.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/scikit-learn/scipy-1.7.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/networkx-3.1-py3-none-any.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/numpy-1.24.4-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_cublas_cu12-12.1.3.1-py3-none-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_cuda_cupti_cu12-12.1.105-py3-none-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_cuda_nvrtc_cu12-12.1.105-py3-none-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_cudnn_cu12-8.9.2.26-py3-none-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_cufft_cu12-11.0.2.54-py3-none-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_curand_cu12-10.3.2.106-py3-none-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_cusolver_cu12-11.4.5.107-py3-none-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_cusparse_cu12-12.1.0.106-py3-none-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_nccl_cu12-2.18.1-py3-none-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/nvidia_nvjitlink_cu12-12.5.40-py3-none-manylinux2014_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/pillow-10.3.0-cp38-cp38-manylinux_2_28_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/sympy-1.12.1-py3-none-any.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/torch-2.1.0-cp38-cp38-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/torchaudio-2.1.0-cp38-cp38-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/torchvision-0.16.0-cp38-cp38-manylinux1_x86_64.whl filter=lfs diff=lfs merge=lfs -text +packages/torch/triton-2.1.0-0-cp38-cp38-manylinux2014_x86_64.manylinux_2_17_x86_64.whl filter=lfs diff=lfs merge=lfs -text +pc_util/build/lib.linux-x86_64-cpython-38/pc_util.cpython-38-x86_64-linux-gnu.so filter=lfs diff=lfs merge=lfs -text +pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query.o filter=lfs diff=lfs merge=lfs -text +pc_util/build/temp.linux-x86_64-cpython-38/src/cluster.o filter=lfs diff=lfs merge=lfs -text +pc_util/build/temp.linux-x86_64-cpython-38/src/group_points.o filter=lfs diff=lfs merge=lfs -text +pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate.o filter=lfs diff=lfs merge=lfs -text +pc_util/build/temp.linux-x86_64-cpython-38/src/pointnet2_api.o filter=lfs diff=lfs merge=lfs -text +pc_util/build/temp.linux-x86_64-cpython-38/src/sampling.o filter=lfs diff=lfs merge=lfs -text +pc_util/dist/pc_util-1.0-py3.8-linux-x86_64.egg filter=lfs diff=lfs merge=lfs -text diff --git a/packages/easydict/easydict-1.13-py3-none-any.whl b/packages/easydict/easydict-1.13-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..58cef95ac8c7bcba2b324f46d42fc167de1f992d Binary files /dev/null and b/packages/easydict/easydict-1.13-py3-none-any.whl differ diff --git a/packages/open3d/MarkupSafe-2.0.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.manylinux_2_12_x86_64.manylinux2010_x86_64.whl b/packages/open3d/MarkupSafe-2.0.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.manylinux_2_12_x86_64.manylinux2010_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..a2e1a20f133608cb1ac59d37d4e50a634316febb Binary files /dev/null and b/packages/open3d/MarkupSafe-2.0.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.manylinux_2_12_x86_64.manylinux2010_x86_64.whl differ diff --git a/packages/open3d/Pillow-8.3.2-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl b/packages/open3d/Pillow-8.3.2-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..75bd095594e0924d16d15e678701bae587c360a8 --- /dev/null +++ b/packages/open3d/Pillow-8.3.2-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:0b9911ec70731711c3b6ebcde26caea620cbdd9dcb73c67b0730c8817f24711b +size 3037045 diff --git a/packages/open3d/Send2Trash-1.8.3-py3-none-any.whl b/packages/open3d/Send2Trash-1.8.3-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..196344b82c518738376938c764b74087943d8e5e Binary files /dev/null and b/packages/open3d/Send2Trash-1.8.3-py3-none-any.whl differ diff --git a/packages/open3d/anyio-3.7.1-py3-none-any.whl b/packages/open3d/anyio-3.7.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..98b7d6dc0dc7297b4cd0863abd094814e850a687 Binary files /dev/null and b/packages/open3d/anyio-3.7.1-py3-none-any.whl differ diff --git a/packages/open3d/argon2_cffi-21.1.0-cp35-abi3-manylinux_2_5_x86_64.manylinux1_x86_64.whl b/packages/open3d/argon2_cffi-21.1.0-cp35-abi3-manylinux_2_5_x86_64.manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..25243be9b446400bcf1a1e89224c98c96f5cf7dd Binary files /dev/null and b/packages/open3d/argon2_cffi-21.1.0-cp35-abi3-manylinux_2_5_x86_64.manylinux1_x86_64.whl differ diff --git a/packages/open3d/asttokens-2.4.1-py2.py3-none-any.whl b/packages/open3d/asttokens-2.4.1-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..3af6a4f59e8f93e9f951b142f0b4c736e901f1d2 Binary files /dev/null and b/packages/open3d/asttokens-2.4.1-py2.py3-none-any.whl differ diff --git a/packages/open3d/attrs-23.2.0-py3-none-any.whl b/packages/open3d/attrs-23.2.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..db9a7314d3559272f2b2a9b326f9b5473dfd21a5 Binary files /dev/null and b/packages/open3d/attrs-23.2.0-py3-none-any.whl differ diff --git a/packages/open3d/backcall-0.2.0-py2.py3-none-any.whl b/packages/open3d/backcall-0.2.0-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..1a384ae953818684e16a393cd181646b35859fc7 Binary files /dev/null and b/packages/open3d/backcall-0.2.0-py2.py3-none-any.whl differ diff --git a/packages/open3d/beautifulsoup4-4.12.3-py3-none-any.whl b/packages/open3d/beautifulsoup4-4.12.3-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..f0f96070ffa9272e2e49383ff22551304bce4ae0 Binary files /dev/null and b/packages/open3d/beautifulsoup4-4.12.3-py3-none-any.whl differ diff --git a/packages/open3d/bleach-6.1.0-py3-none-any.whl b/packages/open3d/bleach-6.1.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..2a5053e63925b8f74672410d206382fa17116934 Binary files /dev/null and b/packages/open3d/bleach-6.1.0-py3-none-any.whl differ diff --git a/packages/open3d/cffi-1.14.6-cp38-cp38-manylinux1_x86_64.whl b/packages/open3d/cffi-1.14.6-cp38-cp38-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..d6903e13e465bc23989777c6fa5f49fded0ec763 Binary files /dev/null and b/packages/open3d/cffi-1.14.6-cp38-cp38-manylinux1_x86_64.whl differ diff --git a/packages/open3d/comm-0.2.2-py3-none-any.whl b/packages/open3d/comm-0.2.2-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..ae19abe6c23e71de9c45aa145f95c06edd0283e5 Binary files /dev/null and b/packages/open3d/comm-0.2.2-py3-none-any.whl differ diff --git a/packages/open3d/cycler-0.12.1-py3-none-any.whl b/packages/open3d/cycler-0.12.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..6478c3f06b70e7486432ecf009461ab2365037b2 Binary files /dev/null and b/packages/open3d/cycler-0.12.1-py3-none-any.whl differ diff --git a/packages/open3d/debugpy-1.8.1-py2.py3-none-any.whl b/packages/open3d/debugpy-1.8.1-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..918322b9f7e76f63a9254a332f96a0bb2d1bb3f3 --- /dev/null +++ b/packages/open3d/debugpy-1.8.1-py2.py3-none-any.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:28acbe2241222b87e255260c76741e1fbf04fdc3b6d094fcf57b6c6f75ce1242 +size 4832569 diff --git a/packages/open3d/decorator-5.1.1-py3-none-any.whl b/packages/open3d/decorator-5.1.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..9ccfbe25f8312938d071ce3020897e0f1ffe64a0 Binary files /dev/null and b/packages/open3d/decorator-5.1.1-py3-none-any.whl differ diff --git a/packages/open3d/defusedxml-0.7.1-py2.py3-none-any.whl b/packages/open3d/defusedxml-0.7.1-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..8e678bf9f15687883935f15648845c735c8d63c0 Binary files /dev/null and b/packages/open3d/defusedxml-0.7.1-py2.py3-none-any.whl differ diff --git a/packages/open3d/entrypoints-0.4-py3-none-any.whl b/packages/open3d/entrypoints-0.4-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..9ed0c065ecdecaa0864ab4e1ce5700b3e1bdb4fc Binary files /dev/null and b/packages/open3d/entrypoints-0.4-py3-none-any.whl differ diff --git a/packages/open3d/exceptiongroup-1.2.1-py3-none-any.whl b/packages/open3d/exceptiongroup-1.2.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..494f2626e8221a08ed97b10c6f15db1ffe1649e3 Binary files /dev/null and b/packages/open3d/exceptiongroup-1.2.1-py3-none-any.whl differ diff --git a/packages/open3d/executing-2.0.1-py2.py3-none-any.whl b/packages/open3d/executing-2.0.1-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..8de030b44d7ad9efbbb45e692eacbd52c864bf67 Binary files /dev/null and b/packages/open3d/executing-2.0.1-py2.py3-none-any.whl differ diff --git a/packages/open3d/fastjsonschema-2.19.1-py3-none-any.whl b/packages/open3d/fastjsonschema-2.19.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..86c72030975fe8e5317f997aeba1af5b2c231a21 Binary files /dev/null and b/packages/open3d/fastjsonschema-2.19.1-py3-none-any.whl differ diff --git a/packages/open3d/fonttools-4.53.0-py3-none-any.whl b/packages/open3d/fonttools-4.53.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..dfc44983f947d7da8b3505440e031129ea4659fd --- /dev/null +++ b/packages/open3d/fonttools-4.53.0-py3-none-any.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:6b4f04b1fbc01a3569d63359f2227c89ab294550de277fd09d8fca6185669fa4 +size 1090184 diff --git a/packages/open3d/idna-3.7-py3-none-any.whl b/packages/open3d/idna-3.7-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..fa4c95b123751094438c2f747e5389c32d870c5b Binary files /dev/null and b/packages/open3d/idna-3.7-py3-none-any.whl differ diff --git a/packages/open3d/importlib_metadata-7.1.0-py3-none-any.whl b/packages/open3d/importlib_metadata-7.1.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..931fbe03dd7f22acc2ced6a9f01f4e9372da9e2b Binary files /dev/null and b/packages/open3d/importlib_metadata-7.1.0-py3-none-any.whl differ diff --git a/packages/open3d/importlib_resources-6.4.0-py3-none-any.whl b/packages/open3d/importlib_resources-6.4.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..8872d14ac9f82ba9fcef4f401118c8df3e13762d Binary files /dev/null and b/packages/open3d/importlib_resources-6.4.0-py3-none-any.whl differ diff --git a/packages/open3d/ipykernel-6.9.1-py3-none-any.whl b/packages/open3d/ipykernel-6.9.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..feb4be1b954e71ff767206b004502bda9b128081 Binary files /dev/null and b/packages/open3d/ipykernel-6.9.1-py3-none-any.whl differ diff --git a/packages/open3d/ipython-8.12.3-py3-none-any.whl b/packages/open3d/ipython-8.12.3-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..4531e46250a67e16b977fa33f63a810f35917d2d Binary files /dev/null and b/packages/open3d/ipython-8.12.3-py3-none-any.whl differ diff --git a/packages/open3d/ipython_genutils-0.2.0-py2.py3-none-any.whl b/packages/open3d/ipython_genutils-0.2.0-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..2b68d9a2ec96236bf60c7bae4bacbef5dda58ee1 Binary files /dev/null and b/packages/open3d/ipython_genutils-0.2.0-py2.py3-none-any.whl differ diff --git a/packages/open3d/ipywidgets-8.1.3-py3-none-any.whl b/packages/open3d/ipywidgets-8.1.3-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..180ccc46c4d3d7c457e51619f5980cc15c4bfdf3 Binary files /dev/null and b/packages/open3d/ipywidgets-8.1.3-py3-none-any.whl differ diff --git a/packages/open3d/jedi-0.19.1-py2.py3-none-any.whl b/packages/open3d/jedi-0.19.1-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..7cf186c84fb3de85ef7ee2ae9bf0c49ed0282563 --- /dev/null +++ b/packages/open3d/jedi-0.19.1-py2.py3-none-any.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:e983c654fe5c02867aef4cdfce5a2fbb4a50adc0af145f70504238f18ef5e7e0 +size 1569361 diff --git a/packages/open3d/jinja2-3.1.4-py3-none-any.whl b/packages/open3d/jinja2-3.1.4-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..fa6a14bfaa53bb44031daf857127636101dd2409 Binary files /dev/null and b/packages/open3d/jinja2-3.1.4-py3-none-any.whl differ diff --git a/packages/open3d/jsonschema-4.17.3-py3-none-any.whl b/packages/open3d/jsonschema-4.17.3-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..e8f512072ee4e8fde12266ef2cf1c574cdce16a4 Binary files /dev/null and b/packages/open3d/jsonschema-4.17.3-py3-none-any.whl differ diff --git a/packages/open3d/jupyter_client-7.2.0-py3-none-any.whl b/packages/open3d/jupyter_client-7.2.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..23b4ac02b3efd859d96b74de2e74f38e33564e1e Binary files /dev/null and b/packages/open3d/jupyter_client-7.2.0-py3-none-any.whl differ diff --git a/packages/open3d/jupyter_core-5.7.2-py3-none-any.whl b/packages/open3d/jupyter_core-5.7.2-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..f75153e1dbd88a8a3635aa454691ffe4d168408f Binary files /dev/null and b/packages/open3d/jupyter_core-5.7.2-py3-none-any.whl differ diff --git a/packages/open3d/jupyter_server-1.24.0-py3-none-any.whl b/packages/open3d/jupyter_server-1.24.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..baabed63b118604a3490afd4b8015f2d9633b5de Binary files /dev/null and b/packages/open3d/jupyter_server-1.24.0-py3-none-any.whl differ diff --git a/packages/open3d/jupyterlab_pygments-0.3.0-py3-none-any.whl b/packages/open3d/jupyterlab_pygments-0.3.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..b82b1b472c4825a6b41f6810a3319e9db2ba592b Binary files /dev/null and b/packages/open3d/jupyterlab_pygments-0.3.0-py3-none-any.whl differ diff --git a/packages/open3d/jupyterlab_widgets-3.0.11-py3-none-any.whl b/packages/open3d/jupyterlab_widgets-3.0.11-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..4ba7de9157fa617abc11063c9c11045ca404b4b7 Binary files /dev/null and b/packages/open3d/jupyterlab_widgets-3.0.11-py3-none-any.whl differ diff --git a/packages/open3d/kiwisolver-1.4.5-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl b/packages/open3d/kiwisolver-1.4.5-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..6814bb300733d95cb0bcf5314ead99b467b0a47d --- /dev/null +++ b/packages/open3d/kiwisolver-1.4.5-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:d2e5a98f0ec99beb3c10e13b387f8db39106d53993f498b295f0c914328b1333 +size 1180026 diff --git a/packages/open3d/matplotlib-3.5.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl b/packages/open3d/matplotlib-3.5.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..db6af9253601126d8605233e16a04c1f77797cf5 --- /dev/null +++ b/packages/open3d/matplotlib-3.5.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:b428076a55fb1c084c76cb93e68006f27d247169f056412607c5c88828d08f88 +size 11266085 diff --git a/packages/open3d/matplotlib_inline-0.1.7-py3-none-any.whl b/packages/open3d/matplotlib_inline-0.1.7-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..6bf8c9828709735aaaf8e6ce6ad345a40a5154fc Binary files /dev/null and b/packages/open3d/matplotlib_inline-0.1.7-py3-none-any.whl differ diff --git a/packages/open3d/mistune-3.0.2-py3-none-any.whl b/packages/open3d/mistune-3.0.2-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..d5863ee31283e31879318edfaaf903af8889ec34 Binary files /dev/null and b/packages/open3d/mistune-3.0.2-py3-none-any.whl differ diff --git a/packages/open3d/nbclassic-1.1.0-py3-none-any.whl b/packages/open3d/nbclassic-1.1.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..909de3eeb22aabc6f9d3bf47c9f4803ed5428f51 --- /dev/null +++ b/packages/open3d/nbclassic-1.1.0-py3-none-any.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:8c0fd6e36e320a18657ff44ed96c3a400f17a903a3744fc322303a515778f2ba +size 9991768 diff --git a/packages/open3d/nbclient-0.10.0-py3-none-any.whl b/packages/open3d/nbclient-0.10.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..6a505f51517376d96b993d6e0657b1ce5e966f28 Binary files /dev/null and b/packages/open3d/nbclient-0.10.0-py3-none-any.whl differ diff --git a/packages/open3d/nbconvert-7.16.4-py3-none-any.whl b/packages/open3d/nbconvert-7.16.4-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..d007e585a8d9d08f438002687a4f982e2fc841df Binary files /dev/null and b/packages/open3d/nbconvert-7.16.4-py3-none-any.whl differ diff --git a/packages/open3d/nbformat-5.10.4-py3-none-any.whl b/packages/open3d/nbformat-5.10.4-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..c9b128bd45368e36adc5a1839a8f0884fc3a704a Binary files /dev/null and b/packages/open3d/nbformat-5.10.4-py3-none-any.whl differ diff --git a/packages/open3d/nest_asyncio-1.6.0-py3-none-any.whl b/packages/open3d/nest_asyncio-1.6.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..d765d6ffead5e15d5b528e8e6f9842cc0c4749d3 Binary files /dev/null and b/packages/open3d/nest_asyncio-1.6.0-py3-none-any.whl differ diff --git a/packages/open3d/notebook-6.5.7-py3-none-any.whl b/packages/open3d/notebook-6.5.7-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..b7ff6a4b195b9bb9f247504105c6e120e20a725f Binary files /dev/null and b/packages/open3d/notebook-6.5.7-py3-none-any.whl differ diff --git a/packages/open3d/notebook_shim-0.2.4-py3-none-any.whl b/packages/open3d/notebook_shim-0.2.4-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..e5a736848f3701af0fea5103883ad224b587db0c Binary files /dev/null and b/packages/open3d/notebook_shim-0.2.4-py3-none-any.whl differ diff --git a/packages/open3d/numpy-1.21.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl b/packages/open3d/numpy-1.21.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..27ec5689d660d3b5759709e65463a1eb95c66bc8 --- /dev/null +++ b/packages/open3d/numpy-1.21.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:dde972a1e11bb7b702ed0e447953e7617723760f420decb97305e66fb4afc54f +size 14092363 diff --git a/packages/open3d/open3d-0.10.0.0-cp38-cp38-manylinux1_x86_64.whl b/packages/open3d/open3d-0.10.0.0-cp38-cp38-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..a70e01bdf3996a0d2d52ac749f03ba2541b330e6 --- /dev/null +++ b/packages/open3d/open3d-0.10.0.0-cp38-cp38-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:505193b84b48a48556a02b3f5c4d2499825fd87383195c5cf01a0464b3b1e353 +size 4664237 diff --git a/packages/open3d/packaging-24.0-py3-none-any.whl b/packages/open3d/packaging-24.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..df5619e735dd0e30356850ef3f74955bbd948cf7 Binary files /dev/null and b/packages/open3d/packaging-24.0-py3-none-any.whl differ diff --git a/packages/open3d/pandocfilters-1.5.1-py2.py3-none-any.whl b/packages/open3d/pandocfilters-1.5.1-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..1c83ca249b13c6e78cdb304996912266857761b5 Binary files /dev/null and b/packages/open3d/pandocfilters-1.5.1-py2.py3-none-any.whl differ diff --git a/packages/open3d/parso-0.8.4-py2.py3-none-any.whl b/packages/open3d/parso-0.8.4-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..19ce715eb5679d95c91b0f382ba3aba0d80b246e Binary files /dev/null and b/packages/open3d/parso-0.8.4-py2.py3-none-any.whl differ diff --git a/packages/open3d/pexpect-4.9.0-py2.py3-none-any.whl b/packages/open3d/pexpect-4.9.0-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..4cc3fbda05fd2d907932fe6b7b8b6e1777ff9e33 Binary files /dev/null and b/packages/open3d/pexpect-4.9.0-py2.py3-none-any.whl differ diff --git a/packages/open3d/pickleshare-0.7.5-py2.py3-none-any.whl b/packages/open3d/pickleshare-0.7.5-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..f7ab1102fa7f0fc9d9ba7e902af47f8d35ebe3b0 Binary files /dev/null and b/packages/open3d/pickleshare-0.7.5-py2.py3-none-any.whl differ diff --git a/packages/open3d/pkgutil_resolve_name-1.3.10-py3-none-any.whl b/packages/open3d/pkgutil_resolve_name-1.3.10-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..164c6fd6cef657d835239345dbe0f6237dc05553 Binary files /dev/null and b/packages/open3d/pkgutil_resolve_name-1.3.10-py3-none-any.whl differ diff --git a/packages/open3d/platformdirs-4.2.2-py3-none-any.whl b/packages/open3d/platformdirs-4.2.2-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..4288a233b8a323b5fada34e575e7e16cbe3fe3c5 Binary files /dev/null and b/packages/open3d/platformdirs-4.2.2-py3-none-any.whl differ diff --git a/packages/open3d/prometheus_client-0.20.0-py3-none-any.whl b/packages/open3d/prometheus_client-0.20.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..d5c78fd78199eb43e592a53d183df5d026087bf7 Binary files /dev/null and b/packages/open3d/prometheus_client-0.20.0-py3-none-any.whl differ diff --git a/packages/open3d/prompt_toolkit-3.0.45-py3-none-any.whl b/packages/open3d/prompt_toolkit-3.0.45-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..f70ee1d0b7f71cd17a251df770af393c8ee5b8b6 Binary files /dev/null and b/packages/open3d/prompt_toolkit-3.0.45-py3-none-any.whl differ diff --git a/packages/open3d/ptyprocess-0.7.0-py2.py3-none-any.whl b/packages/open3d/ptyprocess-0.7.0-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..c70a0021937e53ae36600752729aa0f890da73f1 Binary files /dev/null and b/packages/open3d/ptyprocess-0.7.0-py2.py3-none-any.whl differ diff --git a/packages/open3d/pure_eval-0.2.2-py3-none-any.whl b/packages/open3d/pure_eval-0.2.2-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..1b0d6914b0eb1a155007c6e2e0d16cda2bd168ca Binary files /dev/null and b/packages/open3d/pure_eval-0.2.2-py3-none-any.whl differ diff --git a/packages/open3d/pycparser-2.22-py3-none-any.whl b/packages/open3d/pycparser-2.22-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..b47818710e16179213829ac49a44467f01e35bb4 Binary files /dev/null and b/packages/open3d/pycparser-2.22-py3-none-any.whl differ diff --git a/packages/open3d/pygments-2.18.0-py3-none-any.whl b/packages/open3d/pygments-2.18.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..b58736af4b84c8d95306ef957b82aa7525773a78 --- /dev/null +++ b/packages/open3d/pygments-2.18.0-py3-none-any.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:b8e6aca0523f3ab76fee51799c488e38782ac06eafcf95e7ba832985c8e7b13a +size 1205513 diff --git a/packages/open3d/pyparsing-3.1.2-py3-none-any.whl b/packages/open3d/pyparsing-3.1.2-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..42ee8c776658b273bc0e3873944eda699945d6eb Binary files /dev/null and b/packages/open3d/pyparsing-3.1.2-py3-none-any.whl differ diff --git a/packages/open3d/pyrsistent-0.20.0-py3-none-any.whl b/packages/open3d/pyrsistent-0.20.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..4416bc0e49f0cabb7f175b661efed6df07569b4a Binary files /dev/null and b/packages/open3d/pyrsistent-0.20.0-py3-none-any.whl differ diff --git a/packages/open3d/python_dateutil-2.9.0.post0-py2.py3-none-any.whl b/packages/open3d/python_dateutil-2.9.0.post0-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..b9a14e1bf98404e31ba5398c4d84cbd76dcf8e84 Binary files /dev/null and b/packages/open3d/python_dateutil-2.9.0.post0-py2.py3-none-any.whl differ diff --git a/packages/open3d/pyzmq-20.0.0-cp38-cp38-manylinux1_x86_64.whl b/packages/open3d/pyzmq-20.0.0-cp38-cp38-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..cd043b1d1fc4793bdca2fcdb39edb5406d59b528 --- /dev/null +++ b/packages/open3d/pyzmq-20.0.0-cp38-cp38-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:d92c7f41a53ece82b91703ea433c7d34143248cf0cead33aa11c5fc621c764bf +size 1084742 diff --git a/packages/open3d/six-1.16.0-py2.py3-none-any.whl b/packages/open3d/six-1.16.0-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..fd942658a2f748ba433dd8632abb910a416e184f Binary files /dev/null and b/packages/open3d/six-1.16.0-py2.py3-none-any.whl differ diff --git a/packages/open3d/sniffio-1.3.1-py3-none-any.whl b/packages/open3d/sniffio-1.3.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..04f44e47d69d33e39c88ce206e3f100443083f17 Binary files /dev/null and b/packages/open3d/sniffio-1.3.1-py3-none-any.whl differ diff --git a/packages/open3d/soupsieve-2.5-py3-none-any.whl b/packages/open3d/soupsieve-2.5-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..e1be1281914864efacd06642fbf27970213ee2b2 Binary files /dev/null and b/packages/open3d/soupsieve-2.5-py3-none-any.whl differ diff --git a/packages/open3d/stack_data-0.6.3-py3-none-any.whl b/packages/open3d/stack_data-0.6.3-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..0c3b8375915a50d6e1fc797cb78c99ab7af1736c Binary files /dev/null and b/packages/open3d/stack_data-0.6.3-py3-none-any.whl differ diff --git a/packages/open3d/terminado-0.18.1-py3-none-any.whl b/packages/open3d/terminado-0.18.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..1bc4e4e3bedf78fb008b3d7bbe62ff9a238f9ea9 Binary files /dev/null and b/packages/open3d/terminado-0.18.1-py3-none-any.whl differ diff --git a/packages/open3d/tinycss2-1.3.0-py3-none-any.whl b/packages/open3d/tinycss2-1.3.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..86a3358806d2ec2c8350d6390251880df38f1713 Binary files /dev/null and b/packages/open3d/tinycss2-1.3.0-py3-none-any.whl differ diff --git a/packages/open3d/tornado-6.4-cp38-abi3-manylinux_2_5_x86_64.manylinux1_x86_64.manylinux_2_17_x86_64.manylinux2014_x86_64.whl b/packages/open3d/tornado-6.4-cp38-abi3-manylinux_2_5_x86_64.manylinux1_x86_64.manylinux_2_17_x86_64.manylinux2014_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..82a9d050d4deceb521ec96568b48cc123114e7bd Binary files /dev/null and b/packages/open3d/tornado-6.4-cp38-abi3-manylinux_2_5_x86_64.manylinux1_x86_64.manylinux_2_17_x86_64.manylinux2014_x86_64.whl differ diff --git a/packages/open3d/traitlets-5.14.3-py3-none-any.whl b/packages/open3d/traitlets-5.14.3-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..1fc82c0cfe66a7d782366822cee685b45558da11 Binary files /dev/null and b/packages/open3d/traitlets-5.14.3-py3-none-any.whl differ diff --git a/packages/open3d/typing_extensions-4.12.1-py3-none-any.whl b/packages/open3d/typing_extensions-4.12.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..116314e5b69963d14b5f30b3ba67db0dbed3965f Binary files /dev/null and b/packages/open3d/typing_extensions-4.12.1-py3-none-any.whl differ diff --git a/packages/open3d/wcwidth-0.2.13-py2.py3-none-any.whl b/packages/open3d/wcwidth-0.2.13-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..a315ca15417bf130837ee5df4c92f8981db08965 Binary files /dev/null and b/packages/open3d/wcwidth-0.2.13-py2.py3-none-any.whl differ diff --git a/packages/open3d/webencodings-0.5.1-py2.py3-none-any.whl b/packages/open3d/webencodings-0.5.1-py2.py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..04b75063bf5a2831092b64e040a77afab3648274 Binary files /dev/null and b/packages/open3d/webencodings-0.5.1-py2.py3-none-any.whl differ diff --git a/packages/open3d/websocket_client-1.8.0-py3-none-any.whl b/packages/open3d/websocket_client-1.8.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..a2d768d7937a90e7c5ecbd51c1332b00c9e3cdac Binary files /dev/null and b/packages/open3d/websocket_client-1.8.0-py3-none-any.whl differ diff --git a/packages/open3d/widgetsnbextension-4.0.11-py3-none-any.whl b/packages/open3d/widgetsnbextension-4.0.11-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..2c34a959bb303a48e1d73b658f0aa5c991cb5dd3 --- /dev/null +++ b/packages/open3d/widgetsnbextension-4.0.11-py3-none-any.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:55d4d6949d100e0d08b94948a42efc3ed6dfdc0e9468b2c4b128c9a2ce3a7a36 +size 2335052 diff --git a/packages/open3d/zipp-3.19.1-py3-none-any.whl b/packages/open3d/zipp-3.19.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..b29ccd3137d99754dc8f2cea1e4a353c737d2f01 Binary files /dev/null and b/packages/open3d/zipp-3.19.1-py3-none-any.whl differ diff --git a/packages/scikit-learn/joblib-1.4.2-py3-none-any.whl b/packages/scikit-learn/joblib-1.4.2-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..617a58993064e1d1eb40916f4c093844481dc454 Binary files /dev/null and b/packages/scikit-learn/joblib-1.4.2-py3-none-any.whl differ diff --git a/packages/scikit-learn/numpy-1.21.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl b/packages/scikit-learn/numpy-1.21.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..27ec5689d660d3b5759709e65463a1eb95c66bc8 --- /dev/null +++ b/packages/scikit-learn/numpy-1.21.3-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:dde972a1e11bb7b702ed0e447953e7617723760f420decb97305e66fb4afc54f +size 14092363 diff --git a/packages/scikit-learn/scikit_learn-1.0.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl b/packages/scikit-learn/scikit_learn-1.0.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..2aa57519e2277b50c6a2f81bca0219c8ae75bac9 --- /dev/null +++ b/packages/scikit-learn/scikit_learn-1.0.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:a800665527c1a63f7395a0baae3c89b0d97b54d2c23769c1c9879061bb80bc19 +size 21221584 diff --git a/packages/scikit-learn/scipy-1.7.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl b/packages/scikit-learn/scipy-1.7.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..891015d8698e7d18b03ed981a296ae4cd4bc0e71 --- /dev/null +++ b/packages/scikit-learn/scipy-1.7.1-cp38-cp38-manylinux_2_5_x86_64.manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:90c07ba5f34f33299a428b0d4fa24c30d2ceba44d63f8385b2b05be460819fcb +size 28368779 diff --git a/packages/scikit-learn/threadpoolctl-3.5.0-py3-none-any.whl b/packages/scikit-learn/threadpoolctl-3.5.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..294097f849dc67d11194a3edfb545048562ac1e1 Binary files /dev/null and b/packages/scikit-learn/threadpoolctl-3.5.0-py3-none-any.whl differ diff --git a/packages/torch/MarkupSafe-2.1.5-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl b/packages/torch/MarkupSafe-2.1.5-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..f2adc8feee10834b6d85f66059ee608651e3483b Binary files /dev/null and b/packages/torch/MarkupSafe-2.1.5-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl differ diff --git a/packages/torch/certifi-2024.6.2-py3-none-any.whl b/packages/torch/certifi-2024.6.2-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..df4ec17e8e7392bd572ae0aa94f81af5033b398f Binary files /dev/null and b/packages/torch/certifi-2024.6.2-py3-none-any.whl differ diff --git a/packages/torch/charset_normalizer-3.3.2-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl b/packages/torch/charset_normalizer-3.3.2-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..26f799cd7bb89ab85f6a9fb0c06d5e2bfc7c063c Binary files /dev/null and b/packages/torch/charset_normalizer-3.3.2-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl differ diff --git a/packages/torch/filelock-3.14.0-py3-none-any.whl b/packages/torch/filelock-3.14.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..7f1498fa8684d9050afbec39dcfd45e84da279cd Binary files /dev/null and b/packages/torch/filelock-3.14.0-py3-none-any.whl differ diff --git a/packages/torch/fsspec-2024.6.0-py3-none-any.whl b/packages/torch/fsspec-2024.6.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..8d07fecf3711a434e4ee6da7c50cee223ca73b76 Binary files /dev/null and b/packages/torch/fsspec-2024.6.0-py3-none-any.whl differ diff --git a/packages/torch/idna-3.7-py3-none-any.whl b/packages/torch/idna-3.7-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..fa4c95b123751094438c2f747e5389c32d870c5b Binary files /dev/null and b/packages/torch/idna-3.7-py3-none-any.whl differ diff --git a/packages/torch/jinja2-3.1.4-py3-none-any.whl b/packages/torch/jinja2-3.1.4-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..fa6a14bfaa53bb44031daf857127636101dd2409 Binary files /dev/null and b/packages/torch/jinja2-3.1.4-py3-none-any.whl differ diff --git a/packages/torch/mpmath-1.3.0-py3-none-any.whl b/packages/torch/mpmath-1.3.0-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..267fb143803110c233368618ccefb70cced5ea7f Binary files /dev/null and b/packages/torch/mpmath-1.3.0-py3-none-any.whl differ diff --git a/packages/torch/networkx-3.1-py3-none-any.whl b/packages/torch/networkx-3.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..fb518a90b78c1b3d79e2d37433b458e2b22788ee --- /dev/null +++ b/packages/torch/networkx-3.1-py3-none-any.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:4f33f68cb2afcf86f28a45f43efc27a9386b535d567d2127f8f61d51dec58d36 +size 2072251 diff --git a/packages/torch/numpy-1.24.4-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl b/packages/torch/numpy-1.24.4-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..97c490899149f8a94f0471a8b7139a40858ccee8 --- /dev/null +++ b/packages/torch/numpy-1.24.4-cp38-cp38-manylinux_2_17_x86_64.manylinux2014_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:dd80e219fd4c71fc3699fc1dadac5dcf4fd882bfc6f7ec53d30fa197b8ee22dc +size 17316263 diff --git a/packages/torch/nvidia_cublas_cu12-12.1.3.1-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_cublas_cu12-12.1.3.1-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..4b32e823d19441e92bd6541cf9c011400f18ae3a --- /dev/null +++ b/packages/torch/nvidia_cublas_cu12-12.1.3.1-py3-none-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:ee53ccca76a6fc08fb9701aa95b6ceb242cdaab118c3bb152af4e579af792728 +size 410594774 diff --git a/packages/torch/nvidia_cuda_cupti_cu12-12.1.105-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_cuda_cupti_cu12-12.1.105-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..1405bd2689a274efbff55e2a432994d4efe80bbe --- /dev/null +++ b/packages/torch/nvidia_cuda_cupti_cu12-12.1.105-py3-none-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:e54fde3983165c624cb79254ae9818a456eb6e87a7fd4d56a2352c24ee542d7e +size 14109015 diff --git a/packages/torch/nvidia_cuda_nvrtc_cu12-12.1.105-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_cuda_nvrtc_cu12-12.1.105-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..26a8c22039b6fcbeb8f7f4cab7b99a51049de196 --- /dev/null +++ b/packages/torch/nvidia_cuda_nvrtc_cu12-12.1.105-py3-none-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:339b385f50c309763ca65456ec75e17bbefcbbf2893f462cb8b90584cd27a1c2 +size 23671734 diff --git a/packages/torch/nvidia_cuda_runtime_cu12-12.1.105-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_cuda_runtime_cu12-12.1.105-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..1a6702c4c55199cf8d5f5fe6fec4c9404616998d Binary files /dev/null and b/packages/torch/nvidia_cuda_runtime_cu12-12.1.105-py3-none-manylinux1_x86_64.whl differ diff --git a/packages/torch/nvidia_cudnn_cu12-8.9.2.26-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_cudnn_cu12-8.9.2.26-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..a88608c02f082f14d6bf548dac2c9364f9c66821 --- /dev/null +++ b/packages/torch/nvidia_cudnn_cu12-8.9.2.26-py3-none-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:5ccb288774fdfb07a7e7025ffec286971c06d8d7b4fb162525334616d7629ff9 +size 731725872 diff --git a/packages/torch/nvidia_cufft_cu12-11.0.2.54-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_cufft_cu12-11.0.2.54-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..5a0d80aa82d4d273b54b657318da3045ef7d26bc --- /dev/null +++ b/packages/torch/nvidia_cufft_cu12-11.0.2.54-py3-none-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:794e3948a1aa71fd817c3775866943936774d1c14e7628c74f6f7417224cdf56 +size 121635161 diff --git a/packages/torch/nvidia_curand_cu12-10.3.2.106-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_curand_cu12-10.3.2.106-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..495a490a89408c177e331aba6877c3bcfb7e449f --- /dev/null +++ b/packages/torch/nvidia_curand_cu12-10.3.2.106-py3-none-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:9d264c5036dde4e64f1de8c50ae753237c12e0b1348738169cd0f8a536c0e1e0 +size 56467784 diff --git a/packages/torch/nvidia_cusolver_cu12-11.4.5.107-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_cusolver_cu12-11.4.5.107-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..351d1b4b51368454fc7da0c6ce2818017ba2a379 --- /dev/null +++ b/packages/torch/nvidia_cusolver_cu12-11.4.5.107-py3-none-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:8a7ec542f0412294b15072fa7dab71d31334014a69f953004ea7a118206fe0dd +size 124161928 diff --git a/packages/torch/nvidia_cusparse_cu12-12.1.0.106-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_cusparse_cu12-12.1.0.106-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..a2db4ffcb1e9ec145fb179988a643bb10681fe80 --- /dev/null +++ b/packages/torch/nvidia_cusparse_cu12-12.1.0.106-py3-none-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:f3b50f42cf363f86ab21f720998517a659a48131e8d538dc02f8768237bd884c +size 195958278 diff --git a/packages/torch/nvidia_nccl_cu12-2.18.1-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_nccl_cu12-2.18.1-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..242d116c09532c752f468c54ace5728e9825c598 --- /dev/null +++ b/packages/torch/nvidia_nccl_cu12-2.18.1-py3-none-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:1a6c4acefcbebfa6de320f412bf7866de856e786e0462326ba1bac40de0b5e71 +size 209797524 diff --git a/packages/torch/nvidia_nvjitlink_cu12-12.5.40-py3-none-manylinux2014_x86_64.whl b/packages/torch/nvidia_nvjitlink_cu12-12.5.40-py3-none-manylinux2014_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..1cce34cf52689ad21de34a1e91cc72d7aa1b5b71 --- /dev/null +++ b/packages/torch/nvidia_nvjitlink_cu12-12.5.40-py3-none-manylinux2014_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:d9714f27c1d0f0895cd8915c07a87a1d0029a0aa36acaf9156952ec2a8a12189 +size 21297869 diff --git a/packages/torch/nvidia_nvtx_cu12-12.1.105-py3-none-manylinux1_x86_64.whl b/packages/torch/nvidia_nvtx_cu12-12.1.105-py3-none-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..09ae155c1898474ae8a276a91023386d7332446e Binary files /dev/null and b/packages/torch/nvidia_nvtx_cu12-12.1.105-py3-none-manylinux1_x86_64.whl differ diff --git a/packages/torch/pillow-10.3.0-cp38-cp38-manylinux_2_28_x86_64.whl b/packages/torch/pillow-10.3.0-cp38-cp38-manylinux_2_28_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..299f083945d741f09e0601668f98857fc3c1ce06 --- /dev/null +++ b/packages/torch/pillow-10.3.0-cp38-cp38-manylinux_2_28_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:aa7e402ce11f0885305bfb6afb3434b3cd8f53b563ac065452d9d5654c7b86fd +size 4500528 diff --git a/packages/torch/requests-2.32.3-py3-none-any.whl b/packages/torch/requests-2.32.3-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..23662ce7ee5c52367ac94a50e4d94c6b39645938 Binary files /dev/null and b/packages/torch/requests-2.32.3-py3-none-any.whl differ diff --git a/packages/torch/sympy-1.12.1-py3-none-any.whl b/packages/torch/sympy-1.12.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..88e2b14b093e90eb718a3be90817d16c1916827c --- /dev/null +++ b/packages/torch/sympy-1.12.1-py3-none-any.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:9b2cbc7f1a640289430e13d2a56f02f867a1da0190f2f99d8968c2f74da0e515 +size 5743129 diff --git a/packages/torch/torch-2.1.0-cp38-cp38-manylinux1_x86_64.whl b/packages/torch/torch-2.1.0-cp38-cp38-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..2c1d942a07dd0695ab59c4ea495262253cd98e6f --- /dev/null +++ b/packages/torch/torch-2.1.0-cp38-cp38-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:fb7bf0cc1a3db484eb5d713942a93172f3bac026fcb377a0cd107093d2eba777 +size 670166245 diff --git a/packages/torch/torchaudio-2.1.0-cp38-cp38-manylinux1_x86_64.whl b/packages/torch/torchaudio-2.1.0-cp38-cp38-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..e7c3bd4fc70006fa37796fe8edc116aa6a743875 --- /dev/null +++ b/packages/torch/torchaudio-2.1.0-cp38-cp38-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:1cce2133f7b9f813ff146dffccfdb616b73932cec2fc69aa3a4189b8a8b17f8d +size 3260662 diff --git a/packages/torch/torchvision-0.16.0-cp38-cp38-manylinux1_x86_64.whl b/packages/torch/torchvision-0.16.0-cp38-cp38-manylinux1_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..110d854eebdf555f8ac6d21e6957d769f634e016 --- /dev/null +++ b/packages/torch/torchvision-0.16.0-cp38-cp38-manylinux1_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:eddd91da4603f1dbb340d9aca82344df64605a0897b17014ac8e0b54dd6e5716 +size 6938593 diff --git a/packages/torch/triton-2.1.0-0-cp38-cp38-manylinux2014_x86_64.manylinux_2_17_x86_64.whl b/packages/torch/triton-2.1.0-0-cp38-cp38-manylinux2014_x86_64.manylinux_2_17_x86_64.whl new file mode 100644 index 0000000000000000000000000000000000000000..84d77aaef6b9d4f22a1a78b158c65e0208be3bb3 --- /dev/null +++ b/packages/torch/triton-2.1.0-0-cp38-cp38-manylinux2014_x86_64.manylinux_2_17_x86_64.whl @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:39f6fb6bdccb3e98f3152e3fbea724f1aeae7d749412bbb1fa9c441d474eba26 +size 89160672 diff --git a/packages/torch/typing_extensions-4.12.1-py3-none-any.whl b/packages/torch/typing_extensions-4.12.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..116314e5b69963d14b5f30b3ba67db0dbed3965f Binary files /dev/null and b/packages/torch/typing_extensions-4.12.1-py3-none-any.whl differ diff --git a/packages/torch/urllib3-2.2.1-py3-none-any.whl b/packages/torch/urllib3-2.2.1-py3-none-any.whl new file mode 100644 index 0000000000000000000000000000000000000000..d7cca6a473137e22739248c814ef561b2f8fccc4 Binary files /dev/null and b/packages/torch/urllib3-2.2.1-py3-none-any.whl differ diff --git a/pc_util/build/lib.linux-x86_64-cpython-38/pc_util.cpython-38-x86_64-linux-gnu.so b/pc_util/build/lib.linux-x86_64-cpython-38/pc_util.cpython-38-x86_64-linux-gnu.so new file mode 100644 index 0000000000000000000000000000000000000000..c1474bdd0648360d90278a77494fa1904505d555 --- /dev/null +++ b/pc_util/build/lib.linux-x86_64-cpython-38/pc_util.cpython-38-x86_64-linux-gnu.so @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:a7a125bb537f2b7a9c4869475353d3e766da3677f6099e3ae3656385539f4cd0 +size 14111984 diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/.ninja_deps b/pc_util/build/temp.linux-x86_64-cpython-38/.ninja_deps new file mode 100644 index 0000000000000000000000000000000000000000..d7d76bb946fb5d86f652180696ce1d999fcf3fab Binary files /dev/null and b/pc_util/build/temp.linux-x86_64-cpython-38/.ninja_deps differ diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/.ninja_log b/pc_util/build/temp.linux-x86_64-cpython-38/.ninja_log new file mode 100644 index 0000000000000000000000000000000000000000..bb6895eddc6a4c375de37ddf4d49c619956156d0 --- /dev/null +++ b/pc_util/build/temp.linux-x86_64-cpython-38/.ninja_log @@ -0,0 +1,23 @@ +# ninja log v5 +1 10883 1717452942232322194 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query.o 434c99c569954244 +3 10944 1717452942292322785 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate.o 6c220ff811bc7aca +2 10955 1717452942296322824 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points.o 8a865043d3bfce62 +4 12315 1717452943664336304 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling.o c9a8c1a43975349b +1 14310 1717452945664355980 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster.o 9f93923e9a88ac90 +1 20251 1717452951608414247 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate_gpu.o c822e84ef9b2e8d3 +2 20310 1717452951664414795 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points_gpu.o 7bd3d4a533719a61 +1 20335 1717452951692415069 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query_gpu.o 9818cf16c446eb9e +2 20768 1717452952124419292 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster_gpu.o de1874c4e555d2c1 +4 23553 1717452954908446461 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling_gpu.o 8d44adcb3f955709 +3 24608 1717452955956456671 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/pointnet2_api.o d766420711e09786 +10 6550 1717453891975284457 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling.o 79940462ea0fe7bd +8 7090 1717453892519287692 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster.o 68e2ea0c2e51d619 +8 7115 1717453892539287811 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points.o 91d66b83eaecbbfe +9 8407 1717453893831295493 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate.o 6009748276168eef +7 8552 1717453893975296350 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query.o fdb2ca153416f2da +8 18825 1717453904255357388 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster_gpu.o 42e8c3c50b1a7d79 +9 19018 1717453904435358455 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/pointnet2_api.o bf233b4aa4acd0a6 +7 19339 1717453904771360448 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate_gpu.o 31ebee02a4920f61 +9 19492 1717453904919361326 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points_gpu.o 1defd6433dcc2c0c +10 20684 1717453906111368393 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling_gpu.o 204874c24b66bc +8 21044 1717453906475370550 /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query_gpu.o 6c5c11a3ea6530a5 diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/build.ninja b/pc_util/build/temp.linux-x86_64-cpython-38/build.ninja new file mode 100644 index 0000000000000000000000000000000000000000..fa37399f35cb2c5fe42206c3e71feb7cfa616d07 --- /dev/null +++ b/pc_util/build/temp.linux-x86_64-cpython-38/build.ninja @@ -0,0 +1,43 @@ +ninja_required_version = 1.3 +cxx = c++ +nvcc = /usr/local/cuda-11.8/bin/nvcc + +cflags = -pthread -B /home/tulrfpf/anaconda3/envs/test/compiler_compat -Wl,--sysroot=/ -Wsign-compare -DNDEBUG -g -fwrapv -O3 -Wall -Wstrict-prototypes -fPIC -I/home/tulrfpf/anaconda3/envs/test/lib/python3.8/site-packages/torch/include -I/home/tulrfpf/anaconda3/envs/test/lib/python3.8/site-packages/torch/include/torch/csrc/api/include -I/home/tulrfpf/anaconda3/envs/test/lib/python3.8/site-packages/torch/include/TH -I/home/tulrfpf/anaconda3/envs/test/lib/python3.8/site-packages/torch/include/THC -I/usr/local/cuda-11.8/include -I/home/tulrfpf/anaconda3/envs/test/include/python3.8 -c +post_cflags = -g -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=pc_util -D_GLIBCXX_USE_CXX11_ABI=0 -std=c++14 +cuda_cflags = -I/home/tulrfpf/anaconda3/envs/test/lib/python3.8/site-packages/torch/include -I/home/tulrfpf/anaconda3/envs/test/lib/python3.8/site-packages/torch/include/torch/csrc/api/include -I/home/tulrfpf/anaconda3/envs/test/lib/python3.8/site-packages/torch/include/TH -I/home/tulrfpf/anaconda3/envs/test/lib/python3.8/site-packages/torch/include/THC -I/usr/local/cuda-11.8/include -I/home/tulrfpf/anaconda3/envs/test/include/python3.8 -c +cuda_post_cflags = -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O2 -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=pc_util -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_86,code=compute_86 -gencode=arch=compute_86,code=sm_86 -std=c++14 +cuda_dlink_post_cflags = +ldflags = + +rule compile + command = $cxx -MMD -MF $out.d $cflags -c $in -o $out $post_cflags + depfile = $out.d + deps = gcc + +rule cuda_compile + depfile = $out.d + deps = gcc + command = $nvcc $cuda_cflags -c $in -o $out $cuda_post_cflags + + + + + +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query.o: compile /home/tulrfpf/baseline/pc_util/src/ball_query.cpp +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query_gpu.o: cuda_compile /home/tulrfpf/baseline/pc_util/src/ball_query_gpu.cu +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster.o: compile /home/tulrfpf/baseline/pc_util/src/cluster.cpp +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster_gpu.o: cuda_compile /home/tulrfpf/baseline/pc_util/src/cluster_gpu.cu +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points.o: compile /home/tulrfpf/baseline/pc_util/src/group_points.cpp +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points_gpu.o: cuda_compile /home/tulrfpf/baseline/pc_util/src/group_points_gpu.cu +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate.o: compile /home/tulrfpf/baseline/pc_util/src/interpolate.cpp +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate_gpu.o: cuda_compile /home/tulrfpf/baseline/pc_util/src/interpolate_gpu.cu +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/pointnet2_api.o: compile /home/tulrfpf/baseline/pc_util/src/pointnet2_api.cpp +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling.o: compile /home/tulrfpf/baseline/pc_util/src/sampling.cpp +build /home/tulrfpf/baseline/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling_gpu.o: cuda_compile /home/tulrfpf/baseline/pc_util/src/sampling_gpu.cu + + + + + + + diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query.o new file mode 100644 index 0000000000000000000000000000000000000000..7d9ba96ed88d9d5495012f0d4ab857b94fbcf667 --- /dev/null +++ b/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query.o @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:2230a3197984fe09e443cf91f9f63d5de51e6eb1a3f7d9a1a8e849eeed3cbdf6 +size 3459160 diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query_gpu.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query_gpu.o new file mode 100644 index 0000000000000000000000000000000000000000..f9f95069bf4f446522e75119470bae58986a419c Binary files /dev/null and b/pc_util/build/temp.linux-x86_64-cpython-38/src/ball_query_gpu.o differ diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster.o new file mode 100644 index 0000000000000000000000000000000000000000..eee09a1b1558611b6843f6952862590f89cec933 --- /dev/null +++ b/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster.o @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:e37bf46e1f30cd062baf7a4b317e30530443bd0de5927c665d854fb8c3627377 +size 3415688 diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster_gpu.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster_gpu.o new file mode 100644 index 0000000000000000000000000000000000000000..49499ae0b98bb8b253dd17344afbc195cdc75bd3 Binary files /dev/null and b/pc_util/build/temp.linux-x86_64-cpython-38/src/cluster_gpu.o differ diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points.o new file mode 100644 index 0000000000000000000000000000000000000000..aab8b97d76a5db50af90ab15be718706416dce26 --- /dev/null +++ b/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points.o @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:7ab9356f83871bf414b6652fb9059e516f31d7b09e2d2f3713e839bb18b5c25f +size 3457800 diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points_gpu.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points_gpu.o new file mode 100644 index 0000000000000000000000000000000000000000..4c20525317a9db508a7686736f2fad2c862c9ac2 Binary files /dev/null and b/pc_util/build/temp.linux-x86_64-cpython-38/src/group_points_gpu.o differ diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate.o new file mode 100644 index 0000000000000000000000000000000000000000..7404b6aabe8a23e5bfab3fa3f34e950b2ede4991 --- /dev/null +++ b/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate.o @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:e0c1442b24abd2d8ce54f90e2c4417ca66d81acaa3aa4731c8b13ec0f7fe77a5 +size 3532784 diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate_gpu.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate_gpu.o new file mode 100644 index 0000000000000000000000000000000000000000..13a0125845c398f4161d41e734ab5c431c5dc703 Binary files /dev/null and b/pc_util/build/temp.linux-x86_64-cpython-38/src/interpolate_gpu.o differ diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/pointnet2_api.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/pointnet2_api.o new file mode 100644 index 0000000000000000000000000000000000000000..091c141c99a779f835d1523dd31c46bc2e3947ca --- /dev/null +++ b/pc_util/build/temp.linux-x86_64-cpython-38/src/pointnet2_api.o @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:7740e82490ee2f524d151ded835896c04126a45c43a36f79278a7cba7db113a0 +size 15047816 diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling.o new file mode 100644 index 0000000000000000000000000000000000000000..88dc10b8684a72168eb9376c516d8a0b857020e6 --- /dev/null +++ b/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling.o @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:50e9392a39093f1821267dc9d4ba324c35ce72869d5babc001dc266388b3c98c +size 3260728 diff --git a/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling_gpu.o b/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling_gpu.o new file mode 100644 index 0000000000000000000000000000000000000000..fb678cb99bc49e24d5ebb38c614803295d39d561 Binary files /dev/null and b/pc_util/build/temp.linux-x86_64-cpython-38/src/sampling_gpu.o differ diff --git a/pc_util/dist/pc_util-1.0-py3.8-linux-x86_64.egg b/pc_util/dist/pc_util-1.0-py3.8-linux-x86_64.egg new file mode 100644 index 0000000000000000000000000000000000000000..d0d4d94c02ba410ade587119f6e07ff72019fcf5 --- /dev/null +++ b/pc_util/dist/pc_util-1.0-py3.8-linux-x86_64.egg @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:45b1e24bce755d5204d34ff64a959494d2dab7b92e5c9dc98a3d20b9fe80308a +size 4558194 diff --git a/pc_util/pc_util.egg-info/PKG-INFO b/pc_util/pc_util.egg-info/PKG-INFO new file mode 100644 index 0000000000000000000000000000000000000000..f3db88935116d55f321677a85f58f73ee3008b7b --- /dev/null +++ b/pc_util/pc_util.egg-info/PKG-INFO @@ -0,0 +1,3 @@ +Metadata-Version: 2.1 +Name: pc_util +Version: 1.0 diff --git a/pc_util/pc_util.egg-info/SOURCES.txt b/pc_util/pc_util.egg-info/SOURCES.txt new file mode 100644 index 0000000000000000000000000000000000000000..484fe44b0c2a2c52ad2dc8e2a0bf99d10423766c --- /dev/null +++ b/pc_util/pc_util.egg-info/SOURCES.txt @@ -0,0 +1,16 @@ +setup.py +pc_util.egg-info/PKG-INFO +pc_util.egg-info/SOURCES.txt +pc_util.egg-info/dependency_links.txt +pc_util.egg-info/top_level.txt +src/ball_query.cpp +src/ball_query_gpu.cu +src/cluster.cpp +src/cluster_gpu.cu +src/group_points.cpp +src/group_points_gpu.cu +src/interpolate.cpp +src/interpolate_gpu.cu +src/pointnet2_api.cpp +src/sampling.cpp +src/sampling_gpu.cu \ No newline at end of file diff --git a/pc_util/pc_util.egg-info/dependency_links.txt b/pc_util/pc_util.egg-info/dependency_links.txt new file mode 100644 index 0000000000000000000000000000000000000000..8b137891791fe96927ad78e64b0aad7bded08bdc --- /dev/null +++ b/pc_util/pc_util.egg-info/dependency_links.txt @@ -0,0 +1 @@ + diff --git a/pc_util/pc_util.egg-info/top_level.txt b/pc_util/pc_util.egg-info/top_level.txt new file mode 100644 index 0000000000000000000000000000000000000000..e58dde836a55fd8f5be1df3271a4fed8e22862fa --- /dev/null +++ b/pc_util/pc_util.egg-info/top_level.txt @@ -0,0 +1 @@ +pc_util diff --git a/pc_util/setup.py b/pc_util/setup.py new file mode 100644 index 0000000000000000000000000000000000000000..b3928b124b4f28b81277d25e19348adc9c77887d --- /dev/null +++ b/pc_util/setup.py @@ -0,0 +1,23 @@ +from setuptools import setup +from torch.utils.cpp_extension import BuildExtension, CUDAExtension + +setup( + name='pc_util', + version='1.0', + ext_modules=[ + CUDAExtension('pc_util', [ + 'src/pointnet2_api.cpp', + 'src/ball_query.cpp', + 'src/ball_query_gpu.cu', + 'src/group_points.cpp', + 'src/group_points_gpu.cu', + 'src/interpolate.cpp', + 'src/interpolate_gpu.cu', + 'src/sampling.cpp', + 'src/sampling_gpu.cu', + 'src/cluster.cpp', + 'src/cluster_gpu.cu', + ], extra_compile_args={'cxx': ['-g'], 'nvcc': ['-O2']}) + ], + cmdclass={'build_ext': BuildExtension} +) diff --git a/pc_util/src/ball_query.cpp b/pc_util/src/ball_query.cpp new file mode 100644 index 0000000000000000000000000000000000000000..1c7a0155ca0dcc8744fb70b5c338b1698ef7e3b0 --- /dev/null +++ b/pc_util/src/ball_query.cpp @@ -0,0 +1,84 @@ +#include +#include +// #include +#include +#include +#include "ball_query_gpu.h" + +// extern THCState *state; + +#include +#include +// cudaStream_t stream = at::cuda::getCurrentCUDAStream(); + +#define CHECK_CUDA(x) do { \ + if (!x.type().is_cuda()) { \ + fprintf(stderr, "%s must be CUDA tensor at %s:%d\n", #x, __FILE__, __LINE__); \ + exit(-1); \ + } \ +} while (0) +#define CHECK_CONTIGUOUS(x) do { \ + if (!x.is_contiguous()) { \ + fprintf(stderr, "%s must be contiguous tensor at %s:%d\n", #x, __FILE__, __LINE__); \ + exit(-1); \ + } \ +} while (0) +#define CHECK_INPUT(x) CHECK_CUDA(x);CHECK_CONTIGUOUS(x) + +int ball_query_wrapper_fast(int b, int n, int m, float radius, int nsample, + 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(); + const float *xyz = xyz_tensor.data(); + int *idx = idx_tensor.data(); + + ball_query_kernel_launcher_fast(b, n, m, radius, nsample, new_xyz, xyz, idx); + return 1; +} + + +int ball_center_query_wrapper_fast(int b, int n, int m, float radius, + at::Tensor point_tensor, at::Tensor key_point_tensor, at::Tensor idx_tensor) { + CHECK_INPUT(point_tensor); + CHECK_INPUT(key_point_tensor); + const float *point = point_tensor.data(); + const float *key_point = key_point_tensor.data(); + int *idx = idx_tensor.data(); + + ball_center_query_kernel_launcher_fast(b, n, m, radius, point, key_point, idx); + return 1; +} + + +int knn_query_wrapper_fast(int b, int n, int m, int nsample, + at::Tensor new_xyz_tensor, at::Tensor xyz_tensor, at::Tensor dist2_tensor, at::Tensor idx_tensor) { + CHECK_INPUT(new_xyz_tensor); + CHECK_INPUT(xyz_tensor); + const float *new_xyz = new_xyz_tensor.data(); + const float *xyz = xyz_tensor.data(); + float *dist2 = dist2_tensor.data(); + int *idx = idx_tensor.data(); + + knn_query_kernel_launcher_fast(b, n, m, nsample, new_xyz, xyz, dist2, idx); + return 1; +} + + +int ball_query_wrapper_stack(int B, int M, float radius, int nsample, + at::Tensor new_xyz_tensor, at::Tensor new_xyz_batch_cnt_tensor, + at::Tensor xyz_tensor, at::Tensor xyz_batch_cnt_tensor, at::Tensor idx_tensor) { + CHECK_INPUT(new_xyz_tensor); + CHECK_INPUT(xyz_tensor); + CHECK_INPUT(new_xyz_batch_cnt_tensor); + CHECK_INPUT(xyz_batch_cnt_tensor); + + const float *new_xyz = new_xyz_tensor.data(); + const float *xyz = xyz_tensor.data(); + const int *new_xyz_batch_cnt = new_xyz_batch_cnt_tensor.data(); + const int *xyz_batch_cnt = xyz_batch_cnt_tensor.data(); + int *idx = idx_tensor.data(); + + ball_query_kernel_launcher_stack(B, M, radius, nsample, new_xyz, new_xyz_batch_cnt, xyz, xyz_batch_cnt, idx); + return 1; +} \ No newline at end of file diff --git a/pc_util/src/ball_query_gpu.cu b/pc_util/src/ball_query_gpu.cu new file mode 100644 index 0000000000000000000000000000000000000000..cbd3414fa5e77a9206d83ad209a1b5ee6b14a6b5 --- /dev/null +++ b/pc_util/src/ball_query_gpu.cu @@ -0,0 +1,270 @@ +#include +#include +#include + + +#include "ball_query_gpu.h" +#include "cuda_utils.h" + + +__global__ void ball_query_kernel_fast(int b, int n, int m, float radius, int nsample, + const float *__restrict__ new_xyz, const float *__restrict__ xyz, int *__restrict__ idx) { + // new_xyz: (B, M, 3) + // xyz: (B, N, 3) + // output: + // idx: (B, M, nsample) + int bs_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (bs_idx >= b || pt_idx >= m) return; + + new_xyz += bs_idx * m * 3 + pt_idx * 3; + xyz += bs_idx * n * 3; + idx += bs_idx * m * nsample + pt_idx * nsample; + + float radius2 = radius * radius; + float new_x = new_xyz[0]; + float new_y = new_xyz[1]; + float new_z = new_xyz[2]; + + int cnt = 0; + for (int k = 0; k < n; ++k) { + float x = xyz[k * 3 + 0]; + float y = xyz[k * 3 + 1]; + float z = xyz[k * 3 + 2]; + float d2 = (new_x - x) * (new_x - x) + (new_y - y) * (new_y - y) + (new_z - z) * (new_z - z); + if (d2 < radius2){ + if (cnt == 0){ + for (int l = 0; l < nsample; ++l) { + idx[l] = k; + } + } + idx[cnt] = k; + ++cnt; + if (cnt >= nsample) break; + } + } +} + + +void ball_query_kernel_launcher_fast(int b, int n, int m, float radius, int nsample, \ + const float *new_xyz, const float *xyz, int *idx) { + // new_xyz: (B, M, 3) + // xyz: (B, N, 3) + // output: + // idx: (B, M, nsample) + + cudaError_t err; + + dim3 blocks(DIVUP(m, THREADS_PER_BLOCK), b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + ball_query_kernel_fast<<>>(b, n, m, radius, nsample, new_xyz, xyz, idx); + // cudaDeviceSynchronize(); // for using printf in kernel function + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + +__global__ void ball_center_query_kernel_fast(int b, int n, int m, float radius, \ + const float *__restrict__ point, const float *__restrict__ key_point, int *__restrict__ idx) { + // key_point: (B, M, 3) + // point: (B, N, 3) + // output: + // idx: (B, N) + int bs_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (bs_idx >= b || pt_idx >= n) return; + + point += bs_idx * n * 3 + pt_idx * 3; + key_point += bs_idx * m * 3; + idx += bs_idx * n + pt_idx; + + float radius2 = radius * radius; + float point_x = point[0]; + float point_y = point[1]; + float point_z = point[2]; + + float bestd = 1e8; + for (int k = 0; k < m; ++k) { + float x = key_point[k * 3 + 0]; + float y = key_point[k * 3 + 1]; + float z = key_point[k * 3 + 2]; + if (((x + 1) * (x + 1) + (y + 1) * (y + 1) + (z + 1) * (z + 1)) < 1e-4) break; + float d2 = (point_x - x) * (point_x - x) + (point_y - y) * (point_y - y) + (point_z - z) * (point_z - z); + if (d2 < radius2 && d2 < bestd){ + idx[0] = k; + bestd = d2; + } + } +} + + +void ball_center_query_kernel_launcher_fast(int b, int n, int m, float radius, \ + const float *point, const float *key_point, int *idx) { + // point: (B, n, 3) + // key_point: (B, m, 3) + // output: + // idx: (B, n) + + cudaError_t err; + + dim3 blocks(DIVUP(n, THREADS_PER_BLOCK), b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + ball_center_query_kernel_fast<<>>(b, n, m, radius, point, key_point, idx); + // cudaDeviceSynchronize(); // for using printf in kernel function + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + + + + +__global__ void knn_query_kernel_fast(int b, int n, int m, int nsample, const float *__restrict__ new_xyz, + const float *__restrict__ xyz, float *__restrict__ dist2, int *__restrict__ idx) { + + // new_xyz: (B, M, 3) + // xyz: (B, N, 3) + // output: + // dist2: (B, M, nsample) + // idx: (B, M, nsample) + + int bs_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (bs_idx >= b || pt_idx >= m) return; + + new_xyz += bs_idx * m * 3 + pt_idx * 3; + xyz += bs_idx * n * 3; + dist2 += bs_idx * m * nsample + pt_idx * nsample; + idx += bs_idx * m * nsample + pt_idx * nsample; + + float nx = new_xyz[0]; + float ny = new_xyz[1]; + float nz = new_xyz[2]; + + for (int i = 0; i < n; ++i) { + float x = xyz[i * 3 + 0]; + float y = xyz[i * 3 + 1]; + float z = xyz[i * 3 + 2]; + float d2 = (nx - x) * (nx - x) + (ny - y) * (ny - y) + (nz - z) * (nz - z); + if (d2 < dist2[nsample - 1]) { + dist2[nsample - 1] = d2; + idx[nsample - 1] = i; + for (int j = nsample - 2; j >= 0; j--) { + if (d2 < dist2[j]){ + dist2[j + 1] = dist2[j]; + dist2[j] = d2; + idx[j + 1] = idx[j]; + idx[j] = i; + } + } + } + } +} + + +void knn_query_kernel_launcher_fast(int b, int n, int m, int nsample, \ + const float *new_xyz, const float *xyz, float *dist2, int *idx) { + cudaError_t err; + + dim3 blocks(DIVUP(m, THREADS_PER_BLOCK), b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + knn_query_kernel_fast<<>>(b, n, m, nsample, new_xyz, xyz, dist2, idx); + // cudaDeviceSynchronize(); // for using printf in kernel function + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + + + + + + + +__global__ void ball_query_kernel_stack(int B, int M, float radius, int nsample, \ + const float *new_xyz, const int *new_xyz_batch_cnt, const float *xyz, const int *xyz_batch_cnt, int *idx) { + // :param xyz: (N1 + N2 ..., 3) xyz coordinates of the features + // :param xyz_batch_cnt: (batch_size), [N1, N2, ...] + // :param new_xyz: (M1 + M2 ..., 3) centers of the ball query + // :param new_xyz_batch_cnt: (batch_size), [M1, M2, ...] + // output: + // idx: (M, nsample) + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (pt_idx >= M) return; + + int bs_idx = 0, pt_cnt = new_xyz_batch_cnt[0]; + for (int k = 1; k < B; k++){ + if (pt_idx < pt_cnt) break; + pt_cnt += new_xyz_batch_cnt[k]; + bs_idx = k; + } + + int xyz_batch_start_idx = 0; + for (int k = 0; k < bs_idx; k++) xyz_batch_start_idx += xyz_batch_cnt[k]; + // for (int k = 0; k < bs_idx; k++) new_xyz_batch_start_idx += new_xyz_batch_cnt[k]; + + new_xyz += pt_idx * 3; + xyz += xyz_batch_start_idx * 3; + idx += pt_idx * nsample; + + float radius2 = radius * radius; + float new_x = new_xyz[0]; + float new_y = new_xyz[1]; + float new_z = new_xyz[2]; + int n = xyz_batch_cnt[bs_idx]; + + int cnt = 0; + for (int k = 0; k < n; ++k) { + float x = xyz[k * 3 + 0]; + float y = xyz[k * 3 + 1]; + float z = xyz[k * 3 + 2]; + float d2 = (new_x - x) * (new_x - x) + (new_y - y) * (new_y - y) + (new_z - z) * (new_z - z); + if (d2 < radius2){ + if (cnt == 0){ + for (int l = 0; l < nsample; ++l) { + idx[l] = k; + } + } + idx[cnt] = k; + ++cnt; + if (cnt >= nsample) break; + } + } + if (cnt == 0) idx[0] = -1; +} + + +void ball_query_kernel_launcher_stack(int B, int M, float radius, int nsample, + const float *new_xyz, const int *new_xyz_batch_cnt, const float *xyz, const int *xyz_batch_cnt, int *idx){ + // :param xyz: (N1 + N2 ..., 3) xyz coordinates of the features + // :param xyz_batch_cnt: (batch_size), [N1, N2, ...] + // :param new_xyz: (M1 + M2 ..., 3) centers of the ball query + // :param new_xyz_batch_cnt: (batch_size), [M1, M2, ...] + // output: + // idx: (M, nsample) + + cudaError_t err; + + dim3 blocks(DIVUP(M, THREADS_PER_BLOCK)); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + ball_query_kernel_stack<<>>(B, M, radius, nsample, new_xyz, new_xyz_batch_cnt, xyz, xyz_batch_cnt, idx); + // cudaDeviceSynchronize(); // for using printf in kernel function + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} diff --git a/pc_util/src/ball_query_gpu.h b/pc_util/src/ball_query_gpu.h new file mode 100644 index 0000000000000000000000000000000000000000..17e7e6931f006955b15119e06ecf630efc3106b6 --- /dev/null +++ b/pc_util/src/ball_query_gpu.h @@ -0,0 +1,38 @@ +#ifndef _BALL_QUERY_GPU_H +#define _BALL_QUERY_GPU_H + +#include +#include +#include +#include + +int ball_query_wrapper_fast(int b, int n, int m, float radius, int nsample, + at::Tensor new_xyz_tensor, at::Tensor xyz_tensor, at::Tensor idx_tensor); + +void ball_query_kernel_launcher_fast(int b, int n, int m, float radius, int nsample, + const float *new_xyz, const float *xyz, int *idx); + +int ball_center_query_wrapper_fast(int b, int n, int m, float radius, + at::Tensor point_tensor, at::Tensor key_point_tensor, at::Tensor idx_tensor); + +void ball_center_query_kernel_launcher_fast(int b, int n, int m, float radius, + const float *point, const float *key_point, int *idx); + +int knn_query_wrapper_fast(int b, int n, int m, int nsample, + at::Tensor new_xyz_tensor, at::Tensor xyz_tensor, at::Tensor dist2_tensor, at::Tensor idx_tensor); + +void knn_query_kernel_launcher_fast(int b, int n, int m, int nsample, + const float *new_xyz, const float *xyz, float *dist2, int *idx); + + +int ball_query_wrapper_stack(int B, int M, float radius, int nsample, + at::Tensor new_xyz_tensor, at::Tensor new_xyz_batch_cnt_tensor, + at::Tensor xyz_tensor, at::Tensor xyz_batch_cnt_tensor, at::Tensor idx_tensor); + + +void ball_query_kernel_launcher_stack(int B, int M, float radius, int nsample, + const float *new_xyz, const int *new_xyz_batch_cnt, const float *xyz, const int *xyz_batch_cnt, int *idx); + + + +#endif diff --git a/pc_util/src/cluster.cpp b/pc_util/src/cluster.cpp new file mode 100644 index 0000000000000000000000000000000000000000..2c03d4e25957ce4efea90e583f7acecfcc54a53f --- /dev/null +++ b/pc_util/src/cluster.cpp @@ -0,0 +1,50 @@ +#include +#include +// #include +#include +#include +#include "cluster_gpu.h" + +// extern THCState *state; + +#include +#include +// cudaStream_t stream = at::cuda::getCurrentCUDAStream(); + +#define CHECK_CUDA(x) do { \ + if (!x.type().is_cuda()) { \ + fprintf(stderr, "%s must be CUDA tensor at %s:%d\n", #x, __FILE__, __LINE__); \ + exit(-1); \ + } \ +} while (0) +#define CHECK_CONTIGUOUS(x) do { \ + if (!x.is_contiguous()) { \ + fprintf(stderr, "%s must be contiguous tensor at %s:%d\n", #x, __FILE__, __LINE__); \ + exit(-1); \ + } \ +} while (0) +#define CHECK_INPUT(x) CHECK_CUDA(x);CHECK_CONTIGUOUS(x) + +int dbscan_wrapper_fast(int b, int n, float eps, int min_pts, at::Tensor xyz_tensor, at::Tensor idx_tensor) { + CHECK_INPUT(xyz_tensor); + const float *xyz = xyz_tensor.data(); + int *idx = idx_tensor.data(); + + dbscan_kernel_launcher_fast(b, n, eps, min_pts, xyz, idx); + return 1; +} + + +int cluster_pts_wrapper_fast(int b, int n, int m, at::Tensor xyz_tensor, at::Tensor idx_tensor, + at::Tensor new_xyz_tensor, at::Tensor num_tensor) { + CHECK_INPUT(xyz_tensor); + CHECK_INPUT(idx_tensor); + const float *xyz = xyz_tensor.data(); + const int *idx = idx_tensor.data(); + float *new_xyz = new_xyz_tensor.data(); + int *num = num_tensor.data(); + + cluster_pts_kernel_launcher_fast(b, n, m, xyz, idx, new_xyz, num); + return 1; +} + diff --git a/pc_util/src/cluster_gpu.cu b/pc_util/src/cluster_gpu.cu new file mode 100644 index 0000000000000000000000000000000000000000..36c59c34526f1bfdbcbfa0ff4faaa84d7db1eb2d --- /dev/null +++ b/pc_util/src/cluster_gpu.cu @@ -0,0 +1,192 @@ +#include +#include +#include + + +#include "cluster_gpu.h" +#include "cuda_utils.h" + + +__device__ float get_dis(float x1, float y1, float z1, float x2, float y2, float z2) { + float dis = (x1 - x2) * (x1 - x2) + (y1 - y2) * (y1 - y2) + (z1 - z2) * (z1 - z2); + return sqrt(dis); +} +/* +__device__ void dfs (int i, int c, int n, int min_pts, const int* pts_cnt, const int* pts_adj, int* idx, int label) { + idx[i] = c; + if(pts_cnt[i] < min_pts) return; + + for(int j=0;j= b) return; + + xyz += bs_idx * n * 3; + idx += bs_idx * n; + pts_cnt += bs_idx * n; + pts_stack += bs_idx * n; + pts_adj += bs_idx * n * n; + + for(int i=0;i= min_pts) { + for(int j=0;j>>(b, n, eps, min_pts, xyz, idx, pts_cnt, pts_adj, pts_stack); + // cudaDeviceSynchronize(); // for using printf in kernel function + cudaFree(pts_cnt); + cudaFree(pts_stack); + cudaFree(pts_adj); + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + + +__global__ void cluster_pts_kernel_fast(int b, int n, int m, const float *__restrict__ xyz, const int *__restrict__ idx, + float *__restrict__ new_xyz, int *__restrict__ num) { + int bs_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (bs_idx >= b ) return; + + xyz += bs_idx * n * 3; + idx += bs_idx * n; + new_xyz += bs_idx * m * 3; + num += bs_idx * m; + + for(int i=0;i>>(b, n, m, xyz, idx, new_xyz, num); + // cudaDeviceSynchronize(); // for using printf in kernel function + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + diff --git a/pc_util/src/cluster_gpu.h b/pc_util/src/cluster_gpu.h new file mode 100644 index 0000000000000000000000000000000000000000..0b280b30b28c0e26a8db6e6df2be749ba3d7affd --- /dev/null +++ b/pc_util/src/cluster_gpu.h @@ -0,0 +1,34 @@ +#ifndef _CLUSTER_GPU_H +#define _CLUSTER_GPU_H + +#include +#include +#include +#include + +int dbscan_wrapper_fast(int b, int n, float eps, int min_pts, at::Tensor xyz_tensor, at::Tensor idx_tensor); + +void dbscan_kernel_launcher_fast(int b, int n, float eps, int min_pts, const float *xyz, int *idx); + +int cluster_pts_wrapper_fast(int b, int n, int m, at::Tensor xyz_tensor, at::Tensor idx_tensor, + at::Tensor new_xyz_tensor, at::Tensor num_tensor); + +void cluster_pts_kernel_launcher_fast(int b, int n, int m, const float *xyz, const int *idx, float *new_xyz, int *num); + + +int dbscan_wrapper_stack(int b, int n, float eps, int min_pts, at::Tensor xyz_tensor, at::Tensor xyz_batch_cnt_tensor, + at::Tensor idx_tensor); + + +void dbscan_kernel_launcher_stack(int b, int n, float eps, int min_pts, + const float *xyz, const int *xyz_batch_cnt, int *idx); + +int cluster_pts_wrapper_stack(int B, at::Tensor xyz_tensor, at::Tensor xyz_batch_cnt_tensor, at::Tensor idx_tensor, + at::Tensor new_xyz_tensor, at::Tensor cluster_cnt_tensor); + + +void cluster_pts_kernel_launcher_stack(int B, const float *xyz, const int *xyz_batch_cnt, int *idx, + const float *new_xyz, const int *cluster_cnt); + +#endif + diff --git a/pc_util/src/cuda_utils.h b/pc_util/src/cuda_utils.h new file mode 100644 index 0000000000000000000000000000000000000000..7fe27969179c976a88199bbe962ca4f8d97263a4 --- /dev/null +++ b/pc_util/src/cuda_utils.h @@ -0,0 +1,15 @@ +#ifndef _CUDA_UTILS_H +#define _CUDA_UTILS_H + +#include + +#define TOTAL_THREADS 1024 +#define THREADS_PER_BLOCK 256 +#define DIVUP(m,n) ((m) / (n) + ((m) % (n) > 0)) + +inline int opt_n_threads(int work_size) { + const int pow_2 = std::log(static_cast(work_size)) / std::log(2.0); + + return max(min(1 << pow_2, TOTAL_THREADS), 1); +} +#endif diff --git a/pc_util/src/group_points.cpp b/pc_util/src/group_points.cpp new file mode 100644 index 0000000000000000000000000000000000000000..fb69bc3070bfb582213979bd8d263cb833fd4924 --- /dev/null +++ b/pc_util/src/group_points.cpp @@ -0,0 +1,98 @@ +#include +#include +#include +#include +// #include +#include "group_points_gpu.h" + +// extern THCState *state; + +#include +#include +// cudaStream_t stream = at::cuda::getCurrentCUDAStream(); + +#define CHECK_CUDA(x) do { \ + if (!x.type().is_cuda()) { \ + fprintf(stderr, "%s must be CUDA tensor at %s:%d\n", #x, __FILE__, __LINE__); \ + exit(-1); \ + } \ +} while (0) +#define CHECK_CONTIGUOUS(x) do { \ + if (!x.is_contiguous()) { \ + fprintf(stderr, "%s must be contiguous tensor at %s:%d\n", #x, __FILE__, __LINE__); \ + exit(-1); \ + } \ +} while (0) +#define CHECK_INPUT(x) CHECK_CUDA(x);CHECK_CONTIGUOUS(x) + + + +int group_points_grad_wrapper_fast(int b, int c, int n, int npoints, int nsample, + at::Tensor grad_out_tensor, at::Tensor idx_tensor, at::Tensor grad_points_tensor) { + + float *grad_points = grad_points_tensor.data(); + const int *idx = idx_tensor.data(); + const float *grad_out = grad_out_tensor.data(); + + group_points_grad_kernel_launcher_fast(b, c, n, npoints, nsample, grad_out, idx, grad_points); + return 1; +} + + +int group_points_wrapper_fast(int b, int c, int n, int npoints, int nsample, + at::Tensor points_tensor, at::Tensor idx_tensor, at::Tensor out_tensor) { + + const float *points = points_tensor.data(); + const int *idx = idx_tensor.data(); + float *out = out_tensor.data(); + + group_points_kernel_launcher_fast(b, c, n, npoints, nsample, points, idx, out); + return 1; +} + + + + + + + +int group_points_grad_wrapper_stack(int B, int M, int C, int N, int nsample, + at::Tensor grad_out_tensor, at::Tensor idx_tensor, at::Tensor idx_batch_cnt_tensor, + at::Tensor features_batch_cnt_tensor, at::Tensor grad_features_tensor) { + + CHECK_INPUT(grad_out_tensor); + CHECK_INPUT(idx_tensor); + CHECK_INPUT(idx_batch_cnt_tensor); + CHECK_INPUT(features_batch_cnt_tensor); + CHECK_INPUT(grad_features_tensor); + + const float *grad_out = grad_out_tensor.data(); + const int *idx = idx_tensor.data(); + const int *idx_batch_cnt = idx_batch_cnt_tensor.data(); + const int *features_batch_cnt = features_batch_cnt_tensor.data(); + float *grad_features = grad_features_tensor.data(); + + group_points_grad_kernel_launcher_stack(B, M, C, N, nsample, grad_out, idx, idx_batch_cnt, features_batch_cnt, grad_features); + return 1; +} + + +int group_points_wrapper_stack(int B, int M, int C, int nsample, + at::Tensor features_tensor, at::Tensor features_batch_cnt_tensor, + at::Tensor idx_tensor, at::Tensor idx_batch_cnt_tensor, at::Tensor out_tensor) { + + CHECK_INPUT(features_tensor); + CHECK_INPUT(features_batch_cnt_tensor); + CHECK_INPUT(idx_tensor); + CHECK_INPUT(idx_batch_cnt_tensor); + CHECK_INPUT(out_tensor); + + const float *features = features_tensor.data(); + const int *idx = idx_tensor.data(); + const int *features_batch_cnt = features_batch_cnt_tensor.data(); + const int *idx_batch_cnt = idx_batch_cnt_tensor.data(); + float *out = out_tensor.data(); + + group_points_kernel_launcher_stack(B, M, C, nsample, features, features_batch_cnt, idx, idx_batch_cnt, out); + return 1; +} \ No newline at end of file diff --git a/pc_util/src/group_points_gpu.cu b/pc_util/src/group_points_gpu.cu new file mode 100644 index 0000000000000000000000000000000000000000..14cf2909c0d318d41240f878bd4200ee6c2f1061 --- /dev/null +++ b/pc_util/src/group_points_gpu.cu @@ -0,0 +1,199 @@ +#include +#include + +#include "cuda_utils.h" +#include "group_points_gpu.h" + + +__global__ void group_points_grad_kernel_fast(int b, int c, int n, int npoints, int nsample, + const float *__restrict__ grad_out, const int *__restrict__ idx, float *__restrict__ grad_points) { + // grad_out: (B, C, npoints, nsample) + // idx: (B, npoints, nsample) + // output: + // grad_points: (B, C, N) + int bs_idx = blockIdx.z; + int c_idx = blockIdx.y; + int index = blockIdx.x * blockDim.x + threadIdx.x; + int pt_idx = index / nsample; + if (bs_idx >= b || c_idx >= c || pt_idx >= npoints) return; + + int sample_idx = index % nsample; + grad_out += bs_idx * c * npoints * nsample + c_idx * npoints * nsample + pt_idx * nsample + sample_idx; + idx += bs_idx * npoints * nsample + pt_idx * nsample + sample_idx; + + atomicAdd(grad_points + bs_idx * c * n + c_idx * n + idx[0] , grad_out[0]); +} + +void group_points_grad_kernel_launcher_fast(int b, int c, int n, int npoints, int nsample, + const float *grad_out, const int *idx, float *grad_points) { + // grad_out: (B, C, npoints, nsample) + // idx: (B, npoints, nsample) + // output: + // grad_points: (B, C, N) + cudaError_t err; + dim3 blocks(DIVUP(npoints * nsample, THREADS_PER_BLOCK), c, b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + group_points_grad_kernel_fast<<>>(b, c, n, npoints, nsample, grad_out, idx, grad_points); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + +__global__ void group_points_kernel_fast(int b, int c, int n, int npoints, int nsample, + const float *__restrict__ points, const int *__restrict__ idx, float *__restrict__ out) { + // points: (B, C, N) + // idx: (B, npoints, nsample) + // output: + // out: (B, C, npoints, nsample) + int bs_idx = blockIdx.z; + int c_idx = blockIdx.y; + int index = blockIdx.x * blockDim.x + threadIdx.x; + int pt_idx = index / nsample; + if (bs_idx >= b || c_idx >= c || pt_idx >= npoints) return; + + int sample_idx = index % nsample; + + idx += bs_idx * npoints * nsample + pt_idx * nsample + sample_idx; + int in_idx = bs_idx * c * n + c_idx * n + idx[0]; + int out_idx = bs_idx * c * npoints * nsample + c_idx * npoints * nsample + pt_idx * nsample + sample_idx; + + out[out_idx] = points[in_idx]; +} + + +void group_points_kernel_launcher_fast(int b, int c, int n, int npoints, int nsample, + const float *points, const int *idx, float *out) { + // points: (B, C, N) + // idx: (B, npoints, nsample) + // output: + // out: (B, C, npoints, nsample) + cudaError_t err; + dim3 blocks(DIVUP(npoints * nsample, THREADS_PER_BLOCK), c, b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + group_points_kernel_fast<<>>(b, c, n, npoints, nsample, points, idx, out); + // cudaDeviceSynchronize(); // for using printf in kernel function + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + +__global__ void group_points_grad_kernel_stack(int B, int M, int C, int N, int nsample, + const float *grad_out, const int *idx, const int *idx_batch_cnt, const int *features_batch_cnt, float *grad_features) { + // :param grad_out: (M1 + M2 ..., C, nsample) tensor of the gradients of the output from forward + // :param idx: (M1 + M2 ..., nsample) tensor containing the indicies of features to group with + // :param idx_batch_cnt: (batch_size) [M1 + M2 ...] tensor containing the indicies of features to group with + // :param features_batch_cnt: (batch_size) [N1 + N2 ...] tensor containing the indicies of features to group with + // :return: + // grad_features: (N1 + N2 ..., C) gradient of the features + int index = blockIdx.x * blockDim.x + threadIdx.x; + int sample_idx = index % nsample; + int C_idx = (index / nsample) % C; + int pt_idx = (index / nsample / C); + + if (pt_idx >= M || C_idx >= C || sample_idx >= nsample) return; + + int bs_idx = 0, pt_cnt = idx_batch_cnt[0]; + for (int k = 1; k < B; k++){ + if (pt_idx < pt_cnt) break; + pt_cnt += idx_batch_cnt[k]; + bs_idx = k; + } + + int features_batch_start_idx = 0; + for (int k = 0; k < bs_idx; k++) features_batch_start_idx += features_batch_cnt[k]; + + grad_out += pt_idx * C * nsample + C_idx * nsample + sample_idx; + idx += pt_idx * nsample + sample_idx; + grad_features += (features_batch_start_idx + idx[0]) * C + C_idx; + + atomicAdd(grad_features, grad_out[0]); +} + +void group_points_grad_kernel_launcher_stack(int B, int M, int C, int N, int nsample, + const float *grad_out, const int *idx, const int *idx_batch_cnt, const int *features_batch_cnt, float *grad_features) { + // :param grad_out: (M1 + M2 ..., C, nsample) tensor of the gradients of the output from forward + // :param idx: (M1 + M2 ..., nsample) tensor containing the indicies of features to group with + // :param idx_batch_cnt: (batch_size) [M1 + M2 ...] tensor containing the indicies of features to group with + // :param features_batch_cnt: (batch_size) [N1 + N2 ...] tensor containing the indicies of features to group with + // :return: + // grad_features: (N1 + N2 ..., C) gradient of the features + + cudaError_t err; + // dim3 blocks(DIVUP(npoints * nsample, THREADS_PER_BLOCK), c, b); // blockIdx.x(col), blockIdx.y(row) + dim3 blocks(DIVUP(M * C * nsample, THREADS_PER_BLOCK)); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + group_points_grad_kernel_stack<<>>(B, M, C, N, nsample, grad_out, idx, idx_batch_cnt, features_batch_cnt, grad_features); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + +__global__ void group_points_kernel_stack(int B, int M, int C, int nsample, + const float *features, const int *features_batch_cnt, const int *idx, const int *idx_batch_cnt, float *out) { + // :param features: (N1 + N2 ..., C) tensor of features to group + // :param features_batch_cnt: (batch_size) [N1 + N2 ...] tensor containing the indicies of features to group with + // :param idx: (M1 + M2 ..., nsample) tensor containing the indicies of features to group with + // :param idx_batch_cnt: (batch_size) [M1 + M2 ...] tensor containing the indicies of features to group with + // :return: + // output: (M1 + M2, C, nsample) tensor + int index = blockIdx.x * blockDim.x + threadIdx.x; + int sample_idx = index % nsample; + int C_idx = (index / nsample) % C; + int pt_idx = (index / nsample / C); + + if (pt_idx >= M || C_idx >= C || sample_idx >= nsample) return; + + int bs_idx = 0, pt_cnt = idx_batch_cnt[0]; + for (int k = 1; k < B; k++){ + if (pt_idx < pt_cnt) break; + pt_cnt += idx_batch_cnt[k]; + bs_idx = k; + } + + int features_batch_start_idx = 0; + for (int k = 0; k < bs_idx; k++) features_batch_start_idx += features_batch_cnt[k]; + features += features_batch_start_idx * C; + + idx += pt_idx * nsample + sample_idx; + int in_idx = idx[0] * C + C_idx; + int out_idx = pt_idx * C * nsample + C_idx * nsample + sample_idx; + + out[out_idx] = features[in_idx]; +} + + +void group_points_kernel_launcher_stack(int B, int M, int C, int nsample, + const float *features, const int *features_batch_cnt, const int *idx, const int *idx_batch_cnt, float *out) { + // :param features: (N1 + N2 ..., C) tensor of features to group + // :param features_batch_cnt: (batch_size) [N1 + N2 ...] tensor containing the indicies of features to group with + // :param idx: (M1 + M2 ..., nsample) tensor containing the indicies of features to group with + // :param idx_batch_cnt: (batch_size) [M1 + M2 ...] tensor containing the indicies of features to group with + // :return: + // output: (M1 + M2, C, nsample) tensor + + cudaError_t err; + dim3 blocks(DIVUP(M * C * nsample, THREADS_PER_BLOCK)); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + group_points_kernel_stack<<>>(B, M, C, nsample, features, features_batch_cnt, idx, idx_batch_cnt, out); + // cudaDeviceSynchronize(); // for using printf in kernel function + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} diff --git a/pc_util/src/group_points_gpu.h b/pc_util/src/group_points_gpu.h new file mode 100644 index 0000000000000000000000000000000000000000..a221deffa28db6723195d07089feb7cbe24b24f0 --- /dev/null +++ b/pc_util/src/group_points_gpu.h @@ -0,0 +1,36 @@ +#ifndef _GROUP_POINTS_GPU_H +#define _GROUP_POINTS_GPU_H + +#include +#include +#include +#include + + +int group_points_wrapper_fast(int b, int c, int n, int npoints, int nsample, + at::Tensor points_tensor, at::Tensor idx_tensor, at::Tensor out_tensor); + +void group_points_kernel_launcher_fast(int b, int c, int n, int npoints, int nsample, + const float *points, const int *idx, float *out); + +int group_points_grad_wrapper_fast(int b, int c, int n, int npoints, int nsample, + at::Tensor grad_out_tensor, at::Tensor idx_tensor, at::Tensor grad_points_tensor); + +void group_points_grad_kernel_launcher_fast(int b, int c, int n, int npoints, int nsample, + const float *grad_out, const int *idx, float *grad_points); + +int group_points_wrapper_stack(int B, int M, int C, int nsample, + at::Tensor features_tensor, at::Tensor features_batch_cnt_tensor, + at::Tensor idx_tensor, at::Tensor idx_batch_cnt_tensor, at::Tensor out_tensor); + +void group_points_kernel_launcher_stack(int B, int M, int C, int nsample, + const float *features, const int *features_batch_cnt, const int *idx, const int *idx_batch_cnt, float *out); + +int group_points_grad_wrapper_stack(int B, int M, int C, int N, int nsample, + at::Tensor grad_out_tensor, at::Tensor idx_tensor, at::Tensor idx_batch_cnt_tensor, + at::Tensor features_batch_cnt_tensor, at::Tensor grad_features_tensor); + +void group_points_grad_kernel_launcher_stack(int B, int M, int C, int N, int nsample, + const float *grad_out, const int *idx, const int *idx_batch_cnt, const int *features_batch_cnt, float *grad_features); + +#endif diff --git a/pc_util/src/interpolate.cpp b/pc_util/src/interpolate.cpp new file mode 100644 index 0000000000000000000000000000000000000000..2ccdadae62200791fd9395532128edf06ca6ec2a --- /dev/null +++ b/pc_util/src/interpolate.cpp @@ -0,0 +1,148 @@ +#include +#include +// #include +#include +#include +#include +#include +#include +#include "interpolate_gpu.h" + +// extern THCState *state; + +#include +#include +// cudaStream_t stream = at::cuda::getCurrentCUDAStream(); + +#define CHECK_CUDA(x) do { \ + if (!x.type().is_cuda()) { \ + fprintf(stderr, "%s must be CUDA tensor at %s:%d\n", #x, __FILE__, __LINE__); \ + exit(-1); \ + } \ +} while (0) +#define CHECK_CONTIGUOUS(x) do { \ + if (!x.is_contiguous()) { \ + fprintf(stderr, "%s must be contiguous tensor at %s:%d\n", #x, __FILE__, __LINE__); \ + exit(-1); \ + } \ +} while (0) +#define CHECK_INPUT(x) CHECK_CUDA(x);CHECK_CONTIGUOUS(x) + + +void three_nn_wrapper_fast(int b, int n, int m, at::Tensor unknown_tensor, + at::Tensor known_tensor, at::Tensor dist2_tensor, at::Tensor idx_tensor) { + const float *unknown = unknown_tensor.data(); + const float *known = known_tensor.data(); + float *dist2 = dist2_tensor.data(); + int *idx = idx_tensor.data(); + + three_nn_kernel_launcher_fast(b, n, m, unknown, known, dist2, idx); +} + + +void three_interpolate_wrapper_fast(int b, int c, int m, int n, + at::Tensor points_tensor, + at::Tensor idx_tensor, + at::Tensor weight_tensor, + at::Tensor out_tensor) { + + const float *points = points_tensor.data(); + const float *weight = weight_tensor.data(); + float *out = out_tensor.data(); + const int *idx = idx_tensor.data(); + + + three_interpolate_kernel_launcher_fast(b, c, m, n, points, idx, weight, out); +} + +void three_interpolate_grad_wrapper_fast(int b, int c, int n, int m, + at::Tensor grad_out_tensor, + at::Tensor idx_tensor, + at::Tensor weight_tensor, + at::Tensor grad_points_tensor) { + + const float *grad_out = grad_out_tensor.data(); + const float *weight = weight_tensor.data(); + float *grad_points = grad_points_tensor.data(); + const int *idx = idx_tensor.data(); + + three_interpolate_grad_kernel_launcher_fast(b, c, n, m, grad_out, idx, weight, grad_points); +} + + +void three_nn_wrapper_stack(at::Tensor unknown_tensor, + at::Tensor unknown_batch_cnt_tensor, at::Tensor known_tensor, + at::Tensor known_batch_cnt_tensor, at::Tensor dist2_tensor, at::Tensor idx_tensor){ + // unknown: (N1 + N2 ..., 3) + // unknown_batch_cnt: (batch_size), [N1, N2, ...] + // known: (M1 + M2 ..., 3) + // known_batch_cnt: (batch_size), [M1, M2, ...] + // Return: + // dist: (N1 + N2 ..., 3) l2 distance to the three nearest neighbors + // idx: (N1 + N2 ..., 3) index of the three nearest neighbors + CHECK_INPUT(unknown_tensor); + CHECK_INPUT(unknown_batch_cnt_tensor); + CHECK_INPUT(known_tensor); + CHECK_INPUT(known_batch_cnt_tensor); + CHECK_INPUT(dist2_tensor); + CHECK_INPUT(idx_tensor); + + int batch_size = unknown_batch_cnt_tensor.size(0); + int N = unknown_tensor.size(0); + int M = known_tensor.size(0); + const float *unknown = unknown_tensor.data(); + const int *unknown_batch_cnt = unknown_batch_cnt_tensor.data(); + const float *known = known_tensor.data(); + const int *known_batch_cnt = known_batch_cnt_tensor.data(); + float *dist2 = dist2_tensor.data(); + int *idx = idx_tensor.data(); + + three_nn_kernel_launcher_stack(batch_size, N, M, unknown, unknown_batch_cnt, known, known_batch_cnt, dist2, idx); +} + + +void three_interpolate_wrapper_stack(at::Tensor features_tensor, + at::Tensor idx_tensor, at::Tensor weight_tensor, at::Tensor out_tensor) { + // features_tensor: (M1 + M2 ..., C) + // idx_tensor: [N1 + N2 ..., 3] + // weight_tensor: [N1 + N2 ..., 3] + // Return: + // out_tensor: (N1 + N2 ..., C) + CHECK_INPUT(features_tensor); + CHECK_INPUT(idx_tensor); + CHECK_INPUT(weight_tensor); + CHECK_INPUT(out_tensor); + + int N = out_tensor.size(0); + int channels = features_tensor.size(1); + const float *features = features_tensor.data(); + const float *weight = weight_tensor.data(); + const int *idx = idx_tensor.data(); + float *out = out_tensor.data(); + + three_interpolate_kernel_launcher_stack(N, channels, features, idx, weight, out); +} + + +void three_interpolate_grad_wrapper_stack(at::Tensor grad_out_tensor, at::Tensor idx_tensor, + at::Tensor weight_tensor, at::Tensor grad_features_tensor) { + // grad_out_tensor: (N1 + N2 ..., C) + // idx_tensor: [N1 + N2 ..., 3] + // weight_tensor: [N1 + N2 ..., 3] + // Return: + // grad_features_tensor: (M1 + M2 ..., C) + CHECK_INPUT(grad_out_tensor); + CHECK_INPUT(idx_tensor); + CHECK_INPUT(weight_tensor); + CHECK_INPUT(grad_features_tensor); + + int N = grad_out_tensor.size(0); + int channels = grad_out_tensor.size(1); + const float *grad_out = grad_out_tensor.data(); + const float *weight = weight_tensor.data(); + const int *idx = idx_tensor.data(); + float *grad_features = grad_features_tensor.data(); + + // printf("N=%d, channels=%d\n", N, channels); + three_interpolate_grad_kernel_launcher_stack(N, channels, grad_out, idx, weight, grad_features); +} \ No newline at end of file diff --git a/pc_util/src/interpolate_gpu.cu b/pc_util/src/interpolate_gpu.cu new file mode 100644 index 0000000000000000000000000000000000000000..ff6bcc84c1fb27aa5d9eea749cc672d29de1e26a --- /dev/null +++ b/pc_util/src/interpolate_gpu.cu @@ -0,0 +1,343 @@ +#include +#include +#include + +#include "cuda_utils.h" +#include "interpolate_gpu.h" + + +__global__ void three_nn_kernel_fast(int b, int n, int m, const float *__restrict__ unknown, + const float *__restrict__ known, float *__restrict__ dist2, int *__restrict__ idx) { + // unknown: (B, N, 3) + // known: (B, M, 3) + // output: + // dist2: (B, N, 3) + // idx: (B, N, 3) + + int bs_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (bs_idx >= b || pt_idx >= n) return; + + unknown += bs_idx * n * 3 + pt_idx * 3; + known += bs_idx * m * 3; + dist2 += bs_idx * n * 3 + pt_idx * 3; + idx += bs_idx * n * 3 + pt_idx * 3; + + float ux = unknown[0]; + float uy = unknown[1]; + float uz = unknown[2]; + + double best1 = 1e40, best2 = 1e40, best3 = 1e40; + int besti1 = 0, besti2 = 0, besti3 = 0; + for (int k = 0; k < m; ++k) { + float x = known[k * 3 + 0]; + float y = known[k * 3 + 1]; + float z = known[k * 3 + 2]; + float d = (ux - x) * (ux - x) + (uy - y) * (uy - y) + (uz - z) * (uz - z); + if (d < best1) { + best3 = best2; besti3 = besti2; + best2 = best1; besti2 = besti1; + best1 = d; besti1 = k; + } + else if (d < best2) { + best3 = best2; besti3 = besti2; + best2 = d; besti2 = k; + } + else if (d < best3) { + best3 = d; besti3 = k; + } + } + dist2[0] = best1; dist2[1] = best2; dist2[2] = best3; + idx[0] = besti1; idx[1] = besti2; idx[2] = besti3; +} + + +void three_nn_kernel_launcher_fast(int b, int n, int m, const float *unknown, + const float *known, float *dist2, int *idx) { + // unknown: (B, N, 3) + // known: (B, M, 3) + // output: + // dist2: (B, N, 3) + // idx: (B, N, 3) + + cudaError_t err; + dim3 blocks(DIVUP(n, THREADS_PER_BLOCK), b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + three_nn_kernel_fast<<>>(b, n, m, unknown, known, dist2, idx); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + +__global__ void three_interpolate_kernel_fast(int b, int c, int m, int n, const float *__restrict__ points, + const int *__restrict__ idx, const float *__restrict__ weight, float *__restrict__ out) { + // points: (B, C, M) + // idx: (B, N, 3) + // weight: (B, N, 3) + // output: + // out: (B, C, N) + + int bs_idx = blockIdx.z; + int c_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + + if (bs_idx >= b || c_idx >= c || pt_idx >= n) return; + + weight += bs_idx * n * 3 + pt_idx * 3; + points += bs_idx * c * m + c_idx * m; + idx += bs_idx * n * 3 + pt_idx * 3; + out += bs_idx * c * n + c_idx * n; + + out[pt_idx] = weight[0] * points[idx[0]] + weight[1] * points[idx[1]] + weight[2] * points[idx[2]]; +} + +void three_interpolate_kernel_launcher_fast(int b, int c, int m, int n, + const float *points, const int *idx, const float *weight, float *out) { + // points: (B, C, M) + // idx: (B, N, 3) + // weight: (B, N, 3) + // output: + // out: (B, C, N) + + cudaError_t err; + dim3 blocks(DIVUP(n, THREADS_PER_BLOCK), c, b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + three_interpolate_kernel_fast<<>>(b, c, m, n, points, idx, weight, out); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + +__global__ void three_interpolate_grad_kernel_fast(int b, int c, int n, int m, const float *__restrict__ grad_out, + const int *__restrict__ idx, const float *__restrict__ weight, float *__restrict__ grad_points) { + // grad_out: (B, C, N) + // weight: (B, N, 3) + // output: + // grad_points: (B, C, M) + + int bs_idx = blockIdx.z; + int c_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + + if (bs_idx >= b || c_idx >= c || pt_idx >= n) return; + + grad_out += bs_idx * c * n + c_idx * n + pt_idx; + weight += bs_idx * n * 3 + pt_idx * 3; + grad_points += bs_idx * c * m + c_idx * m; + idx += bs_idx * n * 3 + pt_idx * 3; + + + atomicAdd(grad_points + idx[0], grad_out[0] * weight[0]); + atomicAdd(grad_points + idx[1], grad_out[0] * weight[1]); + atomicAdd(grad_points + idx[2], grad_out[0] * weight[2]); +} + +void three_interpolate_grad_kernel_launcher_fast(int b, int c, int n, int m, const float *grad_out, + const int *idx, const float *weight, float *grad_points) { + // grad_out: (B, C, N) + // weight: (B, N, 3) + // output: + // grad_points: (B, C, M) + + cudaError_t err; + dim3 blocks(DIVUP(n, THREADS_PER_BLOCK), c, b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + three_interpolate_grad_kernel_fast<<>>(b, c, n, m, grad_out, idx, weight, grad_points); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + +__global__ void three_nn_kernel_stack(int batch_size, int N, int M, const float *unknown, + const int *unknown_batch_cnt, const float *known, const int *known_batch_cnt, + float *dist2, int *idx) { + // unknown: (N1 + N2 ..., 3) + // unknown_batch_cnt: (batch_size), [N1, N2, ...] + // known: (M1 + M2 ..., 3) + // known_batch_cnt: (batch_size), [M1, M2, ...] + // Return: + // dist: (N1 + N2 ..., 3) l2 distance to the three nearest neighbors + // idx: (N1 + N2 ..., 3) index of the three nearest neighbors + + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (pt_idx >= N) return; + + int bs_idx = 0, pt_cnt = unknown_batch_cnt[0]; + for (int k = 1; k < batch_size; k++){ + if (pt_idx < pt_cnt) break; + pt_cnt += unknown_batch_cnt[k]; + bs_idx = k; + } + + int cur_num_known_points = known_batch_cnt[bs_idx]; + + int known_batch_start_idx = 0; + for (int k = 0; k < bs_idx; k++) known_batch_start_idx += known_batch_cnt[k]; + + known += known_batch_start_idx * 3; + unknown += pt_idx * 3; + dist2 += pt_idx * 3; + idx += pt_idx * 3; + + float ux = unknown[0]; + float uy = unknown[1]; + float uz = unknown[2]; + + double best1 = 1e40, best2 = 1e40, best3 = 1e40; + int besti1 = 0, besti2 = 0, besti3 = 0; + for (int k = 0; k < cur_num_known_points; ++k) { + float x = known[k * 3 + 0]; + float y = known[k * 3 + 1]; + float z = known[k * 3 + 2]; + float d = (ux - x) * (ux - x) + (uy - y) * (uy - y) + (uz - z) * (uz - z); + if (d < best1) { + best3 = best2; besti3 = besti2; + best2 = best1; besti2 = besti1; + best1 = d; besti1 = k; + } + else if (d < best2) { + best3 = best2; besti3 = besti2; + best2 = d; besti2 = k; + } + else if (d < best3) { + best3 = d; besti3 = k; + } + } + dist2[0] = best1; dist2[1] = best2; dist2[2] = best3; + idx[0] = besti1 + known_batch_start_idx; + idx[1] = besti2 + known_batch_start_idx; + idx[2] = besti3 + known_batch_start_idx; +} + + +void three_nn_kernel_launcher_stack(int batch_size, int N, int M, const float *unknown, + const int *unknown_batch_cnt, const float *known, const int *known_batch_cnt, + float *dist2, int *idx) { + // unknown: (N1 + N2 ..., 3) + // unknown_batch_cnt: (batch_size), [N1, N2, ...] + // known: (M1 + M2 ..., 3) + // known_batch_cnt: (batch_size), [M1, M2, ...] + // Return: + // dist: (N1 + N2 ..., 3) l2 distance to the three nearest neighbors + // idx: (N1 + N2 ..., 3) index of the three nearest neighbors + + cudaError_t err; + dim3 blocks(DIVUP(N, THREADS_PER_BLOCK)); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + three_nn_kernel_stack<<>>( + batch_size, N, M, unknown, unknown_batch_cnt, + known, known_batch_cnt, dist2, idx + ); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + + +__global__ void three_interpolate_kernel_stack(int N, int channels, const float *features, + const int *idx, const float *weight, float *out) { + // features: (M1 + M2 ..., C) + // idx: [N1 + N2 ..., 3] + // weight: [N1 + N2 ..., 3] + // Return: + // out: (N1 + N2 ..., C) + + int c_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (pt_idx >= N || c_idx >= channels) return; + + weight += pt_idx * 3; + idx += pt_idx * 3; + out += pt_idx * channels + c_idx; + + out[0] = weight[0] * features[idx[0] * channels + c_idx] + + weight[1] * features[idx[1] * channels + c_idx] + + weight[2] * features[idx[2] * channels + c_idx]; +} + + + +void three_interpolate_kernel_launcher_stack(int N, int channels, + const float *features, const int *idx, const float *weight, float *out) { + // features: (M1 + M2 ..., C) + // idx: [N1 + N2 ..., 3] + // weight: [N1 + N2 ..., 3] + // Return: + // out: (N1 + N2 ..., C) + + cudaError_t err; + dim3 blocks(DIVUP(N, THREADS_PER_BLOCK), channels); + dim3 threads(THREADS_PER_BLOCK); + three_interpolate_kernel_stack<<>>(N, channels, features, idx, weight, out); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + +__global__ void three_interpolate_grad_kernel_stack(int N, int channels, const float *grad_out, + const int *idx, const float *weight, float *grad_features) { + // grad_out_tensor: (N1 + N2 ..., C) + // idx_tensor: [N1 + N2 ..., 3] + // weight_tensor: [N1 + N2 ..., 3] + // Return: + // grad_features_tensor: (M1 + M2 ..., C) + + int c_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (pt_idx >= N || c_idx >= channels) return; + + grad_out += pt_idx * channels + c_idx; + weight += pt_idx * 3; + idx += pt_idx * 3; + + // printf("pt_idx=%d, c_idx=%d, idx=(%d, %d, %d), grad_out=%f\n", pt_idx, c_idx, idx[0], idx[1], idx[2], grad_out[0]); + + atomicAdd(grad_features + idx[0] * channels + c_idx, grad_out[0] * weight[0]); + atomicAdd(grad_features + idx[1] * channels + c_idx, grad_out[0] * weight[1]); + atomicAdd(grad_features + idx[2] * channels + c_idx, grad_out[0] * weight[2]); +} + + +void three_interpolate_grad_kernel_launcher_stack(int N, int channels, const float *grad_out, + const int *idx, const float *weight, float *grad_features) { + // grad_out_tensor: (N1 + N2 ..., C) + // idx_tensor: [N1 + N2 ..., 3] + // weight_tensor: [N1 + N2 ..., 3] + // Return: + // grad_features_tensor: (M1 + M2 ..., C) + + cudaError_t err; + dim3 blocks(DIVUP(N, THREADS_PER_BLOCK), channels); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + three_interpolate_grad_kernel_stack<<>>( + N, channels, grad_out, idx, weight, grad_features + ); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} \ No newline at end of file diff --git a/pc_util/src/interpolate_gpu.h b/pc_util/src/interpolate_gpu.h new file mode 100644 index 0000000000000000000000000000000000000000..742ce865ff4fbbe03c2cd5f699011f8ff61d3d58 --- /dev/null +++ b/pc_util/src/interpolate_gpu.h @@ -0,0 +1,61 @@ +#ifndef _INTERPOLATE_GPU_H +#define _INTERPOLATE_GPU_H + +#include +#include +#include +#include + + +void three_nn_wrapper_fast(int b, int n, int m, at::Tensor unknown_tensor, + at::Tensor known_tensor, at::Tensor dist2_tensor, at::Tensor idx_tensor); + +void three_nn_kernel_launcher_fast(int b, int n, int m, const float *unknown, + const float *known, float *dist2, int *idx); + + +void three_interpolate_wrapper_fast(int b, int c, int m, int n, at::Tensor points_tensor, + at::Tensor idx_tensor, at::Tensor weight_tensor, at::Tensor out_tensor); + +void three_interpolate_kernel_launcher_fast(int b, int c, int m, int n, + const float *points, const int *idx, const float *weight, float *out); + + +void three_interpolate_grad_wrapper_fast(int b, int c, int n, int m, at::Tensor grad_out_tensor, + at::Tensor idx_tensor, at::Tensor weight_tensor, at::Tensor grad_points_tensor); + +void three_interpolate_grad_kernel_launcher_fast(int b, int c, int n, int m, const float *grad_out, + const int *idx, const float *weight, float *grad_points); + + + +void three_nn_wrapper_stack(at::Tensor unknown_tensor, + at::Tensor unknown_batch_cnt_tensor, at::Tensor known_tensor, + at::Tensor known_batch_cnt_tensor, at::Tensor dist2_tensor, at::Tensor idx_tensor); + + +void three_interpolate_wrapper_stack(at::Tensor features_tensor, + at::Tensor idx_tensor, at::Tensor weight_tensor, at::Tensor out_tensor); + + + +void three_interpolate_grad_wrapper_stack(at::Tensor grad_out_tensor, at::Tensor idx_tensor, + at::Tensor weight_tensor, at::Tensor grad_features_tensor); + + +void three_nn_kernel_launcher_stack(int batch_size, int N, int M, const float *unknown, + const int *unknown_batch_cnt, const float *known, const int *known_batch_cnt, + float *dist2, int *idx); + + +void three_interpolate_kernel_launcher_stack(int N, int channels, + const float *features, const int *idx, const float *weight, float *out); + + + +void three_interpolate_grad_kernel_launcher_stack(int N, int channels, const float *grad_out, + const int *idx, const float *weight, float *grad_features); + + + +#endif diff --git a/pc_util/src/pointnet2_api.cpp b/pc_util/src/pointnet2_api.cpp new file mode 100644 index 0000000000000000000000000000000000000000..fb65a367c216fd228fc3aedce5c0df176bdf9037 --- /dev/null +++ b/pc_util/src/pointnet2_api.cpp @@ -0,0 +1,41 @@ +#include +#include + +#include "ball_query_gpu.h" +#include "group_points_gpu.h" +#include "sampling_gpu.h" +#include "interpolate_gpu.h" +#include "cluster_gpu.h" + + +PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { + m.def("ball_query_wrapper", &ball_query_wrapper_fast, "ball_query_wrapper_fast"); + m.def("ball_center_query_wrapper", &ball_center_query_wrapper_fast, "ball_center_query_wrapper_fast"); + m.def("knn_query_wrapper", &knn_query_wrapper_fast, "knn_query_wrapper_fast"); + + m.def("group_points_wrapper", &group_points_wrapper_fast, "group_points_wrapper_fast"); + m.def("group_points_grad_wrapper", &group_points_grad_wrapper_fast, "group_points_grad_wrapper_fast"); + + m.def("gather_points_wrapper", &gather_points_wrapper_fast, "gather_points_wrapper_fast"); + m.def("gather_points_grad_wrapper", &gather_points_grad_wrapper_fast, "gather_points_grad_wrapper_fast"); + + m.def("furthest_point_sampling_wrapper", &furthest_point_sampling_wrapper, "furthest_point_sampling_wrapper"); + + m.def("three_nn_wrapper", &three_nn_wrapper_fast, "three_nn_wrapper_fast"); + m.def("three_interpolate_wrapper", &three_interpolate_wrapper_fast, "three_interpolate_wrapper_fast"); + m.def("three_interpolate_grad_wrapper", &three_interpolate_grad_wrapper_fast, "three_interpolate_grad_wrapper_fast"); + + m.def("dbscan_wrapper", &dbscan_wrapper_fast, "dbscan_wrapper_fast"); + m.def("cluster_pts_wrapper", &cluster_pts_wrapper_fast, "cluster_pts_wrapper_fast"); + + + m.def("ball_query_wrapper_stack", &ball_query_wrapper_stack, "ball_query_wrapper_stack"); + + m.def("group_points_wrapper_stack", &group_points_wrapper_stack, "group_points_wrapper_stack"); + m.def("group_points_grad_wrapper_stack", &group_points_grad_wrapper_stack, "group_points_grad_wrapper_stack"); + + m.def("three_nn_wrapper_stack", &three_nn_wrapper_stack, "three_nn_wrapper_stack"); + m.def("three_interpolate_wrapper_stack", &three_interpolate_wrapper_stack, "three_interpolate_wrapper_stack"); + m.def("three_interpolate_grad_wrapper_stack", &three_interpolate_grad_wrapper_stack, "three_interpolate_grad_wrapper_stack"); + +} diff --git a/pc_util/src/sampling.cpp b/pc_util/src/sampling.cpp new file mode 100644 index 0000000000000000000000000000000000000000..9ecca40fd43ff02bad1d986ff675ac87a43aafe3 --- /dev/null +++ b/pc_util/src/sampling.cpp @@ -0,0 +1,46 @@ +#include +#include +#include +// #include + +#include "sampling_gpu.h" + +// extern THCState *state; + +#include +#include +// cudaStream_t stream = at::cuda::getCurrentCUDAStream(); + +int gather_points_wrapper_fast(int b, int c, int n, int npoints, + at::Tensor points_tensor, at::Tensor idx_tensor, at::Tensor out_tensor){ + const float *points = points_tensor.data(); + const int *idx = idx_tensor.data(); + float *out = out_tensor.data(); + + gather_points_kernel_launcher_fast(b, c, n, npoints, points, idx, out); + return 1; +} + + +int gather_points_grad_wrapper_fast(int b, int c, int n, int npoints, + at::Tensor grad_out_tensor, at::Tensor idx_tensor, at::Tensor grad_points_tensor) { + + const float *grad_out = grad_out_tensor.data(); + const int *idx = idx_tensor.data(); + float *grad_points = grad_points_tensor.data(); + + gather_points_grad_kernel_launcher_fast(b, c, n, npoints, grad_out, idx, grad_points); + return 1; +} + + +int furthest_point_sampling_wrapper(int b, int c, int n, int m, float w1, float w2, + at::Tensor points_tensor, at::Tensor temp_tensor, at::Tensor idx_tensor) { + + const float *points = points_tensor.data(); + float *temp = temp_tensor.data(); + int *idx = idx_tensor.data(); + + furthest_point_sampling_kernel_launcher(b, c, n, m, w1, w2, points, temp, idx); + return 1; +} diff --git a/pc_util/src/sampling_gpu.cu b/pc_util/src/sampling_gpu.cu new file mode 100644 index 0000000000000000000000000000000000000000..4bf227022814d6cea52b6edd95f3f8ba803b0fa3 --- /dev/null +++ b/pc_util/src/sampling_gpu.cu @@ -0,0 +1,259 @@ +#include +#include + +#include "cuda_utils.h" +#include "sampling_gpu.h" + + +__global__ void gather_points_kernel_fast(int b, int c, int n, int m, + const float *__restrict__ points, const int *__restrict__ idx, float *__restrict__ out) { + // points: (B, C, N) + // idx: (B, M) + // output: + // out: (B, C, M) + + int bs_idx = blockIdx.z; + int c_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (bs_idx >= b || c_idx >= c || pt_idx >= m) return; + + out += bs_idx * c * m + c_idx * m + pt_idx; + idx += bs_idx * m + pt_idx; + points += bs_idx * c * n + c_idx * n; + out[0] = points[idx[0]]; +} + +void gather_points_kernel_launcher_fast(int b, int c, int n, int npoints, + const float *points, const int *idx, float *out) { + // points: (B, C, N) + // idx: (B, npoints) + // output: + // out: (B, C, npoints) + + cudaError_t err; + dim3 blocks(DIVUP(npoints, THREADS_PER_BLOCK), c, b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + gather_points_kernel_fast<<>>(b, c, n, npoints, points, idx, out); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + +__global__ void gather_points_grad_kernel_fast(int b, int c, int n, int m, const float *__restrict__ grad_out, + const int *__restrict__ idx, float *__restrict__ grad_points) { + // grad_out: (B, C, M) + // idx: (B, M) + // output: + // grad_points: (B, C, N) + + int bs_idx = blockIdx.z; + int c_idx = blockIdx.y; + int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; + if (bs_idx >= b || c_idx >= c || pt_idx >= m) return; + + grad_out += bs_idx * c * m + c_idx * m + pt_idx; + idx += bs_idx * m + pt_idx; + grad_points += bs_idx * c * n + c_idx * n; + + atomicAdd(grad_points + idx[0], grad_out[0]); +} + +void gather_points_grad_kernel_launcher_fast(int b, int c, int n, int npoints, + const float *grad_out, const int *idx, float *grad_points) { + // grad_out: (B, C, npoints) + // idx: (B, npoints) + // output: + // grad_points: (B, C, N) + + cudaError_t err; + dim3 blocks(DIVUP(npoints, THREADS_PER_BLOCK), c, b); // blockIdx.x(col), blockIdx.y(row) + dim3 threads(THREADS_PER_BLOCK); + + gather_points_grad_kernel_fast<<>>(b, c, n, npoints, grad_out, idx, grad_points); + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} + + +__device__ void __update(float *__restrict__ dists, int *__restrict__ dists_i, int idx1, int idx2){ + const float v1 = dists[idx1], v2 = dists[idx2]; + const int i1 = dists_i[idx1], i2 = dists_i[idx2]; + dists[idx1] = max(v1, v2); + dists_i[idx1] = v2 > v1 ? i2 : i1; +} + +template +__global__ void furthest_point_sampling_kernel(int b, int c, int n, int m, float w1, float w2, + const float *__restrict__ dataset, float *__restrict__ temp, int *__restrict__ idxs) { + // dataset: (B, N, 3) + // tmp: (B, N) + // output: + // idx: (B, M) + + if (m <= 0) return; + __shared__ float dists[block_size]; + __shared__ int dists_i[block_size]; + + int batch_index = blockIdx.x; + dataset += batch_index * n * c; + temp += batch_index * n; + idxs += batch_index * m; + + int tid = threadIdx.x; + const int stride = block_size; + + int old = 0; + if (threadIdx.x == 0) + idxs[0] = old; + + __syncthreads(); + for (int j = 1; j < m; j++) { + int besti = 0; + float best = -1; + float x1 = dataset[old * c + 0]; + float y1 = dataset[old * c + 1]; + float z1 = dataset[old * c + 2]; + + for (int k = tid; k < n; k += stride) { + float x2, y2, z2; + x2 = dataset[k * c + 0]; + y2 = dataset[k * c + 1]; + z2 = dataset[k * c + 2]; + // float mag = (x2 * x2) + (y2 * y2) + (z2 * z2); + // if (mag <= 1e-3) + // continue; + + float xyz_d = (x2 - x1) * (x2 - x1) + (y2 - y1) * (y2 - y1) + (z2 - z1) * (z2 - z1); + float fea_d = 0; + for (int l = 3; l < c; l++) { + fea_d += (dataset[old * c + l] - dataset[k * c + l]) * (dataset[old * c + l] - dataset[k * c + l]); + } + float d = w1 * xyz_d + w2 * fea_d; + float d2 = min(d, temp[k]); + temp[k] = d2; + besti = d2 > best ? k : besti; + best = d2 > best ? d2 : best; + } + dists[tid] = best; + dists_i[tid] = besti; + __syncthreads(); + + if (block_size >= 1024) { + if (tid < 512) { + __update(dists, dists_i, tid, tid + 512); + } + __syncthreads(); + } + + if (block_size >= 512) { + if (tid < 256) { + __update(dists, dists_i, tid, tid + 256); + } + __syncthreads(); + } + if (block_size >= 256) { + if (tid < 128) { + __update(dists, dists_i, tid, tid + 128); + } + __syncthreads(); + } + if (block_size >= 128) { + if (tid < 64) { + __update(dists, dists_i, tid, tid + 64); + } + __syncthreads(); + } + if (block_size >= 64) { + if (tid < 32) { + __update(dists, dists_i, tid, tid + 32); + } + __syncthreads(); + } + if (block_size >= 32) { + if (tid < 16) { + __update(dists, dists_i, tid, tid + 16); + } + __syncthreads(); + } + if (block_size >= 16) { + if (tid < 8) { + __update(dists, dists_i, tid, tid + 8); + } + __syncthreads(); + } + if (block_size >= 8) { + if (tid < 4) { + __update(dists, dists_i, tid, tid + 4); + } + __syncthreads(); + } + if (block_size >= 4) { + if (tid < 2) { + __update(dists, dists_i, tid, tid + 2); + } + __syncthreads(); + } + if (block_size >= 2) { + if (tid < 1) { + __update(dists, dists_i, tid, tid + 1); + } + __syncthreads(); + } + + old = dists_i[0]; + if (tid == 0) + idxs[j] = old; + } +} + +void furthest_point_sampling_kernel_launcher(int b, int c, int n, int m, float w1, float w2, + const float *dataset, float *temp, int *idxs) { + // dataset: (B, N, 3) + // tmp: (B, N) + // output: + // idx: (B, M) + + cudaError_t err; + unsigned int n_threads = opt_n_threads(n); + + switch (n_threads) { + case 1024: + furthest_point_sampling_kernel<1024><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 512: + furthest_point_sampling_kernel<512><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 256: + furthest_point_sampling_kernel<256><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 128: + furthest_point_sampling_kernel<128><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 64: + furthest_point_sampling_kernel<64><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 32: + furthest_point_sampling_kernel<32><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 16: + furthest_point_sampling_kernel<16><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 8: + furthest_point_sampling_kernel<8><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 4: + furthest_point_sampling_kernel<4><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 2: + furthest_point_sampling_kernel<2><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + case 1: + furthest_point_sampling_kernel<1><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); break; + default: + furthest_point_sampling_kernel<512><<>>(b, c, n, m, w1, w2, dataset, temp, idxs); + } + + err = cudaGetLastError(); + if (cudaSuccess != err) { + fprintf(stderr, "CUDA kernel failed : %s\n", cudaGetErrorString(err)); + exit(-1); + } +} diff --git a/pc_util/src/sampling_gpu.h b/pc_util/src/sampling_gpu.h new file mode 100644 index 0000000000000000000000000000000000000000..f362530a10bb128d077b718f5bbdd80c53da03f8 --- /dev/null +++ b/pc_util/src/sampling_gpu.h @@ -0,0 +1,29 @@ +#ifndef _SAMPLING_GPU_H +#define _SAMPLING_GPU_H + +#include +#include +#include + + +int gather_points_wrapper_fast(int b, int c, int n, int npoints, + at::Tensor points_tensor, at::Tensor idx_tensor, at::Tensor out_tensor); + +void gather_points_kernel_launcher_fast(int b, int c, int n, int npoints, + const float *points, const int *idx, float *out); + + +int gather_points_grad_wrapper_fast(int b, int c, int n, int npoints, + at::Tensor grad_out_tensor, at::Tensor idx_tensor, at::Tensor grad_points_tensor); + +void gather_points_grad_kernel_launcher_fast(int b, int c, int n, int npoints, + const float *grad_out, const int *idx, float *grad_points); + + +int furthest_point_sampling_wrapper(int b, int c, int n, int m, float w1, float w2, + at::Tensor points_tensor, at::Tensor temp_tensor, at::Tensor idx_tensor); + +void furthest_point_sampling_kernel_launcher(int b, int c, int n, int m, float w1, float w2, + const float *dataset, float *temp, int *idxs); + +#endif