completed padding migrations from github

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:

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);

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

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<<l->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";
 

src/Padding.cpp

@@ -0,0 +1,35 @@
+//
+// Created by perseusdg on 03/01/22.
+//
+
+#include <iostream>
+#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;
+    }
+
+}}

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){
+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 = padding[0];
+    int32_t pad_l = pad_w;
-    int32_t pad_r = padding[1];
+    int32_t pad_r = pad_w;
-    int32_t pad_t = padding[2];
+    int32_t pad_t = pad_h;
-    int32_t pad_b = padding[3];
+    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;