yolov4tiny works on CUDNN

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

include/tkDNN/DarknetParser.h

@@ -11,6 +11,7 @@ namespace tk { namespace dnn {
         int channels = 3;
         int batch_normalize=0;
         int groups = 1;
+        int group_id = 0;
         int filters=1;
         int size_x=1;
         int size_y=1;

include/tkDNN/Layer.h

@@ -509,7 +509,7 @@ public:
 class Route : public Layer {
 
 public:
-    Route(Network *net, Layer **layers, int layers_n); 
+    Route(Network *net, Layer **layers, int layers_n, int groups = 1, int group_id = 0); 
     virtual ~Route();
     virtual layerType_t getLayerType() { return LAYER_ROUTE; };
 
@@ -519,6 +519,8 @@ public:
     static const int MAX_LAYERS = 32;
     Layer *layers[MAX_LAYERS];  //ids of layers to be merged
     int layers_n; //number of layers
+    int groups;
+    int group_id;
 };
 
 

src/DarknetParser.cpp

@@ -75,6 +75,8 @@ namespace tk { namespace dnn {
             fields.coords = std::stoi(value);
         else if(name.find("groups") !=  std::string::npos)
             fields.groups = std::stoi(value);
+        else if(name.find("group_id") !=  std::string::npos)
+            fields.group_id = std::stoi(value);
         else if(name.find("scale_x_y") !=  std::string::npos)
             fields.scale_xy = std::stof(value);
         else if(name.find("from") !=  std::string::npos)
@@ -148,7 +150,7 @@ namespace tk { namespace dnn {
                 //std::cout<<"Route to "<<layerIdx<<" "<<netLayers[layerIdx]->getLayerName()<<"\n";
                 layers.push_back(netLayers[layerIdx]);
             }
-            netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size()));
+            netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size(), f.groups, f.group_id));
 
         } else if(f.type == "reorg") {
             netLayers.push_back(new tk::dnn::Reorg(net, f.stride_x));

src/Route.cpp

@@ -5,7 +5,7 @@
 
 namespace tk { namespace dnn {
 
-Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) {
+Route::Route(Network *net, Layer **layers, int layers_n, int groups, int group_id) : Layer(net) {
 
     // copy input layers
     if(layers_n > MAX_LAYERS) {
@@ -15,6 +15,8 @@ Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) {
         this->layers[i] = layers[i];
     }
     this->layers_n = layers_n;
+    this->groups = groups;
+    this->group_id = group_id;
     
     //get dims
     output_dim.l = 1;
@@ -32,6 +34,7 @@ Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) {
         output_dim.c += layers[i]->output_dim.c;
     }
 
+    output_dim.c /= this->groups;
     input_dim = output_dim;
 
     checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
@@ -49,8 +52,9 @@ dnnType* Route::infer(dataDim_t &dim, dnnType* srcData) {
     for(int i=0; i<layers_n; i++) {
         dnnType *input = layers[i]->dstData;
         int in_dim = layers[i]->output_dim.tot();
-        checkCuda( cudaMemcpy(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+        int part_in_dim = in_dim / this->groups;
-        offset += in_dim;
+        checkCuda( cudaMemcpy(dstData + offset, input + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+        offset += part_in_dim;
     }
 
     //update data dimensions    

tests/darknet/cfg/yolo4tiny.cfg

@@ -0,0 +1,281 @@
+[net]
+# Testing
+#batch=1
+#subdivisions=1
+# Training
+batch=64
+subdivisions=1
+width=416
+height=416
+channels=3
+momentum=0.9
+decay=0.0005
+angle=0
+saturation = 1.5
+exposure = 1.5
+hue=.1
+
+learning_rate=0.00261
+burn_in=1000
+max_batches = 500200
+policy=steps
+steps=400000,450000
+scales=.1,.1
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers=-1
+groups=2
+group_id=1
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -1,-2
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -6,-1
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers=-1
+groups=2
+group_id=1
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -1,-2
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -6,-1
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers=-1
+groups=2
+group_id=1
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -1,-2
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -6,-1
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+##################################
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=255
+activation=linear
+
+
+
+[yolo]
+mask = 3,4,5
+anchors = 10,14,  23,27,  37,58,  81,82,  135,169,  344,319
+classes=80
+num=6
+jitter=.3
+scale_x_y = 1.05
+cls_normalizer=1.0
+iou_normalizer=0.07
+iou_loss=ciou
+ignore_thresh = .7
+truth_thresh = 1
+random=0
+resize=1.5
+nms_kind=greedynms
+beta_nms=0.6
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 23
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=255
+activation=linear
+
+[yolo]
+mask = 1,2,3
+anchors = 10,14,  23,27,  37,58,  81,82,  135,169,  344,319
+classes=80
+num=6
+jitter=.3
+scale_x_y = 1.05
+cls_normalizer=1.0
+iou_normalizer=0.07
+iou_loss=ciou
+ignore_thresh = .7
+truth_thresh = 1
+random=0
+resize=1.5
+nms_kind=greedynms
+beta_nms=0.6

tests/darknet/yolo4tiny.cpp

@@ -0,0 +1,33 @@
+#include<iostream>
+#include<vector>
+#include "tkdnn.h"
+#include "test.h"
+#include "DarknetParser.h"
+
+int main() {
+    std::string bin_path  = "yolo4tiny";
+    std::vector<std::string> input_bins = { 
+        bin_path + "/layers/input.bin"
+    };
+    std::vector<std::string> output_bins = {
+        bin_path + "/debug/layer30_out.bin",
+        bin_path + "/debug/layer37_out.bin"
+    };
+    std::string wgs_path  = bin_path + "/layers";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4tiny.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/iRnc4pSqmx78gJs/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, nullptr);
+    net->releaseLayers();
+    delete net;
+    // delete netRT;
+    return ret;
+}