Merge branch 'conv-pad' into 'master'

Convolution and Pooling support arbitrary padding value.

See merge request !255

mace/kernels/conv_2d.h

@@ -178,12 +178,14 @@ void Conv2dKernelFunc(const T *input_ptr,  // batch start
 
 struct Conv2dFunctorBase {
   Conv2dFunctorBase(const int *strides,
-                    const Padding &paddings,
+                    const Padding &padding_type,
+                    const std::vector<int> &paddings,
                     const int *dilations,
                     const ActivationType activation,
                     const float relux_max_limit,
                     const float prelu_alpha)
       : strides_(strides),
+        padding_type_(padding_type),
         paddings_(paddings),
         dilations_(dilations),
         activation_(activation),
@@ -191,7 +193,8 @@ struct Conv2dFunctorBase {
         prelu_alpha_(prelu_alpha) {}
 
   const int *strides_;  // [stride_h, stride_w]
-  const Padding paddings_;
+  const Padding padding_type_;
+  std::vector<int> paddings_;
   const int *dilations_;  // [dilation_h, dilation_w]
   const ActivationType activation_;
   const float relux_max_limit_;
@@ -201,12 +204,14 @@ struct Conv2dFunctorBase {
 template <DeviceType D, typename T>
 struct Conv2dFunctor : Conv2dFunctorBase {
   Conv2dFunctor(const int *strides,
-                const Padding &paddings,
+                const Padding &padding_type,
+                const std::vector<int> &paddings,
                 const int *dilations,
                 const ActivationType activation,
                 const float relux_max_limit,
                 const float prelu_alpha)
       : Conv2dFunctorBase(strides,
+                          padding_type,
                           paddings,
                           dilations,
                           activation,
@@ -223,10 +228,12 @@ struct Conv2dFunctor : Conv2dFunctorBase {
     MACE_CHECK_NOTNULL(output);
 
     std::vector<index_t> output_shape(4);
-    std::vector<int> paddings(2);
+    if (paddings_.empty()) {
+      paddings_.resize(2);
       kernels::CalcNHWCPaddingAndOutputSize(
           input->shape().data(), filter->shape().data(), dilations_, strides_,
-        paddings_, output_shape.data(), paddings.data());
+          padding_type_, output_shape.data(), paddings_.data());
+    }
     output->Resize(output_shape);
 
     int batch = output->dim(0);
@@ -253,13 +260,13 @@ struct Conv2dFunctor : Conv2dFunctorBase {
 
     MACE_CHECK(batch == input_batch, "Input/Output batch size mismatch");
 
-    int padded_height = input_height + paddings[0];
-    int padded_width = input_width + paddings[1];
+    int padded_height = input_height + paddings_[0];
+    int padded_width = input_width + paddings_[1];
 
     Tensor padded_input;
     // Keep this alive during kernel execution
-    if (paddings[0] > 0 || paddings[1] > 0) {
-      ConstructNHWCInputWithPadding(input, paddings.data(), &padded_input);
+    if (paddings_[0] > 0 || paddings_[1] > 0) {
+      ConstructNHWCInputWithPadding(input, paddings_.data(), &padded_input);
       input = &padded_input;
     }
 
@@ -625,12 +632,14 @@ void Conv2dFunctor<DeviceType::NEON, float>::operator()(const Tensor *input,
 template <typename T>
 struct Conv2dFunctor<DeviceType::OPENCL, T> : Conv2dFunctorBase {
   Conv2dFunctor(const int *strides,
-                const Padding &paddings,
+                const Padding &padding_type,
+                const std::vector<int> &paddings,
                 const int *dilations,
                 const ActivationType activation,
                 const float relux_max_limit,
                 const float prelu_alpha)
       : Conv2dFunctorBase(strides,
+                          padding_type,
                           paddings,
                           dilations,
                           activation,

mace/kernels/depthwise_conv2d.h

@@ -237,20 +237,23 @@ void DepthwiseConv2dNoOOBCheckKernel(const T *input_ptr,
 
 struct DepthwiseConv2dFunctorBase {
   DepthwiseConv2dFunctorBase(const int *strides,
-                             const Padding padding,
+                             const Padding padding_type,
+                             const std::vector<int> &paddings,
                              const int *dilations,
                              const ActivationType activation,
                              const float relux_max_limit,
                              const float prelu_alpha)
       : strides_(strides),
-        padding_(padding),
+        padding_type_(padding_type),
+        paddings_(paddings),
         dilations_(dilations),
         activation_(activation),
         relux_max_limit_(relux_max_limit),
         prelu_alpha_(prelu_alpha) {}
 
   const int *strides_;  // [stride_h, stride_w]
-  const Padding padding_;
+  const Padding padding_type_;
+  std::vector<int> paddings_;
   const int *dilations_;  // [dilation_h, dilation_w]
   const ActivationType activation_;
   const float relux_max_limit_;
@@ -260,13 +263,15 @@ struct DepthwiseConv2dFunctorBase {
 template <DeviceType D, typename T>
 struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase {
   DepthwiseConv2dFunctor(const int *strides,
-                         const Padding padding,
+                         const Padding padding_type,
+                         const std::vector<int> &paddings,
                          const int *dilations,
                          const ActivationType activation,
                          const float relux_max_limit,
                          const float prelu_alpha)
       : DepthwiseConv2dFunctorBase(strides,
-                                   padding,
+                                   padding_type,
+                                   paddings,
                                    dilations,
                                    activation,
                                    relux_max_limit,
@@ -289,10 +294,12 @@ struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase {
     fake_filter_shape[3] = 1;
 
     std::vector<index_t> output_shape(4);
-    std::vector<int> paddings(2);
+    if (paddings_.empty()) {
+      paddings_.resize(2);
       kernels::CalcNHWCPaddingAndOutputSize(
           input->shape().data(), fake_filter_shape.data(), dilations_, strides_,
-        padding_, output_shape.data(), paddings.data());
+          padding_type_, output_shape.data(), paddings_.data());
+    }
     auto input_shape = fake_filter_shape;
     output->Resize(output_shape);
 
@@ -322,10 +329,10 @@ struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase {
     MACE_CHECK(batch == input_batch, "Input/Output batch size mismatch");
 
     // The left-upper most offset of the padded input
-    int paddings_top = paddings[0] / 2;
-    int paddings_bottom = paddings[0] - paddings_top;
-    int paddings_left = paddings[1] / 2;
-    int paddings_right = paddings[1] - paddings_left;
+    int paddings_top = paddings_[0] / 2;
+    int paddings_bottom = paddings_[0] - paddings_top;
+    int paddings_left = paddings_[1] / 2;
+    int paddings_right = paddings_[1] - paddings_left;
 
     int padded_h_start = 0 - paddings_top;
     int padded_w_start = 0 - paddings_left;
@@ -413,13 +420,15 @@ template <typename T>
 struct DepthwiseConv2dFunctor<DeviceType::OPENCL, T>
     : DepthwiseConv2dFunctorBase {
   DepthwiseConv2dFunctor(const int *strides,
-                         const Padding padding,
+                         const Padding padding_type,
+                         const std::vector<int> &paddings,
                          const int *dilations,
                          const ActivationType activation,
                          const float relux_max_limit,
                          const float prelu_alpha)
       : DepthwiseConv2dFunctorBase(strides,
-                                   padding,
+                                   padding_type,
+                                   paddings,
                                    dilations,
                                    activation,
                                    relux_max_limit,

mace/kernels/opencl/conv_2d_opencl.cc

@@ -80,10 +80,12 @@ void Conv2dFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
   }
 
   std::vector<index_t> output_shape(4);
-  std::vector<int> paddings(2);
+  if (paddings_.empty()) {
+    paddings_.resize(2);
     kernels::CalcNHWCPaddingAndOutputSize(
         input->shape().data(), filter->shape().data(), dilations_, strides_,
-      paddings_, output_shape.data(), paddings.data());
+        padding_type_, output_shape.data(), paddings_.data());
+  }
 
   std::vector<size_t> output_image_shape;
   CalImage2DShape(output_shape, BufferType::IN_OUT_CHANNEL, output_image_shape);
@@ -93,11 +95,11 @@ void Conv2dFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
       selector[kernel_h - 1] != nullptr &&
       0 < strides_[0] && strides_[0] < 3 ) {
     auto conv2d_func = selector[kernel_h - 1];
-    conv2d_func(&kernel_, input, filter, bias, strides_[0], paddings.data(), dilations_, activation_,
+    conv2d_func(&kernel_, input, filter, bias, strides_[0], paddings_.data(), dilations_, activation_,
                 relux_max_limit_, prelu_alpha_, DataTypeToEnum<T>::value,
                 output, future);
   } else {
-    Conv2dOpencl(&kernel_, input, filter, bias, strides_[0], paddings.data(), dilations_,
+    Conv2dOpencl(&kernel_, input, filter, bias, strides_[0], paddings_.data(), dilations_,
                  activation_, relux_max_limit_, prelu_alpha_,
                  DataTypeToEnum<T>::value, output, future);
   }

mace/kernels/opencl/depthwise_conv_opencl.cc

@@ -140,7 +140,7 @@ void DepthwiseConv2dFunctor<DeviceType::OPENCL, T>::operator()(
                  << " is not implemented yet, using slow version";
     // TODO(heliangliang) The CPU/NEON kernel should map the buffer
     DepthwiseConv2dFunctor<DeviceType::CPU, float>(
-        strides_, padding_, dilations_, activation_, relux_max_limit_,
+        strides_, padding_type_, paddings_, dilations_, activation_, relux_max_limit_,
         prelu_alpha_)(input, filter, bias, output, future);
     return;
   }
@@ -153,16 +153,18 @@ void DepthwiseConv2dFunctor<DeviceType::OPENCL, T>::operator()(
   fake_filter_shape[3] = 1;
 
   std::vector<index_t> output_shape(4);
-  std::vector<int> paddings(2);
+  if (paddings_.empty()) {
+    paddings_.resize(2);
     kernels::CalcNHWCPaddingAndOutputSize(
         input->shape().data(), fake_filter_shape.data(), dilations_, strides_,
-      padding_, output_shape.data(), paddings.data());
+        padding_type_, output_shape.data(), paddings_.data());
+  }
 
   std::vector<size_t> output_image_shape;
   CalImage2DShape(output_shape, BufferType::IN_OUT_CHANNEL, output_image_shape);
   output->ResizeImage(output_shape, output_image_shape);
 
-  DepthwiseConv2d(&kernel_, input, filter, bias, strides_[0], paddings.data(), dilations_,
+  DepthwiseConv2d(&kernel_, input, filter, bias, strides_[0], paddings_.data(), dilations_,
                   activation_, relux_max_limit_, prelu_alpha_,
                   DataTypeToEnum<T>::value, output, future);
 }

mace/kernels/opencl/pooling_opencl.cc

@@ -18,16 +18,18 @@ void PoolingFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
   MACE_CHECK(dilations_[0] == 1 && dilations_[1] == 1)
     << "Pooling opencl kernel not support dilation yet";
   std::vector<index_t> output_shape(4);
-  std::vector<int> paddings(2);
   std::vector<index_t> filter_shape = {
       kernels_[0], kernels_[1],
       input->dim(3), input->dim(3)
   };
 
+  if (paddings_.empty()) {
+    paddings_.resize(2);
     kernels::CalcNHWCPaddingAndOutputSize(
         input->shape().data(), filter_shape.data(),
-      dilations_, strides_, this->padding_,
-      output_shape.data(), paddings.data());
+        dilations_, strides_, this->padding_type_,
+        output_shape.data(), paddings_.data());
+  }
 
   std::vector<size_t> output_image_shape;
   CalImage2DShape(output_shape, BufferType::IN_OUT_CHANNEL, output_image_shape);
@@ -64,8 +66,8 @@ void PoolingFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
     kernel_.setArg(idx++, static_cast<int32_t>(input->dim(1)));
     kernel_.setArg(idx++, static_cast<int32_t>(input->dim(2)));
     kernel_.setArg(idx++, static_cast<int32_t>(out_height));
-    kernel_.setArg(idx++, paddings[0] / 2);
-    kernel_.setArg(idx++, paddings[1] / 2);
+    kernel_.setArg(idx++, paddings_[0] / 2);
+    kernel_.setArg(idx++, paddings_[1] / 2);
     kernel_.setArg(idx++, strides_[0]);
     kernel_.setArg(idx++, kernels_[0]);
     kernel_.setArg(idx++, *(static_cast<cl::Image2D *>(output->buffer())));

mace/kernels/opencl/winograd_transform.cc

@@ -17,10 +17,12 @@ void WinogradTransformFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *i
                                                                  StatsFuture *future) {
   std::vector<index_t> output_shape(4);
   std::vector<index_t> filter_shape = {3, 3, input_tensor->dim(3), 1};
-  std::vector<int> paddings(2);
+  if (paddings_.empty()) {
+    paddings_.resize(2);
     kernels::CalcNHWCPaddingAndOutputSize(
         input_tensor->shape().data(), filter_shape.data(), dilations_.data(),
-      strides_.data(), paddings_, output_shape.data(), paddings.data());
+        strides_.data(), padding_type_, output_shape.data(), paddings_.data());
+  }
 
   const index_t round_h = (output_shape[1] + 1) / 2;
   const index_t round_w = (output_shape[2] + 1) / 2;
@@ -50,8 +52,8 @@ void WinogradTransformFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *i
     kernel_.setArg(idx++, static_cast<uint32_t>(input_tensor->dim(3)));
     kernel_.setArg(idx++, static_cast<uint32_t>(round_h * round_w));
     kernel_.setArg(idx++, static_cast<uint32_t>(round_w));
-    kernel_.setArg(idx++, static_cast<uint32_t>(paddings[0] / 2));
-    kernel_.setArg(idx++, static_cast<uint32_t>(paddings[1] / 2));
+    kernel_.setArg(idx++, static_cast<uint32_t>(paddings_[0] / 2));
+    kernel_.setArg(idx++, static_cast<uint32_t>(paddings_[1] / 2));
   }
 
   const uint32_t gws[2] = {static_cast<uint32_t>(out_width),

mace/kernels/pooling.h

@@ -24,18 +24,21 @@ struct PoolingFunctorBase {
   PoolingFunctorBase(const PoolingType pooling_type,
                      const int *kernels,
                      const int *strides,
-                     const Padding padding,
+                     const Padding padding_type,
+                     const std::vector<int> &paddings,
                      const int *dilations)
       : pooling_type_(pooling_type),
         kernels_(kernels),
         strides_(strides),
-        padding_(padding),
+        padding_type_(padding_type),
+        paddings_(paddings),
         dilations_(dilations) {}
 
   const PoolingType pooling_type_;
   const int *kernels_;
   const int *strides_;
-  const Padding padding_;
+  const Padding padding_type_;
+  std::vector<int> paddings_;
   const int *dilations_;
 };
 
@@ -44,27 +47,31 @@ struct PoolingFunctor : PoolingFunctorBase {
   PoolingFunctor(const PoolingType pooling_type,
                  const int *kernels,
                  const int *strides,
-                 const Padding padding,
+                 const Padding padding_type,
+                 const std::vector<int> &paddings,
                  const int *dilations)
       : PoolingFunctorBase(pooling_type, kernels,
-                           strides, padding,
-                           dilations) {}
+                           strides, padding_type,
+                           paddings, dilations) {}
 
   void operator()(const Tensor *input_tensor,
                   Tensor *output_tensor,
                   StatsFuture *future) {
 
     std::vector<index_t> output_shape(4);
-    std::vector<int> paddings(2);
     std::vector<index_t> filter_shape = {
         kernels_[0], kernels_[1],
         input_tensor->dim(3), input_tensor->dim(3)
     };
 
+    if (paddings_.empty()) {
+      paddings_.resize(2);
       kernels::CalcNHWCPaddingAndOutputSize(
           input_tensor->shape().data(), filter_shape.data(),
-        dilations_, strides_, this->padding_,
-        output_shape.data(), paddings.data());
+          dilations_, strides_, this->padding_type_,
+          output_shape.data(), paddings_.data());
+
+    }
     output_tensor->Resize(output_shape);
 
     Tensor::MappingGuard in_guard(input_tensor);
@@ -92,8 +99,8 @@ struct PoolingFunctor : PoolingFunctorBase {
     int dilation_w = dilations_[1];
 
     // The left-upper most offset of the padded input
-    int padded_h_start = 0 - paddings[0] / 2;
-    int padded_w_start = 0 - paddings[1] / 2;
+    int padded_h_start = 0 - paddings_[0] / 2;
+    int padded_w_start = 0 - paddings_[1] / 2;
 
     if (pooling_type_ == MAX) {
 #pragma omp parallel for collapse(4)
@@ -163,11 +170,12 @@ struct PoolingFunctor<DeviceType::OPENCL, T> : PoolingFunctorBase {
   PoolingFunctor(const PoolingType pooling_type,
                  const int *kernels,
                  const int *strides,
-                 const Padding padding,
+                 const Padding padding_type,
+                 const std::vector<int> &paddings,
                  const int *dilations)
       : PoolingFunctorBase(pooling_type, kernels,
-                           strides, padding,
-                           dilations) {}
+                           strides, padding_type,
+                           paddings, dilations) {}
   void operator()(const Tensor *input_tensor,
                   Tensor *output_tensor,
                   StatsFuture *future);

mace/kernels/winograd_transform.h

@@ -15,18 +15,22 @@ namespace mace {
 namespace kernels {
 
 struct WinogradTransformFunctorBase {
-  WinogradTransformFunctorBase(const Padding &paddings)
-      : strides_({1, 1}), dilations_({1, 1}), paddings_(paddings) {}
+  WinogradTransformFunctorBase(const Padding &padding_type,
+                               const std::vector<int> &paddings)
+      : strides_({1, 1}), dilations_({1, 1}),
+        padding_type_(padding_type), paddings_(paddings) {}
 
   const std::vector<int> strides_;         // [stride_h, stride_w]
   const std::vector<int> dilations_;       // [dilation_h, dilation_w]
-  Padding paddings_;
+  Padding padding_type_;
+  std::vector<int> paddings_;
 };
 
 template<DeviceType D, typename T>
 struct WinogradTransformFunctor : WinogradTransformFunctorBase {
-  WinogradTransformFunctor(const Padding &paddings)
-      : WinogradTransformFunctorBase(paddings) {}
+  WinogradTransformFunctor(const Padding &padding_type,
+                           const std::vector<int> &paddings)
+      : WinogradTransformFunctorBase(padding_type, paddings) {}
 
   void operator()(const Tensor *input,
                   Tensor *output,
@@ -38,8 +42,9 @@ struct WinogradTransformFunctor : WinogradTransformFunctorBase {
 
 template<typename T>
 struct WinogradTransformFunctor<DeviceType::OPENCL, T> : WinogradTransformFunctorBase {
-  WinogradTransformFunctor(const Padding &paddings)
-      : WinogradTransformFunctorBase(paddings) {}
+  WinogradTransformFunctor(const Padding &padding_type,
+                           const std::vector<int> &paddings)
+      : WinogradTransformFunctorBase(padding_type, paddings) {}
 
   void operator()(const Tensor *input,
                   Tensor *output,

mace/ops/addn.h

@@ -26,7 +26,12 @@ class AddNOp : public Operator<D, T> {
     for (int i = 1; i < n; ++i) {
       inputs[i] = this->Input(i);
       MACE_CHECK(inputs[0]->dim_size() == inputs[i]->dim_size());
-      MACE_CHECK(inputs[0]->size() == inputs[i]->size());
+      MACE_CHECK(inputs[0]->size() == inputs[i]->size()) << "Input 0: "
+                                                         << MakeString(inputs[0]->shape())
+                                                         << ", size: " << inputs[0]->size()
+                                                         << ". Input " << i << ": "
+                                                         << MakeString(inputs[i]->shape())
+                                                         << ", size: " << inputs[i]->size();
     }
 
     functor_(inputs, output_tensor, future);

mace/ops/conv_2d.h

@@ -19,7 +19,8 @@ class Conv2dOp : public ConvPool2dOpBase<D, T> {
   Conv2dOp(const OperatorDef &op_def, Workspace *ws)
       : ConvPool2dOpBase<D, T>(op_def, ws),
         functor_(this->strides_.data(),
-                 this->padding_,
+                 this->padding_type_,
+                 this->paddings_,
                  this->dilations_.data(),
                  kernels::ActivationType::NOOP,
                  0.0f,

mace/ops/conv_pool_2d_base.h

@@ -16,14 +16,16 @@ class ConvPool2dOpBase : public Operator<D, T> {
   ConvPool2dOpBase(const OperatorDef &op_def, Workspace *ws)
       : Operator<D, T>(op_def, ws),
         strides_(OperatorBase::GetRepeatedArgument<int>("strides")),
-        padding_(static_cast<Padding>(OperatorBase::GetSingleArgument<int>(
+        padding_type_(static_cast<Padding>(OperatorBase::GetSingleArgument<int>(
             "padding", static_cast<int>(SAME)))),
+        paddings_(OperatorBase::GetRepeatedArgument<int>("padding_values")),
         dilations_(
             OperatorBase::GetRepeatedArgument<int>("dilations", {1, 1})) {}
 
  protected:
   std::vector<int> strides_;
-  Padding padding_;
+  Padding padding_type_;
+  std::vector<int> paddings_;
   std::vector<int> dilations_;
 };
 

mace/ops/depthwise_conv2d.h

@@ -20,7 +20,8 @@ class DepthwiseConv2dOp : public ConvPool2dOpBase<D, T> {
   DepthwiseConv2dOp(const OperatorDef &op_def, Workspace *ws)
       : ConvPool2dOpBase<D, T>(op_def, ws),
         functor_(this->strides_.data(),
-                 this->padding_,
+                 this->padding_type_,
+                 this->paddings_,
                  this->dilations_.data(),
                  kernels::StringToActivationType(
                      OperatorBase::GetSingleArgument<std::string>("activation",

mace/ops/fused_conv_2d.h

@@ -19,7 +19,8 @@ class FusedConv2dOp : public ConvPool2dOpBase<D, T> {
   FusedConv2dOp(const OperatorDef &op_def, Workspace *ws)
       : ConvPool2dOpBase<D, T>(op_def, ws),
         functor_(this->strides_.data(),
-                 this->padding_,
+                 this->padding_type_,
+                 this->paddings_,
                  this->dilations_.data(),
                  kernels::StringToActivationType(
                      OperatorBase::GetSingleArgument<std::string>("activation",

mace/ops/pooling.h

@@ -21,7 +21,7 @@ class PoolingOp : public ConvPool2dOpBase<D, T> {
             static_cast<PoolingType>(OperatorBase::GetSingleArgument<int>(
                 "pooling_type", static_cast<int>(AVG)))),
         functor_(pooling_type_, kernels_.data(), this->strides_.data(),
-                 this->padding_, this->dilations_.data()){};
+                 this->padding_type_, this->paddings_, this->dilations_.data()){};
 
   bool Run(StatsFuture *future) override {
     const Tensor *input = this->Input(INPUT);

mace/ops/winograd_transform.h

@@ -18,7 +18,8 @@ class WinogradTransformOp : public Operator<D, T> {
   WinogradTransformOp(const OperatorDef &op_def, Workspace *ws)
       : Operator<D, T>(op_def, ws),
         functor_(static_cast<Padding>(OperatorBase::GetSingleArgument<int>(
-            "padding", static_cast<int>(VALID)))) {}
+            "padding", static_cast<int>(VALID))),
+        OperatorBase::GetRepeatedArgument<int>("padding_values")) {}
 
   bool Run(StatsFuture *future) override {
     const Tensor *input_tensor = this->Input(INPUT);