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)));