CMakeLists.txt

@@ -18,7 +18,7 @@ if(DEBUG)
 endif()
 
 if(TKDNN_PATH)
-    message("SET TKDNN_PATH:"${TKDNN_PATH})
+    message("SET TKDNN_PATH:${TKDNN_PATH}")
     add_definitions(-DTKDNN_PATH="${TKDNN_PATH}")
 else()
     add_definitions(-DTKDNN_PATH="${CMAKE_CURRENT_SOURCE_DIR}")
@@ -51,11 +51,11 @@ find_package(Eigen3 REQUIRED)
 message("Eigen DIR: " ${EIGEN3_INCLUDE_DIR})
 include_directories(${EIGEN3_INCLUDE_DIR})
 
-find_package(OpenCV REQUIRED)
+find_package(OpenCV 4.5 REQUIRED)
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV")
-# if(OpenCV_CUDA_VERSION)
+if(OpenCV_CUDA_VERSION)
-#     add_compile_definitions(OPENCV_CUDACONTRIB)
+    add_compile_definitions(OPENCV_CUDACONTRIB)
-# endif()
+endif()
 
 # gives problems in cross-compiling, probably malformed cmake config
 find_package(yaml-cpp REQUIRED)
@@ -71,9 +71,9 @@ include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${O
 add_library(tkDNN SHARED ${tkdnn_SRC})
 target_link_libraries(tkDNN ${tkdnn_LIBS})
 
-#static 
+#static
-#add_library(tkDNN_static STATIC ${tkdnn_SRC})
+add_library(tkDNN_static STATIC ${tkdnn_SRC})
-#target_link_libraries(tkDNN_static ${tkdnn_LIBS})
+target_link_libraries(tkDNN_static ${tkdnn_LIBS})
 
 # SMALL NETS
 add_executable(test_simple tests/simple/test_simple.cpp)
@@ -163,7 +163,7 @@ target_link_libraries(seg_demo tkDNN)
 # Install
 #-------------------------------------------------------------------------------
 #if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
-#    set (CMAKE_INSTALL_PREFIX "${CMAKE_BINARY_DIR}/install" 
+#    set (CMAKE_INSTALL_PREFIX "${CMAKE_BINARY_DIR}/install"
 #         CACHE PATH "default install path" FORCE)
 #endif()
 message("install dir:" ${CMAKE_INSTALL_PREFIX})

include/tkDNN/DetectionNN.h

@@ -3,10 +3,10 @@
 
 #include <iostream>
 #include <signal.h>
-#include <stdlib.h>    
+#include <stdlib.h>
 #ifdef __linux__
 #include <unistd.h>
-#endif 
+#endif
 
 #include <mutex>
 #include "utils.h"
@@ -17,7 +17,7 @@
 
 #include "tkdnn.h"
 
-//#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
+#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
 
 #ifdef OPENCV_CUDACONTRIB
 #include <opencv2/cudawarping.hpp>
@@ -37,7 +37,7 @@ class DetectionNN {
 
         cv::Scalar colors[256];
 
-        int nBatches = 1;
+        int nBatches = 2;
 
 #ifdef OPENCV_CUDACONTRIB
         cv::cuda::GpuMat bgr[3];
@@ -57,11 +57,11 @@ class DetectionNN {
         virtual void preprocess(cv::Mat &frame, const int bi=0) = 0;
 
         /**
-         * This method postprocess the output of the NN to obtain the correct 
+         * This method postprocess the output of the NN to obtain the correct
-         * boundig boxes. 
+         * boundig boxes.
-         * 
+         *
          * @param bi batch index
-         * @param mAP set to true only if all the probabilities for a bounding 
+         * @param mAP set to true only if all the probabilities for a bounding
          *            box are needed, as in some cases for the mAP calculation
          */
         virtual void postprocess(const int bi=0,const bool mAP=false) = 0;
@@ -79,25 +79,25 @@ class DetectionNN {
         ~DetectionNN(){};
 
         /**
-         * Method used to initialize the class, allocate memory and compute 
+         * Method used to initialize the class, allocate memory and compute
          * needed data.
-         * 
+         *
          * @param tensor_path path to the rt file of the NN.
          * @param n_classes number of classes for the given dataset.
          * @param n_batches maximum number of batches to use in inference
          * @return true if everything is correct, false otherwise.
          */
         virtual bool init(const std::string& tensor_path, const int n_classes=80, const int n_batches=1, const float conf_thresh=0.3) = 0;
-        
+
         /**
          * This method performs the whole detection of the NN.
-         * 
+         *
          * @param frames frames to run detection on.
          * @param cur_batches number of batches to use in inference
-         * @param save_times if set to true, preprocess, inference and postprocess times 
+         * @param save_times if set to true, preprocess, inference and postprocess times
          *        are saved on a csv file, otherwise not.
          * @param times pointer to the output stream where to write times
-         * @param mAP set to true only if all the probabilities for a bounding 
+         * @param mAP set to true only if all the probabilities for a bounding
          *            box are needed, as in some cases for the mAP calculation
          */
         void update(std::vector<cv::Mat>& frames, const int cur_batches=1, bool save_times=false, std::ofstream *times=nullptr, const bool mAP=false){
@@ -107,14 +107,14 @@ class DetectionNN {
                 FatalError("A batch size greater than nBatches cannot be used");
 
             originalSize.clear();
-            if(TKDNN_VERBOSE) printCenteredTitle(" TENSORRT detection ", '=', 30); 
+            if(TKDNN_VERBOSE) printCenteredTitle(" TENSORRT detection ", '=', 30);
             {
                 TKDNN_TSTART
                 for(int bi=0; bi<cur_batches;++bi){
                     if(!frames[bi].data)
                         FatalError("No image data feed to detection");
                     originalSize.push_back(frames[bi].size());
-                    preprocess(frames[bi], bi);    
+                    preprocess(frames[bi], bi);
                 }
                 TKDNN_TSTOP
                 if(save_times) *times<<t_ns<<";";
@@ -141,11 +141,11 @@ class DetectionNN {
                 TKDNN_TSTOP
                 if(save_times) *times<<t_ns<<"\n";
             }
-        }      
+        }
 
         /**
          * Method to draw bounding boxes and labels on a frame.
-         * 
+         *
          * @param frames original frame to draw bounding box on.
          */
         void draw(std::vector<cv::Mat>& frames) {
@@ -155,11 +155,11 @@ class DetectionNN {
             std::string det_class;
             int baseline = 0;
             float font_scale = 0.5;
-            int thickness = 2;   
+            int thickness = 2;
 
             for(int bi=0; bi<frames.size(); ++bi){
                 // draw dets
-                for(int i=0; i<batchDetected[bi].size(); i++) { 
+                for(int i=0; i<batchDetected[bi].size(); i++) {
                     b           = batchDetected[bi][i];
                     x0   		= b.x;
                     x1   		= b.x + b.w;
@@ -168,11 +168,11 @@ class DetectionNN {
                     det_class 	= classesNames[b.cl];
 
                     // draw rectangle
-                    cv::rectangle(frames[bi], cv::Point(x0, y0), cv::Point(x1, y1), colors[b.cl], 2); 
+                    cv::rectangle(frames[bi], cv::Point(x0, y0), cv::Point(x1, y1), colors[b.cl], 2);
 
                     // draw label
                     cv::Size text_size = getTextSize(det_class, cv::FONT_HERSHEY_SIMPLEX, font_scale, thickness, &baseline);
-                    cv::rectangle(frames[bi], cv::Point(x0, y0), cv::Point((x0 + text_size.width - 2), (y0 - text_size.height - 2)), colors[b.cl], -1);                      
+                    cv::rectangle(frames[bi], cv::Point(x0, y0), cv::Point((x0 + text_size.width - 2), (y0 - text_size.height - 2)), colors[b.cl], -1);
                     cv::putText(frames[bi], det_class, cv::Point(x0, (y0 - (baseline / 2))), cv::FONT_HERSHEY_SIMPLEX, font_scale, cv::Scalar(255, 255, 255), thickness);
                 }
             }

include/tkDNN/DetectionNN3D.h

@@ -4,7 +4,7 @@
 #include <iostream>
 #include <signal.h>
 #include <stdlib.h>
-#ifdef __linux__    
+#ifdef __linux__
 #include <unistd.h>
 #endif
 
@@ -17,7 +17,7 @@
 
 #include "tkdnn.h"
 
-// #define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
+#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
 
 #ifdef OPENCV_CUDACONTRIB
 #include <opencv2/cudawarping.hpp>
@@ -57,11 +57,11 @@ class DetectionNN3D {
         virtual void preprocess(cv::Mat &frame, const int bi=0) = 0;
 
         /**
-         * This method postprocess the output of the NN to obtain the correct 
+         * This method postprocess the output of the NN to obtain the correct
-         * boundig boxes. 
+         * boundig boxes.
-         * 
+         *
          * @param bi batch index
-         * @param mAP set to true only if all the probabilities for a bounding 
+         * @param mAP set to true only if all the probabilities for a bounding
          *            box are needed, as in some cases for the mAP calculation
          */
         virtual void postprocess(const int bi=0,const bool mAP=false) = 0;
@@ -69,7 +69,7 @@ class DetectionNN3D {
     public:
         int classes = 0;
         float confThreshold = 0.3; /*threshold on the confidence of the boxes*/
-        
+
         std::vector<tk::dnn::box3D> detected3D; /*bounding boxes in output*/
         std::vector<std::vector<tk::dnn::box3D>> batchDetected; /*bounding boxes in output*/
         std::vector<double> pre_stats, stats, post_stats, visual_stats; /*keeps track of inference times (ms)*/
@@ -79,29 +79,29 @@ class DetectionNN3D {
         ~DetectionNN3D(){};
 
         /**
-         * Method used to initialize the class, allocate memory and compute 
+         * Method used to initialize the class, allocate memory and compute
          * needed data.
-         * 
+         *
          * @param tensor_path path to the rt file of the NN.
          * @param n_classes number of classes for the given dataset.
          * @param n_batches maximum number of batches to use in inference.
          * @return true if everything is correct, false otherwise.
          */
-        virtual bool init(const std::string& tensor_path, const int n_classes=3, const int n_batches=1, 
+        virtual bool init(const std::string& tensor_path, const int n_classes=3, const int n_batches=1,
                             const float conf_thresh=0.3, const std::vector<cv::Mat>& k_calibs=std::vector<cv::Mat>()) = 0;
 
         /**
          * This method performs the whole detection of the NN.
-         * 
+         *
          * @param frames frames to run detection on.
          * @param cur_batches number of batches to use in inference.
-         * @param save_times if set to true, preprocess, inference and postprocess times 
+         * @param save_times if set to true, preprocess, inference and postprocess times
          *        are saved on a csv file, otherwise not.
          * @param times pointer to the output stream where to write times.
-         * @param mAP set to true only if all the probabilities for a bounding 
+         * @param mAP set to true only if all the probabilities for a bounding
          *            box are needed, as in some cases for the mAP calculation.
          */
-        void update(std::vector<cv::Mat>& frames, const int cur_batches=1, bool save_times=false, 
+        void update(std::vector<cv::Mat>& frames, const int cur_batches=1, bool save_times=false,
                     std::ofstream *times=nullptr, const bool mAP=false){
             if(save_times && times==nullptr)
                 FatalError("save_times set to true, but no valid ofstream given");
@@ -109,17 +109,17 @@ class DetectionNN3D {
                 FatalError("A batch size greater than nBatches cannot be used");
 
             originalSize.clear();
-            if(TKDNN_VERBOSE) printCenteredTitle(" TENSORRT detection ", '=', 30); 
+            if(TKDNN_VERBOSE) printCenteredTitle(" TENSORRT detection ", '=', 30);
             {
                 TKDNN_TSTART
                 for(int bi=0; bi<cur_batches;++bi){
                     if(!frames[bi].data)
                         FatalError("No image data feed to detection");
                     originalSize.push_back(frames[bi].size());
-                    preprocess(frames[bi], bi);   
+                    preprocess(frames[bi], bi);
                 }
                 TKDNN_TSTOP
-                pre_stats.push_back(t_ns); 
+                pre_stats.push_back(t_ns);
                 if(save_times) *times<<t_ns<<";";
             }
 
@@ -149,11 +149,11 @@ class DetectionNN3D {
 
         /**
          * Method to draw bounding boxes and labels on a frame.
-         * 
+         *
          * @param frames original frame to draw bounding box on.
-         */    
+         */
         virtual void draw(std::vector<cv::Mat>& frames){};
-          
+
 };
 
 }}

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. 
+ * BatchStream implements the stream for the INT8 calibrator.
- * It reads the two files .txt with the list of image file names 
+ * It reads the two files .txt with the list of image file names
- * and the list of label 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 
+ * INT8 quantization. It uses a BatchStream stream to scroll through
- * images data. It also implements the calibration cache, a way to 
+ * images data. It also implements the calibration cache, a way to
- * save the calibration process results to reduce the running time: 
+ * 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(); }
+	int getBatchSize() const noexcept override { return mStream.getBatchSize(); }
-	bool getBatch(void* bindings[], const char* names[], int nbBindings) override;
+	bool getBatch(void* bindings[], const char* names[], int nbBindings) noexcept override;
-	const void* readCalibrationCache(size_t& length) override;
+	const void* readCalibrationCache(size_t& length) noexcept override;
-	void writeCalibrationCache(const void* cache, size_t length) 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 {
+// using namespace nvinfer1;
-
-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;
 #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
+namespace tk { namespace dnn {
-{
-public:
-    YoloRT *yolos[16];
-    int n_yolos;
-
-	virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength);
-};
-
-
 
 class NetworkRT {
 
@@ -57,11 +34,11 @@ public:
     nvinfer1::DataType dtRT;
     nvinfer1::IBuilder *builderRT;
     nvinfer1::IRuntime *runtimeRT;
-    nvinfer1::INetworkDefinition *networkRT; 
+    nvinfer1::INetworkDefinition *networkRT;
-#if NV_TENSORRT_MAJOR >= 6  
+#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>
+#ifndef ACTIVATION_LEAKY_RT_H
-#include "../kernels.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>
+#ifndef ACTIVATION_LOGISTIC_RT_H
-#include "../kernels.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;
 
-    }
+    int getNbOutputs() const noexcept override {
-
-    ~ActivationLogisticRT(){
-
-    }
-
-    int getNbOutputs() const 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;
     }
 
-
+    size_t getSerializationSize() const noexcept override {
-    virtual size_t getSerializationSize() override {
         return 1*sizeof(int);
     }
 
-    virtual void serialize(void* buffer) override {
+    void serialize(void* buffer) const noexcept override {
         char *buf = reinterpret_cast<char*>(buffer);
-        tk::dnn::writeBUF(buf, size);
+        writeBUF(buf, size);
     }
 
+    // Extra IPluginV2 overrides
+    bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+        return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+    }
+
+    nvinfer1::IPluginV2 * clone() const noexcept override {
+        auto a = new ActivationLogisticRT(*this);
+        return a;
+    }
+
+    const char* getPluginType() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    void destroy() noexcept override {}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+        mNamespace = pluginNamespace;
+    }
+
+    const char* getPluginNamespace() const noexcept override {
+        return mNamespace.c_str();
+    }
+
+    std::string mNamespace;
     int size;
 };
+
+class ActivationLogisticRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationLogisticRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+        const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+        ActivationLogisticRT *a = new ActivationLogisticRT();
+        a->size = readBUF<int>(buf);
+        assert(buf == bufCheck + serialLength);
+        return a;
+    }
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+        mNamespace = pluginNamespace;
+    }
+
+    const char* getPluginNamespace() const noexcept override {
+        return mNamespace.c_str();
+    }
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // ACTIVATION_LOGISTIC_RT_H

include/tkDNN/pluginsRT/ActivationMishRT.h

@@ -1,61 +1,150 @@
-#include<cassert>
+#ifndef ACTIVATION_MISH_RT_H
-#include "../kernels.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;
 
-	}
+	int getNbOutputs() const noexcept override {
-
-	~ActivationMishRT(){
-
-	}
-
-	int getNbOutputs() const override {
 		return 1;
 	}
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
 		return inputs[0];
 	}
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+	void configureWithFormat(nvinfer1::Dims const * inputDims,
+					int32_t nbInputs,
+					nvinfer1::Dims const * outputDims,
+					int32_t nbOutputs,
+					nvinfer1::DataType type,
+					nvinfer1::PluginFormat format,
+					int32_t maxBatchSize) noexcept override {
 		size = 1;
 		for(int i=0; i<outputDims[0].nbDims; i++)
 			size *= outputDims[0].d[i];
 	}
 
-	int initialize() override {
+	int initialize() noexcept override {
 
 		return 0;
 	}
 
-	virtual void terminate() override {
+	void terminate() noexcept override {
 	}
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+	size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
 		return 0;
 	}
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+	int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
-
+		activationMishForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
-		activationMishForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
 											reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
 		return 0;
 	}
 
-
+	size_t getSerializationSize() const noexcept override {
-	virtual size_t getSerializationSize() override {
 		return 1*sizeof(int);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) const noexcept override {
 		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
-		tk::dnn::writeBUF(buf, size);
+		writeBUF(buf, size);
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ActivationMishRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int size;
 };
+
+class ActivationMishRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationMishRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+        const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+        ActivationMishRT *a = new ActivationMishRT();
+        a->size = readBUF<int>(buf);
+        assert(buf == bufCheck + serialLength);
+        return a;
+    }
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+        mNamespace = pluginNamespace;
+    }
+
+    const char* getPluginNamespace() const noexcept override {
+        return mNamespace.c_str();
+    }
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // ACTIVATION_MISH_RT_H

include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h

@@ -1,63 +1,156 @@
-#include<cassert>
+#ifndef ACTIVATION_RELU_CEILING_RT_H
-#include "../kernels.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 noexcept override {
-
-	int getNbOutputs() const override {
 		return 1;
 	}
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
 		return inputs[0];
 	}
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+	void configureWithFormat(nvinfer1::Dims const * inputDims,
+					int32_t nbInputs,
+					nvinfer1::Dims const * outputDims,
+					int32_t nbOutputs,
+					nvinfer1::DataType type,
+					nvinfer1::PluginFormat format,
+					int32_t maxBatchSize) noexcept override {
 		size = 1;
 		for(int i=0; i<outputDims[0].nbDims; i++)
 			size *= outputDims[0].d[i];
 	}
 
-	int initialize() override {
+	int initialize() noexcept override {
 
 		return 0;
 	}
 
-	virtual void terminate() override {
+	void terminate() noexcept override {
 	}
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+	size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
 		return 0;
 	}
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+	int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
-
+		activationReLUCeilingForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
-		activationReLUCeilingForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
 											reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, ceiling, stream);
 		return 0;
 	}
 
-
+	size_t getSerializationSize() const noexcept override {
-	virtual size_t getSerializationSize() 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);
+		writeBUF(buf, ceiling);
-		tk::dnn::writeBUF(buf, size);
+		writeBUF(buf, size);
 		assert(buf = a + getSerializationSize());
-		
+
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new 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>
+#ifndef ACTIVATION_SIGMOID_RT_H
-#include "../kernels.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;
 
-	}
+	int getNbOutputs() const noexcept override {
-
-	~ActivationSigmoidRT(){
-
-	}
-
-	int getNbOutputs() const override {
 		return 1;
 	}
 
-	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
 		return inputs[0];
 	}
 
-	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+	void configureWithFormat(nvinfer1::Dims const * inputDims,
+					int32_t nbInputs,
+					nvinfer1::Dims const * outputDims,
+					int32_t nbOutputs,
+					nvinfer1::DataType type,
+					nvinfer1::PluginFormat format,
+					int32_t maxBatchSize) noexcept override {
 		size = 1;
 		for(int i=0; i<outputDims[0].nbDims; i++)
 			size *= outputDims[0].d[i];
 	}
 
-	int initialize() override {
+	int initialize() noexcept override {
-
 		return 0;
 	}
 
-	virtual void terminate() override {
+	void terminate() noexcept override {
 	}
 
-	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+	size_t getWorkspaceSize(int maxBatchSize) const noexcept override {
 		return 0;
 	}
 
-	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+	int32_t enqueue(int32_t batchSize, const void* const * inputs, void* const* outputs, void* workspace, cudaStream_t stream) noexcept override {
-
+		activationSIGMOIDForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
-		activationSIGMOIDForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
 											reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
 		return 0;
 	}
 
-
+	size_t getSerializationSize() const noexcept override {
-	virtual size_t getSerializationSize() override {
 		return 1*sizeof(int);
 	}
 
-	virtual void serialize(void* buffer) override {
+	void serialize(void* buffer) const noexcept override {
 		char *buf = reinterpret_cast<char*>(buffer),*a=buf;
-		tk::dnn::writeBUF(buf, size);
+		writeBUF(buf, size);
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ActivationSigmoidRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int size;
 };
+
+class ActivationSigmoidRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ActivationSigmoidRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			ActivationSigmoidRT* a = new ActivationSigmoidRT();
+			a->size = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return a;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // ACTIVATION_SIGMOID_RT_H

include/tkDNN/pluginsRT/DeformableConvRT.h

@@ -1,16 +1,27 @@
-#include<cassert>
+#ifndef DEFORMABLE_CONV_RT_H
+#define DEFORMABLE_CONV_RT_H
+
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
 #include "../kernels.h"
+#include "../buffer_func.h"
+#include "../Layer.h"
 
+#define PLUGIN_NAME "Deformable"
+#define PLUGIN_VERSION "1"
+namespace tk { namespace dnn {
 
-class DeformableConvRT : public IPlugin {
+class DeformableConvRT final : public nvinfer1::IPluginV2 {
-
-
 
 public:
-	DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw, 
+	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 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, 
+					int o_n, int o_c, int o_h, int o_w,
-					tk::dnn::DeformConv2d *deformable = nullptr) {
+					DeformConv2d *deformable = nullptr) {
 		this->chunk_dim = chunk_dim;
 		this->kh = kh;
 		this->kw = kw;
@@ -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 {
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
-		return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w};
+		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(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(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;
 	}
 
-
+	size_t getSerializationSize() const noexcept override {
-	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);
 	}
 
-	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);
+		writeBUF(buf, chunk_dim);
-		tk::dnn::writeBUF(buf, kh);
+		writeBUF(buf, kh);
-		tk::dnn::writeBUF(buf, kw);
+		writeBUF(buf, kw);
-		tk::dnn::writeBUF(buf, sh);
+		writeBUF(buf, sh);
-		tk::dnn::writeBUF(buf, sw);
+		writeBUF(buf, sw);
-		tk::dnn::writeBUF(buf, ph);
+		writeBUF(buf, ph);
-		tk::dnn::writeBUF(buf, pw);
+		writeBUF(buf, pw);
-		tk::dnn::writeBUF(buf, deformableGroup);
+		writeBUF(buf, deformableGroup);
-		tk::dnn::writeBUF(buf, i_n);
+		writeBUF(buf, i_n);
-		tk::dnn::writeBUF(buf, i_c);
+		writeBUF(buf, i_c);
-		tk::dnn::writeBUF(buf, i_h);
+		writeBUF(buf, i_h);
-		tk::dnn::writeBUF(buf, i_w);
+		writeBUF(buf, i_w);
-		tk::dnn::writeBUF(buf, o_n);
+		writeBUF(buf, o_n);
-		tk::dnn::writeBUF(buf, o_c);
+		writeBUF(buf, o_c);
-		tk::dnn::writeBUF(buf, o_h);
+		writeBUF(buf, o_h);
-		tk::dnn::writeBUF(buf, o_w);
+		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; 
+	// Extra IPluginV2 overrides
-	cublasHandle_t handle; 
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new DeformableConvRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
+	cublasStatus_t stat;
+	cublasHandle_t handle;
 	int i_n, i_c, i_h, i_w;
 	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;
-	
+
-	
+	DeformConv2d *defRT;
-	tk::dnn::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 noexcept override {
-
-	int getNbOutputs() const 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 DimsCHW{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1};
+		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;
 	}
 
-
+	size_t getSerializationSize() const noexcept override {
-	virtual size_t getSerializationSize() 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);
+		writeBUF(buf, c);
-		tk::dnn::writeBUF(buf, h);
+		writeBUF(buf, h);
-		tk::dnn::writeBUF(buf, w);
+		writeBUF(buf, w);
-		tk::dnn::writeBUF(buf, rows);
+		writeBUF(buf, rows);
-		tk::dnn::writeBUF(buf, cols);
+		writeBUF(buf, cols);
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new FlattenConcatRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
 	int rows, cols;
-	cublasStatus_t stat; 
+	cublasStatus_t stat;
-	cublasHandle_t handle; 
+	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>
+#ifndef MAX_POOLING_FIXED_SIZE_RT_H
-#include "../kernels.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 {
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
-		return DimsCHW{this->c, this->h, this->w};
+		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;
 	}
 
-
+	size_t getSerializationSize() const noexcept override {
-	virtual size_t getSerializationSize() 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);
+		writeBUF(buf, this->c);
-		tk::dnn::writeBUF(buf, this->h);
+		writeBUF(buf, this->h);
-		tk::dnn::writeBUF(buf, this->w);
+		writeBUF(buf, this->w);
-		tk::dnn::writeBUF(buf, this->n);
+		writeBUF(buf, this->n);
-		tk::dnn::writeBUF(buf, this->stride_H);
+		writeBUF(buf, this->stride_H);
-		tk::dnn::writeBUF(buf, this->stride_W);
+		writeBUF(buf, this->stride_W);
-		tk::dnn::writeBUF(buf, this->winSize);
+		writeBUF(buf, this->winSize);
-		tk::dnn::writeBUF(buf, this->padding);
+		writeBUF(buf, this->padding);
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new MaxPoolFixedSizeRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int n, c, h, w;
 	int stride_H, stride_W;
 	int winSize;
 	int padding;
 };
+
+class MaxPoolFixedSizeRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    MaxPoolFixedSizeRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int cTemp = readBUF<int>(buf);
+			int hTemp = readBUF<int>(buf);
+			int wTemp = readBUF<int>(buf);
+			int nTemp = readBUF<int>(buf);
+			int strideHTemp = readBUF<int>(buf);
+			int strideWTemp = readBUF<int>(buf);
+			int winSizeTemp = readBUF<int>(buf);
+			int paddingTemp = readBUF<int>(buf);
+
+			MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // MAX_POOLING_FIXED_SIZE_RT_H

include/tkDNN/pluginsRT/RegionRT.h

@@ -1,48 +1,62 @@
-#include<cassert>
+#ifndef REGION_RT_H
-#include "../kernels.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 noexcept override {
-
-	int getNbOutputs() const 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, 
+		softmaxForward(	srcData + index, classes, batchSize*num,
-						(c*h*w)/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);
+		writeBUF(buf, classes);
-		tk::dnn::writeBUF(buf, coords);
+		writeBUF(buf, coords);
-		tk::dnn::writeBUF(buf, num);
+		writeBUF(buf, num);
-		tk::dnn::writeBUF(buf, c);
+		writeBUF(buf, c);
-		tk::dnn::writeBUF(buf, h);
+		writeBUF(buf, h);
-		tk::dnn::writeBUF(buf, w);
+		writeBUF(buf, w);
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new RegionRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
-    int classes, coords, num;
+  int classes, coords, num;
 
 	int entry_index(int batch, int location, int entry) {
 		int n =   location / (w*h);
@@ -93,3 +137,57 @@ public:
 	}
 
 };
+
+class RegionRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    RegionRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int classesTemp = readBUF<int>(buf);
+			int coordsTemp = readBUF<int>(buf);
+			int numTemp = readBUF<int>(buf);
+			RegionRT* r = new RegionRT(classesTemp, coordsTemp, numTemp);
+
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // REGION_RT_H

include/tkDNN/pluginsRT/ReorgRT.h

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

include/tkDNN/pluginsRT/ReshapeRT.h

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

include/tkDNN/pluginsRT/ResizeLayerRT.h

@@ -1,68 +1,168 @@
-#include<cassert>
+#ifndef RESIZE_LAYER_RT_H
-#include "../kernels.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 {
+	nvinfer1::Dims getOutputDimensions(int index, const nvinfer1::Dims* inputs, int nbInputDims) noexcept override {
-		return DimsCHW{o_c, o_h, o_w};
+		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]), 
+        resizeForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
-					  reinterpret_cast<dnnType*>(outputs[0]), 
+					  reinterpret_cast<dnnType*>(outputs[0]),
 					  batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream);
 		return 0;
 	}
 
-
+	size_t getSerializationSize() const noexcept override {
-	virtual size_t getSerializationSize() 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);
+		writeBUF(buf, o_c);
-		tk::dnn::writeBUF(buf, o_h);
+		writeBUF(buf, o_h);
-		tk::dnn::writeBUF(buf, o_w);
+		writeBUF(buf, o_w);
 
-		tk::dnn::writeBUF(buf, i_c);
+		writeBUF(buf, i_c);
-		tk::dnn::writeBUF(buf, i_h);
+		writeBUF(buf, i_h);
-		tk::dnn::writeBUF(buf, i_w);
+		writeBUF(buf, i_w);
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ResizeLayerRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int i_c, i_h, i_w, o_c, o_h, o_w;
 };
+
+class ResizeLayerRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ResizeLayerRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int o_cTemp = readBUF<int>(buf);
+			int o_hTemp = readBUF<int>(buf);
+			int o_wTemp = readBUF<int>(buf);
+			ResizeLayerRT* r = new ResizeLayerRT(o_cTemp, o_hTemp, o_wTemp);
+
+			r->i_c = readBUF<int>(buf);
+			r->i_h = readBUF<int>(buf);
+			r->i_w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // RESIZE_LAYER_RT_H

include/tkDNN/pluginsRT/RouteRT.h

@@ -1,7 +1,20 @@
-#include<cassert>
+#ifndef ROUTE_RT_H
-#include "../kernels.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 noexcept override {
-
-	int getNbOutputs() const 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);
+		writeBUF(buf, groups);
-		tk::dnn::writeBUF(buf, group_id);
+		writeBUF(buf, group_id);
-		tk::dnn::writeBUF(buf, in);
+		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);
+		writeBUF(buf, c);
-		tk::dnn::writeBUF(buf, h);
+		writeBUF(buf, h);
-		tk::dnn::writeBUF(buf, w);
+		writeBUF(buf, w);
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new RouteRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	static const int MAX_INPUTS = 4;
 	int in;
 	int c_in[MAX_INPUTS];
 	int c, h, w;
 	int groups, group_id;
 };
+
+class RouteRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    RouteRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int groupsTemp = readBUF<int>(buf);
+			int group_idTemp = readBUF<int>(buf);
+			RouteRT* r = new RouteRT(groupsTemp, group_idTemp);
+			r->in = readBUF<int>(buf);
+			for(int i=0; i<RouteRT::MAX_INPUTS; i++)
+					r->c_in[i] = readBUF<int>(buf);
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // ROUTE_RT_H

include/tkDNN/pluginsRT/ShortcutRT.h

@@ -1,48 +1,64 @@
-#include<cassert>
+#ifndef SHORTCUT_RT_H
-#include "../kernels.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 noexcept override {
-
-	int getNbOutputs() const 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 DimsCHW{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]};
+		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);
+		writeBUF(buf, bc);
-		tk::dnn::writeBUF(buf, bh);
+		writeBUF(buf, bh);
-		tk::dnn::writeBUF(buf, bw);
+		writeBUF(buf, bw);
-		tk::dnn::writeBUF(buf, mul);
+		writeBUF(buf, mul);
-		tk::dnn::writeBUF(buf, c);
+		writeBUF(buf, c);
-		tk::dnn::writeBUF(buf, h);
+		writeBUF(buf, h);
-		tk::dnn::writeBUF(buf, w);
+		writeBUF(buf, w);
 		assert(buf == a + getSerializationSize());
-		
+
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new ShortcutRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
 	int bc, bh, bw;
 	bool mul;
 };
+
+class ShortcutRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    ShortcutRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			dataDim_t bdim;
+			bdim.c = readBUF<int>(buf);
+			bdim.h = readBUF<int>(buf);
+			bdim.w = readBUF<int>(buf);
+			bdim.l = 1;
+
+			ShortcutRT *r = new ShortcutRT(bdim, readBUF<bool>(buf));
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // SHORTCUT_RT_H

include/tkDNN/pluginsRT/UpsampleRT.h

@@ -1,66 +1,160 @@
-#include<cassert>
+#ifndef UPSAMPLE_RT_H
-#include "../kernels.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 noexcept override {
-
-	int getNbOutputs() const 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 DimsCHW(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride);
+		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;
 	}
 
-
+	size_t getSerializationSize() const noexcept override {
-	virtual size_t getSerializationSize() 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);
+		writeBUF(buf, stride);
-		tk::dnn::writeBUF(buf, c);
+		writeBUF(buf, c);
-		tk::dnn::writeBUF(buf, h);
+		writeBUF(buf, h);
-		tk::dnn::writeBUF(buf, w);
+		writeBUF(buf, w);
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new UpsampleRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
 	int c, h, w, stride;
 };
+
+class UpsampleRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    UpsampleRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override {
+			const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+			int strideTemp = readBUF<int>(buf);
+			UpsampleRT* r = new UpsampleRT(strideTemp);
+			r->c = readBUF<int>(buf);
+			r->h = readBUF<int>(buf);
+			r->w = readBUF<int>(buf);
+			assert(buf == bufCheck + serialLength);
+			return r;
+		}
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // UPSAMPLE_RT_H

include/tkDNN/pluginsRT/YoloRT.h

@@ -1,15 +1,26 @@
-#include<cassert>
+#ifndef YOLO_RT_H
+#define YOLO_RT_H
+
+#include <cassert>
+#include <vector>
+
+#include <NvInferRuntimeCommon.h>
+#include <NvInfer.h>
+
+#include "../yoloContainer.h"
 #include "../kernels.h"
+#include "../buffer_func.h"
+#include "../Layer.h"
 
 #define YOLORT_CLASSNAME_W 256
 
-class YoloRT : public IPlugin {
+#define PLUGIN_NAME "Yolo"
-
+#define PLUGIN_VERSION "1"
-
+namespace tk { namespace dnn {
 
+class YoloRT final : public nvinfer1::IPluginV2 {
 public:
-	YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr, int n_masks=3, float scale_xy=1, float nms_thresh=0.45, int nms_kind=0, int new_coords=0) {
+	YoloRT(int classes, int num, Yolo *yolo = nullptr, int n_masks=3, float scale_xy=1, float nms_thresh=0.45, int nms_kind=0, int new_coords=0) {
-
 		this->classes = classes;
 		this->num = num;
 		this->n_masks = n_masks;
@@ -27,38 +38,40 @@ public:
         }
 	}
 
-	~YoloRT(){
+	~YoloRT() = default;
 
-	}
+	int getNbOutputs() const noexcept override {
-
-	int getNbOutputs() const 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;
 	}
 
-
+	size_t getSerializationSize() const noexcept override {
-	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);
 	}
 
-	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;
+		writeBUF(buf, classes); 	//std::cout << "Classes :" << classes << std::endl;
-		tk::dnn::writeBUF(buf, num); 		//std::cout << "Num : " << num << std::endl;
+		writeBUF(buf, num); 		//std::cout << "Num : " << num << std::endl;
-		tk::dnn::writeBUF(buf, n_masks); 	//std::cout << "N_Masks" << n_masks << std::endl;
+		writeBUF(buf, n_masks); 	//std::cout << "N_Masks" << n_masks << std::endl;
-		tk::dnn::writeBUF(buf, scaleXY); 	//std::cout << "ScaleXY :" << scaleXY << std::endl;
+		writeBUF(buf, scaleXY); 	//std::cout << "ScaleXY :" << scaleXY << std::endl;
-		tk::dnn::writeBUF(buf, nms_thresh); //std::cout << "nms_thresh :" << nms_thresh << std::endl;
+		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;
+		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;
+		writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl;
-		tk::dnn::writeBUF(buf, c); 			//std::cout << "C : " << c << std::endl;
+		writeBUF(buf, c); 			//std::cout << "C : " << c << std::endl;
-		tk::dnn::writeBUF(buf, h); 			//std::cout << "H : " << h << std::endl;
+		writeBUF(buf, h); 			//std::cout << "H : " << h << std::endl;
-		tk::dnn::writeBUF(buf, w); 			//std::cout << "C : " << c << std::endl;
+		writeBUF(buf, w); 			//std::cout << "C : " << c << std::endl;
 		for (int i = 0; i < n_masks; i++)
 		{
-			tk::dnn::writeBUF(buf, mask[i]); //std::cout << "mask[i] : " << mask[i] << std::endl;
+			writeBUF(buf, mask[i]); //std::cout << "mask[i] : " << mask[i] << std::endl;
 		}
 		for (int i = 0; i < n_masks * 2 * num; i++)
 		{
-			tk::dnn::writeBUF(buf, bias[i]); //std::cout << "bias[i] : " << bias[i] << std::endl;
+			writeBUF(buf, bias[i]); //std::cout << "bias[i] : " << bias[i] << std::endl;
 		}
 
 		// save classes names
@@ -117,12 +129,42 @@ public:
 			char tmp[YOLORT_CLASSNAME_W];
 			strcpy(tmp, classesNames[i].c_str());
 			for(int j=0; j<YOLORT_CLASSNAME_W; j++) {
-				tk::dnn::writeBUF(buf, tmp[j]);
+				writeBUF(buf, tmp[j]);
 			}
 		}
 		assert(buf == a + getSerializationSize());
 	}
 
+	// Extra IPluginV2 overrides
+	bool supportsFormat(nvinfer1::DataType type, nvinfer1::PluginFormat format) const noexcept override {
+			return (type == nvinfer1::DataType::kFLOAT && format == nvinfer1::PluginFormat::kLINEAR);
+	}
+
+	nvinfer1::IPluginV2 * clone() const noexcept override {
+			auto a = new YoloRT(*this);
+			return a;
+	}
+
+	const char* getPluginType() const noexcept override {
+			return PLUGIN_NAME;
+	}
+
+	const char* getPluginVersion() const noexcept override {
+			return PLUGIN_VERSION;
+	}
+
+	void destroy() noexcept override {}
+
+	void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+	}
+
+	const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+	}
+
+	std::string mNamespace;
+
 	int c, h, w;
     int classes, num, n_masks;
 	float scaleXY;
@@ -139,5 +181,46 @@ public:
 		int loc = location % (w*h);
 		return batch*c*h*w + n*w*h*(4+classes+1) + entry*w*h + loc;
 	}
-
 };
+
+class YoloRTCreator final : public nvinfer1::IPluginCreator {
+public:
+    YoloRTCreator() = default;
+
+    const char* getPluginName() const noexcept override {
+        return PLUGIN_NAME;
+    }
+
+    const char* getPluginVersion() const noexcept override {
+        return PLUGIN_VERSION;
+    }
+
+    const nvinfer1::PluginFieldCollection* getFieldNames() noexcept override {
+        return &mFC;
+    }
+
+    nvinfer1::IPluginV2* createPlugin(const char* name, const nvinfer1::PluginFieldCollection* fc) noexcept override {
+        std::cout << "Create plugin" << std::endl;
+        return nullptr;
+    }
+
+    nvinfer1::IPluginV2* deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept override;
+
+    void setPluginNamespace(const char* pluginNamespace) noexcept override {
+			mNamespace = pluginNamespace;
+		}
+
+    const char* getPluginNamespace() const noexcept override {
+			return mNamespace.c_str();
+		}
+
+private:
+    static nvinfer1::PluginFieldCollection mFC;
+    static std::vector<nvinfer1::PluginField> mPluginAttributes;
+    std::string mNamespace;
+};
+}}
+#undef PLUGIN_NAME
+#undef PLUGIN_VERSION
+
+#endif // YOLO_RT_H

include/tkDNN/yoloContainer.h

@@ -0,0 +1,16 @@
+#ifndef YOLO_CONTAINER_H
+#define YOLO_CONTAINER_H
+
+namespace tk { namespace dnn {
+class YoloRT;
+class YoloContainer
+{
+public:
+    YoloRT *yolos[16];
+    int n_yolos{};
+};
+
+extern YoloContainer yoloContainer;
+}}
+
+#endif // YOLO_CONTAINER_H

src/Int8BatchStream.cpp

@@ -8,14 +8,14 @@
 BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, const std::string& fileimglist, const std::string& filelabellist) {
     mBatchSize = batchSize;
     mMaxBatches = maxBatches;
-    mDims = nvinfer1::DimsNCHW{ dim.n, dim.c, dim.h, dim.w };
+    mDims = nvinfer1::Dims4{ dim.n, dim.c, dim.h, dim.w };
     mHeight = dim.h;
     mWidth = dim.w;
-    mImageSize = mDims.c()*mDims.h()*mDims.w();
+    mImageSize = dim.c*dim.h*dim.w;
     mBatch.resize(mBatchSize*mImageSize, 0);
     mLabels.resize(mBatchSize, 0);
-    mFileBatch.resize(mDims.n()*mImageSize, 0);
+    mFileBatch.resize(dim.n*mImageSize, 0);
-    mFileLabels.resize(mDims.n(), 0);
+    mFileLabels.resize(dim.n, 0);
     mFileImgList = fileimglist;
     readInListFile(fileimglist, mListImg);
     mFileLabelList = filelabellist;
@@ -27,7 +27,7 @@ BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches,
 void BatchStream::reset(int firstBatch) {
     mBatchCount = 0;
     mFileCount = 0;
-    mFileBatchPos = mDims.n();
+    mFileBatchPos = mDims.d[0];
     skip(firstBatch);
 }
 
@@ -37,11 +37,11 @@ bool BatchStream::next() {
         return false;
 
     for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize) {
-        assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.n());
+        assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.d[0]);
-        if (mFileBatchPos == mDims.n() && !update())
+        if (mFileBatchPos == mDims.d[0] && !update())
             return false;
 
-        csize = std::min(mBatchSize - batchPos, mDims.n() - mFileBatchPos);
+        csize = std::min(mBatchSize - batchPos, mDims.d[0] - mFileBatchPos);
         std::copy_n(getFileBatch() + mFileBatchPos * mImageSize, csize * mImageSize, getBatch() + batchPos * mImageSize);
         std::copy_n(getFileLabels() + mFileBatchPos, csize, getLabels() + batchPos);
     }
@@ -50,8 +50,8 @@ bool BatchStream::next() {
 }
 
 void BatchStream::skip(int skipCount) {
-    if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n()) {
+    if (mBatchSize >= mDims.d[0] && mBatchSize%mDims.d[0] == 0 && mFileBatchPos == mDims.d[0]) {
-        mFileCount += skipCount * mBatchSize / mDims.n();
+        mFileCount += skipCount * mBatchSize / mDims.d[0];
         return;
     }
 
@@ -67,7 +67,7 @@ void BatchStream::readInListFile(const std::string& dataFilePath, std::vector<st
     FILE* f = fopen(dataFilePath.c_str(), "r");
     if (!f)
         FatalError("failed to open " + dataFilePath);
-    
+
     char str[512];
     while (fgets(str, 512, f) != NULL) {
         for (int i = 0; str[i] != '\0'; ++i) {
@@ -132,7 +132,7 @@ void BatchStream::readCVimage(std::string inputFileName, std::vector<float>& res
 
 void BatchStream::readLabels(std::string inputFileName, std::vector<float>& ris) {
     std::ifstream is(inputFileName.c_str());
-    
+
     std::string line;
     while (std::getline(is, line))
     {
@@ -153,13 +153,13 @@ bool BatchStream::update() {
     readCVimage(imgFileName, mFileBatch);
     // std::transform(
     //     singleImg_rawData.begin(), singleImg_rawData.end(), mFileBatch.begin(), [](uint8_t val) { return static_cast<float>(val); });
-    
+
     //read label
     mFileLabels.clear();
     readLabels(labelFileName, mFileLabels);
     // std::transform(
     //     singleLabels_rawData.begin(), singleLabels_rawData.end(), mFileLabels.begin(), [](uint8_t val) { return static_cast<float>(val); });
-    
+
     mFileBatchPos = 0;
     return true;
 }

src/Int8Calibrator.cpp

@@ -1,20 +1,20 @@
 #include "Int8Calibrator.h"
 
-Int8EntropyCalibrator::Int8EntropyCalibrator(BatchStream& stream, int firstBatch, 
+Int8EntropyCalibrator::Int8EntropyCalibrator(BatchStream& stream, int firstBatch,
                                              const std::string& calibTableFilePath,
                                              const std::string& inputBlobName,
-                                             bool readCache): 
+                                             bool readCache):
-    mStream(stream), 
+    mStream(stream),
     mCalibTableFilePath(calibTableFilePath),
     mInputBlobName(inputBlobName.c_str()),
     mReadCache(readCache) {
-    nvinfer1::DimsNCHW dims = mStream.getDims();
+    nvinfer1::Dims4 dims = mStream.getDims();
-    mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w();
+    mInputCount = mStream.getBatchSize() * dims.d[1] * dims.d[2] * dims.d[3];
     checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float)));
     mStream.reset(firstBatch);
 }
 
-bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int nbBindings) {
+bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int nbBindings) noexcept {
     if (!mStream.next())
         return false;
 
@@ -24,7 +24,7 @@ bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int
     return true;
 }
 
-const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) {
+const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) noexcept {
     mCalibrationCache.clear();
     assert(!mCalibTableFilePath.empty());
     std::ifstream input(mCalibTableFilePath, std::ios::binary);
@@ -38,7 +38,7 @@ const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) {
     return length ? &mCalibrationCache[0] : nullptr;
 }
 
-void Int8EntropyCalibrator::writeCalibrationCache(const void* cache, size_t length) {
+void Int8EntropyCalibrator::writeCalibrationCache(const void* cache, size_t length) noexcept {
     assert(!mCalibTableFilePath.empty());
     std::ofstream output(mCalibTableFilePath, std::ios::binary);
     output.write(reinterpret_cast<const char*>(cache), length);

src/NetworkRT.cpp

@@ -15,7 +15,7 @@ using namespace nvinfer1;
 
 // Logger for info/warning/errors
 class Logger : public ILogger {
-    void log(Severity severity, const char* msg) override {
+    void log(Severity severity, const char* msg) noexcept override {
 #ifdef DEBUG
         std::cout <<"TENSORRT LOG: "<< msg << std::endl;
 #endif
@@ -24,28 +24,28 @@ class Logger : public ILogger {
 
 namespace tk { namespace dnn {
 
-std::map<Layer*, nvinfer1::ITensor*>tensors; 
+std::map<Layer*, nvinfer1::ITensor*>tensors;
 
 NetworkRT::NetworkRT(Network *net, const char *name) {
 
-    float rt_ver = float(NV_TENSORRT_MAJOR) + 
+    float rt_ver = float(NV_TENSORRT_MAJOR) +
-                   float(NV_TENSORRT_MINOR)/10 + 
+                   float(NV_TENSORRT_MINOR)/10 +
                    float(NV_TENSORRT_PATCH)/100;
     std::cout<<"New NetworkRT (TensorRT v"<<rt_ver<<")\n";
-  
+
     builderRT = createInferBuilder(loggerRT);
     std::cout<<"Float16 support: "<<builderRT->platformHasFastFp16()<<"\n";
     std::cout<<"Int8 support: "<<builderRT->platformHasFastInt8()<<"\n";
 #if NV_TENSORRT_MAJOR >= 5
     std::cout<<"DLAs: "<<builderRT->getNbDLACores()<<"\n";
 #endif
-    networkRT = builderRT->createNetwork();
+    networkRT = builderRT->createNetworkV2(0u);
-#if NV_TENSORRT_MAJOR >= 6                
+#if NV_TENSORRT_MAJOR >= 6
         configRT = builderRT->createBuilderConfig();
 #endif
-    
+
     if(!fileExist(name)) {
-#if NV_TENSORRT_MAJOR >= 6                
+#if NV_TENSORRT_MAJOR >= 6
         // Calibrator life time needs to last until after the engine is built.
         std::unique_ptr<IInt8EntropyCalibrator> calibrator;
 
@@ -53,22 +53,25 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
         configRT->setMinTimingIterations(1);
         configRT->setMaxWorkspaceSize(1 << 30);
         configRT->setFlag(BuilderFlag::kDEBUG);
+#else
+        builderRT->setMaxWorkspaceSize(1 << 30);
 #endif
         //input and dataType
         dataDim_t dim = net->layers[0]->input_dim;
         dtRT = DataType::kFLOAT;
 
         builderRT->setMaxBatchSize(net->maxBatchSize);
-        builderRT->setMaxWorkspaceSize(1 << 30);
 
         if(net->fp16 && builderRT->platformHasFastFp16()) {
             dtRT = DataType::kHALF;
+#if NV_TENSORRT_MAJOR < 6
             builderRT->setHalf2Mode(true);
-#if NV_TENSORRT_MAJOR >= 6                
+#endif
+#if NV_TENSORRT_MAJOR >= 6
             configRT->setFlag(BuilderFlag::kFP16);
 #endif
         }
-#if NV_TENSORRT_MAJOR >= 5
+#if NV_TENSORRT_MAJOR >= 5 && NV_TENSORRT_MAJOR < 8
         if(net->dla && builderRT->getNbDLACores() > 0) {
             dtRT = DataType::kHALF;
             builderRT->setFp16Mode(true);
@@ -77,14 +80,23 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
             builderRT->setDLACore(0);
         }
 #endif
-#if NV_TENSORRT_MAJOR >= 6                
+#if NV_TENSORRT_MAJOR >= 8
+        if(net->dla && builderRT->getNbDLACores() > 0) {
+            dtRT = DataType::kHALF;
+            configRT->setFlag(BuilderFlag::kFP16);
+            configRT->setFlag(BuilderFlag::kGPU_FALLBACK);
+            configRT->setDefaultDeviceType(DeviceType::kDLA);
+            configRT->setDLACore(0);
+        }
+#endif
+#if NV_TENSORRT_MAJOR >= 6
         if(net->int8 && builderRT->platformHasFastInt8()){
             // dtRT = DataType::kINT8;
             // builderRT->setInt8Mode(true);
             configRT->setFlag(BuilderFlag::kINT8);
-            BatchStream calibrationStream(dim, 1, 100,      //TODO: check if 100 images are sufficient to the calibration (or 4951) 
+            BatchStream calibrationStream(dim, 1, 100,      //TODO: check if 100 images are sufficient to the calibration (or 4951)
                                             net->fileImgList, net->fileLabelList);
-            
+
             /* The calibTableFilePath contains the path+filename of the calibration table.
              * Each calibration table can be found in the corresponding network folder (../Test/*).
              * Each network is located in a folder with the same name as the network.
@@ -95,33 +107,33 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
             if(!fileExist((const char *)calib_table_path.c_str()))
                 calib_table_name = "./" + net->networkNameRT.substr(0, net->networkNameRT.find('.')) + "-calibration.table";
 
-            calibrator.reset(new Int8EntropyCalibrator(calibrationStream, 1, 
+            calibrator.reset(new Int8EntropyCalibrator(calibrationStream, 1,
-                                            calib_table_name, 
+                                            calib_table_name,
                                             "data"));
             configRT->setInt8Calibrator(calibrator.get());
         }
 #endif
-        
+
         // add input layer
-        ITensor *input = networkRT->addInput("data", DataType::kFLOAT, 
+        ITensor *input = networkRT->addInput("data", DataType::kFLOAT,
-                        DimsCHW{ dim.c, dim.h, dim.w});
+                        Dims3{ dim.c, dim.h, dim.w});
         checkNULL(input);
 
         //add other layers
         for(int i=0; i<net->num_layers; i++) {
             Layer *l = net->layers[i];
             ILayer *Ilay = convert_layer(input, l);
-#if NV_TENSORRT_MAJOR >= 6                
+#if NV_TENSORRT_MAJOR >= 6
             if(net->int8 && builderRT->platformHasFastInt8())
             {
                 Ilay->setPrecision(DataType::kINT8);
             }
 #endif
             Ilay->setName( (l->getLayerName() + std::to_string(i)).c_str() );
-            
+
             input = Ilay->getOutput(0);
             input->setName( (l->getLayerName() + std::to_string(i) + "_out").c_str() );
-            
+
             if(l->final)
                 networkRT->markOutput(*input);
             tensors[l] = input;
@@ -136,9 +148,9 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
         std::cout<<"Selected maxBatchSize: "<<builderRT->getMaxBatchSize()<<"\n";
         printCudaMemUsage();
         std::cout<<"Building tensorRT cuda engine...\n";
-#if NV_TENSORRT_MAJOR >= 6                
+#if NV_TENSORRT_MAJOR >= 6
         engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT);
-#else 
+#else
         engineRT = builderRT->buildCudaEngine(*networkRT);
         //engineRT = std::shared_ptr<nvinfer1::ICudaEngine>(builderRT->buildCudaEngine(*networkRT));
 #endif
@@ -162,7 +174,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
 
 	// In order to bind the buffers, we need to know the names of the input and output tensors.
 	// note that indices are guaranteed to be less than IEngine::getNbBindings()
-	buf_input_idx = engineRT->getBindingIndex("data"); 
+	buf_input_idx = engineRT->getBindingIndex("data");
     buf_output_idx = engineRT->getBindingIndex("out");
     std::cout<<"input index = "<<buf_input_idx<<" -> output index = "<<buf_output_idx<<"\n";
 
@@ -180,7 +192,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
     output_dim.h = oDim.d[1];
     output_dim.w = oDim.d[2];
     output_dim.print();
-	
+
     // create GPU buffers and a stream
     for(int i=0; i<engineRT->getNbBindings(); i++) {
         Dims dim = engineRT->getBindingDimensions(i);
@@ -261,10 +273,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Dense *l) {
     //std::cout<<"convert Dense\n";
     void *data_b, *bias_b;
     if(dtRT == DataType::kHALF) {
-        data_b     = l->data16_h;    
+        data_b     = l->data16_h;
         bias_b     = l->bias16_h;
     } else {
-        data_b     = l->data_h;    
+        data_b     = l->data_h;
         bias_b     = l->bias_h;
     }
 
@@ -284,7 +296,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
 
     void *data_b, *bias_b, *bias2_b, *power_b, *mean_b, *variance_b, *scales_b;
     if(dtRT == DataType::kHALF) {
-        data_b     = l->data16_h;    
+        data_b     = l->data16_h;
         bias_b     = l->bias16_h;
         bias2_b    = l->bias216_h;
         power_b    = l->power16_h;
@@ -292,7 +304,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
         variance_b = l->variance16_h;
         scales_b   = l->scales16_h;
     } else {
-        data_b     = l->data_h;    
+        data_b     = l->data_h;
         bias_b     = l->bias_h;
         bias2_b    = l->bias2_h;
         power_b    = l->power_h;
@@ -308,14 +320,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
         b = { dtRT, bias_b, l->outputs};
     else{
         if (l->additional_bias)
-            b = { dtRT, bias2_b, l->outputs}; 
+            b = { dtRT, bias2_b, l->outputs};
         else
             b = { dtRT, nullptr, 0}; //on batchnorm bias are added later
     }
 
     ILayer *lRT = nullptr;
     if(!l->deConv) {
-        IConvolutionLayer *lRTconv = networkRT->addConvolution(*input, 
+        IConvolutionLayer *lRTconv = networkRT->addConvolution(*input,
             l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b);
         checkNULL(lRTconv);
         lRTconv->setStride(DimsHW{l->strideH, l->strideW});
@@ -323,14 +335,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
         lRTconv->setNbGroups(l->groups);
         lRT = (ILayer*) lRTconv;
     } else {
-        IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input, 
+        IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input,
             l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b);
         checkNULL(lRTconv);
         lRTconv->setStride(DimsHW{l->strideH, l->strideW});
         lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW});
         lRTconv->setNbGroups(l->groups);
         lRT = (ILayer*) lRTconv;
-        
+
         Dims d = lRTconv->getOutput(0)->getDimensions();
         //std::cout<<"DECONV: "<<d.d[0]<<" "<<d.d[1]<<" "<<d.d[2]<<" "<<d.d[3]<<"\n";
     }
@@ -341,14 +353,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
         Weights shift{dtRT, mean_b, l->outputs};
         Weights scale{dtRT, variance_b, l->outputs};
         // std::cout<<lRT->getNbOutputs()<<std::endl;
-        IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, 
+        IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL,
                     shift, scale, power);
-        
+
         checkNULL(lRT2);
 
         Weights shift2{dtRT, bias_b, l->outputs};
         Weights scale2{dtRT, scales_b, l->outputs};
-        IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, 
+        IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL,
                     shift2, scale2, power);
         checkNULL(lRT3);
 
@@ -368,8 +380,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) {
 
     if(l->pool_mode == tkdnnPoolingMode_t::POOLING_MAX_FIXEDSIZE)
     {
-        IPlugin *plugin = new MaxPoolFixedSizeRT(l->output_dim.c, l->output_dim.h, l->output_dim.w, l->output_dim.n, l->strideH, l->strideW, l->winH, l->winH-1);        
+        auto *plugin = new MaxPoolFixedSizeRT(l->output_dim.c, l->output_dim.h, l->output_dim.w, l->output_dim.n, l->strideH, l->strideW, l->winH, l->winH-1);
-        IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+        auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
         checkNULL(lRT);
         return lRT;
     }
@@ -381,7 +393,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) {
         lRT->setPadding(DimsHW{l->paddingH, l->paddingW});
         lRT->setStride(DimsHW{l->strideH, l->strideW});
         return lRT;
-    }  
+    }
 }
 
 ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
@@ -389,14 +401,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
 
     if(l->act_mode == ACTIVATION_LEAKY) {
         //std::cout<<"New plugin LEAKY\n";
-        
+
-#if NV_TENSORRT_MAJOR < 6                
+#if NV_TENSORRT_MAJOR < 6
         // plugin version
-        IPlugin *plugin = new ActivationLeakyRT(l->slope);
+        auto *plugin = new ActivationLeakyRT(l->slope);
-        IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+        auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
         checkNULL(lRT);
         return lRT;
-#else 
+#else
         IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kLEAKY_RELU);
         lRT->setAlpha(l->slope);
         checkNULL(lRT);
@@ -407,28 +419,35 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
         IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU);
         checkNULL(lRT);
         return lRT;
-    } else if(l->act_mode == CUDNN_ACTIVATION_SIGMOID) {
+    } else if(l->act_mode == CUDNN_ACTIVATION_SIGMOID || l->act_mode == ACTIVATION_LOGISTIC) {
         IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kSIGMOID);
         checkNULL(lRT);
         return lRT;
     }
     else if(l->act_mode == CUDNN_ACTIVATION_CLIPPED_RELU) {
-        IPlugin *plugin = new ActivationReLUCeiling(l->ceiling);
+        IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kCLIP);
-        IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+        lRT->setAlpha(0);
-        checkNULL(lRT);
+        lRT->setBeta(l->ceiling);
-        return lRT;
+
-    } 
-    else if(l->act_mode == ACTIVATION_MISH) {
-        IPlugin *plugin = new ActivationMishRT();
-        IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
         checkNULL(lRT);
         return lRT;
     }
-    else if(l->act_mode == ACTIVATION_LOGISTIC) {
+    else if(l->act_mode == ACTIVATION_MISH) {
-        IPlugin *plugin = new ActivationLogisticRT();
+        // Uncomment this to see if you have better performance
-        IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+        // For older TensorRT or for FP32 this might be better
-        checkNULL(lRT);
+        //auto *plugin = new ActivationMishRT();
-        return lRT;
+        //auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
+
+        // Assemble MISH using 3 layers that are going to be merged by TensorRT
+        IActivationLayer *lRT1 = networkRT->addActivation(*input, ActivationType::kSOFTPLUS);
+        lRT1->setAlpha(1);
+        lRT1->setBeta(1);
+
+        IActivationLayer *lRT2 = networkRT->addActivation(*lRT1->getOutput(0), ActivationType::kTANH);
+        IElementWiseLayer *lRT3 = networkRT->addElementWise(*input, *lRT2->getOutput(0), ElementWiseOperation::kPROD);
+
+        checkNULL(lRT3);
+        return lRT3;
     }
     else {
         FatalError("this Activation mode is not yet implemented");
@@ -447,7 +466,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Softmax *l) {
 
 ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) {
     // std::cout<<"convert route\n";
-    
+
 
 
     ITensor **tens = new ITensor*[l->layers_n];
@@ -460,8 +479,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) {
     }
 
     if(l->groups > 1){
-        IPlugin *plugin = new RouteRT(l->groups, l->group_id);
+        auto *plugin = new RouteRT(l->groups, l->group_id);
-        IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin);
+        auto *lRT = networkRT->addPluginV2(tens, l->layers_n, *plugin);
         checkNULL(lRT);
         return lRT;
     }
@@ -472,8 +491,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) {
 
 ILayer* NetworkRT::convert_layer(ITensor *input, Flatten *l) {
 
-    IPlugin *plugin = new FlattenConcatRT();
+    auto *plugin = new FlattenConcatRT();
-    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
     checkNULL(lRT);
     return lRT;
 }
@@ -481,8 +500,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Flatten *l) {
 ILayer* NetworkRT::convert_layer(ITensor *input, Reshape *l) {
     // std::cout<<"convert Reshape\n";
 
-    IPlugin *plugin = new ReshapeRT(l->output_dim);
+    auto *plugin = new ReshapeRT(l->output_dim);
-    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
     checkNULL(lRT);
     return lRT;
 }
@@ -494,7 +513,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Resize *l) {
     checkNULL(lRT);
     Dims d{};
     lRT->setResizeMode(ResizeMode(l->mode));
-    lRT->setOutputDimensions(DimsCHW{l->output_dim.c, l->output_dim.h, l->output_dim.w});
+    lRT->setOutputDimensions(Dims3{l->output_dim.c, l->output_dim.h, l->output_dim.w});
     return lRT;
 }
 
@@ -502,8 +521,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Reorg *l) {
     //std::cout<<"convert Reorg\n";
 
     //std::cout<<"New plugin REORG\n";
-    IPlugin *plugin = new ReorgRT(l->stride);
+    auto *plugin = new ReorgRT(l->stride);
-    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
     checkNULL(lRT);
     return lRT;
 }
@@ -512,8 +531,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Region *l) {
     //std::cout<<"convert Region\n";
 
     //std::cout<<"New plugin REGION\n";
-    IPlugin *plugin = new RegionRT(l->classes, l->coords, l->num);
+    auto *plugin = new RegionRT(l->classes, l->coords, l->num);
-    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
     checkNULL(lRT);
     return lRT;
 }
@@ -522,10 +541,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
     //std::cout<<"convert Shortcut\n";
 
     //std::cout<<"New plugin Shortcut\n";
-    
+
     ITensor *back_tens = tensors[l->backLayer];
 
-    if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul) 
+    if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul)
     {
         IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM);
         checkNULL(lRT);
@@ -534,11 +553,11 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
     else
     {
         // plugin version
-        IPlugin *plugin = new ShortcutRT(l->backLayer->output_dim, l->mul);
+        auto *plugin = new ShortcutRT(l->backLayer->output_dim, l->mul);
         ITensor **inputs = new ITensor*[2];
         inputs[0] = input;
-        inputs[1] = back_tens; 
+        inputs[1] = back_tens;
-        IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin);
+        auto *lRT = networkRT->addPluginV2(inputs, 2, *plugin);
         checkNULL(lRT);
         return lRT;
     }
@@ -548,8 +567,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Yolo *l) {
     //std::cout<<"convert Yolo\n";
 
     //std::cout<<"New plugin YOLO\n";
-    IPlugin *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY, l->nms_thresh, l->nsm_kind, l->new_coords);
+    auto *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY, l->nms_thresh, l->nsm_kind, l->new_coords);
-    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
     checkNULL(lRT);
     return lRT;
 }
@@ -558,8 +577,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Upsample *l) {
     //std::cout<<"convert Upsample\n";
 
     //std::cout<<"New plugin UPSAMPLE\n";
-    IPlugin *plugin = new UpsampleRT(l->stride);
+    auto *plugin = new UpsampleRT(l->stride);
-    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    auto *lRT = networkRT->addPluginV2(&input, 1, *plugin);
     checkNULL(lRT);
     return lRT;
 }
@@ -574,10 +593,10 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
     inputs[1] = preconv->getOutput(0);
 
     //std::cout<<"New plugin DEFORMABLE\n";
-    IPlugin *plugin = new DeformableConvRT(l->chunk_dim, l->kernelH, l->kernelW, l->strideH, l->strideW, l->paddingH, l->paddingW, 
+    auto *plugin = new DeformableConvRT(l->chunk_dim, l->kernelH, l->kernelW, l->strideH, l->strideW, l->paddingH, l->paddingW,
-                                            l->deformableGroup, l->input_dim.n, l->input_dim.c, l->input_dim.h, l->input_dim.w, 
+                                            l->deformableGroup, l->input_dim.n, l->input_dim.c, l->input_dim.h, l->input_dim.w,
                                             l->output_dim.n, l->output_dim.c, l->output_dim.h, l->output_dim.w, l);
-    IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin);
+    auto *lRT = networkRT->addPluginV2(inputs, 2, *plugin);
     checkNULL(lRT);
     lRT->setName( ("Deformable" + std::to_string(l->id)).c_str() );
     delete[](inputs);
@@ -601,14 +620,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
     Weights shift{dtRT, mean_b, l->outputs};
     Weights scale{dtRT, variance_b, l->outputs};
     //std::cout<<lRT->getNbOutputs()<<std::endl;
-    IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, 
+    IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL,
                 shift, scale, power);
-    
+
     checkNULL(lRT2);
 
     Weights shift2{dtRT, bias_b, l->outputs};
     Weights scale2{dtRT, scales_b, l->outputs};
-    IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, 
+    IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL,
                 shift2, scale2, power);
     checkNULL(lRT3);
 
@@ -646,53 +665,52 @@ bool NetworkRT::deserialize(const char *filename) {
         file.close();
     }
 
-    pluginFactory = new PluginFactory();
     runtimeRT = createInferRuntime(loggerRT);
-    engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory);
+    engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size);
     //if (gieModelStream) delete [] gieModelStream;
 
     return true;
 }
 
+// auto* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
+//     const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
 
+//     std::string name(layerName);
+//     //std::cout<<name<<std::endl;
 
-IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
+// #if NV_TENSORRT_MAJOR < 6
-    const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+//     if(name.find("ActivationLeaky") == 0) {
-
+//         ActivationLeakyRT *a = new ActivationLeakyRT(readBUF<float>(buf));
-    std::string name(layerName);
+//         a->size = readBUF<int>(buf);
-    //std::cout<<name<<std::endl;
+//         assert(buf == bufCheck + serialLength);
-
+//         return a;
-    if(name.find("ActivationLeaky") == 0) {
+//     }
-        ActivationLeakyRT *a = new ActivationLeakyRT(readBUF<float>(buf));
+// #endif
-        a->size = readBUF<int>(buf);
+    /*if(name.find("ActivationMish") == 0) {
-        assert(buf == bufCheck + serialLength);
-        return a;
-    }
-    if(name.find("ActivationMish") == 0) {
         ActivationMishRT *a = new ActivationMishRT();
         a->size = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return a;
-    }
+    }*/
-    if(name.find("ActivationLogistic") == 0) {
+    /*if(name.find("ActivationLogistic") == 0) {
         ActivationLogisticRT *a = new ActivationLogisticRT();
         a->size = readBUF<int>(buf);
         return a;
-    }
+    }*/
-    if(name.find("ActivationLogistic") == 0) {
+    /*if(name.find("ActivationLogistic") == 0) {
         ActivationLogisticRT *a = new ActivationLogisticRT();
         a->size = readBUF<int>(buf);
         return a;
-    }
+    }*/
-    if(name.find("ActivationCReLU") == 0) {
+    /*if(name.find("ActivationCReLU") == 0) {
         float activationReluTemp = readBUF<float>(buf);
         ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp);
         a->size = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return a;
-    }
+    }*/
 
-    if(name.find("Region") == 0) {
+    /*if(name.find("Region") == 0) {
         int classesTemp = readBUF<int>(buf);
         int coordsTemp = readBUF<int>(buf);
         int numTemp = readBUF<int>(buf);
@@ -703,9 +721,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Reorg") == 0) {
+    /*if(name.find("Reorg") == 0) {
         int strideTemp = readBUF<int>(buf);
         ReorgRT *r = new ReorgRT(strideTemp);
         r->c = readBUF<int>(buf);
@@ -713,9 +731,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Shortcut") == 0) {
+    /*if(name.find("Shortcut") == 0) {
         tk::dnn::dataDim_t bdim;
         bdim.c = readBUF<int>(buf);
         bdim.h = readBUF<int>(buf);
@@ -728,9 +746,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(buf);
         return r;
         assert(buf == bufCheck + serialLength);
-    } 
+    }*/
 
-    if(name.find("Pooling") == 0) {
+    /*if(name.find("Pooling") == 0) {
         int cTemp = readBUF<int>(buf);
         int hTemp = readBUF<int>(buf);
         int wTemp = readBUF<int>(buf);
@@ -743,9 +761,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp);
         assert(buf == bufCheck + serialLength);
         return r;
-    }
+    }*/
 
-    if(name.find("Resize") == 0) {
+    /*if(name.find("Resize") == 0) {
         int o_cTemp = readBUF<int>(buf);
         int o_hTemp = readBUF<int>(buf);
         int o_wTemp = readBUF<int>(buf);
@@ -756,10 +774,10 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->i_w = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Flatten") == 0) {
+    /*if(name.find("Flatten") == 0) {
-        FlattenConcatRT *r = new FlattenConcatRT(); 
+        FlattenConcatRT *r = new FlattenConcatRT();
         r->c = readBUF<int>(buf);
         r->h = readBUF<int>(buf);
         r->w = readBUF<int>(buf);
@@ -767,22 +785,22 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->cols = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Reshape") == 0) {
+    /*if(name.find("Reshape") == 0) {
 
         dataDim_t new_dim;
         new_dim.n = readBUF<int>(buf);
         new_dim.c = readBUF<int>(buf);
         new_dim.h = readBUF<int>(buf);
         new_dim.w = readBUF<int>(buf);
-        ReshapeRT *r = new ReshapeRT(new_dim); 
+        ReshapeRT *r = new ReshapeRT(new_dim);
         assert(buf == bufCheck + serialLength);
-        
-        return r;
-    } 
 
-    if(name.find("Yolo") == 0) {
+        return r;
+    }*/
+
+    /*if(name.find("Yolo") == 0) {
 
         int classes_temp = readBUF<int>(buf);
         int num_temp = readBUF<int>(buf);
@@ -792,7 +810,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         int nms_kind_temp = readBUF<int>(buf);
         int new_coords_temp = readBUF<int>(buf);
 
-       YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp);  
+       YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp);
 
 
 
@@ -816,8 +834,8 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
 
         yolos[n_yolos++] = r;
         return r;
-    } 
+    }*/
-    if(name.find("Upsample") == 0) {
+    /*if(name.find("Upsample") == 0) {
         int strideTemp = readBUF<int>(buf);
         UpsampleRT* r = new UpsampleRT(strideTemp);
         r->c = readBUF<int>(buf);
@@ -825,9 +843,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Route") == 0) {
+    /*if(name.find("Route") == 0) {
         int groupsTemp = readBUF<int>(buf);
         int group_idTemp = readBUF<int>(buf);
         RouteRT* r = new RouteRT(groupsTemp, group_idTemp);
@@ -839,9 +857,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->w = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    if(name.find("Deformable") == 0) {
+    /*if(name.find("Deformable") == 0) {
         int chuck_dimTemp = readBUF<int>(buf);
         int khTemp = readBUF<int>(buf);
         int kwTemp = readBUF<int>(buf);
@@ -892,10 +910,10 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         free(aus);
         assert(buf == bufCheck + serialLength);
         return r;
-    } 
+    }*/
 
-    FatalError("Cant deserialize Plugin");
+//     FatalError("Cant deserialize Plugin");
-    return NULL;
+//     return NULL;
-}
+// }
 
 }}

src/Yolo3Detection.cpp

@@ -1,5 +1,6 @@
+#include "yoloContainer.h"
 #include "Yolo3Detection.h"
-
+#include "pluginsRT/YoloRT.h"
 
 namespace tk { namespace dnn {
 
@@ -11,15 +12,15 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
 
     nBatches = n_batches;
     confThreshold = conf_thresh;
-    tk::dnn::dataDim_t idim = netRT->input_dim;    
+    tk::dnn::dataDim_t idim = netRT->input_dim;
     idim.n = nBatches;
 
-    if(netRT->pluginFactory->n_yolos < 2 ) {
+    if(yoloContainer.n_yolos < 2 ) {
         FatalError("this is not yolo3");
     }
 
-    for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
+    for(int i=0; i<yoloContainer.n_yolos; i++) {
-        YoloRT *yRT = netRT->pluginFactory->yolos[i];
+        YoloRT *yRT = yoloContainer.yolos[i];
         classes = yRT->classes;
         num = yRT->num;
         nMasks = yRT->n_masks;
@@ -43,7 +44,7 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
 #endif
     checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*idim.tot()));
 
-    // class colors precompute    
+    // class colors precompute
     for(int c=0; c<classes; c++) {
         int offset = c*123457 % classes;
         float r = getColor(2, offset, classes);
@@ -54,7 +55,7 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
 
     classesNames = getYoloLayer()->classesNames;
     return true;
-} 
+}
 
 void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
 #ifdef OPENCV_CUDACONTRIB
@@ -62,7 +63,7 @@ void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
     orig_img = cv::cuda::GpuMat(frame);
     cv::cuda::resize(orig_img, img_resized, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
 
-    img_resized.convertTo(imagePreproc, CV_32FC3, 1/255.0); 
+    img_resized.convertTo(imagePreproc, CV_32FC3, 1/255.0);
 
     //split channels
     cv::cuda::split(imagePreproc,bgr);//split source
@@ -76,7 +77,7 @@ void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
     }
 #else
     cv::resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
-    frame.convertTo(imagePreproc, CV_32FC3, 1/255.0); 
+    frame.convertTo(imagePreproc, CV_32FC3, 1/255.0);
 
     //split channels
     cv::split(imagePreproc,bgr);//split source
@@ -85,7 +86,7 @@ void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
     for(int i=0; i<netRT->input_dim.c; i++) {
         int idx = i*imagePreproc.rows*imagePreproc.cols;
         int ch = netRT->input_dim.c-1 -i;
-        memcpy((void*)&input[idx + netRT->input_dim.tot()*bi], (void*)bgr[ch].data, imagePreproc.rows*imagePreproc.cols*sizeof(dnnType));     
+        memcpy((void*)&input[idx + netRT->input_dim.tot()*bi], (void*)bgr[ch].data, imagePreproc.rows*imagePreproc.cols*sizeof(dnnType));
     }
     checkCuda(cudaMemcpyAsync(input_d + netRT->input_dim.tot()*bi, input + netRT->input_dim.tot()*bi, netRT->input_dim.tot()*sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
 #endif
@@ -95,8 +96,8 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
 
     //get yolo outputs
     std::vector<float *> rt_out;
-    //dnnType *rt_out[netRT->pluginFactory->n_yolos];
+    //dnnType *rt_out[yoloContainer.n_yolos];
-    for(int i=0; i<netRT->pluginFactory->n_yolos; i++)
+    for(int i=0; i<yoloContainer.n_yolos; i++)
         rt_out.push_back((dnnType*)netRT->buffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi);
 
     float x_ratio =  float(originalSize[bi].width) / float(netRT->input_dim.w);
@@ -104,7 +105,7 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
 
     // compute dets
     nDets = 0;
-    for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
+    for(int i=0; i<yoloContainer.n_yolos; i++) {
         yolo[i]->dstData = rt_out[i];
         yolo[i]->computeDetections(dets, nDets, netRT->input_dim.w, netRT->input_dim.h, confThreshold, yolo[i]->new_coords);
     }
@@ -124,7 +125,7 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
         x1 = x_ratio*x1;
         y0 = y_ratio*y0;
         y1 = y_ratio*y1;
-        
+
         for(int c=0; c<classes; c++) {
             if(dets[j].prob[c] >= confThreshold) {
                 int obj_class = c;
@@ -140,7 +141,7 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
 
                 // FIXME: this shuld be useless
                 // if(mAP)
-                //     for(int c=0; c<classes; c++) 
+                //     for(int c=0; c<classes; c++)
                 //         res.probs.push_back(dets[j].prob[c]);
 
                 detected.push_back(res);
@@ -155,7 +156,7 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
 tk::dnn::Yolo* Yolo3Detection::getYoloLayer(int n) {
     if(n<3)
         return yolo[n];
-    else 
+    else
         return nullptr;
 }
 

src/kernels/activation_leaky.cu

@@ -5,7 +5,7 @@ void activation_leaky(dnnType *input, dnnType *output, int size, float slope) {
 
     int i = blockDim.x*blockIdx.x + threadIdx.x;
 
-    if(i<size) {    
+    if(i<size) {
         if (input[i]>0)
             output[i] = input[i];
         else
@@ -21,7 +21,7 @@ void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, float
 {
     int blocks = (size+255)/256;
     int threads = 256;
-    
+
     activation_leaky<<<blocks, threads, 0, stream>>>(srcData, dstData, size, slope);
 }
 

src/kernels/activation_logistic.cu

@@ -1,11 +1,22 @@
+#include <vector>
+
 #include "kernels.h"
+#include "pluginsRT/ActivationLogisticRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ActivationLogisticRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ActivationLogisticRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ActivationLogisticRTCreator);
+}}
 
 __global__
 void activation_logistic(dnnType *input, dnnType *output, int size) {
 
     int i = blockDim.x*blockIdx.x + threadIdx.x;
 
-    if(i<size) {    
+    if(i<size) {
         output[i] =  1.0f/(1.0f + exp(-input[i]));;
     }
  }
@@ -18,8 +29,6 @@ void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, cud
 {
     int blocks = (size+255)/256;
     int threads = 256;
-    
+
     activation_logistic<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
-}
+}
-
-

src/kernels/activation_mish.cu

@@ -1,41 +1,58 @@
-#include "kernels.h"
 #include <math.h>
+#include "kernels.h"
+#include "pluginsRT/ActivationMishRT.h"
 
 #define MISH_THRESHOLD 20
 
-__device__ 
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ActivationMishRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ActivationMishRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ActivationMishRTCreator);
+}}
+
+__device__
 float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);}
 
-__device__ 
+__device__
 float softplus_kernel(float x, float threshold = 20) {
     if (x > threshold) return x;                // too large
     else if (x < -threshold) return expf(x);    // too small
     return logf(expf(x) + 1);
 }
 
+__device__
+float mish_yashas(float x) {
+    float e = __expf(x);
+    if (x <= -18.0f)
+        return x * e;
 
+    float n = e * e + 2 * e;
+    if (x <= -5.0f)
+        return x * __fdividef(n, n + 2);
 
-__device__ 
+    return x - 2 * __fdividef(x, n + 2);
-float mish_yashas(float x) { 
+}
-    float e = __expf(x); 
+
-    if (x <= -18.0f) 
+__device__ float mish_yashas2(float x)
-        return x * e; 
+{
- 
+    float e = __expf(x);
-    float n = e * e + 2 * e; 
+    float n = e * e + 2 * e;
-    if (x <= -5.0f) 
+    if (x <= -0.6f)
-        return x * __fdividef(n, n + 2); 
+        return x * __fdividef(n, n + 2);
- 
+
-    return x - 2 * __fdividef(x, n + 2); 
+    return x - 2 * __fdividef(x, n + 2);
-} 
+}
 
 // https://github.com/digantamisra98/Mish
 // https://github.com/AlexeyAB/darknet/blob/master/src/activation_kernels.cu
 __global__
 void activation_mish(dnnType *input, dnnType *output, int size) {
     int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
-    if (i < size) 
+    if (i < size)
-        // output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD));    
+        // output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD));
-        output[i] = mish_yashas(input[i]);
+        output[i] = mish_yashas2(input[i]);
 }
 
 /**
@@ -45,6 +62,6 @@ void activationMishForward(dnnType* srcData, dnnType* dstData, int size, cudaStr
 {
     int blocks = (size+255)/256;
     int threads = 256;
-    
+
     activation_mish<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
 }

src/kernels/activation_relu_ceiling.cu

@@ -1,11 +1,20 @@
 #include "kernels.h"
+#include "pluginsRT/ActivationReLUCeilingRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ActivationReLUCeilingCreator::mFC{};
+std::vector<nvinfer1::PluginField> ActivationReLUCeilingCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ActivationReLUCeilingCreator);
+}}
 
 __global__
 void activation_relu_ceiling(dnnType *input, dnnType *output, int size, const float ceiling) {
 
     int i = blockDim.x*blockIdx.x + threadIdx.x;
 
-    if(i<size) {    
+    if(i<size) {
         if (input[i]>0)
         {
             if (input[i]>ceiling)
@@ -26,7 +35,7 @@ void activationReLUCeilingForward(dnnType* srcData, dnnType* dstData, int size,
 {
     int blocks = (size+255)/256;
     int threads = 256;
-    
+
     activation_relu_ceiling<<<blocks, threads, 0, stream>>>(srcData, dstData, size, ceiling);
 }
 

src/kernels/activation_sigmoid.cu

@@ -1,6 +1,14 @@
-#include "kernels.h"
 #include <math.h>
+#include "kernels.h"
+#include "pluginsRT/ActivationSigmoidRT.h"
 
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ActivationSigmoidRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ActivationSigmoidRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ActivationSigmoidRTCreator);
+}}
 
 __global__
 void activation_sigmoid(dnnType *input, dnnType *output, int size) {
@@ -18,6 +26,6 @@ void activationSIGMOIDForward(dnnType* srcData, dnnType* dstData, int size, cuda
 {
     int blocks = (size+255)/256;
     int threads = 256;
-    
+
     activation_sigmoid<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
 }

src/kernels/deformable_conv.cu

@@ -2,10 +2,12 @@
 #include <algorithm>
 #include <cstring>
 #include <string>
-#include <iostream> 
+#include <iostream>
-#include "kernels.h"
 #include <errno.h>
 
+#include "kernels.h"
+#include "pluginsRT/DeformableConvRT.h"
+
 #define CUDA_KERNEL_LOOP(i, n)                          \
   for (int i = blockIdx.x * blockDim.x + threadIdx.x;   \
       i < (n);                                          \
@@ -17,6 +19,13 @@ inline int GET_BLOCKS(const int N)
   return (N + CUDA_NUM_THREADS - 1) / CUDA_NUM_THREADS;
 }
 
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection DeformableConvRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> DeformableConvRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(DeformableConvRTCreator);
+}}
 
 __device__ __host__  float dmcn_im2col_bilinear(const float *bottom_data, const int data_width,
   const int height, const int width, float h, float w) {
@@ -50,7 +59,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n,
   {
     //If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ).
     const int ind_on_w = index / width_col;
-    const int ind_on_w_on_h = ind_on_w / height_col; 
+    const int ind_on_w_on_h = ind_on_w / height_col;
     const int kk = 3 * 3;
     // index index of output matrix
     const int w_col = index % width_col;
@@ -83,10 +92,10 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n,
         const int iter_member = (i * 3 + j);
         // const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col;
         const int data_offset_h_ptr = first_member + s_col2 * iter_member;
-        
+
         // const int data_offset_w_ptr = ((2 * (i * kernel_w + j) + 1) * height_col + h_col) * width_col + w_col;
         const int data_offset_w_ptr = s_col + first_member + s_col2 * iter_member;
-        
+
         // const int data_mask_hw_ptr = ((i * kernel_w + j) * height_col + h_col) * width_col + w_col;
         const int data_mask_hw_ptr = first_member + s_col * iter_member;
 
@@ -127,7 +136,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int
   {
     //If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ).
     const int ind_on_w = index / width_col;
-    const int ind_on_w_on_h = ind_on_w / height_col; 
+    const int ind_on_w_on_h = ind_on_w / height_col;
     const int kk = kernel_h * kernel_w;
     // index index of output matrix
     const int w_col = index % width_col;
@@ -153,7 +162,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int
     const float *data_offset_ptr = data_offset + add_ptr + add_ptr;
 
     const float *data_mask_ptr = data_mask + add_ptr;
-    
+
     #pragma unroll
     for (int i = 0; i < kernel_h; ++i) {
       #pragma unroll
@@ -161,10 +170,10 @@ __global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int
         const int iter_member = (i * kernel_w + j);
         // const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col;
         const int data_offset_h_ptr = first_member + s_col2 * iter_member;
-        
+
         // const int data_offset_w_ptr = ((2 * (i * kernel_w + j) + 1) * height_col + h_col) * width_col + w_col;
         const int data_offset_w_ptr = s_col + first_member + s_col2 * iter_member;
-        
+
         // const int data_mask_hw_ptr = ((i * kernel_w + j) * height_col + h_col) * width_col + w_col;
         const int data_mask_hw_ptr = first_member + s_col * iter_member;
 
@@ -193,7 +202,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int
 
 void modulatedDeformableIm2colCuda(cudaStream_t stream,
   const float* data_im, const float* data_offset, const float* data_mask,
-  const int batch_size, const int channels, const int height_im, const int width_im, 
+  const int batch_size, const int channels, const int height_im, const int width_im,
   const int height_col, const int width_col,
   const int deformable_group, float* data_col) {
   // num_axes should be smaller than block size
@@ -202,9 +211,9 @@ void modulatedDeformableIm2colCuda(cudaStream_t stream,
   modulated_deformable_im2col_gpu_kernel
       <<<GET_BLOCKS(num_kernels), CUDA_NUM_THREADS,
           0, stream>>>(
-      num_kernels, data_im, data_offset, data_mask, height_im, width_im, 
+      num_kernels, data_im, data_offset, data_mask, height_im, width_im,
       batch_size, channels, deformable_group, height_col, width_col, data_col);
-  
+
   cudaError_t err = cudaGetLastError();
   if (err != cudaSuccess)
     FatalError("error in modulatedDeformableIm2colCuda: " + std::string(cudaGetErrorString(err)) + "\n");
@@ -212,9 +221,9 @@ void modulatedDeformableIm2colCuda(cudaStream_t stream,
 
 void modulatedDeformableIm2colCudaGeneralVersion(cudaStream_t stream,
   const float* data_im, const float* data_offset, const float* data_mask,
-  const int batch_size, const int channels, const int height_im, const int width_im, 
+  const int batch_size, const int channels, const int height_im, const int width_im,
   const int height_col, const int width_col, const int kernel_h, const int kenerl_w,
-  const int pad_h, const int pad_w, const int stride_h, const int stride_w, 
+  const int pad_h, const int pad_w, const int stride_h, const int stride_w,
   const int dilation_h, const int dilation_w,
   const int deformable_group, float* data_col) {
   // num_axes should be smaller than block size
@@ -226,13 +235,13 @@ void modulatedDeformableIm2colCudaGeneralVersion(cudaStream_t stream,
       num_kernels, data_im, data_offset, data_mask, height_im, width_im, kernel_h, kenerl_w,
       pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w, channel_per_deformable_group,
       batch_size, channels, deformable_group, height_col, width_col, data_col);
-  
+
   cudaError_t err = cudaGetLastError();
   if (err != cudaSuccess)
     FatalError("error in modulatedDeformableIm2colCudaGeneralVersion: " + std::string(cudaGetErrorString(err)) + "\n");
 }
 
-void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, 
+void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
                          float *input, float *weight,
                          float *bias, float *ones,
                          float *offset, float *mask,
@@ -242,11 +251,11 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
                          const int pad_h, const int pad_w,
                          const int dilation_h, const int dilation_w,
                          const int deformable_group, const int batch_id,
-                         const int in_n, const int in_c, const int in_h, const int in_w, 
+                         const int in_n, const int in_c, const int in_h, const int in_w,
                          const int out_n, const int out_c, const int out_h, const int out_w,
                          const int chunk_dim, cudaStream_t stream)
-{  
+{
-  // stat and handle have be moved out to preserve 2 - 6 milliseconds every 100. 
+  // stat and handle have be moved out to preserve 2 - 6 milliseconds every 100.
   const int batch = batch_id;
   const int channels = in_c;
   const int height = in_h;
@@ -256,22 +265,22 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
 
   const int height_out = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1;
   const int width_out = (width + 2 * pad_w - (dilation_w * (kernel_w - 1) + 1)) / stride_w + 1;
-  
+
   long m = channels_out;
   long n = height_out * width_out;
   long k = 1;
   float alpha = 1.0;
   float beta = 0.0;
-  
+
-  stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, 
+  stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N,
-              n, m, k, &alpha, 
+              n, m, k, &alpha,
-              ones, k, bias, k, 
+              ones, k, bias, k,
               &beta, output + batch * out_c * out_h * out_w, n);
   if (stat != CUBLAS_STATUS_SUCCESS)
     FatalError("CUBLAS initialization failed\n");
 
   modulatedDeformableIm2colCuda(stream,
-                                input + batch * channels * height * width,  
+                                input + batch * channels * height * width,
                                 offset,// + b * 2 * int((float)chunk_dim / batch),
                                 mask,// + b * int((float)chunk_dim / batch),
                                     1, channels, height, width,
@@ -283,15 +292,15 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
   //                                   height_out, width_out, kernel_h, kernel_w,
   //                                   pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w,
   //                                   deformable_group, columns);
-  
+
   //(k * m)  x  (m * n)
   // Y = WC
   k = channels * kernel_h * kernel_w;
   beta = 1.0;
-  
+
-  stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, 
+  stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
-              n, m, k, &alpha, 
+              n, m, k, &alpha,
-              columns, n, weight, k, 
+              columns, n, weight, k,
               &beta, output + batch * out_c * out_h * out_w, n);
 
   if (stat != CUBLAS_STATUS_SUCCESS)

src/kernels/pooling.cu

@@ -1,4 +1,13 @@
 #include "kernels.h"
+#include "pluginsRT/MaxPoolingFixedSizeRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection MaxPoolFixedSizeRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> MaxPoolFixedSizeRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(MaxPoolFixedSizeRTCreator);
+}}
 
 __global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c, int stride_x, int stride_y, int size, int pad, float *input, float *output)
 {
@@ -39,14 +48,14 @@ __global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c
     output[out_index] = max;
 }
 
-void MaxPoolingForward(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride_x, int stride_y, int size, int padding, cudaStream_t stream) 
+void MaxPoolingForward(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride_x, int stride_y, int size, int padding, cudaStream_t stream)
 {
 
     int tot_size = n*c*h*w;
 
     int blocks = (tot_size+255)/256;
     int threads = 256;
-    
+
     forward_maxpool_layer_kernel<<<blocks, threads, 0, stream>>>(tot_size, h, w, c, stride_x, stride_y, size, padding, srcData, dstData);
 }
 

src/kernels/reorg.cu

@@ -1,4 +1,13 @@
 #include "kernels.h"
+#include "pluginsRT/ReorgRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ReorgRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ReorgRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ReorgRTCreator);
+}}
 
 __global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, int stride, int forward, float *out)
 {
@@ -35,14 +44,14 @@ __global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, in
 /**
     reorg function function
 */
-void reorgForward(dnnType* srcData, dnnType* dstData, 
+void reorgForward(dnnType* srcData, dnnType* dstData,
                   int n, int c, int h, int w, int stride, cudaStream_t stream) {
 
     int size = n*c*h*w;
 
     int blocks = (size+255)/256;
     int threads = 256;
-    
+
     reorg_kernel<<<blocks, threads, 0, stream>>>(size, srcData, w, h, c, n, stride, false, dstData);
 }
 

src/kernels/resize.cu

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

src/kernels/scaladd.cu

@@ -1,6 +1,7 @@
-#include "kernels.h"
 #include <math.h>
 
+#include "kernels.h"
+
 __global__ void scal_add_kernel(dnnType* dstData, int size, float alpha, float beta, int inc)
 {
     int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
@@ -11,6 +12,6 @@ void scalAdd(dnnType* dstData, int size, float alpha, float beta, int inc, cudaS
 {
     int blocks = (size+255)/256;
     int threads = 256;
-    
+
     scal_add_kernel<<<blocks, threads, 0, stream>>>(dstData, size, alpha, beta, inc);
 }

src/kernels/shortcut.cu

@@ -1,8 +1,18 @@
 #include "kernels.h"
 #include "assert.h"
 
-__global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stride, int sample, int batch, 
+#include "pluginsRT/ShortcutRT.h"
-                                int w1, int h1, int c1, dnnType *add, 
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection ShortcutRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ShortcutRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ShortcutRTCreator);
+}}
+
+__global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stride, int sample, int batch,
+                                int w1, int h1, int c1, dnnType *add,
                                 int w2, int h2, int c2, float s1, float s2, dnnType *out)
 {
     int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
@@ -21,8 +31,8 @@ __global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stri
     //out[out_index] += add[add_index];
 }
 
-__global__ void shortcut_mul_kernel(int size, int minw, int minh, int minc, int sample, int batch, 
+__global__ void shortcut_mul_kernel(int size, int minw, int minh, int minc, int sample, int batch,
-                                int w1, int h1, int c1, dnnType *mul, 
+                                int w1, int h1, int c1, dnnType *mul,
                                 int w2, int h2, int c2, float s1, float s2, dnnType *out)
 {
     int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
@@ -40,7 +50,7 @@ __global__ void shortcut_mul_kernel(int size, int minw, int minh, int minc, int
 }
 
 void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int w1, int s1,
-                                                         int n2, int c2, int h2, int w2, int s2, 
+                                                         int n2, int c2, int h2, int w2, int s2,
                                                          bool mul, cudaStream_t stream)
 {
     assert(n1 == n2);
@@ -60,10 +70,10 @@ void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1,
         int size = batch * minw * minh * minc;
         int blocks = (size+255)/256;
         int threads = 256;
-        
+
-        shortcut_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, stride, sample, batch, 
+        shortcut_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, stride, sample, batch,
-            w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);  
+            w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);
-    } 
+    }
     else{
         int minw = w1;
         int minh = h1;
@@ -74,7 +84,7 @@ void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1,
         int blocks = (size+255)/256;
         int threads = 256;
 
-        shortcut_mul_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, sample, batch, 
+        shortcut_mul_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, sample, batch,
-            w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);  
+            w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);
     }
 }

src/kernels/static_init.cu

@@ -0,0 +1,33 @@
+#include "pluginsRT/RegionRT.h"
+#include "pluginsRT/RouteRT.h"
+#include "pluginsRT/ReshapeRT.h"
+#include "pluginsRT/FlattenConcatRT.h"
+#include "pluginsRT/YoloRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection RegionRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> RegionRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(RegionRTCreator);
+
+nvinfer1::PluginFieldCollection RouteRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> RouteRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(RouteRTCreator);
+
+nvinfer1::PluginFieldCollection ReshapeRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> ReshapeRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(ReshapeRTCreator);
+
+nvinfer1::PluginFieldCollection FlattenConcatRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> FlattenConcatRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(FlattenConcatRTCreator);
+
+nvinfer1::PluginFieldCollection YoloRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> YoloRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(YoloRTCreator);
+}}

src/kernels/upsample.cu

@@ -1,4 +1,13 @@
 #include "kernels.h"
+#include "pluginsRT/UpsampleRT.h"
+
+// Static class fields initialization
+namespace tk { namespace dnn {
+nvinfer1::PluginFieldCollection UpsampleRTCreator::mFC{};
+std::vector<nvinfer1::PluginField> UpsampleRTCreator::mPluginAttributes;
+
+REGISTER_TENSORRT_PLUGIN(UpsampleRTCreator);
+}}
 
 __global__ void upsample_kernel(size_t N, dnnType *x, int w, int h, int c, int batch, int stride, int forward, float scale, dnnType *out)
 {
@@ -24,8 +33,8 @@ __global__ void upsample_kernel(size_t N, dnnType *x, int w, int h, int c, int b
     else atomicAdd(x+in_index, scale * out[out_index]);
 }
 
-void upsampleForward(dnnType* srcData, dnnType* dstData, 
+void upsampleForward(dnnType* srcData, dnnType* dstData,
-                     int n, int c, int h, int w, int s, int forward, float scale,                                   
+                     int n, int c, int h, int w, int s, int forward, float scale,
                      cudaStream_t stream) {
 
     int size = w*h*c*n*s*s;

src/yoloContainer.cpp

@@ -0,0 +1,41 @@
+#include "yoloContainer.h"
+#include "pluginsRT/YoloRT.h"
+
+namespace tk { namespace dnn {
+YoloContainer yoloContainer;
+
+nvinfer1::IPluginV2* YoloRTCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength) noexcept {
+  const char * buf = reinterpret_cast<const char*>(serialData),*bufCheck = buf;
+  int classes_temp = tk::dnn::readBUF<int>(buf);
+  int num_temp = tk::dnn::readBUF<int>(buf);
+  int n_masks_temp = tk::dnn::readBUF<int>(buf);
+  float scale_xy_temp = tk::dnn::readBUF<float>(buf);
+  float nms_thresh_temp = tk::dnn::readBUF<float>(buf);
+  int nms_kind_temp = tk::dnn::readBUF<int>(buf);
+  int new_coords_temp = tk::dnn::readBUF<int>(buf);
+
+  YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp);
+
+  r->c = tk::dnn::readBUF<int>(buf);
+  r->h = tk::dnn::readBUF<int>(buf);
+  r->w = tk::dnn::readBUF<int>(buf);
+  for(int i=0; i<r->n_masks; i++)
+      r->mask[i] = tk::dnn::readBUF<dnnType>(buf);
+  for(int i=0; i<r->n_masks*2*r->num; i++)
+      r->bias[i] = tk::dnn::readBUF<dnnType>(buf);
+
+  // save classes names
+  r->classesNames.resize(r->classes);
+  for(int i=0; i<r->classes; i++) {
+          char tmp[YOLORT_CLASSNAME_W];
+    for(int j=0; j<YOLORT_CLASSNAME_W; j++)
+      tmp[j] = tk::dnn::readBUF<char>(buf);
+          r->classesNames[i] = std::string(tmp);
+  }
+  assert(buf == bufCheck + serialLength);
+
+  yoloContainer.yolos[yoloContainer.n_yolos++] = r;
+  return r;
+}
+
+}}

tests/mnist/test_mnistRT.cpp

@@ -13,9 +13,9 @@ const char *output_bin   = "mnist/output.bin";
 using namespace nvinfer1;
 
 // Logger for info/warning/errors
-class Logger : public ILogger			
+class Logger : public ILogger
 {
-	void log(Severity severity, const char* msg) override
+	void log(Severity severity, const char* msg) noexcept override
 	{
 		// suppress info-level messages
 		if (severity != Severity::kINFO)
@@ -39,20 +39,20 @@ int main() {
     tk::dnn::Activation l5(&net, CUDNN_ACTIVATION_RELU);
     tk::dnn::Dense      l6(&net, 10, d3_bin);
     tk::dnn::Softmax    l7(&net);
- 
+
     // Load input
     dnnType *data;
     dnnType *input_h;
     readBinaryFile(input_bin, dim.tot(), &input_h, &data);
 
     dim.print(); //print initial dimension
-    
+
     // Inference
     {
         TKDNN_TSTART
         data = net.infer(dim, data);
         TKDNN_TSTOP
-        dim.print();   
+        dim.print();
     }
 
     // Print real test
@@ -61,19 +61,19 @@ int main() {
     dnnType *out_h;
     readBinaryFile(output_bin, dim.tot(), &out_h, &out);
     std::cout<<"Diff: "<<checkResult(dim.tot(), out, data)<<"\n";
- 
+
 
     std::cout<<"\n==== TensorRT ====\n";
 	// create the builder
 	IBuilder* builder = nvinfer1::createInferBuilder(gLogger);
-	INetworkDefinition* network = builder->createNetwork();
+	INetworkDefinition* network = builder->createNetworkV2(0u);
 
     DataType dt = DataType::kFLOAT;
 	//  Create input of shape { 1, 1, 28, 28 } with name referenced by "data"
-	auto input = network->addInput("data", dt, DimsCHW{ 1, 28, 28});
+	auto input = network->addInput("data", dt, Dims3{ 1, 28, 28});
 	assert(input != nullptr);
 
-    tk::dnn::Conv2d *c0 = &l0; 
+    tk::dnn::Conv2d *c0 = &l0;
     Weights w { dt, c0->data_h, c0->inputs*c0->outputs*c0->kernelH*c0->kernelW};
     Weights b { dt, c0->bias_h, c0->outputs};
 	// Add a convolution layer with 20 outputs and a 5x5 filter.
@@ -86,7 +86,7 @@ int main() {
 	assert(pool1 != nullptr);
 	pool1->setStride(DimsHW{2, 2});
 
-    tk::dnn::Conv2d *c1 = &l2; 
+    tk::dnn::Conv2d *c1 = &l2;
     Weights w1 { dt, c1->data_h, c1->inputs*c1->outputs*c1->kernelH*c1->kernelW};
     Weights b1 { dt, c1->bias_h, c1->outputs};
 	// Add a second convolution layer with 50 outputs and a 5x5 filter.
@@ -99,7 +99,7 @@ int main() {
 	assert(pool2 != nullptr);
 	pool2->setStride(DimsHW{2, 2});
 
-    tk::dnn::Dense *d2 = &l4; 
+    tk::dnn::Dense *d2 = &l4;
     Weights w2 { dt, d2->data_h, d2->inputs*d2->outputs};
     Weights b2 { dt, d2->bias_h, d2->outputs};
 	// Add a fully connected layer with 500 outputs.
@@ -110,7 +110,7 @@ int main() {
 	auto relu1 = network->addActivation(*ip1->getOutput(0), ActivationType::kRELU);
 	assert(relu1 != nullptr);
 
-    tk::dnn::Dense *d3 = &l6; 
+    tk::dnn::Dense *d3 = &l6;
     Weights w3 { dt, d3->data_h, d3->inputs*d3->outputs};
     Weights b3 { dt, d3->bias_h, d3->outputs};
 	// Add a second fully connected layer with 20 outputs.
@@ -125,10 +125,21 @@ int main() {
 	network->markOutput(*prob->getOutput(0));
 
 	// Build the engine
-	builder->setMaxBatchSize(1);
+#if NV_TENSORRT_MAJOR >= 6
+  auto config = builder->createBuilderConfig();
+	config->setMaxWorkspaceSize(1 << 20);
+#else
 	builder->setMaxWorkspaceSize(1 << 20);
+#endif
+
+	builder->setMaxBatchSize(1);
+
+#if NV_TENSORRT_MAJOR >= 6
+        auto engine = builder->buildEngineWithConfig(*network, *config);
+#else
+        auto engine = builder->buildCudaEngine(*network);
+#endif
 
-	auto engine = builder->buildCudaEngine(*network);
 	// we don't need the network any more
 	network->destroy();
 
@@ -142,10 +153,10 @@ int main() {
 
 	// In order to bind the buffers, we need to know the names of the input and output tensors.
 	// note that indices are guaranteed to be less than IEngine::getNbBindings()
-	int inputIndex = engine->getBindingIndex("data"); 
+	int inputIndex = engine->getBindingIndex("data");
-    int outputIndex = engine->getBindingIndex("out");
+  int outputIndex = engine->getBindingIndex("out");
 
-    float output[10];
+  float output[10];
 	// create GPU buffers and a stream
 	checkCuda(cudaMalloc(&buffers[inputIndex], 28*28*sizeof(float)));
 	checkCuda(cudaMalloc(&buffers[outputIndex], 10*sizeof(float)));