From f9afee2f3b16536cdcc401c3d84311284db4ff25 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Wed, 30 Oct 2019 00:05:22 +0100 Subject: [PATCH 1/3] deconv layer cudnn --- include/tkDNN/Layer.h | 30 +++++- include/tkDNN/Network.h | 1 + include/tkDNN/utils.h | 2 +- src/Conv2d.cpp | 218 ++++++++++++++++++++++++++-------------- src/LayerWgs.cpp | 10 +- src/Network.cpp | 1 + src/utils.cpp | 37 +++---- 7 files changed, 197 insertions(+), 102 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index bab723d..483b7b5 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -11,6 +11,7 @@ namespace tk { namespace dnn { enum layerType_t { LAYER_DENSE, LAYER_CONV2D, + LAYER_DECONV2D, LAYER_ACTIVATION, LAYER_FLATTEN, LAYER_MULADD, @@ -47,6 +48,7 @@ public: switch(type) { case LAYER_DENSE: return "Dense"; case LAYER_CONV2D: return "Conv2d"; + case LAYER_DECONV2D: return "DeConv2d"; case LAYER_ACTIVATION: return "Activation"; case LAYER_FLATTEN: return "Flatten"; case LAYER_MULADD: return "MulAdd"; @@ -75,7 +77,7 @@ protected: class LayerWgs : public Layer { public: - LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kt, + LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kt, std::string fname_weights, bool batchnorm = false); virtual ~LayerWgs(); @@ -152,7 +154,7 @@ class Conv2d : public LayerWgs { 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); + std::string fname_weights, bool batchnorm = false, bool deConv = false); virtual ~Conv2d(); virtual layerType_t getLayerType() { return LAYER_CONV2D; }; @@ -163,11 +165,33 @@ public: protected: cudnnFilterDescriptor_t filterDesc; cudnnConvolutionDescriptor_t convDesc; - cudnnConvolutionFwdAlgo_t algo; + cudnnConvolutionFwdAlgo_t fwAlgo; + cudnnConvolutionBwdDataAlgo_t bwAlgo; cudnnTensorDescriptor_t biasTensorDesc; + void initCUDNN(bool back = false); + void inferCUDNN(dnnType* srcData, bool back = false); void* workSpace; size_t ws_sizeInBytes; + + bool deConv; +}; + + +/** + Convolutional 2D layer +*/ +class DeConv2d : public Conv2d { + +public: + DeConv2d( Network *net, int out_ch, int kernelH, int kernelW, + int strideH, int strideW, int paddingH, int paddingW, + std::string fname_weights, bool batchnorm = false) : + Conv2d(net, out_ch, kernelH, kernelW, strideH, strideW, paddingH, paddingW, fname_weights, batchnorm, true) {} + virtual ~DeConv2d() {} + virtual layerType_t getLayerType() { return LAYER_DECONV2D; }; + + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); }; diff --git a/include/tkDNN/Network.h b/include/tkDNN/Network.h index 9da5806..c56bd2c 100644 --- a/include/tkDNN/Network.h +++ b/include/tkDNN/Network.h @@ -60,6 +60,7 @@ public: dataDim_t getOutputDim(); bool fp16, dla; + bool dontLoadWeights; }; }} diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index dc34a31..f0503cd 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -90,7 +90,7 @@ void printCenteredTitle(const char *title, char fill, int dim); bool fileExist(const char *fname); -void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0); +void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0, bool skipLoad = false); int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true); void printDeviceVector(int size, dnnType* vec_d, bool device = true); void resize(int size, dnnType **data); diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index 3dfdce6..da1a4a2 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -4,9 +4,121 @@ namespace tk { namespace dnn { +void Conv2d::initCUDNN(bool back) { + + cudnnTensorDescriptor_t srcTensor = srcTensorDesc; + cudnnTensorDescriptor_t dstTensor = dstTensorDesc; + + dataDim_t idim, odim; + if(!back) { + idim = input_dim; + odim = output_dim; + } else { + idim = output_dim; + odim = input_dim; + } + idim.print(); + odim.print(); + + checkCUDNN( cudnnCreateFilterDescriptor(&filterDesc) ); + checkCUDNN( cudnnCreateConvolutionDescriptor(&convDesc) ); + checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) ); + + // input tensor dim + checkCUDNN( cudnnSetTensor4dDescriptor(srcTensor, + net->tensorFormat, net->dataType, idim.n, idim.c, idim.h, idim.w) ); + + checkCUDNN( cudnnSetFilter4dDescriptor(filterDesc, + net->dataType, net->tensorFormat, odim.c, idim.c, + kernelH, kernelW) ); + + checkCUDNN( cudnnSetConvolution2dDescriptor(convDesc, + paddingH, paddingW, // padding + strideH, strideW, // stride + 1,1, // upscale + CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT) ); + + // check dimension of convolution output + dataDim_t tmpdim; + checkCUDNN( cudnnGetConvolution2dForwardOutputDim( + convDesc, srcTensor, filterDesc, + &tmpdim.n, &tmpdim.c, &tmpdim.h, &tmpdim.w) ); + if(odim.n != tmpdim.n || odim.c != tmpdim.c || odim.h != tmpdim.h || odim.w != tmpdim.w) { + std::cout<<"tkdim: "; odim.print(); + std::cout<<"cudnndim: "; tmpdim.print(); + FatalError("Eror conv dimension mismatch"); + } + + checkCUDNN( cudnnSetTensor4dDescriptor(dstTensor, + net->tensorFormat, net->dataType, odim.n, odim.c, odim.h, odim.w) ); + + checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc, + net->tensorFormat, net->dataType, + 1, output_dim.c, 1, 1) ); + + // init workspace + workSpace = NULL; + ws_sizeInBytes = 0; + if(back) { + checkCUDNN( cudnnGetConvolutionBackwardDataAlgorithm(net->cudnnHandle, + filterDesc, dstTensor, convDesc, srcTensor, + CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST, 0, &bwAlgo) ); + checkCUDNN(cudnnGetConvolutionBackwardDataWorkspaceSize(net->cudnnHandle, + filterDesc, dstTensor, convDesc, srcTensor, + bwAlgo, &ws_sizeInBytes)); + + // invert tensors + srcTensorDesc = dstTensor; + dstTensorDesc = srcTensor; + } else { + checkCUDNN( cudnnGetConvolutionForwardAlgorithm(net->cudnnHandle, + srcTensor, filterDesc, convDesc, dstTensor, + CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &fwAlgo) ); + checkCUDNN(cudnnGetConvolutionForwardWorkspaceSize(net->cudnnHandle, + srcTensor, filterDesc, convDesc, dstTensor, + fwAlgo, &ws_sizeInBytes)); + } +} + +void Conv2d::inferCUDNN(dnnType* srcData, bool back) { + + dnnType alpha = dnnType(1); + dnnType beta = dnnType(0); + if(back) { + checkCUDNN(cudnnConvolutionBackwardData(net->cudnnHandle, + &alpha, filterDesc, data_d, + srcTensorDesc, srcData, + convDesc, bwAlgo, workSpace, ws_sizeInBytes, + &beta, dstTensorDesc, dstData)); + } else { + checkCUDNN(cudnnConvolutionForward(net->cudnnHandle, + &alpha, srcTensorDesc, srcData, filterDesc, + data_d, convDesc, fwAlgo, workSpace, ws_sizeInBytes, + &beta, dstTensorDesc, dstData)); + } + + if(!batchnorm) { + // bias + alpha = dnnType(1); + beta = dnnType(0); + checkCUDNN( cudnnAddTensor(net->cudnnHandle, + &alpha, biasTensorDesc, bias_d, + &beta, dstTensorDesc, dstData) ); + } else { + alpha = dnnType(1); + beta = dnnType(0); + cudnnBatchNormalizationForwardInference(net->cudnnHandle, + CUDNN_BATCHNORM_SPATIAL, &alpha, &beta, + dstTensorDesc, dstData, dstTensorDesc, + dstData, biasTensorDesc, //same tensor descriptor as bias + scales_d, bias_d, mean_d, variance_d, + CUDNN_BN_MIN_EPSILON); + } +} + Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, int strideH, int strideW, int paddingH, int paddingW, - std::string fname_weights, bool batchnorm) : + std::string fname_weights, bool batchnorm, bool deConv) : LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, fname_weights, batchnorm) { @@ -17,64 +129,28 @@ Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, this->strideW = strideW; this->paddingH = paddingH; this->paddingW = paddingW; + this->deConv = deConv; - checkCUDNN( cudnnCreateFilterDescriptor(&filterDesc) ); - checkCUDNN( cudnnCreateConvolutionDescriptor(&convDesc) ); - checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) ); - - int n = input_dim.n; - int c = input_dim.c; - int h = input_dim.h; - int w = input_dim.w; - - checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, - net->tensorFormat, net->dataType, n, c, h, w) ); - - checkCUDNN( cudnnSetFilter4dDescriptor(filterDesc, - net->dataType, net->tensorFormat, out_ch, input_dim.c, - kernelH, kernelW) ); - - checkCUDNN( cudnnSetConvolution2dDescriptor(convDesc, - paddingH, paddingW, // padding - strideH, strideW, // stride - 1,1, // upscale - CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT) ); - - // find dimension of convolution output - checkCUDNN( cudnnGetConvolution2dForwardOutputDim( - convDesc, srcTensorDesc, filterDesc, - &n, &c, &h, &w) ); - - checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, - net->tensorFormat, net->dataType, n, c, h, w) ); - - checkCUDNN( cudnnGetConvolutionForwardAlgorithm(net->cudnnHandle, - srcTensorDesc, filterDesc, convDesc, dstTensorDesc, - CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo) ); - - workSpace = NULL; - ws_sizeInBytes = 0; - - checkCUDNN( cudnnGetConvolutionForwardWorkspaceSize(net->cudnnHandle, - srcTensorDesc, filterDesc, convDesc, dstTensorDesc, - algo, &ws_sizeInBytes) ); + 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 * strideH) - 2*paddingH + kernelH -1; + output_dim.w = (input_dim.w * strideW) - 2*paddingW + kernelW -1; + output_dim.l = 1; + } + initCUDNN(deConv); + // allocate warkspace if (ws_sizeInBytes!=0) { checkCuda( cudaMalloc(&workSpace, ws_sizeInBytes) ); } - - checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc, - net->tensorFormat, net->dataType, - 1, out_ch, 1, 1) ); - - - output_dim.n = n; - output_dim.c = c; - output_dim.h = h; - output_dim.w = w; - output_dim.l = 1; - //allocate data for infer result checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); } @@ -93,35 +169,25 @@ Conv2d::~Conv2d() { dnnType* Conv2d::infer(dataDim_t &dim, dnnType* srcData) { + if(deConv) { + FatalError("you must use DeConv class for Deconvolutional layers"); + } // convolution - dnnType alpha = dnnType(1); - dnnType beta = dnnType(0); - checkCUDNN( cudnnConvolutionForward(net->cudnnHandle, - &alpha, srcTensorDesc, srcData, filterDesc, - data_d, convDesc, algo, workSpace, ws_sizeInBytes, - &beta, dstTensorDesc, dstData) ); + inferCUDNN(srcData, false); - if(!batchnorm) { - // bias - alpha = dnnType(1); - beta = dnnType(1); - checkCUDNN( cudnnAddTensor(net->cudnnHandle, - &alpha, biasTensorDesc, bias_d, - &beta, dstTensorDesc, dstData) ); - } else { - float one = 1; - float zero = 0; - cudnnBatchNormalizationForwardInference(net->cudnnHandle, - CUDNN_BATCHNORM_SPATIAL, &one, &zero, - dstTensorDesc, dstData, dstTensorDesc, - dstData, biasTensorDesc, //same tensor descriptor as bias - scales_d, bias_d, mean_d, variance_d, - CUDNN_BN_MIN_EPSILON); - } //update data dimensions dim = output_dim; + return dstData; +} +dnnType* DeConv2d::infer(dataDim_t &dim, dnnType* srcData) { + + // convolution + inferCUDNN(srcData, true); + + //update data dimensions + dim = output_dim; return dstData; } diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index 8219e26..8c739b0 100644 --- a/src/LayerWgs.cpp +++ b/src/LayerWgs.cpp @@ -16,18 +16,18 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs, std::cout<<"Reading weights: I="<dontLoadWeights); seek += inputs*outputs*kh*kw*kl; - readBinaryFile(weights_path.c_str(), outputs, &bias_h, &bias_d, seek); + readBinaryFile(weights_path.c_str(), outputs, &bias_h, &bias_d, seek, net->dontLoadWeights); this->batchnorm = batchnorm; if(batchnorm) { seek += outputs; - readBinaryFile(weights_path.c_str(), outputs, &scales_h, &scales_d, seek); + readBinaryFile(weights_path.c_str(), outputs, &scales_h, &scales_d, seek, net->dontLoadWeights); seek += outputs; - readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek); + readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek, net->dontLoadWeights); seek += outputs; - readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek); + readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek, net->dontLoadWeights); float eps = CUDNN_BN_MIN_EPSILON; diff --git a/src/Network.cpp b/src/Network.cpp index 491e32b..f55ecca 100644 --- a/src/Network.cpp +++ b/src/Network.cpp @@ -17,6 +17,7 @@ Network::Network(dataDim_t input_dim) { <<", CUDNN v"< Date: Wed, 30 Oct 2019 09:40:29 +0100 Subject: [PATCH 2/3] Deconv tensorrt --- include/tkDNN/Layer.h | 3 +-- include/tkDNN/NetworkRT.h | 1 + src/NetworkRT.cpp | 25 +++++++++++++++++-------- 3 files changed, 19 insertions(+), 10 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 483b7b5..6015235 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -161,6 +161,7 @@ public: virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); int kernelH, kernelW, strideH, strideW, paddingH, paddingW; + bool deConv; protected: cudnnFilterDescriptor_t filterDesc; @@ -173,8 +174,6 @@ protected: void inferCUDNN(dnnType* srcData, bool back = false); void* workSpace; size_t ws_sizeInBytes; - - bool deConv; }; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index d30da03..97dd478 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -75,6 +75,7 @@ public: nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeConv2d *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Activation *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Dense *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Pooling *l); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 2b3319c..ce29f01 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -159,7 +159,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { if(type == LAYER_DENSE) return convert_layer(input, (Dense*) l); - if(type == LAYER_CONV2D) + if(type == LAYER_CONV2D || type == LAYER_DECONV2D) return convert_layer(input, (Conv2d*) l); if(type == LAYER_POOLING) return convert_layer(input, (Pooling*) l); @@ -232,13 +232,22 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { else b = { dtRT, nullptr, 0}; //on batchnorm bias are added later - // Add a convolution layer with 20 outputs and a 5x5 filter. - IConvolutionLayer *lRT = networkRT->addConvolution(*input, - l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b); - checkNULL(lRT); - - lRT->setStride(DimsHW{l->strideH, l->strideW}); - lRT->setPadding(DimsHW{l->paddingH, l->paddingW}); + ILayer *lRT = nullptr; + if(!l->deConv) { + IConvolutionLayer *lRTconv = networkRT->addConvolution(*input, + l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b); + checkNULL(lRT); + lRTconv->setStride(DimsHW{l->strideH, l->strideW}); + lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); + lRT = (ILayer*) lRTconv; + } else { + IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input, + l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b); + checkNULL(lRT); + lRTconv->setStride(DimsHW{l->strideH, l->strideW}); + lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); + lRT = (ILayer*) lRTconv; + } if(l->batchnorm) { Weights power{dtRT, power_b, l->outputs}; From d6d93a74f8adfe482dbe49fad0b2772e8754639b Mon Sep 17 00:00:00 2001 From: fbagni Date: Wed, 30 Oct 2019 09:42:54 +0100 Subject: [PATCH 3/3] fix --- src/NetworkRT.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index ce29f01..0f38f87 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -236,14 +236,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { if(!l->deConv) { IConvolutionLayer *lRTconv = networkRT->addConvolution(*input, l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b); - checkNULL(lRT); + checkNULL(lRTconv); lRTconv->setStride(DimsHW{l->strideH, l->strideW}); lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); lRT = (ILayer*) lRTconv; } else { IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input, l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b); - checkNULL(lRT); + checkNULL(lRTconv); lRTconv->setStride(DimsHW{l->strideH, l->strideW}); lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); lRT = (ILayer*) lRTconv;