Added TensorRT8 support #270

Closed
TheExDeus wants to merge 6 commits from feature/tensorrt8_support into master
43 changed files with 2493 additions and 788 deletions
+7 -7
View File
@@ -18,7 +18,7 @@ if(DEBUG)
endif() endif()
if(TKDNN_PATH) if(TKDNN_PATH)
message("SET TKDNN_PATH:"${TKDNN_PATH}) message("SET TKDNN_PATH:${TKDNN_PATH}")
add_definitions(-DTKDNN_PATH="${TKDNN_PATH}") add_definitions(-DTKDNN_PATH="${TKDNN_PATH}")
else() else()
add_definitions(-DTKDNN_PATH="${CMAKE_CURRENT_SOURCE_DIR}") add_definitions(-DTKDNN_PATH="${CMAKE_CURRENT_SOURCE_DIR}")
@@ -51,11 +51,11 @@ find_package(Eigen3 REQUIRED)
message("Eigen DIR: " ${EIGEN3_INCLUDE_DIR}) message("Eigen DIR: " ${EIGEN3_INCLUDE_DIR})
include_directories(${EIGEN3_INCLUDE_DIR}) include_directories(${EIGEN3_INCLUDE_DIR})
find_package(OpenCV REQUIRED) find_package(OpenCV 4.5 REQUIRED)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV")
# if(OpenCV_CUDA_VERSION) if(OpenCV_CUDA_VERSION)
# add_compile_definitions(OPENCV_CUDACONTRIB) add_compile_definitions(OPENCV_CUDACONTRIB)
# endif() endif()
# gives problems in cross-compiling, probably malformed cmake config # gives problems in cross-compiling, probably malformed cmake config
find_package(yaml-cpp REQUIRED) find_package(yaml-cpp REQUIRED)
@@ -72,8 +72,8 @@ add_library(tkDNN SHARED ${tkdnn_SRC})
target_link_libraries(tkDNN ${tkdnn_LIBS}) target_link_libraries(tkDNN ${tkdnn_LIBS})
#static #static
#add_library(tkDNN_static STATIC ${tkdnn_SRC}) add_library(tkDNN_static STATIC ${tkdnn_SRC})
#target_link_libraries(tkDNN_static ${tkdnn_LIBS}) target_link_libraries(tkDNN_static ${tkdnn_LIBS})
# SMALL NETS # SMALL NETS
add_executable(test_simple tests/simple/test_simple.cpp) add_executable(test_simple tests/simple/test_simple.cpp)
+2 -2
View File
@@ -17,7 +17,7 @@
#include "tkdnn.h" #include "tkdnn.h"
//#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib. #define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
#ifdef OPENCV_CUDACONTRIB #ifdef OPENCV_CUDACONTRIB
#include <opencv2/cudawarping.hpp> #include <opencv2/cudawarping.hpp>
@@ -37,7 +37,7 @@ class DetectionNN {
cv::Scalar colors[256]; cv::Scalar colors[256];
int nBatches = 1; int nBatches = 2;
#ifdef OPENCV_CUDACONTRIB #ifdef OPENCV_CUDACONTRIB
cv::cuda::GpuMat bgr[3]; cv::cuda::GpuMat bgr[3];
+1 -1
View File
@@ -17,7 +17,7 @@
#include "tkdnn.h" #include "tkdnn.h"
// #define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib. #define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
#ifdef OPENCV_CUDACONTRIB #ifdef OPENCV_CUDACONTRIB
#include <opencv2/cudawarping.hpp> #include <opencv2/cudawarping.hpp>
+2 -2
View File
@@ -39,7 +39,7 @@ public:
float *getLabels() { return mLabels.data(); } float *getLabels() { return mLabels.data(); }
int getBatchesRead() const { return mBatchCount; } int getBatchesRead() const { return mBatchCount; }
int getBatchSize() const { return mBatchSize; } int getBatchSize() const { return mBatchSize; }
nvinfer1::DimsNCHW getDims() const { return mDims; } nvinfer1::Dims4 getDims() const { return mDims; }
float* getFileBatch() { return &mFileBatch[0]; } float* getFileBatch() { return &mFileBatch[0]; }
float* getFileLabels() { return &mFileLabels[0]; } float* getFileLabels() { return &mFileLabels[0]; }
void readInListFile(const std::string& dataFilePath, std::vector<std::string>& mListIn); void readInListFile(const std::string& dataFilePath, std::vector<std::string>& mListIn);
@@ -55,7 +55,7 @@ private:
int mFileBatchPos{ 0 }; int mFileBatchPos{ 0 };
int mImageSize{ 0 }; int mImageSize{ 0 };
nvinfer1::DimsNCHW mDims; nvinfer1::Dims4 mDims;
std::vector<float> mBatch; std::vector<float> mBatch;
std::vector<float> mLabels; std::vector<float> mLabels;
std::vector<float> mFileBatch; std::vector<float> mFileBatch;
+4 -4
View File
@@ -30,10 +30,10 @@ public:
Int8EntropyCalibrator(BatchStream& stream, int firstBatch, const std::string& calibTableFilePath, Int8EntropyCalibrator(BatchStream& stream, int firstBatch, const std::string& calibTableFilePath,
const std::string& inputBlobName, bool readCache = true); const std::string& inputBlobName, bool readCache = true);
virtual ~Int8EntropyCalibrator() { checkCuda(cudaFree(mDeviceInput)); } virtual ~Int8EntropyCalibrator() { checkCuda(cudaFree(mDeviceInput)); }
int getBatchSize() const override { return mStream.getBatchSize(); } int getBatchSize() const noexcept override { return mStream.getBatchSize(); }
bool getBatch(void* bindings[], const char* names[], int nbBindings) override; bool getBatch(void* bindings[], const char* names[], int nbBindings) noexcept override;
const void* readCalibrationCache(size_t& length) override; const void* readCalibrationCache(size_t& length) noexcept override;
void writeCalibrationCache(const void* cache, size_t length) override; void writeCalibrationCache(const void* cache, size_t length) noexcept override;
private: private:
BatchStream mStream; BatchStream mStream;
+4 -29
View File
@@ -2,28 +2,14 @@
#define NETWORKRT_H #define NETWORKRT_H
#include <string.h> // memcpy #include <string.h> // memcpy
#include <memory>
#include "utils.h" #include "utils.h"
#include "Network.h" #include "Network.h"
#include "Layer.h" #include "Layer.h"
#include "NvInfer.h" #include "NvInfer.h"
#include <memory>
namespace tk { namespace dnn { // using namespace nvinfer1;
template<typename T> void writeBUF(char*& buffer, const T& val)
{
*reinterpret_cast<T*>(buffer) = val;
buffer += sizeof(T);
}
template<typename T> T readBUF(const char*& buffer)
{
T val = *reinterpret_cast<const T*>(buffer);
buffer += sizeof(T);
return val;
}
using namespace nvinfer1;
#include "pluginsRT/ActivationLeakyRT.h" #include "pluginsRT/ActivationLeakyRT.h"
#include "pluginsRT/ActivationLogisticRT.h" #include "pluginsRT/ActivationLogisticRT.h"
#include "pluginsRT/ActivationReLUCeilingRT.h" #include "pluginsRT/ActivationReLUCeilingRT.h"
@@ -40,16 +26,7 @@ using namespace nvinfer1;
#include "pluginsRT/ReshapeRT.h" #include "pluginsRT/ReshapeRT.h"
#include "pluginsRT/MaxPoolingFixedSizeRT.h" #include "pluginsRT/MaxPoolingFixedSizeRT.h"
class PluginFactory : IPluginFactory namespace tk { namespace dnn {
{
public:
YoloRT *yolos[16];
int n_yolos;
virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength);
};
class NetworkRT { class NetworkRT {
@@ -74,8 +51,6 @@ public:
dnnType *output; dnnType *output;
cudaStream_t stream; cudaStream_t stream;
PluginFactory *pluginFactory;
NetworkRT(Network *net, const char *name); NetworkRT(Network *net, const char *name);
virtual ~NetworkRT(); virtual ~NetworkRT();
+20
View File
@@ -0,0 +1,20 @@
#ifndef BUFFER_FUNC_H
#define BUFFER_FUNC_H
namespace tk { namespace dnn {
template<typename T> void writeBUF(char*& buffer, const T& val)
{
*reinterpret_cast<T*>(buffer) = val;
buffer += sizeof(T);
}
template<typename T> T readBUF(const char*& buffer)
{
T val = *reinterpret_cast<const T*>(buffer);
buffer += sizeof(T);
return val;
}
}}
#endif // BUFFER_FUNC_H
+26 -9
View File
@@ -1,7 +1,19 @@
#include<cassert> #ifndef ACTIVATION_LEAKY_RT_H
#include "../kernels.h" #define ACTIVATION_LEAKY_RT_H
class ActivationLeakyRT : public IPlugin { #if NV_TENSORRT_MAJOR < 6
#include <cassert>
#include <vector>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
namespace tk { namespace dnn {
class ActivationLeakyRT final : public IPlugin {
public: public:
ActivationLeakyRT(float s) { ActivationLeakyRT(float s) {
@@ -31,14 +43,14 @@ public:
return 0; return 0;
} }
virtual void terminate() override { void terminate() override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
activationLEAKYForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), activationLEAKYForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, slope, stream); reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, slope, stream);
@@ -46,16 +58,21 @@ public:
} }
virtual size_t getSerializationSize() override { size_t getSerializationSize() override {
return 1*sizeof(int) + 1*sizeof(float); return 1*sizeof(int) + 1*sizeof(float);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, size); writeBUF(buf, size);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
int size; int size;
float slope; float slope;
}; };
}}
#endif
#endif // ACTIVATION_LEAKY_RT_H
+112 -23
View File
@@ -1,60 +1,149 @@
#include<cassert> #ifndef ACTIVATION_LOGISTIC_RT_H
#include "../kernels.h" #define ACTIVATION_LOGISTIC_RT_H
class ActivationLogisticRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "ActivationLogistic"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class ActivationLogisticRT final : public nvinfer1::IPluginV2 {
public: public:
ActivationLogisticRT() { ActivationLogisticRT() = default;
~ActivationLogisticRT() = default;
} int getNbOutputs() const noexcept override {
~ActivationLogisticRT(){
}
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return inputs[0]; return inputs[0];
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
size = 1; size = 1;
for(int i=0; i<outputDims[0].nbDims; i++) for(int i=0; i<outputDims[0].nbDims; i++)
size *= outputDims[0].d[i]; size *= outputDims[0].d[i];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
activationLOGISTICForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), activationLOGISTICForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream); reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 1*sizeof(int); return 1*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer); char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, size); writeBUF(buf, size);
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new ActivationLogisticRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int size; int size;
}; };
class ActivationLogisticRTCreator final : public nvinfer1::IPluginCreator {
public:
ActivationLogisticRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
ActivationLogisticRT *a = new ActivationLogisticRT();
a->size = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return a;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // ACTIVATION_LOGISTIC_RT_H
+111 -22
View File
@@ -1,61 +1,150 @@
#include<cassert> #ifndef ACTIVATION_MISH_RT_H
#include "../kernels.h" #define ACTIVATION_MISH_RT_H
class ActivationMishRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "ActivationMish"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class ActivationMishRT final : public nvinfer1::IPluginV2 {
public: public:
ActivationMishRT() { ActivationMishRT() = default;
~ActivationMishRT() = default;
} int getNbOutputs() const noexcept override {
~ActivationMishRT(){
}
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return inputs[0]; return inputs[0];
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
size = 1; size = 1;
for(int i=0; i<outputDims[0].nbDims; i++) for(int i=0; i<outputDims[0].nbDims; i++)
size *= outputDims[0].d[i]; size *= outputDims[0].d[i];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
activationMishForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), activationMishForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream); reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 1*sizeof(int); return 1*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, size); writeBUF(buf, size);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new ActivationMishRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int size; int size;
}; };
class ActivationMishRTCreator final : public nvinfer1::IPluginCreator {
public:
ActivationMishRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
ActivationMishRT *a = new ActivationMishRT();
a->size = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return a;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // ACTIVATION_MISH_RT_H
+112 -19
View File
@@ -1,63 +1,156 @@
#include<cassert> #ifndef ACTIVATION_RELU_CEILING_RT_H
#include "../kernels.h" #define ACTIVATION_RELU_CEILING_RT_H
class ActivationReLUCeiling : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "ActivationCReLU"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class ActivationReLUCeiling final : public nvinfer1::IPluginV2 {
public: public:
ActivationReLUCeiling(const float ceiling) { ActivationReLUCeiling(const float ceiling) {
this->ceiling = ceiling; this->ceiling = ceiling;
} }
~ActivationReLUCeiling(){ ~ActivationReLUCeiling() = default;
} int getNbOutputs() const noexcept override {
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return inputs[0]; return inputs[0];
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
size = 1; size = 1;
for(int i=0; i<outputDims[0].nbDims; i++) for(int i=0; i<outputDims[0].nbDims; i++)
size *= outputDims[0].d[i]; size *= outputDims[0].d[i];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
activationReLUCeilingForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), activationReLUCeilingForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, ceiling, stream); reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, ceiling, stream);
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 1*sizeof(int) + 1*sizeof(float); return 1*sizeof(int) + 1*sizeof(float);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, ceiling); writeBUF(buf, ceiling);
tk::dnn::writeBUF(buf, size); writeBUF(buf, size);
assert(buf = a + getSerializationSize()); assert(buf = a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new ActivationReLUCeiling(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int size; int size;
float ceiling; float ceiling;
}; };
class ActivationReLUCeilingCreator final : public nvinfer1::IPluginCreator {
public:
ActivationReLUCeilingCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
float activationReluTemp = readBUF<float>(buf);
ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp);
a->size = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return a;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // ACTIVATION_RELU_CEILING_RT_H
+111 -23
View File
@@ -1,61 +1,149 @@
#include<cassert> #ifndef ACTIVATION_SIGMOID_RT_H
#include "../kernels.h" #define ACTIVATION_SIGMOID_RT_H
class ActivationSigmoidRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "ActivationSigmoidRT"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class ActivationSigmoidRT final : public nvinfer1::IPluginV2 {
public: public:
ActivationSigmoidRT() { ActivationSigmoidRT() = default;
~ActivationSigmoidRT() = default;
} int getNbOutputs() const noexcept override {
~ActivationSigmoidRT(){
}
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return inputs[0]; return inputs[0];
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
size = 1; size = 1;
for(int i=0; i<outputDims[0].nbDims; i++) for(int i=0; i<outputDims[0].nbDims; i++)
size *= outputDims[0].d[i]; size *= outputDims[0].d[i];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
activationSIGMOIDForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), activationSIGMOIDForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream); reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 1*sizeof(int); return 1*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, size); writeBUF(buf, size);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new ActivationSigmoidRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int size; int size;
}; };
class ActivationSigmoidRTCreator final : public nvinfer1::IPluginCreator {
public:
ActivationSigmoidRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
ActivationSigmoidRT* a = new ActivationSigmoidRT();
a->size = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return a;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // ACTIVATION_SIGMOID_RT_H
+179 -39
View File
@@ -1,16 +1,27 @@
#ifndef DEFORMABLE_CONV_RT_H
#define DEFORMABLE_CONV_RT_H
#include <cassert> #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h" #include "../kernels.h"
#include "../buffer_func.h"
#include "../Layer.h"
#define PLUGIN_NAME "Deformable"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class DeformableConvRT : public IPlugin { class DeformableConvRT final : public nvinfer1::IPluginV2 {
public: public:
DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw, 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 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, int o_n, int o_c, int o_h, int o_w,
tk::dnn::DeformConv2d *deformable = nullptr) { DeformConv2d *deformable = nullptr) {
this->chunk_dim = chunk_dim; this->chunk_dim = chunk_dim;
this->kh = kh; this->kh = kh;
this->kw = kw; this->kw = kw;
@@ -61,27 +72,34 @@ public:
cublasDestroy(handle); cublasDestroy(handle);
} }
int getNbOutputs() const override { int getNbOutputs() const noexcept override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w}; return nvinfer1::Dims3{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w};
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { } void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
}
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { } void terminate() noexcept override { }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]); dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]); dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
@@ -109,63 +127,92 @@ public:
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 16 * sizeof(int) + chunk_dim * 3 * sizeof(dnnType) + (i_c * o_c * kh * kw * 1 ) * sizeof(dnnType) + 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); o_c * sizeof(dnnType) + height_ones * width_ones * sizeof(dnnType) + dim_ones * sizeof(dnnType);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, chunk_dim); writeBUF(buf, chunk_dim);
tk::dnn::writeBUF(buf, kh); writeBUF(buf, kh);
tk::dnn::writeBUF(buf, kw); writeBUF(buf, kw);
tk::dnn::writeBUF(buf, sh); writeBUF(buf, sh);
tk::dnn::writeBUF(buf, sw); writeBUF(buf, sw);
tk::dnn::writeBUF(buf, ph); writeBUF(buf, ph);
tk::dnn::writeBUF(buf, pw); writeBUF(buf, pw);
tk::dnn::writeBUF(buf, deformableGroup); writeBUF(buf, deformableGroup);
tk::dnn::writeBUF(buf, i_n); writeBUF(buf, i_n);
tk::dnn::writeBUF(buf, i_c); writeBUF(buf, i_c);
tk::dnn::writeBUF(buf, i_h); writeBUF(buf, i_h);
tk::dnn::writeBUF(buf, i_w); writeBUF(buf, i_w);
tk::dnn::writeBUF(buf, o_n); writeBUF(buf, o_n);
tk::dnn::writeBUF(buf, o_c); writeBUF(buf, o_c);
tk::dnn::writeBUF(buf, o_h); writeBUF(buf, o_h);
tk::dnn::writeBUF(buf, o_w); writeBUF(buf, o_w);
dnnType *aus = new dnnType[chunk_dim*2]; dnnType *aus = new dnnType[chunk_dim*2];
checkCuda( cudaMemcpy(aus, offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToHost) ); checkCuda( cudaMemcpy(aus, offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToHost) );
for(int i=0; i<chunk_dim*2; i++) for(int i=0; i<chunk_dim*2; i++)
tk::dnn::writeBUF(buf, aus[i]); writeBUF(buf, aus[i]);
free(aus); free(aus);
aus = new dnnType[chunk_dim]; aus = new dnnType[chunk_dim];
checkCuda( cudaMemcpy(aus, mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToHost) ); checkCuda( cudaMemcpy(aus, mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToHost) );
for(int i=0; i<chunk_dim; i++) for(int i=0; i<chunk_dim; i++)
tk::dnn::writeBUF(buf, aus[i]); writeBUF(buf, aus[i]);
free(aus); free(aus);
aus = new dnnType[(i_c * o_c * kh * kw * 1 )]; aus = new dnnType[(i_c * o_c * kh * kw * 1 )];
checkCuda( cudaMemcpy(aus, data_d, sizeof(dnnType)*(i_c * o_c * kh * kw * 1 ), cudaMemcpyDeviceToHost) ); checkCuda( cudaMemcpy(aus, data_d, sizeof(dnnType)*(i_c * o_c * kh * kw * 1 ), cudaMemcpyDeviceToHost) );
for(int i=0; i<(i_c * o_c * kh * kw * 1 ); i++) for(int i=0; i<(i_c * o_c * kh * kw * 1 ); i++)
tk::dnn::writeBUF(buf, aus[i]); writeBUF(buf, aus[i]);
free(aus); free(aus);
aus = new dnnType[o_c]; aus = new dnnType[o_c];
checkCuda( cudaMemcpy(aus, bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToHost) ); checkCuda( cudaMemcpy(aus, bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToHost) );
for(int i=0; i < o_c; i++) for(int i=0; i < o_c; i++)
tk::dnn::writeBUF(buf, aus[i]); writeBUF(buf, aus[i]);
free(aus); free(aus);
aus = new dnnType[height_ones * width_ones]; aus = new dnnType[height_ones * width_ones];
checkCuda( cudaMemcpy(aus, ones_d1, sizeof(dnnType)*height_ones * width_ones, cudaMemcpyDeviceToHost) ); checkCuda( cudaMemcpy(aus, ones_d1, sizeof(dnnType)*height_ones * width_ones, cudaMemcpyDeviceToHost) );
for(int i=0; i<height_ones * width_ones; i++) for(int i=0; i<height_ones * width_ones; i++)
tk::dnn::writeBUF(buf, aus[i]); writeBUF(buf, aus[i]);
free(aus); free(aus);
aus = new dnnType[dim_ones]; aus = new dnnType[dim_ones];
checkCuda( cudaMemcpy(aus, ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToHost) ); checkCuda( cudaMemcpy(aus, ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToHost) );
for(int i=0; i<dim_ones; i++) for(int i=0; i<dim_ones; i++)
tk::dnn::writeBUF(buf, aus[i]); writeBUF(buf, aus[i]);
free(aus); free(aus);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new DeformableConvRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
cublasStatus_t stat; cublasStatus_t stat;
cublasHandle_t handle; cublasHandle_t handle;
int i_n, i_c, i_h, i_w; int i_n, i_c, i_h, i_w;
@@ -191,6 +238,99 @@ public:
// dnnType *mask_n; // dnnType *mask_n;
// dnnType *output_n; // dnnType *output_n;
DeformConv2d *defRT;
tk::dnn::DeformConv2d *defRT;
}; };
class DeformableConvRTCreator final : public nvinfer1::IPluginCreator {
public:
DeformableConvRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
int chuck_dimTemp = readBUF<int>(buf);
int khTemp = readBUF<int>(buf);
int kwTemp = readBUF<int>(buf);
int shTemp = readBUF<int>(buf);
int swTemp = readBUF<int>(buf);
int phTemp = readBUF<int>(buf);
int pwTemp = readBUF<int>(buf);
int deformableGroupTemp = readBUF<int>(buf);
int i_nTemp = readBUF<int>(buf);
int i_cTemp = readBUF<int>(buf);
int i_hTemp = readBUF<int>(buf);
int i_wTemp = readBUF<int>(buf);
int o_nTemp = readBUF<int>(buf);
int o_cTemp = readBUF<int>(buf);
int o_hTemp = readBUF<int>(buf);
int o_wTemp = readBUF<int>(buf);
DeformableConvRT* r = new DeformableConvRT(chuck_dimTemp, khTemp, kwTemp, shTemp, swTemp, phTemp, pwTemp, deformableGroupTemp, i_nTemp, i_cTemp, i_hTemp, i_wTemp, o_nTemp, o_cTemp, o_hTemp, o_wTemp, nullptr);
dnnType *aus = new dnnType[r->chunk_dim*2];
for(int i=0; i<r->chunk_dim*2; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->offset, aus, sizeof(dnnType)*2*r->chunk_dim, cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[r->chunk_dim];
for(int i=0; i<r->chunk_dim; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->mask, aus, sizeof(dnnType)*r->chunk_dim, cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[(r->i_c * r->o_c * r->kh * r->kw * 1 )];
for(int i=0; i<(r->i_c * r->o_c * r->kh * r->kw * 1 ); i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->data_d, aus, sizeof(dnnType)*(r->i_c * r->o_c * r->kh * r->kw * 1 ), cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[r->o_c];
for(int i=0; i < r->o_c; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->bias2_d, aus, sizeof(dnnType)*r->o_c, cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[r->height_ones * r->width_ones];
for(int i=0; i<r->height_ones * r->width_ones; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->ones_d1, aus, sizeof(dnnType)*r->height_ones * r->width_ones, cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[r->dim_ones];
for(int i=0; i<r->dim_ones; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->ones_d2, aus, sizeof(dnnType)*r->dim_ones, cudaMemcpyHostToDevice) );
free(aus);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // DEFORMABLE_CONV_RT_H
+121 -21
View File
@@ -1,6 +1,20 @@
#include<cassert> #ifndef FLATTEN_CONCAT_RT_H
#define FLATTEN_CONCAT_RT_H
class FlattenConcatRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "Flatten"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class FlattenConcatRT final : public nvinfer1::IPluginV2 {
public: public:
FlattenConcatRT() { FlattenConcatRT() {
@@ -11,19 +25,23 @@ public:
} }
} }
~FlattenConcatRT(){ ~FlattenConcatRT() = default;
} int getNbOutputs() const noexcept override {
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return DimsCHW{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1}; return nvinfer1::Dims3{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1};
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
assert(nbOutputs == 1 && nbInputs ==1); assert(nbOutputs == 1 && nbInputs ==1);
rows = inputDims[0].d[0]; rows = inputDims[0].d[0];
cols = inputDims[0].d[1] * inputDims[0].d[2]; cols = inputDims[0].d[1] * inputDims[0].d[2];
@@ -32,19 +50,19 @@ public:
w = 1; w = 1;
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
checkERROR(cublasDestroy(handle)); checkERROR(cublasDestroy(handle));
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]); dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
@@ -59,23 +77,105 @@ public:
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 5*sizeof(int); return 5*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a = buf; char *buf = reinterpret_cast<char*>(buffer),*a = buf;
tk::dnn::writeBUF(buf, c); writeBUF(buf, c);
tk::dnn::writeBUF(buf, h); writeBUF(buf, h);
tk::dnn::writeBUF(buf, w); writeBUF(buf, w);
tk::dnn::writeBUF(buf, rows); writeBUF(buf, rows);
tk::dnn::writeBUF(buf, cols); writeBUF(buf, cols);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new FlattenConcatRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int c, h, w; int c, h, w;
int rows, cols; int rows, cols;
cublasStatus_t stat; cublasStatus_t stat;
cublasHandle_t handle; cublasHandle_t handle;
}; };
class FlattenConcatRTCreator final : public nvinfer1::IPluginCreator {
public:
FlattenConcatRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
float activationReluTemp = readBUF<float>(buf);
FlattenConcatRT *r = new FlattenConcatRT();
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
r->rows = readBUF<int>(buf);
r->cols = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // FLATTEN_CONCAT_RT_H
+127 -25
View File
@@ -1,7 +1,20 @@
#include<cassert> #ifndef MAX_POOLING_FIXED_SIZE_RT_H
#include "../kernels.h" #define MAX_POOLING_FIXED_SIZE_RT_H
class MaxPoolFixedSizeRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "Pooling"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class MaxPoolFixedSizeRT final : public nvinfer1::IPluginV2 {
public: public:
MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) { MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) {
@@ -15,33 +28,37 @@ public:
this->padding = padding; this->padding = padding;
} }
~MaxPoolFixedSizeRT(){ ~MaxPoolFixedSizeRT() = default;
}
int getNbOutputs() const override { int getNbOutputs() const noexcept override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return DimsCHW{this->c, this->h, this->w}; return nvinfer1::Dims3{this->c, this->h, this->w};
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
//std::cout<<this->n<<" "<<this->c<<" "<<this->h<<" "<<this->w<<" "<<this->stride_H<<" "<<this->stride_W<<" "<<this->winSize<<" "<<this->padding<<std::endl; //std::cout<<this->n<<" "<<this->c<<" "<<this->h<<" "<<this->w<<" "<<this->stride_H<<" "<<this->stride_W<<" "<<this->winSize<<" "<<this->padding<<std::endl;
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]); dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
@@ -49,27 +66,112 @@ public:
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 8*sizeof(int); return 8*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, this->c); writeBUF(buf, this->c);
tk::dnn::writeBUF(buf, this->h); writeBUF(buf, this->h);
tk::dnn::writeBUF(buf, this->w); writeBUF(buf, this->w);
tk::dnn::writeBUF(buf, this->n); writeBUF(buf, this->n);
tk::dnn::writeBUF(buf, this->stride_H); writeBUF(buf, this->stride_H);
tk::dnn::writeBUF(buf, this->stride_W); writeBUF(buf, this->stride_W);
tk::dnn::writeBUF(buf, this->winSize); writeBUF(buf, this->winSize);
tk::dnn::writeBUF(buf, this->padding); writeBUF(buf, this->padding);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new MaxPoolFixedSizeRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int n, c, h, w; int n, c, h, w;
int stride_H, stride_W; int stride_H, stride_W;
int winSize; int winSize;
int padding; int padding;
}; };
class MaxPoolFixedSizeRTCreator final : public nvinfer1::IPluginCreator {
public:
MaxPoolFixedSizeRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
int cTemp = readBUF<int>(buf);
int hTemp = readBUF<int>(buf);
int wTemp = readBUF<int>(buf);
int nTemp = readBUF<int>(buf);
int strideHTemp = readBUF<int>(buf);
int strideWTemp = readBUF<int>(buf);
int winSizeTemp = readBUF<int>(buf);
int paddingTemp = readBUF<int>(buf);
MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // MAX_POOLING_FIXED_SIZE_RT_H
+122 -24
View File
@@ -1,48 +1,62 @@
#include<cassert> #ifndef REGION_RT_H
#include "../kernels.h" #define REGION_RT_H
class RegionRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "Region"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class RegionRT final : public nvinfer1::IPluginV2 {
public: public:
RegionRT(int classes, int coords, int num) { RegionRT(int classes, int coords, int num) {
this->classes = classes; this->classes = classes;
this->coords = coords; this->coords = coords;
this->num = num; this->num = num;
} }
~RegionRT(){ ~RegionRT() = default;
} int getNbOutputs() const noexcept override {
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return inputs[0]; return inputs[0];
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
c = inputDims[0].d[0]; c = inputDims[0].d[0];
h = inputDims[0].d[1]; h = inputDims[0].d[1];
w = inputDims[0].d[2]; w = inputDims[0].d[2];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]); dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
@@ -68,21 +82,51 @@ public:
} }
virtual size_t getSerializationSize() override { size_t getSerializationSize() const noexcept override {
return 6*sizeof(int); return 6*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, classes); writeBUF(buf, classes);
tk::dnn::writeBUF(buf, coords); writeBUF(buf, coords);
tk::dnn::writeBUF(buf, num); writeBUF(buf, num);
tk::dnn::writeBUF(buf, c); writeBUF(buf, c);
tk::dnn::writeBUF(buf, h); writeBUF(buf, h);
tk::dnn::writeBUF(buf, w); writeBUF(buf, w);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new RegionRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int c, h, w; int c, h, w;
int classes, coords, num; int classes, coords, num;
@@ -93,3 +137,57 @@ public:
} }
}; };
class RegionRTCreator final : public nvinfer1::IPluginCreator {
public:
RegionRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
int classesTemp = readBUF<int>(buf);
int coordsTemp = readBUF<int>(buf);
int numTemp = readBUF<int>(buf);
RegionRT* r = new RegionRT(classesTemp, coordsTemp, numTemp);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // REGION_RT_H
+118 -23
View File
@@ -1,64 +1,159 @@
#include<cassert> #ifndef REORG_RT_H
#include "../kernels.h" #define REORG_RT_H
class ReorgRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "Reorg"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class ReorgRT final : public nvinfer1::IPluginV2 {
public: public:
ReorgRT(int stride) { ReorgRT(int stride) {
this->stride = stride; this->stride = stride;
} }
~ReorgRT(){ ~ReorgRT() = default;
} int getNbOutputs() const noexcept override {
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return DimsCHW{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride}; return nvinfer1::Dims3{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride};
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
c = inputDims[0].d[0]; c = inputDims[0].d[0];
h = inputDims[0].d[1]; h = inputDims[0].d[1];
w = inputDims[0].d[2]; w = inputDims[0].d[2];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
reorgForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), reorgForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), reinterpret_cast<dnnType*>(outputs[0]),
batchSize, c, h, w, stride, stream); batchSize, c, h, w, stride, stream);
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 4*sizeof(int); return 4*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, stride); writeBUF(buf, stride);
tk::dnn::writeBUF(buf, c); writeBUF(buf, c);
tk::dnn::writeBUF(buf, h); writeBUF(buf, h);
tk::dnn::writeBUF(buf, w); writeBUF(buf, w);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new ReorgRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int c, h, w, stride; int c, h, w, stride;
}; };
class ReorgRTCreator final : public nvinfer1::IPluginCreator {
public:
ReorgRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
int strideTemp = readBUF<int>(buf);
ReorgRT *r = new ReorgRT(strideTemp);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // REORG_RT_H
+120 -20
View File
@@ -1,6 +1,21 @@
#include<cassert> #ifndef RESHAPE_RT_H
#define RESHAPE_RT_H
class ReshapeRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#include "../Network.h"
#define PLUGIN_NAME "Reshape"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class ReshapeRT final : public nvinfer1::IPluginV2 {
public: public:
ReshapeRT(dataDim_t new_dim) { ReshapeRT(dataDim_t new_dim) {
@@ -10,33 +25,37 @@ public:
w = new_dim.w; w = new_dim.w;
} }
~ReshapeRT(){ ~ReshapeRT() = default;
} int getNbOutputs() const noexcept override {
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return DimsCHW{ c,h,w}; return nvinfer1::Dims3{ c,h,w};
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { virtual void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { virtual size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) 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<const dnnType*>(inputs[0]); dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
@@ -44,19 +63,100 @@ public:
return 0; return 0;
} }
virtual size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 4*sizeof(int); return 4*sizeof(int);
} }
virtual void serialize(void* buffer) override { virtual void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a = buf; char *buf = reinterpret_cast<char*>(buffer),*a = buf;
tk::dnn::writeBUF(buf, n); writeBUF(buf, n);
tk::dnn::writeBUF(buf, c); writeBUF(buf, c);
tk::dnn::writeBUF(buf, h); writeBUF(buf, h);
tk::dnn::writeBUF(buf, w); writeBUF(buf, w);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new ReshapeRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int n, c, h, w; int n, c, h, w;
}; };
class ReshapeRTCreator final : public nvinfer1::IPluginCreator {
public:
ReshapeRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
dataDim_t new_dim;
new_dim.n = readBUF<int>(buf);
new_dim.c = readBUF<int>(buf);
new_dim.h = readBUF<int>(buf);
new_dim.w = readBUF<int>(buf);
ReshapeRT *r = new ReshapeRT(new_dim);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // RESHAPE_RT_H
+122 -22
View File
@@ -1,7 +1,20 @@
#include<cassert> #ifndef RESIZE_LAYER_RT_H
#include "../kernels.h" #define RESIZE_LAYER_RT_H
class ResizeLayerRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "Resize"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class ResizeLayerRT final : public nvinfer1::IPluginV2 {
public: public:
ResizeLayerRT(int c, int h, int w) { ResizeLayerRT(int c, int h, int w) {
@@ -10,35 +23,40 @@ public:
o_w = w; o_w = w;
} }
~ResizeLayerRT(){ ~ResizeLayerRT() = default;
}
int getNbOutputs() const override { int getNbOutputs() const noexcept override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return DimsCHW{o_c, o_h, o_w}; return nvinfer1::Dims3{o_c, o_h, o_w};
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
i_c = inputDims[0].d[0]; i_c = inputDims[0].d[0];
i_h = inputDims[0].d[1]; i_h = inputDims[0].d[1];
i_w = inputDims[0].d[2]; i_w = inputDims[0].d[2];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { 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); // printf("%d %d %d %d %d %d\n", i_c, i_w, i_h, o_c, o_w, o_h);
resizeForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), resizeForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), reinterpret_cast<dnnType*>(outputs[0]),
@@ -46,23 +64,105 @@ public:
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 6*sizeof(int); return 6*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, o_c); writeBUF(buf, o_c);
tk::dnn::writeBUF(buf, o_h); writeBUF(buf, o_h);
tk::dnn::writeBUF(buf, o_w); writeBUF(buf, o_w);
tk::dnn::writeBUF(buf, i_c); writeBUF(buf, i_c);
tk::dnn::writeBUF(buf, i_h); writeBUF(buf, i_h);
tk::dnn::writeBUF(buf, i_w); writeBUF(buf, i_w);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new ResizeLayerRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int i_c, i_h, i_w, o_c, o_h, o_w; int i_c, i_h, i_w, o_c, o_h, o_w;
}; };
class ResizeLayerRTCreator final : public nvinfer1::IPluginCreator {
public:
ResizeLayerRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
int o_cTemp = readBUF<int>(buf);
int o_hTemp = readBUF<int>(buf);
int o_wTemp = readBUF<int>(buf);
ResizeLayerRT* r = new ResizeLayerRT(o_cTemp, o_hTemp, o_wTemp);
r->i_c = readBUF<int>(buf);
r->i_h = readBUF<int>(buf);
r->i_w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // RESIZE_LAYER_RT_H
+125 -25
View File
@@ -1,7 +1,20 @@
#include<cassert> #ifndef ROUTE_RT_H
#include "../kernels.h" #define ROUTE_RT_H
class RouteRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "Route"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class RouteRT final : public nvinfer1::IPluginV2 {
/** /**
THIS IS NOT USED ANYMORE THIS IS NOT USED ANYMORE
@@ -13,21 +26,25 @@ public:
this->group_id = group_id; this->group_id = group_id;
} }
~RouteRT(){ ~RouteRT() = default;
} int getNbOutputs() const noexcept override {
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
int out_c = 0; int out_c = 0;
for(int i=0; i<nbInputDims; i++) out_c += inputs[i].d[0]; for(int i=0; i<nbInputDims; i++) out_c += inputs[i].d[0];
return DimsCHW{out_c/groups, inputs[0].d[1], inputs[0].d[2]}; return nvinfer1::Dims3{out_c/groups, inputs[0].d[1], inputs[0].d[2]};
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
in = nbInputs; in = nbInputs;
c = 0; c = 0;
for(int i=0; i<nbInputs; i++) { for(int i=0; i<nbInputs; i++) {
@@ -39,20 +56,18 @@ public:
c /= groups; c /= groups;
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
for(int b=0; b<batchSize; b++) { for(int b=0; b<batchSize; b++) {
@@ -70,27 +85,112 @@ public:
} }
virtual size_t getSerializationSize() override { size_t getSerializationSize() const noexcept override {
return (6+MAX_INPUTS)*sizeof(int); return (6+MAX_INPUTS)*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, groups); writeBUF(buf, groups);
tk::dnn::writeBUF(buf, group_id); writeBUF(buf, group_id);
tk::dnn::writeBUF(buf, in); writeBUF(buf, in);
for(int i=0; i<MAX_INPUTS; i++) for(int i=0; i<MAX_INPUTS; i++)
tk::dnn::writeBUF(buf, c_in[i]); writeBUF(buf, c_in[i]);
tk::dnn::writeBUF(buf, c); writeBUF(buf, c);
tk::dnn::writeBUF(buf, h); writeBUF(buf, h);
tk::dnn::writeBUF(buf, w); writeBUF(buf, w);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new RouteRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
static const int MAX_INPUTS = 4; static const int MAX_INPUTS = 4;
int in; int in;
int c_in[MAX_INPUTS]; int c_in[MAX_INPUTS];
int c, h, w; int c, h, w;
int groups, group_id; int groups, group_id;
}; };
class RouteRTCreator final : public nvinfer1::IPluginCreator {
public:
RouteRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
int groupsTemp = readBUF<int>(buf);
int group_idTemp = readBUF<int>(buf);
RouteRT* r = new RouteRT(groupsTemp, group_idTemp);
r->in = readBUF<int>(buf);
for(int i=0; i<RouteRT::MAX_INPUTS; i++)
r->c_in[i] = readBUF<int>(buf);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // ROUTE_RT_H
+128 -26
View File
@@ -1,48 +1,64 @@
#include<cassert> #ifndef SHORTCUT_RT_H
#include "../kernels.h" #define SHORTCUT_RT_H
class ShortcutRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#include "../Network.h"
#define PLUGIN_NAME "Shortcut"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class ShortcutRT final : public nvinfer1::IPluginV2 {
public: public:
ShortcutRT(tk::dnn::dataDim_t bdim, bool mul) { ShortcutRT(dataDim_t bdim, bool mul) {
this->bc = bdim.c; this->bc = bdim.c;
this->bh = bdim.h; this->bh = bdim.h;
this->bw = bdim.w; this->bw = bdim.w;
this->mul = mul; this->mul = mul;
} }
~ShortcutRT(){ ~ShortcutRT() = default;
} int getNbOutputs() const noexcept override {
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return DimsCHW{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]}; return nvinfer1::Dims3{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]};
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
c = inputDims[0].d[0]; c = inputDims[0].d[0];
h = inputDims[0].d[1]; h = inputDims[0].d[1];
w = inputDims[0].d[2]; w = inputDims[0].d[2];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]); dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *srcDataBack = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]); dnnType *srcDataBack = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
@@ -54,24 +70,110 @@ public:
} }
virtual size_t getSerializationSize() override { size_t getSerializationSize() const noexcept override {
return 6*sizeof(int) + sizeof(bool); return 6*sizeof(int) + sizeof(bool);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, bc); writeBUF(buf, bc);
tk::dnn::writeBUF(buf, bh); writeBUF(buf, bh);
tk::dnn::writeBUF(buf, bw); writeBUF(buf, bw);
tk::dnn::writeBUF(buf, mul); writeBUF(buf, mul);
tk::dnn::writeBUF(buf, c); writeBUF(buf, c);
tk::dnn::writeBUF(buf, h); writeBUF(buf, h);
tk::dnn::writeBUF(buf, w); writeBUF(buf, w);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new ShortcutRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int c, h, w; int c, h, w;
int bc, bh, bw; int bc, bh, bw;
bool mul; bool mul;
}; };
class ShortcutRTCreator final : public nvinfer1::IPluginCreator {
public:
ShortcutRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
dataDim_t bdim;
bdim.c = readBUF<int>(buf);
bdim.h = readBUF<int>(buf);
bdim.w = readBUF<int>(buf);
bdim.l = 1;
ShortcutRT *r = new ShortcutRT(bdim, readBUF<bool>(buf));
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // SHORTCUT_RT_H
+117 -23
View File
@@ -1,45 +1,60 @@
#include<cassert> #ifndef UPSAMPLE_RT_H
#include "../kernels.h" #define UPSAMPLE_RT_H
class UpsampleRT : public IPlugin { #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../kernels.h"
#include "../buffer_func.h"
#define PLUGIN_NAME "Upsample"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class UpsampleRT final : public nvinfer1::IPluginV2 {
public: public:
UpsampleRT(int stride) { UpsampleRT(int stride) {
this->stride = stride; this->stride = stride;
} }
~UpsampleRT(){ ~UpsampleRT() = default;
} int getNbOutputs() const noexcept override {
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return DimsCHW(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride); return nvinfer1::Dims3(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride);
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
c = inputDims[0].d[0]; c = inputDims[0].d[0];
h = inputDims[0].d[1]; h = inputDims[0].d[1];
w = inputDims[0].d[2]; w = inputDims[0].d[2];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]); dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
@@ -48,19 +63,98 @@ public:
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 4*sizeof(int); return 4*sizeof(int);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, stride); writeBUF(buf, stride);
tk::dnn::writeBUF(buf, c); writeBUF(buf, c);
tk::dnn::writeBUF(buf, h); writeBUF(buf, h);
tk::dnn::writeBUF(buf, w); writeBUF(buf, w);
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new UpsampleRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int c, h, w, stride; int c, h, w, stride;
}; };
class UpsampleRTCreator final : public nvinfer1::IPluginCreator {
public:
UpsampleRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
int strideTemp = readBUF<int>(buf);
UpsampleRT* r = new UpsampleRT(strideTemp);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength);
return r;
}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // UPSAMPLE_RT_H
+117 -34
View File
@@ -1,15 +1,26 @@
#ifndef YOLO_RT_H
#define YOLO_RT_H
#include <cassert> #include <cassert>
#include <vector>
#include <NvInferRuntimeCommon.h>
#include <NvInfer.h>
#include "../yoloContainer.h"
#include "../kernels.h" #include "../kernels.h"
#include "../buffer_func.h"
#include "../Layer.h"
#define YOLORT_CLASSNAME_W 256 #define YOLORT_CLASSNAME_W 256
class YoloRT : public IPlugin { #define PLUGIN_NAME "Yolo"
#define PLUGIN_VERSION "1"
namespace tk { namespace dnn {
class YoloRT final : public nvinfer1::IPluginV2 {
public: 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) { YoloRT(int classes, int num, 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->classes = classes; this->classes = classes;
this->num = num; this->num = num;
this->n_masks = n_masks; this->n_masks = n_masks;
@@ -27,38 +38,40 @@ public:
} }
} }
~YoloRT(){ ~YoloRT() = default;
} int getNbOutputs() const noexcept override {
int getNbOutputs() const override {
return 1; return 1;
} }
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
return inputs[0]; return inputs[0];
} }
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { void configureWithFormat(nvinfer1::Dims const * inputDims,
int32_t nbInputs,
nvinfer1::Dims const * outputDims,
int32_t nbOutputs,
nvinfer1::DataType type,
nvinfer1::PluginFormat format,
int32_t maxBatchSize) noexcept override {
c = inputDims[0].d[0]; c = inputDims[0].d[0];
h = inputDims[0].d[1]; h = inputDims[0].d[1];
w = inputDims[0].d[2]; w = inputDims[0].d[2];
} }
int initialize() override { int initialize() noexcept override {
return 0; return 0;
} }
virtual void terminate() override { void terminate() noexcept override {
} }
virtual size_t getWorkspaceSize(int maxBatchSize) const override { size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
return 0; return 0;
} }
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]); dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
@@ -86,30 +99,29 @@ public:
return 0; return 0;
} }
size_t getSerializationSize() const noexcept override {
virtual size_t getSerializationSize() override {
return 8*sizeof(int) + 2*sizeof(float)+ n_masks*sizeof(dnnType) + num*n_masks*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char); return 8*sizeof(int) + 2*sizeof(float)+ n_masks*sizeof(dnnType) + num*n_masks*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char);
} }
virtual void serialize(void* buffer) override { void serialize(void* buffer) const noexcept override {
char *buf = reinterpret_cast<char*>(buffer),*a=buf; char *buf = reinterpret_cast<char*>(buffer),*a=buf;
tk::dnn::writeBUF(buf, classes); //std::cout << "Classes :" << classes << std::endl; writeBUF(buf, classes); //std::cout << "Classes :" << classes << std::endl;
tk::dnn::writeBUF(buf, num); //std::cout << "Num : " << num << std::endl; writeBUF(buf, num); //std::cout << "Num : " << num << std::endl;
tk::dnn::writeBUF(buf, n_masks); //std::cout << "N_Masks" << n_masks << std::endl; writeBUF(buf, n_masks); //std::cout << "N_Masks" << n_masks << std::endl;
tk::dnn::writeBUF(buf, scaleXY); //std::cout << "ScaleXY :" << scaleXY << std::endl; writeBUF(buf, scaleXY); //std::cout << "ScaleXY :" << scaleXY << std::endl;
tk::dnn::writeBUF(buf, nms_thresh); //std::cout << "nms_thresh :" << nms_thresh << std::endl; writeBUF(buf, nms_thresh); //std::cout << "nms_thresh :" << nms_thresh << std::endl;
tk::dnn::writeBUF(buf, nms_kind); //std::cout << "nms_kind : " << nms_kind << std::endl; writeBUF(buf, nms_kind); //std::cout << "nms_kind : " << nms_kind << std::endl;
tk::dnn::writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl; writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl;
tk::dnn::writeBUF(buf, c); //std::cout << "C : " << c << std::endl; writeBUF(buf, c); //std::cout << "C : " << c << std::endl;
tk::dnn::writeBUF(buf, h); //std::cout << "H : " << h << std::endl; writeBUF(buf, h); //std::cout << "H : " << h << std::endl;
tk::dnn::writeBUF(buf, w); //std::cout << "C : " << c << std::endl; writeBUF(buf, w); //std::cout << "C : " << c << std::endl;
for (int i = 0; i < n_masks; i++) for (int i = 0; i < n_masks; i++)
{ {
tk::dnn::writeBUF(buf, mask[i]); //std::cout << "mask[i] : " << mask[i] << std::endl; writeBUF(buf, mask[i]); //std::cout << "mask[i] : " << mask[i] << std::endl;
} }
for (int i = 0; i < n_masks * 2 * num; i++) for (int i = 0; i < n_masks * 2 * num; i++)
{ {
tk::dnn::writeBUF(buf, bias[i]); //std::cout << "bias[i] : " << bias[i] << std::endl; writeBUF(buf, bias[i]); //std::cout << "bias[i] : " << bias[i] << std::endl;
} }
// save classes names // save classes names
@@ -117,12 +129,42 @@ public:
char tmp[YOLORT_CLASSNAME_W]; char tmp[YOLORT_CLASSNAME_W];
strcpy(tmp, classesNames[i].c_str()); strcpy(tmp, classesNames[i].c_str());
for(int j=0; j<YOLORT_CLASSNAME_W; j++) { for(int j=0; j<YOLORT_CLASSNAME_W; j++) {
tk::dnn::writeBUF(buf, tmp[j]); writeBUF(buf, tmp[j]);
} }
} }
assert(buf == a + getSerializationSize()); assert(buf == a + getSerializationSize());
} }
// Extra IPluginV2 overrides
bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
}
nvinfer1::IPluginV2 * clone() const noexcept override {
auto a = new YoloRT(*this);
return a;
}
const char* getPluginType() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
void destroy() noexcept override {}
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
std::string mNamespace;
int c, h, w; int c, h, w;
int classes, num, n_masks; int classes, num, n_masks;
float scaleXY; float scaleXY;
@@ -139,5 +181,46 @@ public:
int loc = location % (w*h); int loc = location % (w*h);
return batch*c*h*w + n*w*h*(4+classes+1) + entry*w*h + loc; return batch*c*h*w + n*w*h*(4+classes+1) + entry*w*h + loc;
} }
}; };
class YoloRTCreator final : public nvinfer1::IPluginCreator {
public:
YoloRTCreator() = default;
const char* getPluginName() const noexcept override {
return PLUGIN_NAME;
}
const char* getPluginVersion() const noexcept override {
return PLUGIN_VERSION;
}
const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
return &mFC;
}
nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
std::cout << "Create plugin" << std::endl;
return nullptr;
}
nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override;
void setPluginNamespace(const char* pluginNamespace) noexcept override {
mNamespace = pluginNamespace;
}
const char* getPluginNamespace() const noexcept override {
return mNamespace.c_str();
}
private:
static nvinfer1::PluginFieldCollection mFC;
static std::vector<nvinfer1::PluginField> mPluginAttributes;
std::string mNamespace;
};
}}
#undef PLUGIN_NAME
#undef PLUGIN_VERSION
#endif // YOLO_RT_H
+16
View File
@@ -0,0 +1,16 @@
#ifndef YOLO_CONTAINER_H
#define YOLO_CONTAINER_H
namespace tk { namespace dnn {
class YoloRT;
class YoloContainer
{
public:
YoloRT *yolos[16];
int n_yolos{};
};
extern YoloContainer yoloContainer;
}}
#endif // YOLO_CONTAINER_H
+10 -10
View File
@@ -8,14 +8,14 @@
BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, const std::string& fileimglist, const std::string& filelabellist) { BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, const std::string& fileimglist, const std::string& filelabellist) {
mBatchSize = batchSize; mBatchSize = batchSize;
mMaxBatches = maxBatches; mMaxBatches = maxBatches;
mDims = nvinfer1::DimsNCHW{ dim.n, dim.c, dim.h, dim.w }; mDims = nvinfer1::Dims4{ dim.n, dim.c, dim.h, dim.w };
mHeight = dim.h; mHeight = dim.h;
mWidth = dim.w; mWidth = dim.w;
mImageSize = mDims.c()*mDims.h()*mDims.w(); mImageSize = dim.c*dim.h*dim.w;
mBatch.resize(mBatchSize*mImageSize, 0); mBatch.resize(mBatchSize*mImageSize, 0);
mLabels.resize(mBatchSize, 0); mLabels.resize(mBatchSize, 0);
mFileBatch.resize(mDims.n()*mImageSize, 0); mFileBatch.resize(dim.n*mImageSize, 0);
mFileLabels.resize(mDims.n(), 0); mFileLabels.resize(dim.n, 0);
mFileImgList = fileimglist; mFileImgList = fileimglist;
readInListFile(fileimglist, mListImg); readInListFile(fileimglist, mListImg);
mFileLabelList = filelabellist; mFileLabelList = filelabellist;
@@ -27,7 +27,7 @@ BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches,
void BatchStream::reset(int firstBatch) { void BatchStream::reset(int firstBatch) {
mBatchCount = 0; mBatchCount = 0;
mFileCount = 0; mFileCount = 0;
mFileBatchPos = mDims.n(); mFileBatchPos = mDims.d[0];
skip(firstBatch); skip(firstBatch);
} }
@@ -37,11 +37,11 @@ bool BatchStream::next() {
return false; return false;
for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize) { for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize) {
assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.n()); assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.d[0]);
if (mFileBatchPos == mDims.n() && !update()) if (mFileBatchPos == mDims.d[0] && !update())
return false; return false;
csize = std::min(mBatchSize - batchPos, mDims.n() - mFileBatchPos); csize = std::min(mBatchSize - batchPos, mDims.d[0] - mFileBatchPos);
std::copy_n(getFileBatch() + mFileBatchPos * mImageSize, csize * mImageSize, getBatch() + batchPos * mImageSize); std::copy_n(getFileBatch() + mFileBatchPos * mImageSize, csize * mImageSize, getBatch() + batchPos * mImageSize);
std::copy_n(getFileLabels() + mFileBatchPos, csize, getLabels() + batchPos); std::copy_n(getFileLabels() + mFileBatchPos, csize, getLabels() + batchPos);
} }
@@ -50,8 +50,8 @@ bool BatchStream::next() {
} }
void BatchStream::skip(int skipCount) { void BatchStream::skip(int skipCount) {
if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n()) { if (mBatchSize >= mDims.d[0] && mBatchSize%mDims.d[0] == 0 && mFileBatchPos == mDims.d[0]) {
mFileCount += skipCount * mBatchSize / mDims.n(); mFileCount += skipCount * mBatchSize / mDims.d[0];
return; return;
} }
+5 -5
View File
@@ -8,13 +8,13 @@ Int8EntropyCalibrator::Int8EntropyCalibrator(BatchStream& stream, int firstBatch
mCalibTableFilePath(calibTableFilePath), mCalibTableFilePath(calibTableFilePath),
mInputBlobName(inputBlobName.c_str()), mInputBlobName(inputBlobName.c_str()),
mReadCache(readCache) { mReadCache(readCache) {
nvinfer1::DimsNCHW dims = mStream.getDims(); nvinfer1::Dims4 dims = mStream.getDims();
mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w(); mInputCount = mStream.getBatchSize() * dims.d[1] * dims.d[2] * dims.d[3];
checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float))); checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float)));
mStream.reset(firstBatch); mStream.reset(firstBatch);
} }
bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int nbBindings) { bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int nbBindings) noexcept {
if (!mStream.next()) if (!mStream.next())
return false; return false;
@@ -24,7 +24,7 @@ bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int
return true; return true;
} }
const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) { const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) noexcept {
mCalibrationCache.clear(); mCalibrationCache.clear();
assert(!mCalibTableFilePath.empty()); assert(!mCalibTableFilePath.empty());
std::ifstream input(mCalibTableFilePath, std::ios::binary); std::ifstream input(mCalibTableFilePath, std::ios::binary);
@@ -38,7 +38,7 @@ const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) {
return length ? &mCalibrationCache[0] : nullptr; return length ? &mCalibrationCache[0] : nullptr;
} }
void Int8EntropyCalibrator::writeCalibrationCache(const void* cache, size_t length) { void Int8EntropyCalibrator::writeCalibrationCache(const void* cache, size_t length) noexcept {
assert(!mCalibTableFilePath.empty()); assert(!mCalibTableFilePath.empty());
std::ofstream output(mCalibTableFilePath, std::ios::binary); std::ofstream output(mCalibTableFilePath, std::ios::binary);
output.write(reinterpret_cast<const char*>(cache), length); output.write(reinterpret_cast<const char*>(cache), length);
+106 -88
View File
@@ -15,7 +15,7 @@ using namespace nvinfer1;
// Logger for info/warning/errors // Logger for info/warning/errors
class Logger : public ILogger { class Logger : public ILogger {
void log(Severity severity, const char* msg) override { void log(Severity severity, const char* msg) noexcept override {
#ifdef DEBUG #ifdef DEBUG
std::cout <<"TENSORRT LOG: "<< msg << std::endl; std::cout <<"TENSORRT LOG: "<< msg << std::endl;
#endif #endif
@@ -39,7 +39,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
#if NV_TENSORRT_MAJOR >= 5 #if NV_TENSORRT_MAJOR >= 5
std::cout<<"DLAs: "<<builderRT->getNbDLACores()<<"\n"; std::cout<<"DLAs: "<<builderRT->getNbDLACores()<<"\n";
#endif #endif
networkRT = builderRT->createNetwork(); networkRT = builderRT->createNetworkV2(0u);
#if NV_TENSORRT_MAJOR >= 6 #if NV_TENSORRT_MAJOR >= 6
configRT = builderRT->createBuilderConfig(); configRT = builderRT->createBuilderConfig();
#endif #endif
@@ -53,22 +53,25 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
configRT->setMinTimingIterations(1); configRT->setMinTimingIterations(1);
configRT->setMaxWorkspaceSize(1 << 30); configRT->setMaxWorkspaceSize(1 << 30);
configRT->setFlag(BuilderFlag::kDEBUG); configRT->setFlag(BuilderFlag::kDEBUG);
#else
builderRT->setMaxWorkspaceSize(1 << 30);
#endif #endif
//input and dataType //input and dataType
dataDim_t dim = net->layers[0]->input_dim; dataDim_t dim = net->layers[0]->input_dim;
dtRT = DataType::kFLOAT; dtRT = DataType::kFLOAT;
builderRT->setMaxBatchSize(net->maxBatchSize); builderRT->setMaxBatchSize(net->maxBatchSize);
builderRT->setMaxWorkspaceSize(1 << 30);
if(net->fp16 && builderRT->platformHasFastFp16()) { if(net->fp16 && builderRT->platformHasFastFp16()) {
dtRT = DataType::kHALF; dtRT = DataType::kHALF;
#if NV_TENSORRT_MAJOR < 6
builderRT->setHalf2Mode(true); builderRT->setHalf2Mode(true);
#endif
#if NV_TENSORRT_MAJOR >= 6 #if NV_TENSORRT_MAJOR >= 6
configRT->setFlag(BuilderFlag::kFP16); configRT->setFlag(BuilderFlag::kFP16);
#endif #endif
} }
#if NV_TENSORRT_MAJOR >= 5 #if NV_TENSORRT_MAJOR >= 5 && NV_TENSORRT_MAJOR < 8
if(net->dla && builderRT->getNbDLACores() > 0) { if(net->dla && builderRT->getNbDLACores() > 0) {
dtRT = DataType::kHALF; dtRT = DataType::kHALF;
builderRT->setFp16Mode(true); builderRT->setFp16Mode(true);
@@ -77,6 +80,15 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
builderRT->setDLACore(0); builderRT->setDLACore(0);
} }
#endif #endif
#if NV_TENSORRT_MAJOR >= 8
if(net->dla && builderRT->getNbDLACores() > 0) {
dtRT = DataType::kHALF;
configRT->setFlag(BuilderFlag::kFP16);
configRT->setFlag(BuilderFlag::kGPU_FALLBACK);
configRT->setDefaultDeviceType(DeviceType::kDLA);
configRT->setDLACore(0);
}
#endif
#if NV_TENSORRT_MAJOR >= 6 #if NV_TENSORRT_MAJOR >= 6
if(net->int8 && builderRT->platformHasFastInt8()){ if(net->int8 && builderRT->platformHasFastInt8()){
// dtRT = DataType::kINT8; // dtRT = DataType::kINT8;
@@ -104,7 +116,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
// add input layer // add input layer
ITensor *input = networkRT->addInput("data", DataType::kFLOAT, ITensor *input = networkRT->addInput("data", DataType::kFLOAT,
DimsCHW{ dim.c, dim.h, dim.w}); Dims3{ dim.c, dim.h, dim.w});
checkNULL(input); checkNULL(input);
//add other layers //add other layers
@@ -368,8 +380,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) {
if(l->pool_mode == tkdnnPoolingMode_t::POOLING_MAX_FIXEDSIZE) if(l->pool_mode == tkdnnPoolingMode_t::POOLING_MAX_FIXEDSIZE)
{ {
IPlugin *plugin = new MaxPoolFixedSizeRT(l->output_dim.c, l->output_dim.h, l->output_dim.w, l->output_dim.n, l->strideH, l->strideW, l->winH, l->winH-1); auto *plugin = new MaxPoolFixedSizeRT(l->output_dim.c, l->output_dim.h, l->output_dim.w, l->output_dim.n, l->strideH, l->strideW, l->winH, l->winH-1);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
@@ -392,8 +404,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
#if NV_TENSORRT_MAJOR < 6 #if NV_TENSORRT_MAJOR < 6
// plugin version // plugin version
IPlugin *plugin = new ActivationLeakyRT(l->slope); auto *plugin = new ActivationLeakyRT(l->slope);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
#else #else
@@ -407,28 +419,35 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU); IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} else if(l->act_mode == CUDNN_ACTIVATION_SIGMOID) { } else if(l->act_mode == CUDNN_ACTIVATION_SIGMOID || l->act_mode == ACTIVATION_LOGISTIC) {
IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kSIGMOID); IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kSIGMOID);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
else if(l->act_mode == CUDNN_ACTIVATION_CLIPPED_RELU) { else if(l->act_mode == CUDNN_ACTIVATION_CLIPPED_RELU) {
IPlugin *plugin = new ActivationReLUCeiling(l->ceiling); IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kCLIP);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); lRT->setAlpha(0);
lRT->setBeta(l->ceiling);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
else if(l->act_mode == ACTIVATION_MISH) { else if(l->act_mode == ACTIVATION_MISH) {
IPlugin *plugin = new ActivationMishRT(); // Uncomment this to see if you have better performance
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); // For older TensorRT or for FP32 this might be better
checkNULL(lRT); //auto *plugin = new ActivationMishRT();
return lRT; //auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
}
else if(l->act_mode == ACTIVATION_LOGISTIC) { // Assemble MISH using 3 layers that are going to be merged by TensorRT
IPlugin *plugin = new ActivationLogisticRT(); IActivationLayer *lRT1 = networkRT->addActivation(*input, ActivationType::kSOFTPLUS);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); lRT1->setAlpha(1);
checkNULL(lRT); lRT1->setBeta(1);
return lRT;
IActivationLayer *lRT2 = networkRT->addActivation(*lRT1->getOutput(0), ActivationType::kTANH);
IElementWiseLayer *lRT3 = networkRT->addElementWise(*input, *lRT2->getOutput(0), ElementWiseOperation::kPROD);
checkNULL(lRT3);
return lRT3;
} }
else { else {
FatalError("this Activation mode is not yet implemented"); FatalError("this Activation mode is not yet implemented");
@@ -460,8 +479,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) {
} }
if(l->groups > 1){ if(l->groups > 1){
IPlugin *plugin = new RouteRT(l->groups, l->group_id); auto *plugin = new RouteRT(l->groups, l->group_id);
IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin); auto *lRT = networkRT->addPluginV2(tens, l->layers_n, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
@@ -472,8 +491,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) {
ILayer* NetworkRT::convert_layer(ITensor *input, Flatten *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Flatten *l) {
IPlugin *plugin = new FlattenConcatRT(); auto *plugin = new FlattenConcatRT();
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
@@ -481,8 +500,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Flatten *l) {
ILayer* NetworkRT::convert_layer(ITensor *input, Reshape *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Reshape *l) {
// std::cout<<"convert Reshape\n"; // std::cout<<"convert Reshape\n";
IPlugin *plugin = new ReshapeRT(l->output_dim); auto *plugin = new ReshapeRT(l->output_dim);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
@@ -494,7 +513,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Resize *l) {
checkNULL(lRT); checkNULL(lRT);
Dims d{}; Dims d{};
lRT->setResizeMode(ResizeMode(l->mode)); lRT->setResizeMode(ResizeMode(l->mode));
lRT->setOutputDimensions(DimsCHW{l->output_dim.c, l->output_dim.h, l->output_dim.w}); lRT->setOutputDimensions(Dims3{l->output_dim.c, l->output_dim.h, l->output_dim.w});
return lRT; return lRT;
} }
@@ -502,8 +521,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Reorg *l) {
//std::cout<<"convert Reorg\n"; //std::cout<<"convert Reorg\n";
//std::cout<<"New plugin REORG\n"; //std::cout<<"New plugin REORG\n";
IPlugin *plugin = new ReorgRT(l->stride); auto *plugin = new ReorgRT(l->stride);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
@@ -512,8 +531,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Region *l) {
//std::cout<<"convert Region\n"; //std::cout<<"convert Region\n";
//std::cout<<"New plugin REGION\n"; //std::cout<<"New plugin REGION\n";
IPlugin *plugin = new RegionRT(l->classes, l->coords, l->num); auto *plugin = new RegionRT(l->classes, l->coords, l->num);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
@@ -534,11 +553,11 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
else else
{ {
// plugin version // plugin version
IPlugin *plugin = new ShortcutRT(l->backLayer->output_dim, l->mul); auto *plugin = new ShortcutRT(l->backLayer->output_dim, l->mul);
ITensor **inputs = new ITensor*[2]; ITensor **inputs = new ITensor*[2];
inputs[0] = input; inputs[0] = input;
inputs[1] = back_tens; inputs[1] = back_tens;
IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin); auto *lRT = networkRT->addPluginV2(inputs, 2, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
@@ -548,8 +567,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Yolo *l) {
//std::cout<<"convert Yolo\n"; //std::cout<<"convert Yolo\n";
//std::cout<<"New plugin YOLO\n"; //std::cout<<"New plugin YOLO\n";
IPlugin *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY, l->nms_thresh, l->nsm_kind, l->new_coords); auto *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY, l->nms_thresh, l->nsm_kind, l->new_coords);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
@@ -558,8 +577,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Upsample *l) {
//std::cout<<"convert Upsample\n"; //std::cout<<"convert Upsample\n";
//std::cout<<"New plugin UPSAMPLE\n"; //std::cout<<"New plugin UPSAMPLE\n";
IPlugin *plugin = new UpsampleRT(l->stride); auto *plugin = new UpsampleRT(l->stride);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
checkNULL(lRT); checkNULL(lRT);
return lRT; return lRT;
} }
@@ -574,10 +593,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
inputs[1] = preconv->getOutput(0); inputs[1] = preconv->getOutput(0);
//std::cout<<"New plugin DEFORMABLE\n"; //std::cout<<"New plugin DEFORMABLE\n";
IPlugin *plugin = new DeformableConvRT(l->chunk_dim, l->kernelH, l->kernelW, l->strideH, l->strideW, l->paddingH, l->paddingW, auto *plugin = new DeformableConvRT(l->chunk_dim, l->kernelH, l->kernelW, l->strideH, l->strideW, l->paddingH, l->paddingW,
l->deformableGroup, l->input_dim.n, l->input_dim.c, l->input_dim.h, l->input_dim.w, l->deformableGroup, l->input_dim.n, l->input_dim.c, l->input_dim.h, l->input_dim.w,
l->output_dim.n, l->output_dim.c, l->output_dim.h, l->output_dim.w, l); l->output_dim.n, l->output_dim.c, l->output_dim.h, l->output_dim.w, l);
IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin); auto *lRT = networkRT->addPluginV2(inputs, 2, *plugin);
checkNULL(lRT); checkNULL(lRT);
lRT->setName( ("Deformable" + std::to_string(l->id)).c_str() ); lRT->setName( ("Deformable" + std::to_string(l->id)).c_str() );
delete[](inputs); delete[](inputs);
@@ -646,53 +665,52 @@ bool NetworkRT::deserialize(const char *filename) {
file.close(); file.close();
} }
pluginFactory = new PluginFactory();
runtimeRT = createInferRuntime(loggerRT); runtimeRT = createInferRuntime(loggerRT);
engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory); engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size);
//if (gieModelStream) delete [] gieModelStream; //if (gieModelStream) delete [] gieModelStream;
return true; return true;
} }
// auto* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
// const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
// std::string name(layerName);
// //std::cout<<name<<std::endl;
IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) { // #if NV_TENSORRT_MAJOR < 6
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf; // if(name.find("ActivationLeaky") == 0) {
// ActivationLeakyRT *a = new ActivationLeakyRT(readBUF<float>(buf));
std::string name(layerName); // a->size = readBUF<int>(buf);
//std::cout<<name<<std::endl; // assert(buf == bufCheck + serialLength);
// return a;
if(name.find("ActivationLeaky") == 0) { // }
ActivationLeakyRT *a = new ActivationLeakyRT(readBUF<float>(buf)); // #endif
a->size = readBUF<int>(buf); /*if(name.find("ActivationMish") == 0) {
assert(buf == bufCheck + serialLength);
return a;
}
if(name.find("ActivationMish") == 0) {
ActivationMishRT *a = new ActivationMishRT(); ActivationMishRT *a = new ActivationMishRT();
a->size = readBUF<int>(buf); a->size = readBUF<int>(buf);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return a; return a;
} }*/
if(name.find("ActivationLogistic") == 0) { /*if(name.find("ActivationLogistic") == 0) {
ActivationLogisticRT *a = new ActivationLogisticRT(); ActivationLogisticRT *a = new ActivationLogisticRT();
a->size = readBUF<int>(buf); a->size = readBUF<int>(buf);
return a; return a;
} }*/
if(name.find("ActivationLogistic") == 0) { /*if(name.find("ActivationLogistic") == 0) {
ActivationLogisticRT *a = new ActivationLogisticRT(); ActivationLogisticRT *a = new ActivationLogisticRT();
a->size = readBUF<int>(buf); a->size = readBUF<int>(buf);
return a; return a;
} }*/
if(name.find("ActivationCReLU") == 0) { /*if(name.find("ActivationCReLU") == 0) {
float activationReluTemp = readBUF<float>(buf); float activationReluTemp = readBUF<float>(buf);
ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp); ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp);
a->size = readBUF<int>(buf); a->size = readBUF<int>(buf);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return a; return a;
} }*/
if(name.find("Region") == 0) { /*if(name.find("Region") == 0) {
int classesTemp = readBUF<int>(buf); int classesTemp = readBUF<int>(buf);
int coordsTemp = readBUF<int>(buf); int coordsTemp = readBUF<int>(buf);
int numTemp = readBUF<int>(buf); int numTemp = readBUF<int>(buf);
@@ -703,9 +721,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
r->w = readBUF<int>(buf); r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return r; return r;
} }*/
if(name.find("Reorg") == 0) { /*if(name.find("Reorg") == 0) {
int strideTemp = readBUF<int>(buf); int strideTemp = readBUF<int>(buf);
ReorgRT *r = new ReorgRT(strideTemp); ReorgRT *r = new ReorgRT(strideTemp);
r->c = readBUF<int>(buf); r->c = readBUF<int>(buf);
@@ -713,9 +731,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
r->w = readBUF<int>(buf); r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return r; return r;
} }*/
if(name.find("Shortcut") == 0) { /*if(name.find("Shortcut") == 0) {
tk::dnn::dataDim_t bdim; tk::dnn::dataDim_t bdim;
bdim.c = readBUF<int>(buf); bdim.c = readBUF<int>(buf);
bdim.h = readBUF<int>(buf); bdim.h = readBUF<int>(buf);
@@ -728,9 +746,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
r->w = readBUF<int>(buf); r->w = readBUF<int>(buf);
return r; return r;
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
} }*/
if(name.find("Pooling") == 0) { /*if(name.find("Pooling") == 0) {
int cTemp = readBUF<int>(buf); int cTemp = readBUF<int>(buf);
int hTemp = readBUF<int>(buf); int hTemp = readBUF<int>(buf);
int wTemp = readBUF<int>(buf); int wTemp = readBUF<int>(buf);
@@ -743,9 +761,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp); MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return r; return r;
} }*/
if(name.find("Resize") == 0) { /*if(name.find("Resize") == 0) {
int o_cTemp = readBUF<int>(buf); int o_cTemp = readBUF<int>(buf);
int o_hTemp = readBUF<int>(buf); int o_hTemp = readBUF<int>(buf);
int o_wTemp = readBUF<int>(buf); int o_wTemp = readBUF<int>(buf);
@@ -756,9 +774,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
r->i_w = readBUF<int>(buf); r->i_w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return r; return r;
} }*/
if(name.find("Flatten") == 0) { /*if(name.find("Flatten") == 0) {
FlattenConcatRT *r = new FlattenConcatRT(); FlattenConcatRT *r = new FlattenConcatRT();
r->c = readBUF<int>(buf); r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf); r->h = readBUF<int>(buf);
@@ -767,9 +785,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
r->cols = readBUF<int>(buf); r->cols = readBUF<int>(buf);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return r; return r;
} }*/
if(name.find("Reshape") == 0) { /*if(name.find("Reshape") == 0) {
dataDim_t new_dim; dataDim_t new_dim;
new_dim.n = readBUF<int>(buf); new_dim.n = readBUF<int>(buf);
@@ -780,9 +798,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return r; return r;
} }*/
if(name.find("Yolo") == 0) { /*if(name.find("Yolo") == 0) {
int classes_temp = readBUF<int>(buf); int classes_temp = readBUF<int>(buf);
int num_temp = readBUF<int>(buf); int num_temp = readBUF<int>(buf);
@@ -816,8 +834,8 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
yolos[n_yolos++] = r; yolos[n_yolos++] = r;
return r; return r;
} }*/
if(name.find("Upsample") == 0) { /*if(name.find("Upsample") == 0) {
int strideTemp = readBUF<int>(buf); int strideTemp = readBUF<int>(buf);
UpsampleRT* r = new UpsampleRT(strideTemp); UpsampleRT* r = new UpsampleRT(strideTemp);
r->c = readBUF<int>(buf); r->c = readBUF<int>(buf);
@@ -825,9 +843,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
r->w = readBUF<int>(buf); r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return r; return r;
} }*/
if(name.find("Route") == 0) { /*if(name.find("Route") == 0) {
int groupsTemp = readBUF<int>(buf); int groupsTemp = readBUF<int>(buf);
int group_idTemp = readBUF<int>(buf); int group_idTemp = readBUF<int>(buf);
RouteRT* r = new RouteRT(groupsTemp, group_idTemp); RouteRT* r = new RouteRT(groupsTemp, group_idTemp);
@@ -839,9 +857,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
r->w = readBUF<int>(buf); r->w = readBUF<int>(buf);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return r; return r;
} }*/
if(name.find("Deformable") == 0) { /*if(name.find("Deformable") == 0) {
int chuck_dimTemp = readBUF<int>(buf); int chuck_dimTemp = readBUF<int>(buf);
int khTemp = readBUF<int>(buf); int khTemp = readBUF<int>(buf);
int kwTemp = readBUF<int>(buf); int kwTemp = readBUF<int>(buf);
@@ -892,10 +910,10 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
free(aus); free(aus);
assert(buf == bufCheck + serialLength); assert(buf == bufCheck + serialLength);
return r; return r;
} }*/
FatalError("Cant deserialize Plugin"); // FatalError("Cant deserialize Plugin");
return NULL; // return NULL;
} // }
}} }}
+8 -7
View File
@@ -1,5 +1,6 @@
#include "yoloContainer.h"
#include "Yolo3Detection.h" #include "Yolo3Detection.h"
#include "pluginsRT/YoloRT.h"
namespace tk { namespace dnn { namespace tk { namespace dnn {
@@ -14,12 +15,12 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
tk::dnn::dataDim_t idim = netRT->input_dim; tk::dnn::dataDim_t idim = netRT->input_dim;
idim.n = nBatches; idim.n = nBatches;
if(netRT->pluginFactory->n_yolos < 2 ) { if(yoloContainer.n_yolos < 2 ) {
FatalError("this is not yolo3"); FatalError("this is not yolo3");
} }
for(int i=0; i<netRT->pluginFactory->n_yolos; i++) { for(int i=0; i<yoloContainer.n_yolos; i++) {
YoloRT *yRT = netRT->pluginFactory->yolos[i]; YoloRT *yRT = yoloContainer.yolos[i];
classes = yRT->classes; classes = yRT->classes;
num = yRT->num; num = yRT->num;
nMasks = yRT->n_masks; nMasks = yRT->n_masks;
@@ -95,8 +96,8 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
//get yolo outputs //get yolo outputs
std::vector<float *> rt_out; std::vector<float *> rt_out;
//dnnType *rt_out[netRT->pluginFactory->n_yolos]; //dnnType *rt_out[yoloContainer.n_yolos];
for(int i=0; i<netRT->pluginFactory->n_yolos; i++) for(int i=0; i<yoloContainer.n_yolos; i++)
rt_out.push_back((dnnType*)netRT->buffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi); rt_out.push_back((dnnType*)netRT->buffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi);
float x_ratio = float(originalSize[bi].width) / float(netRT->input_dim.w); float x_ratio = float(originalSize[bi].width) / float(netRT->input_dim.w);
@@ -104,7 +105,7 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
// compute dets // compute dets
nDets = 0; nDets = 0;
for(int i=0; i<netRT->pluginFactory->n_yolos; i++) { for(int i=0; i<yoloContainer.n_yolos; i++) {
yolo[i]->dstData = rt_out[i]; yolo[i]->dstData = rt_out[i];
yolo[i]->computeDetections(dets, nDets, netRT->input_dim.w, netRT->input_dim.h, confThreshold, yolo[i]->new_coords); yolo[i]->computeDetections(dets, nDets, netRT->input_dim.w, netRT->input_dim.h, confThreshold, yolo[i]->new_coords);
} }
+11 -2
View File
@@ -1,4 +1,15 @@
#include <vector>
#include "kernels.h" #include "kernels.h"
#include "pluginsRT/ActivationLogisticRT.h"
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection ActivationLogisticRTCreator::mFC{};
std::vector<nvinfer1::PluginField> ActivationLogisticRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(ActivationLogisticRTCreator);
}}
__global__ __global__
void activation_logistic(dnnType *input, dnnType *output, int size) { void activation_logistic(dnnType *input, dnnType *output, int size) {
@@ -21,5 +32,3 @@ void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, cud
activation_logistic<<<blocks, threads, 0, stream>>>(srcData, dstData, size); activation_logistic<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
} }
+21 -4
View File
@@ -1,8 +1,17 @@
#include "kernels.h"
#include <math.h> #include <math.h>
#include "kernels.h"
#include "pluginsRT/ActivationMishRT.h"
#define MISH_THRESHOLD 20 #define MISH_THRESHOLD 20
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection ActivationMishRTCreator::mFC{};
std::vector<nvinfer1::PluginField> ActivationMishRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(ActivationMishRTCreator);
}}
__device__ __device__
float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);} float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);}
@@ -13,8 +22,6 @@ float softplus_kernel(float x, float threshold = 20) {
return logf(expf(x) + 1); return logf(expf(x) + 1);
} }
__device__ __device__
float mish_yashas(float x) { float mish_yashas(float x) {
float e = __expf(x); float e = __expf(x);
@@ -28,6 +35,16 @@ float mish_yashas(float x) {
return x - 2 * __fdividef(x, n + 2); return x - 2 * __fdividef(x, n + 2);
} }
__device__ float mish_yashas2(float x)
{
float e = __expf(x);
float n = e * e + 2 * e;
if (x <= -0.6f)
return x * __fdividef(n, n + 2);
return x - 2 * __fdividef(x, n + 2);
}
// https://github.com/digantamisra98/Mish // https://github.com/digantamisra98/Mish
// https://github.com/AlexeyAB/darknet/blob/master/src/activation_kernels.cu // https://github.com/AlexeyAB/darknet/blob/master/src/activation_kernels.cu
__global__ __global__
@@ -35,7 +52,7 @@ void activation_mish(dnnType *input, dnnType *output, int size) {
int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
if (i < size) if (i < size)
// output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD)); // output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD));
output[i] = mish_yashas(input[i]); output[i] = mish_yashas2(input[i]);
} }
/** /**
+9
View File
@@ -1,4 +1,13 @@
#include "kernels.h" #include "kernels.h"
#include "pluginsRT/ActivationReLUCeilingRT.h"
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection ActivationReLUCeilingCreator::mFC{};
std::vector<nvinfer1::PluginField> ActivationReLUCeilingCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(ActivationReLUCeilingCreator);
}}
__global__ __global__
void activation_relu_ceiling(dnnType *input, dnnType *output, int size, const float ceiling) { void activation_relu_ceiling(dnnType *input, dnnType *output, int size, const float ceiling) {
+9 -1
View File
@@ -1,6 +1,14 @@
#include "kernels.h"
#include <math.h> #include <math.h>
#include "kernels.h"
#include "pluginsRT/ActivationSigmoidRT.h"
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection ActivationSigmoidRTCreator::mFC{};
std::vector<nvinfer1::PluginField> ActivationSigmoidRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(ActivationSigmoidRTCreator);
}}
__global__ __global__
void activation_sigmoid(dnnType *input, dnnType *output, int size) { void activation_sigmoid(dnnType *input, dnnType *output, int size) {
+10 -1
View File
@@ -3,9 +3,11 @@
#include <cstring> #include <cstring>
#include <string> #include <string>
#include <iostream> #include <iostream>
#include "kernels.h"
#include <errno.h> #include <errno.h>
#include "kernels.h"
#include "pluginsRT/DeformableConvRT.h"
#define CUDA_KERNEL_LOOP(i, n) \ #define CUDA_KERNEL_LOOP(i, n) \
for (int i = blockIdx.x * blockDim.x + threadIdx.x; \ for (int i = blockIdx.x * blockDim.x + threadIdx.x; \
i < (n); \ i < (n); \
@@ -17,6 +19,13 @@ inline int GET_BLOCKS(const int N)
return (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS; return (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS;
} }
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection DeformableConvRTCreator::mFC{};
std::vector<nvinfer1::PluginField> DeformableConvRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(DeformableConvRTCreator);
}}
__device__ __host__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width, __device__ __host__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width,
const int height, const int width, float h, float w) { const int height, const int width, float h, float w) {
+9
View File
@@ -1,4 +1,13 @@
#include "kernels.h" #include "kernels.h"
#include "pluginsRT/MaxPoolingFixedSizeRT.h"
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection MaxPoolFixedSizeRTCreator::mFC{};
std::vector<nvinfer1::PluginField> MaxPoolFixedSizeRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(MaxPoolFixedSizeRTCreator);
}}
__global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c, int stride_x, int stride_y, int size, int pad, float *input, float *output) __global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c, int stride_x, int stride_y, int size, int pad, float *input, float *output)
{ {
+9
View File
@@ -1,4 +1,13 @@
#include "kernels.h" #include "kernels.h"
#include "pluginsRT/ReorgRT.h"
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection ReorgRTCreator::mFC{};
std::vector<nvinfer1::PluginField> ReorgRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(ReorgRTCreator);
}}
__global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, int stride, int forward, float *out) __global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, int stride, int forward, float *out)
{ {
+11 -1
View File
@@ -1,6 +1,16 @@
#include "kernels.h"
#include <stdio.h> #include <stdio.h>
#include "kernels.h"
#include "pluginsRT/ResizeLayerRT.h"
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection ResizeLayerRTCreator::mFC{};
std::vector<nvinfer1::PluginField> ResizeLayerRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(ResizeLayerRTCreator);
}}
__global__ void resize_kernel( int size,float *x, int i_w, int i_h, int i_c, __global__ void resize_kernel( int size,float *x, int i_w, int i_h, int i_c,
int o_w, int o_h, int o_c, int batch, float *out) int o_w, int o_h, int o_c, int batch, float *out)
{ {
+2 -1
View File
@@ -1,6 +1,7 @@
#include "kernels.h"
#include <math.h> #include <math.h>
#include "kernels.h"
__global__ void scal_add_kernel(dnnType* dstData, int size, float alpha, float beta, int inc) __global__ void scal_add_kernel(dnnType* dstData, int size, float alpha, float beta, int inc)
{ {
int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+10
View File
@@ -1,6 +1,16 @@
#include "kernels.h" #include "kernels.h"
#include "assert.h" #include "assert.h"
#include "pluginsRT/ShortcutRT.h"
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection ShortcutRTCreator::mFC{};
std::vector<nvinfer1::PluginField> ShortcutRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(ShortcutRTCreator);
}}
__global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stride, int sample, int batch, __global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stride, int sample, int batch,
int w1, int h1, int c1, dnnType *add, int w1, int h1, int c1, dnnType *add,
int w2, int h2, int c2, float s1, float s2, dnnType *out) int w2, int h2, int c2, float s1, float s2, dnnType *out)
+33
View File
@@ -0,0 +1,33 @@
#include "pluginsRT/RegionRT.h"
#include "pluginsRT/RouteRT.h"
#include "pluginsRT/ReshapeRT.h"
#include "pluginsRT/FlattenConcatRT.h"
#include "pluginsRT/YoloRT.h"
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection RegionRTCreator::mFC{};
std::vector<nvinfer1::PluginField> RegionRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(RegionRTCreator);
nvinfer1::PluginFieldCollection RouteRTCreator::mFC{};
std::vector<nvinfer1::PluginField> RouteRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(RouteRTCreator);
nvinfer1::PluginFieldCollection ReshapeRTCreator::mFC{};
std::vector<nvinfer1::PluginField> ReshapeRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(ReshapeRTCreator);
nvinfer1::PluginFieldCollection FlattenConcatRTCreator::mFC{};
std::vector<nvinfer1::PluginField> FlattenConcatRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(FlattenConcatRTCreator);
nvinfer1::PluginFieldCollection YoloRTCreator::mFC{};
std::vector<nvinfer1::PluginField> YoloRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(YoloRTCreator);
}}
+9
View File
@@ -1,4 +1,13 @@
#include "kernels.h" #include "kernels.h"
#include "pluginsRT/UpsampleRT.h"
// Static class fields initialization
namespace tk { namespace dnn {
nvinfer1::PluginFieldCollection UpsampleRTCreator::mFC{};
std::vector<nvinfer1::PluginField> UpsampleRTCreator::mPluginAttributes;
REGISTER_TENSORRT_PLUGIN(UpsampleRTCreator);
}}
__global__ void upsample_kernel(size_t N, dnnType *x, int w, int h, int c, int batch, int stride, int forward, float scale, dnnType *out) __global__ void upsample_kernel(size_t N, dnnType *x, int w, int h, int c, int batch, int stride, int forward, float scale, dnnType *out)
{ {
+41
View File
@@ -0,0 +1,41 @@
#include "yoloContainer.h"
#include "pluginsRT/YoloRT.h"
namespace tk { namespace dnn {
YoloContainer yoloContainer;
nvinfer1::IPluginV2* YoloRTCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept {
const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
int classes_temp = tk::dnn::readBUF<int>(buf);
int num_temp = tk::dnn::readBUF<int>(buf);
int n_masks_temp = tk::dnn::readBUF<int>(buf);
float scale_xy_temp = tk::dnn::readBUF<float>(buf);
float nms_thresh_temp = tk::dnn::readBUF<float>(buf);
int nms_kind_temp = tk::dnn::readBUF<int>(buf);
int new_coords_temp = tk::dnn::readBUF<int>(buf);
YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp);
r->c = tk::dnn::readBUF<int>(buf);
r->h = tk::dnn::readBUF<int>(buf);
r->w = tk::dnn::readBUF<int>(buf);
for(int i=0; i<r->n_masks; i++)
r->mask[i] = tk::dnn::readBUF<dnnType>(buf);
for(int i=0; i<r->n_masks*2*r->num; i++)
r->bias[i] = tk::dnn::readBUF<dnnType>(buf);
// save classes names
r->classesNames.resize(r->classes);
for(int i=0; i<r->classes; i++) {
char tmp[YOLORT_CLASSNAME_W];
for(int j=0; j<YOLORT_CLASSNAME_W; j++)
tmp[j] = tk::dnn::readBUF<char>(buf);
r->classesNames[i] = std::string(tmp);
}
assert(buf == bufCheck + serialLength);
yoloContainer.yolos[yoloContainer.n_yolos++] = r;
return r;
}
}}
+15 -4
View File
@@ -15,7 +15,7 @@ using namespace nvinfer1;
// Logger for info/warning/errors // Logger for info/warning/errors
class Logger : public ILogger class Logger : public ILogger
{ {
void log(Severity severity, const char* msg) override void log(Severity severity, const char* msg) noexcept override
{ {
// suppress info-level messages // suppress info-level messages
if (severity != Severity::kINFO) if (severity != Severity::kINFO)
@@ -66,11 +66,11 @@ int main() {
std::cout<<"\n==== TensorRT ====\n"; std::cout<<"\n==== TensorRT ====\n";
// create the builder // create the builder
IBuilder* builder = nvinfer1::createInferBuilder(gLogger); IBuilder* builder = nvinfer1::createInferBuilder(gLogger);
INetworkDefinition* network = builder->createNetwork(); INetworkDefinition* network = builder->createNetworkV2(0u);
DataType dt = DataType::kFLOAT; DataType dt = DataType::kFLOAT;
// Create input of shape { 1, 1, 28, 28 } with name referenced by "data" // Create input of shape { 1, 1, 28, 28 } with name referenced by "data"
auto input = network->addInput("data", dt, DimsCHW{ 1, 28, 28}); auto input = network->addInput("data", dt, Dims3{ 1, 28, 28});
assert(input != nullptr); assert(input != nullptr);
tk::dnn::Conv2d *c0 = &l0; tk::dnn::Conv2d *c0 = &l0;
@@ -125,10 +125,21 @@ int main() {
network->markOutput(*prob->getOutput(0)); network->markOutput(*prob->getOutput(0));
// Build the engine // Build the engine
builder->setMaxBatchSize(1); #if NV_TENSORRT_MAJOR >= 6
auto config = builder->createBuilderConfig();
config->setMaxWorkspaceSize(1 << 20);
#else
builder->setMaxWorkspaceSize(1 << 20); builder->setMaxWorkspaceSize(1 << 20);
#endif
builder->setMaxBatchSize(1);
#if NV_TENSORRT_MAJOR >= 6
auto engine = builder->buildEngineWithConfig(*network, *config);
#else
auto engine = builder->buildCudaEngine(*network); auto engine = builder->buildCudaEngine(*network);
#endif
// we don't need the network any more // we don't need the network any more
network->destroy(); network->destroy();