Skip to content

GitLab

Commit a6ae8967 authored by traveller59's avatar traveller59
Browse files

spconv v1.1 release: 

1. add cuda hash support for cuda indice generation.
2. use hash table instead of dense table in CPU code.
3. add CPU-only build support.
parent 0757c45b
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 433 additions and 205 deletions
CMakeLists.txt
View file @ a6ae8967
cmake_minimum_required(VERSION 3.13 FATAL_ERROR) cmake_minimum_required(VERSION 3.13 FATAL_ERROR)
project(SparseConv LANGUAGES CXX CUDA VERSION 1.0)
option(SPCONV_BuildTests "Build the unit tests when BUILD_TESTING is enabled." OFF) option(SPCONV_BuildTests "Build the unit tests when BUILD_TESTING is enabled." ON)
set(CMAKE_CXX_EXTENSIONS OFF) # avoid gnu++11 be added to CXX flags option(SPCONV_BuildCUDA "Build cuda code when BUILD_TESTING is enabled." ON)
if (SPCONV_BuildCUDA)
set(CUDA_TOOLKIT_ROOT_DIR "${CMAKE_CUDA_COMPILER}") project(SparseConv LANGUAGES CXX CUDA VERSION 1.0)
get_filename_component(CUDA_TOOLKIT_ROOT_DIR "${CUDA_TOOLKIT_ROOT_DIR}" DIRECTORY)
get_filename_component(CUDA_TOOLKIT_ROOT_DIR "${CUDA_TOOLKIT_ROOT_DIR}" DIRECTORY)
if(WIN32) # true if windows (32 and 64 bit)
set(CUDA_LIB_PATH_HINTS "${CUDA_TOOLKIT_ROOT_DIR}/lib/x64")
add_compile_definitions(TV_WINDOWS)
else() else()
set(CUDA_LIB_PATH_HINTS "${CUDA_TOOLKIT_ROOT_DIR}/lib64") project(SparseConv LANGUAGES CXX VERSION 1.0)
endif() endif()
# set(CMAKE_VERBOSE_MAKEFILE ON)
find_library(CUDA_CUDART NAMES cudart HINTS ${CUDA_LIB_PATH_HINTS})
find_library(CUDA_CUBLAS NAMES cublas HINTS ${CUDA_LIB_PATH_HINTS}) set(CMAKE_CXX_EXTENSIONS OFF) # avoid gnu++11 be added to CXX flags
if(CMAKE_BUILD_TYPE STREQUAL "Debug") if(CMAKE_BUILD_TYPE STREQUAL "Debug")
add_compile_definitions(TV_DEBUG) add_compile_definitions(TV_DEBUG)
endif() endif()
find_package(Torch REQUIRED) find_package(Torch REQUIRED)
torch_cuda_get_nvcc_gencode_flag(NVCC_FLAGS_EXTRA) if (SPCONV_BuildCUDA)
string (REPLACE ";" " " NVCC_FLAGS_EXTRA_STR "${NVCC_FLAGS_EXTRA}") set(CUDA_TOOLKIT_ROOT_DIR "${CMAKE_CUDA_COMPILER}")
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} ${NVCC_FLAGS_EXTRA_STR}") get_filename_component(CUDA_TOOLKIT_ROOT_DIR "${CUDA_TOOLKIT_ROOT_DIR}" DIRECTORY)
get_filename_component(CUDA_TOOLKIT_ROOT_DIR "${CUDA_TOOLKIT_ROOT_DIR}" DIRECTORY)
if(WIN32) # true if windows (32 and 64 bit)
set(CUDA_LIB_PATH_HINTS "${CUDA_TOOLKIT_ROOT_DIR}/lib/x64")
add_compile_definitions(TV_WINDOWS)
else()
set(CUDA_LIB_PATH_HINTS "${CUDA_TOOLKIT_ROOT_DIR}/lib64")
endif()
# set(CMAKE_VERBOSE_MAKEFILE ON)
find_library(CUDA_CUDART NAMES cudart HINTS ${CUDA_LIB_PATH_HINTS})
find_library(CUDA_CUBLAS NAMES cublas HINTS ${CUDA_LIB_PATH_HINTS})
torch_cuda_get_nvcc_gencode_flag(NVCC_FLAGS_EXTRA)
string (REPLACE ";" " " NVCC_FLAGS_EXTRA_STR "${NVCC_FLAGS_EXTRA}")
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} ${NVCC_FLAGS_EXTRA_STR}")
add_compile_definitions(SPCONV_CUDA)
endif()
# add_definitions(-D_GLIBCXX_USE_CXX11_ABI=0) # add_definitions(-D_GLIBCXX_USE_CXX11_ABI=0)
add_compile_definitions(SPCONV_CUDA)
add_subdirectory(third_party/pybind11) add_subdirectory(third_party/pybind11)
set(ALL_LIBS ${CUDA_CUDART} ${CUDA_CUBLAS} ${TORCH_LIBRARIES}) set(ALL_LIBS ${TORCH_LIBRARIES})
set(ALL_INCLUDE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} set(ALL_INCLUDE ${PROJECT_SOURCE_DIR}/include)
${PROJECT_SOURCE_DIR}/include)
if (SPCONV_BuildCUDA)
set(ALL_LIBS ${ALL_LIBS} ${CUDA_CUDART} ${CUDA_CUBLAS})
set(ALL_INCLUDE ${ALL_INCLUDE} ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
add_subdirectory(src/cuhash)
endif()
add_subdirectory(src/spconv) add_subdirectory(src/spconv)
add_subdirectory(src/utils) add_subdirectory(src/utils)
add_subdirectory(src/hash)
if (SPCONV_BuildTests) if (SPCONV_BuildTests)
include(CTest) #adds option BUILD_TESTING (default ON) include(CTest) #adds option BUILD_TESTING (default ON)
......
README.md
View file @ a6ae8967
...@@ -2,15 +2,15 @@ ...@@ -2,15 +2,15 @@
This is a spatially sparse convolution library like [SparseConvNet](https://github.com/facebookresearch/SparseConvNet) but faster and easy to read. This library provide sparse convolution/transposed, submanifold convolution, inverse convolution and sparse maxpool. This is a spatially sparse convolution library like [SparseConvNet](https://github.com/facebookresearch/SparseConvNet) but faster and easy to read. This library provide sparse convolution/transposed, submanifold convolution, inverse convolution and sparse maxpool.
If you need more kinds of spatial layers such as avg pool, please implement it by yourself, I don't have time to do this.
The GPU Indice Generation algorithm is a unofficial implementation of paper [SECOND](http://www.mdpi.com/1424-8220/18/10/3337). That algorithm (don't include GPU SubM indice generation algorithm) may be protected by patent. The GPU Indice Generation algorithm is a unofficial implementation of paper [SECOND](http://www.mdpi.com/1424-8220/18/10/3337). That algorithm (don't include GPU SubM indice generation algorithm) may be protected by patent.
This project only support CUDA 9.0+. If you are using cuda 8.0, please update it to 9.0. This project only support CUDA 9.0+ or CPU only. If you are using cuda 8.0, please update it to 9.0.
This project only support tensors with spatial volume less than ```std::numeric_limits<int>::max()``` (~2e9). if someone really need very large space, open an issue.
## News: ## News:
2019-5-22: spconv v1.1 alpha released, now cuda hash implementation will be default. you can use ```use_hash=False``` to use dense implementation. you may see some message during running, they will be removed in future. 2019-5-24: spconv v1.1 released, now indice generation will use hash table as default (CPU code only support hash table). you can use ```use_hash=False``` to use dense table when using CUDA. In addition, add CPU only build support.
## Install on Ubuntu 16.04/18.04 ## Install on Ubuntu 16.04/18.04
...@@ -20,7 +20,7 @@ This project only support CUDA 9.0+. If you are using cuda 8.0, please update it ...@@ -20,7 +20,7 @@ This project only support CUDA 9.0+. If you are using cuda 8.0, please update it
2. Download cmake >= 3.13.2, then add cmake executables to PATH. 2. Download cmake >= 3.13.2, then add cmake executables to PATH.
3. Ensure you have installed pytorch 1.0 in your environment, run ```python setup.py bdist_wheel``` (don't use ```python setup.py install```). 3. Ensure you have installed pytorch 1.0+ in your environment, run ```python setup.py bdist_wheel``` (don't use ```python setup.py install```).
4. Run ```cd ./dist```, use pip to install generated whl file. 4. Run ```cd ./dist```, use pip to install generated whl file.
...@@ -152,6 +152,20 @@ This implementation use gather-gemm-scatter framework to do sparse convolution. ...@@ -152,6 +152,20 @@ This implementation use gather-gemm-scatter framework to do sparse convolution.
* **Bo Li** - *gpu indice generation idea, owner of patent of the sparse conv gpu indice generation algorithm (don't include subm)* - [prclibo](https://github.com/prclibo) * **Bo Li** - *gpu indice generation idea, owner of patent of the sparse conv gpu indice generation algorithm (don't include subm)* - [prclibo](https://github.com/prclibo)
## Third party libraries
* [CUDPP](https://github.com/cudpp/cudpp): A cuda library. contains a cuda hash implementation.
* [robin-map](https://github.com/Tessil/robin-map): A fast c++ hash library. almost 2x faster than std::unordered_map in this project.
* [pybind11](https://github.com/pybind/pybind11): A head-only python c++ binding library.
* [prettyprint](https://github.com/louisdx/cxx-prettyprint): A head-only library for container print.
## License ## License
This project is licensed under the Apache license 2.0 License - see the [LICENSE.md](LICENSE.md) file for details This project is licensed under the Apache license 2.0 License - see the [LICENSE.md](LICENSE.md) file for details
The [CUDPP](https://github.com/cudpp/cudpp) hash code is licensed under BSD License.
The [robin-map](https://github.com/Tessil/robin-map) code is licensed under MIT license.
\ No newline at end of file
include/hash/debugging.h → include/cuhash/debugging.h
View file @ a6ae8967
...@@ -24,7 +24,7 @@ ...@@ -24,7 +24,7 @@
#include <algorithm> #include <algorithm>
namespace cudahash { namespace cuhash {
//! @name Debugging functions //! @name Debugging functions
/// @{ /// @{
......
include/hash/definitions.h → include/cuhash/definitions.h
View file @ a6ae8967
...@@ -58,7 +58,7 @@ inline void PrintMessage(const char *message, const bool error = false) { ...@@ -58,7 +58,7 @@ inline void PrintMessage(const char *message, const bool error = false) {
//! Prints a message out to the console. //! Prints a message out to the console.
inline void PrintMessage(const char *message, const bool error = false) { inline void PrintMessage(const char *message, const bool error = false) {
if (error) { if (error) {
printf("!!! %s\n", message); printf("cudahash: %s\n", message);
} else { } else {
printf("%s\n", message); printf("%s\n", message);
} }
...@@ -68,7 +68,7 @@ inline void PrintMessage(const char *message, const bool error = false) { ...@@ -68,7 +68,7 @@ inline void PrintMessage(const char *message, const bool error = false) {
/* ------------------------------------------------------------------------- /* -------------------------------------------------------------------------
Hash table constants and definitions. Hash table constants and definitions.
------------------------------------------------------------------------- */ ------------------------------------------------------------------------- */
namespace cudahash { namespace cuhash {
/** /**
* \addtogroup cudpp_hash_data_structures * \addtogroup cudpp_hash_data_structures
...@@ -111,6 +111,6 @@ const float kMinimumSpaceUsages[] = {std::numeric_limits<float>::max(), ...@@ -111,6 +111,6 @@ const float kMinimumSpaceUsages[] = {std::numeric_limits<float>::max(),
/** @} */ // end cudpp_hash_data_structures /** @} */ // end cudpp_hash_data_structures
}; // namespace cudahash }; // namespace cuhash
#endif #endif
include/hash/hash_functions.h → include/cuhash/hash_functions.h
View file @ a6ae8967
...@@ -9,7 +9,7 @@ ...@@ -9,7 +9,7 @@
#include <vector_types.h> #include <vector_types.h>
#include "definitions.h" #include "definitions.h"
namespace cudahash { namespace cuhash {
//! Prime number larger than the largest practical hash table size. //! Prime number larger than the largest practical hash table size.
const unsigned kPrimeDivisor = 4294967291u; const unsigned kPrimeDivisor = 4294967291u;
...@@ -89,6 +89,7 @@ unsigned stash_hash_function(const uint2 stash_constants, ...@@ -89,6 +89,7 @@ unsigned stash_hash_function(const uint2 stash_constants,
return (stash_constants.x ^ key + stash_constants.y) % kStashSize; return (stash_constants.x ^ key + stash_constants.y) % kStashSize;
} }
unsigned generate_random_uint32();
}; // namespace CuckooHashing }; // namespace CuckooHashing
......
include/hash/hash_table.cuh → include/cuhash/hash_table.cuh
View file @ a6ae8967
...@@ -22,7 +22,7 @@ ...@@ -22,7 +22,7 @@
#include <tensorview/tensorview.h> #include <tensorview/tensorview.h>
#include <driver_types.h> #include <driver_types.h>
namespace cudahash { namespace cuhash {
//! Makes an 64-bit Entry out of a key-value pair for the hash table. //! Makes an 64-bit Entry out of a key-value pair for the hash table.
TV_HOST_DEVICE_INLINE Entry make_entry(unsigned key, unsigned value) { TV_HOST_DEVICE_INLINE Entry make_entry(unsigned key, unsigned value) {
......
include/hash/hash_table.h → include/cuhash/hash_table.h
View file @ a6ae8967
...@@ -45,7 +45,7 @@ ...@@ -45,7 +45,7 @@
/* ------------------------------------------------------------------------- /* -------------------------------------------------------------------------
Hash table code. Hash table code.
------------------------------------------------------------------------- */ ------------------------------------------------------------------------- */
namespace cudahash { namespace cuhash {
//! Compute how many thread blocks are required for the given number of threads. //! Compute how many thread blocks are required for the given number of threads.
dim3 ComputeGridDim(unsigned threads); dim3 ComputeGridDim(unsigned threads);
...@@ -113,7 +113,7 @@ class HashTable { ...@@ -113,7 +113,7 @@ class HashTable {
* The input keys are expected to be completely unique. * The input keys are expected to be completely unique.
* To reduce the chance of a failure, increase the space usage or number of * To reduce the chance of a failure, increase the space usage or number of
* functions. * functions.
* Keys are not allowed to be equal to cudahash::kKeyEmpty. * Keys are not allowed to be equal to cuhash::kKeyEmpty.
*/ */
virtual bool Build(const unsigned input_size, virtual bool Build(const unsigned input_size,
const unsigned *d_keys, const unsigned *d_keys,
......
include/hash/mt19937ar.h deleted 100644 → 0
View file @ 0757c45b
void init_genrand(unsigned long s);
void init_by_array(unsigned long init_key[], int key_length);
unsigned long genrand_int32(void);
long genrand_int31(void);
double genrand_real1(void);
double genrand_real2(void);
double genrand_real3(void);
double genrand_res53(void);
include/paramsgrid.h
View file @ a6ae8967
...@@ -12,6 +12,8 @@ ...@@ -12,6 +12,8 @@
// See the License for the specific language governing permissions and // See the License for the specific language governing permissions and
// limitations under the License. // limitations under the License.
// This file is used for c++ unit test, but pytorch jit ops don't support c++ debug build.
#ifndef PARAMS_GRID_H_ #ifndef PARAMS_GRID_H_
#define PARAMS_GRID_H_ #define PARAMS_GRID_H_
#include <tuple> #include <tuple>
......
include/spconv/fused_spconv_ops.h
View file @ a6ae8967
...@@ -15,7 +15,6 @@ ...@@ -15,7 +15,6 @@
#ifndef FUSED_SPARSE_CONV_OP_H_ #ifndef FUSED_SPARSE_CONV_OP_H_
#define FUSED_SPARSE_CONV_OP_H_ #define FUSED_SPARSE_CONV_OP_H_
#include <cuda_runtime_api.h>
#include <spconv/indice.h> #include <spconv/indice.h>
#include <spconv/reordering.h> #include <spconv/reordering.h>
#include <torch/script.h> #include <torch/script.h>
...@@ -84,7 +83,9 @@ torch::Tensor fusedIndiceConvBatchNorm(torch::Tensor features, torch::Tensor fil ...@@ -84,7 +83,9 @@ torch::Tensor fusedIndiceConvBatchNorm(torch::Tensor features, torch::Tensor fil
gatherFtor(tv::CPU(), tv::torch2tv<T>(inputBuffer), gatherFtor(tv::CPU(), tv::torch2tv<T>(inputBuffer),
tv::torch2tv<const T>(features), tv::torch2tv<const T>(features),
tv::torch2tv<const int>(indicePairs).subview(i, inverse), nHot); tv::torch2tv<const int>(indicePairs).subview(i, inverse), nHot);
} else { }
#ifdef SPCONV_CUDA
else if (device == torch::kCUDA) {
functor::SparseGatherFunctor<tv::GPU, T, int> gatherFtor; functor::SparseGatherFunctor<tv::GPU, T, int> gatherFtor;
gatherFtor(tv::TorchGPU(), tv::torch2tv<T>(inputBuffer), gatherFtor(tv::TorchGPU(), tv::torch2tv<T>(inputBuffer),
tv::torch2tv<const T>(features), tv::torch2tv<const T>(features),
...@@ -97,6 +98,11 @@ torch::Tensor fusedIndiceConvBatchNorm(torch::Tensor features, torch::Tensor fil ...@@ -97,6 +98,11 @@ torch::Tensor fusedIndiceConvBatchNorm(torch::Tensor features, torch::Tensor fil
torch::index_select_out(inputBufferBlob, features, 0, torch::index_select_out(inputBufferBlob, features, 0,
indicePairBlob);*/ indicePairBlob);*/
} }
#endif
else {
TV_ASSERT_INVALID_ARG(false, "unknown device type");
}
// totalGatherTime += timer.report() / 1000.0; // totalGatherTime += timer.report() / 1000.0;
torch::mm_out(outputBufferBlob, inputBufferBlob, filters[i]); torch::mm_out(outputBufferBlob, inputBufferBlob, filters[i]);
// totalGEMMTime += timer.report() / 1000.0; // totalGEMMTime += timer.report() / 1000.0;
...@@ -107,7 +113,9 @@ torch::Tensor fusedIndiceConvBatchNorm(torch::Tensor features, torch::Tensor fil ...@@ -107,7 +113,9 @@ torch::Tensor fusedIndiceConvBatchNorm(torch::Tensor features, torch::Tensor fil
tv::torch2tv<const T>(outputBuffer), tv::torch2tv<const T>(outputBuffer),
tv::torch2tv<const int>(indicePairs).subview(i, !inverse), nHot, tv::torch2tv<const int>(indicePairs).subview(i, !inverse), nHot,
true); true);
} else { }
#ifdef SPCONV_CUDA
else if (device == torch::kCUDA) {
functor::SparseScatterAddFunctor<tv::GPU, T, int> scatterFtor; functor::SparseScatterAddFunctor<tv::GPU, T, int> scatterFtor;
scatterFtor(tv::TorchGPU(), tv::torch2tv<T>(output), scatterFtor(tv::TorchGPU(), tv::torch2tv<T>(output),
tv::torch2tv<const T>(outputBuffer), tv::torch2tv<const T>(outputBuffer),
...@@ -115,6 +123,11 @@ torch::Tensor fusedIndiceConvBatchNorm(torch::Tensor features, torch::Tensor fil ...@@ -115,6 +123,11 @@ torch::Tensor fusedIndiceConvBatchNorm(torch::Tensor features, torch::Tensor fil
true); true);
TV_CHECK_CUDA_ERR(); TV_CHECK_CUDA_ERR();
} }
#endif
else {
TV_ASSERT_INVALID_ARG(false, "unknown device type");
}
// totalSAddTime += timer.report() / 1000.0; // totalSAddTime += timer.report() / 1000.0;
} }
// std::cout << "gather time " << totalGatherTime << std::endl; // std::cout << "gather time " << totalGatherTime << std::endl;
......
include/spconv/geometry.h
View file @ a6ae8967
...@@ -18,8 +18,51 @@ ...@@ -18,8 +18,51 @@
#include <iostream> #include <iostream>
#include <limits> #include <limits>
#include <tensorview/tensorview.h> #include <tensorview/tensorview.h>
#include <tsl/robin_map.h>
#include <unordered_map>
namespace spconv { namespace spconv {
namespace detail {
template <typename T> struct ToUnsigned;
template <> struct ToUnsigned<int>{
using type = uint32_t;
};
template <> struct ToUnsigned<long>{
using type = uint64_t;
};
template <typename T> struct FNVInternal;
template <> struct FNVInternal<uint32_t>
{
constexpr static uint32_t defaultOffsetBasis = 0x811C9DC5;
constexpr static uint32_t prime = 0x01000193;
};
template <> struct FNVInternal<uint64_t>
{
constexpr static uint64_t defaultOffsetBasis = 0xcbf29ce484222325;
constexpr static uint64_t prime = 0x100000001b3;
};
}
template <typename T>
using to_unsigned_t = typename detail::ToUnsigned<std::remove_const_t<T>>::type;
template <typename T>
struct FNV1a : detail::FNVInternal<T>{
std::size_t operator()(const T* data, std::size_t size){
to_unsigned_t<T> hash = detail::FNVInternal<T>::defaultOffsetBasis;
for (std::size_t i = 0; i < size; ++i) {
hash *= detail::FNVInternal<T>::prime;
hash ^= static_cast<to_unsigned_t<T>>(data[i]);
}
return hash;
}
};
template <typename Index, unsigned NDim> template <typename Index, unsigned NDim>
TV_HOST_DEVICE Index getValidOutPos(const Index *input_pos, TV_HOST_DEVICE Index getValidOutPos(const Index *input_pos,
const Index *kernelSize, const Index *kernelSize,
...@@ -169,6 +212,7 @@ Index getIndicePairsConv(tv::TensorView<const Index> indicesIn, ...@@ -169,6 +212,7 @@ Index getIndicePairsConv(tv::TensorView<const Index> indicesIn,
std::vector<Index> validPoints_(kernelVolume * (NDim + 1)); std::vector<Index> validPoints_(kernelVolume * (NDim + 1));
Index* validPoints = validPoints_.data(); Index* validPoints = validPoints_.data();
Index *pointPtr = nullptr; Index *pointPtr = nullptr;
tsl::robin_map<Index, Index> hash;
for (int j = 0; j < numActIn; ++j) { for (int j = 0; j < numActIn; ++j) {
batchIdx = indicesIn(j, 0); batchIdx = indicesIn(j, 0);
numValidPoints = getValidOutPos<Index, NDim>( numValidPoints = getValidOutPos<Index, NDim>(
...@@ -179,16 +223,16 @@ Index getIndicePairsConv(tv::TensorView<const Index> indicesIn, ...@@ -179,16 +223,16 @@ Index getIndicePairsConv(tv::TensorView<const Index> indicesIn,
auto offset = pointPtr[NDim]; auto offset = pointPtr[NDim];
auto index = tv::rowArrayIdx<Index, NDim>(pointPtr, outSpatialShape) + auto index = tv::rowArrayIdx<Index, NDim>(pointPtr, outSpatialShape) +
spatialVolume * batchIdx; spatialVolume * batchIdx;
if (gridsOut[index] == -1) { if (hash.find(index) == hash.end()) {
for (unsigned k = 1; k < NDim + 1; ++k) { for (unsigned k = 1; k < NDim + 1; ++k) {
indicesOut(numAct, k) = pointPtr[k - 1]; indicesOut(numAct, k) = pointPtr[k - 1];
} }
indicesOut(numAct, 0) = batchIdx; indicesOut(numAct, 0) = batchIdx;
gridsOut[index] = numAct++; hash[index] = numAct++;
} }
// indicePairs: [K, 2, L] // indicePairs: [K, 2, L]
indicePairs(offset, 0, indiceNum[offset]) = j; indicePairs(offset, 0, indiceNum[offset]) = j;
indicePairs(offset, 1, indiceNum[offset]++) = gridsOut[index]; indicePairs(offset, 1, indiceNum[offset]++) = hash[index];
} }
} }
return numAct; return numAct;
...@@ -220,6 +264,7 @@ Index getIndicePairsDeConv(tv::TensorView<const Index> indicesIn, ...@@ -220,6 +264,7 @@ Index getIndicePairsDeConv(tv::TensorView<const Index> indicesIn,
std::vector<Index> validPoints_(kernelVolume * (NDim + 1)); std::vector<Index> validPoints_(kernelVolume * (NDim + 1));
Index* validPoints = validPoints_.data(); Index* validPoints = validPoints_.data();
Index *pointPtr = nullptr; Index *pointPtr = nullptr;
tsl::robin_map<Index, Index> hash;
for (int j = 0; j < numActIn; ++j) { for (int j = 0; j < numActIn; ++j) {
batchIdx = indicesIn(j, 0); batchIdx = indicesIn(j, 0);
numValidPoints = getValidOutPosTranspose<Index, NDim>( numValidPoints = getValidOutPosTranspose<Index, NDim>(
...@@ -230,16 +275,16 @@ Index getIndicePairsDeConv(tv::TensorView<const Index> indicesIn, ...@@ -230,16 +275,16 @@ Index getIndicePairsDeConv(tv::TensorView<const Index> indicesIn,
auto offset = pointPtr[NDim]; auto offset = pointPtr[NDim];
auto index = tv::rowArrayIdx<Index, NDim>(pointPtr, outSpatialShape) + auto index = tv::rowArrayIdx<Index, NDim>(pointPtr, outSpatialShape) +
spatialVolume * batchIdx; spatialVolume * batchIdx;
if (gridsOut[index] == -1) { if (hash.find(index) == hash.end()) {
for (unsigned k = 1; k < NDim + 1; ++k) { for (unsigned k = 1; k < NDim + 1; ++k) {
indicesOut(numAct, k) = pointPtr[k - 1]; indicesOut(numAct, k) = pointPtr[k - 1];
} }
indicesOut(numAct, 0) = batchIdx; indicesOut(numAct, 0) = batchIdx;
gridsOut[index] = numAct++; hash[index] = numAct++;
} }
// indicePairs: [K, 2, L] // indicePairs: [K, 2, L]
indicePairs(offset, 0, indiceNum[offset]) = j; indicePairs(offset, 0, indiceNum[offset]) = j;
indicePairs(offset, 1, indiceNum[offset]++) = gridsOut[index]; indicePairs(offset, 1, indiceNum[offset]++) = hash[index];
} }
} }
return numAct; return numAct;
...@@ -271,12 +316,13 @@ Index getIndicePairsSubM(tv::TensorView<const Index> indicesIn, ...@@ -271,12 +316,13 @@ Index getIndicePairsSubM(tv::TensorView<const Index> indicesIn,
std::vector<Index> validPoints_(kernelVolume * (NDim + 1)); std::vector<Index> validPoints_(kernelVolume * (NDim + 1));
Index* validPoints = validPoints_.data(); Index* validPoints = validPoints_.data();
Index *pointPtr = nullptr; Index *pointPtr = nullptr;
tsl::robin_map<Index, Index> hash;
for (int j = 0; j < numActIn; ++j) { for (int j = 0; j < numActIn; ++j) {
Index index = 0; Index index = 0;
index = tv::rowArrayIdx<Index, NDim>(indicesIn.data() + j * (NDim + 1) + 1, index = tv::rowArrayIdx<Index, NDim>(indicesIn.data() + j * (NDim + 1) + 1,
outSpatialShape) + outSpatialShape) +
spatialVolume * indicesIn(j, 0); spatialVolume * indicesIn(j, 0);
gridsOut[index] = j; hash[index] = j;
} }
Index index = 0; Index index = 0;
for (int j = 0; j < numActIn; ++j) { for (int j = 0; j < numActIn; ++j) {
...@@ -288,9 +334,9 @@ Index getIndicePairsSubM(tv::TensorView<const Index> indicesIn, ...@@ -288,9 +334,9 @@ Index getIndicePairsSubM(tv::TensorView<const Index> indicesIn,
auto offset = pointPtr[NDim]; auto offset = pointPtr[NDim];
index = tv::rowArrayIdx<Index, NDim>(pointPtr, outSpatialShape) + index = tv::rowArrayIdx<Index, NDim>(pointPtr, outSpatialShape) +
spatialVolume * indicesIn(j, 0); spatialVolume * indicesIn(j, 0);
if (gridsOut[index] > -1) { if (hash.find(index) == hash.end()) {
indicePairs(offset, 0, indiceNum[offset]) = j; indicePairs(offset, 0, indiceNum[offset]) = j;
indicePairs(offset, 1, indiceNum[offset]++) = gridsOut[index]; indicePairs(offset, 1, indiceNum[offset]++) = hash[index];
} }
} }
} }
......
include/spconv/indice.cu.h
View file @ a6ae8967
...@@ -14,18 +14,18 @@ ...@@ -14,18 +14,18 @@
#ifndef INDICE_CU_H_ #ifndef INDICE_CU_H_
#define INDICE_CU_H_ #define INDICE_CU_H_
#include <hash/hash_table.cuh> #include <cuhash/hash_table.cuh>
#include <spconv/geometry.h> #include <spconv/geometry.h>
#include <tensorview/helper_kernel.cu.h> #include <tensorview/helper_kernel.cu.h>
#include <tensorview/tensorview.h> #include <tensorview/tensorview.h>
namespace spconv { namespace spconv {
template <typename Index, typename IndexGrid, unsigned NDim, template <typename Index, typename IndexGrid, unsigned NDim,
int KernelMaxVolume = 256> int KernelMaxVolume = 256, typename Index1D=int>
__global__ void prepareIndicePairsKernel( __global__ void prepareIndicePairsKernel(
tv::TensorView<const Index> indicesIn, tv::TensorView<Index> indicesOut, tv::TensorView<const Index> indicesIn, tv::TensorView<Index> indicesOut,
tv::TensorView<IndexGrid> gridsOut, tv::TensorView<Index> indicePairs, tv::TensorView<IndexGrid> gridsOut, tv::TensorView<Index> indicePairs,
tv::TensorView<Index> indiceNum, tv::TensorView<Index> indicePairUnique, tv::TensorView<Index> indiceNum, tv::TensorView<Index1D> indicePairUnique,
const tv::SimpleVector<Index, NDim> kernelSize, const tv::SimpleVector<Index, NDim> kernelSize,
const tv::SimpleVector<Index, NDim> stride, const tv::SimpleVector<Index, NDim> stride,
const tv::SimpleVector<Index, NDim> padding, const tv::SimpleVector<Index, NDim> padding,
...@@ -151,8 +151,8 @@ __global__ void ...@@ -151,8 +151,8 @@ __global__ void
assignIndicePairsHashKernel(tv::TensorView<Index> indicesOut, int numActIn, assignIndicePairsHashKernel(tv::TensorView<Index> indicesOut, int numActIn,
tv::TensorView<Index> indicePairs, tv::TensorView<Index> indicePairs,
tv::TensorView<Index> indicePairUnique, tv::TensorView<Index> indicePairUnique,
unsigned table_size, const cudahash::Entry *table, unsigned table_size, const cuhash::Entry *table,
cudahash::Functions<kNumHashFunctions> constants, cuhash::Functions<kNumHashFunctions> constants,
uint2 stash_constants, unsigned stash_count) { uint2 stash_constants, unsigned stash_count) {
Index index; Index index;
...@@ -162,9 +162,9 @@ assignIndicePairsHashKernel(tv::TensorView<Index> indicesOut, int numActIn, ...@@ -162,9 +162,9 @@ assignIndicePairsHashKernel(tv::TensorView<Index> indicesOut, int numActIn,
index = indicePairs(i, 1, ix); index = indicePairs(i, 1, ix);
if (index > -1) { if (index > -1) {
auto val = auto val =
cudahash::retrieve((unsigned)(index), table_size, cuhash::retrieve((unsigned)(index), table_size,
table, constants, stash_constants, stash_count); table, constants, stash_constants, stash_count);
assert(val != cudahash::kNotFound); assert(val != cuhash::kNotFound);
indicePairs(i, 1, ix) = (unsigned)val; indicePairs(i, 1, ix) = (unsigned)val;
} }
} }
...@@ -283,8 +283,8 @@ __global__ void getSubMIndicePairsHashKernel( ...@@ -283,8 +283,8 @@ __global__ void getSubMIndicePairsHashKernel(
const tv::SimpleVector<Index, NDim> padding, const tv::SimpleVector<Index, NDim> padding,
const tv::SimpleVector<Index, NDim> dilation, const tv::SimpleVector<Index, NDim> dilation,
const tv::SimpleVector<Index, NDim> outSpatialShape, const tv::SimpleVector<Index, NDim> outSpatialShape,
unsigned table_size, const cudahash::Entry *table, unsigned table_size, const cuhash::Entry *table,
cudahash::Functions<kNumHashFunctions> constants, cuhash::Functions<kNumHashFunctions> constants,
uint2 stash_constants, unsigned stash_count) { uint2 stash_constants, unsigned stash_count) {
auto numActIn = indicesIn.dim(0); auto numActIn = indicesIn.dim(0);
Index spatialVolume = 1; Index spatialVolume = 1;
...@@ -307,9 +307,9 @@ __global__ void getSubMIndicePairsHashKernel( ...@@ -307,9 +307,9 @@ __global__ void getSubMIndicePairsHashKernel(
index = tv::rowArrayIdx<Index, NDim>(pointPtr, outSpatialShape.data()) + index = tv::rowArrayIdx<Index, NDim>(pointPtr, outSpatialShape.data()) +
spatialVolume * indicesIn(ix, 0); spatialVolume * indicesIn(ix, 0);
auto val = auto val =
cudahash::retrieve((unsigned)(index), table_size, cuhash::retrieve((unsigned)(index), table_size,
table, constants, stash_constants, stash_count); table, constants, stash_constants, stash_count);
if (val != cudahash::kNotFound) { if (val != cuhash::kNotFound) {
auto oldNum = atomicAdd(indiceNum.data() + offset, Index(1)); auto oldNum = atomicAdd(indiceNum.data() + offset, Index(1));
indicePairs(offset, 1, oldNum) = val; indicePairs(offset, 1, oldNum) = val;
indicePairs(offset, 0, oldNum) = ix; indicePairs(offset, 0, oldNum) = ix;
......
include/spconv/nms.h
View file @ a6ae8967
...@@ -181,7 +181,7 @@ std::vector<int> rotate_non_max_suppression_cpu(py::array_t<DType> box_corners, ...@@ -181,7 +181,7 @@ std::vector<int> rotate_non_max_suppression_cpu(py::array_t<DType> box_corners,
} }
return keep; return keep;
} }
#ifdef SPCONV_CUDA
constexpr int const threadsPerBlock = sizeof(unsigned long long) * 8; constexpr int const threadsPerBlock = sizeof(unsigned long long) * 8;
template <typename DType> template <typename DType>
...@@ -196,6 +196,7 @@ int non_max_suppression(py::array_t<DType> boxes, py::array_t<int> keep_out, ...@@ -196,6 +196,7 @@ int non_max_suppression(py::array_t<DType> boxes, py::array_t<int> keep_out,
boxes.shape(0), boxes.shape(1), boxes.shape(0), boxes.shape(1),
nms_overlap_thresh, device_id); nms_overlap_thresh, device_id);
} }
#endif
} // namespace spconv } // namespace spconv
#endif #endif
include/spconv/nms_ops.h
View file @ a6ae8967
...@@ -15,7 +15,6 @@ ...@@ -15,7 +15,6 @@
#ifndef NMS_TORCH_OP_H_ #ifndef NMS_TORCH_OP_H_
#define NMS_TORCH_OP_H_ #define NMS_TORCH_OP_H_
#include <cuda_runtime_api.h>
#include <spconv/indice.h> #include <spconv/indice.h>
#include <spconv/reordering.h> #include <spconv/reordering.h>
#include <torch/script.h> #include <torch/script.h>
......
include/spconv/pillar_scatter_ops.h
View file @ a6ae8967
...@@ -15,7 +15,6 @@ ...@@ -15,7 +15,6 @@
#ifndef PILLAR_SCATTER_OP_H_ #ifndef PILLAR_SCATTER_OP_H_
#define PILLAR_SCATTER_OP_H_ #define PILLAR_SCATTER_OP_H_
#include <cuda_runtime_api.h>
#include <spconv/pillar_scatter_functor.h> #include <spconv/pillar_scatter_functor.h>
#include <torch/script.h> #include <torch/script.h>
#include <torch_utils.h> #include <torch_utils.h>
...@@ -28,6 +27,7 @@ template <typename T> ...@@ -28,6 +27,7 @@ template <typename T>
torch::Tensor pointPillarScatter(torch::Tensor features, torch::Tensor coors, torch::Tensor pointPillarScatter(torch::Tensor features, torch::Tensor coors,
torch::Tensor shape) { torch::Tensor shape) {
TV_ASSERT_RT_ERR(shape.device().type() == torch::kCPU, "error"); TV_ASSERT_RT_ERR(shape.device().type() == torch::kCPU, "error");
TV_ASSERT_RT_ERR(features.device().type() == torch::kCUDA, "error");
TV_ASSERT_RT_ERR(shape.dim() == 1, "error"); TV_ASSERT_RT_ERR(shape.dim() == 1, "error");
TV_ASSERT_RT_ERR(shape.size(0) == 4, "error"); TV_ASSERT_RT_ERR(shape.size(0) == 4, "error");
TV_ASSERT_RT_ERR(features.dim() >= 3, "error"); TV_ASSERT_RT_ERR(features.dim() >= 3, "error");
...@@ -42,10 +42,11 @@ torch::Tensor pointPillarScatter(torch::Tensor features, torch::Tensor coors, ...@@ -42,10 +42,11 @@ torch::Tensor pointPillarScatter(torch::Tensor features, torch::Tensor coors,
torch::zeros({shapeData[0], shapeData[1], shapeData[2], shapeData[3]}, torch::zeros({shapeData[0], shapeData[1], shapeData[2], shapeData[3]},
features.options()); features.options());
TV_ASSERT_RT_ERR(shapeData[1] == features.size(1), "error"); TV_ASSERT_RT_ERR(shapeData[1] == features.size(1), "error");
#ifdef SPCONV_CUDA
functor::PointPillarScatter<tv::GPU, T, int> ftor; functor::PointPillarScatter<tv::GPU, T, int> ftor;
ftor(tv::TorchGPU(), tv::torch2tv<T>(canvas), tv::torch2tv<const T>(features.squeeze()), ftor(tv::TorchGPU(), tv::torch2tv<T>(canvas), tv::torch2tv<const T>(features.squeeze()),
tv::torch2tv<const T>(coors.squeeze())); tv::torch2tv<const T>(coors.squeeze()));
#endif
return canvas; return canvas;
} }
......
include/spconv/pool_ops.h
View file @ a6ae8967
...@@ -15,7 +15,6 @@ ...@@ -15,7 +15,6 @@
#ifndef SPARSE_POOL_OP_H_ #ifndef SPARSE_POOL_OP_H_
#define SPARSE_POOL_OP_H_ #define SPARSE_POOL_OP_H_
#include <cuda_runtime_api.h>
#include <spconv/maxpool.h> #include <spconv/maxpool.h>
#include <torch/script.h> #include <torch/script.h>
#include <torch_utils.h> #include <torch_utils.h>
...@@ -44,13 +43,19 @@ torch::Tensor indiceMaxPool(torch::Tensor features, torch::Tensor indicePairs, ...@@ -44,13 +43,19 @@ torch::Tensor indiceMaxPool(torch::Tensor features, torch::Tensor indicePairs,
forwardFtor(tv::CPU(), tv::torch2tv<T>(output), forwardFtor(tv::CPU(), tv::torch2tv<T>(output),
tv::torch2tv<const T>(features), tv::torch2tv<const T>(features),
tv::torch2tv<const int>(indicePairs).subview(i), nHot); tv::torch2tv<const int>(indicePairs).subview(i), nHot);
} else { }
#ifdef SPCONV_CUDA
else if (device == torch::kCUDA) {
functor::SparseMaxPoolForwardFunctor<tv::GPU, T, int> forwardFtor; functor::SparseMaxPoolForwardFunctor<tv::GPU, T, int> forwardFtor;
forwardFtor(tv::TorchGPU(), tv::torch2tv<T>(output), forwardFtor(tv::TorchGPU(), tv::torch2tv<T>(output),
tv::torch2tv<const T>(features), tv::torch2tv<const T>(features),
tv::torch2tv<const int>(indicePairs).subview(i), nHot); tv::torch2tv<const int>(indicePairs).subview(i), nHot);
TV_CHECK_CUDA_ERR(); TV_CHECK_CUDA_ERR();
} }
#endif
else{
TV_ASSERT_INVALID_ARG(false, "unknown device type");
}
// totalTime += timer.report() / 1000.0; // totalTime += timer.report() / 1000.0;
} }
// std::cout << "maxpool forward time " << totalTime << std::endl; // std::cout << "maxpool forward time " << totalTime << std::endl;
...@@ -80,7 +85,9 @@ torch::Tensor indiceMaxPoolBackward(torch::Tensor features, ...@@ -80,7 +85,9 @@ torch::Tensor indiceMaxPoolBackward(torch::Tensor features,
tv::torch2tv<const T>(features), tv::torch2tv<const T>(features),
tv::torch2tv<const T>(outGrad), tv::torch2tv<T>(inputGrad), tv::torch2tv<const T>(outGrad), tv::torch2tv<T>(inputGrad),
tv::torch2tv<const int>(indicePairs).subview(i), nHot); tv::torch2tv<const int>(indicePairs).subview(i), nHot);
} else { }
#ifdef SPCONV_CUDA
else if (device == torch::kCUDA) {
functor::SparseMaxPoolBackwardFunctor<tv::GPU, T, int> backwardFtor; functor::SparseMaxPoolBackwardFunctor<tv::GPU, T, int> backwardFtor;
backwardFtor(tv::TorchGPU(), tv::torch2tv<const T>(outFeatures), backwardFtor(tv::TorchGPU(), tv::torch2tv<const T>(outFeatures),
tv::torch2tv<const T>(features), tv::torch2tv<const T>(features),
...@@ -88,6 +95,11 @@ torch::Tensor indiceMaxPoolBackward(torch::Tensor features, ...@@ -88,6 +95,11 @@ torch::Tensor indiceMaxPoolBackward(torch::Tensor features,
tv::torch2tv<const int>(indicePairs).subview(i), nHot); tv::torch2tv<const int>(indicePairs).subview(i), nHot);
TV_CHECK_CUDA_ERR(); TV_CHECK_CUDA_ERR();
} }
#endif
else{
TV_ASSERT_INVALID_ARG(false, "unknown device type");
}
} }
return inputGrad; return inputGrad;
} }
......
include/spconv/spconv_ops.h
View file @ a6ae8967
This diff is collapsed.
include/tensorview/tensorview.h
View file @ a6ae8967
...@@ -16,17 +16,22 @@ ...@@ -16,17 +16,22 @@
#include <algorithm> #include <algorithm>
#include <cassert> #include <cassert>
#include <cstdlib> #include <cstdlib>
#include <cuda_runtime_api.h>
#include <iostream> #include <iostream>
#include <memory> #include <memory>
// #include <prettyprint.h> // #include <prettyprint.h>
#include <sstream> #include <sstream>
#include <type_traits> #include <type_traits>
#include <vector> #include <vector>
#ifdef SPCONV_CUDA
#include <cuda_runtime_api.h>
#endif
namespace tv { namespace tv {
#ifdef __NVCC__ #ifdef __NVCC__
#define TV_HOST_DEVICE_INLINE __forceinline__ __device__ __host__ #define TV_HOST_DEVICE_INLINE __forceinline__ __device__ __host__
#define TV_DEVICE_INLINE __forceinline__ __device__ #define TV_DEVICE_INLINE __forceinline__ __device__
#define TV_HOST_DEVICE __device__ __host__ #define TV_HOST_DEVICE __device__ __host__
...@@ -113,12 +118,13 @@ void sstream_print(SStream &ss, T val, TArgs... args) { ...@@ -113,12 +118,13 @@ void sstream_print(SStream &ss, T val, TArgs... args) {
} \ } \
} }
#ifdef SPCONV_CUDA
struct GPU { struct GPU {
GPU(cudaStream_t s = 0) : mStream(s) {} GPU(cudaStream_t s = 0) : mStream(s) {}
virtual cudaStream_t getStream() const { return mStream; } virtual cudaStream_t getStream() const { return mStream; }
cudaStream_t mStream = 0; cudaStream_t mStream = 0;
}; };
#endif
struct CPU {}; struct CPU {};
#define TV_MAX_DIM 6 #define TV_MAX_DIM 6
......
include/torch_utils.h
View file @ a6ae8967
...@@ -16,16 +16,19 @@ ...@@ -16,16 +16,19 @@
#include <tensorview/tensorview.h> #include <tensorview/tensorview.h>
#include <torch/script.h> #include <torch/script.h>
#include <ATen/ATen.h> #include <ATen/ATen.h>
#ifdef SPCONV_CUDA
#include <ATen/cuda/CUDAContext.h> #include <ATen/cuda/CUDAContext.h>
#endif
namespace tv { namespace tv {
#ifdef SPCONV_CUDA
struct TorchGPU: public tv::GPU { struct TorchGPU: public tv::GPU {
virtual cudaStream_t getStream() const override { virtual cudaStream_t getStream() const override {
return at::cuda::getCurrentCUDAStream(); return at::cuda::getCurrentCUDAStream();
} }
}; };
#endif
template <typename T> void check_torch_dtype(const torch::Tensor &tensor) { template <typename T> void check_torch_dtype(const torch::Tensor &tensor) {
switch (tensor.type().scalarType()) { switch (tensor.type().scalarType()) {
case at::ScalarType::Double: { case at::ScalarType::Double: {
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment