mmr/tkDNN
Archived
1
0
Fork 0
Code Issues 46 Pull Requests 12 Releases 4 Activity

Adding dilated convolution and support for reorg3d layer #47

Open
mrhosseini wants to merge 8 commits from mrhosseini/master into master
pull from: mrhosseini/master
merge into: mmr:master
Conversation 1 Commits 8 Files Changed 8
+113 -15
8 changed files with 113 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/tkDNN/DarknetParser.h
+1
View File
@@ -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";
include/tkDNN/Layer.h
+7 -1
View File
@@ -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;
@@ -529,13 +534,14 @@ public:
class Reorg : public Layer {
public:
Reorg(Network *net, int stride);
Reorg(Network *net, int stride, bool reorg3d = false);
virtual ~Reorg();
virtual layerType_t getLayerType() { return LAYER_REORG; };
virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
int stride;
bool reorg3d;
};
/**
include/tkDNN/pluginsRT/ReorgRT.h
+12 -4
View File
@@ -4,8 +4,9 @@
class ReorgRT : public IPlugin {
public:
ReorgRT(int stride) {
ReorgRT(int stride, bool reorg3d = false) {
this->stride = stride;
this->reorg3d = reorg3d;
}
~ReorgRT(){
@@ -21,9 +22,15 @@ public:
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
if (reorg3d) {
c = outputDims[0].d[0];
h = outputDims[0].d[1];
w = outputDims[0].d[2];
} else {
c = inputDims[0].d[0];
h = inputDims[0].d[1];
w = inputDims[0].d[2];
}
}
int initialize() override {
@@ -60,4 +67,5 @@ public:
}
int c, h, w, stride;
bool reorg3d;
};
src/Conv2d.cpp
+46 -2
View File
@@ -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/DarknetParser.cpp
+5 -3
View File
@@ -83,6 +83,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, ',');
@@ -111,7 +113,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)
@@ -150,8 +152,8 @@ namespace tk { namespace dnn {
}
netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size()));
} else if(f.type == "reorg") {
netLayers.push_back(new tk::dnn::Reorg(net, f.stride_x));
} else if(f.type == "reorg" || f.type == "reorg3d") {
netLayers.push_back(new tk::dnn::Reorg(net, f.stride_x, f.type == "reorg3d"));
} else if(f.type == "region") {
netLayers.push_back(new tk::dnn::Region(net, f.classes, f.coords, f.num));
src/NetworkRT.cpp
+3 -2
View File
@@ -316,7 +316,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b);
checkNULL(lRTconv);
lRTconv->setStride(DimsHW{l->strideH, l->strideW});
lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW});
lRTconv->setPadding(DimsHW{l->paddingH * l->dilationH, l->paddingW * l->dilationW});
lRTconv->setDilation(DimsHW{l->dilationH, l->dilationW});
lRTconv->setNbGroups(l->groups);
lRT = (ILayer*) lRTconv;
} else {
@@ -480,7 +481,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Reorg *l) {
//std::cout<<"convert Reorg\n";
//std::cout<<"New plugin REORG\n";
IPlugin *plugin = new ReorgRT(l->stride);
IPlugin *plugin = new ReorgRT(l->stride, l->reorg3d);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
return lRT;
src/Reorg.cpp
+6 -3
View File
@@ -5,9 +5,10 @@
namespace tk { namespace dnn {
Reorg::Reorg(Network *net, int stride) : Layer(net) {
Reorg::Reorg(Network *net, int stride, bool reorg3d) : Layer(net) {
this->stride = stride;
this->reorg3d = reorg3d;
output_dim.n = input_dim.n;
output_dim.c = input_dim.c*stride*stride;
@@ -24,8 +25,10 @@ Reorg::~Reorg() {
}
dnnType* Reorg::infer(dataDim_t &dim, dnnType* srcData) {
reorgForward(srcData, dstData, dim.n, dim.c, dim.h, dim.w, stride);
if (reorg3d)
reorgForward(srcData, dstData, output_dim.n, output_dim.c, output_dim.h, output_dim.w, stride);
else
reorgForward(srcData, dstData, dim.n, dim.c, dim.h, dim.w, stride);
dim = output_dim;
return dstData;
tests/darknet/yolo3_tiny_dilation.cpp
+33
View File
@@ -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;
}