From 76c978a2006cf367af20f67c4da02ef7ddfe0d13 Mon Sep 17 00:00:00 2001 From: Harijs Grinbergs <1474290+TheExDeus@users.noreply.github.com> Date: Thu, 28 Oct 2021 01:02:26 +0300 Subject: [PATCH] Added TensorRT8 support - Migrated to TRT8 API, but I didn't really try maintaining old compatiblity. Lowest support now could be TRT6. - Refractored some stuff that are bad C++ practices and made coding really hard, like including headers in namespaces or not using an include guard. - Needed to move the yolo container outside tkdnn object, which means we now only have one and global. Deserialization in TRT8 doesn't happen in your object, but in the plugin itself so it couldn't access the yolo objects. I think the need to hold onto yolo layers itself is flawed and shouldn't be nessesarry. --- include/tkDNN/Int8BatchStream.h | 12 +- include/tkDNN/Int8Calibrator.h | 16 +- include/tkDNN/NetworkRT.h | 43 +-- include/tkDNN/buffer_func.h | 20 ++ include/tkDNN/pluginsRT/ActivationLeakyRT.h | 37 ++- .../tkDNN/pluginsRT/ActivationLogisticRT.h | 135 +++++++-- include/tkDNN/pluginsRT/ActivationMishRT.h | 135 +++++++-- .../tkDNN/pluginsRT/ActivationReLUCeilingRT.h | 135 +++++++-- include/tkDNN/pluginsRT/ActivationSigmoidRT.h | 136 +++++++-- include/tkDNN/pluginsRT/DeformableConvRT.h | 244 ++++++++++++---- include/tkDNN/pluginsRT/FlattenConcatRT.h | 148 ++++++++-- .../tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h | 154 ++++++++-- include/tkDNN/pluginsRT/RegionRT.h | 154 ++++++++-- include/tkDNN/pluginsRT/ReorgRT.h | 145 +++++++-- include/tkDNN/pluginsRT/ReshapeRT.h | 140 +++++++-- include/tkDNN/pluginsRT/ResizeLayerRT.h | 152 ++++++++-- include/tkDNN/pluginsRT/RouteRT.h | 152 ++++++++-- include/tkDNN/pluginsRT/ShortcutRT.h | 156 ++++++++-- include/tkDNN/pluginsRT/UpsampleRT.h | 142 +++++++-- include/tkDNN/pluginsRT/YoloRT.h | 153 +++++++--- include/tkDNN/yoloContainer.h | 16 + src/Int8BatchStream.cpp | 28 +- src/Int8Calibrator.cpp | 16 +- src/NetworkRT.cpp | 275 +++++++++--------- src/Yolo3Detection.cpp | 33 ++- src/kernels/activation_leaky.cu | 4 +- src/kernels/activation_logistic.cu | 19 +- src/kernels/activation_mish.cu | 43 +-- src/kernels/activation_relu_ceiling.cu | 13 +- src/kernels/activation_sigmoid.cu | 12 +- src/kernels/deformable_conv.cu | 69 +++-- src/kernels/pooling.cu | 13 +- src/kernels/reorg.cu | 13 +- src/kernels/resize.cu | 14 +- src/kernels/scaladd.cu | 5 +- src/kernels/shortcut.cu | 32 +- src/kernels/static_init.cu | 33 +++ src/kernels/upsample.cu | 13 +- src/yoloContainer.cpp | 41 +++ tests/mnist/test_mnistRT.cpp | 45 +-- 40 files changed, 2417 insertions(+), 729 deletions(-) create mode 100644 include/tkDNN/buffer_func.h create mode 100644 include/tkDNN/yoloContainer.h create mode 100644 src/kernels/static_init.cu create mode 100644 src/yoloContainer.cpp diff --git a/include/tkDNN/Int8BatchStream.h b/include/tkDNN/Int8BatchStream.h index c39a11c..0fd0281 100644 --- a/include/tkDNN/Int8BatchStream.h +++ b/include/tkDNN/Int8BatchStream.h @@ -12,7 +12,7 @@ #include #include #include -#ifdef __linux__ +#ifdef __linux__ #include #endif @@ -23,9 +23,9 @@ #include "tkdnn.h" /* - * BatchStream implements the stream for the INT8 calibrator. - * It reads the two files .txt with the list of image file names - * and the list of label file names. + * BatchStream implements the stream for the INT8 calibrator. + * It reads the two files .txt with the list of image file names + * and the list of label file names. * It then iterates on images and labels. */ class BatchStream { @@ -39,7 +39,7 @@ public: float *getLabels() { return mLabels.data(); } int getBatchesRead() const { return mBatchCount; } int getBatchSize() const { return mBatchSize; } - nvinfer1::DimsNCHW getDims() const { return mDims; } + nvinfer1::Dims4 getDims() const { return mDims; } float* getFileBatch() { return &mFileBatch[0]; } float* getFileLabels() { return &mFileLabels[0]; } void readInListFile(const std::string& dataFilePath, std::vector& mListIn); @@ -55,7 +55,7 @@ private: int mFileBatchPos{ 0 }; int mImageSize{ 0 }; - nvinfer1::DimsNCHW mDims; + nvinfer1::Dims4 mDims; std::vector mBatch; std::vector mLabels; std::vector mFileBatch; diff --git a/include/tkDNN/Int8Calibrator.h b/include/tkDNN/Int8Calibrator.h index 4a0ea47..bd8df42 100644 --- a/include/tkDNN/Int8Calibrator.h +++ b/include/tkDNN/Int8Calibrator.h @@ -20,20 +20,20 @@ /* * Int8EntropyCalibrator implements the INT8 calibrator to achieve the - * INT8 quantization. It uses a BatchStream stream to scroll through - * images data. It also implements the calibration cache, a way to - * save the calibration process results to reduce the running time: + * INT8 quantization. It uses a BatchStream stream to scroll through + * images data. It also implements the calibration cache, a way to + * save the calibration process results to reduce the running time: * the calibration process takes a long time. */ class Int8EntropyCalibrator : public nvinfer1::IInt8EntropyCalibrator { 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); virtual ~Int8EntropyCalibrator() { checkCuda(cudaFree(mDeviceInput)); } - int getBatchSize() const override { return mStream.getBatchSize(); } - bool getBatch(void* bindings[], const char* names[], int nbBindings) override; - const void* readCalibrationCache(size_t& length) override; - void writeCalibrationCache(const void* cache, size_t length) override; + int getBatchSize() const noexcept override { return mStream.getBatchSize(); } + bool getBatch(void* bindings[], const char* names[], int nbBindings) noexcept override; + const void* readCalibrationCache(size_t& length) noexcept override; + void writeCalibrationCache(const void* cache, size_t length) noexcept override; private: BatchStream mStream; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 9892a24..1b44a0a 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -2,28 +2,14 @@ #define NETWORKRT_H #include // memcpy +#include + #include "utils.h" #include "Network.h" #include "Layer.h" #include "NvInfer.h" -#include -namespace tk { namespace dnn { - -template void writeBUF(char*& buffer, const T& val) -{ - *reinterpret_cast(buffer) = val; - buffer += sizeof(T); -} - -template T readBUF(const char*& buffer) -{ - T val = *reinterpret_cast(buffer); - buffer += sizeof(T); - return val; -} - -using namespace nvinfer1; +// using namespace nvinfer1; #include "pluginsRT/ActivationLeakyRT.h" #include "pluginsRT/ActivationLogisticRT.h" #include "pluginsRT/ActivationReLUCeilingRT.h" @@ -40,16 +26,7 @@ using namespace nvinfer1; #include "pluginsRT/ReshapeRT.h" #include "pluginsRT/MaxPoolingFixedSizeRT.h" -class PluginFactory : IPluginFactory -{ -public: - YoloRT *yolos[16]; - int n_yolos; - - virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength); -}; - - +namespace tk { namespace dnn { class NetworkRT { @@ -57,11 +34,11 @@ public: nvinfer1::DataType dtRT; nvinfer1::IBuilder *builderRT; nvinfer1::IRuntime *runtimeRT; - nvinfer1::INetworkDefinition *networkRT; -#if NV_TENSORRT_MAJOR >= 6 + nvinfer1::INetworkDefinition *networkRT; +#if NV_TENSORRT_MAJOR >= 6 nvinfer1::IBuilderConfig *configRT; #endif - + nvinfer1::ICudaEngine *engineRT; nvinfer1::IExecutionContext *contextRT; @@ -74,8 +51,6 @@ public: dnnType *output; cudaStream_t stream; - PluginFactory *pluginFactory; - NetworkRT(Network *net, const char *name); virtual ~NetworkRT(); @@ -89,7 +64,7 @@ public: int getBuffersN() { if(engineRT != nullptr) return engineRT->getNbBindings(); - else + else return 0; } @@ -97,7 +72,7 @@ public: Do inference */ dnnType* infer(dataDim_t &dim, dnnType* data); - void enqueue(int batchSize = 1); + void enqueue(int batchSize = 1); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l); diff --git a/include/tkDNN/buffer_func.h b/include/tkDNN/buffer_func.h new file mode 100644 index 0000000..e697d34 --- /dev/null +++ b/include/tkDNN/buffer_func.h @@ -0,0 +1,20 @@ +#ifndef BUFFER_FUNC_H +#define BUFFER_FUNC_H + +namespace tk { namespace dnn { + +template void writeBUF(char*& buffer, const T& val) +{ + *reinterpret_cast(buffer) = val; + buffer += sizeof(T); +} + +template T readBUF(const char*& buffer) +{ + T val = *reinterpret_cast(buffer); + buffer += sizeof(T); + return val; +} +}} + +#endif // BUFFER_FUNC_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationLeakyRT.h b/include/tkDNN/pluginsRT/ActivationLeakyRT.h index 330ed37..ffa564e 100644 --- a/include/tkDNN/pluginsRT/ActivationLeakyRT.h +++ b/include/tkDNN/pluginsRT/ActivationLeakyRT.h @@ -1,7 +1,19 @@ -#include -#include "../kernels.h" +#ifndef ACTIVATION_LEAKY_RT_H +#define ACTIVATION_LEAKY_RT_H -class ActivationLeakyRT : public IPlugin { +#if NV_TENSORRT_MAJOR < 6 + +#include +#include + +#include + +#include "../kernels.h" +#include "../buffer_func.h" + +namespace tk { namespace dnn { + +class ActivationLeakyRT final : public IPlugin { public: ActivationLeakyRT(float s) { @@ -31,31 +43,36 @@ public: 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; } - 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(inputs[0]), + activationLEAKYForward((dnnType*)reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), batchSize*size, slope, stream); return 0; } - virtual size_t getSerializationSize() override { + size_t getSerializationSize() override { return 1*sizeof(int) + 1*sizeof(float); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, size); + writeBUF(buf, size); assert(buf == a + getSerializationSize()); } int size; float slope; }; +}} + +#endif + +#endif // ACTIVATION_LEAKY_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationLogisticRT.h b/include/tkDNN/pluginsRT/ActivationLogisticRT.h index 063931f..48b5bbc 100644 --- a/include/tkDNN/pluginsRT/ActivationLogisticRT.h +++ b/include/tkDNN/pluginsRT/ActivationLogisticRT.h @@ -1,60 +1,149 @@ -#include -#include "../kernels.h" +#ifndef ACTIVATION_LOGISTIC_RT_H +#define ACTIVATION_LOGISTIC_RT_H -class ActivationLogisticRT : public IPlugin { +#include +#include + +#include +#include + +#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: - ActivationLogisticRT() { + ActivationLogisticRT() = default; + ~ActivationLogisticRT() = default; - } - - ~ActivationLogisticRT(){ - - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { 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]; } - 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; for(int i=0; i(inputs[0]), reinterpret_cast(outputs[0]), batchSize*size, stream); return 0; } - - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 1*sizeof(int); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer); - tk::dnn::writeBUF(buf, size); + writeBUF(buf, size); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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; }; + +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(serialData),*bufCheck = buf; + ActivationLogisticRT *a = new ActivationLogisticRT(); + a->size = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // ACTIVATION_LOGISTIC_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationMishRT.h b/include/tkDNN/pluginsRT/ActivationMishRT.h index 5d660af..6f821d7 100644 --- a/include/tkDNN/pluginsRT/ActivationMishRT.h +++ b/include/tkDNN/pluginsRT/ActivationMishRT.h @@ -1,61 +1,150 @@ -#include -#include "../kernels.h" +#ifndef ACTIVATION_MISH_RT_H +#define ACTIVATION_MISH_RT_H -class ActivationMishRT : public IPlugin { +#include +#include + +#include +#include + +#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: - ActivationMishRT() { + ActivationMishRT() = default; + ~ActivationMishRT() = default; - } - - ~ActivationMishRT(){ - - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { 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]; } - 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; for(int i=0; i(inputs[0]), + int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override { + activationMishForward((dnnType*)reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), batchSize*size, stream); return 0; } - - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 1*sizeof(int); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, size); + writeBUF(buf, size); assert(buf == a + getSerializationSize()); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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; }; + +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(serialData),*bufCheck = buf; + ActivationMishRT *a = new ActivationMishRT(); + a->size = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // ACTIVATION_MISH_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h index 50ceb81..fa2c782 100644 --- a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h +++ b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h @@ -1,63 +1,156 @@ -#include -#include "../kernels.h" +#ifndef ACTIVATION_RELU_CEILING_RT_H +#define ACTIVATION_RELU_CEILING_RT_H -class ActivationReLUCeiling : public IPlugin { +#include +#include + +#include +#include + +#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: ActivationReLUCeiling(const float ceiling) { this->ceiling = ceiling; } - ~ActivationReLUCeiling(){ + ~ActivationReLUCeiling() = default; - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { 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]; } - 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; for(int i=0; i(inputs[0]), + int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override { + activationReLUCeilingForward((dnnType*)reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), batchSize*size, ceiling, stream); return 0; } - - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 1*sizeof(int) + 1*sizeof(float); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, ceiling); - tk::dnn::writeBUF(buf, size); + writeBUF(buf, ceiling); + writeBUF(buf, size); assert(buf = a + getSerializationSize()); - + } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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; 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(serialData),*bufCheck = buf; + float activationReluTemp = readBUF(buf); + ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp); + a->size = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // ACTIVATION_RELU_CEILING_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ActivationSigmoidRT.h b/include/tkDNN/pluginsRT/ActivationSigmoidRT.h index bcc58c7..0731759 100644 --- a/include/tkDNN/pluginsRT/ActivationSigmoidRT.h +++ b/include/tkDNN/pluginsRT/ActivationSigmoidRT.h @@ -1,61 +1,149 @@ -#include -#include "../kernels.h" +#ifndef ACTIVATION_SIGMOID_RT_H +#define ACTIVATION_SIGMOID_RT_H -class ActivationSigmoidRT : public IPlugin { +#include +#include + +#include +#include + +#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: - ActivationSigmoidRT() { + ActivationSigmoidRT() = default; + ~ActivationSigmoidRT() = default; - } - - ~ActivationSigmoidRT(){ - - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { 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]; } - 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; for(int i=0; i(inputs[0]), + int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override { + activationSIGMOIDForward((dnnType*)reinterpret_cast(inputs[0]), reinterpret_cast(outputs[0]), batchSize*size, stream); return 0; } - - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 1*sizeof(int); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, size); + writeBUF(buf, size); assert(buf == a + getSerializationSize()); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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; }; + +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(serialData),*bufCheck = buf; + ActivationSigmoidRT* a = new ActivationSigmoidRT(); + a->size = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // ACTIVATION_SIGMOID_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index 711198f..13abba3 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -1,16 +1,27 @@ -#include +#ifndef DEFORMABLE_CONV_RT_H +#define DEFORMABLE_CONV_RT_H + +#include +#include + +#include +#include + #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: - DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw, - int deformableGroup, int i_n, int i_c, int i_h, int i_w, - int o_n, int o_c, int o_h, int o_w, - tk::dnn::DeformConv2d *deformable = nullptr) { + DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw, + int deformableGroup, int i_n, int i_c, int i_h, int i_w, + int o_n, int o_c, int o_h, int o_w, + DeformConv2d *deformable = nullptr) { this->chunk_dim = chunk_dim; this->kh = kh; this->kw = kw; @@ -30,7 +41,7 @@ public: height_ones = (i_h + 2 * ph - (1 * (kh - 1) + 1)) / sh + 1; width_ones = (i_w + 2 * pw - (1 * (kw - 1) + 1)) / sw + 1; dim_ones = i_c * kh * kw * 1 * height_ones * width_ones; - + checkCuda( cudaMalloc(&data_d, i_c * o_c * kh * kw * 1 * sizeof(dnnType))); checkCuda( cudaMalloc(&bias2_d, o_c*sizeof(dnnType))); checkCuda( cudaMalloc(&ones_d1, height_ones * width_ones * sizeof(dnnType))); @@ -61,38 +72,45 @@ public: cublasDestroy(handle); } - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { - return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w}; + nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override { + 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; } - 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; } - 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(inputs[0]); dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]); // split conv2d outputs into offset to mask for(int b=0; b(buffer),*a=buf; - tk::dnn::writeBUF(buf, chunk_dim); - tk::dnn::writeBUF(buf, kh); - tk::dnn::writeBUF(buf, kw); - tk::dnn::writeBUF(buf, sh); - tk::dnn::writeBUF(buf, sw); - tk::dnn::writeBUF(buf, ph); - tk::dnn::writeBUF(buf, pw); - tk::dnn::writeBUF(buf, deformableGroup); - tk::dnn::writeBUF(buf, i_n); - tk::dnn::writeBUF(buf, i_c); - tk::dnn::writeBUF(buf, i_h); - tk::dnn::writeBUF(buf, i_w); - tk::dnn::writeBUF(buf, o_n); - tk::dnn::writeBUF(buf, o_c); - tk::dnn::writeBUF(buf, o_h); - tk::dnn::writeBUF(buf, o_w); + writeBUF(buf, chunk_dim); + writeBUF(buf, kh); + writeBUF(buf, kw); + writeBUF(buf, sh); + writeBUF(buf, sw); + writeBUF(buf, ph); + writeBUF(buf, pw); + writeBUF(buf, deformableGroup); + writeBUF(buf, i_n); + writeBUF(buf, i_c); + writeBUF(buf, i_h); + writeBUF(buf, i_w); + writeBUF(buf, o_n); + writeBUF(buf, o_c); + writeBUF(buf, o_h); + writeBUF(buf, o_w); dnnType *aus = new dnnType[chunk_dim*2]; checkCuda( cudaMemcpy(aus, offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToHost) ); for(int i=0; i(serialData),*bufCheck = buf; + int chuck_dimTemp = readBUF(buf); + int khTemp = readBUF(buf); + int kwTemp = readBUF(buf); + int shTemp = readBUF(buf); + int swTemp = readBUF(buf); + int phTemp = readBUF(buf); + int pwTemp = readBUF(buf); + int deformableGroupTemp = readBUF(buf); + int i_nTemp = readBUF(buf); + int i_cTemp = readBUF(buf); + int i_hTemp = readBUF(buf); + int i_wTemp = readBUF(buf); + int o_nTemp = readBUF(buf); + int o_cTemp = readBUF(buf); + int o_hTemp = readBUF(buf); + int o_wTemp = readBUF(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; ichunk_dim*2; i++) + aus[i] = readBUF(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; ichunk_dim; i++) + aus[i] = readBUF(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(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(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; iheight_ones * r->width_ones; i++) + aus[i] = readBUF(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; idim_ones; i++) + aus[i] = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // DEFORMABLE_CONV_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/FlattenConcatRT.h b/include/tkDNN/pluginsRT/FlattenConcatRT.h index 51aa1ab..8c1173d 100644 --- a/include/tkDNN/pluginsRT/FlattenConcatRT.h +++ b/include/tkDNN/pluginsRT/FlattenConcatRT.h @@ -1,6 +1,20 @@ -#include +#ifndef FLATTEN_CONCAT_RT_H +#define FLATTEN_CONCAT_RT_H -class FlattenConcatRT : public IPlugin { +#include +#include + +#include +#include + +#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: FlattenConcatRT() { @@ -11,19 +25,23 @@ public: } } - ~FlattenConcatRT(){ + ~FlattenConcatRT() = default; - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { - return DimsCHW{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1}; + nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override { + 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); rows = inputDims[0].d[0]; cols = inputDims[0].d[1] * inputDims[0].d[2]; @@ -32,24 +50,24 @@ public: w = 1; } - int initialize() override { + int initialize() noexcept override { return 0; } - virtual void terminate() override { + void terminate() noexcept override { checkERROR(cublasDestroy(handle)); } - virtual size_t getWorkspaceSize(int maxBatchSize) const override { + size_t getWorkspaceSize(int maxBatchSize) const noexcept override { 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(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - checkERROR( cublasSetStream(handle, stream) ); + checkERROR( cublasSetStream(handle, stream) ); for(int i=0; i(buffer),*a = buf; - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); - tk::dnn::writeBUF(buf, rows); - tk::dnn::writeBUF(buf, cols); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); + writeBUF(buf, rows); + writeBUF(buf, cols); assert(buf == a + getSerializationSize()); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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 rows, cols; - cublasStatus_t stat; - cublasHandle_t handle; + cublasStatus_t stat; + 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(serialData),*bufCheck = buf; + float activationReluTemp = readBUF(buf); + FlattenConcatRT *r = new FlattenConcatRT(); + r->c = readBUF(buf); + r->h = readBUF(buf); + r->w = readBUF(buf); + r->rows = readBUF(buf); + r->cols = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // FLATTEN_CONCAT_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h index 0899a34..69427e4 100644 --- a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h +++ b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h @@ -1,11 +1,24 @@ -#include -#include "../kernels.h" +#ifndef MAX_POOLING_FIXED_SIZE_RT_H +#define MAX_POOLING_FIXED_SIZE_RT_H -class MaxPoolFixedSizeRT : public IPlugin { +#include +#include + +#include +#include + +#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: MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) { - this->c = c; + this->c = c; this->h = h; this->w = w; this->n = n; @@ -15,33 +28,37 @@ public: this->padding = padding; } - ~MaxPoolFixedSizeRT(){ - } + ~MaxPoolFixedSizeRT() = default; - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { - return DimsCHW{this->c, this->h, this->w}; + nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override { + 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; } - 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; } - 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<n<<" "<c<<" "<h<<" "<w<<" "<stride_H<<" "<stride_W<<" "<winSize<<" "<padding<(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); @@ -49,27 +66,112 @@ public: return 0; } - - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 8*sizeof(int); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, this->c); - tk::dnn::writeBUF(buf, this->h); - tk::dnn::writeBUF(buf, this->w); - tk::dnn::writeBUF(buf, this->n); - tk::dnn::writeBUF(buf, this->stride_H); - tk::dnn::writeBUF(buf, this->stride_W); - tk::dnn::writeBUF(buf, this->winSize); - tk::dnn::writeBUF(buf, this->padding); + writeBUF(buf, this->c); + writeBUF(buf, this->h); + writeBUF(buf, this->w); + writeBUF(buf, this->n); + writeBUF(buf, this->stride_H); + writeBUF(buf, this->stride_W); + writeBUF(buf, this->winSize); + writeBUF(buf, this->padding); assert(buf == a + getSerializationSize()); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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 stride_H, stride_W; int winSize; 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(serialData),*bufCheck = buf; + int cTemp = readBUF(buf); + int hTemp = readBUF(buf); + int wTemp = readBUF(buf); + int nTemp = readBUF(buf); + int strideHTemp = readBUF(buf); + int strideWTemp = readBUF(buf); + int winSizeTemp = readBUF(buf); + int paddingTemp = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // MAX_POOLING_FIXED_SIZE_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h index 8487652..0ddc61e 100644 --- a/include/tkDNN/pluginsRT/RegionRT.h +++ b/include/tkDNN/pluginsRT/RegionRT.h @@ -1,48 +1,62 @@ -#include -#include "../kernels.h" +#ifndef REGION_RT_H +#define REGION_RT_H -class RegionRT : public IPlugin { +#include +#include + +#include +#include + +#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: RegionRT(int classes, int coords, int num) { - this->classes = classes; this->coords = coords; this->num = num; } - ~RegionRT(){ + ~RegionRT() = default; - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { 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]; } - 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]; h = inputDims[0].d[1]; w = inputDims[0].d[2]; } - int initialize() override { - + int initialize() noexcept override { 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; } - 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(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); @@ -52,7 +66,7 @@ public: for(int n = 0; n < num; ++n){ int index = entry_index(b, n*w*h, 0); activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); - + index = entry_index(b, n*w*h, coords); activationLOGISTICForward(srcData + index, dstData + index, w*h, stream); } @@ -60,31 +74,61 @@ public: //softmax start int index = entry_index(0, 0, coords + 1); - softmaxForward( srcData + index, classes, batchSize*num, - (c*h*w)/num, + softmaxForward( srcData + index, classes, batchSize*num, + (c*h*w)/num, w*h, 1, w*h, 1, dstData + index, stream); return 0; } - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 6*sizeof(int); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, classes); - tk::dnn::writeBUF(buf, coords); - tk::dnn::writeBUF(buf, num); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); + writeBUF(buf, classes); + writeBUF(buf, coords); + writeBUF(buf, num); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); assert(buf == a + getSerializationSize()); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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 classes, coords, num; + int classes, coords, num; int entry_index(int batch, int location, int entry) { int n = location / (w*h); @@ -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(serialData),*bufCheck = buf; + int classesTemp = readBUF(buf); + int coordsTemp = readBUF(buf); + int numTemp = readBUF(buf); + RegionRT* r = new RegionRT(classesTemp, coordsTemp, numTemp); + + r->c = readBUF(buf); + r->h = readBUF(buf); + r->w = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // REGION_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ReorgRT.h b/include/tkDNN/pluginsRT/ReorgRT.h index c1b529a..bbda446 100644 --- a/include/tkDNN/pluginsRT/ReorgRT.h +++ b/include/tkDNN/pluginsRT/ReorgRT.h @@ -1,64 +1,159 @@ -#include -#include "../kernels.h" +#ifndef REORG_RT_H +#define REORG_RT_H -class ReorgRT : public IPlugin { +#include +#include + +#include +#include + +#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: ReorgRT(int stride) { this->stride = stride; } - ~ReorgRT(){ + ~ReorgRT() = default; - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { - return DimsCHW{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride}; + nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override { + 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]; h = inputDims[0].d[1]; w = inputDims[0].d[2]; } - int initialize() override { - + int initialize() noexcept override { 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; } - virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { - - reorgForward((dnnType*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), + int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override { + reorgForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize, c, h, w, stride, stream); return 0; } - - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 4*sizeof(int); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, stride); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); + writeBUF(buf, stride); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); assert(buf == a + getSerializationSize()); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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; }; + +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(serialData),*bufCheck = buf; + int strideTemp = readBUF(buf); + ReorgRT *r = new ReorgRT(strideTemp); + r->c = readBUF(buf); + r->h = readBUF(buf); + r->w = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // REORG_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ReshapeRT.h b/include/tkDNN/pluginsRT/ReshapeRT.h index 37017c7..18b065d 100644 --- a/include/tkDNN/pluginsRT/ReshapeRT.h +++ b/include/tkDNN/pluginsRT/ReshapeRT.h @@ -1,6 +1,21 @@ -#include +#ifndef RESHAPE_RT_H +#define RESHAPE_RT_H -class ReshapeRT : public IPlugin { +#include +#include + +#include +#include + +#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: ReshapeRT(dataDim_t new_dim) { @@ -10,33 +25,37 @@ public: w = new_dim.w; } - ~ReshapeRT(){ + ~ReshapeRT() = default; - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { - return DimsCHW{ c,h,w}; + nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override { + 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; } - 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; } - 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(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); @@ -44,19 +63,100 @@ public: return 0; } - - virtual size_t getSerializationSize() override { + virtual size_t getSerializationSize() const noexcept override { return 4*sizeof(int); } - virtual void serialize(void* buffer) override { + virtual void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a = buf; - tk::dnn::writeBUF(buf, n); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); + writeBUF(buf, n); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); assert(buf == a + getSerializationSize()); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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; }; + +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(serialData),*bufCheck = buf; + dataDim_t new_dim; + new_dim.n = readBUF(buf); + new_dim.c = readBUF(buf); + new_dim.h = readBUF(buf); + new_dim.w = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // RESHAPE_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ResizeLayerRT.h b/include/tkDNN/pluginsRT/ResizeLayerRT.h index cde52bf..bc5e121 100644 --- a/include/tkDNN/pluginsRT/ResizeLayerRT.h +++ b/include/tkDNN/pluginsRT/ResizeLayerRT.h @@ -1,68 +1,168 @@ -#include -#include "../kernels.h" +#ifndef RESIZE_LAYER_RT_H +#define RESIZE_LAYER_RT_H -class ResizeLayerRT : public IPlugin { +#include +#include + +#include +#include + +#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: ResizeLayerRT(int c, int h, int w) { o_c = c; o_h = h; - o_w = w; + o_w = w; } - ~ResizeLayerRT(){ - } + ~ResizeLayerRT() = default; - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { - return DimsCHW{o_c, o_h, o_w}; + nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override { + 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_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; } - 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; } - 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); - resizeForward((dnnType*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), + resizeForward((dnnType*)reinterpret_cast(inputs[0]), + reinterpret_cast(outputs[0]), batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream); return 0; } - - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 6*sizeof(int); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, o_c); - tk::dnn::writeBUF(buf, o_h); - tk::dnn::writeBUF(buf, o_w); + writeBUF(buf, o_c); + writeBUF(buf, o_h); + writeBUF(buf, o_w); - tk::dnn::writeBUF(buf, i_c); - tk::dnn::writeBUF(buf, i_h); - tk::dnn::writeBUF(buf, i_w); + writeBUF(buf, i_c); + writeBUF(buf, i_h); + writeBUF(buf, i_w); assert(buf == a + getSerializationSize()); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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; }; + +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(serialData),*bufCheck = buf; + int o_cTemp = readBUF(buf); + int o_hTemp = readBUF(buf); + int o_wTemp = readBUF(buf); + ResizeLayerRT* r = new ResizeLayerRT(o_cTemp, o_hTemp, o_wTemp); + + r->i_c = readBUF(buf); + r->i_h = readBUF(buf); + r->i_w = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // RESIZE_LAYER_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h index 5a8c170..edb5d96 100644 --- a/include/tkDNN/pluginsRT/RouteRT.h +++ b/include/tkDNN/pluginsRT/RouteRT.h @@ -1,7 +1,20 @@ -#include -#include "../kernels.h" +#ifndef ROUTE_RT_H +#define ROUTE_RT_H -class RouteRT : public IPlugin { +#include +#include + +#include +#include + +#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 @@ -13,25 +26,29 @@ public: this->group_id = group_id; } - ~RouteRT(){ + ~RouteRT() = default; - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { 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; for(int i=0; i(outputs[0]); for(int b=0; b(buffer),*a=buf; - tk::dnn::writeBUF(buf, groups); - tk::dnn::writeBUF(buf, group_id); - tk::dnn::writeBUF(buf, in); + writeBUF(buf, groups); + writeBUF(buf, group_id); + writeBUF(buf, in); for(int i=0; i(serialData),*bufCheck = buf; + int groupsTemp = readBUF(buf); + int group_idTemp = readBUF(buf); + RouteRT* r = new RouteRT(groupsTemp, group_idTemp); + r->in = readBUF(buf); + for(int i=0; ic_in[i] = readBUF(buf); + r->c = readBUF(buf); + r->h = readBUF(buf); + r->w = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // ROUTE_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h index 04091ac..22f73f2 100644 --- a/include/tkDNN/pluginsRT/ShortcutRT.h +++ b/include/tkDNN/pluginsRT/ShortcutRT.h @@ -1,48 +1,64 @@ -#include -#include "../kernels.h" +#ifndef SHORTCUT_RT_H +#define SHORTCUT_RT_H -class ShortcutRT : public IPlugin { +#include +#include + +#include +#include + +#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: - ShortcutRT(tk::dnn::dataDim_t bdim, bool mul) { + ShortcutRT(dataDim_t bdim, bool mul) { this->bc = bdim.c; this->bh = bdim.h; this->bw = bdim.w; this->mul = mul; } - ~ShortcutRT(){ + ~ShortcutRT() = default; - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { - return DimsCHW{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]}; + nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override { + 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]; h = inputDims[0].d[1]; w = inputDims[0].d[2]; } - int initialize() override { - + int initialize() noexcept override { 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; } - 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(inputs[0]); dnnType *srcDataBack = (dnnType*)reinterpret_cast(inputs[1]); dnnType *dstData = reinterpret_cast(outputs[0]); @@ -54,24 +70,110 @@ public: } - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 6*sizeof(int) + sizeof(bool); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, bc); - tk::dnn::writeBUF(buf, bh); - tk::dnn::writeBUF(buf, bw); - tk::dnn::writeBUF(buf, mul); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); + writeBUF(buf, bc); + writeBUF(buf, bh); + writeBUF(buf, bw); + writeBUF(buf, mul); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); assert(buf == a + getSerializationSize()); - + } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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 bc, bh, bw; 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(serialData),*bufCheck = buf; + dataDim_t bdim; + bdim.c = readBUF(buf); + bdim.h = readBUF(buf); + bdim.w = readBUF(buf); + bdim.l = 1; + + ShortcutRT *r = new ShortcutRT(bdim, readBUF(buf)); + r->c = readBUF(buf); + r->h = readBUF(buf); + r->w = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // SHORTCUT_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/UpsampleRT.h b/include/tkDNN/pluginsRT/UpsampleRT.h index a11d7b4..0b91988 100644 --- a/include/tkDNN/pluginsRT/UpsampleRT.h +++ b/include/tkDNN/pluginsRT/UpsampleRT.h @@ -1,66 +1,160 @@ -#include -#include "../kernels.h" +#ifndef UPSAMPLE_RT_H +#define UPSAMPLE_RT_H -class UpsampleRT : public IPlugin { +#include +#include + +#include +#include + +#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: UpsampleRT(int stride) { this->stride = stride; } - ~UpsampleRT(){ + ~UpsampleRT() = default; - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { return 1; } - Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { - return DimsCHW(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride); + nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override { + 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]; h = inputDims[0].d[1]; w = inputDims[0].d[2]; } - int initialize() override { - + int initialize() noexcept override { 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; } - 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(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); - + fill(dstData, batchSize*c*h*w*stride*stride, 0.0, stream); upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream); return 0; } - - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 4*sizeof(int); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, stride); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); + writeBUF(buf, stride); + writeBUF(buf, c); + writeBUF(buf, h); + writeBUF(buf, w); assert(buf == a + getSerializationSize()); } + // Extra IPluginV2 overrides + bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override { + return true; + } + + 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; }; + +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(serialData),*bufCheck = buf; + int strideTemp = readBUF(buf); + UpsampleRT* r = new UpsampleRT(strideTemp); + r->c = readBUF(buf); + r->h = readBUF(buf); + r->w = readBUF(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 mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // UPSAMPLE_RT_H \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index 5ffe39c..ae1eb40 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -1,15 +1,26 @@ -#include +#ifndef YOLO_RT_H +#define YOLO_RT_H + +#include +#include + +#include +#include + +#include "../yoloContainer.h" #include "../kernels.h" +#include "../buffer_func.h" +#include "../Layer.h" #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: - 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->num = num; this->n_masks = n_masks; @@ -27,38 +38,40 @@ public: } } - ~YoloRT(){ + ~YoloRT() = default; - } - - int getNbOutputs() const override { + int getNbOutputs() const noexcept override { 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]; } - 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]; h = inputDims[0].d[1]; w = inputDims[0].d[2]; } - int initialize() override { - + int initialize() noexcept override { 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; } - 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(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); @@ -86,30 +99,29 @@ public: return 0; } - - virtual size_t getSerializationSize() override { + size_t getSerializationSize() const noexcept override { return 8*sizeof(int) + 2*sizeof(float)+ n_masks*sizeof(dnnType) + num*n_masks*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char); } - virtual void serialize(void* buffer) override { + void serialize(void* buffer) const noexcept override { char *buf = reinterpret_cast(buffer),*a=buf; - tk::dnn::writeBUF(buf, classes); //std::cout << "Classes :" << classes << std::endl; - tk::dnn::writeBUF(buf, num); //std::cout << "Num : " << num << std::endl; - tk::dnn::writeBUF(buf, n_masks); //std::cout << "N_Masks" << n_masks << std::endl; - tk::dnn::writeBUF(buf, scaleXY); //std::cout << "ScaleXY :" << scaleXY << std::endl; - tk::dnn::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; - tk::dnn::writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl; - tk::dnn::writeBUF(buf, c); //std::cout << "C : " << c << std::endl; - tk::dnn::writeBUF(buf, h); //std::cout << "H : " << h << std::endl; - tk::dnn::writeBUF(buf, w); //std::cout << "C : " << c << std::endl; + writeBUF(buf, classes); //std::cout << "Classes :" << classes << std::endl; + writeBUF(buf, num); //std::cout << "Num : " << num << std::endl; + writeBUF(buf, n_masks); //std::cout << "N_Masks" << n_masks << std::endl; + writeBUF(buf, scaleXY); //std::cout << "ScaleXY :" << scaleXY << std::endl; + writeBUF(buf, nms_thresh); //std::cout << "nms_thresh :" << nms_thresh << std::endl; + writeBUF(buf, nms_kind); //std::cout << "nms_kind : " << nms_kind << std::endl; + writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl; + writeBUF(buf, c); //std::cout << "C : " << c << std::endl; + writeBUF(buf, h); //std::cout << "H : " << h << std::endl; + writeBUF(buf, w); //std::cout << "C : " << c << std::endl; 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++) { - 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 @@ -117,12 +129,42 @@ public: char tmp[YOLORT_CLASSNAME_W]; strcpy(tmp, classesNames[i].c_str()); for(int j=0; j mPluginAttributes; + std::string mNamespace; +}; +}} +#undef PLUGIN_NAME +#undef PLUGIN_VERSION + +#endif // YOLO_RT_H \ No newline at end of file diff --git a/include/tkDNN/yoloContainer.h b/include/tkDNN/yoloContainer.h new file mode 100644 index 0000000..39e830f --- /dev/null +++ b/include/tkDNN/yoloContainer.h @@ -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 \ No newline at end of file diff --git a/src/Int8BatchStream.cpp b/src/Int8BatchStream.cpp index fdc1db2..042f637 100644 --- a/src/Int8BatchStream.cpp +++ b/src/Int8BatchStream.cpp @@ -8,14 +8,14 @@ BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, const std::string& fileimglist, const std::string& filelabellist) { mBatchSize = batchSize; 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; mWidth = dim.w; - mImageSize = mDims.c()*mDims.h()*mDims.w(); + mImageSize = dim.c*dim.h*dim.w; mBatch.resize(mBatchSize*mImageSize, 0); mLabels.resize(mBatchSize, 0); - mFileBatch.resize(mDims.n()*mImageSize, 0); - mFileLabels.resize(mDims.n(), 0); + mFileBatch.resize(dim.n*mImageSize, 0); + mFileLabels.resize(dim.n, 0); mFileImgList = fileimglist; readInListFile(fileimglist, mListImg); mFileLabelList = filelabellist; @@ -27,7 +27,7 @@ BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, void BatchStream::reset(int firstBatch) { mBatchCount = 0; mFileCount = 0; - mFileBatchPos = mDims.n(); + mFileBatchPos = mDims.d[0]; skip(firstBatch); } @@ -37,11 +37,11 @@ bool BatchStream::next() { return false; for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize) { - assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.n()); - if (mFileBatchPos == mDims.n() && !update()) + assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.d[0]); + if (mFileBatchPos == mDims.d[0] && !update()) 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(getFileLabels() + mFileBatchPos, csize, getLabels() + batchPos); } @@ -50,8 +50,8 @@ bool BatchStream::next() { } void BatchStream::skip(int skipCount) { - if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n()) { - mFileCount += skipCount * mBatchSize / mDims.n(); + if (mBatchSize >= mDims.d[0] && mBatchSize%mDims.d[0] == 0 && mFileBatchPos == mDims.d[0]) { + mFileCount += skipCount * mBatchSize / mDims.d[0]; return; } @@ -67,7 +67,7 @@ void BatchStream::readInListFile(const std::string& dataFilePath, std::vector& res void BatchStream::readLabels(std::string inputFileName, std::vector& ris) { std::ifstream is(inputFileName.c_str()); - + std::string line; while (std::getline(is, line)) { @@ -153,13 +153,13 @@ bool BatchStream::update() { readCVimage(imgFileName, mFileBatch); // std::transform( // singleImg_rawData.begin(), singleImg_rawData.end(), mFileBatch.begin(), [](uint8_t val) { return static_cast(val); }); - + //read label mFileLabels.clear(); readLabels(labelFileName, mFileLabels); // std::transform( // singleLabels_rawData.begin(), singleLabels_rawData.end(), mFileLabels.begin(), [](uint8_t val) { return static_cast(val); }); - + mFileBatchPos = 0; return true; } diff --git a/src/Int8Calibrator.cpp b/src/Int8Calibrator.cpp index 773a9d8..1ead6fa 100644 --- a/src/Int8Calibrator.cpp +++ b/src/Int8Calibrator.cpp @@ -1,20 +1,20 @@ #include "Int8Calibrator.h" -Int8EntropyCalibrator::Int8EntropyCalibrator(BatchStream& stream, int firstBatch, +Int8EntropyCalibrator::Int8EntropyCalibrator(BatchStream& stream, int firstBatch, const std::string& calibTableFilePath, const std::string& inputBlobName, - bool readCache): - mStream(stream), + bool readCache): + mStream(stream), mCalibTableFilePath(calibTableFilePath), mInputBlobName(inputBlobName.c_str()), mReadCache(readCache) { - nvinfer1::DimsNCHW dims = mStream.getDims(); - mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w(); + nvinfer1::Dims4 dims = mStream.getDims(); + mInputCount = mStream.getBatchSize() * dims.d[1] * dims.d[2] * dims.d[3]; checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float))); 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()) return false; @@ -24,7 +24,7 @@ bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int return true; } -const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) { +const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) noexcept { mCalibrationCache.clear(); assert(!mCalibTableFilePath.empty()); std::ifstream input(mCalibTableFilePath, std::ios::binary); @@ -38,7 +38,7 @@ const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) { 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()); std::ofstream output(mCalibTableFilePath, std::ios::binary); output.write(reinterpret_cast(cache), length); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 6ac7235..e3365e3 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -15,7 +15,7 @@ using namespace nvinfer1; // Logger for info/warning/errors class Logger : public ILogger { - void log(Severity severity, const char* msg) override { + void log(Severity severity, const char* msg) noexcept override { #ifdef DEBUG std::cout <<"TENSORRT LOG: "<< msg << std::endl; #endif @@ -24,28 +24,28 @@ class Logger : public ILogger { namespace tk { namespace dnn { -std::maptensors; +std::maptensors; NetworkRT::NetworkRT(Network *net, const char *name) { - float rt_ver = float(NV_TENSORRT_MAJOR) + - float(NV_TENSORRT_MINOR)/10 + + float rt_ver = float(NV_TENSORRT_MAJOR) + + float(NV_TENSORRT_MINOR)/10 + float(NV_TENSORRT_PATCH)/100; std::cout<<"New NetworkRT (TensorRT v"<platformHasFastFp16()<<"\n"; std::cout<<"Int8 support: "<platformHasFastInt8()<<"\n"; #if NV_TENSORRT_MAJOR >= 5 std::cout<<"DLAs: "<getNbDLACores()<<"\n"; #endif - networkRT = builderRT->createNetwork(); -#if NV_TENSORRT_MAJOR >= 6 + networkRT = builderRT->createNetworkV2(0u); +#if NV_TENSORRT_MAJOR >= 6 configRT = builderRT->createBuilderConfig(); #endif - + if(!fileExist(name)) { -#if NV_TENSORRT_MAJOR >= 6 +#if NV_TENSORRT_MAJOR >= 6 // Calibrator life time needs to last until after the engine is built. std::unique_ptr calibrator; @@ -53,22 +53,25 @@ NetworkRT::NetworkRT(Network *net, const char *name) { configRT->setMinTimingIterations(1); configRT->setMaxWorkspaceSize(1 << 30); configRT->setFlag(BuilderFlag::kDEBUG); +#else + builderRT->setMaxWorkspaceSize(1 << 30); #endif //input and dataType dataDim_t dim = net->layers[0]->input_dim; dtRT = DataType::kFLOAT; builderRT->setMaxBatchSize(net->maxBatchSize); - builderRT->setMaxWorkspaceSize(1 << 30); if(net->fp16 && builderRT->platformHasFastFp16()) { dtRT = DataType::kHALF; +#if NV_TENSORRT_MAJOR < 6 builderRT->setHalf2Mode(true); -#if NV_TENSORRT_MAJOR >= 6 +#endif +#if NV_TENSORRT_MAJOR >= 6 configRT->setFlag(BuilderFlag::kFP16); #endif } -#if NV_TENSORRT_MAJOR >= 5 +#if NV_TENSORRT_MAJOR >= 5 && NV_TENSORRT_MAJOR < 8 if(net->dla && builderRT->getNbDLACores() > 0) { dtRT = DataType::kHALF; builderRT->setFp16Mode(true); @@ -77,14 +80,23 @@ NetworkRT::NetworkRT(Network *net, const char *name) { builderRT->setDLACore(0); } #endif -#if NV_TENSORRT_MAJOR >= 6 +#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(net->int8 && builderRT->platformHasFastInt8()){ // dtRT = DataType::kINT8; // builderRT->setInt8Mode(true); configRT->setFlag(BuilderFlag::kINT8); - BatchStream calibrationStream(dim, 1, 100, //TODO: check if 100 images are sufficient to the calibration (or 4951) + BatchStream calibrationStream(dim, 1, 100, //TODO: check if 100 images are sufficient to the calibration (or 4951) net->fileImgList, net->fileLabelList); - + /* The calibTableFilePath contains the path+filename of the calibration table. * Each calibration table can be found in the corresponding network folder (../Test/*). * Each network is located in a folder with the same name as the network. @@ -95,33 +107,33 @@ NetworkRT::NetworkRT(Network *net, const char *name) { if(!fileExist((const char *)calib_table_path.c_str())) calib_table_name = "./" + net->networkNameRT.substr(0, net->networkNameRT.find('.')) + "-calibration.table"; - calibrator.reset(new Int8EntropyCalibrator(calibrationStream, 1, - calib_table_name, + calibrator.reset(new Int8EntropyCalibrator(calibrationStream, 1, + calib_table_name, "data")); configRT->setInt8Calibrator(calibrator.get()); } #endif - + // add input layer - ITensor *input = networkRT->addInput("data", DataType::kFLOAT, - DimsCHW{ dim.c, dim.h, dim.w}); + ITensor *input = networkRT->addInput("data", DataType::kFLOAT, + Dims3{ dim.c, dim.h, dim.w}); checkNULL(input); //add other layers for(int i=0; inum_layers; i++) { Layer *l = net->layers[i]; ILayer *Ilay = convert_layer(input, l); -#if NV_TENSORRT_MAJOR >= 6 +#if NV_TENSORRT_MAJOR >= 6 if(net->int8 && builderRT->platformHasFastInt8()) { Ilay->setPrecision(DataType::kINT8); } #endif Ilay->setName( (l->getLayerName() + std::to_string(i)).c_str() ); - + input = Ilay->getOutput(0); input->setName( (l->getLayerName() + std::to_string(i) + "_out").c_str() ); - + if(l->final) networkRT->markOutput(*input); tensors[l] = input; @@ -136,9 +148,9 @@ NetworkRT::NetworkRT(Network *net, const char *name) { std::cout<<"Selected maxBatchSize: "<getMaxBatchSize()<<"\n"; printCudaMemUsage(); std::cout<<"Building tensorRT cuda engine...\n"; -#if NV_TENSORRT_MAJOR >= 6 +#if NV_TENSORRT_MAJOR >= 6 engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT); -#else +#else engineRT = builderRT->buildCudaEngine(*networkRT); //engineRT = std::shared_ptr(builderRT->buildCudaEngine(*networkRT)); #endif @@ -162,7 +174,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { // In order to bind the buffers, we need to know the names of the input and output tensors. // note that indices are guaranteed to be less than IEngine::getNbBindings() - buf_input_idx = engineRT->getBindingIndex("data"); + buf_input_idx = engineRT->getBindingIndex("data"); buf_output_idx = engineRT->getBindingIndex("out"); std::cout<<"input index = "< output index = "<getNbBindings(); i++) { Dims dim = engineRT->getBindingDimensions(i); @@ -261,10 +273,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Dense *l) { //std::cout<<"convert Dense\n"; void *data_b, *bias_b; if(dtRT == DataType::kHALF) { - data_b = l->data16_h; + data_b = l->data16_h; bias_b = l->bias16_h; } else { - data_b = l->data_h; + data_b = l->data_h; bias_b = l->bias_h; } @@ -284,7 +296,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { void *data_b, *bias_b, *bias2_b, *power_b, *mean_b, *variance_b, *scales_b; if(dtRT == DataType::kHALF) { - data_b = l->data16_h; + data_b = l->data16_h; bias_b = l->bias16_h; bias2_b = l->bias216_h; power_b = l->power16_h; @@ -292,7 +304,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { variance_b = l->variance16_h; scales_b = l->scales16_h; } else { - data_b = l->data_h; + data_b = l->data_h; bias_b = l->bias_h; bias2_b = l->bias2_h; power_b = l->power_h; @@ -308,14 +320,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { b = { dtRT, bias_b, l->outputs}; else{ if (l->additional_bias) - b = { dtRT, bias2_b, l->outputs}; + b = { dtRT, bias2_b, l->outputs}; else b = { dtRT, nullptr, 0}; //on batchnorm bias are added later } ILayer *lRT = nullptr; if(!l->deConv) { - IConvolutionLayer *lRTconv = networkRT->addConvolution(*input, + IConvolutionLayer *lRTconv = networkRT->addConvolution(*input, l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b); checkNULL(lRTconv); lRTconv->setStride(DimsHW{l->strideH, l->strideW}); @@ -323,14 +335,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { lRTconv->setNbGroups(l->groups); lRT = (ILayer*) lRTconv; } else { - IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input, + IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input, l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b); checkNULL(lRTconv); lRTconv->setStride(DimsHW{l->strideH, l->strideW}); lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); lRTconv->setNbGroups(l->groups); lRT = (ILayer*) lRTconv; - + Dims d = lRTconv->getOutput(0)->getDimensions(); //std::cout<<"DECONV: "<outputs}; Weights scale{dtRT, variance_b, l->outputs}; // std::cout<getNbOutputs()<addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, + IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, shift, scale, power); - + checkNULL(lRT2); Weights shift2{dtRT, bias_b, l->outputs}; Weights scale2{dtRT, scales_b, l->outputs}; - IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, + IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, shift2, scale2, power); checkNULL(lRT3); @@ -368,8 +380,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) { 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); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + 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); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -381,7 +393,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) { lRT->setPadding(DimsHW{l->paddingH, l->paddingW}); lRT->setStride(DimsHW{l->strideH, l->strideW}); return lRT; - } + } } ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { @@ -389,14 +401,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { if(l->act_mode == ACTIVATION_LEAKY) { //std::cout<<"New plugin LEAKY\n"; - -#if NV_TENSORRT_MAJOR < 6 + +#if NV_TENSORRT_MAJOR < 6 // plugin version - IPlugin *plugin = new ActivationLeakyRT(l->slope); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new ActivationLeakyRT(l->slope); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; -#else +#else IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kLEAKY_RELU); lRT->setAlpha(l->slope); checkNULL(lRT); @@ -413,20 +425,20 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { return lRT; } else if(l->act_mode == CUDNN_ACTIVATION_CLIPPED_RELU) { - IPlugin *plugin = new ActivationReLUCeiling(l->ceiling); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new ActivationReLUCeiling(l->ceiling); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; - } + } else if(l->act_mode == ACTIVATION_MISH) { - IPlugin *plugin = new ActivationMishRT(); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new ActivationMishRT(); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } else if(l->act_mode == ACTIVATION_LOGISTIC) { - IPlugin *plugin = new ActivationLogisticRT(); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new ActivationLogisticRT(); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -447,7 +459,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Softmax *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) { // std::cout<<"convert route\n"; - + ITensor **tens = new ITensor*[l->layers_n]; @@ -460,8 +472,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) { } if(l->groups > 1){ - IPlugin *plugin = new RouteRT(l->groups, l->group_id); - IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin); + auto *plugin = new RouteRT(l->groups, l->group_id); + auto *lRT = networkRT->addPluginV2(tens, l->layers_n, *plugin); checkNULL(lRT); return lRT; } @@ -472,8 +484,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Flatten *l) { - IPlugin *plugin = new FlattenConcatRT(); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new FlattenConcatRT(); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -481,8 +493,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Flatten *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Reshape *l) { // std::cout<<"convert Reshape\n"; - IPlugin *plugin = new ReshapeRT(l->output_dim); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new ReshapeRT(l->output_dim); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -494,7 +506,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Resize *l) { checkNULL(lRT); Dims d{}; 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; } @@ -502,8 +514,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Reorg *l) { //std::cout<<"convert Reorg\n"; //std::cout<<"New plugin REORG\n"; - IPlugin *plugin = new ReorgRT(l->stride); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new ReorgRT(l->stride); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -512,8 +524,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Region *l) { //std::cout<<"convert Region\n"; //std::cout<<"New plugin REGION\n"; - IPlugin *plugin = new RegionRT(l->classes, l->coords, l->num); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new RegionRT(l->classes, l->coords, l->num); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -522,10 +534,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) { //std::cout<<"convert Shortcut\n"; //std::cout<<"New plugin Shortcut\n"; - + ITensor *back_tens = tensors[l->backLayer]; - if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul) + if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul) { IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM); checkNULL(lRT); @@ -534,11 +546,11 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) { else { // 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]; inputs[0] = input; - inputs[1] = back_tens; - IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin); + inputs[1] = back_tens; + auto *lRT = networkRT->addPluginV2(inputs, 2, *plugin); checkNULL(lRT); return lRT; } @@ -548,8 +560,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Yolo *l) { //std::cout<<"convert 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); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY, l->nms_thresh, l->nsm_kind, l->new_coords); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -558,8 +570,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Upsample *l) { //std::cout<<"convert Upsample\n"; //std::cout<<"New plugin UPSAMPLE\n"; - IPlugin *plugin = new UpsampleRT(l->stride); - IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + auto *plugin = new UpsampleRT(l->stride); + auto *lRT = networkRT->addPluginV2(&input, 1, *plugin); checkNULL(lRT); return lRT; } @@ -574,10 +586,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { inputs[1] = preconv->getOutput(0); //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, - l->deformableGroup, l->input_dim.n, l->input_dim.c, l->input_dim.h, l->input_dim.w, + 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->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); lRT->setName( ("Deformable" + std::to_string(l->id)).c_str() ); delete[](inputs); @@ -601,14 +613,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { Weights shift{dtRT, mean_b, l->outputs}; Weights scale{dtRT, variance_b, l->outputs}; //std::cout<getNbOutputs()<addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, + IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, shift, scale, power); - + checkNULL(lRT2); Weights shift2{dtRT, bias_b, l->outputs}; Weights scale2{dtRT, scales_b, l->outputs}; - IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, + IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, shift2, scale2, power); checkNULL(lRT3); @@ -646,53 +658,52 @@ bool NetworkRT::deserialize(const char *filename) { file.close(); } - pluginFactory = new PluginFactory(); runtimeRT = createInferRuntime(loggerRT); - engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory); + engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size); //if (gieModelStream) delete [] gieModelStream; return true; } +// auto* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) { +// const char * buf = reinterpret_cast(serialData),*bufCheck = buf; +// std::string name(layerName); +// //std::cout<(serialData),*bufCheck = buf; - - std::string name(layerName); - //std::cout<(buf)); - a->size = readBUF(buf); - assert(buf == bufCheck + serialLength); - return a; - } - if(name.find("ActivationMish") == 0) { +// #if NV_TENSORRT_MAJOR < 6 +// if(name.find("ActivationLeaky") == 0) { +// ActivationLeakyRT *a = new ActivationLeakyRT(readBUF(buf)); +// a->size = readBUF(buf); +// assert(buf == bufCheck + serialLength); +// return a; +// } +// #endif + /*if(name.find("ActivationMish") == 0) { ActivationMishRT *a = new ActivationMishRT(); a->size = readBUF(buf); assert(buf == bufCheck + serialLength); return a; - } - if(name.find("ActivationLogistic") == 0) { + }*/ + /*if(name.find("ActivationLogistic") == 0) { ActivationLogisticRT *a = new ActivationLogisticRT(); a->size = readBUF(buf); return a; - } - if(name.find("ActivationLogistic") == 0) { + }*/ + /*if(name.find("ActivationLogistic") == 0) { ActivationLogisticRT *a = new ActivationLogisticRT(); a->size = readBUF(buf); return a; - } - if(name.find("ActivationCReLU") == 0) { + }*/ + /*if(name.find("ActivationCReLU") == 0) { float activationReluTemp = readBUF(buf); ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp); a->size = readBUF(buf); assert(buf == bufCheck + serialLength); return a; - } + }*/ - if(name.find("Region") == 0) { + /*if(name.find("Region") == 0) { int classesTemp = readBUF(buf); int coordsTemp = readBUF(buf); int numTemp = readBUF(buf); @@ -703,9 +714,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->w = readBUF(buf); assert(buf == bufCheck + serialLength); return r; - } + }*/ - if(name.find("Reorg") == 0) { + /*if(name.find("Reorg") == 0) { int strideTemp = readBUF(buf); ReorgRT *r = new ReorgRT(strideTemp); r->c = readBUF(buf); @@ -713,9 +724,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->w = readBUF(buf); assert(buf == bufCheck + serialLength); return r; - } + }*/ - if(name.find("Shortcut") == 0) { + /*if(name.find("Shortcut") == 0) { tk::dnn::dataDim_t bdim; bdim.c = readBUF(buf); bdim.h = readBUF(buf); @@ -728,9 +739,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->w = readBUF(buf); return r; assert(buf == bufCheck + serialLength); - } + }*/ - if(name.find("Pooling") == 0) { + /*if(name.find("Pooling") == 0) { int cTemp = readBUF(buf); int hTemp = readBUF(buf); int wTemp = readBUF(buf); @@ -743,9 +754,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp); assert(buf == bufCheck + serialLength); return r; - } + }*/ - if(name.find("Resize") == 0) { + /*if(name.find("Resize") == 0) { int o_cTemp = readBUF(buf); int o_hTemp = readBUF(buf); int o_wTemp = readBUF(buf); @@ -756,10 +767,10 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->i_w = readBUF(buf); assert(buf == bufCheck + serialLength); return r; - } + }*/ - if(name.find("Flatten") == 0) { - FlattenConcatRT *r = new FlattenConcatRT(); + /*if(name.find("Flatten") == 0) { + FlattenConcatRT *r = new FlattenConcatRT(); r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); @@ -767,22 +778,22 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->cols = readBUF(buf); assert(buf == bufCheck + serialLength); return r; - } + }*/ - if(name.find("Reshape") == 0) { + /*if(name.find("Reshape") == 0) { dataDim_t new_dim; new_dim.n = readBUF(buf); new_dim.c = readBUF(buf); new_dim.h = readBUF(buf); new_dim.w = readBUF(buf); - ReshapeRT *r = new ReshapeRT(new_dim); + ReshapeRT *r = new ReshapeRT(new_dim); assert(buf == bufCheck + serialLength); - - return r; - } - if(name.find("Yolo") == 0) { + return r; + }*/ + + /*if(name.find("Yolo") == 0) { int classes_temp = readBUF(buf); int num_temp = readBUF(buf); @@ -792,7 +803,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa int nms_kind_temp = readBUF(buf); int new_coords_temp = readBUF(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); + YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp); @@ -816,8 +827,8 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa yolos[n_yolos++] = r; return r; - } - if(name.find("Upsample") == 0) { + }*/ + /*if(name.find("Upsample") == 0) { int strideTemp = readBUF(buf); UpsampleRT* r = new UpsampleRT(strideTemp); r->c = readBUF(buf); @@ -825,9 +836,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->w = readBUF(buf); assert(buf == bufCheck + serialLength); return r; - } + }*/ - if(name.find("Route") == 0) { + /*if(name.find("Route") == 0) { int groupsTemp = readBUF(buf); int group_idTemp = readBUF(buf); RouteRT* r = new RouteRT(groupsTemp, group_idTemp); @@ -839,9 +850,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->w = readBUF(buf); assert(buf == bufCheck + serialLength); return r; - } + }*/ - if(name.find("Deformable") == 0) { + /*if(name.find("Deformable") == 0) { int chuck_dimTemp = readBUF(buf); int khTemp = readBUF(buf); int kwTemp = readBUF(buf); @@ -892,10 +903,10 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa free(aus); assert(buf == bufCheck + serialLength); return r; - } + }*/ - FatalError("Cant deserialize Plugin"); - return NULL; -} +// FatalError("Cant deserialize Plugin"); +// return NULL; +// } }} diff --git a/src/Yolo3Detection.cpp b/src/Yolo3Detection.cpp index 0c638e6..01e95b4 100644 --- a/src/Yolo3Detection.cpp +++ b/src/Yolo3Detection.cpp @@ -1,5 +1,6 @@ +#include "yoloContainer.h" #include "Yolo3Detection.h" - +#include "pluginsRT/YoloRT.h" namespace tk { namespace dnn { @@ -11,15 +12,15 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c nBatches = n_batches; confThreshold = conf_thresh; - tk::dnn::dataDim_t idim = netRT->input_dim; + tk::dnn::dataDim_t idim = netRT->input_dim; idim.n = nBatches; - if(netRT->pluginFactory->n_yolos < 2 ) { + if(yoloContainer.n_yolos < 2 ) { FatalError("this is not yolo3"); } - for(int i=0; ipluginFactory->n_yolos; i++) { - YoloRT *yRT = netRT->pluginFactory->yolos[i]; + for(int i=0; iclasses; num = yRT->num; nMasks = yRT->n_masks; @@ -43,7 +44,7 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c #endif checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*idim.tot())); - // class colors precompute + // class colors precompute for(int c=0; cclassesNames; return true; -} +} void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){ #ifdef OPENCV_CUDACONTRIB @@ -62,7 +63,7 @@ void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){ orig_img = cv::cuda::GpuMat(frame); cv::cuda::resize(orig_img, img_resized, cv::Size(netRT->input_dim.w, netRT->input_dim.h)); - img_resized.convertTo(imagePreproc, CV_32FC3, 1/255.0); + img_resized.convertTo(imagePreproc, CV_32FC3, 1/255.0); //split channels cv::cuda::split(imagePreproc,bgr);//split source @@ -76,7 +77,7 @@ void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){ } #else cv::resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h)); - frame.convertTo(imagePreproc, CV_32FC3, 1/255.0); + frame.convertTo(imagePreproc, CV_32FC3, 1/255.0); //split channels cv::split(imagePreproc,bgr);//split source @@ -85,7 +86,7 @@ void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){ for(int i=0; iinput_dim.c; i++) { int idx = i*imagePreproc.rows*imagePreproc.cols; int ch = netRT->input_dim.c-1 -i; - memcpy((void*)&input[idx + netRT->input_dim.tot()*bi], (void*)bgr[ch].data, imagePreproc.rows*imagePreproc.cols*sizeof(dnnType)); + memcpy((void*)&input[idx + netRT->input_dim.tot()*bi], (void*)bgr[ch].data, imagePreproc.rows*imagePreproc.cols*sizeof(dnnType)); } checkCuda(cudaMemcpyAsync(input_d + netRT->input_dim.tot()*bi, input + netRT->input_dim.tot()*bi, netRT->input_dim.tot()*sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream)); #endif @@ -95,8 +96,8 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){ //get yolo outputs std::vector rt_out; - //dnnType *rt_out[netRT->pluginFactory->n_yolos]; - for(int i=0; ipluginFactory->n_yolos; i++) + //dnnType *rt_out[yoloContainer.n_yolos]; + for(int i=0; ibuffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi); 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 nDets = 0; - for(int i=0; ipluginFactory->n_yolos; i++) { + for(int i=0; idstData = rt_out[i]; yolo[i]->computeDetections(dets, nDets, netRT->input_dim.w, netRT->input_dim.h, confThreshold, yolo[i]->new_coords); } @@ -124,7 +125,7 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){ x1 = x_ratio*x1; y0 = y_ratio*y0; y1 = y_ratio*y1; - + for(int c=0; c= confThreshold) { int obj_class = c; @@ -140,7 +141,7 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){ // FIXME: this shuld be useless // if(mAP) - // for(int c=0; c0) output[i] = input[i]; else @@ -21,7 +21,7 @@ void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, float { int blocks = (size+255)/256; int threads = 256; - + activation_leaky<<>>(srcData, dstData, size, slope); } diff --git a/src/kernels/activation_logistic.cu b/src/kernels/activation_logistic.cu index a1ca770..46fa964 100644 --- a/src/kernels/activation_logistic.cu +++ b/src/kernels/activation_logistic.cu @@ -1,11 +1,22 @@ +#include + #include "kernels.h" +#include "pluginsRT/ActivationLogisticRT.h" + +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection ActivationLogisticRTCreator::mFC{}; +std::vector ActivationLogisticRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(ActivationLogisticRTCreator); +}} __global__ void activation_logistic(dnnType *input, dnnType *output, int size) { int i = blockDim.x*blockIdx.x + threadIdx.x; - if(i>>(srcData, dstData, size); -} - - +} \ No newline at end of file diff --git a/src/kernels/activation_mish.cu b/src/kernels/activation_mish.cu index 8900061..75a39c9 100644 --- a/src/kernels/activation_mish.cu +++ b/src/kernels/activation_mish.cu @@ -1,40 +1,47 @@ -#include "kernels.h" #include +#include "kernels.h" +#include "pluginsRT/ActivationMishRT.h" #define MISH_THRESHOLD 20 -__device__ +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection ActivationMishRTCreator::mFC{}; +std::vector ActivationMishRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(ActivationMishRTCreator); +}} + +__device__ float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);} -__device__ +__device__ float softplus_kernel(float x, float threshold = 20) { if (x > threshold) return x; // too large else if (x < -threshold) return expf(x); // too small return logf(expf(x) + 1); } +__device__ +float mish_yashas(float x) { + float e = __expf(x); + if (x <= -18.0f) + return x * e; + float n = e * e + 2 * e; + if (x <= -5.0f) + return x * __fdividef(n, n + 2); -__device__ -float mish_yashas(float x) { - float e = __expf(x); - if (x <= -18.0f) - return x * e; - - float n = e * e + 2 * e; - if (x <= -5.0f) - return x * __fdividef(n, n + 2); - - return x - 2 * __fdividef(x, n + 2); -} + return x - 2 * __fdividef(x, n + 2); +} // https://github.com/digantamisra98/Mish // https://github.com/AlexeyAB/darknet/blob/master/src/activation_kernels.cu __global__ void activation_mish(dnnType *input, dnnType *output, int size) { int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; - if (i < size) - // output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD)); + if (i < size) + // output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD)); output[i] = mish_yashas(input[i]); } @@ -45,6 +52,6 @@ void activationMishForward(dnnType* srcData, dnnType* dstData, int size, cudaStr { int blocks = (size+255)/256; int threads = 256; - + activation_mish<<>>(srcData, dstData, size); } \ No newline at end of file diff --git a/src/kernels/activation_relu_ceiling.cu b/src/kernels/activation_relu_ceiling.cu index 72ccf4f..255c3b2 100644 --- a/src/kernels/activation_relu_ceiling.cu +++ b/src/kernels/activation_relu_ceiling.cu @@ -1,11 +1,20 @@ #include "kernels.h" +#include "pluginsRT/ActivationReLUCeilingRT.h" + +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection ActivationReLUCeilingCreator::mFC{}; +std::vector ActivationReLUCeilingCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(ActivationReLUCeilingCreator); +}} __global__ void activation_relu_ceiling(dnnType *input, dnnType *output, int size, const float ceiling) { int i = blockDim.x*blockIdx.x + threadIdx.x; - if(i0) { if (input[i]>ceiling) @@ -26,7 +35,7 @@ void activationReLUCeilingForward(dnnType* srcData, dnnType* dstData, int size, { int blocks = (size+255)/256; int threads = 256; - + activation_relu_ceiling<<>>(srcData, dstData, size, ceiling); } diff --git a/src/kernels/activation_sigmoid.cu b/src/kernels/activation_sigmoid.cu index 400a948..6b6bf98 100644 --- a/src/kernels/activation_sigmoid.cu +++ b/src/kernels/activation_sigmoid.cu @@ -1,6 +1,14 @@ -#include "kernels.h" #include +#include "kernels.h" +#include "pluginsRT/ActivationSigmoidRT.h" +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection ActivationSigmoidRTCreator::mFC{}; +std::vector ActivationSigmoidRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(ActivationSigmoidRTCreator); +}} __global__ void activation_sigmoid(dnnType *input, dnnType *output, int size) { @@ -18,6 +26,6 @@ void activationSIGMOIDForward(dnnType* srcData, dnnType* dstData, int size, cuda { int blocks = (size+255)/256; int threads = 256; - + activation_sigmoid<<>>(srcData, dstData, size); } \ No newline at end of file diff --git a/src/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index 4dbc552..f807caf 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -2,10 +2,12 @@ #include #include #include -#include -#include "kernels.h" +#include #include +#include "kernels.h" +#include "pluginsRT/DeformableConvRT.h" + #define CUDA_KERNEL_LOOP(i, n) \ for (int i = blockIdx.x * blockDim.x + threadIdx.x; \ i < (n); \ @@ -17,6 +19,13 @@ inline int GET_BLOCKS(const int N) return (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS; } +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection DeformableConvRTCreator::mFC{}; +std::vector DeformableConvRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(DeformableConvRTCreator); +}} __device__ __host__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width, const int height, const int width, float h, float w) { @@ -50,7 +59,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n, { //If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ). const int ind_on_w = index / width_col; - const int ind_on_w_on_h = ind_on_w / height_col; + const int ind_on_w_on_h = ind_on_w / height_col; const int kk = 3 * 3; // index index of output matrix const int w_col = index % width_col; @@ -83,10 +92,10 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n, const int iter_member = (i * 3 + j); // const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col; const int data_offset_h_ptr = first_member + s_col2 * iter_member; - + // const int data_offset_w_ptr = ((2 * (i * kernel_w + j) + 1) * height_col + h_col) * width_col + w_col; const int data_offset_w_ptr = s_col + first_member + s_col2 * iter_member; - + // const int data_mask_hw_ptr = ((i * kernel_w + j) * height_col + h_col) * width_col + w_col; const int data_mask_hw_ptr = first_member + s_col * iter_member; @@ -127,7 +136,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int { //If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ). const int ind_on_w = index / width_col; - const int ind_on_w_on_h = ind_on_w / height_col; + const int ind_on_w_on_h = ind_on_w / height_col; const int kk = kernel_h * kernel_w; // index index of output matrix const int w_col = index % width_col; @@ -153,7 +162,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int const float *data_offset_ptr = data_offset + add_ptr + add_ptr; const float *data_mask_ptr = data_mask + add_ptr; - + #pragma unroll for (int i = 0; i < kernel_h; ++i) { #pragma unroll @@ -161,10 +170,10 @@ __global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int const int iter_member = (i * kernel_w + j); // const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col; const int data_offset_h_ptr = first_member + s_col2 * iter_member; - + // const int data_offset_w_ptr = ((2 * (i * kernel_w + j) + 1) * height_col + h_col) * width_col + w_col; const int data_offset_w_ptr = s_col + first_member + s_col2 * iter_member; - + // const int data_mask_hw_ptr = ((i * kernel_w + j) * height_col + h_col) * width_col + w_col; const int data_mask_hw_ptr = first_member + s_col * iter_member; @@ -193,7 +202,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int void modulatedDeformableIm2colCuda(cudaStream_t stream, const float* data_im, const float* data_offset, const float* data_mask, - const int batch_size, const int channels, const int height_im, const int width_im, + const int batch_size, const int channels, const int height_im, const int width_im, const int height_col, const int width_col, const int deformable_group, float* data_col) { // num_axes should be smaller than block size @@ -202,9 +211,9 @@ void modulatedDeformableIm2colCuda(cudaStream_t stream, modulated_deformable_im2col_gpu_kernel <<>>( - num_kernels, data_im, data_offset, data_mask, height_im, width_im, + num_kernels, data_im, data_offset, data_mask, height_im, width_im, batch_size, channels, deformable_group, height_col, width_col, data_col); - + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) FatalError("error in modulatedDeformableIm2colCuda: " + std::string(cudaGetErrorString(err)) + "\n"); @@ -212,9 +221,9 @@ void modulatedDeformableIm2colCuda(cudaStream_t stream, void modulatedDeformableIm2colCudaGeneralVersion(cudaStream_t stream, const float* data_im, const float* data_offset, const float* data_mask, - const int batch_size, const int channels, const int height_im, const int width_im, + const int batch_size, const int channels, const int height_im, const int width_im, const int height_col, const int width_col, const int kernel_h, const int kenerl_w, - const int pad_h, const int pad_w, const int stride_h, const int stride_w, + const int pad_h, const int pad_w, const int stride_h, const int stride_w, const int dilation_h, const int dilation_w, const int deformable_group, float* data_col) { // num_axes should be smaller than block size @@ -226,13 +235,13 @@ void modulatedDeformableIm2colCudaGeneralVersion(cudaStream_t stream, num_kernels, data_im, data_offset, data_mask, height_im, width_im, kernel_h, kenerl_w, pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w, channel_per_deformable_group, batch_size, channels, deformable_group, height_col, width_col, data_col); - + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) FatalError("error in modulatedDeformableIm2colCudaGeneralVersion: " + std::string(cudaGetErrorString(err)) + "\n"); } -void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, +void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, float *input, float *weight, float *bias, float *ones, float *offset, float *mask, @@ -242,11 +251,11 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, const int pad_h, const int pad_w, const int dilation_h, const int dilation_w, const int deformable_group, const int batch_id, - const int in_n, const int in_c, const int in_h, const int in_w, + const int in_n, const int in_c, const int in_h, const int in_w, const int out_n, const int out_c, const int out_h, const int out_w, const int chunk_dim, cudaStream_t stream) -{ - // stat and handle have be moved out to preserve 2 - 6 milliseconds every 100. +{ + // stat and handle have be moved out to preserve 2 - 6 milliseconds every 100. const int batch = batch_id; const int channels = in_c; const int height = in_h; @@ -256,22 +265,22 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, const int height_out = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1; const int width_out = (width + 2 * pad_w - (dilation_w * (kernel_w - 1) + 1)) / stride_w + 1; - + long m = channels_out; long n = height_out * width_out; long k = 1; float alpha = 1.0; float beta = 0.0; - - stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, - n, m, k, &alpha, - ones, k, bias, k, + + stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, + n, m, k, &alpha, + ones, k, bias, k, &beta, output + batch * out_c * out_h * out_w, n); if (stat != CUBLAS_STATUS_SUCCESS) FatalError("CUBLAS initialization failed\n"); modulatedDeformableIm2colCuda(stream, - input + batch * channels * height * width, + input + batch * channels * height * width, offset,// + b * 2 * int((float)chunk_dim / batch), mask,// + b * int((float)chunk_dim / batch), 1, channels, height, width, @@ -283,15 +292,15 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, // height_out, width_out, kernel_h, kernel_w, // pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w, // deformable_group, columns); - + //(k * m) x (m * n) // Y = WC k = channels * kernel_h * kernel_w; beta = 1.0; - - stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, - n, m, k, &alpha, - columns, n, weight, k, + + stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, + n, m, k, &alpha, + columns, n, weight, k, &beta, output + batch * out_c * out_h * out_w, n); if (stat != CUBLAS_STATUS_SUCCESS) diff --git a/src/kernels/pooling.cu b/src/kernels/pooling.cu index cd80a6d..9298ad9 100644 --- a/src/kernels/pooling.cu +++ b/src/kernels/pooling.cu @@ -1,4 +1,13 @@ #include "kernels.h" +#include "pluginsRT/MaxPoolingFixedSizeRT.h" + +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection MaxPoolFixedSizeRTCreator::mFC{}; +std::vector 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) { @@ -39,14 +48,14 @@ __global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c output[out_index] = max; } -void MaxPoolingForward(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride_x, int stride_y, int size, int padding, cudaStream_t stream) +void MaxPoolingForward(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride_x, int stride_y, int size, int padding, cudaStream_t stream) { int tot_size = n*c*h*w; int blocks = (tot_size+255)/256; int threads = 256; - + forward_maxpool_layer_kernel<<>>(tot_size, h, w, c, stride_x, stride_y, size, padding, srcData, dstData); } diff --git a/src/kernels/reorg.cu b/src/kernels/reorg.cu index dc98d0d..666d558 100644 --- a/src/kernels/reorg.cu +++ b/src/kernels/reorg.cu @@ -1,4 +1,13 @@ #include "kernels.h" +#include "pluginsRT/ReorgRT.h" + +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection ReorgRTCreator::mFC{}; +std::vector 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) { @@ -35,14 +44,14 @@ __global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, in /** reorg function function */ -void reorgForward(dnnType* srcData, dnnType* dstData, +void reorgForward(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride, cudaStream_t stream) { int size = n*c*h*w; int blocks = (size+255)/256; int threads = 256; - + reorg_kernel<<>>(size, srcData, w, h, c, n, stride, false, dstData); } diff --git a/src/kernels/resize.cu b/src/kernels/resize.cu index 6059ef5..1662af1 100644 --- a/src/kernels/resize.cu +++ b/src/kernels/resize.cu @@ -1,7 +1,17 @@ -#include "kernels.h" #include -__global__ void resize_kernel( int size,float *x, int i_w, int i_h, int i_c, +#include "kernels.h" +#include "pluginsRT/ResizeLayerRT.h" + +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection ResizeLayerRTCreator::mFC{}; +std::vector ResizeLayerRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(ResizeLayerRTCreator); +}} + +__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 id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; diff --git a/src/kernels/scaladd.cu b/src/kernels/scaladd.cu index 53bcd7a..0fec513 100644 --- a/src/kernels/scaladd.cu +++ b/src/kernels/scaladd.cu @@ -1,6 +1,7 @@ -#include "kernels.h" #include +#include "kernels.h" + __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; @@ -11,6 +12,6 @@ void scalAdd(dnnType* dstData, int size, float alpha, float beta, int inc, cudaS { int blocks = (size+255)/256; int threads = 256; - + scal_add_kernel<<>>(dstData, size, alpha, beta, inc); } \ No newline at end of file diff --git a/src/kernels/shortcut.cu b/src/kernels/shortcut.cu index c2d99bd..cb7489c 100644 --- a/src/kernels/shortcut.cu +++ b/src/kernels/shortcut.cu @@ -1,8 +1,18 @@ #include "kernels.h" #include "assert.h" -__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, +#include "pluginsRT/ShortcutRT.h" + +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection ShortcutRTCreator::mFC{}; +std::vector ShortcutRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(ShortcutRTCreator); +}} + +__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 w2, int h2, int c2, float s1, float s2, dnnType *out) { int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; @@ -21,8 +31,8 @@ __global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stri //out[out_index] += add[add_index]; } -__global__ void shortcut_mul_kernel(int size, int minw, int minh, int minc, int sample, int batch, - int w1, int h1, int c1, dnnType *mul, +__global__ void shortcut_mul_kernel(int size, int minw, int minh, int minc, int sample, int batch, + int w1, int h1, int c1, dnnType *mul, int w2, int h2, int c2, float s1, float s2, dnnType *out) { int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; @@ -40,7 +50,7 @@ __global__ void shortcut_mul_kernel(int size, int minw, int minh, int minc, int } void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int w1, int s1, - int n2, int c2, int h2, int w2, int s2, + int n2, int c2, int h2, int w2, int s2, bool mul, cudaStream_t stream) { assert(n1 == n2); @@ -60,10 +70,10 @@ void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int size = batch * minw * minh * minc; int blocks = (size+255)/256; int threads = 256; - - shortcut_kernel<<>>(size, minw, minh, minc, stride, sample, batch, - w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData); - } + + shortcut_kernel<<>>(size, minw, minh, minc, stride, sample, batch, + w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData); + } else{ int minw = w1; int minh = h1; @@ -74,7 +84,7 @@ void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int blocks = (size+255)/256; int threads = 256; - shortcut_mul_kernel<<>>(size, minw, minh, minc, sample, batch, - w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData); + shortcut_mul_kernel<<>>(size, minw, minh, minc, sample, batch, + w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData); } } diff --git a/src/kernels/static_init.cu b/src/kernels/static_init.cu new file mode 100644 index 0000000..37b8af2 --- /dev/null +++ b/src/kernels/static_init.cu @@ -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 RegionRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(RegionRTCreator); + +nvinfer1::PluginFieldCollection RouteRTCreator::mFC{}; +std::vector RouteRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(RouteRTCreator); + +nvinfer1::PluginFieldCollection ReshapeRTCreator::mFC{}; +std::vector ReshapeRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(ReshapeRTCreator); + +nvinfer1::PluginFieldCollection FlattenConcatRTCreator::mFC{}; +std::vector FlattenConcatRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(FlattenConcatRTCreator); + +nvinfer1::PluginFieldCollection YoloRTCreator::mFC{}; +std::vector YoloRTCreator::mPluginAttributes; + +REGISTER_TENSORRT_PLUGIN(YoloRTCreator); +}} \ No newline at end of file diff --git a/src/kernels/upsample.cu b/src/kernels/upsample.cu index 3b21fa7..840e5b4 100644 --- a/src/kernels/upsample.cu +++ b/src/kernels/upsample.cu @@ -1,4 +1,13 @@ #include "kernels.h" +#include "pluginsRT/UpsampleRT.h" + +// Static class fields initialization +namespace tk { namespace dnn { +nvinfer1::PluginFieldCollection UpsampleRTCreator::mFC{}; +std::vector 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) { @@ -24,8 +33,8 @@ __global__ void upsample_kernel(size_t N, dnnType *x, int w, int h, int c, int b else atomicAdd(x+in_index, scale * out[out_index]); } -void upsampleForward(dnnType* srcData, dnnType* dstData, - int n, int c, int h, int w, int s, int forward, float scale, +void upsampleForward(dnnType* srcData, dnnType* dstData, + int n, int c, int h, int w, int s, int forward, float scale, cudaStream_t stream) { int size = w*h*c*n*s*s; diff --git a/src/yoloContainer.cpp b/src/yoloContainer.cpp new file mode 100644 index 0000000..d3e2366 --- /dev/null +++ b/src/yoloContainer.cpp @@ -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(serialData),*bufCheck = buf; + int classes_temp = tk::dnn::readBUF(buf); + int num_temp = tk::dnn::readBUF(buf); + int n_masks_temp = tk::dnn::readBUF(buf); + float scale_xy_temp = tk::dnn::readBUF(buf); + float nms_thresh_temp = tk::dnn::readBUF(buf); + int nms_kind_temp = tk::dnn::readBUF(buf); + int new_coords_temp = tk::dnn::readBUF(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(buf); + r->h = tk::dnn::readBUF(buf); + r->w = tk::dnn::readBUF(buf); + for(int i=0; in_masks; i++) + r->mask[i] = tk::dnn::readBUF(buf); + for(int i=0; in_masks*2*r->num; i++) + r->bias[i] = tk::dnn::readBUF(buf); + + // save classes names + r->classesNames.resize(r->classes); + for(int i=0; iclasses; i++) { + char tmp[YOLORT_CLASSNAME_W]; + for(int j=0; j(buf); + r->classesNames[i] = std::string(tmp); + } + assert(buf == bufCheck + serialLength); + + yoloContainer.yolos[yoloContainer.n_yolos++] = r; + return r; +} + +}} \ No newline at end of file diff --git a/tests/mnist/test_mnistRT.cpp b/tests/mnist/test_mnistRT.cpp index 1b4e9c3..f21ef06 100644 --- a/tests/mnist/test_mnistRT.cpp +++ b/tests/mnist/test_mnistRT.cpp @@ -13,9 +13,9 @@ const char *output_bin = "mnist/output.bin"; using namespace nvinfer1; // 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 if (severity != Severity::kINFO) @@ -39,20 +39,20 @@ int main() { tk::dnn::Activation l5(&net, CUDNN_ACTIVATION_RELU); tk::dnn::Dense l6(&net, 10, d3_bin); tk::dnn::Softmax l7(&net); - + // Load input dnnType *data; dnnType *input_h; readBinaryFile(input_bin, dim.tot(), &input_h, &data); dim.print(); //print initial dimension - + // Inference { TKDNN_TSTART data = net.infer(dim, data); TKDNN_TSTOP - dim.print(); + dim.print(); } // Print real test @@ -61,19 +61,19 @@ int main() { dnnType *out_h; readBinaryFile(output_bin, dim.tot(), &out_h, &out); std::cout<<"Diff: "<createNetwork(); + INetworkDefinition* network = builder->createNetworkV2(0u); DataType dt = DataType::kFLOAT; // 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); - tk::dnn::Conv2d *c0 = &l0; + tk::dnn::Conv2d *c0 = &l0; Weights w { dt, c0->data_h, c0->inputs*c0->outputs*c0->kernelH*c0->kernelW}; Weights b { dt, c0->bias_h, c0->outputs}; // Add a convolution layer with 20 outputs and a 5x5 filter. @@ -86,7 +86,7 @@ int main() { assert(pool1 != nullptr); pool1->setStride(DimsHW{2, 2}); - tk::dnn::Conv2d *c1 = &l2; + tk::dnn::Conv2d *c1 = &l2; Weights w1 { dt, c1->data_h, c1->inputs*c1->outputs*c1->kernelH*c1->kernelW}; Weights b1 { dt, c1->bias_h, c1->outputs}; // Add a second convolution layer with 50 outputs and a 5x5 filter. @@ -99,7 +99,7 @@ int main() { assert(pool2 != nullptr); pool2->setStride(DimsHW{2, 2}); - tk::dnn::Dense *d2 = &l4; + tk::dnn::Dense *d2 = &l4; Weights w2 { dt, d2->data_h, d2->inputs*d2->outputs}; Weights b2 { dt, d2->bias_h, d2->outputs}; // Add a fully connected layer with 500 outputs. @@ -110,7 +110,7 @@ int main() { auto relu1 = network->addActivation(*ip1->getOutput(0), ActivationType::kRELU); assert(relu1 != nullptr); - tk::dnn::Dense *d3 = &l6; + tk::dnn::Dense *d3 = &l6; Weights w3 { dt, d3->data_h, d3->inputs*d3->outputs}; Weights b3 { dt, d3->bias_h, d3->outputs}; // Add a second fully connected layer with 20 outputs. @@ -125,10 +125,21 @@ int main() { network->markOutput(*prob->getOutput(0)); // Build the engine - builder->setMaxBatchSize(1); +#if NV_TENSORRT_MAJOR >= 6 + auto config = builder->createBuilderConfig(); + config->setMaxWorkspaceSize(1 << 20); +#else 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); +#endif - auto engine = builder->buildCudaEngine(*network); // we don't need the network any more network->destroy(); @@ -142,10 +153,10 @@ int main() { // In order to bind the buffers, we need to know the names of the input and output tensors. // note that indices are guaranteed to be less than IEngine::getNbBindings() - int inputIndex = engine->getBindingIndex("data"); - int outputIndex = engine->getBindingIndex("out"); + int inputIndex = engine->getBindingIndex("data"); + int outputIndex = engine->getBindingIndex("out"); - float output[10]; + float output[10]; // create GPU buffers and a stream checkCuda(cudaMalloc(&buffers[inputIndex], 28*28*sizeof(float))); checkCuda(cudaMalloc(&buffers[outputIndex], 10*sizeof(float)));