added dilated convolution

include/tkDNN/DarknetParser.h

@@ -24,6 +24,7 @@ namespace tk { namespace dnn {
         int pad = 0;
         int coords = 4;
         float scale_xy = 1;
+        int dilation = 1;
         std::vector<int> layers;
         std::string activation = "linear";
 
@@ -115,6 +116,8 @@ namespace tk { namespace dnn {
             auto vec = fromStringToIntVec(value, ',');
             fields.n_mask = vec.size();
         }
+        else if(name.find("dilation") !=  std::string::npos)
+            fields.dilation = std::stoi(value);
         else if(name.find("layers") !=  std::string::npos)
             fields.layers = fromStringToIntVec(value, ',');
 
@@ -143,7 +146,7 @@ namespace tk { namespace dnn {
             std::string wgs = wgs_path + "/c" + std::to_string(netLayers.size()) + ".bin";
             //printf("%d (%d,%d) (%d,%d) (%d,%d) %s %d %d\n", f.filters, f.size_x, f.size_y, f.stride_x, f.stride_y, f.padding_x, f.padding_y, wgs.c_str(), f.batch_normalize, f.groups);
             tk::dnn::Conv2d *l= new tk::dnn::Conv2d(net, f.filters, f.size_x, f.size_y, f.stride_x, 
-                                f.stride_y, f.padding_x, f.padding_y, wgs, f.batch_normalize, false, f.groups);
+                                f.stride_y, f.padding_x, f.padding_y, f.dilation, f.dilation, wgs, f.batch_normalize, false, f.groups);
             netLayers.push_back(l);
         } else if(f.type == "maxpool") {
             if(f.stride_x == 1 && f.stride_y == 1)

include/tkDNN/Layer.h

@@ -261,6 +261,10 @@ public:
     Conv2d( Network *net, int out_ch, int kernelH, int kernelW, 
                 int strideH, int strideW, int paddingH, int paddingW,
                 std::string fname_weights, bool batchnorm = false, bool deConv = false, int groups = 1, bool additional_bias=false);
+    Conv2d( Network *net, int out_ch, int kernelH, int kernelW,
+                int strideH, int strideW, int paddingH, int paddingW,
+                int dilationW, int dilationH,
+                std::string fname_weights, bool batchnorm = false, bool deConv = false, int groups = 1, bool additional_bias=false);
     virtual ~Conv2d();
     virtual layerType_t getLayerType() { return LAYER_CONV2D; };
 
@@ -269,6 +273,7 @@ public:
     int kernelH, kernelW, strideH, strideW, paddingH, paddingW;
     bool deConv, additional_bias;
     int groups;
+    int dilationW, dilationH;
 
 protected:
     cudnnFilterDescriptor_t filterDesc;

src/Conv2d.cpp

@@ -31,9 +31,9 @@ void Conv2d::initCUDNN(bool back) {
                                            kernelH, kernelW) );
 
     checkCUDNN( cudnnSetConvolution2dDescriptor(convDesc,
-                                                paddingH, paddingW, // padding
+                                                paddingH * dilationH, paddingW * dilationW, // padding
                                                 strideH, strideW, // stride
-                                                1,1, // upscale
+                                                dilationH, dilationW, // upscale
                                                 CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT) );
 
     checkCUDNN(  cudnnSetConvolutionGroupCount(convDesc,
@@ -146,6 +146,50 @@ Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW,
     this->groups = groups;
     this->additional_bias = additional_bias;
 
+    this->dilationW = 1;
+    this->dilationH = 1;
+
+    if(!deConv) {
+        output_dim.n = input_dim.n;
+        output_dim.c = out_ch;
+        output_dim.h = (input_dim.h + 2 * paddingH - kernelH) / strideH + 1;
+        output_dim.w = (input_dim.w + 2 * paddingW - kernelW) / strideW + 1;
+        output_dim.l = 1;
+    } else {
+        output_dim.n = input_dim.n;
+        output_dim.c = out_ch;
+        output_dim.h = ((input_dim.h-1) * strideH) - 2*paddingH + kernelH;
+        output_dim.w = ((input_dim.w-1) * strideW) - 2*paddingW + kernelW;
+        output_dim.l = 1;
+    }
+    initCUDNN(deConv);
+
+    // allocate warkspace
+    if (ws_sizeInBytes!=0) {
+        checkCuda( cudaMalloc(&workSpace, ws_sizeInBytes) );
+    }
+
+    //allocate data for infer result
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
+}
+
+Conv2d::Conv2d(Network *net, int out_ch, int kernelH, int kernelW, int strideH, int strideW, int paddingH, int paddingW, int dilationX, int dilationY, std::string fname_weights, bool batchnorm, bool deConv, int groups, bool additional_bias)
+    :LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1,
+              fname_weights, batchnorm, additional_bias, deConv, groups)
+{
+    this->dilationW = dilationX;
+    this->dilationH = dilationY;
+
+    this->kernelH = kernelH;
+    this->kernelW = kernelW;
+    this->strideH = strideH;
+    this->strideW = strideW;
+    this->paddingH = paddingH;
+    this->paddingW = paddingW;
+    this->deConv = deConv;
+    this->groups = groups;
+    this->additional_bias = additional_bias;
+
     if(!deConv) {
         output_dim.n = input_dim.n;
         output_dim.c = out_ch;

src/NetworkRT.cpp

@@ -317,6 +317,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
         checkNULL(lRTconv);
         lRTconv->setStride(DimsHW{l->strideH, l->strideW});
         lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW});
+        //std::cout << "dilation: " << l->dilationH << ", " << l->dilationW << std::endl;
+        lRTconv->setDilation(DimsHW{l->dilationH, l->dilationW}); //mrho
         lRTconv->setNbGroups(l->groups);
         lRT = (ILayer*) lRTconv;
     } else {

tests/darknet/yolo3_tiny_dilation.cpp

@@ -0,0 +1,33 @@
+#include<iostream>
+#include<vector>
+#include "tkdnn.h"
+#include "test.h"
+#include "DarknetParser.h"
+
+int main() {
+    std::string bin_path  = "dilation";
+    std::vector<std::string> input_bins = {
+        bin_path + "/layers/input.bin"
+    };
+    std::vector<std::string> output_bins = {
+        bin_path + "/debug/layer33_out.bin",
+        bin_path + "/debug/layer45_out.bin",
+    };
+    std::string wgs_path  = bin_path + "/layers";
+    std::string cfg_path  = "config.cfg";
+    std::string name_path = "config.names";
+    //downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/LMcSHtWaLeps8yN/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;
+}