include/tkDNN/Layer.h

@@ -19,6 +19,7 @@ enum layerType_t {
     LAYER_ACTIVATION_CRELU,
     LAYER_ACTIVATION_LEAKY,
     LAYER_ACTIVATION_MISH,
+    LAYER_ACTIVATION_SWISH,
     LAYER_ACTIVATION_LOGISTIC,
     LAYER_FLATTEN,
     LAYER_RESHAPE,
@@ -75,6 +76,7 @@ public:
             case LAYER_ACTIVATION_CRELU:    return "ActivationCReLU";
             case LAYER_ACTIVATION_LEAKY:    return "ActivationLeaky";
             case LAYER_ACTIVATION_MISH:     return "ActivationMish";
+            case LAYER_ACTIVATION_SWISH:    return "ActivationSwish";
             case LAYER_ACTIVATION_LOGISTIC: return "ActivationLogistic";
             case LAYER_FLATTEN:             return "Flatten";
             case LAYER_RESHAPE:             return "Reshape";
@@ -223,7 +225,8 @@ typedef enum {
     ACTIVATION_ELU     = 100,
     ACTIVATION_LEAKY   = 101,
     ACTIVATION_MISH   = 102,
-    ACTIVATION_LOGISTIC   = 103
+    ACTIVATION_LOGISTIC   = 103,
+    ACTIVATION_SWISH = 104
 } tkdnnActivationMode_t;
 
 /**
@@ -245,6 +248,8 @@ public:
             return LAYER_ACTIVATION_LEAKY;
         else if (act_mode == ACTIVATION_MISH)
             return LAYER_ACTIVATION_MISH;
+        else if (act_mode == ACTIVATION_SWISH)
+            return LAYER_ACTIVATION_SWISH;
         else if (act_mode == ACTIVATION_LOGISTIC)
             return LAYER_ACTIVATION_LOGISTIC;
         else

include/tkDNN/NetworkRT.h

@@ -7,10 +7,11 @@
 #include "Layer.h"
 #include "NvInfer.h"
 #include <memory>
-#include <tkDNN/kernels.h>
+#include <kernels.h>
 #include <pluginsRT/ActivationLeakyRT.h>
 #include <pluginsRT/ActivationLogisticRT.h>
 #include <pluginsRT/ActivationMishRT.h>
+#include <pluginsRT/ActivationSwishRT.h>
 #include <pluginsRT/ActivationReLUCeilingRT.h>
 #include <pluginsRT/DeformableConvRT.h>
 #include <pluginsRT/FlattenConcatRT.h>
@@ -30,85 +31,85 @@
 
 namespace tk { namespace dnn {
 
-class NetworkRT {
+        class NetworkRT {
 
-public:
-    nvinfer1::DataType dtRT;
-    nvinfer1::IBuilder *builderRT;
-    nvinfer1::IRuntime *runtimeRT;
-    nvinfer1::INetworkDefinition *networkRT; 
-#if NV_TENSORRT_MAJOR >= 6  
-    nvinfer1::IBuilderConfig *configRT;
+        public:
+            nvinfer1::DataType dtRT;
+            nvinfer1::IBuilder *builderRT;
+            nvinfer1::IRuntime *runtimeRT;
+            nvinfer1::INetworkDefinition *networkRT;
+#if NV_TENSORRT_MAJOR >= 6
+            nvinfer1::IBuilderConfig *configRT;
 #endif
-    
-    nvinfer1::ICudaEngine *engineRT;
-    nvinfer1::IExecutionContext *contextRT;
 
-    const static int MAX_BUFFERS_RT = 10;
-    void* buffersRT[MAX_BUFFERS_RT];
-    dataDim_t buffersDIM[MAX_BUFFERS_RT];
-    int buf_input_idx, buf_output_idx;
-    bool builderActive = false;
-    dataDim_t input_dim, output_dim;
-    dnnType *output;
-    cudaStream_t stream;
+            nvinfer1::ICudaEngine *engineRT;
+            nvinfer1::IExecutionContext *contextRT;
 
-    std::vector<nvinfer1::YoloRT*> yolo_plugins; // yolo layers in network
+            const static int MAX_BUFFERS_RT = 10;
+            void* buffersRT[MAX_BUFFERS_RT];
+            dataDim_t buffersDIM[MAX_BUFFERS_RT];
+            int buf_input_idx, buf_output_idx;
+            bool builderActive = false;
+            dataDim_t input_dim, output_dim;
+            dnnType *output;
+            cudaStream_t stream;
 
-    NetworkRT(Network *net, const char *name);
-    virtual ~NetworkRT();
+            std::vector<nvinfer1::YoloRT*> yolo_plugins; // yolo layers in network
 
-    int getMaxBatchSize() {
-        if(engineRT != nullptr)
-            return engineRT->getMaxBatchSize();
-        else
-            return 0;
-    }
+            NetworkRT(Network *net, const char *name);
+            virtual ~NetworkRT();
 
-    int getBuffersN() {
-        if(engineRT != nullptr)
-            return engineRT->getNbBindings();
-        else 
-            return 0;
-    }
+            int getMaxBatchSize() {
+                if(engineRT != nullptr)
+                    return engineRT->getMaxBatchSize();
+                else
+                    return 0;
+            }
 
-    /**
-        Do inference
-    */
-    dnnType* infer(dataDim_t &dim, dnnType* data);
-    void enqueue(int batchSize = 1);    
+            int getBuffersN() {
+                if(engineRT != nullptr)
+                    return engineRT->getNbBindings();
+                else
+                    return 0;
+            }
 
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Activation *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Dense *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Pooling *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Softmax *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Route *l);
-    nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Flatten *l);
-    nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reshape *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Resize *l);
-    nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reorg *l);
-    nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Region *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l);
-    nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Yolo *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Upsample *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeformConv2d *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input,Padding *l);
-    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor* input,MulAdd *l);
+            /**
+                Do inference
+            */
+            dnnType* infer(dataDim_t &dim, dnnType* data);
+            void enqueue(int batchSize = 1);
+
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Activation *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Dense *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Pooling *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Softmax *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Route *l);
+            nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Flatten *l);
+            nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reshape *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Resize *l);
+            nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Reorg *l);
+            nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Region *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l);
+            nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Yolo *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Upsample *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeformConv2d *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input,Padding *l);
+            nvinfer1::ILayer* convert_layer(nvinfer1::ITensor* input,MulAdd *l);
 
 #if NV_TENSORRT_MAJOR > 5 && NV_TENSORRT_MAJOR < 8
-    bool serialize(const char *filename);
+            bool serialize(const char *filename);
 #else
-    bool serialize(const char *filename,nvinfer1::IHostMemory *ptr);
+            bool serialize(const char *filename,nvinfer1::IHostMemory *ptr);
 #endif
 
-    bool deserialize(const char *filename);
-    void destroy();
+            bool deserialize(const char *filename);
+            void destroy();
 
 
 
-};
+        };
 
-}}
-#endif //NETWORKRT_H
+    }}
+#endif //NETWORKRT_H

include/tkDNN/kernels.h

@@ -9,6 +9,7 @@ void activationReLUCeilingForward(dnnType *srcData, dnnType *dstData, int size,
 void activationLOGISTICForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0));
 void activationSIGMOIDForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0));
 void activationMishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream= cudaStream_t(0));
+void activationSwishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream= cudaStream_t(0));
 
 void fill(dnnType *data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0));
 

include/tkDNN/pluginsRT/ActivationSwishRT.h

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

src/Activation.cpp

@@ -52,6 +52,10 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
     else if(act_mode == ACTIVATION_MISH) {
         activationMishForward(srcData, dstData, dim.tot());
 
+    }
+    else if(act_mode == ACTIVATION_SWISH) {
+        activationSwishForward(srcData, dstData, dim.tot());
+
     }
     else if(act_mode == ACTIVATION_LOGISTIC) {
         activationLOGISTICForward(srcData, dstData, dim.tot());

src/DarknetParser.cpp

@@ -197,6 +197,7 @@ namespace tk { namespace dnn {
             if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU);
             else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY;
             else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH;
+            else if(f.activation == "swish") act = tk::dnn::ACTIVATION_SWISH;
             else if(f.activation == "logistic") act = tk::dnn::ACTIVATION_LOGISTIC;
             else { FatalError("activation not supported: " + f.activation); }
             netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act);

src/NetworkRT.cpp

@@ -254,7 +254,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
         return convert_layer(input, (Conv2d*) l);
     if(type == LAYER_POOLING)
         return convert_layer(input, (Pooling*) l);
-    if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_LOGISTIC)
+    if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_SWISH || type == LAYER_ACTIVATION_LOGISTIC)
         return convert_layer(input, (Activation*) l);
     if(type == LAYER_SOFTMAX)
         return convert_layer(input, (Softmax*) l);
@@ -646,8 +646,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
         checkNULL(lRT);
         return lRT;
     }
-    else if(l->act_mode == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU){
-        IActivationLayer *lRT = networkRT->addActivation(*input,ActivationType::kELU);
+    else if(l->act_mode == CUDNN_ACTIVATION_ELU || l->act_mode == ACTIVATION_ELU) {
+        IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kELU);
+    }
+    else if(l->act_mode == ACTIVATION_SWISH) {
+        IPluginV2 *plugin = new ActivationSwishRT();
+        ILayer *lRT = networkRT->addPluginV2(&input, 1, *plugin);
         checkNULL(lRT);
         return lRT;
     }
@@ -1025,18 +1029,17 @@ bool NetworkRT::deserialize(const char *filename) {
 }
 
 #if NV_TENSORRT_MAJOR > 7
-void NetworkRT::destroy() {
-    delete contextRT;
-    if(builderActive) {
-        delete engineRT;
-        delete builderRT;
-    }
-}
+        void NetworkRT::destroy() {
+            delete contextRT;
+            if(builderActive) {
+                delete engineRT;
+                delete builderRT;
+            }
+        }
 #elif NV_TENSORRT_MAJOR <=7
-void NetworkRT::destroy() {
+        void NetworkRT::destroy() {
 
 }
 #endif
 
-
-}}
+}}

src/kernels/activation_swish.cu

@@ -0,0 +1,23 @@
+#include "kernels.h"
+
+__global__
+void activation_swish(dnnType *input, dnnType *output, int size) {
+
+    int i = blockDim.x*blockIdx.x + threadIdx.x;
+
+    if(i<size) {    
+        output[i] =  input[i] * 1.0f/(1.0f + exp(-input[i]));;
+    }
+ }
+
+
+/**
+    Swish activation function
+*/
+void activationSwishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    activation_swish<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+}

src/pluginsRT/ActivationSwishRT.cpp

@@ -0,0 +1,133 @@
+//
+// Created by Adam on 4/11/2022
+//
+#include <tkDNN/pluginsRT/ActivationSwishRT.h>
+using namespace  nvinfer1;
+std::vector<PluginField> ActivationSwishRTPluginCreator::mPluginAttributes;
+PluginFieldCollection ActivationSwishRTPluginCreator::mFC{};
+
+ActivationSwishRT::ActivationSwishRT() {
+
+}
+
+ActivationSwishRT::~ActivationSwishRT() {
+
+}
+
+ActivationSwishRT::ActivationSwishRT(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 ActivationSwishRT::getNbOutputs() const NOEXCEPT  { return 1; }
+
+Dims ActivationSwishRT::getOutputDimensions(int index, const Dims *inputs, int nbInputDims) NOEXCEPT  { return inputs[0]; }
+
+void ActivationSwishRT::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 ActivationSwishRT::initialize() NOEXCEPT  { return 0; }
+
+void ActivationSwishRT::terminate() NOEXCEPT {}
+
+size_t ActivationSwishRT::getWorkspaceSize(int maxBatchSize) const NOEXCEPT  { return 0; }
+
+#if NV_TENSORRT_MAJOR > 7
+int ActivationSwishRT::enqueue(int batchSize, const void *const *inputs, void *const *outputs, void *workspace,
+            cudaStream_t stream) NOEXCEPT {
+    activationSwishForward((dnnType *) reinterpret_cast<const dnnType *>(inputs[0]),
+    reinterpret_cast<dnnType *>(outputs[0]), batchSize * size, stream);
+    return 0;
+}
+#elif NV_TENSORRT_MAJOR == 7
+int32_t ActivationSwishRT::enqueue(int32_t batchSize, const void *const *inputs, void **outputs, void *workspace,
+                                  cudaStream_t stream) {
+    activationSwishForward((dnnType *) reinterpret_cast<const dnnType *>(inputs[0]),
+                          reinterpret_cast<dnnType *>(outputs[0]), batchSize * size, stream);
+    return 0;
+}
+#endif
+
+size_t ActivationSwishRT::getSerializationSize() const NOEXCEPT  {
+    return 1 * sizeof(int);
+}
+
+void ActivationSwishRT::serialize(void *buffer) const NOEXCEPT  {
+    char *buf = reinterpret_cast<char *>(buffer), *a = buf;
+    writeBUF(buf, size);
+    assert(buf == a + getSerializationSize());
+}
+
+const char* ActivationSwishRT::getPluginType() const NOEXCEPT  {
+    return "ActivationSwishRT_tkDNN";
+}
+
+const char *ActivationSwishRT::getPluginVersion() const NOEXCEPT  {
+    return "1";
+}
+
+bool ActivationSwishRT::supportsFormat(DataType type, PluginFormat format) const NOEXCEPT {
+        return (type == DataType::kFLOAT && format == PluginFormat::kLINEAR);
+}
+
+const char *ActivationSwishRT::getPluginNamespace() const NOEXCEPT  {
+    return mPluginNamespace.c_str();
+}
+
+void ActivationSwishRT::setPluginNamespace(const char *plguinNamespace) NOEXCEPT  {
+    mPluginNamespace = plguinNamespace;
+}
+
+IPluginV2 *ActivationSwishRT::clone() const NOEXCEPT {
+    auto *p = new ActivationSwishRT();
+    p->setPluginNamespace(mPluginNamespace.c_str());
+    return p;
+}
+
+
+
+ActivationSwishRTPluginCreator::ActivationSwishRTPluginCreator() {
+    mPluginAttributes.clear();
+    mFC.nbFields = mPluginAttributes.size();
+    mFC.fields = mPluginAttributes.data();
+}
+
+void ActivationSwishRTPluginCreator::setPluginNamespace(const char *pluginNamespace) NOEXCEPT {
+    mPluginNamespace = pluginNamespace;
+}
+
+const char *ActivationSwishRTPluginCreator::getPluginNamespace() const NOEXCEPT  {
+    return mPluginNamespace.c_str();
+}
+
+IPluginV2 *ActivationSwishRTPluginCreator::deserializePlugin(const char *name, const void *serialData, size_t serialLength) NOEXCEPT  {
+    auto *pluginObj = new ActivationSwishRT(serialData, serialLength);
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+IPluginV2 *ActivationSwishRTPluginCreator::createPlugin(const char *name, const PluginFieldCollection *fc) NOEXCEPT  {
+    const PluginField *fields = fc->fields;
+    auto *pluginObj = new ActivationSwishRT();
+    pluginObj->setPluginNamespace(mPluginNamespace.c_str());
+    return pluginObj;
+}
+
+const char *ActivationSwishRTPluginCreator::getPluginName() const NOEXCEPT  {
+    return "ActivationSwishRT_tkDNN";
+}
+
+const char *ActivationSwishRTPluginCreator::getPluginVersion() const NOEXCEPT{
+    return "1";
+}
+
+const PluginFieldCollection *ActivationSwishRTPluginCreator::getFieldNames() NOEXCEPT  {
+    return &mFC;
+}
+

tests/darknet/yolov4-csp-swish.cpp

@@ -0,0 +1,40 @@
+//
+// Created by Adam T. Cuellar on 9/15/21.
+//
+
+#include<iostream>
+#include<vector>
+#include "tkdnn.h"
+#include "test.h"
+#include "DarknetParser.h"
+
+int main() {
+    std::string bin_path  = "yolov4-csp-swish";
+    std::vector<std::string> input_bins = {
+            bin_path + "/layers/input.bin"
+    };
+    std::vector<std::string> output_bins = {
+            bin_path + "/debug/layer167_out.bin",
+            bin_path + "/debug/layer171_out.bin",
+            bin_path + "/debug/layer175_out.bin"
+    };
+    std::string wgs_path  = bin_path + "/layers";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolov4-csp-swish.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
+    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download");
+
+
+
+    // parse darknet network
+    tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
+    net->print();
+
+    //convert network to tensorRT
+    tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str()));
+
+    int ret = testInference(input_bins, output_bins, net, netRT);
+    net->releaseLayers();
+    delete net;
+    delete netRT;
+    return ret;
+}