From c51675bf44274148e6bee2aef2d3068e53bb32b3 Mon Sep 17 00:00:00 2001 From: MohammadReza Hosseini Date: Sun, 14 Jun 2020 18:49:10 +0430 Subject: [PATCH] added dilated convolution --- include/tkDNN/DarknetParser.h | 5 ++- include/tkDNN/Layer.h | 5 +++ src/Conv2d.cpp | 48 +++++++++++++++++++++++++-- src/NetworkRT.cpp | 2 ++ tests/darknet/yolo3_tiny_dilation.cpp | 33 ++++++++++++++++++ 5 files changed, 90 insertions(+), 3 deletions(-) create mode 100644 tests/darknet/yolo3_tiny_dilation.cpp diff --git a/include/tkDNN/DarknetParser.h b/include/tkDNN/DarknetParser.h index b4c9253..31830fb 100644 --- a/include/tkDNN/DarknetParser.h +++ b/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 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) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 9bd8432..1380c1d 100644 --- a/include/tkDNN/Layer.h +++ b/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; diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index 4704c66..43dcfd7 100644 --- a/src/Conv2d.cpp +++ b/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; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 6e86de1..af0b5ab 100644 --- a/src/NetworkRT.cpp +++ b/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 { diff --git a/tests/darknet/yolo3_tiny_dilation.cpp b/tests/darknet/yolo3_tiny_dilation.cpp new file mode 100644 index 0000000..d4a1814 --- /dev/null +++ b/tests/darknet/yolo3_tiny_dilation.cpp @@ -0,0 +1,33 @@ +#include +#include +#include "tkdnn.h" +#include "test.h" +#include "DarknetParser.h" + +int main() { + std::string bin_path = "dilation"; + std::vector input_bins = { + bin_path + "/layers/input.bin" + }; + std::vector 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; +}