From ba022663f15a75be0f820a9b7dcd9a3bdf59142e Mon Sep 17 00:00:00 2001 From: perseusdg <43143075+perseusdg@users.noreply.github.com> Date: Mon, 3 Jan 2022 19:49:32 +0530 Subject: [PATCH] completed padding migrations from github --- include/tkDNN/Layer.h | 27 ++++++++++++++++++++++++++- include/tkDNN/NetworkRT.h | 1 + include/tkDNN/kernels.h | 3 +-- src/NetworkRT.cpp | 11 +++++++++++ src/Padding.cpp | 35 +++++++++++++++++++++++++++++++++++ src/kernels/padding.cu | 10 +++++----- 6 files changed, 79 insertions(+), 8 deletions(-) create mode 100644 src/Padding.cpp diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 5273c83..2917733 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -31,7 +31,8 @@ enum layerType_t { LAYER_SHORTCUT, LAYER_UPSAMPLE, LAYER_REGION, - LAYER_YOLO + LAYER_YOLO, + LAYER_PADDING }; #define TKDNN_BN_MIN_EPSILON 1e-5 @@ -87,6 +88,7 @@ public: case LAYER_UPSAMPLE: return "Upsample"; case LAYER_REGION: return "Region"; case LAYER_YOLO: return "Yolo"; + case LAYER_PADDING: return "Padding"; default: return "unknown"; } } @@ -520,9 +522,32 @@ protected: bool poolOn3d; }; +/** + * Padding Layers + * tkDNN supports reflection,constant and zero padding + */ + +typedef enum { + PADDING_MODE_CONSTANT = 0, + PADDING_MODE_ZERO = 1, + PADDING_MODE_REFLECTION = 2 +} tkdnnPaddingMode_t; + +class Padding : public Layer { +public: + Padding(Network *net,int32_t pad_h,int32_t pad_w,tkdnnPaddingMode_t padding_mode); + virtual ~Padding(); + virtual layerType_t getLayerType(){return LAYER_PADDING ;}; + virtual dnnType* infer(dataDim_t& dim,dnnType* srcData); + int32_t paddingH,paddingW; + tkdnnPaddingMode_t padding_mode; + +}; + /** Softmax layer */ + class Softmax : public Layer { public: diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 58a545a..571d127 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -95,6 +95,7 @@ public: nvinfer1::IPluginV2Layer* convert_layer(nvinfer1::ITensor *input, Yolo *l); nvinfer1::IResizeLayer* convert_layer(nvinfer1::ITensor *input, Upsample *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeformConv2d *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input,Padding *l); #if NV_TENSORRT_MAJOR > 5 && NV_TENSORRT_MAJOR < 8 bool serialize(const char *filename); diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index a954fe7..7c9170f 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -49,7 +49,6 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, void scalAdd(dnnType* dstData, int size, float alpha, float beta, int inc, cudaStream_t stream = cudaStream_t(0)); -void reflection_pad2d_out_forward(int32_t padding[4], float* srcData, float* dstData, int32_t input_h, int32_t input_w, int32_t plane_dim, int32_t n_batch, cudaStream_t cudaStream = cudaStream_t(0)); - +void reflection_pad2d_out_forward(int32_t pad_h,int32_t pad_w,float *srcData,float *dstData,int32_t input_h,int32_t input_w,int32_t plane_dim,int32_t n_batch,cudaStream_t cudaStream = cudaStream_t(0)); #endif //KERNELS_H diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index ac1f23c..b3aa263 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -275,6 +275,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { return convert_layer(input, (Upsample*) l); if(type == LAYER_DEFORMCONV2D) return convert_layer(input, (DeformConv2d*) l); + if(type == LAYER_PADDING) + return convert_layer(input, (Padding*) l); std::cout<getLayerName()<<"\n"; FatalError("Layer not implemented in tensorRT"); @@ -453,6 +455,15 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) { } } +ILayer* NetworkRT::convert_layer(ITensor *input,Padding *l){ + auto *lRT = networkRT->addSlice(*input,Dims3{0,0,0},Dims3{l->output_dim.c,l->output_dim.h,l->output_dim.w},Dims3{0,0,0}); + if(l->padding_mode == PADDING_MODE_REFLECTION){ + lRT->setMode(SliceMode::kREFLECT); + } + checkNULL(lRT); + return lRT; +} + ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { //std::cout<<"convert Activation\n"; diff --git a/src/Padding.cpp b/src/Padding.cpp new file mode 100644 index 0000000..91857ab --- /dev/null +++ b/src/Padding.cpp @@ -0,0 +1,35 @@ +// +// Created by perseusdg on 03/01/22. +// + +#include +#include "Layer.h" +#include "kernels.h" + +namespace tk{ namespace dnn { + Padding::Padding(Network *net, int32_t pad_h, int32_t pad_w, tkdnnPaddingMode_t padding_mode) : Layer(net) { + this->paddingH = pad_h; + this->paddingW = pad_w; + this->padding_mode = padding_mode; + output_dim.c = input_dim.c; + output_dim.n = input_dim.n; + output_dim.h = input_dim.h + 2 * (this->paddingH); + output_dim.w = input_dim.w + 2 * (this->paddingW); + checkCuda(cudaMalloc(&dstData,output_dim.tot()*sizeof(dnnType))); + } + + Padding::~Padding() { + checkCuda(cudaFree(dstData)); + } + dnnType* Padding::infer(dataDim_t &dim, float *srcData) { + fill(dstData,output_dim.tot(),0.0); + if(padding_mode == tkdnnPaddingMode_t::PADDING_MODE_REFLECTION) + { + reflection_pad2d_out_forward(paddingH, paddingW, srcData, dstData, input_dim.h, input_dim.w, input_dim.c, + input_dim.n); + } + dim = output_dim; + return dstData; + } + +}} diff --git a/src/kernels/padding.cu b/src/kernels/padding.cu index 731f684..1b84ce2 100644 --- a/src/kernels/padding.cu +++ b/src/kernels/padding.cu @@ -44,11 +44,11 @@ int32_t ceilDiv(int32_t a,int32_t b){ } -void reflection_pad2d_out_forward(int32_t padding[4],float *srcData,float *dstData,int32_t input_h,int32_t input_w,int32_t plane_dim,int32_t n_batch,cudaStream_t cudaStream){ - int32_t pad_l = padding[0]; - int32_t pad_r = padding[1]; - int32_t pad_t = padding[2]; - int32_t pad_b = padding[3]; +void reflection_pad2d_out_forward(int32_t pad_h,int32_t pad_w,float *srcData,float *dstData,int32_t input_h,int32_t input_w,int32_t plane_dim,int32_t n_batch,cudaStream_t cudaStream){ + int32_t pad_l = pad_w; + int32_t pad_r = pad_w; + int32_t pad_t = pad_h; + int32_t pad_b = pad_w; int32_t output_h = input_h + pad_t + pad_b; int32_t output_w = input_w + pad_l + pad_r; int32_t size_y = plane_dim;