From 65ba5c9844d3e1c8ac497e3968f7e7a7b1548d36 Mon Sep 17 00:00:00 2001 From: Harshvardhan Chandirasekar Date: Sun, 5 Sep 2021 01:06:56 +0530 Subject: [PATCH] tkDNN can now deserialize tensorrt-8 engine (both through test_* and trtexec) but demo has issues in yolo::computeDetections --- CMakeLists.txt | 10 +- include/tkDNN/NetworkRT.h | 53 +-- include/tkDNN/pluginsRT/ActivationLeakyRT.h | 162 ++----- .../tkDNN/pluginsRT/ActivationLogisticRT.h | 160 ++----- include/tkDNN/pluginsRT/ActivationMishRT.h | 148 ++---- .../tkDNN/pluginsRT/ActivationReLUCeilingRT.h | 166 ++----- include/tkDNN/pluginsRT/DeformableConvRT.h | 421 ++++-------------- include/tkDNN/pluginsRT/FlattenConcatRT.h | 181 +++----- .../tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h | 215 +++------ include/tkDNN/pluginsRT/RegionRT.h | 223 +++------- include/tkDNN/pluginsRT/ReorgRT.h | 178 +++----- include/tkDNN/pluginsRT/ReshapeRT.h | 183 +++----- include/tkDNN/pluginsRT/ResizeLayerRT.h | 206 +++------ include/tkDNN/pluginsRT/RouteRT.h | 215 +++------ include/tkDNN/pluginsRT/ShortcutRT.h | 204 +++------ include/tkDNN/pluginsRT/UpsampleRT.h | 189 +++----- include/tkDNN/pluginsRT/YoloRT.h | 323 ++++---------- include/tkDNN/utils.h | 13 + src/DarknetParser.cpp | 9 +- src/Yolo.cpp | 5 +- src/pluginsRT/ActivationLeakyRT.cpp | 138 ++++++ src/pluginsRT/ActivationLogisticRT.cpp | 126 ++++++ src/pluginsRT/ActivationMishRT.cpp | 121 +++++ src/pluginsRT/ActivationReLUCeilingRT.cpp | 129 ++++++ src/pluginsRT/DeformableConvRT.cpp | 331 ++++++++++++++ src/pluginsRT/FlattenConcatRT.cpp | 163 +++++++ src/pluginsRT/MaxPoolingSizeRT.cpp | 179 ++++++++ src/pluginsRT/RegionRT.cpp | 174 ++++++++ src/pluginsRT/ReorgRT.cpp | 151 +++++++ src/pluginsRT/ReshapeRT.cpp | 152 +++++++ src/pluginsRT/ResizeLayerRT.cpp | 163 +++++++ src/pluginsRT/RouteRT.cpp | 183 ++++++++ src/pluginsRT/ShortcutRT.cpp | 167 +++++++ src/pluginsRT/UpsampleRT.cpp | 155 +++++++ src/pluginsRT/YoloRT.cpp | 253 +++++++++++ 35 files changed, 3527 insertions(+), 2322 deletions(-) create mode 100644 src/pluginsRT/ActivationLeakyRT.cpp create mode 100644 src/pluginsRT/ActivationLogisticRT.cpp create mode 100644 src/pluginsRT/ActivationMishRT.cpp create mode 100644 src/pluginsRT/ActivationReLUCeilingRT.cpp create mode 100644 src/pluginsRT/DeformableConvRT.cpp create mode 100644 src/pluginsRT/FlattenConcatRT.cpp create mode 100644 src/pluginsRT/MaxPoolingSizeRT.cpp create mode 100644 src/pluginsRT/RegionRT.cpp create mode 100644 src/pluginsRT/ReorgRT.cpp create mode 100644 src/pluginsRT/ReshapeRT.cpp create mode 100644 src/pluginsRT/ResizeLayerRT.cpp create mode 100644 src/pluginsRT/RouteRT.cpp create mode 100644 src/pluginsRT/ShortcutRT.cpp create mode 100644 src/pluginsRT/UpsampleRT.cpp create mode 100644 src/pluginsRT/YoloRT.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 9200b8c..f3471e2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -31,7 +31,7 @@ endif() find_package(CUDA 9.0 REQUIRED) SET(CUDA_SEPARABLE_COMPILATION ON) #set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_30 --compiler-options '-fPIC'") -set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32 -G -g) +set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32) find_package(CUDNN REQUIRED) @@ -39,8 +39,10 @@ include_directories(${CUDNN_INCLUDE_DIR}) # compile -file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu" ) +file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu" "src/pluginsRT/*.cpp") cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS}) +cuda_add_library(kernels SHARED ${tkdnn_CUSRC}) +target_link_libraries(kernels ${CUDA_CUBLAS_LIBRAY} ${CUDA_LIBRARIES} ${CUDNN_LIBRARIES}) @@ -64,11 +66,11 @@ find_package(yaml-cpp REQUIRED) # Build Libraries #------------------------------------------------------------------------------- file(GLOB tkdnn_SRC "src/*.cpp") -set(tkdnn_LIBS ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDNN_LIBRARIES} ${OpenCV_LIBS} yaml-cpp) +set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDNN_LIBRARIES} ${OpenCV_LIBS} yaml-cpp) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES}) -cuda_add_library(tkDNN SHARED ${tkdnn_SRC} ${tkdnn_CUSRC}) +add_library(tkDNN SHARED ${tkdnn_SRC}) target_link_libraries(tkDNN ${tkdnn_LIBS} ${CUDA_CUBLAS_LIBRAY}) #static diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 95ecb5e..94a60f3 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -8,48 +8,25 @@ #include "NvInfer.h" #include #include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include namespace tk { namespace dnn { -template void writeBUF(char*& buffer, const T& val) -{ - *reinterpret_cast(buffer) = val; - buffer += sizeof(T); -} - -template T readBUF(const char*& buffer) -{ - T val = *reinterpret_cast(buffer); - buffer += sizeof(T); - return val; -} - using namespace nvinfer1; -#include "pluginsRT/ActivationLeakyRT.h" -#include "pluginsRT/ActivationLogisticRT.h" -#include "pluginsRT/ActivationReLUCeilingRT.h" -#include "pluginsRT/ActivationMishRT.h" -#include "pluginsRT/ReorgRT.h" -#include "pluginsRT/RegionRT.h" -#include "pluginsRT/RouteRT.h" -#include "pluginsRT/ShortcutRT.h" -#include "pluginsRT/YoloRT.h" -#include "pluginsRT/UpsampleRT.h" -#include "pluginsRT/ResizeLayerRT.h" -#include "pluginsRT/DeformableConvRT.h" -#include "pluginsRT/FlattenConcatRT.h" -#include "pluginsRT/ReshapeRT.h" -#include "pluginsRT/MaxPoolingFixedSizeRT.h" - -/* -class PluginFactory : IPlugin -{ -public: - YoloRT *yolos[16]; - int n_yolos; - - virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength); -};*/ diff --git a/include/tkDNN/pluginsRT/ActivationLeakyRT.h b/include/tkDNN/pluginsRT/ActivationLeakyRT.h index e73ca2b..ae395ef 100644 --- a/include/tkDNN/pluginsRT/ActivationLeakyRT.h +++ b/include/tkDNN/pluginsRT/ActivationLeakyRT.h @@ -1,148 +1,84 @@ +#include "NvInfer.h" #include "../kernels.h" #include +#include -class ActivationLeakyRT : public IPluginV2 { +namespace nvinfer1 { + class ActivationLeakyRT : public IPluginV2 { -public: - ActivationLeakyRT(float s) { slope = s; } + public: + explicit ActivationLeakyRT(float s); - ActivationLeakyRT(const void *data, size_t length) - { - std::cout<<"DESERIALIZE LEAKYRT"<(data),*bufCheck = buf; - slope = readBUF(buf); - size = readBUF(buf); - assert(buf == bufCheck + length); + ActivationLeakyRT(const void *data, size_t length); - } + ~ActivationLeakyRT(); - ~ActivationLeakyRT() {} + int getNbOutputs() const NOEXCEPT override; - int getNbOutputs() const NOEXCEPT override { return 1; } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override; - Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override { - return inputs[0]; - } + void + configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override; - void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override - { - assert(type == DataType::kFLOAT && format == PluginFormat::kLINEAR); - size = 1; - for (int i = 0; i < outputDims[0].nbDims; i++) - size *= outputDims[0].d[i]; - } + int initialize() NOEXCEPT override; - int initialize() NOEXCEPT override { return 0; } + void terminate() NOEXCEPT override {} - virtual void terminate() NOEXCEPT override {} + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { - return 0; - } + int enqueue(int batchSize, void const *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT override; - virtual int enqueue(int batchSize, void const *const *inputs, void *const *outputs, void *workspace, - cudaStream_t stream) NOEXCEPT override { - activationLEAKYForward( - (dnnType *) reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), batchSize * size, slope, - stream); - return 0; - } + size_t getSerializationSize() const NOEXCEPT override; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 1 * sizeof(int) + 1 * sizeof(float); - } + void serialize(void *buffer) const NOEXCEPT override; - virtual void serialize(void *buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer), *a = buf; - tk::dnn::writeBUF(buf, size); - assert(buf == a + getSerializationSize()); - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override; - bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override { - return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR); - } + const char *getPluginType() const NOEXCEPT override; - const char *getPluginType() const NOEXCEPT override { - return "ActivationLeakyRT_tkDNN"; - } + const char *getPluginVersion() const NOEXCEPT override; - const char *getPluginVersion() const NOEXCEPT override { - return "1"; - } + void destroy() NOEXCEPT override; - void destroy() NOEXCEPT override { delete this; } + const char *getPluginNamespace() const NOEXCEPT override; - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } + IPluginV2 *clone() const NOEXCEPT override; - IPluginV2* clone() const NOEXCEPT override { - ActivationLeakyRT *p = new ActivationLeakyRT(slope); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + int size; + float slope; - int size; - float slope; + private: + std::string mPluginNamespace; + }; -private: - std::string mPluginNamespace; -}; + class ActivationLeakyRTPluginCreator : public IPluginCreator { + public: + ActivationLeakyRTPluginCreator(); -class ActivationLeakyRTPluginCreator : public IPluginCreator { -public: - ActivationLeakyRTPluginCreator() { - mPluginAttributes.emplace_back( - PluginField("slope", nullptr, PluginFieldType::kFLOAT32, 1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override; - IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override { - ActivationLeakyRT *pluginObj = new ActivationLeakyRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } + const char *getPluginNamespace() const NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; - IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override { - const PluginField *fields = fc->fields; - assert(fc->nbFields == 1); - assert(fields[0].type == PluginFieldType::kFLOAT32); - float slope = *(static_cast(fields[0].data)); - ActivationLeakyRT *pluginObj = new ActivationLeakyRT(slope); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } + const char *getPluginName() const NOEXCEPT override ; - const char *getPluginName() const NOEXCEPT override{ - return "ActivationLeakyRT_tkDNN"; - } + const char *getPluginVersion() const NOEXCEPT override ; - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } + const PluginFieldCollection *getFieldNames() NOEXCEPT override; - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; -REGISTER_TENSORRT_PLUGIN(ActivationLeakyRTPluginCreator); \ No newline at end of file + REGISTER_TENSORRT_PLUGIN(ActivationLeakyRTPluginCreator); +}; \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationLogisticRT.h b/include/tkDNN/pluginsRT/ActivationLogisticRT.h index e40e7b4..9646efc 100644 --- a/include/tkDNN/pluginsRT/ActivationLogisticRT.h +++ b/include/tkDNN/pluginsRT/ActivationLogisticRT.h @@ -1,147 +1,83 @@ #include #include "../kernels.h" +#include +#include +#include -class ActivationLogisticRT : public IPluginV2 { +namespace nvinfer1 { -public: - ActivationLogisticRT() { + class ActivationLogisticRT : public IPluginV2 { - } + public: + ActivationLogisticRT() ; - ActivationLogisticRT(const void *data, size_t length) - { - const char* buf = reinterpret_cast(data),*bufCheck = buf; - size = readBUF(buf); - assert(buf == bufCheck + length); + ActivationLogisticRT(const void *data, size_t length) ; - } + ~ActivationLogisticRT() ; - ~ActivationLogisticRT(){ + int getNbOutputs() const NOEXCEPT override ; - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return inputs[0]; - } + int initialize() NOEXCEPT override ; - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override { - size = 1; - for(int i=0; i(inputs[0]), - reinterpret_cast(outputs[0]), batchSize*size, stream); - return 0; - } + const char *getPluginVersion() const NOEXCEPT override ; + void destroy() NOEXCEPT override ; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 1*sizeof(int); - } + const char *getPluginNamespace() const NOEXCEPT override ; - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer); - tk::dnn::writeBUF(buf, size); - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; - const char *getPluginType() const NOEXCEPT override { - return "ActivationLogisticRT_tkDNN"; - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; - const char *getPluginVersion() const NOEXCEPT override { - return "1"; - } + IPluginV2 *clone() const NOEXCEPT override ; - void destroy() NOEXCEPT override { delete this; } + int size; - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } + private: + std::string mPluginNamespace; + }; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } + class ActivationLogisticRTPluginCreator : public IPluginCreator { + public: + ActivationLogisticRTPluginCreator() ; - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - //todo assert; - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; - IPluginV2 *clone() const NOEXCEPT override{ - ActivationLogisticRT *p = new ActivationLogisticRT(); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; - int size; + const char *getPluginNamespace() const NOEXCEPT override ; -private: - std::string mPluginNamespace; -}; + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; -class ActivationLogisticRTPluginCreator : public IPluginCreator{ -public: - ActivationLogisticRTPluginCreator(){ - mPluginAttributes.clear(); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } + const char *getPluginVersion() const NOEXCEPT override ; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; - IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override { - ActivationLogisticRT *pluginObj = new ActivationLogisticRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } + const char *getPluginName() const NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; - IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override { - ActivationLogisticRT *pluginObj = new ActivationLogisticRT(); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ActivationLogisticRT_tkDNN"; - } - -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ActivationLogisticRTPluginCreator); \ No newline at end of file + REGISTER_TENSORRT_PLUGIN(ActivationLogisticRTPluginCreator); +}; \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationMishRT.h b/include/tkDNN/pluginsRT/ActivationMishRT.h index 65153d0..3c9f15a 100644 --- a/include/tkDNN/pluginsRT/ActivationMishRT.h +++ b/include/tkDNN/pluginsRT/ActivationMishRT.h @@ -1,135 +1,79 @@ #include #include "../kernels.h" +#include +#include -class ActivationMishRT : public IPluginV2 { +namespace nvinfer1 { + class ActivationMishRT : public IPluginV2 { -public: - ActivationMishRT() {} + public: + ActivationMishRT() ; - ~ActivationMishRT() {} + ~ActivationMishRT() ; - ActivationMishRT(const void *data, size_t length) { - std::cout<<"DESERIALIZE MISH"<(data), *bufCheck = buf; - size = readBUF(buf); - assert(buf == bufCheck + length); - } + ActivationMishRT(const void *data, size_t length) ; - int getNbOutputs() const NOEXCEPT override { return 1; } + int getNbOutputs() const NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override { return inputs[0]; } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, - PluginFormat format, int maxBatchSize) NOEXCEPT override { - assert(format == PluginFormat::kLINEAR); - size = 1; - for (int i = 0; i < outputDims[0].nbDims; i++) - size *= outputDims[0].d[i]; - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override ; - int initialize() NOEXCEPT override { return 0; } + int initialize() NOEXCEPT override ; - virtual void terminate() NOEXCEPT override {} + void terminate() NOEXCEPT override ; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0; } + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - virtual int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, - cudaStream_t stream) NOEXCEPT override { - activationMishForward((dnnType *) reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), batchSize * size, stream); - return 0; - } + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ; + size_t getSerializationSize() const NOEXCEPT override ; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 1 * sizeof(int); - } + void serialize(void *buffer) const NOEXCEPT override ; - virtual void serialize(void *buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer), *a = buf; - tk::dnn::writeBUF(buf, size); - assert(buf == a + getSerializationSize()); - } + const char *getPluginType() const NOEXCEPT override ; - const char *getPluginType() const NOEXCEPT override { - return "ActivationMishRT_tkDNN"; - } + const char *getPluginVersion() const NOEXCEPT override ; - const char *getPluginVersion() const NOEXCEPT override { - return "1"; - } + void destroy() NOEXCEPT override { delete this; } - void destroy() NOEXCEPT override { delete this; } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - } + const char *getPluginNamespace() const NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } + void setPluginNamespace(const char *plguinNamespace) NOEXCEPT override ; - void setPluginNamespace(const char *plguinNamespace) NOEXCEPT override { - mPluginNamespace = plguinNamespace; - } + IPluginV2 *clone() const NOEXCEPT override ; - IPluginV2 *clone() const NOEXCEPT override { - ActivationMishRT *p = new ActivationMishRT(); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + int size; + private: + std::string mPluginNamespace; + }; - int size; -private: - std::string mPluginNamespace; -}; + class ActivationMishRTPluginCreator : public IPluginCreator { + public: + ActivationMishRTPluginCreator() ; -class ActivationMishRTPluginCreator : public IPluginCreator { -public: - ActivationMishRTPluginCreator() { - mPluginAttributes.clear(); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + const char *getPluginNamespace() const NOEXCEPT override ; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; - IPluginV2 *deserializePlugin(const char* name,const void* serialData,size_t serialLength) NOEXCEPT override{ - ActivationMishRT *pluginObj = new ActivationMishRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } + const char *getPluginName() const NOEXCEPT override ; - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - ActivationMishRT *pluginObj = new ActivationMishRT(); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } + const char *getPluginVersion() const NOEXCEPT override ; - const char *getPluginName() const NOEXCEPT override{ - return "ActivationMishRT_tkDNN"; - } + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ActivationMishRTPluginCreator); \ No newline at end of file + REGISTER_TENSORRT_PLUGIN(ActivationMishRTPluginCreator); +}; \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h index 8d295df..b5c186c 100644 --- a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h +++ b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h @@ -1,150 +1,78 @@ #include #include "../kernels.h" +#include +#include +#include +namespace nvinfer1 { + class ActivationReLUCeiling : public IPluginV2 { -class ActivationReLUCeiling : public IPluginV2 { + public: + explicit ActivationReLUCeiling(const float ceiling) ; -public: - ActivationReLUCeiling(const float ceiling) { - this->ceiling = ceiling; - } + ~ActivationReLUCeiling() ; - ~ActivationReLUCeiling() { + ActivationReLUCeiling(const void *data, size_t length) ; - } + int getNbOutputs() const NOEXCEPT override ; - ActivationReLUCeiling(const void *data, size_t length) { - std::cout<<"RELU CEILING DESERIALIZE"<(data), *bufCheck = buf; - ceiling = readBUF(buf); - size = readBUF(buf); - assert(buf == bufCheck + length); - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override { - return inputs[0]; - } + int initialize() NOEXCEPT override ; - void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, - PluginFormat format, int maxBatchSize) NOEXCEPT override { - assert(type == DataType::kFLOAT && format == PluginFormat::kLINEAR); - size = 1; - for (int i = 0; i < outputDims[0].nbDims; i++) - size *= outputDims[0].d[i]; - } + void terminate() NOEXCEPT override ; - int initialize() NOEXCEPT override { return 0; } + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - virtual void terminate() NOEXCEPT override {} + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { - return 0; - } + size_t getSerializationSize() const NOEXCEPT override ; - virtual int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, - cudaStream_t stream) NOEXCEPT override { - activationReLUCeilingForward((dnnType *) reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), batchSize * size, ceiling, stream); - return 0; - } + void serialize(void *buffer) const NOEXCEPT override ; + IPluginV2 *clone() const NOEXCEPT override ; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 1 * sizeof(int) + 1 * sizeof(float); - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; - virtual void serialize(void *buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer), *a = buf; - tk::dnn::writeBUF(buf, ceiling); - tk::dnn::writeBUF(buf, size); - assert(buf = a + getSerializationSize()); + void destroy() NOEXCEPT override ; - } + const char *getPluginType() const NOEXCEPT override ; - IPluginV2 *clone() const NOEXCEPT override { - ActivationReLUCeiling *p = new ActivationReLUCeiling(ceiling); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + const char *getPluginVersion() const NOEXCEPT override ; - bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override { - return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR); - } + const char *getPluginNamespace() const NOEXCEPT override ; - void destroy() NOEXCEPT override { delete this; }; + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + int size; + float ceiling; + private: + std::string mPluginNamespace; + }; - const char *getPluginType() const NOEXCEPT override { - return "ActivationReLUCeilingRT_tkDNN"; - } + class ActivationReLUCeilingPluginCreator : public IPluginCreator { + public: + ActivationReLUCeilingPluginCreator() ; - const char *getPluginVersion() const NOEXCEPT override { - return "1"; - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } + const char *getPluginNamespace() const NOEXCEPT override ; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; - int size; - float ceiling; -private: - std::string mPluginNamespace; -}; + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; -class ActivationReLUCeilingPluginCreator : public IPluginCreator { -public: - ActivationReLUCeilingPluginCreator() { - mPluginAttributes.emplace_back(PluginField("ceiling", nullptr, PluginFieldType::kFLOAT32, 1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } + const char *getPluginName() const NOEXCEPT override ; + const char *getPluginVersion() const NOEXCEPT override ; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override { - mPluginNamespace = pluginNamespace; - } + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } + public: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; - IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override { - ActivationReLUCeiling *pluginObj = new ActivationReLUCeiling(serialData, serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - float ceiling = *(static_cast(fields[0].data)); - ActivationReLUCeiling *pluginObj = new ActivationReLUCeiling(ceiling); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ActivationReLUCeilingRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -public: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ActivationReLUCeilingPluginCreator); + REGISTER_TENSORRT_PLUGIN(ActivationReLUCeilingPluginCreator); +}; \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index 7076379..adc5554 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -1,360 +1,115 @@ +#ifndef _DEFORMABLECONVRT_PLUGIN_H +#define _DEFORMABLECONVRT_PLUGIN_H + +#include +#include #include #include "../kernels.h" +#include -class DeformableConvRT : public IPluginV2 { +namespace nvinfer1 { + class DeformableConvRT : public IPluginV2 { + public: + DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw, + int deformableGroup, int i_n, int i_c, int i_h, int i_w, + int o_n, int o_c, int o_h, int o_w, + tk::dnn::DeformConv2d *deformable = nullptr); -public: - DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw, - int deformableGroup, int i_n, int i_c, int i_h, int i_w, - int o_n, int o_c, int o_h, int o_w, - tk::dnn::DeformConv2d *deformable = nullptr) { - this->chunk_dim = chunk_dim; - this->kh = kh; - this->kw = kw; - this->sh = sh; - this->sw = sw; - this->ph = ph; - this->pw = pw; - this->deformableGroup = deformableGroup; - this->i_n = i_n; - this->i_c = i_c; - this->i_h = i_h; - this->i_w = i_w; - this->o_n = o_n; - this->o_c = o_c; - this->o_h = o_h; - this->o_w = o_w; - this->defRT = deformable; + ~DeformableConvRT(); - height_ones = (i_h + 2 * ph - (1 * (kh - 1) + 1)) / sh + 1; - width_ones = (i_w + 2 * pw - (1 * (kw - 1) + 1)) / sw + 1; - dim_ones = i_c * kh * kw * 1 * height_ones * width_ones; - - checkCuda( cudaMalloc(&data_d, i_c * o_c * kh * kw * 1 * sizeof(dnnType))); - checkCuda( cudaMalloc(&bias2_d, o_c*sizeof(dnnType))); - checkCuda( cudaMalloc(&ones_d1, height_ones * width_ones * sizeof(dnnType))); - checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType))); - checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); - checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType))); - if(deformable != nullptr) { - checkCuda( cudaMemcpy(data_d, deformable->data_d, sizeof(dnnType)*i_c * o_c * kh * kw * 1, cudaMemcpyDeviceToDevice) ); - checkCuda( cudaMemcpy(bias2_d, deformable->bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToDevice) ); - checkCuda( cudaMemcpy(ones_d1, deformable->ones_d1, sizeof(dnnType)*height_ones*width_ones, cudaMemcpyDeviceToDevice) ); - checkCuda( cudaMemcpy(offset, deformable->offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToDevice) ); - checkCuda( cudaMemcpy(mask, deformable->mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToDevice) ); - checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) ); - } - stat = cublasCreate(&handle); - if (stat != CUBLAS_STATUS_SUCCESS) - FatalError("CUBLAS initialization failed\n"); - } + DeformableConvRT(const void *data, size_t length) ; - ~DeformableConvRT() { - checkCuda( cudaFree(data_d) ); - checkCuda( cudaFree(bias2_d) ); - checkCuda( cudaFree(ones_d1) ); - checkCuda( cudaFree(offset) ); - checkCuda( cudaFree(mask) ); - checkCuda( cudaFree(ones_d2) ); - cublasDestroy(handle); - } + int getNbOutputs() const NOEXCEPT override ; - DeformableConvRT(const void *data,size_t length){ - const char* buf = reinterpret_cast(data),*bufCheck = buf; - chunk_dim = readBUF(buf); - kh = readBUF(buf); - kw = readBUF(buf); - sh = readBUF(buf); - sw = readBUF(buf); - ph = readBUF(buf); - pw = readBUF(buf); - deformableGroup = readBUF(buf); - i_n = readBUF(buf); - i_c = readBUF(buf); - i_h = readBUF(buf); - i_w = readBUF(buf); - o_n = readBUF(buf); - o_c = readBUF(buf); - o_h = readBUF(buf); - o_w = readBUF(buf); - dnnType *aus = new dnnType[chunk_dim*2]; - for(int i=0;i(buf); - checkCuda(cudaMemcpy(offset,aus,sizeof(dnnType)*2*chunk_dim,cudaMemcpyHostToDevice)); - free(aus); + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - aus = new dnnType[chunk_dim]; - for(int i=0;i(buf); - checkCuda(cudaMemcpy(mask,aus,sizeof(dnnType)*chunk_dim,cudaMemcpyHostToDevice)); - free(aus); + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override ; - aus = new dnnType[i_c*o_c*kh*kw*1]; - for(int i=0;i<(i_c*o_c*kh*kw*1);i++) - aus[i] = readBUF(buf); - checkCuda(cudaMemcpy(data_d,aus,sizeof(dnnType)*(i_c*o_c*kh*kw*1),cudaMemcpyHostToDevice)); - free(aus); + int initialize() NOEXCEPT override ; - aus = new dnnType[o_c]; - for(int i=0; i < o_c; i++) - aus[i] = readBUF(buf); - checkCuda( cudaMemcpy(bias2_d, aus, sizeof(dnnType)*o_c, cudaMemcpyHostToDevice) ); - free(aus); + void terminate() NOEXCEPT override ; - aus = new dnnType[height_ones * width_ones]; - for(int i=0; i(buf); - checkCuda( cudaMemcpy(ones_d1, aus, sizeof(dnnType)*height_ones * width_ones, cudaMemcpyHostToDevice) ); - free(aus); + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - aus = new dnnType[dim_ones]; - for(int i=0; i(buf); - checkCuda( cudaMemcpy(ones_d2, aus, sizeof(dnnType)*dim_ones, cudaMemcpyHostToDevice) ); - free(aus); + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT override; - assert(buf == bufCheck + length); - } + size_t getSerializationSize() const NOEXCEPT override ; - int getNbOutputs() const NOEXCEPT override {return 1;} + void serialize(void *buffer) const NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w}; - } + void destroy() NOEXCEPT override ; - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; - int initialize() NOEXCEPT override {return 0;} + const char *getPluginNamespace() const NOEXCEPT override ; - virtual void terminate() NOEXCEPT override { } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + const char *getPluginType() const NOEXCEPT override ; - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { - dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); - dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]); + const char *getPluginVersion() const NOEXCEPT override ; - // split conv2d outputs into offset to mask - for(int b=0; b(outputs[0]), ones_d2, - kh, kw, - sh, sw, - ph, pw, - 1, 1, - deformableGroup, b, - i_n, i_c, i_h, i_w, - o_n, o_c, o_h, o_w, - chunk_dim); - } - return 0; - } - - virtual size_t getSerializationSize() const NOEXCEPT override { - return 16 * sizeof(int) + chunk_dim * 3 * sizeof(dnnType) + (i_c * o_c * kh * kw * 1 ) * sizeof(dnnType) + - o_c * sizeof(dnnType) + height_ones * width_ones * sizeof(dnnType) + dim_ones * sizeof(dnnType); - } - - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, chunk_dim); - tk::dnn::writeBUF(buf, kh); - tk::dnn::writeBUF(buf, kw); - tk::dnn::writeBUF(buf, sh); - tk::dnn::writeBUF(buf, sw); - tk::dnn::writeBUF(buf, ph); - tk::dnn::writeBUF(buf, pw); - tk::dnn::writeBUF(buf, deformableGroup); - tk::dnn::writeBUF(buf, i_n); - tk::dnn::writeBUF(buf, i_c); - tk::dnn::writeBUF(buf, i_h); - tk::dnn::writeBUF(buf, i_w); - tk::dnn::writeBUF(buf, o_n); - tk::dnn::writeBUF(buf, o_c); - tk::dnn::writeBUF(buf, o_h); - tk::dnn::writeBUF(buf, o_w); - dnnType *aus = new dnnType[chunk_dim*2]; - checkCuda( cudaMemcpy(aus, offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToHost) ); - for(int i=0; isetPluginNamespace(mPluginNamespace.c_str()); - return p; - } + IPluginV2 *clone() const NOEXCEPT override ; - cublasStatus_t stat; - cublasHandle_t handle; - int i_n, i_c, i_h, i_w; - int o_n, o_c, o_h, o_w; - int size; - int chunk_dim; - int kh, kw; - int sh, sw; - int ph, pw; - int deformableGroup; - int height_ones; - int width_ones; - int dim_ones; - - dnnType *data_d; - dnnType *bias2_d; - dnnType *ones_d1; - dnnType * offset; - dnnType * mask; - dnnType *ones_d2; - // dnnType *input_n; - // dnnType *offset_n; - // dnnType *mask_n; - // dnnType *output_n; - - - tk::dnn::DeformConv2d *defRT; + cublasStatus_t stat; + cublasHandle_t handle; + int i_n, i_c, i_h, i_w; + int o_n, o_c, o_h, o_w; + int size; + int chunk_dim; + int kh, kw; + int sh, sw; + int ph, pw; + int deformableGroup; + int height_ones; + int width_ones; + int dim_ones; + + dnnType *data_d; + dnnType *bias2_d; + dnnType *ones_d1; + dnnType *offset; + dnnType *mask; + dnnType *ones_d2; + // dnnType *input_n; + // dnnType *offset_n; + // dnnType *mask_n; + // dnnType *output_n; + + + tk::dnn::DeformConv2d *defRT; + + private: + std::string mPluginNamespace; + }; + + class DeformableConvRTPluginCreator : public IPluginCreator { + public: + DeformableConvRTPluginCreator(); + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + + const char *getPluginNamespace() const NOEXCEPT override ; + + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; + + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; + + const char *getPluginName() const NOEXCEPT override ; + + const char *getPluginVersion() const NOEXCEPT override ; + + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; -private: - std::string mPluginNamespace; }; - -class DeformableConvRTPluginCreator : public IPluginCreator{ -public: - DeformableConvRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("chunk_dim",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("kh",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("kw",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("sh",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("sw",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("ph",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("pw",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("deformableGroup",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_n",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("i_w",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_n",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("defRT",nullptr,PluginFieldType::kUNKNOWN,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override { - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - DeformableConvRT *pluginObj = new DeformableConvRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - int chunk_dim = *(static_cast(fields[0].data)); - int kh = *(static_cast(fields[1].data)); - int kw = *(static_cast(fields[2].data)); - int sh = *(static_cast(fields[3].data)); - int sw = *(static_cast(fields[4].data)); - int ph = *(static_cast(fields[5].data)); - int pw = *(static_cast(fields[6].data)); - int deformableGroup = *(static_cast(fields[7].data)); - int i_n = *(static_cast(fields[8].data)); - int i_c = *(static_cast(fields[9].data)); - int i_h = *(static_cast(fields[10].data)); - int i_w = *(static_cast(fields[11].data)); - int o_n = *(static_cast(fields[12].data)); - int o_c = *(static_cast(fields[13].data)); - int o_h = *(static_cast(fields[14].data)); - int o_w = *(static_cast(fields[14].data)); - DeformConv2d *defRT = const_cast(static_cast(fields[15].data)); - DeformableConvRT *pluginObj = new DeformableConvRT(chunk_dim,kh,kw,sh,sw,ph,pw,deformableGroup,i_n,i_c,i_h,i_w,o_n,o_c,o_h,o_w,defRT); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "DeformableConvRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - - +#endif diff --git a/include/tkDNN/pluginsRT/FlattenConcatRT.h b/include/tkDNN/pluginsRT/FlattenConcatRT.h index 3c6f3ee..7d655d3 100644 --- a/include/tkDNN/pluginsRT/FlattenConcatRT.h +++ b/include/tkDNN/pluginsRT/FlattenConcatRT.h @@ -1,162 +1,81 @@ #include +#include +#include +#include +namespace nvinfer1 { + class FlattenConcatRT : public IPluginV2 { -class FlattenConcatRT : public IPluginV2 { + public: + FlattenConcatRT() ; -public: - FlattenConcatRT() { - stat = cublasCreate(&handle); - if (stat != CUBLAS_STATUS_SUCCESS) { - printf ("CUBLAS initialization failed\n"); - return; - } - } + FlattenConcatRT(const void *data, size_t length) ; - FlattenConcatRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck=buf; - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - rows = readBUF(buf); - cols = readBUF(buf); - assert(buf == bufCheck + length); - } + ~FlattenConcatRT() ; - ~FlattenConcatRT(){ + int getNbOutputs() const NOEXCEPT override ; - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1}; - } + int initialize() NOEXCEPT override ; - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { - assert(nbOutputs == 1 && nbInputs ==1); - rows = inputDims[0].d[0]; - cols = inputDims[0].d[1] * inputDims[0].d[2]; - c = inputDims[0].d[0] * inputDims[0].d[1] * inputDims[0].d[2]; - h = 1; - w = 1; - } + void terminate() NOEXCEPT override ; - int initialize() NOEXCEPT override {return 0;} + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - virtual void terminate() NOEXCEPT override { checkERROR(cublasDestroy(handle));} + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override {return 0;} + size_t getSerializationSize() const NOEXCEPT override ; - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { - dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); - dnnType *dstData = reinterpret_cast(outputs[0]); - checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + void serialize(void *buffer) const NOEXCEPT override ; - checkERROR( cublasSetStream(handle, stream) ); - for(int i=0; i(buffer),*a = buf; - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); - tk::dnn::writeBUF(buf, rows); - tk::dnn::writeBUF(buf, cols); - assert(buf == a + getSerializationSize()); - } + const char *getPluginType() const NOEXCEPT override ; - void destroy() NOEXCEPT override{delete this;} + const char *getPluginVersion() const NOEXCEPT override; - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - } + const char *getPluginNamespace() const NOEXCEPT override ; - const char *getPluginType() const NOEXCEPT override{ - return "FlattenConcatRT_tkDNN"; - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } + IPluginV2 *clone() const NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } + int c, h, w; + int rows, cols; + cublasStatus_t stat; + cublasHandle_t handle; + private: + std::string mPluginNamespace; + }; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } + class FlattenConcatRTPluginCreator : public IPluginCreator { + public: + FlattenConcatRTPluginCreator() ; - IPluginV2 *clone() const NOEXCEPT override { - FlattenConcatRT *p = new FlattenConcatRT(); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; - int c, h, w; - int rows, cols; - cublasStatus_t stat; - cublasHandle_t handle; -private: - std::string mPluginNamespace; -}; + const char *getPluginNamespace() const NOEXCEPT override ; -class FlattenConcatRTPluginCreator : public IPluginCreator{ -public: - FlattenConcatRTPluginCreator(){ - mPluginAttributes.clear(); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } + const char *getPluginName() const NOEXCEPT override ; - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - FlattenConcatRT *pluginObj = new FlattenConcatRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } + const char *getPluginVersion() const NOEXCEPT override; - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - FlattenConcatRT *pluginObj = new FlattenConcatRT(); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; - const char *getPluginName() const NOEXCEPT override{ - return "FlattenConcatRT_tkDNN"; - } + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(FlattenConcatRTPluginCreator); \ No newline at end of file + REGISTER_TENSORRT_PLUGIN(FlattenConcatRTPluginCreator); +}; \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h index ef5a8f7..14f714e 100644 --- a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h +++ b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h @@ -1,182 +1,89 @@ #include #include "../kernels.h" +#include +#include +#include -class MaxPoolFixedSizeRT : public IPluginV2 { +namespace nvinfer1 { + class MaxPoolFixedSizeRT : public IPluginV2 { -public: - MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) { - this->c = c; - this->h = h; - this->w = w; - this->n = n; - this->stride_H = strideH; - this->stride_W = strideW; - this->winSize = winSize; - this->padding = padding; - } + public: + MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) ; - MaxPoolFixedSizeRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck = buf; - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - n = readBUF(buf); - stride_H = readBUF(buf); - stride_W = readBUF(buf); - winSize = readBUF(buf); - padding = readBUF(buf); - assert(buf == bufCheck + length); - } + MaxPoolFixedSizeRT(const void *data, size_t length) ; - ~MaxPoolFixedSizeRT(){ - } + ~MaxPoolFixedSizeRT() ; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + int getNbOutputs() const NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{this->c, this->h, this->w}; - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override ; - int initialize() NOEXCEPT override {return 0;} + int initialize() NOEXCEPT override ; - virtual void terminate() NOEXCEPT override {} + void terminate() NOEXCEPT override ; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { - - //std::cout<n<<" "<c<<" "<h<<" "<w<<" "<stride_H<<" "<stride_W<<" "<winSize<<" "<padding<(inputs[0]); - dnnType *dstData = reinterpret_cast(outputs[0]); - MaxPoolingForward(srcData, dstData, batchSize, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding, stream); - return 0; - } + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT override ; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 8*sizeof(int); - } + size_t getSerializationSize() const NOEXCEPT override ; - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer),*a=buf; + void serialize(void *buffer) const NOEXCEPT override ; - tk::dnn::writeBUF(buf, this->c); - tk::dnn::writeBUF(buf, this->h); - tk::dnn::writeBUF(buf, this->w); - tk::dnn::writeBUF(buf, this->n); - tk::dnn::writeBUF(buf, this->stride_H); - tk::dnn::writeBUF(buf, this->stride_W); - tk::dnn::writeBUF(buf, this->winSize); - tk::dnn::writeBUF(buf, this->padding); - assert(buf == a + getSerializationSize()); - } + void destroy() NOEXCEPT override ; - void destroy() NOEXCEPT override{delete this;} + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - //todo assert - } + const char *getPluginNamespace() const NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - const char *getPluginType() const NOEXCEPT override{ - return "MaxPoolingFixedSizeRT_tkDNN"; - } + const char *getPluginType() const NOEXCEPT override ; - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } + const char *getPluginVersion() const NOEXCEPT override ; - IPluginV2 *clone() const NOEXCEPT override{ - MaxPoolFixedSizeRT *p = new MaxPoolFixedSizeRT(c,h,w,n,stride_H,stride_W,winSize,padding); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + IPluginV2 *clone() const NOEXCEPT override ; - int n, c, h, w; - int stride_H, stride_W; - int winSize; - int padding; + int n, c, h, w; + int stride_H, stride_W; + int winSize; + int padding; -private: - std::string mPluginNamespace; + private: + std::string mPluginNamespace; + }; + + class MaxPoolFixedSizeRTPluginCreator : public IPluginCreator { + public: + MaxPoolFixedSizeRTPluginCreator() ; + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + + const char *getPluginNamespace() const NOEXCEPT override ; + + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; + + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; + + const char *getPluginName() const NOEXCEPT override ; + + const char *getPluginVersion() const NOEXCEPT override ; + + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; + + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + + }; + + REGISTER_TENSORRT_PLUGIN(MaxPoolFixedSizeRTPluginCreator); }; - -class MaxPoolFixedSizeRTPluginCreator : public IPluginCreator{ -public: - MaxPoolFixedSizeRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("w",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("n",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("stride_H",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("stride_W",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("winSize",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("padding",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - MaxPoolFixedSizeRT *pluginObj = new MaxPoolFixedSizeRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - //todo assert - int c = *(static_cast(fields[0].data)); - int h = *(static_cast(fields[1].data)); - int w = *(static_cast(fields[2].data)); - int n = *(static_cast(fields[3].data)); - int stride_H = *(static_cast(fields[4].data)); - int stride_W = *(static_cast(fields[5].data)); - int winSize = *(static_cast(fields[6].data)); - int padding = *(static_cast(fields[7].data)); - MaxPoolFixedSizeRT *pluginObj = new MaxPoolFixedSizeRT(c,h,w,n,stride_H,stride_W,winSize,padding); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "MaxPoolingFixedSizeRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; - -}; - -REGISTER_TENSORRT_PLUGIN(MaxPoolFixedSizeRTPluginCreator); diff --git a/include/tkDNN/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h index e187d6c..cdcb515 100644 --- a/include/tkDNN/pluginsRT/RegionRT.h +++ b/include/tkDNN/pluginsRT/RegionRT.h @@ -1,188 +1,95 @@ +#ifndef _REGIONRT_PLUGIN_H +#define _REGIONRT_PLUGIN_H #include #include "../kernels.h" +#include +#include +#include -class RegionRT : public IPluginV2 { +namespace nvinfer1 { + class RegionRT : public IPluginV2 { -public: - RegionRT(int classes, int coords, int num) { - this->classes = classes; - this->coords = coords; - this->num = num; - } + public: + RegionRT(int classes, int coords, int num); - ~RegionRT(){ + ~RegionRT() ; - } + RegionRT(const void *data, size_t length) ; - RegionRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck=buf; - classes = readBUF(buf); - coords = readBUF(buf); - num = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck+length); - } + int getNbOutputs() const NOEXCEPT override ; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return inputs[0]; - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override ; - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override { - c = inputDims[0].d[0]; - h = inputDims[0].d[1]; - w = inputDims[0].d[2]; - } - - int initialize() NOEXCEPT override { return 0; } + int initialize() NOEXCEPT override ; - virtual void terminate() NOEXCEPT override { } + void terminate() NOEXCEPT override ; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0; } + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { - - dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); - dnnType *dstData = reinterpret_cast(outputs[0]); - - checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - - for (int b = 0; b < batchSize; ++b){ - for(int n = 0; n < num; ++n){ - int index = entry_index(b, n*w*h, 0); - activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); - - index = entry_index(b, n*w*h, coords); - activationLOGISTICForward(srcData + index, dstData + index, w*h, stream); - } - } - - //softmax start - int index = entry_index(0, 0, coords + 1); - softmaxForward( srcData + index, classes, batchSize*num, - (c*h*w)/num, - w*h, 1, w*h, 1, dstData + index, stream); - - return 0; - } + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT override ; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 6*sizeof(int); - } + size_t getSerializationSize() const NOEXCEPT override ; - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, classes); - tk::dnn::writeBUF(buf, coords); - tk::dnn::writeBUF(buf, num); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); - assert(buf == a + getSerializationSize()); - } + void serialize(void *buffer) const NOEXCEPT override ; - const char *getPluginType() const NOEXCEPT override{ - return "RegionRT_tkDNN"; - } + const char *getPluginType() const NOEXCEPT override ; - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } + const char *getPluginVersion() const NOEXCEPT override ; - void destroy() NOEXCEPT override {delete this;} + void destroy() NOEXCEPT override ; - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } + const char *getPluginNamespace() const NOEXCEPT override ; - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; - IPluginV2* clone() const NOEXCEPT override{ - RegionRT *p = new RegionRT(classes,coords,num); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + IPluginV2 *clone() const NOEXCEPT override ; + int c, h, w; + int classes, coords, num; - int c, h, w; - int classes, coords, num; + int entry_index(int batch, int location, int entry) { + int n = location / (w * h); + int loc = location % (w * h); + return batch * c * h * w + n * w * h * (coords + classes + 1) + entry * w * h + loc; + } - int entry_index(int batch, int location, int entry) { - int n = location / (w*h); - int loc = location % (w*h); - return batch*c*h*w + n*w*h*(coords+classes+1) + entry*w*h + loc; - } + private: + std::string mPluginNamespace; + }; -private: - std::string mPluginNamespace; + class RegionRTPluginCreator : public IPluginCreator { + public: + RegionRTPluginCreator(); + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + + const char *getPluginNamespace() const NOEXCEPT override ; + + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; + + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; + + const char *getPluginName() const NOEXCEPT override ; + + const char *getPluginVersion() const NOEXCEPT override ; + + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; + + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; + + REGISTER_TENSORRT_PLUGIN(RegionRTPluginCreator); }; -class RegionRTPluginCreator : public IPluginCreator{ -public: - RegionRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("coords",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - IPluginV2 *deserializePlugin(const char* name,const void *serialData,size_t serialLength) NOEXCEPT override{ - RegionRT *pluginObj = new RegionRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - assert(fc->nbFields == 3); - assert(fields[0].type == PluginFieldType::kINT32); - assert(fields[1].type == PluginFieldType::kINT32); - assert(fields[2].type == PluginFieldType::kINT32); - int classes = *(static_cast(fields[0].data)); - int coords = *(static_cast(fields[1].data)); - int num = *(static_cast(fields[2].data)); - RegionRT *pluginObj = new RegionRT(classes,coords,num); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - - } - - const char *getPluginName() const NOEXCEPT override{ - return "RegionRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(RegionRTPluginCreator); +#endif diff --git a/include/tkDNN/pluginsRT/ReorgRT.h b/include/tkDNN/pluginsRT/ReorgRT.h index 7605fde..a780615 100644 --- a/include/tkDNN/pluginsRT/ReorgRT.h +++ b/include/tkDNN/pluginsRT/ReorgRT.h @@ -1,143 +1,83 @@ #include #include "../kernels.h" +#include +#include -class ReorgRT : public IPluginV2 { +namespace nvinfer1 { + class ReorgRT : public IPluginV2 { -public: - ReorgRT(int stride) { - this->stride = stride; - } + public: + ReorgRT(int stride); - ~ReorgRT(){ + ~ReorgRT(); - } + ReorgRT(const void *data, size_t length); - ReorgRT(const void* data,size_t length){ - const char* buf = reinterpret_cast(data),*bufCheck = buf; - stride = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck + length); - } + int getNbOutputs() const NOEXCEPT override; - int getNbOutputs() const NOEXCEPT override {return 1;} + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride}; - } + void + configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override; - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override { - c = inputDims[0].d[0]; - h = inputDims[0].d[1]; - w = inputDims[0].d[2]; - } + int initialize() NOEXCEPT override; - int initialize() NOEXCEPT override { return 0;} + void terminate() NOEXCEPT override; - virtual void terminate() NOEXCEPT override {} + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} - - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { - - reorgForward((dnnType*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), - batchSize, c, h, w, stride, stream); - return 0; - } + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT override; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 4*sizeof(int); - } + size_t getSerializationSize() const NOEXCEPT override; - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, stride); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); - assert(buf == a + getSerializationSize()); - } - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{return true;} + void serialize(void *buffer) const NOEXCEPT override; - const char *getPluginType() const NOEXCEPT override{ - return "ReorgRT_tkDNN"; - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override; - const char* getPluginVersion() const NOEXCEPT override{ - return "1"; - } - void destroy() NOEXCEPT override{ delete this;} + const char *getPluginType() const NOEXCEPT override; - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } + const char *getPluginVersion() const NOEXCEPT override; - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } + void destroy() NOEXCEPT override; - IPluginV2* clone() const NOEXCEPT override{ - ReorgRT *p = new ReorgRT(stride); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + const char *getPluginNamespace() const NOEXCEPT override; - int c, h, w, stride; -private: - std::string mPluginNamespace; + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; + + IPluginV2 *clone() const NOEXCEPT override; + + int c, h, w, stride; + private: + std::string mPluginNamespace; + }; + + class ReorgRTPluginCreator : public IPluginCreator { + public: + ReorgRTPluginCreator(); + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; + + const char *getPluginNamespace() const NOEXCEPT override; + + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override; + + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override; + + const char *getPluginName() const NOEXCEPT override; + + const char *getPluginVersion() const NOEXCEPT override; + + const PluginFieldCollection *getFieldNames() NOEXCEPT override; + + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; + + REGISTER_TENSORRT_PLUGIN(ReorgRTPluginCreator); }; -class ReorgRTPluginCreator : public IPluginCreator{ -public: - ReorgRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2* deserializePlugin(const char* name,const void* serialData,size_t serialLength) NOEXCEPT override{ - ReorgRT *pluginObj = new ReorgRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection* fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - assert(fc->nbFields == 1); - assert(fields[0].type == PluginFieldType::kINT32); - int stride = *(static_cast(fields[0].data)); - ReorgRT *pluginObj = new ReorgRT(stride); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ReorgRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ReorgRTPluginCreator); - diff --git a/include/tkDNN/pluginsRT/ReshapeRT.h b/include/tkDNN/pluginsRT/ReshapeRT.h index 17037da..8d3fdb5 100644 --- a/include/tkDNN/pluginsRT/ReshapeRT.h +++ b/include/tkDNN/pluginsRT/ReshapeRT.h @@ -1,148 +1,85 @@ +#ifndef _RESHAPERT_PLUGIN_H +#define _RESHAPERT_PLUGIN_H + #include +#include +#include +#include +using namespace tk::dnn; -class ReshapeRT : public IPluginV2 { -public: - ReshapeRT(dataDim_t newDim) { - new_dim = newDim; - n = new_dim.n; - c = new_dim.c; - h = new_dim.h; - w = new_dim.w; - } +namespace nvinfer1 { + class ReshapeRT : public IPluginV2 { - ReshapeRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck = buf; - new_dim.n = readBUF(buf); - new_dim.c = readBUF(buf); - new_dim.h = readBUF(buf); - new_dim.w = readBUF(buf); - assert(buf == bufCheck + length); - } + public: + explicit ReshapeRT(dataDim_t newDim) ; - ~ReshapeRT(){ + ReshapeRT(const void *data, size_t length) ; - } + ~ReshapeRT() ; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + int getNbOutputs() const NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{ c,h,w}; - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - void configureWithFormat (const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override { - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override ; - int initialize() NOEXCEPT override {return 0;} + int initialize() NOEXCEPT override ; - virtual void terminate() NOEXCEPT override {} + void terminate() NOEXCEPT override ; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { - dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); - dnnType *dstData = reinterpret_cast(outputs[0]); + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ; - checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - return 0; - } + size_t getSerializationSize() const NOEXCEPT override ; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 4*sizeof(int); - } + void serialize(void *buffer) const NOEXCEPT override ; - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer),*a = buf; - tk::dnn::writeBUF(buf, n); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); - assert(buf == a + getSerializationSize()); - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - //todo assert - } + const char *getPluginType() const NOEXCEPT override ; - const char *getPluginType() const NOEXCEPT override{ - return "1"; - } + const char *getPluginVersion() const NOEXCEPT override ; - const char *getPluginVersion() const NOEXCEPT override{ - return "ReshapeRT_tkDNN"; - } + void destroy() NOEXCEPT override ; - void destroy() NOEXCEPT override{delete this;} + const char *getPluginNamespace() const NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } + IPluginV2 *clone() const NOEXCEPT override ; + int n, c, h, w; + dataDim_t new_dim; + private: + std::string mPluginNamespace; + }; - IPluginV2 *clone() const NOEXCEPT override{ - ReshapeRT *p = new ReshapeRT(new_dim); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + class ReshapeRTPluginCreator : public IPluginCreator { + public: + ReshapeRTPluginCreator() ; - int n, c, h, w; - dataDim_t new_dim; -private: - std::string mPluginNamespace; + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + + const char *getPluginNamespace() const NOEXCEPT override ; + + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; + + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; + + const char *getPluginName() const NOEXCEPT override ; + + const char *getPluginVersion() const NOEXCEPT override ; + + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; + + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; + + REGISTER_TENSORRT_PLUGIN(ReshapeRTPluginCreator); }; - -class ReshapeRTPluginCreator : public IPluginCreator{ -public: - ReshapeRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("new_dim",nullptr,PluginFieldType::kUNKNOWN,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char* name,const void *serialData,size_t serialLength) NOEXCEPT override{ - ReshapeRT *pluginObj = new ReshapeRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - dataDim_t newDim = *(static_cast(fields[0].data)); - ReshapeRT *pluginObj = new ReshapeRT(newDim); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ReshapeRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ReshapeRTPluginCreator); +#endif \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ResizeLayerRT.h b/include/tkDNN/pluginsRT/ResizeLayerRT.h index 374b7d8..7316048 100644 --- a/include/tkDNN/pluginsRT/ResizeLayerRT.h +++ b/include/tkDNN/pluginsRT/ResizeLayerRT.h @@ -1,165 +1,87 @@ #include #include "../kernels.h" +#include +#include +#include -class ResizeLayerRT : public IPluginV2 { +namespace nvinfer1 { -public: - ResizeLayerRT(int c, int h, int w) { - o_c = c; - o_h = h; - o_w = w; - } + class ResizeLayerRT : public IPluginV2 { - ResizeLayerRT(const void *data,size_t length){ - const char *buf = reinterpret_cast(data),*bufCheck = buf; - o_c = readBUF(buf); - o_h = readBUF(buf); - o_w = readBUF(buf); - i_c = readBUF(buf); - i_h = readBUF(buf); - i_w = readBUF(buf); - assert(buf == bufCheck + length); - } + public: + ResizeLayerRT(int c, int h, int w) ; - ~ResizeLayerRT(){ - } + ResizeLayerRT(const void *data, size_t length) ; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + ~ResizeLayerRT() ; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{o_c, o_h, o_w}; - } + int getNbOutputs() const NOEXCEPT override ; - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { - i_c = inputDims[0].d[0]; - i_h = inputDims[0].d[1]; - i_w = inputDims[0].d[2]; - } - - int initialize() NOEXCEPT override {return 0;} - - virtual void terminate() NOEXCEPT override {} - - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} - - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { - // printf("%d %d %d %d %d %d\n", i_c, i_w, i_h, o_c, o_w, o_h); - resizeForward((dnnType*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), - batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream); - return 0; - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 6*sizeof(int); - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override ; - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer),*a=buf; + int initialize() NOEXCEPT override ; - tk::dnn::writeBUF(buf, o_c); - tk::dnn::writeBUF(buf, o_h); - tk::dnn::writeBUF(buf, o_w); + void terminate() NOEXCEPT override ; - tk::dnn::writeBUF(buf, i_c); - tk::dnn::writeBUF(buf, i_h); - tk::dnn::writeBUF(buf, i_w); - assert(buf == a + getSerializationSize()); - } + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ; - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - //todo assert - } + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT override ; - const char *getPluginType() const NOEXCEPT override{ - return "ResizeLayerRT_tkDNN"; - } - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - void destroy() NOEXCEPT override{delete this;} + size_t getSerializationSize() const NOEXCEPT override ; - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } + void serialize(void *buffer) const NOEXCEPT override ; - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - IPluginV2 *clone() const NOEXCEPT override{ - ResizeLayerRT *p = new ResizeLayerRT(o_c,o_h,o_w); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ; - int i_c, i_h, i_w, o_c, o_h, o_w; + const char *getPluginType() const NOEXCEPT override ; -private: - std::string mPluginNamespace; + const char *getPluginVersion() const NOEXCEPT override ; + + void destroy() NOEXCEPT override ; + + const char *getPluginNamespace() const NOEXCEPT override ; + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + + IPluginV2 *clone() const NOEXCEPT override ; + + int i_c, i_h, i_w, o_c, o_h, o_w; + + private: + std::string mPluginNamespace; + }; + + class ResizeLayerRTPluginCreator : public IPluginCreator { + public: + ResizeLayerRTPluginCreator() ; + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + + const char *getPluginNamespace() const NOEXCEPT override ; + + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; + + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; + + const char *getPluginName() const NOEXCEPT override ; + + const char *getPluginVersion() const NOEXCEPT override ; + + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; + + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + + }; + + REGISTER_TENSORRT_PLUGIN(ResizeLayerRTPluginCreator); }; -class ResizeLayerRTPluginCreator : public IPluginCreator{ -public: - ResizeLayerRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - ResizeLayerRT *pluginObj = new ResizeLayerRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - assert(fc->nbFields == 3); - assert(fields[0].type == PluginFieldType::kINT32); - assert(fields[1].type == PluginFieldType::kINT32); - assert(fields[2].type == PluginFieldType::kINT32); - int oc = *(static_cast(fields[0].data)); - int oh = *(static_cast(fields[1].data)); - int ow = *(static_cast(fields[2].data)); - ResizeLayerRT *pluginObj = new ResizeLayerRT(oc,oh,ow); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ResizeLayerRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - - -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; - -}; - -REGISTER_TENSORRT_PLUGIN(ResizeLayerRTPluginCreator); - diff --git a/include/tkDNN/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h index 2ceede3..2b9bc67 100644 --- a/include/tkDNN/pluginsRT/RouteRT.h +++ b/include/tkDNN/pluginsRT/RouteRT.h @@ -1,181 +1,86 @@ #include #include "../kernels.h" +#include +#include -class RouteRT : public IPluginV2 { +namespace nvinfer1 { + class RouteRT : public IPluginV2 { - /** - THIS IS NOT USED ANYMORE - */ + /** + THIS IS NOT USED ANYMORE + */ -public: - RouteRT(int groups, int group_id) { - this->groups = groups; - this->group_id = group_id; - } + public: + RouteRT(int groups, int group_id) ; - ~RouteRT(){ + ~RouteRT() ; - } + RouteRT(const void *data, size_t length) ; - RouteRT(const void* data,size_t length){ - const char* buf = reinterpret_cast(data),*bufCheck = buf; - groups = readBUF(buf); - group_id = readBUF(buf); - in = readBUF(buf); - for(int i=0;i (buf); - } - c= readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck + length); - } + int getNbOutputs() const NOEXCEPT override ; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - int out_c = 0; - for(int i=0; i(outputs[0]); - for(int b=0; b(inputs[i]); - int in_dim = c_in[i]*h*w; - int part_in_dim = in_dim / this->groups; - checkCuda( cudaMemcpyAsync(dstData + b*c*w*h + offset, input + b*c*w*h*groups + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) ); - offset += part_in_dim; - } - } - return 0; - } + size_t getSerializationSize() const NOEXCEPT override ; - virtual size_t getSerializationSize() const NOEXCEPT override { - return (6+MAX_INPUTS)*sizeof(int); - } + void serialize(void *buffer) const NOEXCEPT override ; - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, groups); - tk::dnn::writeBUF(buf, group_id); - tk::dnn::writeBUF(buf, in); - for(int i=0; isetPluginNamespace(mPluginNamespace.c_str()); - return p; - } - static const int MAX_INPUTS = 4; - int in; - int c_in[MAX_INPUTS]; - int c, h, w; - int groups, group_id; -private: - std::string mPluginNamespace; + class RouteRTPluginCreator : public IPluginCreator { + public: + RouteRTPluginCreator() ; + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ; + + const char *getPluginNamespace() const NOEXCEPT override ; + + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ; + + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ; + + const char *getPluginName() const NOEXCEPT override ; + + const char *getPluginVersion() const NOEXCEPT override ; + + const PluginFieldCollection *getFieldNames() NOEXCEPT override ; + + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; + + REGISTER_TENSORRT_PLUGIN(RouteRTPluginCreator); }; - -class RouteRTPluginCreator : public IPluginCreator{ -public: - RouteRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("groups",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("group_id",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char* name,const void* serialData,size_t serialLength) NOEXCEPT override{ - RouteRT *pluginObj = new RouteRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - assert(fc->nbFields == 2); - assert(fields[0].type == PluginFieldType::kINT32); - assert(fields[1].type == PluginFieldType::kINT32); - int groups = *(static_cast(fields[0].data)); - int group_id = *(static_cast(fields[1].data)); - RouteRT *pluginObj = new RouteRT(groups,group_id); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "RouteRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(RouteRTPluginCreator); diff --git a/include/tkDNN/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h index 9f3d343..7069180 100644 --- a/include/tkDNN/pluginsRT/ShortcutRT.h +++ b/include/tkDNN/pluginsRT/ShortcutRT.h @@ -1,166 +1,94 @@ +#ifndef _SHORTCUTRT_PLUGIN_H +#define _SHORTCUTRT_PLUGIN_H + #include #include "../kernels.h" +#include +#include +#include -class ShortcutRT : public IPluginV2 { +namespace nvinfer1 { -public: - ShortcutRT(tk::dnn::dataDim_t bdim, bool mul) { - bDim = bdim; - this->bc = bDim.c; - this->bh = bDim.h; - this->bw = bDim.w; - this->mul = mul; - } + class ShortcutRT : public IPluginV2 { - ~ShortcutRT(){} + public: + ShortcutRT(tk::dnn::dataDim_t bdim, bool mul); - ShortcutRT(const void* data,size_t length){ - const char* buf =reinterpret_cast(data),*bufCheck = buf; - bDim.c = readBUF(buf); - bDim.h = readBUF(buf); - bDim.w = readBUF(buf); - bDim.l = 1; - mul = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck + length); - } + ~ShortcutRT(); - int getNbOutputs() const NOEXCEPT override {return 1;} + ShortcutRT(const void *data, size_t length); - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]}; - } + int getNbOutputs() const NOEXCEPT override; - void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { - c = inputDims[0].d[0]; - h = inputDims[0].d[1]; - w = inputDims[0].d[2]; - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override; - int initialize() NOEXCEPT override {return 0;} + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override; - virtual void terminate() NOEXCEPT override {} + int initialize() NOEXCEPT override; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + void terminate() NOEXCEPT override; - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; - dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); - dnnType *srcDataBack = (dnnType*)reinterpret_cast(inputs[1]); - dnnType *dstData = reinterpret_cast(outputs[0]); - - checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, mul, stream); - - return 0; - } + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT override; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 6*sizeof(int) + sizeof(bool); - } + size_t getSerializationSize() const NOEXCEPT override; - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, bc); - tk::dnn::writeBUF(buf, bh); - tk::dnn::writeBUF(buf, bw); - tk::dnn::writeBUF(buf, mul); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); - assert(buf == a + getSerializationSize()); - - } + void serialize(void *buffer) const NOEXCEPT override; - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - return true; - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override; - const char* getPluginType() const NOEXCEPT override{ - return "1"; - } + const char *getPluginType() const NOEXCEPT override; - const char* getPluginVersion() const NOEXCEPT override{ - return "ShortcutRT_tkDNN"; - } + const char *getPluginVersion() const NOEXCEPT override; - void destroy() NOEXCEPT override{delete this;} + void destroy() NOEXCEPT override; - const char* getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } + const char *getPluginNamespace() const NOEXCEPT override; - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; - IPluginV2 *clone() const NOEXCEPT override{ - ShortcutRT *p = new ShortcutRT(bDim,mul); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } + IPluginV2 *clone() const NOEXCEPT override; + + int c, h, w; + int bc, bh, bw; + bool mul; + tk::dnn::dataDim_t bDim; + private: + std::string mPluginNamespace; + }; + + + class ShortcutRTPluginCreator : public IPluginCreator { + public: + ShortcutRTPluginCreator(); + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; + + const char *getPluginNamespace() const NOEXCEPT override; + + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override; + + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override; + + const char *getPluginName() const NOEXCEPT override; + + const char *getPluginVersion() const NOEXCEPT override; + + const PluginFieldCollection *getFieldNames() NOEXCEPT override; + + public: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; + + REGISTER_TENSORRT_PLUGIN(ShortcutRTPluginCreator); - int c, h, w; - int bc, bh, bw; - bool mul; - tk::dnn::dataDim_t bDim; -private: - std::string mPluginNamespace; }; - -class ShortcutRTPluginCreator : public IPluginCreator { -public: - ShortcutRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("bDim",nullptr,PluginFieldType::kUNKNOWN,1)); - mPluginAttributes.emplace_back(PluginField("mul",nullptr,PluginFieldType::kUNKNOWN,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - ShortcutRT *pluginObj = new ShortcutRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char *name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - //todo assert - tk::dnn::dataDim_t bdim = *(static_cast(fields[0].data)); - bool mul = *(static_cast(fields[1].data)); - ShortcutRT *pluginObj = new ShortcutRT(bdim,mul); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "ShortcutRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } -public: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(ShortcutRTPluginCreator); \ No newline at end of file +#endif \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/UpsampleRT.h b/include/tkDNN/pluginsRT/UpsampleRT.h index 9d7a62c..030fe5a 100644 --- a/include/tkDNN/pluginsRT/UpsampleRT.h +++ b/include/tkDNN/pluginsRT/UpsampleRT.h @@ -1,150 +1,87 @@ +#ifndef _UPSAMPLERT_PLUGIN_H +#define _UPSAMPLERT_PLUGIN_H + #include #include "../kernels.h" +#include +#include + +namespace nvinfer1 { + + class UpsampleRT : public IPluginV2 { + + public: + explicit UpsampleRT(int stride); + + UpsampleRT(const void *data, size_t length); + + ~UpsampleRT(); + + int getNbOutputs() const NOEXCEPT override; + + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override; + + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override; + + int initialize() NOEXCEPT override; + + void terminate() NOEXCEPT override; + + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; + + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT override; -class UpsampleRT : public IPluginV2 { + size_t getSerializationSize() const NOEXCEPT override; -public: - UpsampleRT(int stride) { - this->stride = stride; - } + void serialize(void *buffer) const NOEXCEPT override; - UpsampleRT(const void *data,size_t length){ - const char* buf = reinterpret_cast(data),*bufCheck=buf; - stride = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - assert(buf == bufCheck + length); - } + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override; + const char *getPluginType() const NOEXCEPT override; - ~UpsampleRT(){} + const char *getPluginVersion() const NOEXCEPT override; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + void destroy() NOEXCEPT override; - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) NOEXCEPT override { - return Dims3(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride); - } + const char *getPluginNamespace() const NOEXCEPT override; - void configureWithFormat (const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs,DataType type,PluginFormat format,int maxBatchSize) NOEXCEPT override { - c = inputDims[0].d[0]; - h = inputDims[0].d[1]; - w = inputDims[0].d[2]; - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; - int initialize() NOEXCEPT override {return 0;} + IPluginV2 *clone() const NOEXCEPT override ; - virtual void terminate() NOEXCEPT override {} + int c, h, w, stride; + private: + std::string mPluginNamespace; + }; - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { return 0;} + class UpsampleRTPluginCreator : public IPluginCreator { + public: + UpsampleRTPluginCreator(); - virtual int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) NOEXCEPT override { + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; - dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); - dnnType *dstData = reinterpret_cast(outputs[0]); - - fill(dstData, batchSize*c*h*w*stride*stride, 0.0, stream); - upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream); - return 0; - } + const char *getPluginNamespace() const NOEXCEPT override; + IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override; - virtual size_t getSerializationSize() const NOEXCEPT override { return 4*sizeof(int);} + IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override; - virtual void serialize(void* buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, stride); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); - assert(buf == a + getSerializationSize()); - } + const char *getPluginName() const NOEXCEPT override; - bool supportsFormat(DataType type,PluginFormat format) const NOEXCEPT override{ - //todo assert - return true; - } + const char *getPluginVersion() const NOEXCEPT override; - const char *getPluginType() const NOEXCEPT override{ - return "1"; - } + const PluginFieldCollection *getFieldNames() NOEXCEPT override; - const char *getPluginVersion() const NOEXCEPT override{ - static const char* UPSAMPLE_RT_PLUGIN = "UpsampleRT_TRT"; - return UPSAMPLE_RT_PLUGIN; - } + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; - void destroy() NOEXCEPT override{delete this;} - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - IPluginV2* clone() const NOEXCEPT override{ - UpsampleRT *p = new UpsampleRT(stride); - p->setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - - int c, h, w, stride; -private: - std::string mPluginNamespace; -}; - -class UpsampleRTPluginCreator : public IPluginCreator{ -public: - UpsampleRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char* pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char* name,const void* serialData,size_t serialLength) NOEXCEPT override{ - UpsampleRT *pluginObj = new UpsampleRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - int stride = *(static_cast(fields[0].data)); - UpsampleRT *pluginObj = new UpsampleRT(stride); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - static const char* UPSAMPLE_RT_PLUGIN = "UpsampleRT_TRT"; - return UPSAMPLE_RT_PLUGIN; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(UpsampleRTPluginCreator); + REGISTER_TENSORRT_PLUGIN(UpsampleRTPluginCreator); + }; +#endif \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index 7ae7b35..0d895e0 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -1,282 +1,109 @@ +#ifndef _YOLORT_PLUGIN_H +#define _YOLORT_PLUGIN_H + #include #include #include "../kernels.h" +#include +#include + #define YOLORT_CLASSNAME_W 256 +namespace nvinfer1 { + class YoloRT : public IPluginV2 { -class YoloRT : public IPluginV2 { + public: + YoloRT(int classes, int num, tk::dnn::Yolo *Yolo = nullptr, int n_masks = 3, float scale_xy = 1, + float nms_thresh = 0.45, int nms_kind = 0, int new_coords = 0); -public: - YoloRT(int classes, int num, tk::dnn::Yolo *Yolo = nullptr, int n_masks = 3, float scale_xy = 1, - float nms_thresh = 0.45, int nms_kind = 0, int new_coords = 0) { - this->yolo = Yolo; - this->classes = classes; - this->num = num; - this->n_masks = n_masks; - this->scaleXY = scale_xy; - this->nms_thresh = nms_thresh; - this->nms_kind = nms_kind; - this->new_coords = new_coords; + YoloRT(const void *data, size_t length); - mask = new dnnType[n_masks]; - bias = new dnnType[num * n_masks * 2]; - if (yolo != nullptr) { - memcpy(mask, yolo->mask_h, sizeof(dnnType) * n_masks); - memcpy(bias, yolo->bias_h, sizeof(dnnType) * num * n_masks * 2); - classesNames = yolo->classesNames; - } - } - - YoloRT(const void *data,size_t length){ - std::vector maskTemp,biasTemp; - std::cout<<"LENGTH : "<(data),*bufCheck = buf; - classes = readBUF(buf); - num = readBUF(buf); - n_masks = readBUF(buf); - scaleXY = readBUF(buf); - nms_thresh = readBUF(buf); - nms_kind = readBUF(buf); - new_coords = readBUF(buf); - c = readBUF(buf); - h = readBUF(buf); - w = readBUF(buf); - for(int i=0;i(buf)); - std::cout<(buf)); - std::cout<(buf); - classesNames[1] = std::string(tmp); - } - assert(buf == bufCheck + length); - - } - - ~YoloRT() { - - } + ~YoloRT(); + int getNbOutputs() const NOEXCEPT override; - int getNbOutputs() const NOEXCEPT override { - return 1; - } + Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override; - Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override { - return inputs[0]; - } + void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT override; - void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, - PluginFormat format, int maxBatchSize) NOEXCEPT override { - c = inputDims[0].d[0]; - h = inputDims[0].d[1]; - w = inputDims[0].d[2]; - } + int initialize() NOEXCEPT override; - int initialize() NOEXCEPT override { + void terminate() NOEXCEPT override; - return 0; - } + size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override; - virtual void terminate() NOEXCEPT override { - } - - virtual size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override { - return 0; - } - - virtual int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, - cudaStream_t stream) NOEXCEPT override { - - dnnType *srcData = (dnnType *) reinterpret_cast(inputs[0]); - dnnType *dstData = reinterpret_cast(outputs[0]); - - checkCuda(cudaMemcpyAsync(dstData, srcData, batchSize * c * h * w * sizeof(dnnType), cudaMemcpyDeviceToDevice, - stream)); + int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT override; - for (int b = 0; b < batchSize; ++b) { - for (int n = 0; n < n_masks; ++n) { - int index = entry_index(b, n * w * h, 0); - if (new_coords == 1) { - if (this->scaleXY != 1) - scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1); - } else { - activationLOGISTICForward(srcData + index, dstData + index, 2 * w * h, stream); //x,y + size_t getSerializationSize() const NOEXCEPT override; - if (this->scaleXY != 1) - scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1); + bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override; - index = entry_index(b, n * w * h, 4); - activationLOGISTICForward(srcData + index, dstData + index, (1 + classes) * w * h, stream); - } - } + void serialize(void *buffer) const NOEXCEPT override; + + const char *getPluginType() const NOEXCEPT override; + + const char *getPluginVersion() const NOEXCEPT override; + + void destroy() NOEXCEPT override; + + const char *getPluginNamespace() const NOEXCEPT override; + + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; + + IPluginV2 *clone() const NOEXCEPT override; + + tk::dnn::Yolo *yolo; + int c, h, w; + int classes, num, n_masks; + float scaleXY; + float nms_thresh; + int nms_kind; + int new_coords; + int NUM = 0; + std::vector classesNames; + + dnnType *mask; + dnnType *bias; + + int entry_index(int batch, int location, int entry) { + int n = location / (w * h); + int loc = location % (w * h); + return batch * c * h * w + n * w * h * (4 + classes + 1) + entry * w * h + loc; } - //std::cout<<"YOLO END\n"; - return 0; - } + private: + std::string mPluginNamespace; + }; - virtual size_t getSerializationSize() const NOEXCEPT override { - return 8 * sizeof(int) + 2 * sizeof(float) + n_masks * sizeof(dnnType) + num * n_masks * 2 * sizeof(dnnType) + - YOLORT_CLASSNAME_W * classes * sizeof(char); - } + class YoloRTPluginCreator : public IPluginCreator { + public: + YoloRTPluginCreator(); - bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override { - return true; //todo implement proper supportsFormat - } + void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override; - virtual void serialize(void *buffer) const NOEXCEPT override { - char *buf = reinterpret_cast(buffer), *a = buf; - tk::dnn::writeBUF(buf, classes); //std::cout << "Classes :" << classes << std::endl; - tk::dnn::writeBUF(buf, num); //std::cout << "Num : " << num << std::endl; - std::cout<setPluginNamespace(mPluginNamespace.c_str()); - return p; - } - - Yolo *yolo; - int c, h, w; - int classes, num, n_masks; - float scaleXY; - float nms_thresh; - int nms_kind; - int new_coords; - int NUM=0; - std::vector classesNames; - - dnnType *mask; - dnnType *bias; - - int entry_index(int batch, int location, int entry) { - int n = location / (w * h); - int loc = location % (w * h); - return batch * c * h * w + n * w * h * (4 + classes + 1) + entry * w * h + loc; - } - -private: - std::string mPluginNamespace; + private: + static PluginFieldCollection mFC; + static std::vector mPluginAttributes; + std::string mPluginNamespace; + }; + REGISTER_TENSORRT_PLUGIN(YoloRTPluginCreator); }; - -class YoloRTPluginCreator : public IPluginCreator{ -public: - YoloRTPluginCreator(){ - mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("yolo",nullptr,PluginFieldType::kUNKNOWN,1)); - mPluginAttributes.emplace_back(PluginField("numMasks",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("scaleXY",nullptr,PluginFieldType::kFLOAT32,1)); - mPluginAttributes.emplace_back(PluginField("nmsThresh",nullptr,PluginFieldType::kFLOAT32,1)); - mPluginAttributes.emplace_back(PluginField("nmsKind",nullptr,PluginFieldType::kINT32,1)); - mPluginAttributes.emplace_back(PluginField("newCoords",nullptr,PluginFieldType::kINT32,1)); - mFC.nbFields = mPluginAttributes.size(); - mFC.fields = mPluginAttributes.data(); - } - - void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override{ - mPluginNamespace = pluginNamespace; - } - - const char *getPluginNamespace() const NOEXCEPT override{ - return mPluginNamespace.c_str(); - } - - IPluginV2 *deserializePlugin(const char *name,const void *serialData,size_t serialLength) NOEXCEPT override{ - YoloRT *pluginObj = new YoloRT(serialData,serialLength); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - IPluginV2 *createPlugin(const char* name,const PluginFieldCollection *fc) NOEXCEPT override{ - const PluginField *fields = fc->fields; - //todo assert - int classes = *(static_cast(fields[0].data)); - int num = *(static_cast(fields[1].data)); - Yolo *yoloTemp = const_cast(static_cast(fields[2].data)); - int numMasks = *(static_cast(fields[3].data)); - float scaleXY = *(static_cast(fields[4].data)); - float nmsThresh = *(static_cast(fields[5].data)); - int nmsKind = *(static_cast(fields[6].data)); - int newCoords = *(static_cast(fields[7].data)); - YoloRT *pluginObj = new YoloRT(classes,num,yoloTemp,numMasks,scaleXY,nmsThresh,nmsKind,newCoords); - pluginObj->setPluginNamespace(mPluginNamespace.c_str()); - return pluginObj; - } - - const char *getPluginName() const NOEXCEPT override{ - return "YoloRT_tkDNN"; - } - - const char *getPluginVersion() const NOEXCEPT override{ - return "1"; - } - - const PluginFieldCollection *getFieldNames() NOEXCEPT override{ - return &mFC; - } - -private: - PluginFieldCollection mFC; - std::vector mPluginAttributes; - std::string mPluginNamespace; -}; - -REGISTER_TENSORRT_PLUGIN(YoloRTPluginCreator); +#endif \ No newline at end of file diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index 017cf1d..1219ec3 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -34,6 +34,19 @@ #define dnnType float +template void writeBUF(char*& buffer, const T& val) +{ + *reinterpret_cast(buffer) = val; + buffer += sizeof(T); +} + +template T readBUF(const char*& buffer) +{ + T val = *reinterpret_cast(buffer); + buffer += sizeof(T); + return val; +} + // Colored output #define COL_END "\033[0m" diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp index a370b91..3333afd 100644 --- a/src/DarknetParser.cpp +++ b/src/DarknetParser.cpp @@ -17,8 +17,8 @@ namespace tk { namespace dnn { if(sep == std::string::npos) return false; - name = line.substr(0, sep); - value = line.substr(sep+1, line.size() - (sep+1)); + name = line.substr(0, sep); + value = line.substr(sep+1, line.size() - (sep+1)); return true; } @@ -322,7 +322,7 @@ namespace tk { namespace dnn { // skip empty lines if(line.empty()) continue; - if(count > lineNo && count <=20){ + if(count > lineNo && count <=lineNo+30){ divideNameAndValue(line,name,value); if(name == "mask "){ maskTemp = fromStringToFloatVec(value,','); @@ -353,9 +353,8 @@ namespace tk { namespace dnn { if(name == "beta_nms"){ nmsThreshTemp = std::stof(value); } - - count++; } + count++; } mask = maskTemp; anchors = anchorsTemp; diff --git a/src/Yolo.cpp b/src/Yolo.cpp index deffbeb..35abb19 100644 --- a/src/Yolo.cpp +++ b/src/Yolo.cpp @@ -133,10 +133,11 @@ void correct_yolo_boxes(Yolo::detection *dets, int n, int w, int h, int netw, in } } -int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh, int new_coords) { +int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh, int newCoords) { if(predictions == nullptr) predictions = new dnnType[output_dim.tot()]; + checkCuda(cudaDeviceSynchronize()); checkCuda( cudaMemcpy(predictions, dstData, output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost)); int lw = output_dim.w; @@ -157,7 +158,7 @@ int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int net if(objectness <= thresh) continue; int box_index = entry_index(0, n*lw*lh + i, 0, classes, input_dim, output_dim); - dets[count].bbox = get_yolo_box(predictions, bias_h, mask_h[n], box_index, col, row, lw, lh, netw, neth, lw*lh, new_coords); + dets[count].bbox = get_yolo_box(predictions, bias_h, mask_h[n], box_index, col, row, lw, lh, netw, neth, lw*lh, newCoords); dets[count].objectness = objectness; dets[count].classes = classes; for(j = 0; j < classes; ++j){ diff --git a/src/pluginsRT/ActivationLeakyRT.cpp b/src/pluginsRT/ActivationLeakyRT.cpp new file mode 100644 index 0000000..77063f6 --- /dev/null +++ b/src/pluginsRT/ActivationLeakyRT.cpp @@ -0,0 +1,138 @@ +#include +using namespace nvinfer1; + + +std::vector ActivationLeakyRTPluginCreator::mPluginAttributes; +PluginFieldCollection ActivationLeakyRTPluginCreator::mFC{}; + +ActivationLeakyRT::ActivationLeakyRT(float s) { + slope = s; +} + +ActivationLeakyRT::ActivationLeakyRT(const void *data, size_t length) { + std::cout << "DESERIALIZE LEAKYRT" << std::endl; + const char *buf = reinterpret_cast(data), *bufCheck = buf; + slope = readBUF(buf); + size = readBUF(buf); + assert(buf == bufCheck + length); +} + +ActivationLeakyRT::~ActivationLeakyRT() {} + +int ActivationLeakyRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims ActivationLeakyRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return inputs[0]; +} + +void ActivationLeakyRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, + DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT { + assert(type == DataType::kFLOAT && format == PluginFormat::kLINEAR); + size = 1; + for (int i = 0; i < outputDims[0].nbDims; i++) + size *= outputDims[0].d[i]; +} +int ActivationLeakyRT::initialize() NOEXCEPT { + return 0; +} + +size_t ActivationLeakyRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +int ActivationLeakyRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + activationLEAKYForward( + (dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, slope, + stream); + return 0; + +} + +size_t ActivationLeakyRT::getSerializationSize() const NOEXCEPT { + return 1 * sizeof(int) + 1 * sizeof(float); +} + +void ActivationLeakyRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer), *a = buf; + writeBUF(buf, size); + assert(buf == a + getSerializationSize()); +} + +bool ActivationLeakyRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR); +} + +const char *ActivationLeakyRT::getPluginType() const NOEXCEPT { + return "ActivationLeakyRT_tkDNN"; +} + +const char *ActivationLeakyRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +void ActivationLeakyRT::destroy() NOEXCEPT { + delete this; +} + +const char *ActivationLeakyRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); + +} + +void ActivationLeakyRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2* ActivationLeakyRT::clone() const NOEXCEPT { + auto *p = new ActivationLeakyRT(slope); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + +ActivationLeakyRTPluginCreator::ActivationLeakyRTPluginCreator() { + mPluginAttributes.emplace_back( + PluginField("slope", nullptr, PluginFieldType::kFLOAT32, 1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ActivationLeakyRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2* ActivationLeakyRTPluginCreator::deserializePlugin(const char *name, const void *serialData,size_t serialLength) NOEXCEPT { + auto *pluginObj = new ActivationLeakyRT(serialData, serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char* ActivationLeakyRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2* ActivationLeakyRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + assert(fc->nbFields == 1); + assert(fields[0].type == PluginFieldType::kFLOAT32); + float slope = *(static_cast(fields[0].data)); + auto *pluginObj = new ActivationLeakyRT(slope); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char* ActivationLeakyRTPluginCreator::getPluginName() const NOEXCEPT { + return "ActivationLeakyRT_tkDNN"; +} + +const char* ActivationLeakyRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection* ActivationLeakyRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + diff --git a/src/pluginsRT/ActivationLogisticRT.cpp b/src/pluginsRT/ActivationLogisticRT.cpp new file mode 100644 index 0000000..9966708 --- /dev/null +++ b/src/pluginsRT/ActivationLogisticRT.cpp @@ -0,0 +1,126 @@ +#include +using namespace nvinfer1; +std::vector ActivationLogisticRTPluginCreator::mPluginAttributes; +PluginFieldCollection ActivationLogisticRTPluginCreator::mFC{}; + +ActivationLogisticRT::ActivationLogisticRT() {} + +ActivationLogisticRT::ActivationLogisticRT(const void *data, size_t length) { + const char *buf = reinterpret_cast(data), *bufCheck = buf; + size = readBUF(buf); + assert(buf == bufCheck + length); +} + +ActivationLogisticRT::~ActivationLogisticRT() {} + +int ActivationLogisticRT::getNbOutputs() const NOEXCEPT { +return 1; +} + +Dims ActivationLogisticRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return inputs[0]; +} + +void ActivationLogisticRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, + int nbOutputs, DataType type, PluginFormat format, + int maxBatchSize) NOEXCEPT { + size = 1; + for (int i = 0; i < outputDims[0].nbDims; i++) + size *= outputDims[0].d[i]; +} + +int ActivationLogisticRT::initialize() NOEXCEPT { + return 0; +} + +void ActivationLogisticRT::terminate() NOEXCEPT {} + +size_t ActivationLogisticRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +int ActivationLogisticRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + activationLOGISTICForward((dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, stream); + return 0; +} + +size_t ActivationLogisticRT::getSerializationSize() const NOEXCEPT { + return 1 * sizeof(int); +} + +void ActivationLogisticRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer); + writeBUF(buf, size); +} + +const char* ActivationLogisticRT::getPluginType() const NOEXCEPT { + return "ActivationLogisticRT_tkDNN"; +} + +const char* ActivationLogisticRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +void ActivationLogisticRT::destroy() NOEXCEPT { + delete this; +} + +const char* ActivationLogisticRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ActivationLogisticRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +bool ActivationLogisticRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; + //todo assert +} + +IPluginV2* ActivationLogisticRT::clone() const NOEXCEPT { + auto *p = new ActivationLogisticRT(); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + +ActivationLogisticRTPluginCreator::ActivationLogisticRTPluginCreator() { + mPluginAttributes.clear(); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ActivationLogisticRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2* ActivationLogisticRTPluginCreator::deserializePlugin(const char *name, const void *serialData, + size_t serialLength) NOEXCEPT { + auto *pluginObj = new ActivationLogisticRT(serialData, serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char* ActivationLogisticRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2* ActivationLogisticRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + auto *pluginObj = new ActivationLogisticRT(); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char* ActivationLogisticRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection* ActivationLogisticRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + +const char *ActivationLogisticRTPluginCreator::getPluginName() const NOEXCEPT { + return "ActivationLogisticRT_tkDNN"; +} \ No newline at end of file diff --git a/src/pluginsRT/ActivationMishRT.cpp b/src/pluginsRT/ActivationMishRT.cpp new file mode 100644 index 0000000..7797df9 --- /dev/null +++ b/src/pluginsRT/ActivationMishRT.cpp @@ -0,0 +1,121 @@ +// +// Created by perseusdg on 9/4/21. +// +#include +using namespace nvinfer1; +std::vector ActivationMishRTPluginCreator::mPluginAttributes; +PluginFieldCollection ActivationMishRTPluginCreator::mFC{}; + +ActivationMishRT::ActivationMishRT() { + +} + +ActivationMishRT::~ActivationMishRT() { + +} + +ActivationMishRT::ActivationMishRT(const void *data, size_t length) { + const char *buf = reinterpret_cast(data), *bufCheck = buf; + size = readBUF(buf); + assert(buf == bufCheck + length); +} + +int ActivationMishRT::getNbOutputs() const NOEXCEPT { return 1; } + +Dims ActivationMishRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { return inputs[0]; } + +void ActivationMishRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT { +assert(format == PluginFormat::kLINEAR); +size = 1; +for (int i = 0; i < outputDims[0].nbDims; i++) +size *= outputDims[0].d[i]; +} + +int ActivationMishRT::initialize() NOEXCEPT { return 0; } + +void ActivationMishRT::terminate() NOEXCEPT {} + +size_t ActivationMishRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } + +int ActivationMishRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + activationMishForward((dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, stream); + return 0; +} +size_t ActivationMishRT::getSerializationSize() const NOEXCEPT { + return 1 * sizeof(int); +} + +void ActivationMishRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer), *a = buf; + writeBUF(buf, size); + assert(buf == a + getSerializationSize()); +} + +const char* ActivationMishRT::getPluginType() const NOEXCEPT { + return "ActivationMishRT_tkDNN"; +} + +const char *ActivationMishRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +bool ActivationMishRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; +} + +const char *ActivationMishRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ActivationMishRT::setPluginNamespace(const char *plguinNamespace) NOEXCEPT { + mPluginNamespace = plguinNamespace; +} + +IPluginV2 *ActivationMishRT::clone() const NOEXCEPT { + auto *p = new ActivationMishRT(); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + +ActivationMishRTPluginCreator::ActivationMishRTPluginCreator() { + mPluginAttributes.clear(); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ActivationMishRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *ActivationMishRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *ActivationMishRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new ActivationMishRT(serialData, serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *ActivationMishRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + auto *pluginObj = new ActivationMishRT(); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *ActivationMishRTPluginCreator::getPluginName() const NOEXCEPT { + return "ActivationMishRT_tkDNN"; +} + +const char *ActivationMishRTPluginCreator::getPluginVersion() const NOEXCEPT{ + return "1"; +} + +const PluginFieldCollection *ActivationMishRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + diff --git a/src/pluginsRT/ActivationReLUCeilingRT.cpp b/src/pluginsRT/ActivationReLUCeilingRT.cpp new file mode 100644 index 0000000..80b5a74 --- /dev/null +++ b/src/pluginsRT/ActivationReLUCeilingRT.cpp @@ -0,0 +1,129 @@ +#include +using namespace nvinfer1; + +std::vector ActivationReLUCeilingPluginCreator::mPluginAttributes; +PluginFieldCollection ActivationReLUCeilingPluginCreator::mFC{}; + +ActivationReLUCeiling::ActivationReLUCeiling(const float ceiling) { + this->ceiling = ceiling; +} + +ActivationReLUCeiling::~ActivationReLUCeiling() { + +} + +ActivationReLUCeiling::ActivationReLUCeiling(const void *data, size_t length) { + const char *buf = reinterpret_cast(data), *bufCheck = buf; + ceiling = readBUF(buf); + size = readBUF(buf); + assert(buf == bufCheck + length); +} + +int ActivationReLUCeiling::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims ActivationReLUCeiling::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { return inputs[0]; } + +void ActivationReLUCeiling::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT { + assert(type == DataType::kFLOAT && format == PluginFormat::kLINEAR); + size = 1; + for (int i = 0; i < outputDims[0].nbDims; i++) + size *= outputDims[0].d[i]; +} + +int ActivationReLUCeiling::initialize() NOEXCEPT { return 0; } + +void ActivationReLUCeiling::terminate() NOEXCEPT {} + +size_t ActivationReLUCeiling::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +int ActivationReLUCeiling::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT { + activationReLUCeilingForward((dnnType *) reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize * size, ceiling, stream); + return 0; +} + +size_t ActivationReLUCeiling::getSerializationSize() const NOEXCEPT { + return 1 * sizeof(int) + 1 * sizeof(float); +} + +void ActivationReLUCeiling::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer), *a = buf; + writeBUF(buf, ceiling); + writeBUF(buf, size); + assert(buf = a + getSerializationSize()); +} + +IPluginV2 *ActivationReLUCeiling::clone() const NOEXCEPT { + auto *p = new ActivationReLUCeiling(ceiling); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + +bool ActivationReLUCeiling::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR); +} + +void ActivationReLUCeiling::destroy() NOEXCEPT { delete this; } + +const char *ActivationReLUCeiling::getPluginType() const NOEXCEPT { + return "ActivationReLUCeilingRT_tkDNN"; +} + +const char *ActivationReLUCeiling::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const char *ActivationReLUCeiling::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ActivationReLUCeiling::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +ActivationReLUCeilingPluginCreator::ActivationReLUCeilingPluginCreator() { + mPluginAttributes.emplace_back(PluginField("ceiling", nullptr, PluginFieldType::kFLOAT32, 1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ActivationReLUCeilingPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *ActivationReLUCeilingPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *ActivationReLUCeilingPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new ActivationReLUCeiling(serialData, serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *ActivationReLUCeilingPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + float ceiling = *(static_cast(fields[0].data)); + auto *pluginObj = new ActivationReLUCeiling(ceiling); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *ActivationReLUCeilingPluginCreator::getPluginName() const NOEXCEPT { + return "ActivationReLUCeilingRT_tkDNN"; +} + +const char *ActivationReLUCeilingPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *ActivationReLUCeilingPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + diff --git a/src/pluginsRT/DeformableConvRT.cpp b/src/pluginsRT/DeformableConvRT.cpp new file mode 100644 index 0000000..417af8e --- /dev/null +++ b/src/pluginsRT/DeformableConvRT.cpp @@ -0,0 +1,331 @@ +#include +using namespace nvinfer1; +using namespace tk::dnn; + +std::vector DeformableConvRTPluginCreator::mPluginAttributes; +PluginFieldCollection DeformableConvRTPluginCreator::mFC{}; + + +DeformableConvRT::DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw, int deformableGroup, + int i_n, int i_c, int i_h, int i_w, int o_n, int o_c, int o_h, int o_w, + tk::dnn::DeformConv2d *deformable) { + this->chunk_dim = chunk_dim; + this->kh = kh; + this->kw = kw; + this->sh = sh; + this->sw = sw; + this->ph = ph; + this->pw = pw; + this->deformableGroup = deformableGroup; + this->i_n = i_n; + this->i_c = i_c; + this->i_h = i_h; + this->i_w = i_w; + this->o_n = o_n; + this->o_c = o_c; + this->o_h = o_h; + this->o_w = o_w; + this->defRT = deformable; + + height_ones = (i_h + 2 * ph - (1 * (kh - 1) + 1)) / sh + 1; + width_ones = (i_w + 2 * pw - (1 * (kw - 1) + 1)) / sw + 1; + dim_ones = i_c * kh * kw * 1 * height_ones * width_ones; + + checkCuda( cudaMalloc(&data_d, i_c * o_c * kh * kw * 1 * sizeof(dnnType))); + checkCuda( cudaMalloc(&bias2_d, o_c*sizeof(dnnType))); + checkCuda( cudaMalloc(&ones_d1, height_ones * width_ones * sizeof(dnnType))); + checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType))); + checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); + checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType))); + if(deformable != nullptr) { + checkCuda( cudaMemcpy(data_d, deformable->data_d, sizeof(dnnType)*i_c * o_c * kh * kw * 1, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(bias2_d, deformable->bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(ones_d1, deformable->ones_d1, sizeof(dnnType)*height_ones*width_ones, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(offset, deformable->offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(mask, deformable->mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToDevice) ); + checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) ); + } + stat = cublasCreate(&handle); + if (stat != CUBLAS_STATUS_SUCCESS) + FatalError("CUBLAS initialization failed\n"); + +} + +DeformableConvRT::~DeformableConvRT() { + checkCuda( cudaFree(data_d) ); + checkCuda( cudaFree(bias2_d) ); + checkCuda( cudaFree(ones_d1) ); + checkCuda( cudaFree(offset) ); + checkCuda( cudaFree(mask) ); + checkCuda( cudaFree(ones_d2) ); + cublasDestroy(handle); +} + +DeformableConvRT::DeformableConvRT(const void *data, size_t length) { + const char* buf = reinterpret_cast(data),*bufCheck = buf; + chunk_dim = readBUF(buf); + kh = readBUF(buf); + kw = readBUF(buf); + sh = readBUF(buf); + sw = readBUF(buf); + ph = readBUF(buf); + pw = readBUF(buf); + deformableGroup = readBUF(buf); + i_n = readBUF(buf); + i_c = readBUF(buf); + i_h = readBUF(buf); + i_w = readBUF(buf); + o_n = readBUF(buf); + o_c = readBUF(buf); + o_h = readBUF(buf); + o_w = readBUF(buf); + dnnType *aus = new dnnType[chunk_dim*2]; + for(int i=0;i(buf); + checkCuda(cudaMemcpy(offset,aus,sizeof(dnnType)*2*chunk_dim,cudaMemcpyHostToDevice)); + free(aus); + + aus = new dnnType[chunk_dim]; + for(int i=0;i(buf); + checkCuda(cudaMemcpy(mask,aus,sizeof(dnnType)*chunk_dim,cudaMemcpyHostToDevice)); + free(aus); + + aus = new dnnType[i_c*o_c*kh*kw*1]; + for(int i=0;i<(i_c*o_c*kh*kw*1);i++) + aus[i] = readBUF(buf); + checkCuda(cudaMemcpy(data_d,aus,sizeof(dnnType)*(i_c*o_c*kh*kw*1),cudaMemcpyHostToDevice)); + free(aus); + + aus = new dnnType[o_c]; + for(int i=0; i < o_c; i++) + aus[i] = readBUF(buf); + checkCuda( cudaMemcpy(bias2_d, aus, sizeof(dnnType)*o_c, cudaMemcpyHostToDevice) ); + free(aus); + + aus = new dnnType[height_ones * width_ones]; + for(int i=0; i(buf); + checkCuda( cudaMemcpy(ones_d1, aus, sizeof(dnnType)*height_ones * width_ones, cudaMemcpyHostToDevice) ); + free(aus); + + aus = new dnnType[dim_ones]; + for(int i=0; i(buf); + checkCuda( cudaMemcpy(ones_d2, aus, sizeof(dnnType)*dim_ones, cudaMemcpyHostToDevice) ); + free(aus); + + assert(buf == bufCheck + length); + +} + +int DeformableConvRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims DeformableConvRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w}; +} + +void DeformableConvRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {} + +int DeformableConvRT::initialize() NOEXCEPT { + return 0; +} + +void DeformableConvRT::terminate() NOEXCEPT {} + +size_t DeformableConvRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {return 0;} + +int DeformableConvRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]); + + // split conv2d outputs into offset to mask + for(int b=0; b(outputs[0]), ones_d2, + kh, kw, + sh, sw, + ph, pw, + 1, 1, + deformableGroup, b, + i_n, i_c, i_h, i_w, + o_n, o_c, o_h, o_w, + chunk_dim); + } + return 0; +} + +size_t DeformableConvRT::getSerializationSize() const NOEXCEPT { + return 16 * sizeof(int) + chunk_dim * 3 * sizeof(dnnType) + (i_c * o_c * kh * kw * 1 ) * sizeof(dnnType) + + o_c * sizeof(dnnType) + height_ones * width_ones * sizeof(dnnType) + dim_ones * sizeof(dnnType); +} + +void DeformableConvRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf, chunk_dim); + writeBUF(buf, kh); + writeBUF(buf, kw); + writeBUF(buf, sh); + writeBUF(buf, sw); + writeBUF(buf, ph); + writeBUF(buf, pw); + writeBUF(buf, deformableGroup); + writeBUF(buf, i_n); + writeBUF(buf, i_c); + writeBUF(buf, i_h); + writeBUF(buf, i_w); + writeBUF(buf, o_n); + writeBUF(buf, o_c); + writeBUF(buf, o_h); + writeBUF(buf, o_w); + dnnType *aus = new dnnType[chunk_dim*2]; + checkCuda( cudaMemcpy(aus, offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToHost) ); + for(int i=0; isetPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +DeformableConvRTPluginCreator::DeformableConvRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("chunk_dim",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("kh",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("kw",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("sh",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("sw",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("ph",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("pw",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("deformableGroup",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("i_n",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("i_c",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("i_h",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("i_w",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("o_n",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("defRT",nullptr,PluginFieldType::kUNKNOWN,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void DeformableConvRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *DeformableConvRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *DeformableConvRTPluginCreator::deserializePlugin(const char *name, const void *serialData, + size_t serialLength) NOEXCEPT { + auto *pluginObj = new DeformableConvRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *DeformableConvRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + int chunk_dim = *(static_cast(fields[0].data)); + int kh = *(static_cast(fields[1].data)); + int kw = *(static_cast(fields[2].data)); + int sh = *(static_cast(fields[3].data)); + int sw = *(static_cast(fields[4].data)); + int ph = *(static_cast(fields[5].data)); + int pw = *(static_cast(fields[6].data)); + int deformableGroup = *(static_cast(fields[7].data)); + int i_n = *(static_cast(fields[8].data)); + int i_c = *(static_cast(fields[9].data)); + int i_h = *(static_cast(fields[10].data)); + int i_w = *(static_cast(fields[11].data)); + int o_n = *(static_cast(fields[12].data)); + int o_c = *(static_cast(fields[13].data)); + int o_h = *(static_cast(fields[14].data)); + int o_w = *(static_cast(fields[14].data)); + auto *defRT = const_cast(static_cast(fields[15].data)); + auto *pluginObj = new DeformableConvRT(chunk_dim,kh,kw,sh,sw,ph,pw,deformableGroup,i_n,i_c,i_h,i_w,o_n,o_c,o_h,o_w,defRT); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *DeformableConvRTPluginCreator::getPluginName() const NOEXCEPT { + return "DeformableConvRT_tkDNN"; +} + +const char *DeformableConvRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *DeformableConvRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + diff --git a/src/pluginsRT/FlattenConcatRT.cpp b/src/pluginsRT/FlattenConcatRT.cpp new file mode 100644 index 0000000..51cebc5 --- /dev/null +++ b/src/pluginsRT/FlattenConcatRT.cpp @@ -0,0 +1,163 @@ +#include +using namespace nvinfer1; + +std::vector FlattenConcatRTPluginCreator::mPluginAttributes; +PluginFieldCollection FlattenConcatRTPluginCreator::mFC{}; + +FlattenConcatRT::FlattenConcatRT() { + stat = cublasCreate(&handle); + if (stat != CUBLAS_STATUS_SUCCESS) { + printf ("CUBLAS initialization failed\n"); + return; + } +} + +FlattenConcatRT::FlattenConcatRT(const void *data, size_t length) { + const char *buf = reinterpret_cast(data),*bufCheck=buf; + c = readBUF(buf); + h = readBUF(buf); + w = readBUF(buf); + rows = readBUF(buf); + cols = readBUF(buf); + assert(buf == bufCheck + length); +} + +FlattenConcatRT::~FlattenConcatRT() {} + +int FlattenConcatRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims FlattenConcatRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1}; +} + +void FlattenConcatRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, + DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT { + assert(nbOutputs == 1 && nbInputs ==1); + rows = inputDims[0].d[0]; + cols = inputDims[0].d[1] * inputDims[0].d[2]; + c = inputDims[0].d[0] * inputDims[0].d[1] * inputDims[0].d[2]; + h = 1; + w = 1; +} + +int FlattenConcatRT::initialize() NOEXCEPT { + return 0; +} + +void FlattenConcatRT::terminate() NOEXCEPT { + checkERROR(cublasDestroy(handle)); +} + +size_t FlattenConcatRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +int FlattenConcatRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + + checkERROR( cublasSetStream(handle, stream) ); + for(int i=0; i(buffer),*a = buf; + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); + writeBUF(buf, rows); + writeBUF(buf, cols); + assert(buf == a + getSerializationSize()); +} + +void FlattenConcatRT::destroy() NOEXCEPT { + delete this; +} + +bool FlattenConcatRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; +} + +const char *FlattenConcatRT::getPluginType() const NOEXCEPT { + return "FlattenConcatRT_tkDNN"; +} + +const char *FlattenConcatRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const char *FlattenConcatRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void FlattenConcatRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2 *FlattenConcatRT::clone() const NOEXCEPT { + auto *p = new FlattenConcatRT(); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +FlattenConcatRTPluginCreator::FlattenConcatRTPluginCreator() { + mPluginAttributes.clear(); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void FlattenConcatRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *FlattenConcatRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *FlattenConcatRTPluginCreator::deserializePlugin(const char *name, const void *serialData, + size_t serialLength) NOEXCEPT { + auto *pluginObj = new FlattenConcatRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *FlattenConcatRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + auto *pluginObj = new FlattenConcatRT(); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *FlattenConcatRTPluginCreator::getPluginName() const NOEXCEPT { + return "FlattenConcatRT_tkDNN"; +} + +const char *FlattenConcatRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *FlattenConcatRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + diff --git a/src/pluginsRT/MaxPoolingSizeRT.cpp b/src/pluginsRT/MaxPoolingSizeRT.cpp new file mode 100644 index 0000000..4b95d7f --- /dev/null +++ b/src/pluginsRT/MaxPoolingSizeRT.cpp @@ -0,0 +1,179 @@ +#include +using namespace nvinfer1; + +std::vector MaxPoolFixedSizeRTPluginCreator::mPluginAttributes; +PluginFieldCollection MaxPoolFixedSizeRTPluginCreator::mFC{}; + +MaxPoolFixedSizeRT::MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding){ + this->c = c; + this->h = h; + this->w = w; + this->n = n; + this->stride_H = strideH; + this->stride_W = strideW; + this->winSize = winSize; + this->padding = padding; +} + +MaxPoolFixedSizeRT::MaxPoolFixedSizeRT(const void *data, size_t length) { + const char *buf = reinterpret_cast(data),*bufCheck = buf; + c = readBUF(buf); + h = readBUF(buf); + w = readBUF(buf); + n = readBUF(buf); + stride_H = readBUF(buf); + stride_W = readBUF(buf); + winSize = readBUF(buf); + padding = readBUF(buf); + assert(buf == bufCheck + length); +} + +MaxPoolFixedSizeRT::~MaxPoolFixedSizeRT() { + +} + +int MaxPoolFixedSizeRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims MaxPoolFixedSizeRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3{this->c, this->h, this->w}; +} + +void MaxPoolFixedSizeRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {} + +int MaxPoolFixedSizeRT::initialize() NOEXCEPT { + return 0; +} + +void MaxPoolFixedSizeRT::terminate() NOEXCEPT { + +} + +size_t MaxPoolFixedSizeRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +int MaxPoolFixedSizeRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + MaxPoolingForward(srcData, dstData, batchSize, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding, stream); + return 0; +} + +size_t MaxPoolFixedSizeRT::getSerializationSize() const NOEXCEPT { + return 8*sizeof(int); +} + +void MaxPoolFixedSizeRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf, this->c); + writeBUF(buf, this->h); + writeBUF(buf, this->w); + writeBUF(buf, this->n); + writeBUF(buf, this->stride_H); + writeBUF(buf, this->stride_W); + writeBUF(buf, this->winSize); + writeBUF(buf, this->padding); + assert(buf == a + getSerializationSize()); +} + +void MaxPoolFixedSizeRT::destroy() NOEXCEPT { +delete this; +} + +bool MaxPoolFixedSizeRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; + //todo assert +} + +const char *MaxPoolFixedSizeRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void MaxPoolFixedSizeRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *MaxPoolFixedSizeRT::getPluginType() const NOEXCEPT { + return "MaxPoolingFixedSizeRT_tkDNN"; +} + +const char *MaxPoolFixedSizeRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +IPluginV2 *MaxPoolFixedSizeRT::clone() const NOEXCEPT { + auto *p = new MaxPoolFixedSizeRT(c,h,w,n,stride_H,stride_W,winSize,padding); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +MaxPoolFixedSizeRTPluginCreator::MaxPoolFixedSizeRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("c",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("h",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("w",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("n",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("stride_H",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("stride_W",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("winSize",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("padding",nullptr,PluginFieldType::kINT32,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void MaxPoolFixedSizeRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *MaxPoolFixedSizeRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *MaxPoolFixedSizeRTPluginCreator::deserializePlugin(const char *name, const void *serialData,size_t serialLength) NOEXCEPT { + auto *pluginObj = new MaxPoolFixedSizeRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *MaxPoolFixedSizeRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + //todo assert + int c = *(static_cast(fields[0].data)); + int h = *(static_cast(fields[1].data)); + int w = *(static_cast(fields[2].data)); + int n = *(static_cast(fields[3].data)); + int stride_H = *(static_cast(fields[4].data)); + int stride_W = *(static_cast(fields[5].data)); + int winSize = *(static_cast(fields[6].data)); + int padding = *(static_cast(fields[7].data)); + auto *pluginObj = new MaxPoolFixedSizeRT(c,h,w,n,stride_H,stride_W,winSize,padding); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *MaxPoolFixedSizeRTPluginCreator::getPluginName() const NOEXCEPT { + return "MaxPoolingFixedSizeRT_tkDNN"; +} + +const char *MaxPoolFixedSizeRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *MaxPoolFixedSizeRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + + + + + diff --git a/src/pluginsRT/RegionRT.cpp b/src/pluginsRT/RegionRT.cpp new file mode 100644 index 0000000..de1dd3b --- /dev/null +++ b/src/pluginsRT/RegionRT.cpp @@ -0,0 +1,174 @@ +#include +using namespace nvinfer1; +std::vector RegionRTPluginCreator::mPluginAttributes; +PluginFieldCollection RegionRTPluginCreator::mFC{}; + +RegionRT::RegionRT(int classes, int coords, int num) { + this->classes = classes; + this->coords = coords; + this->num = num; +} + +RegionRT::~RegionRT() {} + +RegionRT::RegionRT(const void *data, size_t length) { + const char *buf = reinterpret_cast(data),*bufCheck=buf; + classes = readBUF(buf); + coords = readBUF(buf); + num = readBUF(buf); + c = readBUF(buf); + h = readBUF(buf); + w = readBUF(buf); + assert(buf == bufCheck+length); +} + +int RegionRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims RegionRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return inputs[0]; +} + +void RegionRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT { + c = inputDims[0].d[0]; + h = inputDims[0].d[1]; + w = inputDims[0].d[2]; +} + +int RegionRT::initialize() NOEXCEPT {return 0;} + +void RegionRT::terminate() NOEXCEPT {} + +size_t RegionRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } + +int RegionRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + + for (int b = 0; b < batchSize; ++b){ + for(int n = 0; n < num; ++n){ + int index = entry_index(b, n*w*h, 0); + activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); + + index = entry_index(b, n*w*h, coords); + activationLOGISTICForward(srcData + index, dstData + index, w*h, stream); + } + } + + //softmax start + int index = entry_index(0, 0, coords + 1); + softmaxForward( srcData + index, classes, batchSize*num, + (c*h*w)/num, + w*h, 1, w*h, 1, dstData + index, stream); + + return 0; +} + +size_t RegionRT::getSerializationSize() const NOEXCEPT { + return 6*sizeof(int); +} + +void RegionRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf, classes); + writeBUF(buf, coords); + writeBUF(buf, num); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); + assert(buf == a + getSerializationSize()); +} + +const char *RegionRT::getPluginType() const NOEXCEPT { + return "RegionRT_tkDNN"; +} + +const char *RegionRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +void RegionRT::destroy() NOEXCEPT { delete this; } + +const char *RegionRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void RegionRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +bool RegionRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; +} + +IPluginV2 *RegionRT::clone() const NOEXCEPT { + auto *p = new RegionRT(classes,coords,num); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +RegionRTPluginCreator::RegionRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("coords",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void RegionRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *RegionRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *RegionRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new RegionRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *RegionRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + assert(fc->nbFields == 3); + assert(fields[0].type == PluginFieldType::kINT32); + assert(fields[1].type == PluginFieldType::kINT32); + assert(fields[2].type == PluginFieldType::kINT32); + int classes = *(static_cast(fields[0].data)); + int coords = *(static_cast(fields[1].data)); + int num = *(static_cast(fields[2].data)); + RegionRT *pluginObj = new RegionRT(classes,coords,num); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *RegionRTPluginCreator::getPluginName() const NOEXCEPT { + return "RegionRT_tkDNN"; +} + +const char *RegionRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *RegionRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + + + + + diff --git a/src/pluginsRT/ReorgRT.cpp b/src/pluginsRT/ReorgRT.cpp new file mode 100644 index 0000000..d3dd12a --- /dev/null +++ b/src/pluginsRT/ReorgRT.cpp @@ -0,0 +1,151 @@ +#include +using namespace nvinfer1; + +std::vector ReorgRTPluginCreator::mPluginAttributes; +PluginFieldCollection ReorgRTPluginCreator::mFC{}; + +ReorgRT::ReorgRT(int stride) { + this->stride = stride; +} + +ReorgRT::~ReorgRT() {} + +ReorgRT::ReorgRT(const void *data, size_t length) { + const char* buf = reinterpret_cast(data),*bufCheck = buf; + stride = readBUF(buf); + c = readBUF(buf); + h = readBUF(buf); + w = readBUF(buf); + assert(buf == bufCheck + length); +} + +int ReorgRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims ReorgRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride}; +} + +void ReorgRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT { + c = inputDims[0].d[0]; + h = inputDims[0].d[1]; + w = inputDims[0].d[2]; +} + +int ReorgRT::initialize() NOEXCEPT { + return 0; +} + +void ReorgRT::terminate() NOEXCEPT {} + +size_t ReorgRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +int ReorgRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT { + reorgForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), + batchSize, c, h, w, stride, stream); + return 0; +} + +size_t ReorgRT::getSerializationSize() const NOEXCEPT { + return 4*sizeof(int); +} + +void ReorgRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf, stride); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); + assert(buf == a + getSerializationSize()); +} + +bool ReorgRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; +} + +const char *ReorgRT::getPluginType() const NOEXCEPT { + return "ReorgRT_tkDNN"; +} + +const char *ReorgRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +void ReorgRT::destroy() NOEXCEPT { + delete this; +} + +const char *ReorgRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ReorgRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2 *ReorgRT::clone() const NOEXCEPT { + auto *p = new ReorgRT(stride); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +ReorgRTPluginCreator::ReorgRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ReorgRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *ReorgRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *ReorgRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new ReorgRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *ReorgRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + assert(fc->nbFields == 1); + assert(fields[0].type == PluginFieldType::kINT32); + int stride = *(static_cast(fields[0].data)); + auto *pluginObj = new ReorgRT(stride); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *ReorgRTPluginCreator::getPluginName() const NOEXCEPT { + return "ReorgRT_tkDNN"; +} + +const char *ReorgRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *ReorgRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + + + + + + + diff --git a/src/pluginsRT/ReshapeRT.cpp b/src/pluginsRT/ReshapeRT.cpp new file mode 100644 index 0000000..5e23ab3 --- /dev/null +++ b/src/pluginsRT/ReshapeRT.cpp @@ -0,0 +1,152 @@ +#include +using namespace nvinfer1; + +std::vector ReshapeRTPluginCreator::mPluginAttributes; +PluginFieldCollection ReshapeRTPluginCreator::mFC{}; + +ReshapeRT::ReshapeRT(dataDim_t newDim) { + new_dim = newDim; + n = new_dim.n; + c = new_dim.c; + h = new_dim.h; + w = new_dim.w; +} + +ReshapeRT::ReshapeRT(const void *data, size_t length) { + const char *buf = reinterpret_cast(data),*bufCheck = buf; + new_dim.n = readBUF(buf); + new_dim.c = readBUF(buf); + new_dim.h = readBUF(buf); + new_dim.w = readBUF(buf); + assert(buf == bufCheck + length); +} + +ReshapeRT::~ReshapeRT() {} + +int ReshapeRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims ReshapeRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3{ c,h,w} ; +} + +void ReshapeRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {} + +int ReshapeRT::initialize() NOEXCEPT { + return 0; +} + +void ReshapeRT::terminate() NOEXCEPT {} + +size_t ReshapeRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +int ReshapeRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + return 0; +} + +size_t ReshapeRT::getSerializationSize() const NOEXCEPT { + return 4*sizeof(int); +} + +void ReshapeRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a = buf; + writeBUF(buf, n); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); + assert(buf == a + getSerializationSize()); +} + +bool ReshapeRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; + //todo assert +} + +const char *ReshapeRT::getPluginType() const NOEXCEPT { + return "ReshapeRT_tkDNN"; +} + +const char *ReshapeRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +void ReshapeRT::destroy() NOEXCEPT { + delete this; +} + +const char *ReshapeRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ReshapeRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2 *ReshapeRT::clone() const NOEXCEPT { + auto *p = new ReshapeRT(new_dim); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +ReshapeRTPluginCreator::ReshapeRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("new_dim",nullptr,PluginFieldType::kUNKNOWN,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ReshapeRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *ReshapeRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *ReshapeRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new ReshapeRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *ReshapeRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + dataDim_t newDim = *(static_cast(fields[0].data)); + ReshapeRT *pluginObj = new ReshapeRT(newDim); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *ReshapeRTPluginCreator::getPluginName() const NOEXCEPT { + return "ReshapeRT_tkDNN"; +} + +const char *ReshapeRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *ReshapeRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + + + + + + + diff --git a/src/pluginsRT/ResizeLayerRT.cpp b/src/pluginsRT/ResizeLayerRT.cpp new file mode 100644 index 0000000..12a3ff9 --- /dev/null +++ b/src/pluginsRT/ResizeLayerRT.cpp @@ -0,0 +1,163 @@ +#include +using namespace nvinfer1; + +std::vector ResizeLayerRTPluginCreator::mPluginAttributes; +PluginFieldCollection ResizeLayerRTPluginCreator::mFC{}; + + +ResizeLayerRT::ResizeLayerRT(int c, int h, int w) { + o_c = c; + o_h = h; + o_w = w; +} + +ResizeLayerRT::ResizeLayerRT(const void *data, size_t length) { + const char *buf = reinterpret_cast(data),*bufCheck = buf; + o_c = readBUF(buf); + o_h = readBUF(buf); + o_w = readBUF(buf); + i_c = readBUF(buf); + i_h = readBUF(buf); + i_w = readBUF(buf); + assert(buf == bufCheck + length); +} + +ResizeLayerRT::~ResizeLayerRT() {} + +int ResizeLayerRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims ResizeLayerRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3{o_c, o_h, o_w}; +} + +void ResizeLayerRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, + DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT { + i_c = inputDims[0].d[0]; + i_h = inputDims[0].d[1]; + i_w = inputDims[0].d[2]; +} + +int ResizeLayerRT::initialize() NOEXCEPT { + return 0; +} + +void ResizeLayerRT::terminate() NOEXCEPT {} + +size_t ResizeLayerRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } + +int ResizeLayerRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + resizeForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), + batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream); + return 0; +} + +size_t ResizeLayerRT::getSerializationSize() const NOEXCEPT { + return 6*sizeof(int); +} + +void ResizeLayerRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf, o_c); + writeBUF(buf, o_h); + writeBUF(buf, o_w); + writeBUF(buf, i_c); + writeBUF(buf, i_h); + writeBUF(buf, i_w); + assert(buf == a + getSerializationSize()); +} + +bool ResizeLayerRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; + //todo assert +} + +const char *ResizeLayerRT::getPluginType() const NOEXCEPT { + return "ResizeLayerRT_tkDNN"; +} + +const char *ResizeLayerRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +void ResizeLayerRT::destroy() NOEXCEPT { + delete this; +} + +const char *ResizeLayerRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ResizeLayerRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2 *ResizeLayerRT::clone() const NOEXCEPT { + auto *p = new ResizeLayerRT(o_c,o_h,o_w); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +ResizeLayerRTPluginCreator::ResizeLayerRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ResizeLayerRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *ResizeLayerRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *ResizeLayerRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new ResizeLayerRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *ResizeLayerRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + assert(fc->nbFields == 3); + assert(fields[0].type == PluginFieldType::kINT32); + assert(fields[1].type == PluginFieldType::kINT32); + assert(fields[2].type == PluginFieldType::kINT32); + int oc = *(static_cast(fields[0].data)); + int oh = *(static_cast(fields[1].data)); + int ow = *(static_cast(fields[2].data)); + auto *pluginObj = new ResizeLayerRT(oc,oh,ow); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *ResizeLayerRTPluginCreator::getPluginName() const NOEXCEPT { + return "ResizeLayerRT_tkDNN"; +} + +const char *ResizeLayerRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *ResizeLayerRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + + + + + diff --git a/src/pluginsRT/RouteRT.cpp b/src/pluginsRT/RouteRT.cpp new file mode 100644 index 0000000..87233a5 --- /dev/null +++ b/src/pluginsRT/RouteRT.cpp @@ -0,0 +1,183 @@ +#include +using namespace nvinfer1; + +std::vector RouteRTPluginCreator::mPluginAttributes; +PluginFieldCollection RouteRTPluginCreator::mFC{}; + + +RouteRT::RouteRT(int groups, int group_id) { + this->groups = groups; + this->group_id = group_id; +} + +RouteRT::~RouteRT() {} + +RouteRT::RouteRT(const void *data, size_t length) { + const char* buf = reinterpret_cast(data),*bufCheck = buf; + groups = readBUF(buf); + group_id = readBUF(buf); + in = readBUF(buf); + for(int i=0;i (buf); + } + c= readBUF(buf); + h = readBUF(buf); + w = readBUF(buf); + assert(buf == bufCheck + length); +} + +int RouteRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims RouteRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + int out_c = 0; + for(int i=0; i(outputs[0]); + for(int b=0; b(inputs[i]); + int in_dim = c_in[i]*h*w; + int part_in_dim = in_dim / this->groups; + checkCuda( cudaMemcpyAsync(dstData + b*c*w*h + offset, input + b*c*w*h*groups + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) ); + offset += part_in_dim; + } + } + return 0; +} + +size_t RouteRT::getSerializationSize() const NOEXCEPT { + return (6+MAX_INPUTS)*sizeof(int); +} + +void RouteRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf, groups); + writeBUF(buf, group_id); + writeBUF(buf, in); + for(int i=0; isetPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +RouteRTPluginCreator::RouteRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("groups",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("group_id",nullptr,PluginFieldType::kINT32,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void RouteRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *RouteRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *RouteRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new RouteRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *RouteRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + assert(fc->nbFields == 2); + assert(fields[0].type == PluginFieldType::kINT32); + assert(fields[1].type == PluginFieldType::kINT32); + int groups = *(static_cast(fields[0].data)); + int group_id = *(static_cast(fields[1].data)); + RouteRT *pluginObj = new RouteRT(groups,group_id); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *RouteRTPluginCreator::getPluginName() const NOEXCEPT { + return "RouteRT_tkDNN"; +} + +const char *RouteRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *RouteRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + + + + + diff --git a/src/pluginsRT/ShortcutRT.cpp b/src/pluginsRT/ShortcutRT.cpp new file mode 100644 index 0000000..eb35bd6 --- /dev/null +++ b/src/pluginsRT/ShortcutRT.cpp @@ -0,0 +1,167 @@ +#include +using namespace nvinfer1; + +std::vector ShortcutRTPluginCreator::mPluginAttributes; +PluginFieldCollection ShortcutRTPluginCreator::mFC{}; + +ShortcutRT::ShortcutRT(tk::dnn::dataDim_t bdim, bool mul) { + bDim = bdim; + this->bc = bDim.c; + this->bh = bDim.h; + this->bw = bDim.w; + this->mul = mul; +} + +ShortcutRT::~ShortcutRT() {} + +ShortcutRT::ShortcutRT(const void *data, size_t length) { + const char* buf =reinterpret_cast(data),*bufCheck = buf; + bDim.c = readBUF(buf); + bDim.h = readBUF(buf); + bDim.w = readBUF(buf); + bDim.l = 1; + mul = readBUF(buf); + c = readBUF(buf); + h = readBUF(buf); + w = readBUF(buf); + assert(buf == bufCheck + length); +} + +int ShortcutRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims ShortcutRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]}; +} + +void ShortcutRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, + DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT { + c = inputDims[0].d[0]; + h = inputDims[0].d[1]; + w = inputDims[0].d[2]; +} + +int ShortcutRT::initialize() NOEXCEPT { + return 0; +} + +void ShortcutRT::terminate() NOEXCEPT {} + +size_t ShortcutRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; } + +int ShortcutRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *srcDataBack = (dnnType*)reinterpret_cast(inputs[1]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, mul, stream); + + return 0; +} + +size_t ShortcutRT::getSerializationSize() const NOEXCEPT { + return 6*sizeof(int) + sizeof(bool); +} + +void ShortcutRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf, bc); + writeBUF(buf, bh); + writeBUF(buf, bw); + writeBUF(buf, mul); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); + assert(buf == a + getSerializationSize()); +} + +bool ShortcutRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; +} + +const char *ShortcutRT::getPluginType() const NOEXCEPT { + return "ShortcutRT_tkDNN"; +} + +const char *ShortcutRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +void ShortcutRT::destroy() NOEXCEPT { + delete this; +} + +const char *ShortcutRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void ShortcutRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2 *ShortcutRT::clone() const NOEXCEPT { + auto *p = new ShortcutRT(bDim,mul); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +ShortcutRTPluginCreator::ShortcutRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("bDim",nullptr,PluginFieldType::kUNKNOWN,1)); + mPluginAttributes.emplace_back(PluginField("mul",nullptr,PluginFieldType::kUNKNOWN,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void ShortcutRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *ShortcutRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *ShortcutRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new ShortcutRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *ShortcutRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + //todo assert + tk::dnn::dataDim_t bdim = *(static_cast(fields[0].data)); + bool mul = *(static_cast(fields[1].data)); + auto *pluginObj = new ShortcutRT(bdim,mul); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *ShortcutRTPluginCreator::getPluginName() const NOEXCEPT { + return "ShortcutRT_tkDNN"; +} + +const char *ShortcutRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *ShortcutRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + + + + + + + diff --git a/src/pluginsRT/UpsampleRT.cpp b/src/pluginsRT/UpsampleRT.cpp new file mode 100644 index 0000000..044405b --- /dev/null +++ b/src/pluginsRT/UpsampleRT.cpp @@ -0,0 +1,155 @@ +#include +using namespace nvinfer1; + +std::vector UpsampleRTPluginCreator::mPluginAttributes; +PluginFieldCollection UpsampleRTPluginCreator::mFC{}; + +UpsampleRT::UpsampleRT(int stride) { + this->stride = stride; +} + +UpsampleRT::UpsampleRT(const void *data, size_t length) { + const char* buf = reinterpret_cast(data),*bufCheck=buf; + stride = readBUF(buf); + c = readBUF(buf); + h = readBUF(buf); + w = readBUF(buf); + assert(buf == bufCheck + length); +} + +UpsampleRT::~UpsampleRT() {} + +int UpsampleRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims UpsampleRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return Dims3(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride); +} + +void UpsampleRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, + DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT { + c = inputDims[0].d[0]; + h = inputDims[0].d[1]; + w = inputDims[0].d[2]; +} + +int UpsampleRT::initialize() NOEXCEPT { + return 0; +} + +void UpsampleRT::terminate() NOEXCEPT {} + +size_t UpsampleRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +int UpsampleRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + fill(dstData, batchSize*c*h*w*stride*stride, 0.0, stream); + upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream); + return 0; +} + +size_t UpsampleRT::getSerializationSize() const NOEXCEPT { + return 4*sizeof(int); +} + +void UpsampleRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer),*a=buf; + writeBUF(buf, stride); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); + assert(buf == a + getSerializationSize()); +} + +bool UpsampleRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; +} + +const char *UpsampleRT::getPluginType() const NOEXCEPT { + return "Upsample_tkDNN"; +} + +const char *UpsampleRT::getPluginVersion() const NOEXCEPT { + return "1"; +} + +void UpsampleRT::destroy() NOEXCEPT { + delete this; +} + +const char *UpsampleRT::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +void UpsampleRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +IPluginV2 *UpsampleRT::clone() const NOEXCEPT { + auto *p = new UpsampleRT(stride); + p->setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +UpsampleRTPluginCreator::UpsampleRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void UpsampleRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *UpsampleRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *UpsampleRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new UpsampleRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *UpsampleRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + int stride = *(static_cast(fields[0].data)); + auto *pluginObj = new UpsampleRT(stride); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *UpsampleRTPluginCreator::getPluginName() const NOEXCEPT { + return "Upsample_tkDNN"; +} + +const char *UpsampleRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *UpsampleRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + + + + + + + + + diff --git a/src/pluginsRT/YoloRT.cpp b/src/pluginsRT/YoloRT.cpp new file mode 100644 index 0000000..7f615b4 --- /dev/null +++ b/src/pluginsRT/YoloRT.cpp @@ -0,0 +1,253 @@ +#include +using namespace nvinfer1; + +std::vector YoloRTPluginCreator::mPluginAttributes; +PluginFieldCollection YoloRTPluginCreator::mFC{}; + +YoloRT::YoloRT(int classes, int num, tk::dnn::Yolo *Yolo, int n_masks, float scale_xy, float nms_thresh, int nms_kind, + int new_coords) { + this->yolo = Yolo; + this->classes = classes; + this->num = num; + this->n_masks = n_masks; + this->scaleXY = scale_xy; + this->nms_thresh = nms_thresh; + this->nms_kind = nms_kind; + this->new_coords = new_coords; + + mask = new dnnType[n_masks]; + bias = new dnnType[num * n_masks * 2]; + if (yolo != nullptr) { + memcpy(mask, yolo->mask_h, sizeof(dnnType) * n_masks); + memcpy(bias, yolo->bias_h, sizeof(dnnType) * num * n_masks * 2); + classesNames = yolo->classesNames; + } +} + +YoloRT::YoloRT(const void *data, size_t length) { + std::vector maskTemp,biasTemp; + std::cout<<"LENGTH : "<(data),*bufCheck = buf; + classes = readBUF(buf); + num = readBUF(buf); + n_masks = readBUF(buf); + scaleXY = readBUF(buf); + nms_thresh = readBUF(buf); + nms_kind = readBUF(buf); + new_coords = readBUF(buf); + c = readBUF(buf); + h = readBUF(buf); + w = readBUF(buf); + for(int i=0;i(buf)); + std::cout<(buf)); + std::cout<(buf); + classesNames[1] = std::string(tmp); + } + assert(buf == bufCheck + length); +} + +YoloRT::~YoloRT() {} + +int YoloRT::getNbOutputs() const NOEXCEPT { + return 1; +} + +Dims YoloRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { + return inputs[0]; +} + +void YoloRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, + PluginFormat format, int maxBatchSize) NOEXCEPT { + c = inputDims[0].d[0]; + h = inputDims[0].d[1]; + w = inputDims[0].d[2]; +} + +int YoloRT::initialize() NOEXCEPT { + return 0; +} + +void YoloRT::terminate() NOEXCEPT {} + +size_t YoloRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { + return 0; +} + +int YoloRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, + cudaStream_t stream) NOEXCEPT { + dnnType *srcData = (dnnType *) reinterpret_cast(inputs[0]); + dnnType *dstData = reinterpret_cast(outputs[0]); + + checkCuda(cudaMemcpyAsync(dstData, srcData, batchSize * c * h * w * sizeof(dnnType), cudaMemcpyDeviceToDevice, + stream)); + + + for (int b = 0; b < batchSize; ++b) { + for (int n = 0; n < n_masks; ++n) { + int index = entry_index(b, n * w * h, 0); + if (new_coords == 1) { + if (this->scaleXY != 1) + scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1); + } else { + activationLOGISTICForward(srcData + index, dstData + index, 2 * w * h, stream); //x,y + + if (this->scaleXY != 1) + scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1); + + index = entry_index(b, n * w * h, 4); + activationLOGISTICForward(srcData + index, dstData + index, (1 + classes) * w * h, stream); + } + } + } + + //std::cout<<"YOLO END\n"; + return 0; +} + +size_t YoloRT::getSerializationSize() const NOEXCEPT { + return 8 * sizeof(int) + 2 * sizeof(float) + n_masks * sizeof(dnnType) + num * n_masks * 2 * sizeof(dnnType) + + YOLORT_CLASSNAME_W * classes * sizeof(char); +} + +bool YoloRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT { + return true; +} + +void YoloRT::serialize(void *buffer) const NOEXCEPT { + char *buf = reinterpret_cast(buffer), *a = buf; + writeBUF(buf, classes); //std::cout << "Classes :" << classes << std::endl; + writeBUF(buf, num); //std::cout << "Num : " << num << std::endl; + std::cout<setPluginNamespace(mPluginNamespace.c_str()); + return p; +} + + +YoloRTPluginCreator::YoloRTPluginCreator() { + mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("yolo",nullptr,PluginFieldType::kUNKNOWN,1)); + mPluginAttributes.emplace_back(PluginField("numMasks",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("scaleXY",nullptr,PluginFieldType::kFLOAT32,1)); + mPluginAttributes.emplace_back(PluginField("nmsThresh",nullptr,PluginFieldType::kFLOAT32,1)); + mPluginAttributes.emplace_back(PluginField("nmsKind",nullptr,PluginFieldType::kINT32,1)); + mPluginAttributes.emplace_back(PluginField("newCoords",nullptr,PluginFieldType::kINT32,1)); + mFC.nbFields = mPluginAttributes.size(); + mFC.fields = mPluginAttributes.data(); +} + +void YoloRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT { + mPluginNamespace = pluginNamespace; +} + +const char *YoloRTPluginCreator::getPluginNamespace() const NOEXCEPT { + return mPluginNamespace.c_str(); +} + +IPluginV2 *YoloRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT { + auto *pluginObj = new YoloRT(serialData,serialLength); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +IPluginV2 *YoloRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT { + const PluginField *fields = fc->fields; + //todo assert + int classes = *(static_cast(fields[0].data)); + int num = *(static_cast(fields[1].data)); + Yolo *yoloTemp = const_cast(static_cast(fields[2].data)); + int numMasks = *(static_cast(fields[3].data)); + float scaleXY = *(static_cast(fields[4].data)); + float nmsThresh = *(static_cast(fields[5].data)); + int nmsKind = *(static_cast(fields[6].data)); + int newCoords = *(static_cast(fields[7].data)); + YoloRT *pluginObj = new YoloRT(classes,num,yoloTemp,numMasks,scaleXY,nmsThresh,nmsKind,newCoords); + pluginObj->setPluginNamespace(mPluginNamespace.c_str()); + return pluginObj; +} + +const char *YoloRTPluginCreator::getPluginName() const NOEXCEPT { + return "YoloRT_tkDNN"; +} + +const char *YoloRTPluginCreator::getPluginVersion() const NOEXCEPT { + return "1"; +} + +const PluginFieldCollection *YoloRTPluginCreator::getFieldNames() NOEXCEPT { + return &mFC; +} + + + + + + + + + + + + + +