tkDNN can now deserialize tensorrt-8 engine (both through test_* and trtexec)

but demo has issues in yolo::computeDetections

src/DarknetParser.cpp

@@ -17,8 +17,8 @@ namespace tk { namespace dnn {
         if(sep == std::string::npos)
             return false;
         
-        name   = line.substr(0, sep);   
-        value  = line.substr(sep+1, line.size() - (sep+1));   
+        name   = line.substr(0, sep);
+        value  = line.substr(sep+1, line.size() - (sep+1));
         return true;
     }
 
@@ -322,7 +322,7 @@ namespace tk { namespace dnn {
             // skip empty lines
             if(line.empty())
                 continue;
-            if(count > lineNo && count <=20){
+            if(count > lineNo && count <=lineNo+30){
                     divideNameAndValue(line,name,value);
                     if(name == "mask "){
                         maskTemp = fromStringToFloatVec(value,',');
@@ -353,9 +353,8 @@ namespace tk { namespace dnn {
                     if(name == "beta_nms"){
                         nmsThreshTemp = std::stof(value);
                     }
-
-                count++;
             }
+            count++;
         }
         mask = maskTemp;
         anchors = anchorsTemp;

src/Yolo.cpp

@@ -133,10 +133,11 @@ void correct_yolo_boxes(Yolo::detection *dets, int n, int w, int h, int netw, in
     }
 }
 
-int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh, int new_coords) {
+int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh, int newCoords) {
 
     if(predictions == nullptr)
         predictions = new dnnType[output_dim.tot()];
+    checkCuda(cudaDeviceSynchronize());
     checkCuda( cudaMemcpy(predictions, dstData, output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost));
 
     int lw = output_dim.w;
@@ -157,7 +158,7 @@ int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int net
             if(objectness <= thresh) continue;
             int box_index  = entry_index(0, n*lw*lh + i, 0, classes, input_dim, output_dim);
             
-            dets[count].bbox = get_yolo_box(predictions, bias_h, mask_h[n], box_index, col, row, lw, lh, netw, neth, lw*lh, new_coords);
+            dets[count].bbox = get_yolo_box(predictions, bias_h, mask_h[n], box_index, col, row, lw, lh, netw, neth, lw*lh, newCoords);
             dets[count].objectness = objectness;
             dets[count].classes = classes;
             for(j = 0; j < classes; ++j){

src/pluginsRT/ActivationLeakyRT.cpp

@@ -0,0 +1,138 @@
+#include <tkDNN/pluginsRT/ActivationLeakyRT.h>
+using namespace nvinfer1;
+
+
+std::vector<PluginField> ActivationLeakyRTPluginCreator::mPluginAttributes;
+PluginFieldCollection ActivationLeakyRTPluginCreator::mFC{};
+
+ActivationLeakyRT::ActivationLeakyRT(float s) {
+    slope = s;
+}
+
+ActivationLeakyRT::ActivationLeakyRT(const void *data, size_t length) {
+    std::cout << "DESERIALIZE LEAKYRT" << std::endl;
+    const char *buf = reinterpret_cast<const char *>(data), *bufCheck = buf;
+    slope = readBUF<float>(buf);
+    size = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+ActivationLeakyRT::~ActivationLeakyRT() {}
+
+int ActivationLeakyRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims ActivationLeakyRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return inputs[0];
+}
+
+void ActivationLeakyRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,
+                                            DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {
+    assert(type == DataType::kFLOAT && format == PluginFormat::kLINEAR);
+    size = 1;
+    for (int i = 0; i < outputDims[0].nbDims; i++)
+        size *= outputDims[0].d[i];
+}
+int ActivationLeakyRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+size_t ActivationLeakyRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int ActivationLeakyRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                               cudaStream_t stream) NOEXCEPT {
+    activationLEAKYForward(
+            (dnnType *) reinterpret_cast<const dnnType *>(inputs[0]),
+            reinterpret_cast<dnnType *>(outputs[0]), batchSize * size, slope,
+            stream);
+    return 0;
+
+}
+
+size_t ActivationLeakyRT::getSerializationSize() const NOEXCEPT {
+    return 1 * sizeof(int) + 1 * sizeof(float);
+}
+
+void ActivationLeakyRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char *>(buffer), *a = buf;
+    writeBUF(buf, size);
+    assert(buf == a + getSerializationSize());
+}
+
+bool ActivationLeakyRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR);
+}
+
+const char *ActivationLeakyRT::getPluginType() const NOEXCEPT {
+    return "ActivationLeakyRT_tkDNN";
+}
+
+const char *ActivationLeakyRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void ActivationLeakyRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+const char *ActivationLeakyRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+
+}
+
+void ActivationLeakyRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2* ActivationLeakyRT::clone() const NOEXCEPT {
+    auto *p = new ActivationLeakyRT(slope);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+ActivationLeakyRTPluginCreator::ActivationLeakyRTPluginCreator() {
+    mPluginAttributes.emplace_back(
+            PluginField("slope", nullptr, PluginFieldType::kFLOAT32, 1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void ActivationLeakyRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2* ActivationLeakyRTPluginCreator::deserializePlugin(const char *name, const void *serialData,size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new ActivationLeakyRT(serialData, serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char* ActivationLeakyRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2* ActivationLeakyRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    assert(fc->nbFields == 1);
+    assert(fields[0].type == PluginFieldType::kFLOAT32);
+    float slope = *(static_cast<const float *>(fields[0].data));
+    auto *pluginObj = new ActivationLeakyRT(slope);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char* ActivationLeakyRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "ActivationLeakyRT_tkDNN";
+}
+
+const char* ActivationLeakyRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection* ActivationLeakyRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+

src/pluginsRT/ActivationLogisticRT.cpp

@@ -0,0 +1,126 @@
+#include <tkDNN/pluginsRT/ActivationLogisticRT.h>
+using namespace nvinfer1;
+std::vector<PluginField> ActivationLogisticRTPluginCreator::mPluginAttributes;
+PluginFieldCollection ActivationLogisticRTPluginCreator::mFC{};
+
+ActivationLogisticRT::ActivationLogisticRT() {}
+
+ActivationLogisticRT::ActivationLogisticRT(const void *data, size_t length) {
+    const char *buf = reinterpret_cast<const char *>(data), *bufCheck = buf;
+    size = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+ActivationLogisticRT::~ActivationLogisticRT() {}
+
+int ActivationLogisticRT::getNbOutputs() const NOEXCEPT {
+return 1;
+}
+
+Dims ActivationLogisticRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return inputs[0];
+}
+
+void ActivationLogisticRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims,
+                                               int nbOutputs, DataType type, PluginFormat format,
+                                               int maxBatchSize) NOEXCEPT {
+    size = 1;
+    for (int i = 0; i < outputDims[0].nbDims; i++)
+        size *= outputDims[0].d[i];
+}
+
+int ActivationLogisticRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void ActivationLogisticRT::terminate() NOEXCEPT {}
+
+size_t ActivationLogisticRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int ActivationLogisticRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                                  cudaStream_t stream) NOEXCEPT {
+    activationLOGISTICForward((dnnType *) reinterpret_cast<const dnnType *>(inputs[0]),
+                              reinterpret_cast<dnnType *>(outputs[0]), batchSize * size, stream);
+    return 0;
+}
+
+size_t ActivationLogisticRT::getSerializationSize() const NOEXCEPT {
+    return 1 * sizeof(int);
+}
+
+void ActivationLogisticRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char *>(buffer);
+    writeBUF(buf, size);
+}
+
+const char* ActivationLogisticRT::getPluginType() const NOEXCEPT {
+    return "ActivationLogisticRT_tkDNN";
+}
+
+const char* ActivationLogisticRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void ActivationLogisticRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+const char* ActivationLogisticRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void ActivationLogisticRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+bool ActivationLogisticRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+    //todo assert
+}
+
+IPluginV2* ActivationLogisticRT::clone() const NOEXCEPT {
+    auto *p = new ActivationLogisticRT();
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+ActivationLogisticRTPluginCreator::ActivationLogisticRTPluginCreator() {
+    mPluginAttributes.clear();
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void ActivationLogisticRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2* ActivationLogisticRTPluginCreator::deserializePlugin(const char *name, const void *serialData,
+                                                                size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new ActivationLogisticRT(serialData, serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char* ActivationLogisticRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2* ActivationLogisticRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    auto *pluginObj = new ActivationLogisticRT();
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char* ActivationLogisticRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection* ActivationLogisticRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+const char *ActivationLogisticRTPluginCreator::getPluginName() const NOEXCEPT  {
+    return "ActivationLogisticRT_tkDNN";
+}

src/pluginsRT/ActivationMishRT.cpp

@@ -0,0 +1,121 @@
+//
+// Created by perseusdg on 9/4/21.
+//
+#include <tkDNN/pluginsRT/ActivationMishRT.h>
+using namespace  nvinfer1;
+std::vector<PluginField> ActivationMishRTPluginCreator::mPluginAttributes;
+PluginFieldCollection ActivationMishRTPluginCreator::mFC{};
+
+ActivationMishRT::ActivationMishRT() {
+
+}
+
+ActivationMishRT::~ActivationMishRT() {
+
+}
+
+ActivationMishRT::ActivationMishRT(const void *data, size_t length) {
+    const char *buf = reinterpret_cast<const char *>(data), *bufCheck = buf;
+    size = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+int ActivationMishRT::getNbOutputs() const NOEXCEPT  { return 1; }
+
+Dims ActivationMishRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT  { return inputs[0]; }
+
+void ActivationMishRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,
+                         PluginFormat format, int maxBatchSize) NOEXCEPT {
+assert(format == PluginFormat::kLINEAR);
+size = 1;
+for (int i = 0; i < outputDims[0].nbDims; i++)
+size *= outputDims[0].d[i];
+}
+
+int ActivationMishRT::initialize() NOEXCEPT  { return 0; }
+
+void ActivationMishRT::terminate() NOEXCEPT {}
+
+size_t ActivationMishRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT  { return 0; }
+
+int ActivationMishRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+            cudaStream_t stream) NOEXCEPT {
+    activationMishForward((dnnType *) reinterpret_cast<const dnnType *>(inputs[0]),
+    reinterpret_cast<dnnType *>(outputs[0]), batchSize * size, stream);
+    return 0;
+}
+size_t ActivationMishRT::getSerializationSize() const NOEXCEPT  {
+    return 1 * sizeof(int);
+}
+
+void ActivationMishRT::serialize(void *buffer) const NOEXCEPT  {
+    char *buf = reinterpret_cast<char *>(buffer), *a = buf;
+    writeBUF(buf, size);
+    assert(buf == a + getSerializationSize());
+}
+
+const char* ActivationMishRT::getPluginType() const NOEXCEPT  {
+    return "ActivationMishRT_tkDNN";
+}
+
+const char *ActivationMishRT::getPluginVersion() const NOEXCEPT  {
+    return "1";
+}
+
+bool ActivationMishRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+}
+
+const char *ActivationMishRT::getPluginNamespace() const NOEXCEPT  {
+    return mPluginNamespace.c_str();
+}
+
+void ActivationMishRT::setPluginNamespace(const char *plguinNamespace) NOEXCEPT  {
+    mPluginNamespace = plguinNamespace;
+}
+
+IPluginV2 *ActivationMishRT::clone() const NOEXCEPT {
+    auto *p = new ActivationMishRT();
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+ActivationMishRTPluginCreator::ActivationMishRTPluginCreator() {
+    mPluginAttributes.clear();
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void ActivationMishRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *ActivationMishRTPluginCreator::getPluginNamespace() const NOEXCEPT  {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *ActivationMishRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT  {
+    auto *pluginObj = new ActivationMishRT(serialData, serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *ActivationMishRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT  {
+    const PluginField *fields = fc->fields;
+    auto *pluginObj = new ActivationMishRT();
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *ActivationMishRTPluginCreator::getPluginName() const NOEXCEPT  {
+    return "ActivationMishRT_tkDNN";
+}
+
+const char *ActivationMishRTPluginCreator::getPluginVersion() const NOEXCEPT{
+    return "1";
+}
+
+const PluginFieldCollection *ActivationMishRTPluginCreator::getFieldNames() NOEXCEPT  {
+    return &mFC;
+}
+

src/pluginsRT/ActivationReLUCeilingRT.cpp

@@ -0,0 +1,129 @@
+#include <tkDNN/pluginsRT/ActivationReLUCeilingRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> ActivationReLUCeilingPluginCreator::mPluginAttributes;
+PluginFieldCollection ActivationReLUCeilingPluginCreator::mFC{};
+
+ActivationReLUCeiling::ActivationReLUCeiling(const float ceiling) {
+    this->ceiling = ceiling;
+}
+
+ActivationReLUCeiling::~ActivationReLUCeiling() {
+
+}
+
+ActivationReLUCeiling::ActivationReLUCeiling(const void *data, size_t length) {
+    const char *buf = reinterpret_cast<const char *>(data), *bufCheck = buf;
+    ceiling = readBUF<float>(buf);
+    size = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+int ActivationReLUCeiling::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims ActivationReLUCeiling::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT { return inputs[0]; }
+
+void ActivationReLUCeiling::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {
+    assert(type == DataType::kFLOAT && format == PluginFormat::kLINEAR);
+    size = 1;
+    for (int i = 0; i < outputDims[0].nbDims; i++)
+        size *= outputDims[0].d[i];
+}
+
+int ActivationReLUCeiling::initialize() NOEXCEPT  { return 0; }
+
+void ActivationReLUCeiling::terminate() NOEXCEPT {}
+
+size_t ActivationReLUCeiling::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int ActivationReLUCeiling::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT  {
+    activationReLUCeilingForward((dnnType *) reinterpret_cast<const dnnType *>(inputs[0]),
+    reinterpret_cast<dnnType *>(outputs[0]), batchSize * size, ceiling, stream);
+    return 0;
+}
+
+size_t ActivationReLUCeiling::getSerializationSize() const NOEXCEPT {
+    return 1 * sizeof(int) + 1 * sizeof(float);
+}
+
+void ActivationReLUCeiling::serialize(void *buffer) const NOEXCEPT  {
+    char *buf = reinterpret_cast<char *>(buffer), *a = buf;
+    writeBUF(buf, ceiling);
+    writeBUF(buf, size);
+    assert(buf = a + getSerializationSize());
+}
+
+IPluginV2 *ActivationReLUCeiling::clone() const NOEXCEPT {
+    auto *p = new ActivationReLUCeiling(ceiling);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+bool ActivationReLUCeiling::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT  {
+    return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR);
+}
+
+void ActivationReLUCeiling::destroy() NOEXCEPT { delete this; }
+
+const char *ActivationReLUCeiling::getPluginType() const NOEXCEPT {
+    return "ActivationReLUCeilingRT_tkDNN";
+}
+
+const char *ActivationReLUCeiling::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const char *ActivationReLUCeiling::getPluginNamespace() const NOEXCEPT  {
+    return mPluginNamespace.c_str();
+}
+
+void ActivationReLUCeiling::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+ActivationReLUCeilingPluginCreator::ActivationReLUCeilingPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("ceiling", nullptr, PluginFieldType::kFLOAT32, 1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void ActivationReLUCeilingPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT  {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *ActivationReLUCeilingPluginCreator::getPluginNamespace() const NOEXCEPT  {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *ActivationReLUCeilingPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new ActivationReLUCeiling(serialData, serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *ActivationReLUCeilingPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT  {
+    const PluginField *fields = fc->fields;
+    float ceiling = *(static_cast<const float *>(fields[0].data));
+    auto *pluginObj = new ActivationReLUCeiling(ceiling);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *ActivationReLUCeilingPluginCreator::getPluginName() const NOEXCEPT  {
+    return "ActivationReLUCeilingRT_tkDNN";
+}
+
+const char *ActivationReLUCeilingPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *ActivationReLUCeilingPluginCreator::getFieldNames() NOEXCEPT  {
+    return &mFC;
+}
+
+
+

src/pluginsRT/DeformableConvRT.cpp

@@ -0,0 +1,331 @@
+#include <tkDNN/pluginsRT/DeformableConvRT.h>
+using namespace nvinfer1;
+using namespace tk::dnn;
+
+std::vector<PluginField> DeformableConvRTPluginCreator::mPluginAttributes;
+PluginFieldCollection DeformableConvRTPluginCreator::mFC{};
+
+
+DeformableConvRT::DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw, int deformableGroup,
+                                   int i_n, int i_c, int i_h, int i_w, int o_n, int o_c, int o_h, int o_w,
+                                   tk::dnn::DeformConv2d *deformable) {
+    this->chunk_dim = chunk_dim;
+    this->kh = kh;
+    this->kw = kw;
+    this->sh = sh;
+    this->sw = sw;
+    this->ph = ph;
+    this->pw = pw;
+    this->deformableGroup = deformableGroup;
+    this->i_n = i_n;
+    this->i_c = i_c;
+    this->i_h = i_h;
+    this->i_w = i_w;
+    this->o_n = o_n;
+    this->o_c = o_c;
+    this->o_h = o_h;
+    this->o_w = o_w;
+    this->defRT = deformable;
+
+    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)));
+    checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType)));
+    checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType)));
+    checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType)));
+    if(deformable != nullptr) {
+        checkCuda( cudaMemcpy(data_d, deformable->data_d, sizeof(dnnType)*i_c * o_c * kh * kw * 1, cudaMemcpyDeviceToDevice) );
+        checkCuda( cudaMemcpy(bias2_d, deformable->bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToDevice) );
+        checkCuda( cudaMemcpy(ones_d1, deformable->ones_d1, sizeof(dnnType)*height_ones*width_ones, cudaMemcpyDeviceToDevice) );
+        checkCuda( cudaMemcpy(offset, deformable->offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToDevice) );
+        checkCuda( cudaMemcpy(mask, deformable->mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToDevice) );
+        checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) );
+    }
+    stat = cublasCreate(&handle);
+    if (stat != CUBLAS_STATUS_SUCCESS)
+        FatalError("CUBLAS initialization failed\n");
+
+}
+
+DeformableConvRT::~DeformableConvRT() {
+    checkCuda( cudaFree(data_d) );
+    checkCuda( cudaFree(bias2_d) );
+    checkCuda( cudaFree(ones_d1) );
+    checkCuda( cudaFree(offset) );
+    checkCuda( cudaFree(mask) );
+    checkCuda( cudaFree(ones_d2) );
+    cublasDestroy(handle);
+}
+
+DeformableConvRT::DeformableConvRT(const void *data, size_t length) {
+    const char* buf = reinterpret_cast<const char*>(data),*bufCheck = buf;
+    chunk_dim = readBUF<int>(buf);
+    kh = readBUF<int>(buf);
+    kw = readBUF<int>(buf);
+    sh = readBUF<int>(buf);
+    sw = readBUF<int>(buf);
+    ph = readBUF<int>(buf);
+    pw = readBUF<int>(buf);
+    deformableGroup = readBUF<int>(buf);
+    i_n = readBUF<int>(buf);
+    i_c = readBUF<int>(buf);
+    i_h = readBUF<int>(buf);
+    i_w = readBUF<int>(buf);
+    o_n = readBUF<int>(buf);
+    o_c = readBUF<int>(buf);
+    o_h = readBUF<int>(buf);
+    o_w = readBUF<int>(buf);
+    dnnType *aus = new dnnType[chunk_dim*2];
+    for(int i=0;i<chunk_dim*2;i++)
+        aus[i] = readBUF<dnnType>(buf);
+    checkCuda(cudaMemcpy(offset,aus,sizeof(dnnType)*2*chunk_dim,cudaMemcpyHostToDevice));
+    free(aus);
+
+    aus = new dnnType[chunk_dim];
+    for(int i=0;i<chunk_dim;i++)
+        aus[i] = readBUF<dnnType>(buf);
+    checkCuda(cudaMemcpy(mask,aus,sizeof(dnnType)*chunk_dim,cudaMemcpyHostToDevice));
+    free(aus);
+
+    aus = new dnnType[i_c*o_c*kh*kw*1];
+    for(int i=0;i<(i_c*o_c*kh*kw*1);i++)
+        aus[i] = readBUF<dnnType>(buf);
+    checkCuda(cudaMemcpy(data_d,aus,sizeof(dnnType)*(i_c*o_c*kh*kw*1),cudaMemcpyHostToDevice));
+    free(aus);
+
+    aus = new dnnType[o_c];
+    for(int i=0; i < o_c; i++)
+        aus[i] = readBUF<dnnType>(buf);
+    checkCuda( cudaMemcpy(bias2_d, aus, sizeof(dnnType)*o_c, cudaMemcpyHostToDevice) );
+    free(aus);
+
+    aus = new dnnType[height_ones * width_ones];
+    for(int i=0; i<height_ones * width_ones; i++)
+        aus[i] = readBUF<dnnType>(buf);
+    checkCuda( cudaMemcpy(ones_d1, aus, sizeof(dnnType)*height_ones * width_ones, cudaMemcpyHostToDevice) );
+    free(aus);
+
+    aus = new dnnType[dim_ones];
+    for(int i=0; i<dim_ones; i++)
+        aus[i] = readBUF<dnnType>(buf);
+    checkCuda( cudaMemcpy(ones_d2, aus, sizeof(dnnType)*dim_ones, cudaMemcpyHostToDevice) );
+    free(aus);
+
+    assert(buf == bufCheck + length);
+
+}
+
+int DeformableConvRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims DeformableConvRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return Dims3{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w};
+}
+
+void DeformableConvRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {}
+
+int DeformableConvRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void DeformableConvRT::terminate() NOEXCEPT {}
+
+size_t DeformableConvRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {return 0;}
+
+int DeformableConvRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                              cudaStream_t stream) NOEXCEPT {
+    dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+    dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
+
+    // split conv2d outputs into offset to mask
+    for(int b=0; b<batchSize; b++) {
+        checkCuda(cudaMemcpy(offset, output_conv + b * 3 * chunk_dim, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+        checkCuda(cudaMemcpy(mask, output_conv + b * 3 * chunk_dim + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+        // kernel sigmoid
+        activationSIGMOIDForward(mask, mask, chunk_dim);
+        // deformable convolution
+        dcnV2CudaForward(stat, handle,
+                         srcData, data_d,
+                         bias2_d, ones_d1,
+                         offset, mask,
+                         reinterpret_cast<dnnType*>(outputs[0]), ones_d2,
+                         kh, kw,
+                         sh, sw,
+                         ph, pw,
+                         1, 1,
+                         deformableGroup, b,
+                         i_n, i_c, i_h, i_w,
+                         o_n, o_c, o_h, o_w,
+                         chunk_dim);
+    }
+    return 0;
+}
+
+size_t DeformableConvRT::getSerializationSize() const NOEXCEPT {
+    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);
+}
+
+void DeformableConvRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    writeBUF(buf, chunk_dim);
+    writeBUF(buf, kh);
+    writeBUF(buf, kw);
+    writeBUF(buf, sh);
+    writeBUF(buf, sw);
+    writeBUF(buf, ph);
+    writeBUF(buf, pw);
+    writeBUF(buf, deformableGroup);
+    writeBUF(buf, i_n);
+    writeBUF(buf, i_c);
+    writeBUF(buf, i_h);
+    writeBUF(buf, i_w);
+    writeBUF(buf, o_n);
+    writeBUF(buf, o_c);
+    writeBUF(buf, o_h);
+    writeBUF(buf, o_w);
+    dnnType *aus = new dnnType[chunk_dim*2];
+    checkCuda( cudaMemcpy(aus, offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToHost) );
+    for(int i=0; i<chunk_dim*2; i++)
+        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++)
+        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++)
+        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++)
+        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++)
+        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++)
+        writeBUF(buf, aus[i]);
+    free(aus);
+    assert(buf == a + getSerializationSize());
+}
+
+void DeformableConvRT::destroy() NOEXCEPT { delete this; }
+
+bool DeformableConvRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+    //todo assert
+}
+
+const char *DeformableConvRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void DeformableConvRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *DeformableConvRT::getPluginType() const NOEXCEPT {
+    return "DeformableConvRT_tkDNN";
+}
+
+const char *DeformableConvRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+IPluginV2 *DeformableConvRT::clone() const NOEXCEPT {
+    auto *p = new DeformableConvRT(chunk_dim,kh,kw,sh,sw,ph,pw,deformableGroup,i_n,i_c,i_h,i_w,o_n,o_c,o_h,o_w,defRT);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+DeformableConvRTPluginCreator::DeformableConvRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("chunk_dim",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("kh",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("kw",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("sh",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("sw",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("ph",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("pw",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("deformableGroup",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("i_n",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("i_c",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("i_h",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("i_w",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("o_n",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("defRT",nullptr,PluginFieldType::kUNKNOWN,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void DeformableConvRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *DeformableConvRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *DeformableConvRTPluginCreator::deserializePlugin(const char *name, const void *serialData,
+                                                            size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new DeformableConvRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *DeformableConvRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    int chunk_dim = *(static_cast<const int *>(fields[0].data));
+    int kh = *(static_cast<const int *>(fields[1].data));
+    int kw = *(static_cast<const int *>(fields[2].data));
+    int sh = *(static_cast<const int *>(fields[3].data));
+    int sw = *(static_cast<const int *>(fields[4].data));
+    int ph = *(static_cast<const int *>(fields[5].data));
+    int pw = *(static_cast<const int *>(fields[6].data));
+    int deformableGroup = *(static_cast<const int *>(fields[7].data));
+    int i_n = *(static_cast<const int *>(fields[8].data));
+    int i_c = *(static_cast<const int *>(fields[9].data));
+    int i_h = *(static_cast<const int *>(fields[10].data));
+    int i_w = *(static_cast<const int *>(fields[11].data));
+    int o_n = *(static_cast<const int *>(fields[12].data));
+    int o_c = *(static_cast<const int *>(fields[13].data));
+    int o_h = *(static_cast<const int *>(fields[14].data));
+    int o_w = *(static_cast<const int *>(fields[14].data));
+    auto *defRT = const_cast<DeformConv2d *>(static_cast<const DeformConv2d *>(fields[15].data));
+    auto *pluginObj = new DeformableConvRT(chunk_dim,kh,kw,sh,sw,ph,pw,deformableGroup,i_n,i_c,i_h,i_w,o_n,o_c,o_h,o_w,defRT);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *DeformableConvRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "DeformableConvRT_tkDNN";
+}
+
+const char *DeformableConvRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *DeformableConvRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+

src/pluginsRT/FlattenConcatRT.cpp

@@ -0,0 +1,163 @@
+#include <tkDNN/pluginsRT/FlattenConcatRT.h>
+using namespace  nvinfer1;
+
+std::vector<PluginField> FlattenConcatRTPluginCreator::mPluginAttributes;
+PluginFieldCollection FlattenConcatRTPluginCreator::mFC{};
+
+FlattenConcatRT::FlattenConcatRT() {
+    stat = cublasCreate(&handle);
+    if (stat != CUBLAS_STATUS_SUCCESS) {
+        printf ("CUBLAS initialization failed\n");
+        return;
+    }
+}
+
+FlattenConcatRT::FlattenConcatRT(const void *data, size_t length) {
+    const char *buf = reinterpret_cast<const char *>(data),*bufCheck=buf;
+    c = readBUF<int>(buf);
+    h = readBUF<int>(buf);
+    w = readBUF<int>(buf);
+    rows = readBUF<int>(buf);
+    cols = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+FlattenConcatRT::~FlattenConcatRT() {}
+
+int FlattenConcatRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims FlattenConcatRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return Dims3{ inputs[0].d[0] * inputs[0].d[1] * inputs[0].d[2], 1, 1};
+}
+
+void FlattenConcatRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,
+                                          DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {
+    assert(nbOutputs == 1 && nbInputs ==1);
+    rows = inputDims[0].d[0];
+    cols = inputDims[0].d[1] * inputDims[0].d[2];
+    c = inputDims[0].d[0] * inputDims[0].d[1] * inputDims[0].d[2];
+    h = 1;
+    w = 1;
+}
+
+int FlattenConcatRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void FlattenConcatRT::terminate() NOEXCEPT {
+    checkERROR(cublasDestroy(handle));
+}
+
+size_t FlattenConcatRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int FlattenConcatRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                             cudaStream_t stream) NOEXCEPT {
+    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) );
+    for(int i=0; i<batchSize; i++) {
+        float const alpha(1.0);
+        float const beta(0.0);
+        int offset = i*rows*cols;
+        checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, rows, cols, &alpha, srcData + offset, cols, &beta, srcData + offset, rows, dstData + offset, rows ));
+    }
+    return 0;
+}
+
+size_t FlattenConcatRT::getSerializationSize() const NOEXCEPT {
+    return 5*sizeof(int);
+}
+
+void FlattenConcatRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a = buf;
+    writeBUF(buf, c);
+    writeBUF(buf, h);
+    writeBUF(buf, w);
+    writeBUF(buf, rows);
+    writeBUF(buf, cols);
+    assert(buf == a + getSerializationSize());
+}
+
+void FlattenConcatRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+bool FlattenConcatRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+}
+
+const char *FlattenConcatRT::getPluginType() const NOEXCEPT {
+    return "FlattenConcatRT_tkDNN";
+}
+
+const char *FlattenConcatRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const char *FlattenConcatRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void FlattenConcatRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2 *FlattenConcatRT::clone() const NOEXCEPT {
+    auto *p = new FlattenConcatRT();
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+FlattenConcatRTPluginCreator::FlattenConcatRTPluginCreator() {
+    mPluginAttributes.clear();
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void FlattenConcatRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *FlattenConcatRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *FlattenConcatRTPluginCreator::deserializePlugin(const char *name, const void *serialData,
+                                                           size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new FlattenConcatRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *FlattenConcatRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    auto *pluginObj = new FlattenConcatRT();
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *FlattenConcatRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "FlattenConcatRT_tkDNN";
+}
+
+const char *FlattenConcatRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *FlattenConcatRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+

src/pluginsRT/MaxPoolingSizeRT.cpp

@@ -0,0 +1,179 @@
+#include <tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> MaxPoolFixedSizeRTPluginCreator::mPluginAttributes;
+PluginFieldCollection MaxPoolFixedSizeRTPluginCreator::mFC{};
+
+MaxPoolFixedSizeRT::MaxPoolFixedSizeRT(int c, int h, int w, int n, int strideH, int strideW, int winSize, int padding){
+    this->c = c;
+    this->h = h;
+    this->w = w;
+    this->n = n;
+    this->stride_H = strideH;
+    this->stride_W = strideW;
+    this->winSize = winSize;
+    this->padding = padding;
+}
+
+MaxPoolFixedSizeRT::MaxPoolFixedSizeRT(const void *data, size_t length) {
+    const char *buf = reinterpret_cast<const char*>(data),*bufCheck = buf;
+    c = readBUF<int>(buf);
+    h = readBUF<int>(buf);
+    w = readBUF<int>(buf);
+    n = readBUF<int>(buf);
+    stride_H = readBUF<int>(buf);
+    stride_W = readBUF<int>(buf);
+    winSize  = readBUF<int>(buf);
+    padding = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+MaxPoolFixedSizeRT::~MaxPoolFixedSizeRT() {
+
+}
+
+int MaxPoolFixedSizeRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims MaxPoolFixedSizeRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return Dims3{this->c, this->h, this->w};
+}
+
+void MaxPoolFixedSizeRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {}
+
+int MaxPoolFixedSizeRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void MaxPoolFixedSizeRT::terminate() NOEXCEPT {
+
+}
+
+size_t MaxPoolFixedSizeRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int MaxPoolFixedSizeRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                                cudaStream_t stream) NOEXCEPT {
+    dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+    dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+    MaxPoolingForward(srcData, dstData, batchSize, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding, stream);
+    return 0;
+}
+
+size_t MaxPoolFixedSizeRT::getSerializationSize() const NOEXCEPT {
+    return 8*sizeof(int);
+}
+
+void MaxPoolFixedSizeRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    writeBUF(buf, this->c);
+    writeBUF(buf, this->h);
+    writeBUF(buf, this->w);
+    writeBUF(buf, this->n);
+    writeBUF(buf, this->stride_H);
+    writeBUF(buf, this->stride_W);
+    writeBUF(buf, this->winSize);
+    writeBUF(buf, this->padding);
+    assert(buf == a + getSerializationSize());
+}
+
+void MaxPoolFixedSizeRT::destroy() NOEXCEPT {
+delete this;
+}
+
+bool MaxPoolFixedSizeRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+    //todo assert
+}
+
+const char *MaxPoolFixedSizeRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void MaxPoolFixedSizeRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *MaxPoolFixedSizeRT::getPluginType() const NOEXCEPT {
+    return "MaxPoolingFixedSizeRT_tkDNN";
+}
+
+const char *MaxPoolFixedSizeRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+IPluginV2 *MaxPoolFixedSizeRT::clone() const NOEXCEPT {
+    auto *p = new MaxPoolFixedSizeRT(c,h,w,n,stride_H,stride_W,winSize,padding);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+MaxPoolFixedSizeRTPluginCreator::MaxPoolFixedSizeRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("c",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("h",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("w",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("n",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("stride_H",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("stride_W",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("winSize",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("padding",nullptr,PluginFieldType::kINT32,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void MaxPoolFixedSizeRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *MaxPoolFixedSizeRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *MaxPoolFixedSizeRTPluginCreator::deserializePlugin(const char *name, const void *serialData,size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new MaxPoolFixedSizeRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *MaxPoolFixedSizeRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    //todo assert
+    int c = *(static_cast<const int *>(fields[0].data));
+    int h = *(static_cast<const int *>(fields[1].data));
+    int w = *(static_cast<const int *>(fields[2].data));
+    int n = *(static_cast<const int *>(fields[3].data));
+    int stride_H = *(static_cast<const int *>(fields[4].data));
+    int stride_W = *(static_cast<const int *>(fields[5].data));
+    int winSize = *(static_cast<const int *>(fields[6].data));
+    int padding = *(static_cast<const int *>(fields[7].data));
+    auto *pluginObj = new MaxPoolFixedSizeRT(c,h,w,n,stride_H,stride_W,winSize,padding);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *MaxPoolFixedSizeRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "MaxPoolingFixedSizeRT_tkDNN";
+}
+
+const char *MaxPoolFixedSizeRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *MaxPoolFixedSizeRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+

src/pluginsRT/RegionRT.cpp

@@ -0,0 +1,174 @@
+#include <tkDNN/pluginsRT/RegionRT.h>
+using namespace nvinfer1;
+std::vector<PluginField> RegionRTPluginCreator::mPluginAttributes;
+PluginFieldCollection RegionRTPluginCreator::mFC{};
+
+RegionRT::RegionRT(int classes, int coords, int num) {
+    this->classes = classes;
+    this->coords = coords;
+    this->num = num;
+}
+
+RegionRT::~RegionRT() {}
+
+RegionRT::RegionRT(const void *data, size_t length) {
+    const char *buf = reinterpret_cast<const char*>(data),*bufCheck=buf;
+    classes = readBUF<int>(buf);
+    coords = readBUF<int>(buf);
+    num = readBUF<int>(buf);
+    c = readBUF<int>(buf);
+    h = readBUF<int>(buf);
+    w = readBUF<int>(buf);
+    assert(buf == bufCheck+length);
+}
+
+int RegionRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims RegionRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return inputs[0];
+}
+
+void RegionRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,
+                              PluginFormat format, int maxBatchSize) NOEXCEPT {
+    c = inputDims[0].d[0];
+    h = inputDims[0].d[1];
+    w = inputDims[0].d[2];
+}
+
+int RegionRT::initialize() NOEXCEPT {return 0;}
+
+void RegionRT::terminate() NOEXCEPT {}
+
+size_t RegionRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; }
+
+int RegionRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                      cudaStream_t stream) NOEXCEPT {
+    dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+    dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+
+    checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
+
+    for (int b = 0; b < batchSize; ++b){
+        for(int n = 0; n < num; ++n){
+            int index = entry_index(b, n*w*h, 0);
+            activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
+
+            index = entry_index(b, n*w*h, coords);
+            activationLOGISTICForward(srcData + index, dstData + index, w*h, stream);
+        }
+    }
+
+    //softmax start
+    int index = entry_index(0, 0, coords + 1);
+    softmaxForward(	srcData + index, classes, batchSize*num,
+                       (c*h*w)/num,
+                       w*h, 1, w*h, 1, dstData + index, stream);
+
+    return 0;
+}
+
+size_t RegionRT::getSerializationSize() const NOEXCEPT {
+    return 6*sizeof(int);
+}
+
+void RegionRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    writeBUF(buf, classes);
+    writeBUF(buf, coords);
+    writeBUF(buf, num);
+    writeBUF(buf, c);
+    writeBUF(buf, h);
+    writeBUF(buf, w);
+    assert(buf == a + getSerializationSize());
+}
+
+const char *RegionRT::getPluginType() const NOEXCEPT {
+    return "RegionRT_tkDNN";
+}
+
+const char *RegionRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void RegionRT::destroy() NOEXCEPT { delete this; }
+
+const char *RegionRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void RegionRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+bool RegionRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+}
+
+IPluginV2 *RegionRT::clone() const NOEXCEPT {
+    auto *p = new RegionRT(classes,coords,num);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+RegionRTPluginCreator::RegionRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("coords",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void RegionRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *RegionRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *RegionRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new RegionRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *RegionRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    assert(fc->nbFields == 3);
+    assert(fields[0].type == PluginFieldType::kINT32);
+    assert(fields[1].type == PluginFieldType::kINT32);
+    assert(fields[2].type == PluginFieldType::kINT32);
+    int classes = *(static_cast<const int*>(fields[0].data));
+    int coords = *(static_cast<const int*>(fields[1].data));
+    int num = *(static_cast<const int*>(fields[2].data));
+    RegionRT *pluginObj = new RegionRT(classes,coords,num);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *RegionRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "RegionRT_tkDNN";
+}
+
+const char *RegionRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *RegionRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+

src/pluginsRT/ReorgRT.cpp

@@ -0,0 +1,151 @@
+#include <tkDNN/pluginsRT/ReorgRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> ReorgRTPluginCreator::mPluginAttributes;
+PluginFieldCollection ReorgRTPluginCreator::mFC{};
+
+ReorgRT::ReorgRT(int stride) {
+    this->stride = stride;
+}
+
+ReorgRT::~ReorgRT() {}
+
+ReorgRT::ReorgRT(const void *data, size_t length) {
+    const char* buf  = reinterpret_cast<const char*>(data),*bufCheck = buf;
+    stride = readBUF<int>(buf);
+    c = readBUF<int>(buf);
+    h = readBUF<int>(buf);
+    w = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+int ReorgRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims ReorgRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return Dims3{inputs[0].d[0]*stride*stride, inputs[0].d[1]/stride, inputs[0].d[2]/stride};
+}
+
+void ReorgRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {
+    c = inputDims[0].d[0];
+    h = inputDims[0].d[1];
+    w = inputDims[0].d[2];
+}
+
+int ReorgRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void ReorgRT::terminate() NOEXCEPT {}
+
+size_t ReorgRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int ReorgRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,cudaStream_t stream) NOEXCEPT {
+    reorgForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
+                 reinterpret_cast<dnnType*>(outputs[0]),
+                 batchSize, c, h, w, stride, stream);
+    return 0;
+}
+
+size_t ReorgRT::getSerializationSize() const NOEXCEPT {
+    return 4*sizeof(int);
+}
+
+void ReorgRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    writeBUF(buf, stride);
+    writeBUF(buf, c);
+    writeBUF(buf, h);
+    writeBUF(buf, w);
+    assert(buf == a + getSerializationSize());
+}
+
+bool ReorgRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+}
+
+const char *ReorgRT::getPluginType() const NOEXCEPT {
+    return "ReorgRT_tkDNN";
+}
+
+const char *ReorgRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void ReorgRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+const char *ReorgRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void ReorgRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2 *ReorgRT::clone() const NOEXCEPT {
+    auto *p = new ReorgRT(stride);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+ReorgRTPluginCreator::ReorgRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void ReorgRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *ReorgRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *ReorgRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new ReorgRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *ReorgRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    assert(fc->nbFields == 1);
+    assert(fields[0].type == PluginFieldType::kINT32);
+    int stride = *(static_cast<const int *>(fields[0].data));
+    auto *pluginObj = new ReorgRT(stride);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *ReorgRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "ReorgRT_tkDNN";
+}
+
+const char *ReorgRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *ReorgRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+
+
+

src/pluginsRT/ReshapeRT.cpp

@@ -0,0 +1,152 @@
+#include <tkDNN/pluginsRT/ReshapeRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> ReshapeRTPluginCreator::mPluginAttributes;
+PluginFieldCollection ReshapeRTPluginCreator::mFC{};
+
+ReshapeRT::ReshapeRT(dataDim_t newDim) {
+    new_dim = newDim;
+    n = new_dim.n;
+    c = new_dim.c;
+    h = new_dim.h;
+    w = new_dim.w;
+}
+
+ReshapeRT::ReshapeRT(const void *data, size_t length) {
+    const char *buf = reinterpret_cast<const char*>(data),*bufCheck = buf;
+    new_dim.n = readBUF<int>(buf);
+    new_dim.c = readBUF<int>(buf);
+    new_dim.h = readBUF<int>(buf);
+    new_dim.w = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+ReshapeRT::~ReshapeRT() {}
+
+int ReshapeRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims ReshapeRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return Dims3{ c,h,w} ;
+}
+
+void ReshapeRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {}
+
+int ReshapeRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void ReshapeRT::terminate() NOEXCEPT {}
+
+size_t ReshapeRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int ReshapeRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                       cudaStream_t stream) NOEXCEPT {
+    dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+    dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+
+    checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
+    return 0;
+}
+
+size_t ReshapeRT::getSerializationSize() const NOEXCEPT {
+    return 4*sizeof(int);
+}
+
+void ReshapeRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a = buf;
+    writeBUF(buf, n);
+    writeBUF(buf, c);
+    writeBUF(buf, h);
+    writeBUF(buf, w);
+    assert(buf == a + getSerializationSize());
+}
+
+bool ReshapeRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+    //todo assert
+}
+
+const char *ReshapeRT::getPluginType() const NOEXCEPT {
+    return "ReshapeRT_tkDNN";
+}
+
+const char *ReshapeRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void ReshapeRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+const char *ReshapeRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void ReshapeRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2 *ReshapeRT::clone() const NOEXCEPT {
+    auto *p = new ReshapeRT(new_dim);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+ReshapeRTPluginCreator::ReshapeRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("new_dim",nullptr,PluginFieldType::kUNKNOWN,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void ReshapeRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *ReshapeRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *ReshapeRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new ReshapeRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *ReshapeRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    dataDim_t newDim = *(static_cast<const dataDim_t *>(fields[0].data));
+    ReshapeRT *pluginObj = new ReshapeRT(newDim);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *ReshapeRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "ReshapeRT_tkDNN";
+}
+
+const char *ReshapeRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *ReshapeRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+
+
+

src/pluginsRT/ResizeLayerRT.cpp

@@ -0,0 +1,163 @@
+#include <tkDNN/pluginsRT/ResizeLayerRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> ResizeLayerRTPluginCreator::mPluginAttributes;
+PluginFieldCollection ResizeLayerRTPluginCreator::mFC{};
+
+
+ResizeLayerRT::ResizeLayerRT(int c, int h, int w) {
+    o_c = c;
+    o_h = h;
+    o_w = w;
+}
+
+ResizeLayerRT::ResizeLayerRT(const void *data, size_t length) {
+    const char *buf = reinterpret_cast<const char*>(data),*bufCheck = buf;
+    o_c = readBUF<int>(buf);
+    o_h = readBUF<int>(buf);
+    o_w = readBUF<int>(buf);
+    i_c = readBUF<int>(buf);
+    i_h = readBUF<int>(buf);
+    i_w = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+ResizeLayerRT::~ResizeLayerRT() {}
+
+int ResizeLayerRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims ResizeLayerRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return Dims3{o_c, o_h, o_w};
+}
+
+void ResizeLayerRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,
+                                        DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {
+    i_c = inputDims[0].d[0];
+    i_h = inputDims[0].d[1];
+    i_w = inputDims[0].d[2];
+}
+
+int ResizeLayerRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void ResizeLayerRT::terminate() NOEXCEPT {}
+
+size_t ResizeLayerRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; }
+
+int ResizeLayerRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                           cudaStream_t stream) NOEXCEPT {
+    resizeForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
+                  reinterpret_cast<dnnType*>(outputs[0]),
+                  batchSize, i_c, i_h, i_w, o_c, o_h, o_w, stream);
+    return 0;
+}
+
+size_t ResizeLayerRT::getSerializationSize() const NOEXCEPT {
+    return 6*sizeof(int);
+}
+
+void ResizeLayerRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    writeBUF(buf, o_c);
+    writeBUF(buf, o_h);
+    writeBUF(buf, o_w);
+    writeBUF(buf, i_c);
+    writeBUF(buf, i_h);
+    writeBUF(buf, i_w);
+    assert(buf == a + getSerializationSize());
+}
+
+bool ResizeLayerRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+    //todo assert
+}
+
+const char *ResizeLayerRT::getPluginType() const NOEXCEPT {
+    return "ResizeLayerRT_tkDNN";
+}
+
+const char *ResizeLayerRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void ResizeLayerRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+const char *ResizeLayerRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void ResizeLayerRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2 *ResizeLayerRT::clone() const NOEXCEPT {
+    auto *p = new ResizeLayerRT(o_c,o_h,o_w);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+ResizeLayerRTPluginCreator::ResizeLayerRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("o_c",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("o_h",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("o_w",nullptr,PluginFieldType::kINT32,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void ResizeLayerRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *ResizeLayerRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *ResizeLayerRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new ResizeLayerRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *ResizeLayerRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    assert(fc->nbFields == 3);
+    assert(fields[0].type == PluginFieldType::kINT32);
+    assert(fields[1].type == PluginFieldType::kINT32);
+    assert(fields[2].type == PluginFieldType::kINT32);
+    int oc = *(static_cast<const int *>(fields[0].data));
+    int oh = *(static_cast<const int *>(fields[1].data));
+    int ow = *(static_cast<const int *>(fields[2].data));
+    auto *pluginObj = new ResizeLayerRT(oc,oh,ow);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *ResizeLayerRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "ResizeLayerRT_tkDNN";
+}
+
+const char *ResizeLayerRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *ResizeLayerRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+

src/pluginsRT/RouteRT.cpp

@@ -0,0 +1,183 @@
+#include <tkDNN/pluginsRT/RouteRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> RouteRTPluginCreator::mPluginAttributes;
+PluginFieldCollection RouteRTPluginCreator::mFC{};
+
+
+RouteRT::RouteRT(int groups, int group_id) {
+    this->groups = groups;
+    this->group_id = group_id;
+}
+
+RouteRT::~RouteRT() {}
+
+RouteRT::RouteRT(const void *data, size_t length) {
+    const char* buf = reinterpret_cast<const char*>(data),*bufCheck = buf;
+    groups = readBUF<int>(buf);
+    group_id = readBUF<int>(buf);
+    in = readBUF<int>(buf);
+    for(int i=0;i <MAX_INPUTS;i++){
+        c_in[i] = readBUF<int>(buf);
+    }
+    c= readBUF<int>(buf);
+    h = readBUF<int>(buf);
+    w = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+int RouteRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims RouteRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    int out_c = 0;
+    for(int i=0; i<nbInputDims; i++) out_c += inputs[i].d[0];
+    return Dims3{out_c/groups, inputs[0].d[1], inputs[0].d[2]};
+}
+
+void
+RouteRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,
+                             PluginFormat format, int maxBatchSize) NOEXCEPT {
+    in = nbInputs;
+    c = 0;
+    for(int i=0; i<nbInputs; i++) {
+        c_in[i] = inputDims[i].d[0];
+        c += inputDims[i].d[0];
+    }
+    h = inputDims[0].d[1];
+    w = inputDims[0].d[2];
+    c /= groups;
+}
+
+int RouteRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void RouteRT::terminate() NOEXCEPT {}
+
+size_t RouteRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int RouteRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                     cudaStream_t stream) NOEXCEPT {
+    dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+    for(int b=0; b<batchSize; b++) {
+        int offset = 0;
+        for(int i=0; i<in; i++) {
+            dnnType *input = (dnnType*)reinterpret_cast<const dnnType*>(inputs[i]);
+            int in_dim = c_in[i]*h*w;
+            int part_in_dim = in_dim / this->groups;
+            checkCuda( cudaMemcpyAsync(dstData + b*c*w*h + offset, input + b*c*w*h*groups + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) );
+            offset += part_in_dim;
+        }
+    }
+    return 0;
+}
+
+size_t RouteRT::getSerializationSize() const NOEXCEPT {
+    return (6+MAX_INPUTS)*sizeof(int);
+}
+
+void RouteRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    writeBUF(buf, groups);
+    writeBUF(buf, group_id);
+    writeBUF(buf, in);
+    for(int i=0; i<MAX_INPUTS; i++)
+        writeBUF(buf, c_in[i]);
+    writeBUF(buf, c);
+    writeBUF(buf, h);
+    writeBUF(buf, w);
+    assert(buf == a + getSerializationSize());
+}
+
+const char *RouteRT::getPluginType() const NOEXCEPT {
+    return "RouteRT_tkDNN";
+}
+
+const char *RouteRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void RouteRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+const char *RouteRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void RouteRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+bool RouteRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+    //todo assert
+}
+
+IPluginV2 *RouteRT::clone() const NOEXCEPT {
+    auto *p = new RouteRT(groups,group_id);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+RouteRTPluginCreator::RouteRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("groups",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("group_id",nullptr,PluginFieldType::kINT32,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void RouteRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *RouteRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *RouteRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new RouteRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *RouteRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    assert(fc->nbFields == 2);
+    assert(fields[0].type == PluginFieldType::kINT32);
+    assert(fields[1].type == PluginFieldType::kINT32);
+    int groups = *(static_cast<const int *>(fields[0].data));
+    int group_id = *(static_cast<const int *>(fields[1].data));
+    RouteRT *pluginObj = new RouteRT(groups,group_id);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *RouteRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "RouteRT_tkDNN";
+}
+
+const char *RouteRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *RouteRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+

src/pluginsRT/ShortcutRT.cpp

@@ -0,0 +1,167 @@
+#include <tkDNN/pluginsRT/ShortcutRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> ShortcutRTPluginCreator::mPluginAttributes;
+PluginFieldCollection ShortcutRTPluginCreator::mFC{};
+
+ShortcutRT::ShortcutRT(tk::dnn::dataDim_t bdim, bool mul) {
+    bDim = bdim;
+    this->bc = bDim.c;
+    this->bh = bDim.h;
+    this->bw = bDim.w;
+    this->mul = mul;
+}
+
+ShortcutRT::~ShortcutRT() {}
+
+ShortcutRT::ShortcutRT(const void *data, size_t length) {
+    const char* buf =reinterpret_cast<const char*>(data),*bufCheck = buf;
+    bDim.c = readBUF<int>(buf);
+    bDim.h = readBUF<int>(buf);
+    bDim.w = readBUF<int>(buf);
+    bDim.l = 1;
+    mul = readBUF<bool>(buf);
+    c = readBUF<int>(buf);
+    h = readBUF<int>(buf);
+    w = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+int ShortcutRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims ShortcutRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return Dims3{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]};
+}
+
+void ShortcutRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,
+                                     DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {
+    c = inputDims[0].d[0];
+    h = inputDims[0].d[1];
+    w = inputDims[0].d[2];
+}
+
+int ShortcutRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void ShortcutRT::terminate() NOEXCEPT {}
+
+size_t ShortcutRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT { return 0; }
+
+int ShortcutRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                        cudaStream_t stream) NOEXCEPT {
+    dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+    dnnType *srcDataBack = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
+    dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+
+    checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
+    shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, mul, stream);
+
+    return 0;
+}
+
+size_t ShortcutRT::getSerializationSize() const NOEXCEPT {
+    return 6*sizeof(int) + sizeof(bool);
+}
+
+void ShortcutRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    writeBUF(buf, bc);
+    writeBUF(buf, bh);
+    writeBUF(buf, bw);
+    writeBUF(buf, mul);
+    writeBUF(buf, c);
+    writeBUF(buf, h);
+    writeBUF(buf, w);
+    assert(buf == a + getSerializationSize());
+}
+
+bool ShortcutRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+}
+
+const char *ShortcutRT::getPluginType() const NOEXCEPT {
+    return "ShortcutRT_tkDNN";
+}
+
+const char *ShortcutRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void ShortcutRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+const char *ShortcutRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void ShortcutRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2 *ShortcutRT::clone() const NOEXCEPT {
+    auto *p = new ShortcutRT(bDim,mul);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+ShortcutRTPluginCreator::ShortcutRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("bDim",nullptr,PluginFieldType::kUNKNOWN,1));
+    mPluginAttributes.emplace_back(PluginField("mul",nullptr,PluginFieldType::kUNKNOWN,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void ShortcutRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *ShortcutRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *ShortcutRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new ShortcutRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *ShortcutRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    //todo assert
+    tk::dnn::dataDim_t bdim = *(static_cast<const tk::dnn::dataDim_t *>(fields[0].data));
+    bool mul = *(static_cast<const bool *>(fields[1].data));
+    auto *pluginObj = new ShortcutRT(bdim,mul);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *ShortcutRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "ShortcutRT_tkDNN";
+}
+
+const char *ShortcutRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *ShortcutRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+
+
+

src/pluginsRT/UpsampleRT.cpp

@@ -0,0 +1,155 @@
+#include <tkDNN/pluginsRT/UpsampleRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> UpsampleRTPluginCreator::mPluginAttributes;
+PluginFieldCollection UpsampleRTPluginCreator::mFC{};
+
+UpsampleRT::UpsampleRT(int stride) {
+    this->stride = stride;
+}
+
+UpsampleRT::UpsampleRT(const void *data, size_t length) {
+    const char* buf = reinterpret_cast<const char*>(data),*bufCheck=buf;
+    stride = readBUF<int>(buf);
+    c = readBUF<int>(buf);
+    h = readBUF<int>(buf);
+    w = readBUF<int>(buf);
+    assert(buf == bufCheck + length);
+}
+
+UpsampleRT::~UpsampleRT() {}
+
+int UpsampleRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims UpsampleRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return Dims3(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride);
+}
+
+void UpsampleRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs,
+                                     DataType type, PluginFormat format, int maxBatchSize) NOEXCEPT {
+    c = inputDims[0].d[0];
+    h = inputDims[0].d[1];
+    w = inputDims[0].d[2];
+}
+
+int UpsampleRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void UpsampleRT::terminate() NOEXCEPT {}
+
+size_t UpsampleRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int UpsampleRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                        cudaStream_t stream) NOEXCEPT {
+    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 UpsampleRT::getSerializationSize() const NOEXCEPT {
+    return 4*sizeof(int);
+}
+
+void UpsampleRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char*>(buffer),*a=buf;
+    writeBUF(buf, stride);
+    writeBUF(buf, c);
+    writeBUF(buf, h);
+    writeBUF(buf, w);
+    assert(buf == a + getSerializationSize());
+}
+
+bool UpsampleRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+}
+
+const char *UpsampleRT::getPluginType() const NOEXCEPT {
+    return "Upsample_tkDNN";
+}
+
+const char *UpsampleRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void UpsampleRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+const char *UpsampleRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void UpsampleRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2 *UpsampleRT::clone() const NOEXCEPT {
+    auto *p = new UpsampleRT(stride);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+UpsampleRTPluginCreator::UpsampleRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("stride",nullptr,PluginFieldType::kINT32,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void UpsampleRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *UpsampleRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *UpsampleRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new UpsampleRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *UpsampleRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    int stride = *(static_cast<const int *>(fields[0].data));
+    auto *pluginObj = new UpsampleRT(stride);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *UpsampleRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "Upsample_tkDNN";
+}
+
+const char *UpsampleRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *UpsampleRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+

src/pluginsRT/YoloRT.cpp

@@ -0,0 +1,253 @@
+#include <tkDNN/pluginsRT/YoloRT.h>
+using namespace nvinfer1;
+
+std::vector<PluginField> YoloRTPluginCreator::mPluginAttributes;
+PluginFieldCollection YoloRTPluginCreator::mFC{};
+
+YoloRT::YoloRT(int classes, int num, tk::dnn::Yolo *Yolo, int n_masks, float scale_xy, float nms_thresh, int nms_kind,
+               int new_coords) {
+    this->yolo = Yolo;
+    this->classes = classes;
+    this->num = num;
+    this->n_masks = n_masks;
+    this->scaleXY = scale_xy;
+    this->nms_thresh = nms_thresh;
+    this->nms_kind = nms_kind;
+    this->new_coords = new_coords;
+
+    mask = new dnnType[n_masks];
+    bias = new dnnType[num * n_masks * 2];
+    if (yolo != nullptr) {
+        memcpy(mask, yolo->mask_h, sizeof(dnnType) * n_masks);
+        memcpy(bias, yolo->bias_h, sizeof(dnnType) * num * n_masks * 2);
+        classesNames = yolo->classesNames;
+    }
+}
+
+YoloRT::YoloRT(const void *data, size_t length) {
+    std::vector<float> maskTemp,biasTemp;
+    std::cout<<"LENGTH : "<<length<<std::endl;
+    const char* buf = reinterpret_cast<const char*>(data),*bufCheck = buf;
+    classes = readBUF<int>(buf);
+    num = readBUF<int>(buf);
+    n_masks = readBUF<int>(buf);
+    scaleXY = readBUF<float>(buf);
+    nms_thresh = readBUF<float>(buf);
+    nms_kind = readBUF<int>(buf);
+    new_coords = readBUF<int>(buf);
+    c = readBUF<int>(buf);
+    h = readBUF<int>(buf);
+    w = readBUF<int>(buf);
+    for(int i=0;i<n_masks;i++){
+        maskTemp.push_back(readBUF<dnnType>(buf));
+        std::cout<<maskTemp[i]<<std::endl;
+    }
+    for(int i=0;i<n_masks*2*num;i++){
+        biasTemp.push_back(readBUF<dnnType>(buf));
+        std::cout<<biasTemp[i]<<std::endl;
+    }
+    mask = maskTemp.data();
+    bias = biasTemp.data();
+    classesNames.resize(classes);
+    for(int i=0;i<classes;i++){
+        char tmp[YOLORT_CLASSNAME_W];
+        for(int j=0;j<YOLORT_CLASSNAME_W;j++)
+            tmp[j] = readBUF<char>(buf);
+        classesNames[1] = std::string(tmp);
+    }
+    assert(buf == bufCheck + length);
+}
+
+YoloRT::~YoloRT() {}
+
+int YoloRT::getNbOutputs() const NOEXCEPT {
+    return 1;
+}
+
+Dims YoloRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT {
+    return inputs[0];
+}
+
+void YoloRT::configureWithFormat(const Dims *inputDims, int nbInputs, const Dims *outputDims, int nbOutputs, DataType type,
+                            PluginFormat format, int maxBatchSize) NOEXCEPT {
+    c = inputDims[0].d[0];
+    h = inputDims[0].d[1];
+    w = inputDims[0].d[2];
+}
+
+int YoloRT::initialize() NOEXCEPT {
+    return 0;
+}
+
+void YoloRT::terminate() NOEXCEPT {}
+
+size_t YoloRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT {
+    return 0;
+}
+
+int YoloRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+                    cudaStream_t stream) NOEXCEPT {
+    dnnType *srcData = (dnnType *) reinterpret_cast<const dnnType *>(inputs[0]);
+    dnnType *dstData = reinterpret_cast<dnnType *>(outputs[0]);
+
+    checkCuda(cudaMemcpyAsync(dstData, srcData, batchSize * c * h * w * sizeof(dnnType), cudaMemcpyDeviceToDevice,
+                              stream));
+
+
+    for (int b = 0; b < batchSize; ++b) {
+        for (int n = 0; n < n_masks; ++n) {
+            int index = entry_index(b, n * w * h, 0);
+            if (new_coords == 1) {
+                if (this->scaleXY != 1)
+                    scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1);
+            } else {
+                activationLOGISTICForward(srcData + index, dstData + index, 2 * w * h, stream); //x,y
+
+                if (this->scaleXY != 1)
+                    scalAdd(dstData + index, 2 * w * h, this->scaleXY, -0.5 * (this->scaleXY - 1), 1);
+
+                index = entry_index(b, n * w * h, 4);
+                activationLOGISTICForward(srcData + index, dstData + index, (1 + classes) * w * h, stream);
+            }
+        }
+    }
+
+    //std::cout<<"YOLO END\n";
+    return 0;
+}
+
+size_t YoloRT::getSerializationSize() const NOEXCEPT {
+    return 8 * sizeof(int) + 2 * sizeof(float) + n_masks * sizeof(dnnType) + num * n_masks * 2 * sizeof(dnnType) +
+           YOLORT_CLASSNAME_W * classes * sizeof(char);
+}
+
+bool YoloRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+    return true;
+}
+
+void YoloRT::serialize(void *buffer) const NOEXCEPT {
+    char *buf = reinterpret_cast<char *>(buffer), *a = buf;
+    writeBUF(buf, classes);    //std::cout << "Classes :" << classes << std::endl;
+    writeBUF(buf, num);        //std::cout << "Num : " << num << std::endl;
+    std::cout<<num<<std::endl;
+    writeBUF(buf, n_masks);    //std::cout << "N_Masks" << n_masks << std::endl;
+    writeBUF(buf, scaleXY);    //std::cout << "ScaleXY :" << scaleXY << std::endl;
+    writeBUF(buf, nms_thresh); //std::cout << "nms_thresh :" << nms_thresh << std::endl;
+    writeBUF(buf, nms_kind);    //std::cout << "nms_kind : " << nms_kind << std::endl;
+    writeBUF(buf, new_coords); //std::cout << "new_coords : " << new_coords << std::endl;
+    writeBUF(buf, c);            //std::cout << "C : " << c << std::endl;
+    writeBUF(buf, h);            //std::cout << "H : " << h << std::endl;
+    writeBUF(buf, w);            //std::cout << "C : " << c << std::endl;
+    for (int i = 0; i < n_masks; i++) {
+        writeBUF(buf, mask[i]); //std::cout << "mask[i] : " << mask[i] << std::endl;
+    }
+    for (int i = 0; i < n_masks * 2 * num; i++) {
+        writeBUF(buf, bias[i]); //std::cout << "bias[i] : " << bias[i] << std::endl;
+    }
+
+    // save classes names
+    for (int i = 0; i < classes; i++) {
+        char tmp[YOLORT_CLASSNAME_W];
+        strcpy(tmp, classesNames[i].c_str());
+        for (int j = 0; j < YOLORT_CLASSNAME_W; j++) {
+            writeBUF(buf, tmp[j]);
+        }
+    }
+    assert(buf == a + getSerializationSize());
+}
+
+const char *YoloRT::getPluginType() const NOEXCEPT {
+    return "YoloRT_tkDNN";
+}
+
+const char *YoloRT::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+void YoloRT::destroy() NOEXCEPT {
+    delete this;
+}
+
+const char *YoloRT::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+void YoloRT::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+IPluginV2 *YoloRT::clone() const NOEXCEPT {
+    auto *p = new YoloRT(classes, num,yolo, n_masks, scaleXY, nms_thresh, nms_kind, new_coords);
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+YoloRTPluginCreator::YoloRTPluginCreator() {
+    mPluginAttributes.emplace_back(PluginField("classes",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("num",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("yolo",nullptr,PluginFieldType::kUNKNOWN,1));
+    mPluginAttributes.emplace_back(PluginField("numMasks",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("scaleXY",nullptr,PluginFieldType::kFLOAT32,1));
+    mPluginAttributes.emplace_back(PluginField("nmsThresh",nullptr,PluginFieldType::kFLOAT32,1));
+    mPluginAttributes.emplace_back(PluginField("nmsKind",nullptr,PluginFieldType::kINT32,1));
+    mPluginAttributes.emplace_back(PluginField("newCoords",nullptr,PluginFieldType::kINT32,1));
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void YoloRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *YoloRTPluginCreator::getPluginNamespace() const NOEXCEPT {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *YoloRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT {
+    auto *pluginObj = new YoloRT(serialData,serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *YoloRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT {
+    const PluginField *fields = fc->fields;
+    //todo assert
+    int classes = *(static_cast<const int *>(fields[0].data));
+    int num = *(static_cast<const int *>(fields[1].data));
+    Yolo *yoloTemp = const_cast<Yolo *>(static_cast<const Yolo *>(fields[2].data));
+    int numMasks = *(static_cast<const int*>(fields[3].data));
+    float scaleXY = *(static_cast<const float *>(fields[4].data));
+    float nmsThresh = *(static_cast<const float *>(fields[5].data));
+    int nmsKind = *(static_cast<const int *>(fields[6].data));
+    int newCoords = *(static_cast<const int *>(fields[7].data));
+    YoloRT *pluginObj = new YoloRT(classes,num,yoloTemp,numMasks,scaleXY,nmsThresh,nmsKind,newCoords);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *YoloRTPluginCreator::getPluginName() const NOEXCEPT {
+    return "YoloRT_tkDNN";
+}
+
+const char *YoloRTPluginCreator::getPluginVersion() const NOEXCEPT {
+    return "1";
+}
+
+const PluginFieldCollection *YoloRTPluginCreator::getFieldNames() NOEXCEPT {
+    return &mFC;
+}
+
+
+
+
+
+
+
+
+
+
+
+
+
+