Added TensorRT8 support

- Migrated to TRT8 API, but I didn't really try maintaining old compatiblity. Lowest support now could be TRT6.
- Refractored some stuff that are bad C++ practices and made coding really hard, like including headers in namespaces or not using an include guard.
- Needed to move the yolo container outside tkdnn object, which means we now only have one and global. Deserialization in TRT8 doesn't happen in your object, but in the plugin itself so it couldn't access the yolo objects. I think the need to hold onto yolo layers itself is flawed and shouldn't be nessesarry.

include/tkDNN/Int8BatchStream.h

@@ -12,7 +12,7 @@
 #include <iomanip>
 #include <signal.h>
 #include <stdlib.h>
-#ifdef __linux__    
+#ifdef __linux__
 #include <unistd.h>
 #endif
 
@@ -23,9 +23,9 @@
 #include "tkdnn.h"
 
 /*
- * BatchStream implements the stream for the INT8 calibrator. 
- * It reads the two files .txt with the list of image file names 
- * and the list of label file names. 
+ * BatchStream implements the stream for the INT8 calibrator.
+ * It reads the two files .txt with the list of image file names
+ * and the list of label file names.
  * It then iterates on images and labels.
  */
 class BatchStream {
@@ -39,7 +39,7 @@ public:
 	float *getLabels() { return mLabels.data(); }
 	int getBatchesRead() const { return mBatchCount; }
 	int getBatchSize() const { return mBatchSize; }
-	nvinfer1::DimsNCHW getDims() const { return mDims; }
+	nvinfer1::Dims4 getDims() const { return mDims; }
 	float* getFileBatch() { return &mFileBatch[0]; }
 	float* getFileLabels() { return &mFileLabels[0]; }
 	void readInListFile(const std::string& dataFilePath, std::vector<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

@@ -20,20 +20,20 @@
 
 /*
  * Int8EntropyCalibrator implements the INT8 calibrator to achieve the
- * INT8 quantization. It uses a BatchStream stream to scroll through 
- * images data. It also implements the calibration cache, a way to 
- * save the calibration process results to reduce the running time: 
+ * INT8 quantization. It uses a BatchStream stream to scroll through
+ * images data. It also implements the calibration cache, a way to
+ * save the calibration process results to reduce the running time:
  * the calibration process takes a long time.
  */
 class Int8EntropyCalibrator : public nvinfer1::IInt8EntropyCalibrator {
 public:
-	Int8EntropyCalibrator(BatchStream& stream, int firstBatch, const std::string& calibTableFilePath, 
+	Int8EntropyCalibrator(BatchStream& stream, int firstBatch, const std::string& calibTableFilePath,
 							const std::string& inputBlobName, bool readCache = true);
 	virtual ~Int8EntropyCalibrator() { checkCuda(cudaFree(mDeviceInput)); }
-	int getBatchSize() const override { return mStream.getBatchSize(); }
-	bool getBatch(void* bindings[], const char* names[], int nbBindings) override;
-	const void* readCalibrationCache(size_t& length) override;
-	void writeCalibrationCache(const void* cache, size_t length) override;
+	int getBatchSize() const noexcept override { return mStream.getBatchSize(); }
+	bool getBatch(void* bindings[], const char* names[], int nbBindings) noexcept override;
+	const void* readCalibrationCache(size_t& length) noexcept override;
+	void writeCalibrationCache(const void* cache, size_t length) noexcept override;
 
 private:
 	BatchStream mStream;

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 {
 
@@ -57,11 +34,11 @@ public:
     nvinfer1::DataType dtRT;
     nvinfer1::IBuilder *builderRT;
     nvinfer1::IRuntime *runtimeRT;
-    nvinfer1::INetworkDefinition *networkRT; 
-#if NV_TENSORRT_MAJOR >= 6  
+    nvinfer1::INetworkDefinition *networkRT;
+#if NV_TENSORRT_MAJOR >= 6
     nvinfer1::IBuilderConfig *configRT;
 #endif
-    
+
     nvinfer1::ICudaEngine *engineRT;
     nvinfer1::IExecutionContext *contextRT;
 
@@ -74,8 +51,6 @@ public:
     dnnType *output;
     cudaStream_t stream;
 
-    PluginFactory *pluginFactory;
-
     NetworkRT(Network *net, const char *name);
     virtual ~NetworkRT();
 
@@ -89,7 +64,7 @@ public:
     int getBuffersN() {
         if(engineRT != nullptr)
             return engineRT->getNbBindings();
-        else 
+        else
             return 0;
     }
 
@@ -97,7 +72,7 @@ public:
         Do inference
     */
     dnnType* infer(dataDim_t &dim, dnnType* data);
-    void enqueue(int batchSize = 1);    
+    void enqueue(int batchSize = 1);
 
     nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l);
     nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l);

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,31 +43,36 @@ public:
 		return 0;
 	}
 
-	virtual void terminate() override {
+	void terminate() override {
 	}
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+	size_t getWorkspaceSize(int maxBatchSize) const override {
 		return 0;
 	}
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+	int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
 
-		activationLEAKYForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
+		activationLEAKYForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
 											reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, slope, stream);
 		return 0;
 	}
 
 
-	virtual size_t getSerializationSize() override {
+	size_t getSerializationSize() override {
 		return 1*sizeof(int) + 1*sizeof(float);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) override {
 		char *buf = reinterpret_cast<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 true;
+    }
+
+    nvinfer1::IPluginV2 * clone() const noexcept override {
+        auto a = new ActivationLogisticRT(*this);
+        return a;
+    }
+
+    const char* getPluginType() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    void destroy() noexcept override {}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+        mNamespace = pluginNamespace;
+    }
+
+    const char* getPluginNamespace() const noexcept override {
+        return mNamespace.c_str();
+    }
+
+    std::string mNamespace;
     int size;
 };
+
+class ActivationLogisticRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationLogisticRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+        const char * buf = reinterpret_cast<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 {
-
-		activationMishForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
+	int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
+		activationMishForward((dnnType*)reinterpret_cast<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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ActivationMishRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int size;
 };
+
+class ActivationMishRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationMishRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+        const char * buf = reinterpret_cast<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 {
-
-		activationReLUCeilingForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
+	int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
+		activationReLUCeilingForward((dnnType*)reinterpret_cast<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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ActivationReLUCeiling(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int size;
 	float ceiling;
 };
+
+class ActivationReLUCeilingCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationReLUCeilingCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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 {
-
-		activationSIGMOIDForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
+	int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
+		activationSIGMOIDForward((dnnType*)reinterpret_cast<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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ActivationSigmoidRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int size;
 };
+
+class ActivationSigmoidRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationSigmoidRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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) {
+	DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw,
+					int deformableGroup, int i_n, int i_c, int i_h, int i_w,
+					int o_n, int o_c, int o_h, int o_w,
+					DeformConv2d *deformable = nullptr) {
 		this->chunk_dim = chunk_dim;
 		this->kh = kh;
 		this->kw = kw;
@@ -30,7 +41,7 @@ public:
 		height_ones = (i_h + 2 * ph - (1 * (kh - 1) + 1)) / sh + 1;
 		width_ones = (i_w + 2 * pw - (1 * (kw - 1) + 1)) / sw + 1;
 		dim_ones = i_c * kh * kw * 1 * height_ones * width_ones;
-		
+
 		checkCuda( cudaMalloc(&data_d, i_c * o_c * kh * kw * 1 * sizeof(dnnType)));
 		checkCuda( cudaMalloc(&bias2_d, o_c*sizeof(dnnType)));
 		checkCuda( cudaMalloc(&ones_d1, height_ones * width_ones * sizeof(dnnType)));
@@ -61,38 +72,45 @@ public:
 		cublasDestroy(handle);
 	}
 
-	int getNbOutputs() const override {
+	int getNbOutputs() const noexcept override {
 		return 1;
 	}
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w};
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
+		return nvinfer1::Dims3{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w};
 	}
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { }
+	void configureWithFormat(nvinfer1::Dims const * inputDims,
+					int32_t nbInputs,
+					nvinfer1::Dims const * outputDims,
+					int32_t nbOutputs,
+					nvinfer1::DataType type,
+					nvinfer1::PluginFormat format,
+					int32_t maxBatchSize) noexcept override {
+	}
 
-	int initialize() override {
+	int initialize() noexcept override {
 		return 0;
 	}
 
-	virtual void terminate() override { }
+	void terminate() noexcept override { }
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+	size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
 		return 0;
 	}
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+	int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
 		dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
 		dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
 
 		// split conv2d outputs into offset to mask
 		for(int b=0; b<batchSize; b++) {
-			checkCuda(cudaMemcpy(offset, output_conv + b * 3 * chunk_dim, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); 
-			checkCuda(cudaMemcpy(mask, output_conv + b * 3 * chunk_dim + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); 
+			checkCuda(cudaMemcpy(offset, output_conv + b * 3 * chunk_dim, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+			checkCuda(cudaMemcpy(mask, output_conv + b * 3 * chunk_dim + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
 			// kernel sigmoid
 			activationSIGMOIDForward(mask, mask, chunk_dim);
 			// deformable convolution
-			dcnV2CudaForward(stat, handle, 
+			dcnV2CudaForward(stat, handle,
 								srcData, data_d,
 								bias2_d, ones_d1,
 								offset, mask,
@@ -109,65 +127,94 @@ 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());
 	}
 
-	cublasStatus_t stat; 
-	cublasHandle_t handle; 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new 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;
 	int o_n, o_c, o_h, o_w;
 	int size;
@@ -179,9 +226,9 @@ public:
 	int height_ones;
 	int width_ones;
 	int dim_ones;
-		
+
 	dnnType *data_d;
-    dnnType *bias2_d;
+  dnnType *bias2_d;
 	dnnType *ones_d1;
 	dnnType * offset;
 	dnnType * mask;
@@ -190,7 +237,100 @@ public:
 	// dnnType *offset_n;
 	// 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,24 +50,24 @@ public:
 		w = 1;
 	}
 
-	int initialize() override {
+	int initialize() noexcept override {
 		return 0;
 	}
 
-	virtual void terminate() override {
+	void terminate() noexcept override {
 		checkERROR(cublasDestroy(handle));
 	}
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+	size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
 		return 0;
 	}
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+	int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
 		dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
 		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
 		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
 
-		checkERROR( cublasSetStream(handle, stream) );	
+		checkERROR( cublasSetStream(handle, stream) );
 		for(int i=0; i<batchSize; i++) {
 			float const alpha(1.0);
 			float const beta(0.0);
@@ -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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new FlattenConcatRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
 	int rows, cols;
-	cublasStatus_t stat; 
-	cublasHandle_t handle; 
+	cublasStatus_t stat;
+	cublasHandle_t handle;
 };
+
+class FlattenConcatRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    FlattenConcatRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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,11 +1,24 @@
-#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) {
-		this->c = c;	
+		this->c = c;
 		this->h = h;
 		this->w = w;
 		this->n = n;
@@ -15,33 +28,37 @@ public:
 		this->padding = padding;
 	}
 
-	~MaxPoolFixedSizeRT(){
-	}
+	~MaxPoolFixedSizeRT() = default;
 
-	int getNbOutputs() const override {
+	int getNbOutputs() const noexcept override {
 		return 1;
 	}
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return DimsCHW{this->c, this->h, this->w};
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
+		return nvinfer1::Dims3{this->c, this->h, this->w};
 	}
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {        
+	void configureWithFormat(nvinfer1::Dims const * inputDims,
+					int32_t nbInputs,
+					nvinfer1::Dims const * outputDims,
+					int32_t nbOutputs,
+					nvinfer1::DataType type,
+					nvinfer1::PluginFormat format,
+					int32_t maxBatchSize) noexcept override {
 	}
 
-	int initialize() override {
+	int initialize() noexcept override {
 		return 0;
 	}
 
-	virtual void terminate() override {
+	void terminate() noexcept override {
 	}
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+	size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
 		return 0;
 	}
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
-
+	int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
 		//std::cout<<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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new MaxPoolFixedSizeRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int n, c, h, w;
 	int stride_H, stride_W;
 	int winSize;
 	int padding;
 };
+
+class MaxPoolFixedSizeRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    MaxPoolFixedSizeRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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]);
 
@@ -52,7 +66,7 @@ public:
 			for(int n = 0; n < num; ++n){
 				int index = entry_index(b, n*w*h, 0);
 				activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
-				
+
 				index = entry_index(b, n*w*h, coords);
 				activationLOGISTICForward(srcData + index, dstData + index, w*h, stream);
 			}
@@ -60,31 +74,61 @@ public:
 
 		//softmax start
 		int index = entry_index(0, 0, coords + 1);
-		softmaxForward(	srcData + index, classes, batchSize*num, 
-						(c*h*w)/num, 
+		softmaxForward(	srcData + index, classes, batchSize*num,
+						(c*h*w)/num,
 						w*h, 1, w*h, 1, dstData + index, stream);
 
 		return 0;
 	}
 
 
-	virtual size_t getSerializationSize() override {
+	size_t getSerializationSize() const noexcept override {
 		return 6*sizeof(int);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) const noexcept override {
 		char *buf = reinterpret_cast<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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new RegionRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
-    int classes, coords, num;
+  int classes, coords, num;
 
 	int entry_index(int batch, int location, int entry) {
 		int n =   location / (w*h);
@@ -93,3 +137,57 @@ public:
 	}
 
 };
+
+class RegionRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    RegionRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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 {
-
-		reorgForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
-					  reinterpret_cast<dnnType*>(outputs[0]), 
+  int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
+		reorgForward((dnnType*)reinterpret_cast<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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ReorgRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w, stride;
 };
+
+class ReorgRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ReorgRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ReshapeRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int n, c, h, w;
 };
+
+class ReshapeRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ReshapeRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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,68 +1,168 @@
-#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) {
 		o_c = c;
 		o_h = h;
-		o_w = w;	
+		o_w = w;
 	}
 
-	~ResizeLayerRT(){
-	}
+	~ResizeLayerRT() = default;
 
-	int getNbOutputs() const override {
+	int getNbOutputs() const noexcept override {
 		return 1;
 	}
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return DimsCHW{o_c, o_h, o_w};
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
+		return nvinfer1::Dims3{o_c, o_h, o_w};
 	}
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+	void configureWithFormat(nvinfer1::Dims const * inputDims,
+					int32_t nbInputs,
+					nvinfer1::Dims const * outputDims,
+					int32_t nbOutputs,
+					nvinfer1::DataType type,
+					nvinfer1::PluginFormat format,
+					int32_t maxBatchSize) noexcept override {
 		i_c = inputDims[0].d[0];
 		i_h = inputDims[0].d[1];
-		i_w = inputDims[0].d[2];        
+		i_w = inputDims[0].d[2];
 	}
 
-	int initialize() override {
+	int initialize() noexcept override {
 		return 0;
 	}
 
-	virtual void terminate() override {
+	void terminate() noexcept override {
 	}
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+	size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
 		return 0;
 	}
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+	int enqueue(int batchSize, const void*const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
     	// printf("%d %d %d %d %d %d\n", i_c, i_w, i_h, o_c, o_w, o_h);
-        resizeForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
-					  reinterpret_cast<dnnType*>(outputs[0]), 
+        resizeForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
+					  reinterpret_cast<dnnType*>(outputs[0]),
 					  batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream);
 		return 0;
 	}
 
-
-	virtual size_t getSerializationSize() override {
+	size_t getSerializationSize() const noexcept override {
 		return 6*sizeof(int);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) const noexcept override {
 		char *buf = reinterpret_cast<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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ResizeLayerRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int i_c, i_h, i_w, o_c, o_h, o_w;
 };
+
+class ResizeLayerRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ResizeLayerRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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,25 +26,29 @@ public:
 		this->group_id = group_id;
 	}
 
-	~RouteRT(){
+	~RouteRT() = default;
 
-	}
-
-	int getNbOutputs() const override {
+	int getNbOutputs() const noexcept override {
 		return 1;
 	}
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
 		int out_c = 0;
 		for(int i=0; i<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++) {
-			c_in[i] = inputDims[i].d[0]; 
+			c_in[i] = inputDims[i].d[0];
 			c += inputDims[i].d[0];
 		}
 		h = inputDims[0].d[1];
@@ -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 true;
+	}
+
+	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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ShortcutRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
 	int bc, bh, bw;
 	bool mul;
 };
+
+class ShortcutRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ShortcutRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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,66 +1,160 @@
-#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]);
-	    
+
 		fill(dstData, batchSize*c*h*w*stride*stride, 0.0, stream);
     	upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream);
 		return 0;
 	}
 
-
-	virtual size_t getSerializationSize() override {
+	size_t getSerializationSize() const noexcept override {
 		return 4*sizeof(int);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) const noexcept override {
 		char *buf = reinterpret_cast<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 true;
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new UpsampleRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int c, h, w, stride;
 };
+
+class UpsampleRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    UpsampleRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<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 true;
+	}
+
+	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