Merge of perseusdg-tensorrt8 inside tkDNN

Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>

include/tkDNN/pluginsRT/ActivationLeakyRT.h

@@ -1,61 +1,46 @@
-#include<cassert>
+#include "NvInfer.h"
 #include "../kernels.h"
+#include <cassert>
+#include <vector>
 
-class ActivationLeakyRT : public IPluginV2 {
+namespace nvinfer1 {
+    class ActivationLeakyRT : public IPluginV2 {
 
-public:
-	ActivationLeakyRT() {
+    public:
+        explicit ActivationLeakyRT(float s);
 
+        ActivationLeakyRT(const void *data, size_t length);
 
-	}
+        ~ActivationLeakyRT();
 
-	~ActivationLeakyRT(){
+        int getNbOutputs() const NOEXCEPT override;
 
-	}
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
 
-	int getNbOutputs() const noexcept override {
-		return 1;
-	}
+        void
+        configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,
+                            PluginFormat format, int maxBatchSize) NOEXCEPT override;
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) noexcept override {
-		return inputs[0];
-	}
+        int initialize() NOEXCEPT override;
 
-	//void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-	//	size = 1;
-	//	for(int i=0; i<outputDims[0].nbDims; i++)
-	//		size *= outputDims[0].d[i];
-	//}
+        void terminate() NOEXCEPT override {}
 
-	int initialize() noexcept override {
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
 
-		return 0;
-	}
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, void const *const *inputs, void *const *outputs, void *workspace,
+                    cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
-	virtual void terminate() noexcept override {
-	}
+        size_t getSerializationSize() const NOEXCEPT override;
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
-		return 0;
-	}
+        void serialize(void *buffer) const NOEXCEPT override;
 
-	virtual int enqueue(int32_t batchSize, void const *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) noexcept override {
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
 
-		activationLEAKYForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
-											reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
-		return 0;
-	}
+        const char *getPluginType() const NOEXCEPT override;
 
-
-	virtual size_t getSerializationSize() const noexcept override {
-		return 1*sizeof(int);
-	}
-
-	virtual void serialize(void* buffer) const noexcept override {
-		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
-		tk::dnn::writeBUF(buf, size);
-		assert(buf == a + getSerializationSize());
-	}
-
-	int size;
-};
+    REGISTER_TENSORRT_PLUGIN(ActivationLeakyRTPluginCreator);
+};

include/tkDNN/pluginsRT/ActivationLogisticRT.h

@@ -1,60 +1,88 @@
 #include<cassert>
 #include "../kernels.h"
+#include <NvInfer.h>
+#include <vector>
+#include <utils.h>
 
-class ActivationLogisticRT : public IPlugin {
+namespace nvinfer1 {
 
-public:
-    ActivationLogisticRT() {
+    class ActivationLogisticRT : public IPluginV2 {
+
+    public:
+        ActivationLogisticRT() ;
+
+        ActivationLogisticRT(const void *data, size_t length) ;
+
+        ~ActivationLogisticRT() ;
+
+        int getNbOutputs() const NOEXCEPT override ;
+
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+        void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,
+                            PluginFormat format, int maxBatchSize) NOEXCEPT override ;
+
+        int initialize() NOEXCEPT override ;
+
+        void terminate() NOEXCEPT override ;
+
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
+
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                            cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
 
-    }
+        size_t getSerializationSize() const NOEXCEPT override ;
 
-    ~ActivationLogisticRT(){
+        void serialize(void *buffer) const NOEXCEPT override ;
 
-    }
+        const char *getPluginType() const NOEXCEPT override ;
 
-    int getNbOutputs() const override {
-        return 1;
-    }
+        const char *getPluginVersion() const NOEXCEPT override ;
 
-    Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-        return inputs[0];
-    }
+        void destroy() NOEXCEPT override ;
 
-    void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-        size = 1;
-        for(int i=0; i<outputDims[0].nbDims; i++)
-            size *= outputDims[0].d[i];
-    }
+        const char *getPluginNamespace() const NOEXCEPT override ;
 
-    int initialize() override {
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
 
-        return 0;
-    }
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
 
-    virtual void terminate() override {
-    }
+        IPluginV2 *clone() const NOEXCEPT override ;
 
-    virtual size_t getWorkspaceSize(int maxBatchSize) const override {
-        return 0;
-    }
+        int size;
 
-    virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+    private:
+        std::string mPluginNamespace;
+    };
 
-        activationLOGISTICForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
-                                  reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
-        return 0;
-    }
+    class ActivationLogisticRTPluginCreator : public IPluginCreator {
+    public:
+        ActivationLogisticRTPluginCreator() ;
 
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
 
-    virtual size_t getSerializationSize() override {
-        return 1*sizeof(int);
-    }
+        IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
 
-    virtual void serialize(void* buffer) override {
-        char *buf = reinterpret_cast<char*>(buffer);
-        tk::dnn::writeBUF(buf, size);
-    }
+        const char *getPluginNamespace() const NOEXCEPT override ;
 
-    int size;
-};
+        IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(ActivationLogisticRTPluginCreator);
+};

include/tkDNN/pluginsRT/ActivationMishRT.h

@@ -1,61 +1,82 @@
 #include<cassert>
 #include "../kernels.h"
+#include <NvInfer.h>
+#include <vector>
 
-class ActivationMishRT : public IPlugin {
+namespace nvinfer1 {
+    class ActivationMishRT : public IPluginV2 {
 
-public:
-	ActivationMishRT() {
+    public:
+        ActivationMishRT() ;
+
+        ~ActivationMishRT() ;
+
+        ActivationMishRT(const void *data, size_t length) ;
 
 
-	}
+        int getNbOutputs() const NOEXCEPT override ;
 
-	~ActivationMishRT(){
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
 
-	}
+        void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,
+                            PluginFormat format, int maxBatchSize) NOEXCEPT override ;
 
-	int getNbOutputs() const override {
-		return 1;
-	}
+        int initialize() NOEXCEPT override ;
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return inputs[0];
-	}
+        void terminate() NOEXCEPT override ;
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		size = 1;
-		for(int i=0; i<outputDims[0].nbDims; i++)
-			size *= outputDims[0].d[i];
-	}
+         size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
-	int initialize() override {
+        size_t getSerializationSize() const NOEXCEPT override ;
 
-		return 0;
-	}
+        void serialize(void *buffer) const NOEXCEPT override ;
 
-	virtual void terminate() override {
-	}
+        const char *getPluginType() const NOEXCEPT override ;
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
-		return 0;
-	}
+        const char *getPluginVersion() const NOEXCEPT override ;
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+        void destroy() NOEXCEPT override { delete this; }
 
-		activationMishForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
-											reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
-		return 0;
-	}
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
 
+        const char *getPluginNamespace() const NOEXCEPT override ;
 
-	virtual size_t getSerializationSize() override {
-		return 1*sizeof(int);
-	}
+        void setPluginNamespace(const char *plguinNamespace) NOEXCEPT override ;
 
-	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
-		tk::dnn::writeBUF(buf, size);
-		assert(buf == a + getSerializationSize());
-	}
+        IPluginV2 *clone() const NOEXCEPT override ;
 
-	int size;
-};
+        int size;
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class ActivationMishRTPluginCreator : public IPluginCreator {
+    public:
+        ActivationMishRTPluginCreator() ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(ActivationMishRTPluginCreator);
+};

include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h

@@ -1,63 +1,81 @@
 #include<cassert>
 #include "../kernels.h"
+#include <NvInfer.h>
+#include <vector>
+#include <utils.h>
 
-class ActivationReLUCeiling : public IPlugin {
+namespace nvinfer1 {
+    class ActivationReLUCeiling : public IPluginV2 {
 
-public:
-	ActivationReLUCeiling(const float ceiling) {
-		this->ceiling = ceiling;
-	}
+    public:
+        explicit ActivationReLUCeiling(const float ceiling) ;
 
-	~ActivationReLUCeiling(){
+        ~ActivationReLUCeiling() ;
 
-	}
+        ActivationReLUCeiling(const void *data, size_t length) ;
 
-	int getNbOutputs() const override {
-		return 1;
-	}
+        int getNbOutputs() const NOEXCEPT override ;
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return inputs[0];
-	}
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		size = 1;
-		for(int i=0; i<outputDims[0].nbDims; i++)
-			size *= outputDims[0].d[i];
-	}
+        void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override ;
 
-	int initialize() override {
+        int initialize() NOEXCEPT override ;
 
-		return 0;
-	}
+        void terminate() NOEXCEPT override ;
 
-	virtual void terminate() override {
-	}
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
-		return 0;
-	}
+        size_t getSerializationSize() const NOEXCEPT override ;
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+        void serialize(void *buffer) const NOEXCEPT override ;
 
-		activationReLUCeilingForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
-											reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, ceiling, stream);
-		return 0;
-	}
+        IPluginV2 *clone() const NOEXCEPT override ;
 
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
 
-	virtual size_t getSerializationSize() override {
-		return 1*sizeof(int) +  1*sizeof(float);
-	}
+        void destroy() NOEXCEPT override ;
 
-	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
-		tk::dnn::writeBUF(buf, ceiling);
-		tk::dnn::writeBUF(buf, size);
-		assert(buf = a + getSerializationSize());
-		
-	}
+        const char *getPluginType() const NOEXCEPT override ;
 
-	int size;
-	float ceiling;
-};
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+        int size;
+        float ceiling;
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class ActivationReLUCeilingPluginCreator : public IPluginCreator {
+    public:
+        ActivationReLUCeilingPluginCreator() ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+    public:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(ActivationReLUCeilingPluginCreator);
+};

include/tkDNN/pluginsRT/DeformableConvRT.h

@@ -1,196 +1,137 @@
+#ifndef _DEFORMABLECONVRT_PLUGIN_H
+#define _DEFORMABLECONVRT_PLUGIN_H
+
+#include <NvInfer.h>
+#include <vector>
 #include<cassert>
 #include "../kernels.h"
+#include <tkdnn.h>
+
+namespace nvinfer1 {
+    class DeformableConvRT : public IPluginV2Ext {
 
 
-class DeformableConvRT : public IPlugin {
+    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,std::vector<dnnType> data_H,std::vector<dnnType> bias2_H,
+                         std::vector<dnnType> ones_d1_h,std::vector<dnnType> ones_d2_h,std::vector<dnnType> offsetH,std::vector<dnnType> maskH,int height_ones,
+                         int width_ones,int dim_ones);
+
+        ~DeformableConvRT();
+
+        DeformableConvRT(const void *data, size_t length) ;
+
+        int getNbOutputs() const NOEXCEPT override ;
+
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+        int initialize() NOEXCEPT override ;
+
+        void terminate() NOEXCEPT override ;
+
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                            cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+        size_t getSerializationSize() const NOEXCEPT override ;
+
+        void serialize(void *buffer) const NOEXCEPT override ;
+
+        void destroy() NOEXCEPT override ;
+
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginType() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        IPluginV2Ext *clone() const NOEXCEPT override ;
+
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+                              int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+                              bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+                              int32_t maxBatchSize) NOEXCEPT override;
+
+        void detachFromContext() NOEXCEPT override;
 
 
+        cublasStatus_t stat;
+        cublasHandle_t handle{nullptr};
+        int i_n, i_c, i_h, i_w;
+        int o_n, o_c, o_h, o_w;
+        int size;
+        int chunk_dim;
+        int kh, kw;
+        int sh, sw;
+        int ph, pw;
+        int deformableGroup;
+        int height_ones;
+        int width_ones;
+        int dim_ones;
 
-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) {
-		this->chunk_dim = chunk_dim;
-		this->kh = kh;
-		this->kw = kw;
-		this->sh = sh;
-		this->sw = sw;
-		this->ph = ph;
-		this->pw = pw;
-		this->deformableGroup = deformableGroup;
-		this->i_n = i_n;
-		this->i_c = i_c;
-		this->i_h = i_h;
-		this->i_w = i_w;
-		this->o_n = o_n;
-		this->o_c = o_c;
-		this->o_h = o_h;
-		this->o_w = o_w;
-		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;
-		std::cout<<i_c * o_c * kh * kw * 1<<"\n";
-		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)));
-		checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType)));
-		checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType)));
-		checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType)));
-		if(deformable != nullptr) {
-			this->defRT = deformable;
-			checkCuda( cudaMemcpy(data_d, deformable->data_d, sizeof(dnnType)*i_c * o_c * kh * kw * 1, cudaMemcpyDeviceToDevice) );
-			checkCuda( cudaMemcpy(bias2_d, deformable->bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToDevice) );
-            checkCuda( cudaMemcpy(ones_d1, deformable->ones_d1, sizeof(dnnType)*height_ones*width_ones, cudaMemcpyDeviceToDevice) );
-            checkCuda( cudaMemcpy(offset, deformable->offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToDevice) );
-            checkCuda( cudaMemcpy(mask, deformable->mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToDevice) );
-            checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) );
-		}
-		stat = cublasCreate(&handle);
-		if (stat != CUBLAS_STATUS_SUCCESS)
-			FatalError("CUBLAS initialization failed\n");
-	}
-
-	~DeformableConvRT() {
-		checkCuda( cudaFree(data_d) );
-		checkCuda( cudaFree(bias2_d) );
-		checkCuda( cudaFree(ones_d1) );
-		checkCuda( cudaFree(offset) );
-		checkCuda( cudaFree(mask) );
-		checkCuda( cudaFree(ones_d2) );
-		cublasDestroy(handle);
-	}
-
-	int getNbOutputs() const 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};
-	}
-
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { }
-
-	int initialize() override {
-		return 0;
-	}
-
-	virtual void terminate() override { }
-
-	virtual 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 {
-		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)); 
-			// kernel sigmoid
-			activationSIGMOIDForward(mask, mask, chunk_dim);
-			// deformable convolution
-			dcnV2CudaForward(stat, handle, 
-								srcData, data_d,
-								bias2_d, ones_d1,
-								offset, mask,
-								reinterpret_cast<dnnType*>(outputs[0]), ones_d2,
-								kh, kw,
-								sh, sw,
-								ph, pw,
-								1, 1,
-								deformableGroup, b,
-								i_n, i_c, i_h, i_w,
-								o_n, o_c, o_h, o_w,
-								chunk_dim);
-		}
-		return 0;
-	}
+        std::vector<dnnType> data_d_v;
+        std::vector<dnnType> bias2_d_v;
+        std::vector<dnnType> ones_d1_v;
+        std::vector<dnnType> offset_v;
+        std::vector<dnnType> mask_v;
+        std::vector<dnnType> ones_d2_v;
+        dnnType* data_d;
+        dnnType* bias2_d;
+        dnnType* ones_d1;
+        dnnType* offset;
+        dnnType* mask;
+        dnnType* ones_d2;
+        // dnnType *input_n;
+        // dnnType *offset_n;
+        // dnnType *mask_n;
+        // dnnType *output_n;
 
 
-	virtual size_t getSerializationSize() 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);
-	}
+        tk::dnn::DeformConv2d *defRT;
 
-	virtual void serialize(void* buffer) 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);
-		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]);
-		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]);
-		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]);
-		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]);
-		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]);
-		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]);
-		free(aus);
-		assert(buf == a + getSerializationSize());
-	}
+    private:
+        std::string mPluginNamespace;
+    };
 
-	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;
-	int chunk_dim;
-	int kh, kw;
-	int sh, sw;
-	int ph, pw;
-	int deformableGroup;
-	int height_ones;
-	int width_ones;
-	int dim_ones;
-		
-	dnnType *data_d;
-    dnnType *bias2_d;
-	dnnType *ones_d1;
-	dnnType * offset;
-	dnnType * mask;
-	dnnType *ones_d2;
-	// dnnType *input_n;
-	// dnnType *offset_n;
-	// dnnType *mask_n;
-	// dnnType *output_n;
-	
-	
-	tk::dnn::DeformConv2d *defRT;
+    class DeformableConvRTPluginCreator : public IPluginCreator {
+    public:
+        DeformableConvRTPluginCreator();
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+    REGISTER_TENSORRT_PLUGIN(DeformableConvRTPluginCreator);
 };
+#endif

include/tkDNN/pluginsRT/FlattenConcatRT.h

@@ -1,81 +1,96 @@
 #include<cassert>
+#include <NvInfer.h>
+#include <vector>
+#include <utils.h>
+namespace nvinfer1 {
+    class FlattenConcatRT : public IPluginV2Ext {
 
-class FlattenConcatRT : public IPlugin {
+    public:
+        FlattenConcatRT(int c,int h,int w,int rows,int cols) ;
 
-public:
-	FlattenConcatRT() {
-		stat = cublasCreate(&handle);
-		if (stat != CUBLAS_STATUS_SUCCESS) {
-			printf ("CUBLAS initialization failed\n");
-			return;
-  		}
-	}
+        FlattenConcatRT(const void *data, size_t length) ;
 
-	~FlattenConcatRT(){
+        ~FlattenConcatRT() ;
 
-	}
+        int getNbOutputs() const NOEXCEPT override ;
 
-	int getNbOutputs() const override {
-		return 1;
-	}
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
 
-	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};
-	}
+        int initialize() NOEXCEPT override ;
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		assert(nbOutputs == 1 && nbInputs ==1);
-		rows = inputDims[0].d[0];
-		cols = inputDims[0].d[1] * inputDims[0].d[2];
-		c = inputDims[0].d[0] * inputDims[0].d[1] * inputDims[0].d[2];
-		h = 1;
-		w = 1;
-	}
+        void terminate() NOEXCEPT override ;
 
-	int initialize() override {
-		return 0;
-	}
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
 
-	virtual void terminate() override {
-		checkERROR(cublasDestroy(handle));
-	}
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR <= 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
-		return 0;
-	}
+        size_t getSerializationSize() const NOEXCEPT override ;
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) 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));
+        void serialize(void *buffer) const NOEXCEPT override ;
 
-		checkERROR( cublasSetStream(handle, stream) );	
-		for(int i=0; i<batchSize; i++) {
-			float const alpha(1.0);
-			float const beta(0.0);
-			int offset = i*rows*cols;
-			checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, rows, cols, &alpha, srcData + offset, cols, &beta, srcData + offset, rows, dstData + offset, rows ));
-		}
-		return 0;
-	}
+        void destroy() NOEXCEPT override ;
 
+        const char *getPluginType() const NOEXCEPT override ;
 
-	virtual size_t getSerializationSize() override {
-		return 5*sizeof(int);
-	}
+        const char *getPluginVersion() const NOEXCEPT override;
 
-	virtual void serialize(void* buffer) 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);
-		assert(buf == a + getSerializationSize());
-	}
+        const char *getPluginNamespace() const NOEXCEPT override ;
 
-	int c, h, w;
-	int rows, cols;
-	cublasStatus_t stat; 
-	cublasHandle_t handle; 
-};
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        IPluginV2Ext *clone() const NOEXCEPT override ;
+
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+                            int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+                            bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+                            int32_t maxBatchSize) NOEXCEPT override;
+
+        void detachFromContext() NOEXCEPT override;
+
+        bool supportsFormat (DataType type, PluginFormat format) const NOEXCEPT override;
+
+        int c, h, w;
+        int rows, cols;
+        cublasHandle_t handle{nullptr};
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class FlattenConcatRTPluginCreator : public IPluginCreator {
+    public:
+        FlattenConcatRTPluginCreator() ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(FlattenConcatRTPluginCreator);
+};

include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h

@@ -1,75 +1,105 @@
 #include<cassert>
 #include "../kernels.h"
-
-class MaxPoolFixedSizeRT : public IPlugin {
-
-public:
-	MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) {
-		this->c = c;	
-		this->h = h;
-		this->w = w;
-		this->n = n;
-		this->stride_H = strideH;
-		this->stride_W = strideW;
-		this->winSize = winSize;
-		this->padding = padding;
-	}
-
-	~MaxPoolFixedSizeRT(){
-	}
-
-	int getNbOutputs() const override {
-		return 1;
-	}
-
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return DimsCHW{this->c, this->h, this->w};
-	}
-
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {        
-	}
-
-	int initialize() override {
-		return 0;
-	}
-
-	virtual void terminate() override {
-	}
-
-	virtual 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 {
-
-		//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]);
-		MaxPoolingForward(srcData, dstData, batchSize, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding, stream);
-		return 0;
-	}
+#include <NvInfer.h>
+#include <vector>
+#include <utils.h>
 
 
-	virtual size_t getSerializationSize() override {
-		return 8*sizeof(int);
-	}
+namespace nvinfer1 {
+    class MaxPoolFixedSizeRT : public IPluginV2Ext {
 
-	virtual void serialize(void* buffer) override {
-		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    public:
+        MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding) ;
 
-		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);
-		assert(buf == a + getSerializationSize());
-	}
+        MaxPoolFixedSizeRT(const void *data, size_t length) ;
 
-	int n, c, h, w;
-	int stride_H, stride_W;
-	int winSize;
-	int padding;
+        ~MaxPoolFixedSizeRT() ;
+
+        int getNbOutputs() const NOEXCEPT override ;
+
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+        int initialize() NOEXCEPT override ;
+
+        void terminate() NOEXCEPT override ;
+
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                            cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR <= 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+
+        size_t getSerializationSize() const NOEXCEPT override ;
+
+        void serialize(void *buffer) const NOEXCEPT override ;
+
+        void destroy() NOEXCEPT override ;
+
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginType() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        IPluginV2Ext *clone() const NOEXCEPT override ;
+
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+                              int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+                              bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+                              int32_t maxBatchSize) NOEXCEPT override;
+
+        void detachFromContext() NOEXCEPT override;
+
+
+        int n, c, h, w;
+        int stride_H, stride_W;
+        int winSize;
+        int padding;
+
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class MaxPoolFixedSizeRTPluginCreator : public IPluginCreator {
+    public:
+        MaxPoolFixedSizeRTPluginCreator() ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+
+    };
+
+    REGISTER_TENSORRT_PLUGIN(MaxPoolFixedSizeRTPluginCreator);
 };

include/tkDNN/pluginsRT/RegionRT.h

@@ -1,95 +1,110 @@
+#ifndef _REGIONRT_PLUGIN_H
+#define _REGIONRT_PLUGIN_H
 #include<cassert>
 #include "../kernels.h"
+#include <NvInfer.h>
+#include <vector>
+#include <utils.h>
 
-class RegionRT : public IPlugin {
+namespace nvinfer1 {
+    class RegionRT : public IPluginV2Ext {
 
-public:
-	RegionRT(int classes, int coords, int num) {
+    public:
+        RegionRT(int classes, int coords, int num,int c,int h,int w);
 
-		this->classes = classes;
-		this->coords = coords;
-		this->num = num;
-	}
+        ~RegionRT() ;
 
-	~RegionRT(){
+        RegionRT(const void *data, size_t length) ;
 
-	}
+        int getNbOutputs() const NOEXCEPT override ;
 
-	int getNbOutputs() const override {
-		return 1;
-	}
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return inputs[0];
-	}
-
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		c = inputDims[0].d[0];
-		h = inputDims[0].d[1];
-		w = inputDims[0].d[2];
-	}
-
-	int initialize() override {
-
-		return 0;
-	}
-
-	virtual void terminate() override {
-	}
-
-	virtual 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 {
-
-		dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
-		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
-
-		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
-
-		for (int b = 0; b < batchSize; ++b){
-			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);
-			}
-		}
-
-		//softmax start
-		int index = entry_index(0, 0, coords + 1);
-		softmaxForward(	srcData + index, classes, batchSize*num, 
-						(c*h*w)/num, 
-						w*h, 1, w*h, 1, dstData + index, stream);
-
-		return 0;
-	}
+        int initialize() NOEXCEPT override ;
 
 
-	virtual size_t getSerializationSize() override {
-		return 6*sizeof(int);
-	}
+        void terminate() NOEXCEPT override ;
 
-	virtual void serialize(void* buffer) 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);
-		assert(buf == a + getSerializationSize());
-	}
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
 
-	int c, h, w;
-    int classes, coords, num;
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                    cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
-	int entry_index(int batch, int location, int entry) {
-		int n =   location / (w*h);
-		int loc = location % (w*h);
-		return batch*c*h*w + n*w*h*(coords+classes+1) + entry*w*h + loc;
-	}
 
+        size_t getSerializationSize() const NOEXCEPT override ;
+
+        void serialize(void *buffer) const NOEXCEPT override ;
+
+        const char *getPluginType() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        void destroy() NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+                              int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+                              bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+                              int32_t maxBatchSize) NOEXCEPT override;
+
+        void detachFromContext() NOEXCEPT override;
+
+        IPluginV2Ext *clone() const NOEXCEPT override ;
+        int c, h, w;
+        int classes, coords, num;
+
+        int entry_index(int batch, int location, int entry) {
+            int n = location / (w * h);
+            int loc = location % (w * h);
+            return batch * c * h * w + n * w * h * (coords + classes + 1) + entry * w * h + loc;
+        }
+
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class RegionRTPluginCreator : public IPluginCreator {
+    public:
+        RegionRTPluginCreator();
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(RegionRTPluginCreator);
 };
+
+#endif

include/tkDNN/pluginsRT/ReorgRT.h

@@ -1,64 +1,98 @@
 #include<cassert>
 #include "../kernels.h"
+#include <NvInfer.h>
+#include <vector>
 
-class ReorgRT : public IPlugin {
+namespace nvinfer1 {
+    class ReorgRT : public IPluginV2Ext {
 
-public:
-	ReorgRT(int stride) {
-		this->stride = stride;
-	}
+    public:
+        ReorgRT(int stride,int c,int h,int w);
 
-	~ReorgRT(){
+        ~ReorgRT();
 
-	}
+        ReorgRT(const void *data, size_t length);
 
-	int getNbOutputs() const override {
-		return 1;
-	}
+        int getNbOutputs() const NOEXCEPT override;
 
-	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};
-	}
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		c = inputDims[0].d[0];
-		h = inputDims[0].d[1];
-		w = inputDims[0].d[2];
-	}
+        int initialize() NOEXCEPT override;
 
-	int initialize() override {
+        void terminate() NOEXCEPT override;
 
-		return 0;
-	}
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
 
-	virtual void terminate() override {
-	}
-
-	virtual 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 {
-
-		reorgForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
-					  reinterpret_cast<dnnType*>(outputs[0]), 
-					  batchSize, c, h, w, stride, stream);
-		return 0;
-	}
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                    cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
 
-	virtual size_t getSerializationSize() override {
-		return 4*sizeof(int);
-	}
+        size_t getSerializationSize() const NOEXCEPT override;
 
-	virtual void serialize(void* buffer) 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);
-		assert(buf == a + getSerializationSize());
-	}
+        void serialize(void *buffer) const NOEXCEPT override;
 
-	int c, h, w, stride;
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
+
+        const char *getPluginType() const NOEXCEPT override;
+
+        const char *getPluginVersion() const NOEXCEPT override;
+
+        void destroy() NOEXCEPT override;
+
+        const char *getPluginNamespace() const NOEXCEPT override;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+        IPluginV2Ext *clone() const NOEXCEPT override;
+
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+                              int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+                              bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+                              int32_t maxBatchSize) NOEXCEPT override;
+
+        void detachFromContext() NOEXCEPT override;
+
+        int c, h, w, stride;
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class ReorgRTPluginCreator : public IPluginCreator {
+    public:
+        ReorgRTPluginCreator();
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+        const char *getPluginNamespace() const NOEXCEPT override;
+
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override;
+
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override;
+
+        const char *getPluginName() const NOEXCEPT override;
+
+        const char *getPluginVersion() const NOEXCEPT override;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(ReorgRTPluginCreator);
 };
+

include/tkDNN/pluginsRT/ReshapeRT.h

@@ -1,62 +1,102 @@
+#ifndef _RESHAPERT_PLUGIN_H
+#define _RESHAPERT_PLUGIN_H
+
 #include<cassert>
-
-class ReshapeRT : public IPlugin {
-
-public:
-	ReshapeRT(dataDim_t new_dim) {
-		n = new_dim.n;
-		c = new_dim.c;
-		h = new_dim.h;
-		w = new_dim.w;
-	}
-
-	~ReshapeRT(){
-
-	}
-
-	int getNbOutputs() const override {
-		return 1;
-	}
-
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return DimsCHW{ c,h,w};
-	}
-
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-	}
-
-	int initialize() override {
-		return 0;
-	}
-
-	virtual void terminate() override {
-	}
-
-	virtual 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 {
-		dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
-		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
-
-		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
-		return 0;
-	}
+#include <NvInfer.h>
+#include <vector>
+#include <tkdnn.h>
+using namespace tk::dnn;
 
 
-	virtual size_t getSerializationSize() override {
-		return 4*sizeof(int);
-	}
+namespace nvinfer1 {
+    class ReshapeRT : public IPluginV2Ext {
 
-	virtual void serialize(void* buffer) 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);
-		assert(buf == a + getSerializationSize());
-	}
+    public:
+        ReshapeRT(int n,int c,int h,int w) ;
 
-	int n, c, h, w;
+        ReshapeRT(const void *data, size_t length) ;
+
+        ~ReshapeRT() ;
+
+        int getNbOutputs() const NOEXCEPT override ;
+
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
+
+
+        int initialize() NOEXCEPT override ;
+
+        void terminate() NOEXCEPT override ;
+
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace, cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+        size_t getSerializationSize() const NOEXCEPT override ;
+
+        void serialize(void *buffer) const NOEXCEPT override ;
+
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+        const char *getPluginType() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        void destroy() NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        IPluginV2Ext *clone() const NOEXCEPT override ;
+
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+                              int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+                              bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+                              int32_t maxBatchSize) NOEXCEPT override;
+
+        void detachFromContext() NOEXCEPT override;
+
+        int n, c, h, w;
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class ReshapeRTPluginCreator : public IPluginCreator {
+    public:
+        ReshapeRTPluginCreator() ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(ReshapeRTPluginCreator);
 };
+#endif

include/tkDNN/pluginsRT/ResizeLayerRT.h

@@ -1,68 +1,104 @@
 #include<cassert>
 #include "../kernels.h"
+#include <NvInfer.h>
+#include <vector>
+#include <utils.h>
 
-class ResizeLayerRT : public IPlugin {
+namespace nvinfer1 {
 
-public:
-	ResizeLayerRT(int c, int h, int w) {
-		o_c = c;
-		o_h = h;
-		o_w = w;	
-	}
+    class ResizeLayerRT : public IPluginV2Ext {
 
-	~ResizeLayerRT(){
-	}
+    public:
+        ResizeLayerRT(int oc, int oh, int ow,int ic,int ih,int iw) ;
 
-	int getNbOutputs() const override {
-		return 1;
-	}
+        ResizeLayerRT(const void *data, size_t length) ;
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return DimsCHW{o_c, o_h, o_w};
-	}
+        ~ResizeLayerRT() ;
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		i_c = inputDims[0].d[0];
-		i_h = inputDims[0].d[1];
-		i_w = inputDims[0].d[2];        
-	}
+        int getNbOutputs() const NOEXCEPT override ;
 
-	int initialize() override {
-		return 0;
-	}
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
 
-	virtual void terminate() override {
-	}
+        int initialize() NOEXCEPT override ;
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
-		return 0;
-	}
+        void terminate() NOEXCEPT override ;
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) 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]), 
-					  batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream);
-		return 0;
-	}
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
+
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                            cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR <= 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+        size_t getSerializationSize() const NOEXCEPT override ;
+
+        void serialize(void *buffer) const NOEXCEPT override ;
+
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+        const char *getPluginType() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        void destroy() NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        IPluginV2Ext *clone() const NOEXCEPT override ;
+
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+                              int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+                              bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+                              int32_t maxBatchSize) NOEXCEPT override;
+
+        void detachFromContext() NOEXCEPT override;
+
+        int i_c, i_h, i_w, o_c, o_h, o_w;
+
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class ResizeLayerRTPluginCreator : public IPluginCreator {
+    public:
+        ResizeLayerRTPluginCreator() ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
 
 
-	virtual size_t getSerializationSize() override {
-		return 6*sizeof(int);
-	}
 
-	virtual void serialize(void* buffer) 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);
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
 
-		tk::dnn::writeBUF(buf, i_c);
-		tk::dnn::writeBUF(buf, i_h);
-		tk::dnn::writeBUF(buf, i_w);
-		assert(buf == a + getSerializationSize());
-	}
+    };
 
-	int i_c, i_h, i_w, o_c, o_h, o_w;
+    REGISTER_TENSORRT_PLUGIN(ResizeLayerRTPluginCreator);
 };
+

include/tkDNN/pluginsRT/RouteRT.h

@@ -1,96 +1,90 @@
 #include<cassert>
 #include "../kernels.h"
+#include <vector>
+#include <NvInfer.h>
 
-class RouteRT : public IPlugin {
+namespace nvinfer1 {
+    class RouteRT : public IPluginV2 {
 
-	/**
-		THIS IS NOT USED ANYMORE
-	*/
+        /**
+            THIS IS NOT USED ANYMORE
+        */
 
-public:
-	RouteRT(int groups, int group_id) {
-		this->groups = groups;
-		this->group_id = group_id;
-	}
+    public:
+        RouteRT(int groups, int group_id) ;
 
-	~RouteRT(){
+        ~RouteRT() ;
 
-	}
+        RouteRT(const void *data, size_t length) ;
 
-	int getNbOutputs() const override {
-		return 1;
-	}
+        int getNbOutputs() const NOEXCEPT override ;
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) 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]};
-	}
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override ;
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		in = nbInputs;
-		c = 0;
-		for(int i=0; i<nbInputs; i++) {
-			c_in[i] = inputDims[i].d[0]; 
-			c += inputDims[i].d[0];
-		}
-		h = inputDims[0].d[1];
-		w = inputDims[0].d[2];
-		c /= groups;
-	}
+        void configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,PluginFormat format, int maxBatchSize) NOEXCEPT override ;
 
-	int initialize() override {
+        int initialize() NOEXCEPT override ;
 
-		return 0;
-	}
+        void terminate() NOEXCEPT override ;
 
-	virtual void terminate() override {
-	}
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override ;
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
-		return 0;
-	}
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT override ;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
-		
-		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+        size_t getSerializationSize() const NOEXCEPT override ;
 
-		for(int b=0; b<batchSize; b++) {
-			int offset = 0;
-			for(int i=0; i<in; i++) {
-				dnnType *input = (dnnType*)reinterpret_cast<const dnnType*>(inputs[i]);
-				int in_dim = c_in[i]*h*w;
-				int part_in_dim = in_dim / this->groups;
-				checkCuda( cudaMemcpyAsync(dstData + b*c*w*h + offset, input + b*c*w*h*groups + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) );
-				offset += part_in_dim;
-			}
-		}
+        void serialize(void *buffer) const NOEXCEPT override ;
 
-		return 0;
-	}
+        const char *getPluginType() const NOEXCEPT override ;
 
+        const char *getPluginVersion() const NOEXCEPT override ;
 
-	virtual size_t getSerializationSize() override {
-		return (6+MAX_INPUTS)*sizeof(int);
-	}
+        void destroy() NOEXCEPT override ;
 
-	virtual void serialize(void* buffer) 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);
-		for(int i=0; i<MAX_INPUTS; i++)
-			tk::dnn::writeBUF(buf, c_in[i]);
+        const char *getPluginNamespace() const NOEXCEPT override ;
 
-		tk::dnn::writeBUF(buf, c);
-		tk::dnn::writeBUF(buf, h);
-		tk::dnn::writeBUF(buf, w);
-		assert(buf == a + getSerializationSize());
-	}
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
 
-	static const int MAX_INPUTS = 4;
-	int in;
-	int c_in[MAX_INPUTS];
-	int c, h, w;
-	int groups, group_id;
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override ;
+
+        IPluginV2 *clone() const NOEXCEPT override ;
+
+        static const int MAX_INPUTS = 4;
+        int in;
+        int c_in[MAX_INPUTS];
+        int c, h, w;
+        int groups, group_id;
+    private:
+        std::string mPluginNamespace;
+    };
+
+    class RouteRTPluginCreator : public IPluginCreator {
+    public:
+        RouteRTPluginCreator() ;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override ;
+
+        const char *getPluginNamespace() const NOEXCEPT override ;
+
+        IPluginV2 *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override ;
+
+        IPluginV2 *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override ;
+
+        const char *getPluginName() const NOEXCEPT override ;
+
+        const char *getPluginVersion() const NOEXCEPT override ;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override ;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(RouteRTPluginCreator);
 };

include/tkDNN/pluginsRT/ShortcutRT.h

@@ -1,75 +1,109 @@
+#ifndef _SHORTCUTRT_PLUGIN_H
+#define _SHORTCUTRT_PLUGIN_H
+
 #include<cassert>
 #include "../kernels.h"
-
-class ShortcutRT : public IPlugin {
-
-public:
-	ShortcutRT(tk::dnn::dataDim_t bdim) {
-		this->bc = bdim.c;
-		this->bh = bdim.h;
-		this->bw = bdim.w;
-	}
-
-	~ShortcutRT(){
-
-	}
-
-	int getNbOutputs() const 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]};
-	}
-
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		c = inputDims[0].d[0];
-		h = inputDims[0].d[1];
-		w = inputDims[0].d[2];
-	}
-
-	int initialize() override {
-
-		return 0;
-	}
-
-	virtual void terminate() override {
-	}
-
-	virtual 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 {
-
-		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]);
-
-		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
-		for(int b=0; b < batchSize; ++b)
-			shortcutForward(srcDataBack + b*bc*bh*bw, dstData + b*c*h*w, 1, c, h, w, 1, 1, bc, bh, bw, 1, stream);
-
-		return 0;
-	}
+#include <NvInfer.h>
+#include <vector>
+#include <tkdnn.h>
 
 
-	virtual size_t getSerializationSize() override {
-		return 6*sizeof(int);
-	}
+namespace nvinfer1 {
 
-	virtual void serialize(void* buffer) 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, c);
-		tk::dnn::writeBUF(buf, h);
-		tk::dnn::writeBUF(buf, w);
-		assert(buf == a + getSerializationSize());
-		
-	}
+    class ShortcutRT : public IPluginV2Ext {
+
+    public:
+        ShortcutRT(int bc,int bh,int bw,int c,int h,int w ,bool mul);
+
+        ~ShortcutRT();
+
+        ShortcutRT(const void *data, size_t length);
+
+        int getNbOutputs() const NOEXCEPT override;
+
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims, int32_t nbOutputs,
+                              DataType const *inputTypes, DataType const *outputTypes, bool const *inputIsBroadcast,
+                              bool const *outputIsBroadcast, PluginFormat floatFormat, int32_t maxBatchSize) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch (int32_t outputIndex, bool const *inputIsBroadcasted, int32_t nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch (int32_t inputIndex) const NOEXCEPT override;
+
+        void attachToContext (cudnnContext *, cublasContext *, IGpuAllocator *) NOEXCEPT override;
+
+        void detachFromContext () NOEXCEPT override;
+
+        DataType getOutputDataType(int32_t index, nvinfer1::DataType const *inputTypes, int32_t nbInputs) const NOEXCEPT override;
+
+        int initialize() NOEXCEPT override;
+
+        void terminate() NOEXCEPT override;
+
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
+
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                    cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
+
+
+        size_t getSerializationSize() const NOEXCEPT override;
+
+        void serialize(void *buffer) const NOEXCEPT override;
+
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
+
+        const char *getPluginType() const NOEXCEPT override;
+
+        const char *getPluginVersion() const NOEXCEPT override;
+
+        void destroy() NOEXCEPT override;
+
+        const char *getPluginNamespace() const NOEXCEPT override;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+        IPluginV2Ext *clone() const NOEXCEPT override;
+
+        int c, h, w;
+        int bc, bh, bw,bl;
+        bool mul;
+        tk::dnn::dataDim_t bDim;
+    private:
+        std::string mPluginNamespace;
+    };
+
+
+    class ShortcutRTPluginCreator : public IPluginCreator {
+    public:
+        ShortcutRTPluginCreator();
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+        const char *getPluginNamespace() const NOEXCEPT override;
+
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override;
+
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override;
+
+        const char *getPluginName() const NOEXCEPT override;
+
+        const char *getPluginVersion() const NOEXCEPT override;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override;
+
+    public:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(ShortcutRTPluginCreator);
 
-	int c, h, w;
-	int bc, bh, bw;
 };
+
+#endif

include/tkDNN/pluginsRT/UpsampleRT.h

@@ -1,66 +1,103 @@
+#ifndef _UPSAMPLERT_PLUGIN_H
+#define _UPSAMPLERT_PLUGIN_H
+
 #include<cassert>
 #include "../kernels.h"
+#include <NvInfer.h>
+#include <vector>
 
-class UpsampleRT : public IPlugin {
+namespace nvinfer1 {
 
-public:
-	UpsampleRT(int stride) {
-		this->stride = stride;
-	}
+    class UpsampleRT : public IPluginV2Ext {
 
-	~UpsampleRT(){
+    public:
+        UpsampleRT(int stride,int c,int h,int w);
 
-	}
+        UpsampleRT(const void *data, size_t length);
 
-	int getNbOutputs() const override {
-		return 1;
-	}
+        ~UpsampleRT();
 
-	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);
-	}
+        int getNbOutputs() const NOEXCEPT override;
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		c = inputDims[0].d[0];
-		h = inputDims[0].d[1];
-		w = inputDims[0].d[2];
-	}
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
 
-	int initialize() override {
+        int initialize() NOEXCEPT override;
 
-		return 0;
-	}
+        void terminate() NOEXCEPT override;
 
-	virtual void terminate() override {
-	}
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
 
-	virtual 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 {
-
-		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;
-	}
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                            cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
 
-	virtual size_t getSerializationSize() override {
-		return 4*sizeof(int);
-	}
+        size_t getSerializationSize() const NOEXCEPT override;
 
-	virtual void serialize(void* buffer) 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);
-		assert(buf == a + getSerializationSize());
-	}
+        void serialize(void *buffer) const NOEXCEPT override;
 
-	int c, h, w, stride;
-};
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
+
+        const char *getPluginType() const NOEXCEPT override;
+
+        const char *getPluginVersion() const NOEXCEPT override;
+
+        void destroy() NOEXCEPT override;
+
+        const char *getPluginNamespace() const NOEXCEPT override;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+        IPluginV2Ext *clone() const NOEXCEPT override ;
+
+        bool isOutputBroadcastAcrossBatch (int32_t outputIndex, bool const *inputIsBroadcasted, int32_t nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch (int32_t inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims, int32_t nbOutputs,
+                              DataType const *inputTypes, DataType const *outputTypes, bool const *inputIsBroadcast,
+                              bool const *outputIsBroadcast, PluginFormat floatFormat, int32_t maxBatchSize) NOEXCEPT override;
+
+        void attachToContext (cudnnContext *, cublasContext *, IGpuAllocator *) NOEXCEPT override;
+
+        void detachFromContext () NOEXCEPT override;
+
+        DataType getOutputDataType (int32_t index, nvinfer1::DataType const *inputTypes, int32_t nbInputs) const NOEXCEPT override;
+
+
+        int c, h, w, stride;
+    private:
+        std::string mPluginNamespace;
+    };
+
+        class UpsampleRTPluginCreator : public IPluginCreator {
+        public:
+            UpsampleRTPluginCreator();
+
+            void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+            const char *getPluginNamespace() const NOEXCEPT override;
+
+            IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override;
+
+            IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override;
+
+            const char *getPluginName() const NOEXCEPT override;
+
+            const char *getPluginVersion() const NOEXCEPT override;
+
+            const PluginFieldCollection *getFieldNames() NOEXCEPT override;
+
+        private:
+            static PluginFieldCollection mFC;
+            static std::vector<PluginField> mPluginAttributes;
+            std::string mPluginNamespace;
+        };
+
+        REGISTER_TENSORRT_PLUGIN(UpsampleRTPluginCreator);
+    };
+
+#endif

include/tkDNN/pluginsRT/YoloRT.h

@@ -1,143 +1,124 @@
+#ifndef _YOLORT_PLUGIN_H
+#define _YOLORT_PLUGIN_H
+
 #include<cassert>
+#include <vector>
 #include "../kernels.h"
+#include <NvInfer.h>
+#include <tkdnn.h>
 
 #define YOLORT_CLASSNAME_W 256
 
-class YoloRT : public IPlugin {
+namespace nvinfer1 {
+    class YoloRT : public IPluginV2Ext {
+
+    public:
+        YoloRT(int classes, int num,int c,int h,int w, int n_masks = 3, float scale_xy = 1,
+               float nms_thresh = 0.45, int nms_kind = 0, int new_coords = 0);
+
+        YoloRT(const void *data, size_t length);
+
+        ~YoloRT();
 
 
+        int getNbOutputs() const NOEXCEPT override;
 
-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) {
+        Dims getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT override;
 
-		this->classes = classes;
-		this->num = num;
-		this->n_masks = n_masks;
-		this->scaleXY = scale_xy;
-		this->nms_thresh = nms_thresh;
-		this->nms_kind = nms_kind;
-		this->new_coords = new_coords;
+        int initialize() NOEXCEPT override;
 
-        mask = new dnnType[n_masks];
-        bias = new dnnType[num*n_masks*2];
-        if(yolo != nullptr) {
-            memcpy(mask, yolo->mask_h, sizeof(dnnType)*n_masks);
-            memcpy(bias, yolo->bias_h, sizeof(dnnType)*num*n_masks*2);
-			classesNames = yolo->classesNames;
-        }
-	}
+        void terminate() NOEXCEPT override;
 
-	~YoloRT(){
-
-	}
-
-	int getNbOutputs() const override {
-		return 1;
-	}
-
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
-		return inputs[0];
-	}
-
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
-		c = inputDims[0].d[0];
-		h = inputDims[0].d[1];
-		w = inputDims[0].d[2];
-	}
-
-	int initialize() override {
-
-		return 0;
-	}
-
-	virtual void terminate() override {
-	}
-
-	virtual 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 {
-
-		dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
-		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
-
-		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
+        size_t getWorkspaceSize(int maxBatchSize) const NOEXCEPT override;
 
 
-        for (int b = 0; b < batchSize; ++b){
-            for(int n = 0; n < n_masks; ++n){
-                int index = entry_index(b, n*w*h, 0);
-                if (new_coords == 1){
-                    if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
-                }
-                else{
-                    activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); //x,y
+#if NV_TENSORRT_MAJOR > 7
+        int enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                            cudaStream_t stream) NOEXCEPT override;
+#elif NV_TENSORRT_MAJOR == 7
+        int32_t enqueue (int32_t batchSize, const void *const *inputs, void **outputs, void *workspace, cudaStream_t stream) override;
+#endif
 
-                    if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
 
-                    index = entry_index(b, n*w*h, 4);
-                    activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
-                }
-            }
+        size_t getSerializationSize() const NOEXCEPT override;
+
+        bool supportsFormat(DataType type, PluginFormat format) const NOEXCEPT override;
+
+        void serialize(void *buffer) const NOEXCEPT override;
+
+        const char *getPluginType() const NOEXCEPT override;
+
+        const char *getPluginVersion() const NOEXCEPT override;
+
+        void destroy() NOEXCEPT override;
+
+        const char *getPluginNamespace() const NOEXCEPT override;
+
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
+
+        IPluginV2Ext *clone() const NOEXCEPT override;
+
+        DataType getOutputDataType(int index, const nvinfer1::DataType* inputTypes, int nbInputs) const NOEXCEPT override;
+
+        void attachToContext(cudnnContext* cudnnContext, cublasContext* cublasContext, IGpuAllocator* gpuAllocator) NOEXCEPT override;
+
+        bool isOutputBroadcastAcrossBatch(int outputIndex, const bool* inputIsBroadcasted, int nbInputs) const NOEXCEPT override;
+
+        bool canBroadcastInputAcrossBatch(int inputIndex) const NOEXCEPT override;
+
+        void configurePlugin (Dims const *inputDims, int32_t nbInputs, Dims const *outputDims,
+                              int32_t nbOutputs, DataType const *inputTypes, DataType const *outputTypes,
+                              bool const *inputIsBroadcast, bool const *outputIsBroadcast, PluginFormat floatFormat,
+                              int32_t maxBatchSize) NOEXCEPT override;
+
+        void detachFromContext() NOEXCEPT override;
+
+
+        int c, h, w;
+        int classes, num, n_masks;
+        float scaleXY;
+        float nms_thresh;
+        int nms_kind;
+        int new_coords;
+        int NUM = 0;
+        std::vector<std::string> classesNames;
+
+
+        int entry_index(int batch, int location, int entry) {
+            int n = location / (w * h);
+            int loc = location % (w * h);
+            return batch * c * h * w + n * w * h * (4 + classes + 1) + entry * w * h + loc;
         }
 
-		//std::cout<<"YOLO END\n";
-		return 0;
-	}
+    private:
+        std::string mPluginNamespace;
 
+    };
 
-	virtual size_t getSerializationSize() override {
-		return 8*sizeof(int) + 2*sizeof(float)+ n_masks*sizeof(dnnType) + num*n_masks*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char);
-	}
+    class YoloRTPluginCreator : public IPluginCreator {
+    public:
+        YoloRTPluginCreator();
 
-	virtual void serialize(void* buffer) 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;
-		for (int i = 0; i < n_masks; i++)
-		{
-			tk::dnn::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;
-		}
+        void setPluginNamespace(const char *pluginNamespace) NOEXCEPT override;
 
-		// save classes names
-		for(int i=0; i<classes; i++) {
-			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]);
-			}
-		}
-		assert(buf == a + getSerializationSize());
-	}
+        const char *getPluginNamespace() const NOEXCEPT override;
 
-	int c, h, w;
-    int classes, num, n_masks;
-	float scaleXY;
-	float nms_thresh;
-	int nms_kind;
-	int new_coords;
-	std::vector<std::string> classesNames;
+        IPluginV2Ext *deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT override;
 
-    dnnType *mask;
-    dnnType *bias;
+        IPluginV2Ext *createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT override;
 
-	int entry_index(int batch, int location, int entry) {
-		int n =   location / (w*h);
-		int loc = location % (w*h);
-		return batch*c*h*w + n*w*h*(4+classes+1) + entry*w*h + loc;
-	}
+        const char *getPluginName() const NOEXCEPT override;
 
+        const char *getPluginVersion() const NOEXCEPT override;
+
+        const PluginFieldCollection *getFieldNames() NOEXCEPT override;
+
+    private:
+        static PluginFieldCollection mFC;
+        static std::vector<PluginField> mPluginAttributes;
+        std::string mPluginNamespace;
+    };
+
+    REGISTER_TENSORRT_PLUGIN(YoloRTPluginCreator);
 };
+#endif