added yolo4x and yolo4-csp from the github repo and download file corrections

include/tkDNN/Layer.h

@@ -19,6 +19,7 @@ enum layerType_t {
     LAYER_ACTIVATION_CRELU,
     LAYER_ACTIVATION_LEAKY,
     LAYER_ACTIVATION_MISH,
+    LAYER_ACTIVATION_LOGISTIC,
     LAYER_FLATTEN,
     LAYER_RESHAPE,
     LAYER_MULADD,
@@ -68,6 +69,7 @@ public:
             case LAYER_ACTIVATION_CRELU:    return "ActivationCReLU";
             case LAYER_ACTIVATION_LEAKY:    return "ActivationLeaky";
             case LAYER_ACTIVATION_MISH:     return "ActivationMish";
+            case LAYER_ACTIVATION_LOGISTIC: return "ActivationLogistic";
             case LAYER_FLATTEN:             return "Flatten";
             case LAYER_RESHAPE:             return "Reshape";
             case LAYER_MULADD:              return "MulAdd";
@@ -212,7 +214,8 @@ public:
 typedef enum {
     ACTIVATION_ELU     = 100,
     ACTIVATION_LEAKY   = 101,
-    ACTIVATION_MISH   = 102
+    ACTIVATION_MISH   = 102,
+    ACTIVATION_LOGISTIC   = 103
 } tkdnnActivationMode_t;
 
 /**
@@ -233,6 +236,8 @@ public:
             return LAYER_ACTIVATION_LEAKY;
         else if (act_mode == ACTIVATION_MISH)
             return LAYER_ACTIVATION_MISH;
+        else if (act_mode == ACTIVATION_LOGISTIC)
+            return LAYER_ACTIVATION_LOGISTIC;
         else
             return LAYER_ACTIVATION;
          };

include/tkDNN/NetworkRT.h

@@ -27,6 +27,7 @@ using namespace nvinfer1;
 #include "pluginsRT/ActivationLeakyRT.h"
 #include "pluginsRT/ActivationReLUCeilingRT.h"
 #include "pluginsRT/ActivationMishRT.h"
+#include "pluginsRT/ActivationLogisticRT.h"
 #include "pluginsRT/ReorgRT.h"
 #include "pluginsRT/RegionRT.h"
 #include "pluginsRT/RouteRT.h"

include/tkDNN/pluginsRT/ActivationLogisticRT.h

@@ -0,0 +1,60 @@
+#include<cassert>
+#include "../kernels.h"
+
+class ActivationLogisticRT : public IPlugin {
+
+public:
+    ActivationLogisticRT() {
+
+
+    }
+
+    ~ActivationLogisticRT(){
+
+    }
+
+    int getNbOutputs() const override {
+        return 1;
+    }
+
+    Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
+        return inputs[0];
+    }
+
+    void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+        size = 1;
+        for(int i=0; i<outputDims[0].nbDims; i++)
+            size *= outputDims[0].d[i];
+    }
+
+    int initialize() override {
+
+        return 0;
+    }
+
+    virtual void terminate() override {
+    }
+
+    virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+        return 0;
+    }
+
+    virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+
+        activationLOGISTICForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
+                                  reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
+        return 0;
+    }
+
+
+    virtual size_t getSerializationSize() override {
+        return 1*sizeof(int);
+    }
+
+    virtual void serialize(void* buffer) override {
+        char *buf = reinterpret_cast<char*>(buffer);
+        tk::dnn::writeBUF(buf, size);
+    }
+
+    int size;
+};

include/tkDNN/pluginsRT/YoloRT.h

@@ -64,20 +64,23 @@ public:
 
 		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)
-					activationLOGISTICForward(srcData + index, dstData + index, 4*w*h, stream); //x,y,w,h
-				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);
+        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
 
-				index = entry_index(b, n*w*h, 4);
+                    if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
-				activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
+
-			}
+                    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;

src/Activation.cpp

@@ -52,7 +52,11 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
     else if(act_mode == ACTIVATION_MISH) {
         activationMishForward(srcData, dstData, dim.tot());
 
-    } else {
+    }
+    else if(act_mode == ACTIVATION_LOGISTIC) {
+        activationLOGISTICForward(srcData, dstData, dim.tot());
+
+    }else {
         dnnType alpha = dnnType(1);
         dnnType beta  = dnnType(0);
         checkCUDNN( cudnnActivationForward(net->cudnnHandle,

src/DarknetParser.cpp

@@ -187,6 +187,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 == "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

@@ -229,7 +229,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)
+    if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_LOGISTIC)
         return convert_layer(input, (Activation*) l);
     if(type == LAYER_SOFTMAX)
         return convert_layer(input, (Softmax*) l);
@@ -424,6 +424,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
         checkNULL(lRT);
         return lRT;
     }
+    else if(l->act_mode == ACTIVATION_LOGISTIC) {
+        IPlugin *plugin = new ActivationLogisticRT();
+        IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+        checkNULL(lRT);
+        return lRT;
+    }
     else {
         FatalError("this Activation mode is not yet implemented");
         return NULL;
@@ -660,6 +666,11 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         assert(buf == bufCheck + serialLength);
         return a;
     }
+    if(name.find("ActivationLogistic") == 0) {
+        ActivationLogisticRT *a = new ActivationLogisticRT();
+        a->size = readBUF<int>(buf);
+        return a;
+    }
     if(name.find("ActivationCReLU") == 0) {
         float activationReluTemp = readBUF<float>(buf);
         //ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF<float>(buf));
@@ -784,17 +795,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         float nms_thresh_temp = readBUF<float>(buf);
         int nms_kind_temp = readBUF<int>(buf);
         int new_coords_temp = readBUF<int>(buf);
-        std::cout << classes_temp << ":" << num_temp << ":" << ":" << n_masks_temp << ":" << nms_thresh_temp << ":" << nms_kind_temp << ":" << new_coords_temp << std::endl;
 
        YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp);  
 
-       /* std::cout << "classes : " << r->classes;
-        std::cout << "num : " << r->num;
-        std::cout << "n_masks : " << r->n_masks;
-        std::cout << "scalexy : " << r->scaleXY;
-        std::cout << "nms_thresh : " << r->nms_thresh;
-        std::cout << "nms_kind : " << r->nms_kind;
-        std::cout << "new_coords : " << r->new_coords;*/
 
 
         r->c = readBUF<int>(buf);

src/Yolo.cpp

@@ -73,8 +73,8 @@ Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j,
         b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h;
     }
     else{
-        b.x = (i + x[index + 0 * stride] * 2 - 0.5) / lw;
+        b.x = (i + x[index + 0 * stride] ) / lw;
-        b.y = (j + x[index + 1 * stride] * 2 - 0.5) / lh;
+        b.y = (j + x[index + 1 * stride] ) / lh;
         b.w = x[index + 2 * stride] * x[index + 2 * stride] * 4 * biases[2 * n] / w;
         b.h = x[index + 3 * stride] * x[index + 3 * stride] * 4 * biases[2 * n + 1] / h;
     }
@@ -88,15 +88,18 @@ dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) {
     for (int b = 0; b < dim.n; ++b){
         for(int n = 0; n < n_masks; ++n){
             int index = entry_index(b, n*dim.w*dim.h, 0, classes, input_dim, output_dim);
-            if (new_coords == 1)
+            std::cout<<"new_coords"<<new_coords<<std::endl;
-                activationLOGISTICForward(srcData + index, dstData + index, 4*dim.w*dim.h);
+            if (new_coords == 1){
-            else
+                if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
+            }
+            else{
                 activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
 
-            if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
+                if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
 
-            index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim);
+                index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim);
-            activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h);
+                activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h);
+            }
         }
     }
 

tests/darknet/cfg/yolo4x.cfg

@@ -5,8 +5,8 @@
 # Training
 batch=64
 subdivisions=8
-width=672
+width=640
-height=672
+height=640
 channels=3
 momentum=0.949
 decay=0.0005
@@ -15,7 +15,7 @@ saturation = 1.5
 exposure = 1.5
 hue=.1
 
-learning_rate=0.00261
+learning_rate=0.001
 burn_in=1000
 max_batches = 500500
 policy=steps
@@ -26,6 +26,8 @@ mosaic=1
 
 letter_box=1
 
+#optimized_memory=1
+
 [convolutional]
 batch_normalize=1
 filters=32
@@ -1131,6 +1133,7 @@ size=1
 stride=1
 pad=1
 activation=mish
+stopbackward=800
 
 ##########################
 
@@ -1147,7 +1150,7 @@ size=1
 stride=1
 pad=1
 filters=255
-activation=linear
+activation=logistic
 
 
 [yolo]
@@ -1156,6 +1159,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4
 classes=80
 num=9
 jitter=.1
+scale_x_y = 2.0
 objectness_smooth=0
 ignore_thresh = .7
 truth_thresh = 1
@@ -1169,6 +1173,7 @@ iou_loss=ciou
 nms_kind=diounms
 beta_nms=0.6
 new_coords=1
+max_delta=5
 
 [route]
 layers = -4
@@ -1275,7 +1280,7 @@ size=1
 stride=1
 pad=1
 filters=255
-activation=linear
+activation=logistic
 
 
 [yolo]
@@ -1284,6 +1289,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4
 classes=80
 num=9
 jitter=.1
+scale_x_y = 2.0
 objectness_smooth=1
 ignore_thresh = .7
 truth_thresh = 1
@@ -1297,6 +1303,7 @@ iou_loss=ciou
 nms_kind=diounms
 beta_nms=0.6
 new_coords=1
+max_delta=5
 
 [route]
 layers = -4
@@ -1403,7 +1410,7 @@ size=1
 stride=1
 pad=1
 filters=255
-activation=linear
+activation=logistic
 
 
 [yolo]
@@ -1412,6 +1419,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4
 classes=80
 num=9
 jitter=.1
+scale_x_y = 2.0
 objectness_smooth=1
 ignore_thresh = .7
 truth_thresh = 1
@@ -1425,3 +1433,4 @@ iou_loss=ciou
 nms_kind=diounms
 beta_nms=0.6
 new_coords=1
+max_delta=2

tests/darknet/yolo4-csp.cpp

@@ -0,0 +1,36 @@
+#include<iostream>
+#include<vector>
+#include "tkdnn.h"
+#include "test.h"
+#include "DarknetParser.h"
+
+int main() {
+    std::string bin_path  = "yolo4-csp";
+    std::vector<std::string> input_bins = {
+            bin_path + "/layers/input.bin"
+    };
+    std::vector<std::string> output_bins = {
+            bin_path + "/debug/layer144_out.bin",
+            bin_path + "/debug/layer159_out.bin",
+            bin_path + "/debug/layer174_out.bin"
+    };
+    std::string wgs_path  = bin_path + "/layers";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4-csp.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/AfzHE4BfTeEm2gH/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;
+}

tests/darknet/yolo4x.cpp

@@ -17,7 +17,7 @@ int main() {
     std::string wgs_path  = bin_path + "/layers";
     std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4x.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/BLPpiAigZJLorQD/download");
+    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download");