From c5bc6a5a9676ba81598f993de9da926bdea78684 Mon Sep 17 00:00:00 2001 From: liuqi Date: Thu, 1 Mar 2018 18:12:21 +0800 Subject: [PATCH] Convolution support arbitrary padding value. --- mace/kernels/conv_2d.h | 33 ++++++++++------ mace/kernels/depthwise_conv2d.h | 39 +++++++++++-------- mace/kernels/opencl/conv_2d_opencl.cc | 14 ++++--- mace/kernels/opencl/depthwise_conv_opencl.cc | 14 ++++--- mace/kernels/opencl/pooling_opencl.cc | 16 ++++---- mace/kernels/opencl/winograd_transform.cc | 14 ++++--- mace/kernels/pooling.h | 40 ++++++++++++-------- mace/kernels/winograd_transform.h | 19 ++++++---- mace/ops/addn.h | 7 +++- mace/ops/conv_2d.h | 3 +- mace/ops/conv_pool_2d_base.h | 6 ++- mace/ops/depthwise_conv2d.h | 3 +- mace/ops/fused_conv_2d.h | 3 +- mace/ops/pooling.h | 2 +- mace/ops/winograd_transform.h | 3 +- 15 files changed, 133 insertions(+), 83 deletions(-) diff --git a/mace/kernels/conv_2d.h b/mace/kernels/conv_2d.h index ae605016..a457edfb 100644 --- a/mace/kernels/conv_2d.h +++ b/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 &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 paddings_; const int *dilations_; // [dilation_h, dilation_w] const ActivationType activation_; const float relux_max_limit_; @@ -201,12 +204,14 @@ struct Conv2dFunctorBase { template struct Conv2dFunctor : Conv2dFunctorBase { Conv2dFunctor(const int *strides, - const Padding &paddings, + const Padding &padding_type, + const std::vector &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 output_shape(4); - std::vector paddings(2); - kernels::CalcNHWCPaddingAndOutputSize( - input->shape().data(), filter->shape().data(), dilations_, strides_, - paddings_, output_shape.data(), paddings.data()); + if (paddings_.empty()) { + paddings_.resize(2); + kernels::CalcNHWCPaddingAndOutputSize( + input->shape().data(), filter->shape().data(), dilations_, strides_, + 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::operator()(const Tensor *input, template struct Conv2dFunctor : Conv2dFunctorBase { Conv2dFunctor(const int *strides, - const Padding &paddings, + const Padding &padding_type, + const std::vector &paddings, const int *dilations, const ActivationType activation, const float relux_max_limit, const float prelu_alpha) : Conv2dFunctorBase(strides, + padding_type, paddings, dilations, activation, diff --git a/mace/kernels/depthwise_conv2d.h b/mace/kernels/depthwise_conv2d.h index ea848f35..65e8b500 100644 --- a/mace/kernels/depthwise_conv2d.h +++ b/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 &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 paddings_; const int *dilations_; // [dilation_h, dilation_w] const ActivationType activation_; const float relux_max_limit_; @@ -260,13 +263,15 @@ struct DepthwiseConv2dFunctorBase { template struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase { DepthwiseConv2dFunctor(const int *strides, - const Padding padding, + const Padding padding_type, + const std::vector &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 output_shape(4); - std::vector paddings(2); - kernels::CalcNHWCPaddingAndOutputSize( - input->shape().data(), fake_filter_shape.data(), dilations_, strides_, - padding_, output_shape.data(), paddings.data()); + if (paddings_.empty()) { + paddings_.resize(2); + kernels::CalcNHWCPaddingAndOutputSize( + input->shape().data(), fake_filter_shape.data(), dilations_, strides_, + 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 struct DepthwiseConv2dFunctor : DepthwiseConv2dFunctorBase { DepthwiseConv2dFunctor(const int *strides, - const Padding padding, + const Padding padding_type, + const std::vector &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, diff --git a/mace/kernels/opencl/conv_2d_opencl.cc b/mace/kernels/opencl/conv_2d_opencl.cc index 5a6d93cb..48e775e1 100644 --- a/mace/kernels/opencl/conv_2d_opencl.cc +++ b/mace/kernels/opencl/conv_2d_opencl.cc @@ -80,10 +80,12 @@ void Conv2dFunctor::operator()(const Tensor *input, } std::vector output_shape(4); - std::vector paddings(2); - kernels::CalcNHWCPaddingAndOutputSize( - input->shape().data(), filter->shape().data(), dilations_, strides_, - paddings_, output_shape.data(), paddings.data()); + if (paddings_.empty()) { + paddings_.resize(2); + kernels::CalcNHWCPaddingAndOutputSize( + input->shape().data(), filter->shape().data(), dilations_, strides_, + padding_type_, output_shape.data(), paddings_.data()); + } std::vector output_image_shape; CalImage2DShape(output_shape, BufferType::IN_OUT_CHANNEL, output_image_shape); @@ -93,11 +95,11 @@ void Conv2dFunctor::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::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::value, output, future); } diff --git a/mace/kernels/opencl/depthwise_conv_opencl.cc b/mace/kernels/opencl/depthwise_conv_opencl.cc index 4365c02f..440b18b4 100644 --- a/mace/kernels/opencl/depthwise_conv_opencl.cc +++ b/mace/kernels/opencl/depthwise_conv_opencl.cc @@ -140,7 +140,7 @@ void DepthwiseConv2dFunctor::operator()( << " is not implemented yet, using slow version"; // TODO(heliangliang) The CPU/NEON kernel should map the buffer DepthwiseConv2dFunctor( - 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::operator()( fake_filter_shape[3] = 1; std::vector output_shape(4); - std::vector paddings(2); - kernels::CalcNHWCPaddingAndOutputSize( - input->shape().data(), fake_filter_shape.data(), dilations_, strides_, - padding_, output_shape.data(), paddings.data()); + if (paddings_.empty()) { + paddings_.resize(2); + kernels::CalcNHWCPaddingAndOutputSize( + input->shape().data(), fake_filter_shape.data(), dilations_, strides_, + padding_type_, output_shape.data(), paddings_.data()); + } std::vector 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::value, output, future); } diff --git a/mace/kernels/opencl/pooling_opencl.cc b/mace/kernels/opencl/pooling_opencl.cc index 8a4274e8..39bb994e 100644 --- a/mace/kernels/opencl/pooling_opencl.cc +++ b/mace/kernels/opencl/pooling_opencl.cc @@ -18,16 +18,18 @@ void PoolingFunctor::operator()(const Tensor *input, MACE_CHECK(dilations_[0] == 1 && dilations_[1] == 1) << "Pooling opencl kernel not support dilation yet"; std::vector output_shape(4); - std::vector paddings(2); std::vector filter_shape = { kernels_[0], kernels_[1], input->dim(3), input->dim(3) }; - kernels::CalcNHWCPaddingAndOutputSize( - input->shape().data(), filter_shape.data(), - dilations_, strides_, this->padding_, - output_shape.data(), paddings.data()); + if (paddings_.empty()) { + paddings_.resize(2); + kernels::CalcNHWCPaddingAndOutputSize( + input->shape().data(), filter_shape.data(), + dilations_, strides_, this->padding_type_, + output_shape.data(), paddings_.data()); + } std::vector output_image_shape; CalImage2DShape(output_shape, BufferType::IN_OUT_CHANNEL, output_image_shape); @@ -64,8 +66,8 @@ void PoolingFunctor::operator()(const Tensor *input, kernel_.setArg(idx++, static_cast(input->dim(1))); kernel_.setArg(idx++, static_cast(input->dim(2))); kernel_.setArg(idx++, static_cast(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(output->buffer()))); diff --git a/mace/kernels/opencl/winograd_transform.cc b/mace/kernels/opencl/winograd_transform.cc index bacbdb63..0396af97 100644 --- a/mace/kernels/opencl/winograd_transform.cc +++ b/mace/kernels/opencl/winograd_transform.cc @@ -17,10 +17,12 @@ void WinogradTransformFunctor::operator()(const Tensor *i StatsFuture *future) { std::vector output_shape(4); std::vector filter_shape = {3, 3, input_tensor->dim(3), 1}; - std::vector paddings(2); - kernels::CalcNHWCPaddingAndOutputSize( - input_tensor->shape().data(), filter_shape.data(), dilations_.data(), - strides_.data(), paddings_, output_shape.data(), paddings.data()); + if (paddings_.empty()) { + paddings_.resize(2); + kernels::CalcNHWCPaddingAndOutputSize( + input_tensor->shape().data(), filter_shape.data(), dilations_.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::operator()(const Tensor *i kernel_.setArg(idx++, static_cast(input_tensor->dim(3))); kernel_.setArg(idx++, static_cast(round_h * round_w)); kernel_.setArg(idx++, static_cast(round_w)); - kernel_.setArg(idx++, static_cast(paddings[0] / 2)); - kernel_.setArg(idx++, static_cast(paddings[1] / 2)); + kernel_.setArg(idx++, static_cast(paddings_[0] / 2)); + kernel_.setArg(idx++, static_cast(paddings_[1] / 2)); } const uint32_t gws[2] = {static_cast(out_width), diff --git a/mace/kernels/pooling.h b/mace/kernels/pooling.h index d9a28c54..4c62ba4e 100644 --- a/mace/kernels/pooling.h +++ b/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 &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 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 &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 output_shape(4); - std::vector paddings(2); std::vector filter_shape = { kernels_[0], kernels_[1], input_tensor->dim(3), input_tensor->dim(3) }; - kernels::CalcNHWCPaddingAndOutputSize( - input_tensor->shape().data(), filter_shape.data(), - dilations_, strides_, this->padding_, - output_shape.data(), paddings.data()); + if (paddings_.empty()) { + paddings_.resize(2); + kernels::CalcNHWCPaddingAndOutputSize( + input_tensor->shape().data(), filter_shape.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 : PoolingFunctorBase { PoolingFunctor(const PoolingType pooling_type, const int *kernels, const int *strides, - const Padding padding, + const Padding padding_type, + const std::vector &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); diff --git a/mace/kernels/winograd_transform.h b/mace/kernels/winograd_transform.h index d6ba1e62..639138ee 100644 --- a/mace/kernels/winograd_transform.h +++ b/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 &paddings) + : strides_({1, 1}), dilations_({1, 1}), + padding_type_(padding_type), paddings_(paddings) {} const std::vector strides_; // [stride_h, stride_w] const std::vector dilations_; // [dilation_h, dilation_w] - Padding paddings_; + Padding padding_type_; + std::vector paddings_; }; template struct WinogradTransformFunctor : WinogradTransformFunctorBase { - WinogradTransformFunctor(const Padding &paddings) - : WinogradTransformFunctorBase(paddings) {} + WinogradTransformFunctor(const Padding &padding_type, + const std::vector &paddings) + : WinogradTransformFunctorBase(padding_type, paddings) {} void operator()(const Tensor *input, Tensor *output, @@ -38,8 +42,9 @@ struct WinogradTransformFunctor : WinogradTransformFunctorBase { template struct WinogradTransformFunctor : WinogradTransformFunctorBase { - WinogradTransformFunctor(const Padding &paddings) - : WinogradTransformFunctorBase(paddings) {} + WinogradTransformFunctor(const Padding &padding_type, + const std::vector &paddings) + : WinogradTransformFunctorBase(padding_type, paddings) {} void operator()(const Tensor *input, Tensor *output, diff --git a/mace/ops/addn.h b/mace/ops/addn.h index 723fca60..9adc3341 100644 --- a/mace/ops/addn.h +++ b/mace/ops/addn.h @@ -26,7 +26,12 @@ class AddNOp : public Operator { 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); diff --git a/mace/ops/conv_2d.h b/mace/ops/conv_2d.h index 79af4894..b7347b5e 100644 --- a/mace/ops/conv_2d.h +++ b/mace/ops/conv_2d.h @@ -19,7 +19,8 @@ class Conv2dOp : public ConvPool2dOpBase { Conv2dOp(const OperatorDef &op_def, Workspace *ws) : ConvPool2dOpBase(op_def, ws), functor_(this->strides_.data(), - this->padding_, + this->padding_type_, + this->paddings_, this->dilations_.data(), kernels::ActivationType::NOOP, 0.0f, diff --git a/mace/ops/conv_pool_2d_base.h b/mace/ops/conv_pool_2d_base.h index c1c8d3d7..c3db95ab 100644 --- a/mace/ops/conv_pool_2d_base.h +++ b/mace/ops/conv_pool_2d_base.h @@ -16,14 +16,16 @@ class ConvPool2dOpBase : public Operator { ConvPool2dOpBase(const OperatorDef &op_def, Workspace *ws) : Operator(op_def, ws), strides_(OperatorBase::GetRepeatedArgument("strides")), - padding_(static_cast(OperatorBase::GetSingleArgument( + padding_type_(static_cast(OperatorBase::GetSingleArgument( "padding", static_cast(SAME)))), + paddings_(OperatorBase::GetRepeatedArgument("padding_values")), dilations_( OperatorBase::GetRepeatedArgument("dilations", {1, 1})) {} protected: std::vector strides_; - Padding padding_; + Padding padding_type_; + std::vector paddings_; std::vector dilations_; }; diff --git a/mace/ops/depthwise_conv2d.h b/mace/ops/depthwise_conv2d.h index 04ad7f1b..82d6e899 100644 --- a/mace/ops/depthwise_conv2d.h +++ b/mace/ops/depthwise_conv2d.h @@ -20,7 +20,8 @@ class DepthwiseConv2dOp : public ConvPool2dOpBase { DepthwiseConv2dOp(const OperatorDef &op_def, Workspace *ws) : ConvPool2dOpBase(op_def, ws), functor_(this->strides_.data(), - this->padding_, + this->padding_type_, + this->paddings_, this->dilations_.data(), kernels::StringToActivationType( OperatorBase::GetSingleArgument("activation", diff --git a/mace/ops/fused_conv_2d.h b/mace/ops/fused_conv_2d.h index c11cb778..5c5957f1 100644 --- a/mace/ops/fused_conv_2d.h +++ b/mace/ops/fused_conv_2d.h @@ -19,7 +19,8 @@ class FusedConv2dOp : public ConvPool2dOpBase { FusedConv2dOp(const OperatorDef &op_def, Workspace *ws) : ConvPool2dOpBase(op_def, ws), functor_(this->strides_.data(), - this->padding_, + this->padding_type_, + this->paddings_, this->dilations_.data(), kernels::StringToActivationType( OperatorBase::GetSingleArgument("activation", diff --git a/mace/ops/pooling.h b/mace/ops/pooling.h index 43250852..2e4aed62 100644 --- a/mace/ops/pooling.h +++ b/mace/ops/pooling.h @@ -21,7 +21,7 @@ class PoolingOp : public ConvPool2dOpBase { static_cast(OperatorBase::GetSingleArgument( "pooling_type", static_cast(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); diff --git a/mace/ops/winograd_transform.h b/mace/ops/winograd_transform.h index f2cc5f10..71d8a527 100644 --- a/mace/ops/winograd_transform.h +++ b/mace/ops/winograd_transform.h @@ -18,7 +18,8 @@ class WinogradTransformOp : public Operator { WinogradTransformOp(const OperatorDef &op_def, Workspace *ws) : Operator(op_def, ws), functor_(static_cast(OperatorBase::GetSingleArgument( - "padding", static_cast(VALID)))) {} + "padding", static_cast(VALID))), + OperatorBase::GetRepeatedArgument("padding_values")) {} bool Run(StatsFuture *future) override { const Tensor *input_tensor = this->Input(INPUT); -- GitLab