add third_party

third_party/3DSparseConvolution/.gitignore

@@ -0,0 +1,13 @@
+objs
+tool/__pycache__/
+tool/deploy/__pycache__/
+workspace/lib
+workspace/*.pth
+workspace/pro
+nuscenes
+tool/*.so
+workspace/centerpoint/output.zyx.dense
+workspace/centerpoint/out_dense.py.fp16.tensor
+workspace/bevfusion/output.zyx.dense
+workspace/bevfusion/output.xyz.dense
+src/pybind11.hpp

third_party/3DSparseConvolution/.vscode/settings.json

@@ -0,0 +1,5 @@
+{
+    "files.associations": {
+        "functional": "cpp"
+    }
+}

third_party/3DSparseConvolution/Makefile

@@ -0,0 +1,96 @@
+# SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: MIT
+#
+# Permission is hereby granted, free of charge, to any person obtaining a
+# copy of this software and associated documentation files (the "Software"),
+# to deal in the Software without restriction, including without limitation
+# the rights to use, copy, modify, merge, publish, distribute, sublicense,
+# and/or sell copies of the Software, and to permit persons to whom the
+# Software is furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+# THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+# FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+# DEALINGS IN THE SOFTWARE.
+
+optimize  := -O3
+cc        := g++
+stdcpp    := c++11
+arch      := $(shell arch)
+pwd       := $(abspath .)
+cppstrict := -Wall -Werror -Wextra -Werror=all-warnings
+custrict  := -Werror=all-warnings
+nvcc      := $(CUDA_HOME)/bin/nvcc
+cuda_arch := -gencode arch=compute_80,code=sm_80 \
+			 -gencode arch=compute_86,code=sm_86
+libspconv := libspconv-1.1.0
+
+ifeq ($(arch), aarch64)
+cuda_arch += -gencode arch=compute_87,code=sm_87
+endif
+
+include_paths     := -Isrc -I$(libspconv)/include -I$(CUDA_HOME)/include
+cpp_compile_flags := -std=$(stdcpp) $(strict) -g $(optimize) -fPIC -fopenmp -pthread $(include_paths)
+cu_compile_flags  := -std=$(stdcpp) $(custrict) -g $(optimize) $(cuda_arch) -Xcompiler "$(cpp_compile_flags)"
+link_flags        := -L$(libspconv)/lib/$(arch) -L$(CUDA_HOME)/lib64 -lcudart -lstdc++ -ldl -lspconv -lprotobuf -pthread \
+					 -fopenmp -Wl,-rpath='$$ORIGIN' -Wl,-rpath=$(pwd)/$(libspconv)/lib/$(arch)
+
+ifeq ($(MAKECMDGOALS), pyscn)
+python_include := $(shell python3 -c "import sysconfig;print(sysconfig.get_path('include'))")
+python_soname  := $(shell python3 -c "import sysconfig;import re;print(re.sub('lib|.so|.a', '', sysconfig.get_config_var('LIBRARY')))")
+python_libpath := $(shell python3 -c "import sysconfig;print(sysconfig.get_config_var('LIBDIR'))")
+pybind_include := $(shell cd ../.. && pwd)/dependencies/pybind11/include
+$(info Use Python Include: $(python_include))
+$(info Use Python SO Name: $(python_soname))
+$(info Use Python Library: $(python_libpath))
+$(info Use PyBind11: $(pybind_include))
+cpp_compile_flags += -I$(python_include) -I$(pybind_include)
+cu_compile_flags  += -I$(python_include) -I$(pybind_include)
+link_flags        += -L$(python_libpath) -l$(python_soname)
+endif
+
+all   : pro pyscn
+pro   : workspace/pro
+pyscn : tool/pyscn.so
+
+int8      : pro
+	@cd workspace && ./pro int8
+
+fp16      : pro
+	@cd workspace && ./pro fp16
+
+memint8      : pro
+	@cd workspace && ./pro memint8
+
+memfp16      : pro
+	@cd workspace && ./pro memfp16
+
+tool/pyscn.so : objs/pyscn.o objs/onnx-parser.o objs/onnx/onnx-ml.pb.o objs/onnx/onnx-operators-ml.pb.o
+	@echo Link $@
+	@$(cc) -shared -o $@ $^ $(link_flags)
+	@echo You can run \"python tool/pytest.py\" to test
+
+workspace/pro : objs/main.o objs/voxelization.uo objs/onnx-parser.o objs/onnx/onnx-ml.pb.o objs/onnx/onnx-operators-ml.pb.o
+	@echo Link $@
+	@$(cc) -o $@ $^ $(link_flags)
+
+objs/%.o : src/%.cpp
+	@echo Compile CXX $<
+	@mkdir -p $(dir $@)
+	@$(cc) -c $< -o $@ $(cpp_compile_flags)
+
+objs/%.uo : src/%.cu
+	@echo Compile CUDA $<
+	@mkdir -p $(dir $@)
+	@$(nvcc) -c $< -o $@ $(cu_compile_flags)
+
+clean:
+	@rm -rf objs workspace/pro tool/pyscn.so
+
+.PHONY: pyscn pro clean

third_party/3DSparseConvolution/README.md

@@ -0,0 +1,101 @@
+# 3D Sparse Convolution Network
+A tiny inference engine for [3d sparse convolutional networks](https://github.com/tianweiy/CenterPoint/blob/master/det3d/models/backbones/scn.py) using int8/fp16.
+![title](/assets/3dsparse_conv.png)
+
+## News
+- (8/18/2023) The libspconv.so 1.1.0 is released now!
+  - Open sourced the onnx parser part. libprotobuf will be configured by yourself.
+  - Multiple outputs are supported.
+
+## Model && Data
+This demo uses lidar data from [nuScenes Dataset](https://www.nuscenes.org/).
+Onnx model can be converted from checkpoint and config below using given script.
+|  Dataset  |  Checkpoint  | Config |
+| --------- | ------------ | ------ |
+|  nuScenes | [epoch_20.pth](https://mitprod-my.sharepoint.com/:f:/g/personal/tianweiy_mit_edu/EhgzjwV2EghOnHFKyRgSadoBr2kUo7yPu52N-I3dG3c5dA?e=a9MdhX) | [nusc_centerpoint_voxelnet_0075voxel_fix_bn_z](https://github.com/tianweiy/CenterPoint/blob/master/configs/nusc/voxelnet/nusc_centerpoint_voxelnet_0075voxel_fix_bn_z.py) |
+
+## Accuracy on nuScenes Validation
+|         **Model**        |  **3D Inference** | **Precision** | **mAP** | **NDS** |                   **Description**                  |
+|:------------------------:|:-----------------:|:-------------:|:-------:|:-------:|:--------------------------------------------------:|
+| centerpoint.scn.PTQ.onnx |     libspconv.so     |      INT8     | 59.15  |  66.45  | PTQ Model, spconv.so INT8 Inference                |
+| centerpoint.scn.PTQ.pth | PyTorch FakeQuant |      INT8     |   59.08  |   66.2  | PTQ Model, PyTorch FP16 Inference + FakeQuant-INT8 |
+|     [scn.nuscenes.pth](https://github.com/tianweiy/CenterPoint/blob/master/configs/nusc/README.md)     |      PyTorch      |      FP16     |   59.6  |   66.8  | From CenterPoint official, Validation              |
+|     centerpoint.scn.onnx    |     libspconv.so     |      FP16     |   59.5  |  66.71  | From CenterPoint official & Inference by spconv.so |
+
+## Memory Usage
+|     **Model**     | **Precision** | **Memory** |
+|:-----------------:|:-------------:|:----------:|
+| centerpoint.scn.PTQ.onnx |      FP16     |    422MB   |
+| centerpoint.scn.PTQ.onnx |      INT8     |    426MB   |
+
+## Export ONNX
+1. Download and configure the CenterPoint environment from https://github.com/tianweiy/CenterPoint
+2. Export SCN ONNX
+```
+$ cp -r tool/centerpoint-export path/to/CenterPoint
+$ cd path/to/CenterPoint
+$ python centerpoint-export/export-scn.py --ckpt=epoch_20.pth --save-onnx=scn.nuscenes.onnx
+$ cp scn.nuscenes.onnx path/to/3DSparseConvolution/workspace/
+```
+
+3. ## Compile && Run
+- Build and run test
+```
+$ sudo apt-get install libprotobuf-dev
+$ cd path/to/3DSparseConvolution
+->>>>>> modify main.cpp:80 to scn.nuscenes.onnx
+$ make fp16 -j
+🙌 Output.shape: 1 x 256 x 180 x 180
+[PASSED 🤗], libspconv version is 1.0.0
+To verify the results, you can execute the following command.
+Verify Result:
+  python tool/compare.py workspace/centerpoint/out_dense.torch.fp16.tensor workspace/centerpoint/output.zyx.dense --detail
+[PASSED].
+```
+
+- Verify output
+```
+$ python tool/compare.py workspace/centerpoint/out_dense.torch.fp16.tensor workspace/centerpoint/output.zyx.dense --detail
+================ Compare Information =================
+ CPP     Tensor: 1 x 256 x 180 x 180, float16 : workspace/centerpoint/out_dense.torch.fp16.tensor
+ PyTorch Tensor: 1 x 256 x 180 x 180, float16 : workspace/centerpoint/output.zyx.dense
+[absdiff]: max:0.19891357421875, sum:1438.463379, std:0.001725, mean:0.000173
+CPP:   absmax:3.066406, min:0.000000, std:0.034445, mean:0.003252
+Torch: absmax:3.054688, min:0.000000, std:0.034600, mean:0.003279
+[absdiff > m75% --- 0.149185]: 0.000 %, 2
+[absdiff > m50% --- 0.099457]: 0.000 %, 17
+[absdiff > m25% --- 0.049728]: 0.010 %, 846
+[absdiff > 0]: 2.140 %, 177539
+[absdiff = 0]: 97.860 %, 8116861
+[cosine]: 99.876 %
+======================================================
+```
+
+## For Python
+```
+$ make pyscn -j
+Use Python Include: /usr/include/python3.8
+Use Python SO Name: python3.8
+Use Python Library: /usr/lib
+Compile CXX src/pyscn.cpp
+Link tool/pyscn.so
+You can run "python tool/pytest.py" to test
+
+$ python tool/pytest.py
+[PASSED 🤗].
+To verify result:
+  python tool/compare.py workspace/centerpoint/out_dense.py.fp16.tensor workspace/centerpoint/out_dense.torch.fp16.tensor --detail
+```
+
+## Performance on ORIN
+- Summary performance using 6019 data from nuscenes
+![](workspace/perf.png)
+
+## Note
+- The current version supports compute arch are required sm_80, sm_86, and sm_87..
+- Supported operators:
+  - SparseConvolution, Add, Relu, Add&Relu, ScatterDense, Reshape and ScatterDense&Transpose.
+- Supported SparseConvolution:
+  - SpatiallySparseConvolution and SubmanifoldSparseConvolution.
+- Supported properties of SparseConvolution:
+  - activation, kernel_size, dilation, stride, padding, rulebook, subm, output_bound, precision and output_precision.

third_party/3DSparseConvolution/libspconv-1.1.0/include/spconv/check.hpp

@@ -0,0 +1,90 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: LicenseRef-NvidiaProprietary
+ *
+ * NVIDIA CORPORATION, its affiliates and licensors retain all intellectual
+ * property and proprietary rights in and to this material, related
+ * documentation and any modifications thereto. Any use, reproduction,
+ * disclosure or distribution of this material and related documentation
+ * without an express license agreement from NVIDIA CORPORATION or
+ * its affiliates is strictly prohibited.
+ */
+
+#ifndef __SPCONV_CHECK_HPP__
+#define __SPCONV_CHECK_HPP__
+
+#include <assert.h>
+#include <cuda_runtime.h>
+#include <stdarg.h>
+#include <stdio.h>
+
+#include <string>
+
+namespace spconv {
+
+#if DEBUG
+#define checkRuntime(call) spconv::check_runtime(call, #call, __LINE__, __FILE__)
+#define checkKernel(...)                                                                \
+  [&] {                                                                                 \
+    __VA_ARGS__;                                                                        \
+    checkRuntime(cudaStreamSynchronize(nullptr));                                       \
+    return spconv::check_runtime(cudaGetLastError(), #__VA_ARGS__, __LINE__, __FILE__); \
+  }()
+#define dprintf printf
+#else
+#define checkRuntime(call) spconv::check_runtime(call, #call, __LINE__, __FILE__)
+#define checkKernel(...)                                                            \
+  do {                                                                              \
+    __VA_ARGS__;                                                                    \
+    spconv::check_runtime(cudaPeekAtLastError(), #__VA_ARGS__, __LINE__, __FILE__); \
+  } while (0)
+#define dprintf(...)
+#endif
+
+#define Assertf(cond, fmt, ...)                                                                 \
+  do {                                                                                          \
+    if (!(cond)) {                                                                              \
+      fprintf(stderr, "Assert failed 💀. %s in file %s:%d, message: " fmt "\n", #cond, __FILE__, \
+              __LINE__, __VA_ARGS__);                                                           \
+      abort();                                                                                  \
+    }                                                                                           \
+  } while (false)
+#define Asserts(cond, s)                                                                      \
+  do {                                                                                        \
+    if (!(cond)) {                                                                            \
+      fprintf(stderr, "Assert failed 💀. %s in file %s:%d, message: " s "\n", #cond, __FILE__, \
+              __LINE__);                                                                      \
+      abort();                                                                                \
+    }                                                                                         \
+  } while (false)
+#define Assert(cond)                                                                     \
+  do {                                                                                   \
+    if (!(cond)) {                                                                       \
+      fprintf(stderr, "Assert failed 💀. %s in file %s:%d\n", #cond, __FILE__, __LINE__); \
+      abort();                                                                           \
+    }                                                                                    \
+  } while (false)
+
+static inline std::string format(const char *fmt, ...) {
+  char buffer[2048];
+  va_list vl;
+  va_start(vl, fmt);
+  vsnprintf(buffer, sizeof(buffer), fmt, vl);
+  return buffer;
+}
+
+static inline bool check_runtime(cudaError_t e, const char *call, int line, const char *file) {
+  if (e != cudaSuccess) {
+    fprintf(stderr,
+            "CUDA Runtime error %s # %s, code = %s [ %d ] in file "
+            "%s:%d\n",
+            call, cudaGetErrorString(e), cudaGetErrorName(e), e, file, line);
+    abort();
+    return false;
+  }
+  return true;
+}
+
+};  // namespace spconv
+
+#endif  // #ifndef __SPCONV_CHECK_HPP__