CMakeLists.txt

@@ -18,7 +18,7 @@ if(DEBUG)
 endif()
 
 if(TKDNN_PATH)
-    message("SET TKDNN_PATH:"${TKDNN_PATH})
+    message("SET TKDNN_PATH:${TKDNN_PATH}")
     add_definitions(-DTKDNN_PATH="${TKDNN_PATH}")
 else()
     add_definitions(-DTKDNN_PATH="${CMAKE_CURRENT_SOURCE_DIR}")
@@ -51,11 +51,11 @@ find_package(Eigen3 REQUIRED)
 message("Eigen DIR: " ${EIGEN3_INCLUDE_DIR})
 include_directories(${EIGEN3_INCLUDE_DIR})
 
-find_package(OpenCV REQUIRED)
+find_package(OpenCV 4.5 REQUIRED)
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV")
-# if(OpenCV_CUDA_VERSION)
-#     add_compile_definitions(OPENCV_CUDACONTRIB)
-# endif()
+if(OpenCV_CUDA_VERSION)
+    add_compile_definitions(OPENCV_CUDACONTRIB)
+endif()
 
 # gives problems in cross-compiling, probably malformed cmake config
 find_package(yaml-cpp REQUIRED)
@@ -72,8 +72,8 @@ add_library(tkDNN SHARED ${tkdnn_SRC})
 target_link_libraries(tkDNN ${tkdnn_LIBS})
 
 #static
-#add_library(tkDNN_static STATIC ${tkdnn_SRC})
-#target_link_libraries(tkDNN_static ${tkdnn_LIBS})
+add_library(tkDNN_static STATIC ${tkdnn_SRC})
+target_link_libraries(tkDNN_static ${tkdnn_LIBS})
 
 # SMALL NETS
 add_executable(test_simple tests/simple/test_simple.cpp)

include/tkDNN/DetectionNN.h

@@ -17,7 +17,7 @@
 
 #include "tkdnn.h"
 
-//#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
+#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
 
 #ifdef OPENCV_CUDACONTRIB
 #include <opencv2/cudawarping.hpp>
@@ -37,7 +37,7 @@ class DetectionNN {
 
         cv::Scalar colors[256];
 
-        int nBatches = 1;
+        int nBatches = 2;
 
 #ifdef OPENCV_CUDACONTRIB
         cv::cuda::GpuMat bgr[3];

include/tkDNN/DetectionNN3D.h

@@ -17,7 +17,7 @@
 
 #include "tkdnn.h"
 
-// #define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
+#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
 
 #ifdef OPENCV_CUDACONTRIB
 #include <opencv2/cudawarping.hpp>

include/tkDNN/Int8BatchStream.h

@@ -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<std::string>& mListIn);
@@ -55,7 +55,7 @@ private:
 	int mFileBatchPos{ 0 };
 	int mImageSize{ 0 };
 
-	nvinfer1::DimsNCHW mDims;
+	nvinfer1::Dims4 mDims;
 	std::vector<float> mBatch;
 	std::vector<float> mLabels;
 	std::vector<float> mFileBatch;

include/tkDNN/Int8Calibrator.h

@@ -30,10 +30,10 @@ public:
 	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;

include/tkDNN/NetworkRT.h

@@ -2,28 +2,14 @@
 #define NETWORKRT_H
 
 #include <string.h> // memcpy
+#include <memory>
+
 #include "utils.h"
 #include "Network.h"
 #include "Layer.h"
 #include "NvInfer.h"
-#include <memory>
 
-namespace tk { namespace dnn {
-
-template<typename T> void writeBUF(char*& buffer, const T& val)
-{
-    *reinterpret_cast<T*>(buffer) = val;
-    buffer += sizeof(T);
-}
-
-template<typename T> T readBUF(const char*& buffer)
-{
-    T val = *reinterpret_cast<const T*>(buffer);
-    buffer += sizeof(T);
-    return val;
-}
-
-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 {
 
@@ -74,8 +51,6 @@ public:
     dnnType *output;
     cudaStream_t stream;
 
-    PluginFactory *pluginFactory;
-
     NetworkRT(Network *net, const char *name);
     virtual ~NetworkRT();
 

include/tkDNN/buffer_func.h

@@ -0,0 +1,20 @@
+#ifndef BUFFER_FUNC_H
+#define BUFFER_FUNC_H
+
+namespace tk { namespace dnn {
+
+template<typename T> void writeBUF(char*& buffer, const T& val)
+{
+    *reinterpret_cast<T*>(buffer) = val;
+    buffer += sizeof(T);
+}
+
+template<typename T> T readBUF(const char*& buffer)
+{
+    T val = *reinterpret_cast<const T*>(buffer);
+    buffer += sizeof(T);
+    return val;
+}
+}}
+
+#endif // BUFFER_FUNC_H

include/tkDNN/pluginsRT/ActivationLeakyRT.h

@@ -1,7 +1,19 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef ACTIVATION_LEAKY_RT_H
+#define ACTIVATION_LEAKY_RT_H
 
-class ActivationLeakyRT : public IPlugin {
+#if NV_TENSORRT_MAJOR < 6
+
+#include <cassert>
+#include <vector>
+
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+namespace tk { namespace dnn {
+
+class ActivationLeakyRT final : public IPlugin {
 
 public:
 	ActivationLeakyRT(float s) {
@@ -31,14 +43,14 @@ 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<const dnnType*>(inputs[0]),
 											reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, slope, stream);
@@ -46,16 +58,21 @@ public:
 	}
 
 
-	virtual size_t getSerializationSize() override {
+	size_t getSerializationSize() override {
 		return 1*sizeof(int) + 1*sizeof(float);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) override {
 		char *buf = reinterpret_cast<char*>(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

include/tkDNN/pluginsRT/ActivationLogisticRT.h

@@ -1,60 +1,149 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef ACTIVATION_LOGISTIC_RT_H
+#define ACTIVATION_LOGISTIC_RT_H
 
-class ActivationLogisticRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "ActivationLogistic"
+#define PLUGIN_VERSION "1"
+
+namespace tk { namespace dnn {
+
+class ActivationLogisticRT final : public nvinfer1::IPluginV2 {
 
 public:
-    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<outputDims[0].nbDims; i++)
             size *= outputDims[0].d[i];
     }
 
-    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 {
         activationLOGISTICForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
                                   reinterpret_cast<dnnType*>(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<char*>(buffer);
-        tk::dnn::writeBUF(buf, size);
+        writeBUF(buf, size);
     }
 
+    // Extra IPluginV2 overrides
+    bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+        return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+    }
+
+    nvinfer1::IPluginV2 * clone() const noexcept override {
+        auto a = new ActivationLogisticRT(*this);
+        return a;
+    }
+
+    const char* getPluginType() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    void destroy() noexcept override {}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+        mNamespace = pluginNamespace;
+    }
+
+    const char* getPluginNamespace() const noexcept override {
+        return mNamespace.c_str();
+    }
+
+    std::string mNamespace;
     int size;
 };
+
+class ActivationLogisticRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationLogisticRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+        const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+        ActivationLogisticRT *a = new ActivationLogisticRT();
+        a->size = readBUF<int>(buf);
+        assert(buf == bufCheck + serialLength);
+        return a;
+    }
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+        mNamespace = pluginNamespace;
+    }
+
+    const char* getPluginNamespace() const noexcept override {
+        return mNamespace.c_str();
+    }
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // ACTIVATION_LOGISTIC_RT_H

include/tkDNN/pluginsRT/ActivationMishRT.h

@@ -1,61 +1,150 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef ACTIVATION_MISH_RT_H
+#define ACTIVATION_MISH_RT_H
 
-class ActivationMishRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "ActivationMish"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class ActivationMishRT final : public nvinfer1::IPluginV2 {
 
 public:
-	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<outputDims[0].nbDims; i++)
 			size *= outputDims[0].d[i];
 	}
 
-	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 {
 		activationMishForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
 											reinterpret_cast<dnnType*>(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<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ActivationMishRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int size;
 };
+
+class ActivationMishRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationMishRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+        const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+        ActivationMishRT *a = new ActivationMishRT();
+        a->size = readBUF<int>(buf);
+        assert(buf == bufCheck + serialLength);
+        return a;
+    }
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+        mNamespace = pluginNamespace;
+    }
+
+    const char* getPluginNamespace() const noexcept override {
+        return mNamespace.c_str();
+    }
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // ACTIVATION_MISH_RT_H

include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h

@@ -1,63 +1,156 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef ACTIVATION_RELU_CEILING_RT_H
+#define ACTIVATION_RELU_CEILING_RT_H
 
-class ActivationReLUCeiling : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "ActivationCReLU"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class ActivationReLUCeiling final : public nvinfer1::IPluginV2 {
 
 public:
 	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<outputDims[0].nbDims; i++)
 			size *= outputDims[0].d[i];
 	}
 
-	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 {
 		activationReLUCeilingForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
 											reinterpret_cast<dnnType*>(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<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ActivationReLUCeiling(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int size;
 	float ceiling;
 };
+
+class ActivationReLUCeilingCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationReLUCeilingCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			float activationReluTemp = readBUF<float>(buf);
+			ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp);
+			a->size = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return a;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+        mNamespace = pluginNamespace;
+    }
+
+    const char* getPluginNamespace() const noexcept override {
+        return mNamespace.c_str();
+    }
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // ACTIVATION_RELU_CEILING_RT_H

include/tkDNN/pluginsRT/ActivationSigmoidRT.h

@@ -1,61 +1,149 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef ACTIVATION_SIGMOID_RT_H
+#define ACTIVATION_SIGMOID_RT_H
 
-class ActivationSigmoidRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "ActivationSigmoidRT"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class ActivationSigmoidRT final : public nvinfer1::IPluginV2 {
 
 public:
-	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<outputDims[0].nbDims; i++)
 			size *= outputDims[0].d[i];
 	}
 
-	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 {
 		activationSIGMOIDForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
 											reinterpret_cast<dnnType*>(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<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ActivationSigmoidRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int size;
 };
+
+class ActivationSigmoidRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationSigmoidRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			ActivationSigmoidRT* a = new ActivationSigmoidRT();
+			a->size = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return a;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // ACTIVATION_SIGMOID_RT_H

include/tkDNN/pluginsRT/DeformableConvRT.h

@@ -1,16 +1,27 @@
-#include<cassert>
+#ifndef DEFORMABLE_CONV_RT_H
+#define DEFORMABLE_CONV_RT_H
+
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
 #include "../kernels.h"
+#include "../buffer_func.h"
+#include "../Layer.h"
 
+#define PLUGIN_NAME "Deformable"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
 
-class DeformableConvRT : public IPlugin {
-
-
+class DeformableConvRT final : public nvinfer1::IPluginV2 {
 
 public:
 	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) {
+					DeformConv2d *deformable = nullptr) {
 		this->chunk_dim = chunk_dim;
 		this->kh = kh;
 		this->kw = kw;
@@ -61,27 +72,34 @@ 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<const dnnType*>(inputs[0]);
 		dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
 
@@ -109,63 +127,92 @@ public:
 		return 0;
 	}
 
-
-	virtual size_t getSerializationSize() override {
+	size_t getSerializationSize() const noexcept override {
 		return 16 * sizeof(int) + chunk_dim * 3 * sizeof(dnnType) + (i_c * o_c * kh * kw * 1 ) * sizeof(dnnType) +
 			   o_c * sizeof(dnnType) + height_ones * width_ones * sizeof(dnnType) + dim_ones * sizeof(dnnType);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) const noexcept override {
 		char *buf = reinterpret_cast<char*>(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<chunk_dim*2; i++)
-    		tk::dnn::writeBUF(buf, aus[i]);
+    		writeBUF(buf, aus[i]);
 		free(aus);
 		aus = new dnnType[chunk_dim];
 		checkCuda( cudaMemcpy(aus, mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToHost) );
         for(int i=0; i<chunk_dim; i++)
-    		tk::dnn::writeBUF(buf, aus[i]);
+    		writeBUF(buf, aus[i]);
 		free(aus);
 		aus = new dnnType[(i_c * o_c * kh * kw * 1 )];
 		checkCuda( cudaMemcpy(aus, data_d, sizeof(dnnType)*(i_c * o_c * kh * kw * 1 ), cudaMemcpyDeviceToHost) );
         for(int i=0; i<(i_c * o_c * kh * kw * 1 ); i++)
-    		tk::dnn::writeBUF(buf, aus[i]);
+    		writeBUF(buf, aus[i]);
 		free(aus);
 		aus = new dnnType[o_c];
 		checkCuda( cudaMemcpy(aus, bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToHost) );
         for(int i=0; i < o_c; i++)
-    		tk::dnn::writeBUF(buf, aus[i]);
+    		writeBUF(buf, aus[i]);
 		free(aus);
 		aus = new dnnType[height_ones * width_ones];
 		checkCuda( cudaMemcpy(aus, ones_d1, sizeof(dnnType)*height_ones * width_ones, cudaMemcpyDeviceToHost) );
         for(int i=0; i<height_ones * width_ones; i++)
-    		tk::dnn::writeBUF(buf, aus[i]);
+    		writeBUF(buf, aus[i]);
 		free(aus);
 		aus = new dnnType[dim_ones];
 		checkCuda( cudaMemcpy(aus, ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToHost) );
         for(int i=0; i<dim_ones; i++)
-    		tk::dnn::writeBUF(buf, aus[i]);
+    		writeBUF(buf, aus[i]);
 		free(aus);
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new DeformableConvRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	cublasStatus_t stat;
 	cublasHandle_t handle;
 	int i_n, i_c, i_h, i_w;
@@ -181,7 +228,7 @@ public:
 	int dim_ones;
 
 	dnnType *data_d;
-    dnnType *bias2_d;
+  dnnType *bias2_d;
 	dnnType *ones_d1;
 	dnnType * offset;
 	dnnType * mask;
@@ -191,6 +238,99 @@ public:
 	// dnnType *mask_n;
 	// dnnType *output_n;
 
-	
-	tk::dnn::DeformConv2d *defRT;
+	DeformConv2d *defRT;
 };
+
+class DeformableConvRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    DeformableConvRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int chuck_dimTemp = readBUF<int>(buf);
+			int khTemp = readBUF<int>(buf);
+			int kwTemp = readBUF<int>(buf);
+			int shTemp = readBUF<int>(buf);
+			int swTemp = readBUF<int>(buf);
+			int phTemp = readBUF<int>(buf);
+			int pwTemp = readBUF<int>(buf);
+			int deformableGroupTemp = readBUF<int>(buf);
+			int i_nTemp = readBUF<int>(buf);
+			int i_cTemp = readBUF<int>(buf);
+			int i_hTemp = readBUF<int>(buf);
+			int i_wTemp = readBUF<int>(buf);
+			int o_nTemp = readBUF<int>(buf);
+			int o_cTemp = readBUF<int>(buf);
+			int o_hTemp = readBUF<int>(buf);
+			int o_wTemp = readBUF<int>(buf);
+
+			DeformableConvRT* r = new DeformableConvRT(chuck_dimTemp, khTemp, kwTemp, shTemp, swTemp, phTemp, pwTemp, deformableGroupTemp, i_nTemp, i_cTemp, i_hTemp, i_wTemp, o_nTemp, o_cTemp, o_hTemp, o_wTemp, nullptr);
+			dnnType *aus = new dnnType[r->chunk_dim*2];
+			for(int i=0; i<r->chunk_dim*2; i++)
+			aus[i] = readBUF<dnnType>(buf);
+			checkCuda( cudaMemcpy(r->offset, aus, sizeof(dnnType)*2*r->chunk_dim, cudaMemcpyHostToDevice) );
+					free(aus);
+			aus = new dnnType[r->chunk_dim];
+			for(int i=0; i<r->chunk_dim; i++)
+							aus[i] = readBUF<dnnType>(buf);
+			checkCuda( cudaMemcpy(r->mask, aus, sizeof(dnnType)*r->chunk_dim, cudaMemcpyHostToDevice) );
+					free(aus);
+			aus = new dnnType[(r->i_c * r->o_c * r->kh * r->kw * 1 )];
+			for(int i=0; i<(r->i_c * r->o_c * r->kh * r->kw * 1 ); i++)
+					aus[i] = readBUF<dnnType>(buf);
+			checkCuda( cudaMemcpy(r->data_d, aus, sizeof(dnnType)*(r->i_c * r->o_c * r->kh * r->kw * 1 ), cudaMemcpyHostToDevice) );
+					free(aus);
+			aus = new dnnType[r->o_c];
+			for(int i=0; i < r->o_c; i++)
+					aus[i] = readBUF<dnnType>(buf);
+			checkCuda( cudaMemcpy(r->bias2_d, aus, sizeof(dnnType)*r->o_c, cudaMemcpyHostToDevice) );
+					free(aus);
+			aus = new dnnType[r->height_ones * r->width_ones];
+			for(int i=0; i<r->height_ones * r->width_ones; i++)
+					aus[i] = readBUF<dnnType>(buf);
+			checkCuda( cudaMemcpy(r->ones_d1, aus, sizeof(dnnType)*r->height_ones * r->width_ones, cudaMemcpyHostToDevice) );
+					free(aus);
+			aus = new dnnType[r->dim_ones];
+			for(int i=0; i<r->dim_ones; i++)
+					aus[i] = readBUF<dnnType>(buf);
+			checkCuda( cudaMemcpy(r->ones_d2, aus, sizeof(dnnType)*r->dim_ones, cudaMemcpyHostToDevice) );
+					free(aus);
+					assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // DEFORMABLE_CONV_RT_H

include/tkDNN/pluginsRT/FlattenConcatRT.h

@@ -1,6 +1,20 @@
-#include<cassert>
+#ifndef FLATTEN_CONCAT_RT_H
+#define FLATTEN_CONCAT_RT_H
 
-class FlattenConcatRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "Flatten"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class FlattenConcatRT final : public nvinfer1::IPluginV2 {
 
 public:
 	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,19 +50,19 @@ 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<const dnnType*>(inputs[0]);
 		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
 		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
@@ -59,23 +77,105 @@ public:
 		return 0;
 	}
 
-
-	virtual size_t getSerializationSize() override {
+	size_t getSerializationSize() const noexcept override {
 		return 5*sizeof(int);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) const noexcept override {
 		char *buf = reinterpret_cast<char*>(buffer),*a = buf;
-		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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new FlattenConcatRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
 	int rows, cols;
 	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<const char*>(serialData),*bufCheck = buf;
+			float activationReluTemp = readBUF<float>(buf);
+			FlattenConcatRT *r = new FlattenConcatRT();
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			r->rows = readBUF<int>(buf);
+			r->cols = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // FLATTEN_CONCAT_RT_H

include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h

@@ -1,7 +1,20 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef MAX_POOLING_FIXED_SIZE_RT_H
+#define MAX_POOLING_FIXED_SIZE_RT_H
 
-class MaxPoolFixedSizeRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "Pooling"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class MaxPoolFixedSizeRT final : public nvinfer1::IPluginV2 {
 
 public:
 	MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) {
@@ -15,33 +28,37 @@ public:
 		this->padding = padding;
 	}
 
-	~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<<this->n<<"  "<<this->c<<"  "<<this->h<<"  "<<this->w<<"  "<<this->stride_H<<"  "<<this->stride_W<<"  "<<this->winSize<<"  "<<this->padding<<std::endl;
 		dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
 		dnnType *dstData = reinterpret_cast<dnnType*>(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<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new MaxPoolFixedSizeRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int n, c, h, w;
 	int 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<const char*>(serialData),*bufCheck = buf;
+			int cTemp = readBUF<int>(buf);
+			int hTemp = readBUF<int>(buf);
+			int wTemp = readBUF<int>(buf);
+			int nTemp = readBUF<int>(buf);
+			int strideHTemp = readBUF<int>(buf);
+			int strideWTemp = readBUF<int>(buf);
+			int winSizeTemp = readBUF<int>(buf);
+			int paddingTemp = readBUF<int>(buf);
+
+			MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // MAX_POOLING_FIXED_SIZE_RT_H

include/tkDNN/pluginsRT/RegionRT.h

@@ -1,48 +1,62 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef REGION_RT_H
+#define REGION_RT_H
 
-class RegionRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "Region"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class RegionRT final : public nvinfer1::IPluginV2 {
 
 public:
 	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<const dnnType*>(inputs[0]);
 		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
 
@@ -68,23 +82,53 @@ public:
 	}
 
 
-	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<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new RegionRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
-    int 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<const char*>(serialData),*bufCheck = buf;
+			int classesTemp = readBUF<int>(buf);
+			int coordsTemp = readBUF<int>(buf);
+			int numTemp = readBUF<int>(buf);
+			RegionRT* r = new RegionRT(classesTemp, coordsTemp, numTemp);
+
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // REGION_RT_H

include/tkDNN/pluginsRT/ReorgRT.h

@@ -1,64 +1,159 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef REORG_RT_H
+#define REORG_RT_H
 
-class ReorgRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "Reorg"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class ReorgRT final : public nvinfer1::IPluginV2 {
 
 public:
 	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 {
-
+  int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
 		reorgForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
 					  reinterpret_cast<dnnType*>(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<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ReorgRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w, stride;
 };
+
+class ReorgRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ReorgRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int strideTemp = readBUF<int>(buf);
+			ReorgRT *r = new ReorgRT(strideTemp);
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // REORG_RT_H

include/tkDNN/pluginsRT/ReshapeRT.h

@@ -1,6 +1,21 @@
-#include<cassert>
+#ifndef RESHAPE_RT_H
+#define RESHAPE_RT_H
 
-class ReshapeRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+#include "../Network.h"
+
+#define PLUGIN_NAME "Reshape"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class ReshapeRT final : public nvinfer1::IPluginV2 {
 
 public:
 	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<const dnnType*>(inputs[0]);
 		dnnType *dstData = reinterpret_cast<dnnType*>(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<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ReshapeRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int n, c, h, w;
 };
+
+class ReshapeRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ReshapeRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			dataDim_t new_dim;
+			new_dim.n = readBUF<int>(buf);
+			new_dim.c = readBUF<int>(buf);
+			new_dim.h = readBUF<int>(buf);
+			new_dim.w = readBUF<int>(buf);
+			ReshapeRT *r = new ReshapeRT(new_dim);
+			assert(buf == bufCheck + serialLength);
+
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // RESHAPE_RT_H

include/tkDNN/pluginsRT/ResizeLayerRT.h

@@ -1,7 +1,20 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef RESIZE_LAYER_RT_H
+#define RESIZE_LAYER_RT_H
 
-class ResizeLayerRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "Resize"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class ResizeLayerRT final : public nvinfer1::IPluginV2 {
 
 public:
 	ResizeLayerRT(int c, int h, int w) {
@@ -10,35 +23,40 @@ public:
 		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];
 	}
 
-	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<const dnnType*>(inputs[0]),
 					  reinterpret_cast<dnnType*>(outputs[0]),
@@ -46,23 +64,105 @@ public:
 		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<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ResizeLayerRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int i_c, i_h, i_w, o_c, o_h, o_w;
 };
+
+class ResizeLayerRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ResizeLayerRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int o_cTemp = readBUF<int>(buf);
+			int o_hTemp = readBUF<int>(buf);
+			int o_wTemp = readBUF<int>(buf);
+			ResizeLayerRT* r = new ResizeLayerRT(o_cTemp, o_hTemp, o_wTemp);
+
+			r->i_c = readBUF<int>(buf);
+			r->i_h = readBUF<int>(buf);
+			r->i_w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // RESIZE_LAYER_RT_H

include/tkDNN/pluginsRT/RouteRT.h

@@ -1,7 +1,20 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef ROUTE_RT_H
+#define ROUTE_RT_H
 
-class RouteRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "Route"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class RouteRT final : public nvinfer1::IPluginV2 {
 
 	/**
 		THIS IS NOT USED ANYMORE
@@ -13,21 +26,25 @@ 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<nbInputDims; i++) out_c += inputs[i].d[0];
-		return DimsCHW{out_c/groups, inputs[0].d[1], inputs[0].d[2]};
+		return nvinfer1::Dims3{out_c/groups, inputs[0].d[1], inputs[0].d[2]};
 	}
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+	void configureWithFormat(nvinfer1::Dims const * inputDims,
+					int32_t nbInputs,
+					nvinfer1::Dims const * outputDims,
+					int32_t nbOutputs,
+					nvinfer1::DataType type,
+					nvinfer1::PluginFormat format,
+					int32_t maxBatchSize) noexcept override {
 		in = nbInputs;
 		c = 0;
 		for(int i=0; i<nbInputs; i++) {
@@ -39,20 +56,18 @@ public:
 		c /= groups;
 	}
 
-	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 *dstData = reinterpret_cast<dnnType*>(outputs[0]);
 
 		for(int b=0; b<batchSize; b++) {
@@ -70,27 +85,112 @@ public:
 	}
 
 
-	virtual size_t getSerializationSize() override {
+	size_t getSerializationSize() const noexcept override {
 		return (6+MAX_INPUTS)*sizeof(int);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) const noexcept override {
 		char *buf = reinterpret_cast<char*>(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<MAX_INPUTS; i++)
-			tk::dnn::writeBUF(buf, c_in[i]);
+			writeBUF(buf, c_in[i]);
 
-		tk::dnn::writeBUF(buf, c);
-		tk::dnn::writeBUF(buf, h);
-		tk::dnn::writeBUF(buf, w);
+		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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new RouteRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	static const int MAX_INPUTS = 4;
 	int in;
 	int c_in[MAX_INPUTS];
 	int c, h, w;
 	int groups, group_id;
 };
+
+class RouteRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    RouteRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int groupsTemp = readBUF<int>(buf);
+			int group_idTemp = readBUF<int>(buf);
+			RouteRT* r = new RouteRT(groupsTemp, group_idTemp);
+			r->in = readBUF<int>(buf);
+			for(int i=0; i<RouteRT::MAX_INPUTS; i++)
+					r->c_in[i] = readBUF<int>(buf);
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // ROUTE_RT_H

include/tkDNN/pluginsRT/ShortcutRT.h

@@ -1,48 +1,64 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef SHORTCUT_RT_H
+#define SHORTCUT_RT_H
 
-class ShortcutRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+#include "../Network.h"
+
+#define PLUGIN_NAME "Shortcut"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class ShortcutRT final : public nvinfer1::IPluginV2 {
 
 public:
-	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<const dnnType*>(inputs[0]);
 		dnnType *srcDataBack = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
 		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
@@ -54,24 +70,110 @@ public:
 	}
 
 
-	virtual size_t getSerializationSize() override {
+	size_t getSerializationSize() const noexcept override {
 		return 6*sizeof(int) + sizeof(bool);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) const noexcept override {
 		char *buf = reinterpret_cast<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ShortcutRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
 	int 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<const char*>(serialData),*bufCheck = buf;
+			dataDim_t bdim;
+			bdim.c = readBUF<int>(buf);
+			bdim.h = readBUF<int>(buf);
+			bdim.w = readBUF<int>(buf);
+			bdim.l = 1;
+
+			ShortcutRT *r = new ShortcutRT(bdim, readBUF<bool>(buf));
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // SHORTCUT_RT_H

include/tkDNN/pluginsRT/UpsampleRT.h

@@ -1,45 +1,60 @@
-#include<cassert>
-#include "../kernels.h"
+#ifndef UPSAMPLE_RT_H
+#define UPSAMPLE_RT_H
 
-class UpsampleRT : public IPlugin {
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../kernels.h"
+#include "../buffer_func.h"
+
+#define PLUGIN_NAME "Upsample"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
+
+class UpsampleRT final : public nvinfer1::IPluginV2 {
 
 public:
 	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<const dnnType*>(inputs[0]);
 		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
 
@@ -48,19 +63,98 @@ public:
 		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<char*>(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 (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new UpsampleRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int c, h, w, stride;
 };
+
+class UpsampleRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    UpsampleRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int strideTemp = readBUF<int>(buf);
+			UpsampleRT* r = new UpsampleRT(strideTemp);
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // UPSAMPLE_RT_H

include/tkDNN/pluginsRT/YoloRT.h

@@ -1,15 +1,26 @@
-#include<cassert>
+#ifndef YOLO_RT_H
+#define YOLO_RT_H
+
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#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<const dnnType*>(inputs[0]);
 		dnnType *dstData = reinterpret_cast<dnnType*>(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<char*>(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<YOLORT_CLASSNAME_W; j++) {
-				tk::dnn::writeBUF(buf, tmp[j]);
+				writeBUF(buf, tmp[j]);
 			}
 		}
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new YoloRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
     int classes, num, n_masks;
 	float scaleXY;
@@ -139,5 +181,46 @@ public:
 		int loc = location % (w*h);
 		return batch*c*h*w + n*w*h*(4+classes+1) + entry*w*h + loc;
 	}
-
 };
+
+class YoloRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    YoloRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override;
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // YOLO_RT_H

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

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

src/Int8Calibrator.cpp

@@ -8,13 +8,13 @@ Int8EntropyCalibrator::Int8EntropyCalibrator(BatchStream& stream, int firstBatch
     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<const char*>(cache), length);

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
@@ -39,7 +39,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
 #if NV_TENSORRT_MAJOR >= 5
     std::cout<<"DLAs: "<<builderRT->getNbDLACores()<<"\n";
 #endif
-    networkRT = builderRT->createNetwork();
+    networkRT = builderRT->createNetworkV2(0u);
 #if NV_TENSORRT_MAJOR >= 6
         configRT = builderRT->createBuilderConfig();
 #endif
@@ -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);
+#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,6 +80,15 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
             builderRT->setDLACore(0);
         }
 #endif
+#if NV_TENSORRT_MAJOR >= 8
+        if(net->dla && builderRT->getNbDLACores() > 0) {
+            dtRT = DataType::kHALF;
+            configRT->setFlag(BuilderFlag::kFP16);
+            configRT->setFlag(BuilderFlag::kGPU_FALLBACK);
+            configRT->setDefaultDeviceType(DeviceType::kDLA);
+            configRT->setDLACore(0);
+        }
+#endif
 #if NV_TENSORRT_MAJOR >= 6
         if(net->int8 && builderRT->platformHasFastInt8()){
             // dtRT = DataType::kINT8;
@@ -104,7 +116,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
 
         // add input layer
         ITensor *input = networkRT->addInput("data", DataType::kFLOAT,
-                        DimsCHW{ dim.c, dim.h, dim.w});
+                        Dims3{ dim.c, dim.h, dim.w});
         checkNULL(input);
 
         //add other layers
@@ -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;
     }
@@ -392,8 +404,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
 
 #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
@@ -407,28 +419,35 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
         IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU);
         checkNULL(lRT);
         return lRT;
-    } else if(l->act_mode == CUDNN_ACTIVATION_SIGMOID) {
+    } else if(l->act_mode == CUDNN_ACTIVATION_SIGMOID || l->act_mode == ACTIVATION_LOGISTIC) {
         IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kSIGMOID);
         checkNULL(lRT);
         return lRT;
     }
     else if(l->act_mode == CUDNN_ACTIVATION_CLIPPED_RELU) {
-        IPlugin *plugin = new ActivationReLUCeiling(l->ceiling);
-        IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+        IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kCLIP);
+        lRT->setAlpha(0);
+        lRT->setBeta(l->ceiling);
+
         checkNULL(lRT);
         return lRT;
     }
     else if(l->act_mode == ACTIVATION_MISH) {
-        IPlugin *plugin = new ActivationMishRT();
-        IPluginLayer *lRT = networkRT->addPlugin(&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);
-        checkNULL(lRT);
-        return lRT;
+        // Uncomment this to see if you have better performance
+        // For older TensorRT or for FP32 this might be better
+        //auto *plugin = new ActivationMishRT();
+        //auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
+
+        // Assemble MISH using 3 layers that are going to be merged by TensorRT
+        IActivationLayer *lRT1 = networkRT->addActivation(*input, ActivationType::kSOFTPLUS);
+        lRT1->setAlpha(1);
+        lRT1->setBeta(1);
+
+        IActivationLayer *lRT2 = networkRT->addActivation(*lRT1->getOutput(0), ActivationType::kTANH);
+        IElementWiseLayer *lRT3 = networkRT->addElementWise(*input, *lRT2->getOutput(0), ElementWiseOperation::kPROD);
+
+        checkNULL(lRT3);
+        return lRT3;
     }
     else {
         FatalError("this Activation mode is not yet implemented");
@@ -460,8 +479,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 +491,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 +500,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 +513,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 +521,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 +531,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;
 }
@@ -534,11 +553,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);
+        auto *lRT = networkRT->addPluginV2(inputs, 2, *plugin);
         checkNULL(lRT);
         return lRT;
     }
@@ -548,8 +567,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 +577,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 +593,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, 
+    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);
@@ -646,53 +665,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<const char*>(serialData),*bufCheck = buf;
 
+//     std::string name(layerName);
+//     //std::cout<<name<<std::endl;
 
-IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
-    const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
-
-    std::string name(layerName);
-    //std::cout<<name<<std::endl;
-
-    if(name.find("ActivationLeaky") == 0) {
-        ActivationLeakyRT *a = new ActivationLeakyRT(readBUF<float>(buf));
-        a->size = readBUF<int>(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<float>(buf));
+//         a->size = readBUF<int>(buf);
+//         assert(buf == bufCheck + serialLength);
+//         return a;
+//     }
+// #endif
+    /*if(name.find("ActivationMish") == 0) {
         ActivationMishRT *a = new ActivationMishRT();
         a->size = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return a;
-    }
-    if(name.find("ActivationLogistic") == 0) {
+    }*/
+    /*if(name.find("ActivationLogistic") == 0) {
         ActivationLogisticRT *a = new ActivationLogisticRT();
         a->size = readBUF<int>(buf);
         return a;
-    }
-    if(name.find("ActivationLogistic") == 0) {
+    }*/
+    /*if(name.find("ActivationLogistic") == 0) {
         ActivationLogisticRT *a = new ActivationLogisticRT();
         a->size = readBUF<int>(buf);
         return a;
-    }
-    if(name.find("ActivationCReLU") == 0) {
+    }*/
+    /*if(name.find("ActivationCReLU") == 0) {
         float activationReluTemp = readBUF<float>(buf);
         ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp);
         a->size = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return a;
-    }
+    }*/
 
-    if(name.find("Region") == 0) {
+    /*if(name.find("Region") == 0) {
         int classesTemp = readBUF<int>(buf);
         int coordsTemp = readBUF<int>(buf);
         int numTemp = readBUF<int>(buf);
@@ -703,9 +721,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Reorg") == 0) {
+    /*if(name.find("Reorg") == 0) {
         int strideTemp = readBUF<int>(buf);
         ReorgRT *r = new ReorgRT(strideTemp);
         r->c = readBUF<int>(buf);
@@ -713,9 +731,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(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<int>(buf);
         bdim.h = readBUF<int>(buf);
@@ -728,9 +746,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(buf);
         return r;
         assert(buf == bufCheck + serialLength);
-    } 
+    }*/
 
-    if(name.find("Pooling") == 0) {
+    /*if(name.find("Pooling") == 0) {
         int cTemp = readBUF<int>(buf);
         int hTemp = readBUF<int>(buf);
         int wTemp = readBUF<int>(buf);
@@ -743,9 +761,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp);
         assert(buf == bufCheck + serialLength);
         return r;
-    }
+    }*/
 
-    if(name.find("Resize") == 0) {
+    /*if(name.find("Resize") == 0) {
         int o_cTemp = readBUF<int>(buf);
         int o_hTemp = readBUF<int>(buf);
         int o_wTemp = readBUF<int>(buf);
@@ -756,9 +774,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->i_w = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Flatten") == 0) {
+    /*if(name.find("Flatten") == 0) {
         FlattenConcatRT *r = new FlattenConcatRT();
         r->c = readBUF<int>(buf);
         r->h = readBUF<int>(buf);
@@ -767,9 +785,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->cols = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Reshape") == 0) {
+    /*if(name.find("Reshape") == 0) {
 
         dataDim_t new_dim;
         new_dim.n = readBUF<int>(buf);
@@ -780,9 +798,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         assert(buf == bufCheck + serialLength);
 
         return r;
-    } 
+    }*/
 
-    if(name.find("Yolo") == 0) {
+    /*if(name.find("Yolo") == 0) {
 
         int classes_temp = readBUF<int>(buf);
         int num_temp = readBUF<int>(buf);
@@ -816,8 +834,8 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
 
         yolos[n_yolos++] = r;
         return r;
-    } 
-    if(name.find("Upsample") == 0) {
+    }*/
+    /*if(name.find("Upsample") == 0) {
         int strideTemp = readBUF<int>(buf);
         UpsampleRT* r = new UpsampleRT(strideTemp);
         r->c = readBUF<int>(buf);
@@ -825,9 +843,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Route") == 0) {
+    /*if(name.find("Route") == 0) {
         int groupsTemp = readBUF<int>(buf);
         int group_idTemp = readBUF<int>(buf);
         RouteRT* r = new RouteRT(groupsTemp, group_idTemp);
@@ -839,9 +857,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Deformable") == 0) {
+    /*if(name.find("Deformable") == 0) {
         int chuck_dimTemp = readBUF<int>(buf);
         int khTemp = readBUF<int>(buf);
         int kwTemp = readBUF<int>(buf);
@@ -892,10 +910,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;
+// }
 
 }}

src/Yolo3Detection.cpp

@@ -1,5 +1,6 @@
+#include "yoloContainer.h"
 #include "Yolo3Detection.h"
-
+#include "pluginsRT/YoloRT.h"
 
 namespace tk { namespace dnn {
 
@@ -14,12 +15,12 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
     tk::dnn::dataDim_t idim = netRT->input_dim;
     idim.n = nBatches;
 
-    if(netRT->pluginFactory->n_yolos < 2 ) {
+    if(yoloContainer.n_yolos < 2 ) {
         FatalError("this is not yolo3");
     }
 
-    for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
-        YoloRT *yRT = netRT->pluginFactory->yolos[i];
+    for(int i=0; i<yoloContainer.n_yolos; i++) {
+        YoloRT *yRT = yoloContainer.yolos[i];
         classes = yRT->classes;
         num = yRT->num;
         nMasks = yRT->n_masks;
@@ -95,8 +96,8 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
 
     //get yolo outputs
     std::vector<float *> rt_out;
-    //dnnType *rt_out[netRT->pluginFactory->n_yolos];
-    for(int i=0; i<netRT->pluginFactory->n_yolos; i++)
+    //dnnType *rt_out[yoloContainer.n_yolos];
+    for(int i=0; i<yoloContainer.n_yolos; i++)
         rt_out.push_back((dnnType*)netRT->buffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi);
 
     float x_ratio =  float(originalSize[bi].width) / float(netRT->input_dim.w);
@@ -104,7 +105,7 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
 
     // compute dets
     nDets = 0;
-    for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
+    for(int i=0; i<yoloContainer.n_yolos; i++) {
         yolo[i]->dstData = rt_out[i];
         yolo[i]->computeDetections(dets, nDets, netRT->input_dim.w, netRT->input_dim.h, confThreshold, yolo[i]->new_coords);
     }

src/kernels/activation_logistic.cu

@@ -1,4 +1,15 @@
+#include <vector>
+
 #include "kernels.h"
+#include "pluginsRT/ActivationLogisticRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ActivationLogisticRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ActivationLogisticRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ActivationLogisticRTCreator);
+}}
 
 __global__
 void activation_logistic(dnnType *input, dnnType *output, int size) {
@@ -21,5 +32,3 @@ void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, cud
 
     activation_logistic<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
 }
-
-

src/kernels/activation_mish.cu

@@ -1,8 +1,17 @@
-#include "kernels.h"
 #include <math.h>
+#include "kernels.h"
+#include "pluginsRT/ActivationMishRT.h"
 
 #define MISH_THRESHOLD 20
 
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ActivationMishRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ActivationMishRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ActivationMishRTCreator);
+}}
+
 __device__
 float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);}
 
@@ -13,8 +22,6 @@ float softplus_kernel(float x, float threshold = 20) {
     return logf(expf(x) + 1);
 }
 
-
-
 __device__
 float mish_yashas(float x) {
     float e = __expf(x);
@@ -28,6 +35,16 @@ float mish_yashas(float x) {
     return x - 2 * __fdividef(x, n + 2);
 }
 
+__device__ float mish_yashas2(float x)
+{
+    float e = __expf(x);
+    float n = e * e + 2 * e;
+    if (x <= -0.6f)
+        return x * __fdividef(n, n + 2);
+
+    return x - 2 * __fdividef(x, n + 2);
+}
+
 // https://github.com/digantamisra98/Mish
 // https://github.com/AlexeyAB/darknet/blob/master/src/activation_kernels.cu
 __global__
@@ -35,7 +52,7 @@ void activation_mish(dnnType *input, dnnType *output, int size) {
     int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
     if (i < size)
         // output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD));
-        output[i] = mish_yashas(input[i]);
+        output[i] = mish_yashas2(input[i]);
 }
 
 /**

src/kernels/activation_relu_ceiling.cu

@@ -1,4 +1,13 @@
 #include "kernels.h"
+#include "pluginsRT/ActivationReLUCeilingRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ActivationReLUCeilingCreator::mFC{};
+std::vector<nvinfer1::PluginField> ActivationReLUCeilingCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ActivationReLUCeilingCreator);
+}}
 
 __global__
 void activation_relu_ceiling(dnnType *input, dnnType *output, int size, const float ceiling) {

src/kernels/activation_sigmoid.cu

@@ -1,6 +1,14 @@
-#include "kernels.h"
 #include <math.h>
+#include "kernels.h"
+#include "pluginsRT/ActivationSigmoidRT.h"
 
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ActivationSigmoidRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ActivationSigmoidRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ActivationSigmoidRTCreator);
+}}
 
 __global__
 void activation_sigmoid(dnnType *input, dnnType *output, int size) {

src/kernels/deformable_conv.cu

@@ -3,9 +3,11 @@
 #include <cstring>
 #include <string>
 #include <iostream>
-#include "kernels.h"
 #include <errno.h>
 
+#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<nvinfer1::PluginField> 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) {

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<nvinfer1::PluginField> MaxPoolFixedSizeRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(MaxPoolFixedSizeRTCreator);
+}}
 
 __global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c, int stride_x, int stride_y, int size, int pad, float *input, float *output)
 {

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<nvinfer1::PluginField> ReorgRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ReorgRTCreator);
+}}
 
 __global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, int stride, int forward, float *out)
 {

src/kernels/resize.cu

@@ -1,6 +1,16 @@
-#include "kernels.h"
 #include <stdio.h>
 
+#include "kernels.h"
+#include "pluginsRT/ResizeLayerRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ResizeLayerRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ResizeLayerRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ResizeLayerRTCreator);
+}}
+
 __global__ void resize_kernel(  int size,float *x, int i_w, int i_h, int i_c,
                                 int o_w, int o_h, int o_c, int batch, float *out)
 {

src/kernels/scaladd.cu

@@ -1,6 +1,7 @@
-#include "kernels.h"
 #include <math.h>
 
+#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;

src/kernels/shortcut.cu

@@ -1,6 +1,16 @@
 #include "kernels.h"
 #include "assert.h"
 
+#include "pluginsRT/ShortcutRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ShortcutRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ShortcutRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ShortcutRTCreator);
+}}
+
 __global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stride, int sample, int batch,
                                 int w1, int h1, int c1, dnnType *add,
                                 int w2, int h2, int c2, float s1, float s2, dnnType *out)

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<nvinfer1::PluginField> RegionRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(RegionRTCreator);
+
+nvinfer1::PluginFieldCollection RouteRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> RouteRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(RouteRTCreator);
+
+nvinfer1::PluginFieldCollection ReshapeRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ReshapeRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ReshapeRTCreator);
+
+nvinfer1::PluginFieldCollection FlattenConcatRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> FlattenConcatRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(FlattenConcatRTCreator);
+
+nvinfer1::PluginFieldCollection YoloRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> YoloRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(YoloRTCreator);
+}}

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<nvinfer1::PluginField> UpsampleRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(UpsampleRTCreator);
+}}
 
 __global__ void upsample_kernel(size_t N, dnnType *x, int w, int h, int c, int batch, int stride, int forward, float scale, dnnType *out)
 {

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<const char*>(serialData),*bufCheck = buf;
+  int classes_temp = tk::dnn::readBUF<int>(buf);
+  int num_temp = tk::dnn::readBUF<int>(buf);
+  int n_masks_temp = tk::dnn::readBUF<int>(buf);
+  float scale_xy_temp = tk::dnn::readBUF<float>(buf);
+  float nms_thresh_temp = tk::dnn::readBUF<float>(buf);
+  int nms_kind_temp = tk::dnn::readBUF<int>(buf);
+  int new_coords_temp = tk::dnn::readBUF<int>(buf);
+
+  YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp);
+
+  r->c = tk::dnn::readBUF<int>(buf);
+  r->h = tk::dnn::readBUF<int>(buf);
+  r->w = tk::dnn::readBUF<int>(buf);
+  for(int i=0; i<r->n_masks; i++)
+      r->mask[i] = tk::dnn::readBUF<dnnType>(buf);
+  for(int i=0; i<r->n_masks*2*r->num; i++)
+      r->bias[i] = tk::dnn::readBUF<dnnType>(buf);
+
+  // save classes names
+  r->classesNames.resize(r->classes);
+  for(int i=0; i<r->classes; i++) {
+          char tmp[YOLORT_CLASSNAME_W];
+    for(int j=0; j<YOLORT_CLASSNAME_W; j++)
+      tmp[j] = tk::dnn::readBUF<char>(buf);
+          r->classesNames[i] = std::string(tmp);
+  }
+  assert(buf == bufCheck + serialLength);
+
+  yoloContainer.yolos[yoloContainer.n_yolos++] = r;
+  return r;
+}
+
+}}

tests/mnist/test_mnistRT.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
 	{
 		// suppress info-level messages
 		if (severity != Severity::kINFO)
@@ -66,11 +66,11 @@ int main() {
     std::cout<<"\n==== TensorRT ====\n";
 	// create the builder
 	IBuilder* builder = nvinfer1::createInferBuilder(gLogger);
-	INetworkDefinition* network = builder->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;
@@ -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();
 
@@ -143,9 +154,9 @@ 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 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)));