COMPMID-3163: Remove padding from NEDepthwiseConvolutionLayerNativeKernel
Change-Id: Ibbd6bee5c6a4ce4f212b207d17a65b9c33bcfa78
Signed-off-by: Sang-Hoon Park <sang-hoon.park@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4106
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
diff --git a/arm_compute/core/NEON/kernels/NEDepthwiseConvolutionLayerNativeKernel.h b/arm_compute/core/NEON/kernels/NEDepthwiseConvolutionLayerNativeKernel.h
index 2e29234..335a70f 100644
--- a/arm_compute/core/NEON/kernels/NEDepthwiseConvolutionLayerNativeKernel.h
+++ b/arm_compute/core/NEON/kernels/NEDepthwiseConvolutionLayerNativeKernel.h
@@ -25,7 +25,7 @@
#define ARM_COMPUTE_NEDEPTHWISECONVOLUTIONLAYERNATIVEKERNEL_H
#include "arm_compute/core/NEON/INEKernel.h"
-#include "arm_compute/core/utils/misc/Requires.h"
+#include "arm_compute/core/utils/misc/Traits.h"
#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
#include <arm_neon.h>
@@ -92,18 +92,18 @@
// Inherited methods overridden:
void run(const Window &window, const ThreadInfo &info) override;
- BorderSize border_size() const override;
private:
- template < typename T, typename TW, int S, typename std::enable_if < std::is_same<T, float>::value
-#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
- || std::is_same<T, float16_t>::value
-#endif // __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
- ,
- int >::type = 0 >
+ template <typename T>
+ using FloatEnalber = typename std::enable_if<arm_compute::utils::traits::is_floating_point<T>::value, int>::type;
+
+ template <typename T, typename TW, FloatEnalber<T> = 0>
void run_depthwise(const Window &window, bool has_biases);
- template < typename T, typename TW, int S, REQUIRES_TA(std::is_same<T, uint8_t>::value || std::is_same<T, int8_t>::value) >
+ template <typename T>
+ using Quantized8bitEnalber = typename std::enable_if < std::is_same<T, uint8_t>::value || std::is_same<T, int8_t>::value, int >::type;
+
+ template <typename T, typename TW, Quantized8bitEnalber<T> = 0>
void run_depthwise(const Window &window, bool has_biases);
/** Common signature for all the specialised depthwise convolution native functions
@@ -113,7 +113,6 @@
using DepthwiseFunctionPtr = void (NEDepthwiseConvolutionLayerNativeKernel::*)(const Window &window, bool has_biases);
DepthwiseFunctionPtr _func;
- BorderSize _border_size;
const ITensor *_input;
const ITensor *_weights;
const ITensor *_biases;
diff --git a/arm_compute/core/Window.h b/arm_compute/core/Window.h
index 2ba5440..150320a 100644
--- a/arm_compute/core/Window.h
+++ b/arm_compute/core/Window.h
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2016-2019 Arm Limited.
+ * Copyright (c) 2016-2020 Arm Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -45,6 +45,8 @@
static constexpr size_t DimY = 1;
/** Alias for dimension 2 also known as Z dimension */
static constexpr size_t DimZ = 2;
+ /** Alias for dimension 3 also known as W dimension */
+ static constexpr size_t DimW = 3;
/** Default constructor: create a window containing a single element. */
constexpr Window()
diff --git a/arm_compute/runtime/NEON/functions/NEDepthwiseConvolutionLayer.h b/arm_compute/runtime/NEON/functions/NEDepthwiseConvolutionLayer.h
index 3b75bb1..c6b98ed 100644
--- a/arm_compute/runtime/NEON/functions/NEDepthwiseConvolutionLayer.h
+++ b/arm_compute/runtime/NEON/functions/NEDepthwiseConvolutionLayer.h
@@ -249,7 +249,6 @@
private:
NEDepthwiseConvolutionLayerNativeKernel _depthwise_conv_kernel;
- NEFillBorderKernel _fill_border;
NEPermute _permute_input;
NEPermute _permute_weights;
NEPermute _permute_output;
diff --git a/src/core/NEON/kernels/NEDepthwiseConvolutionLayerNativeKernel.cpp b/src/core/NEON/kernels/NEDepthwiseConvolutionLayerNativeKernel.cpp
index fb47879..0a34ee6 100644
--- a/src/core/NEON/kernels/NEDepthwiseConvolutionLayerNativeKernel.cpp
+++ b/src/core/NEON/kernels/NEDepthwiseConvolutionLayerNativeKernel.cpp
@@ -36,46 +36,103 @@
{
namespace
{
-void pad_vectors(std::vector<int> &mult, std::vector<int> &shift, int vec_size)
+constexpr auto data_layout = DataLayout::NHWC;
+const size_t batch_idx = get_data_layout_dimension_index(data_layout, DataLayoutDimension::BATCHES);
+const size_t width_idx = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH);
+const size_t height_idx = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT);
+const size_t channel_idx = get_data_layout_dimension_index(data_layout, DataLayoutDimension::CHANNEL);
+
+constexpr auto dim_manual_loop = Window::Dimension(0, 0, 0);
+constexpr auto dim_single_unit_step = Window::Dimension(0, 1, 1);
+constexpr size_t vector_size = 8;
+
+struct DepthwiseConvolutionRunInfo
{
- ARM_COMPUTE_ERROR_ON(mult.size() != shift.size());
- while(mult.size() % vec_size != 0)
+public:
+ const size_t num_read_elements_per_iteration;
+ const uint32_t x_start;
+ const uint32_t x_end;
+ const uint32_t x_step;
+ const uint32_t x_leftover_start;
+ const size_t input_stride_y;
+ const size_t input_stride_z;
+ const size_t input_max_offset;
+ const size_t weights_width;
+ const size_t weights_height;
+ const size_t weights_stride_y;
+ const size_t weights_stride_z;
+ const size_t conv_stride_x;
+ const size_t conv_stride_y;
+ const size_t conv_pad_left;
+ const size_t conv_pad_top;
+ const size_t input_height;
+ const size_t input_width;
+ const size_t input_depth;
+
+ DepthwiseConvolutionRunInfo(const ITensorInfo &input, const ITensorInfo &weights, const PadStrideInfo &conv_info, const Window &w, uint32_t depth_multiplier = 1)
+ : num_read_elements_per_iteration((depth_multiplier == 1 ? (vector_size / element_size_from_data_type(input.data_type())) : 1)),
+ x_start(w.x().start()),
+ x_end(w.x().end()),
+ x_step(static_cast<uint32_t>(num_read_elements_per_iteration * depth_multiplier)),
+ x_leftover_start(std::max(static_cast<int32_t>(w.x().end()) - static_cast<int32_t>(x_step) + 1, int32_t(0))),
+ input_stride_y(input.strides_in_bytes().y()),
+ input_stride_z(input.strides_in_bytes().z()),
+ input_max_offset(input.strides_in_bytes().z() * input.dimension(height_idx) - (input.padding().bottom + input.padding().top) * input.strides_in_bytes().y()),
+ weights_width(weights.dimension(width_idx)),
+ weights_height(weights.dimension(height_idx)),
+ weights_stride_y(weights.strides_in_bytes().y()),
+ weights_stride_z(weights.strides_in_bytes().z()),
+ conv_stride_x(conv_info.stride().first),
+ conv_stride_y(conv_info.stride().second),
+ conv_pad_left(conv_info.pad_left()),
+ conv_pad_top(conv_info.pad_top()),
+ input_height(input.dimension(height_idx)),
+ input_width(input.dimension(width_idx)),
+ input_depth(input.dimension(channel_idx))
{
- mult.push_back(0);
- shift.push_back(0);
}
+};
+
+inline bool is_valid_input_region(int32_t base_w, uint32_t base_h, uint32_t w, uint32_t h, const DepthwiseConvolutionRunInfo &run_info, const Size2D &dilation)
+{
+ const int32_t current_h = base_h + h * dilation.y();
+ const bool is_valid_h = current_h >= 0 && current_h < static_cast<int32_t>(run_info.input_height);
+
+ const int32_t current_w = base_w + w * dilation.x();
+ const bool is_valid_w = current_w >= 0 && current_w < static_cast<int32_t>(run_info.input_width);
+
+ return is_valid_h && is_valid_w;
}
-template <typename T, int S>
+template <typename T>
void depthwise_loop_multiplier1_fp(const ITensor *input, const ITensor *weights, const ITensor *biases, ITensor *output, const PadStrideInfo &conv_info,
const Size2D &dilation, const Window &window, bool has_biases)
{
- using VectorType = typename wrapper::traits::neon_vector<T, S>::type;
- using TagType = typename wrapper::traits::neon_vector<T, S>::tag_type;
+ constexpr auto element_per_vector = vector_size / sizeof(T);
+ using VectorType = typename wrapper::traits::neon_vector<T, element_per_vector>::type;
+ using TagType = typename wrapper::traits::neon_vector<T, element_per_vector>::tag_type;
- const size_t input_stride_y = input->info()->strides_in_bytes().y();
- const size_t input_stride_z = input->info()->strides_in_bytes().z();
- const size_t input_max_offset = input->info()->strides_in_bytes().z() * input->info()->dimension(2) - (input->info()->padding().bottom + input->info()->padding().top) *
- input->info()->strides_in_bytes().y();
- const size_t weights_width = weights->info()->dimension(1);
- const size_t weights_height = weights->info()->dimension(2);
- const size_t weights_stride_y = weights->info()->strides_in_bytes().y();
- const size_t weights_stride_z = weights->info()->strides_in_bytes().z();
- const size_t conv_stride_x = conv_info.stride().first;
- const size_t conv_stride_y = conv_info.stride().second;
- const size_t conv_pad_left = conv_info.pad_left();
- const size_t conv_pad_top = conv_info.pad_top();
+ const auto run_info = DepthwiseConvolutionRunInfo(*input->info(), *weights->info(), conv_info, window);
+
+ const VectorType zero_vector = wrapper::vdup_n(static_cast<T>(0), TagType{});
+
+ Window execution_window = window;
+ execution_window.set(Window::DimX, dim_single_unit_step);
Window win_input = window;
- win_input.set(Window::DimY, Window::Dimension(0, 0, 0));
- win_input.set(Window::DimZ, Window::Dimension(0, 0, 0));
+ win_input.set(Window::DimX, dim_manual_loop);
+ win_input.set(Window::DimY, dim_manual_loop);
+ win_input.set(Window::DimZ, dim_manual_loop);
Window win_weights = win_input;
- win_weights.set(3, Window::Dimension(0, 0, 0));
+ win_weights.set(Window::DimW, dim_manual_loop);
+
+ Window win_output = window;
+ win_output.set(Window::DimX, dim_manual_loop);
Iterator input_it(input, win_input);
Iterator weights_it(weights, win_weights);
- Iterator output_it(output, window);
+ Iterator output_it(output, win_output);
Iterator biases_it{};
if(has_biases)
@@ -83,38 +140,80 @@
biases_it = Iterator(biases, win_weights);
}
- execute_window_loop(window, [&](const Coordinates & id)
+ execute_window_loop(execution_window, [&](const Coordinates & id)
{
- VectorType acc = wrapper::vdup_n(static_cast<T>(0), TagType{});
+ const int32_t input_y = id.y() * run_info.conv_stride_x - run_info.conv_pad_left;
+ const int32_t input_z = id.z() * run_info.conv_stride_y - run_info.conv_pad_top;
+ const int64_t base_input_offset = input_y * run_info.input_stride_y + input_z * run_info.input_stride_z;
- const int input_y = id.y() * conv_stride_x - conv_pad_left;
- const int input_z = id.z() * conv_stride_y - conv_pad_top;
- int input_offset = input_y * input_stride_y + input_z * input_stride_z;
+ auto const base_weights_ptr = weights_it.ptr();
+ uint32_t x = run_info.x_start;
- auto weights_ptr = weights_it.ptr();
- for(size_t h = 0; h < weights_height; ++h)
+ for(; x < run_info.x_leftover_start; x += run_info.x_step)
{
- int offs = input_offset;
- for(size_t w = 0; w < weights_width; ++w)
- {
- const auto input_vals = wrapper::vload(reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), input_max_offset)));
- const auto weights_vals = wrapper::vload(reinterpret_cast<T *>(weights_ptr + w * weights_stride_y));
+ VectorType acc = zero_vector;
+ auto weights_ptr = base_weights_ptr;
+ int64_t input_offset = base_input_offset;
- acc = wrapper::vmla(acc, weights_vals, input_vals);
- offs += dilation.x() * input_stride_y;
+ for(uint32_t h = 0; h < run_info.weights_height; ++h)
+ {
+ int64_t offs = input_offset + x * sizeof(T);
+ for(uint32_t w = 0; w < run_info.weights_width; ++w)
+ {
+ const bool is_valid_region = is_valid_input_region(input_y, input_z, w, h, run_info, dilation);
+ const auto input_vals = is_valid_region ?
+ wrapper::vload(reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), run_info.input_max_offset))) :
+ zero_vector;
+ const auto weights_vals = wrapper::vload(reinterpret_cast<T *>(weights_ptr + w * run_info.weights_stride_y) + x);
+ acc = wrapper::vmla(acc, weights_vals, input_vals);
+
+ offs += dilation.x() * run_info.input_stride_y;
+ }
+
+ weights_ptr += run_info.weights_stride_z;
+ input_offset += dilation.y() * run_info.input_stride_z;
}
- weights_ptr += weights_stride_z;
- input_offset += dilation.y() * input_stride_z;
+ if(has_biases)
+ {
+ const auto biases_vals = wrapper::vload(reinterpret_cast<T *>(biases_it.ptr()) + x);
+ acc = wrapper::vadd(acc, biases_vals);
+ }
+
+ wrapper::vstore(reinterpret_cast<T *>(output_it.ptr()) + x, acc);
}
- if(has_biases)
+ for(; x < run_info.x_end; ++x)
{
- const auto biases_vals = wrapper::vload(reinterpret_cast<T *>(biases_it.ptr()));
- acc = wrapper::vadd(acc, biases_vals);
- }
+ auto acc_scalar = T{ 0 };
+ auto weights_ptr = base_weights_ptr;
+ int64_t input_offset = base_input_offset;
- wrapper::vstore(reinterpret_cast<T *>(output_it.ptr()), acc);
+ for(size_t h = 0; h < run_info.weights_height; ++h)
+ {
+ int64_t offs = input_offset + x * sizeof(T);
+ for(size_t w = 0; w < run_info.weights_width; ++w)
+ {
+ const bool is_valid_region = is_valid_input_region(input_y, input_z, w, h, run_info, dilation);
+ const auto input_vals = is_valid_region ? *reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), run_info.input_max_offset)) : 0;
+ const auto weights_vals = *(reinterpret_cast<T *>(weights_ptr + w * run_info.weights_stride_y) + x);
+
+ acc_scalar += (input_vals * weights_vals);
+
+ offs += dilation.x() * run_info.input_stride_y;
+ }
+
+ weights_ptr += run_info.weights_stride_z;
+ input_offset += dilation.y() * run_info.input_stride_z;
+ }
+
+ if(has_biases)
+ {
+ const auto biases_vals = *(reinterpret_cast<T *>(biases_it.ptr()) + x);
+ acc_scalar += biases_vals;
+ }
+ *(reinterpret_cast<T *>(output_it.ptr()) + x) = acc_scalar;
+ }
},
input_it, weights_it, biases_it, output_it);
}
@@ -123,31 +222,28 @@
void depthwise_loop_generic_fp(const ITensor *input, const ITensor *weights, const ITensor *biases, ITensor *output, const PadStrideInfo &conv_info,
const Size2D &dilation, unsigned int depth_multiplier, const Window &window, bool has_biases)
{
- const size_t input_stride_y = input->info()->strides_in_bytes().y();
- const size_t input_stride_z = input->info()->strides_in_bytes().z();
- const size_t input_max_offset = input->info()->strides_in_bytes().z() * input->info()->dimension(2) - (input->info()->padding().bottom + input->info()->padding().top) *
- input->info()->strides_in_bytes().y();
- const size_t weights_width = weights->info()->dimension(1);
- const size_t weights_height = weights->info()->dimension(2);
- const size_t weights_stride_y = weights->info()->strides_in_bytes().y();
- const size_t weights_stride_z = weights->info()->strides_in_bytes().z();
- const size_t conv_stride_x = conv_info.stride().first;
- const size_t conv_stride_y = conv_info.stride().second;
- const size_t conv_pad_left = conv_info.pad_left();
- const size_t conv_pad_top = conv_info.pad_top();
+ const auto run_info = DepthwiseConvolutionRunInfo(*input->info(), *weights->info(), conv_info, window, depth_multiplier);
- Window win_input = window;
- win_input.set(Window::DimY, Window::Dimension(0, 0, 0));
- win_input.set(Window::DimZ, Window::Dimension(0, 0, 0));
+ Window execution_window = window;
+ execution_window.set(Window::DimX, Window::Dimension(0, run_info.input_depth, 1));
- Window win_weights = win_input;
- win_weights.set(3, Window::Dimension(0, 0, 0));
+ Window win_input = execution_window;
+ win_input.set(Window::DimX, Window::Dimension(0, run_info.input_depth, 1));
+ win_input.set(Window::DimY, dim_manual_loop);
+ win_input.set(Window::DimZ, dim_manual_loop);
- win_input.set_dimension_step(Window::DimX, 1);
+ Window win_weights = window;
+ win_weights.set_dimension_step(Window::DimX, run_info.x_step);
+ win_weights.set(Window::DimY, dim_manual_loop);
+ win_weights.set(Window::DimZ, dim_manual_loop);
+ win_weights.set(Window::DimW, dim_manual_loop);
+
+ Window win_output = window;
+ win_output.set_dimension_step(Window::DimX, run_info.x_step);
Iterator input_it(input, win_input);
Iterator weights_it(weights, win_weights);
- Iterator output_it(output, window);
+ Iterator output_it(output, win_output);
Iterator biases_it{};
if(has_biases)
@@ -155,33 +251,34 @@
biases_it = Iterator(biases, win_weights);
}
- execute_window_loop(window, [&](const Coordinates & id)
+ execute_window_loop(execution_window, [&](const Coordinates & id)
{
std::vector<T> acc(depth_multiplier, static_cast<T>(0));
- const int input_y = id.y() * conv_stride_x - conv_pad_left;
- const int input_z = id.z() * conv_stride_y - conv_pad_top;
- int input_offset = input_y * input_stride_y + input_z * input_stride_z;
+ const int input_y = id.y() * run_info.conv_stride_x - run_info.conv_pad_left;
+ const int input_z = id.z() * run_info.conv_stride_y - run_info.conv_pad_top;
+ int input_offset = input_y * run_info.input_stride_y + input_z * run_info.input_stride_z;
auto weights_ptr = weights_it.ptr();
- for(size_t h = 0; h < weights_height; ++h)
+ for(size_t h = 0; h < run_info.weights_height; ++h)
{
int offs = input_offset;
- for(size_t w = 0; w < weights_width; ++w)
+ for(size_t w = 0; w < run_info.weights_width; ++w)
{
- const auto input_val = *(reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), input_max_offset)));
+ const bool is_valid_region = is_valid_input_region(input_y, input_z, w, h, run_info, dilation);
+ const auto input_val = is_valid_region ? *(reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), run_info.input_max_offset))) : T(0);
for(size_t m = 0; m < depth_multiplier; ++m)
{
- const auto weights_val = *(reinterpret_cast<T *>(weights_ptr + m * sizeof(T) + w * weights_stride_y));
+ const auto weights_val = *(reinterpret_cast<T *>(weights_ptr + m * sizeof(T) + w * run_info.weights_stride_y));
acc.at(m) = support::cpp11::fma(weights_val, input_val, acc.at(m));
}
- offs += dilation.x() * input_stride_y;
+ offs += dilation.x() * run_info.input_stride_y;
}
- weights_ptr += weights_stride_z;
- input_offset += dilation.y() * input_stride_z;
+ weights_ptr += run_info.weights_stride_z;
+ input_offset += dilation.y() * run_info.input_stride_z;
}
if(has_biases)
@@ -203,41 +300,43 @@
input_it, weights_it, biases_it, output_it);
}
-template <typename T, typename TW, int S>
+template <typename T, typename TW>
void depthwise_loop_multiplier1_quantized(const ITensor *input, const ITensor *weights, const ITensor *biases, ITensor *output, const PadStrideInfo &conv_info,
const Size2D &dilation, std::vector<int> output_multiplier, std::vector<int> output_shift, const Window &window, bool has_biases)
{
- using VectorType = typename wrapper::traits::neon_vector<T, S>::type;
- using TagType = typename wrapper::traits::neon_vector<T, S>::tag_type;
+ constexpr auto element_per_vector = vector_size / sizeof(T);
+ using VectorType = typename wrapper::traits::neon_vector<T, element_per_vector>::type;
+ using TagType = typename wrapper::traits::neon_vector<T, element_per_vector>::tag_type;
+ using AccType = int32_t;
+ using AccArrayType = std::array<AccType, element_per_vector>;
- const size_t input_stride_y = input->info()->strides_in_bytes().y();
- const size_t input_stride_z = input->info()->strides_in_bytes().z();
- const size_t input_max_offset = input->info()->strides_in_bytes().z() * input->info()->dimension(2) - (input->info()->padding().bottom + input->info()->padding().top) *
- input->info()->strides_in_bytes().y();
- const size_t weights_width = weights->info()->dimension(1);
- const size_t weights_height = weights->info()->dimension(2);
- const size_t weights_stride_y = weights->info()->strides_in_bytes().y();
- const size_t weights_stride_z = weights->info()->strides_in_bytes().z();
- const size_t conv_stride_x = conv_info.stride().first;
- const size_t conv_stride_y = conv_info.stride().second;
- const size_t conv_pad_left = conv_info.pad_left();
- const size_t conv_pad_top = conv_info.pad_top();
+ const auto out_of_bound_value = PixelValue(static_cast<uint64_t>(0), input->info()->data_type(), input->info()->quantization_info()).get<T>();
+ const auto out_of_bound_vector = wrapper::vdup_n(static_cast<T>(out_of_bound_value), TagType{});
+
+ const auto run_info = DepthwiseConvolutionRunInfo(*input->info(), *weights->info(), conv_info, window);
const int32_t input_qoffset = input->info()->quantization_info().uniform().offset;
const int32_t weights_qoffset = weights->info()->quantization_info().uniform().offset;
const int32_t output_qoffset = output->info()->quantization_info().uniform().offset;
- const int32_t k_offset = weights_width * weights_height * input_qoffset * weights_qoffset;
+ const int32_t k_offset = run_info.weights_width * run_info.weights_height * input_qoffset * weights_qoffset;
+
+ Window execution_window = window;
+ execution_window.set(Window::DimX, dim_single_unit_step);
Window win_input = window;
- win_input.set(Window::DimY, Window::Dimension(0, 0, 0));
- win_input.set(Window::DimZ, Window::Dimension(0, 0, 0));
+ win_input.set(Window::DimX, dim_manual_loop);
+ win_input.set(Window::DimY, dim_manual_loop);
+ win_input.set(Window::DimZ, dim_manual_loop);
Window win_weights = win_input;
- win_weights.set(3, Window::Dimension(0, 0, 0));
+ win_weights.set(Window::DimW, dim_manual_loop);
+
+ Window win_output = window;
+ win_output.set(Window::DimX, dim_manual_loop);
Iterator input_it(input, win_input);
Iterator weights_it(weights, win_weights);
- Iterator output_it(output, window);
+ Iterator output_it(output, win_output);
Iterator biases_it{};
if(has_biases)
@@ -245,65 +344,134 @@
biases_it = Iterator(biases, win_weights);
}
- execute_window_loop(window, [&](const Coordinates & id)
+ execute_window_loop(execution_window, [&](const Coordinates & id)
{
- std::vector<int32_t> acc(S, 0);
- std::vector<int32_t> in_sum(S, 0);
- std::vector<int32_t> we_sum(S, 0);
+ const int32_t input_y = id.y() * run_info.conv_stride_x - run_info.conv_pad_left;
+ const int32_t input_z = id.z() * run_info.conv_stride_y - run_info.conv_pad_top;
+ const int64_t base_input_offset = input_y * run_info.input_stride_y + input_z * run_info.input_stride_z;
+ auto const base_weights_ptr = weights_it.ptr();
+ size_t x = run_info.x_start;
- const int input_y = id.y() * conv_stride_x - conv_pad_left;
- const int input_z = id.z() * conv_stride_y - conv_pad_top;
- int input_offset = input_y * input_stride_y + input_z * input_stride_z;
-
- auto weights_ptr = weights_it.ptr();
- for(size_t h = 0; h < weights_height; ++h)
+ for(; x < run_info.x_leftover_start; x += run_info.x_step)
{
- int offs = input_offset;
- for(size_t w = 0; w < weights_width; ++w)
- {
- const auto input_vals = wrapper::vload(reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), input_max_offset)));
- const auto weights_vals = wrapper::vload(reinterpret_cast<TW *>(weights_ptr + w * weights_stride_y));
+ AccArrayType acc{};
+ AccArrayType in_sum{};
+ AccArrayType we_sum{};
- for(int i = 0; i < S; ++i)
+ auto weights_ptr = base_weights_ptr;
+ auto input_offset = base_input_offset;
+
+ for(size_t h = 0; h < run_info.weights_height; ++h)
+ {
+ int64_t offs = input_offset + x * sizeof(T);
+ for(size_t w = 0; w < run_info.weights_width; ++w)
{
- acc.at(i) += input_vals[i] * weights_vals[i];
- in_sum.at(i) += input_vals[i];
- we_sum.at(i) += weights_vals[i];
+ const bool is_valid_region = is_valid_input_region(input_y, input_z, w, h, run_info, dilation);
+ const auto input_vals = is_valid_region ?
+ wrapper::vload(reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), run_info.input_max_offset))) :
+ out_of_bound_vector;
+ const auto weights_vals = wrapper::vload(reinterpret_cast<TW *>(weights_ptr + w * run_info.weights_stride_y) + x);
+
+ for(size_t i = 0; i < run_info.x_step; ++i)
+ {
+ acc.at(i) += input_vals[i] * weights_vals[i];
+ in_sum.at(i) += input_vals[i];
+ we_sum.at(i) += weights_vals[i];
+ }
+
+ offs += dilation.x() * run_info.input_stride_y;
}
- offs += dilation.x() * input_stride_y;
+ weights_ptr += run_info.weights_stride_z;
+ input_offset += dilation.y() * run_info.input_stride_z;
}
- weights_ptr += weights_stride_z;
- input_offset += dilation.y() * input_stride_z;
+ VectorType out_vals = wrapper::vdup_n(static_cast<T>(0), TagType{});
+ for(size_t i = 0; i < run_info.x_step; ++i)
+ {
+ acc.at(i) -= in_sum.at(i) * weights_qoffset;
+ acc.at(i) -= we_sum.at(i) * input_qoffset;
+ acc.at(i) += k_offset;
+
+ if(has_biases)
+ {
+ acc.at(i) += *(reinterpret_cast<int32_t *>(biases_it.ptr() + i * sizeof(int32_t)) + x);
+ }
+
+ const int32_t out_mul = output_multiplier.at(x + i);
+ const int32_t out_shift = output_shift.at(x + i);
+ if(out_shift < 0)
+ {
+ acc.at(i) = saturating_doubling_high_mul(acc.at(i) * (1 << (-out_shift)), out_mul) + output_qoffset;
+ }
+ else
+ {
+ acc.at(i) = rounding_divide_by_exp2(saturating_doubling_high_mul(acc.at(i), out_mul), out_shift) + output_qoffset;
+ }
+ out_vals[i] = static_cast<T>(utility::clamp<AccType, T>(acc.at(i)));
+ }
+
+ wrapper::vstore(reinterpret_cast<T *>(output_it.ptr()) + x, out_vals);
}
- VectorType out_vals = wrapper::vdup_n(static_cast<T>(0), TagType{});
- for(int i = 0; i < S; ++i)
+ // left-over
+ for(; x < run_info.x_end; ++x)
{
- acc.at(i) -= in_sum.at(i) * weights_qoffset;
- acc.at(i) -= we_sum.at(i) * input_qoffset;
- acc.at(i) += k_offset;
+ AccType acc = 0;
+ AccType in_sum = 0;
+ AccType we_sum = 0;
+
+ auto weights_ptr = base_weights_ptr;
+ auto input_offset = base_input_offset;
+
+ for(size_t h = 0; h < run_info.weights_height; ++h)
+ {
+ int64_t offs = input_offset + x * sizeof(T);
+ for(size_t w = 0; w < run_info.weights_width; ++w)
+ {
+ const bool is_valid_region = is_valid_input_region(input_y, input_z, w, h, run_info, dilation);
+ const auto input_val = is_valid_region ?
+ *reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), run_info.input_max_offset)) :
+ out_of_bound_value;
+ const auto weights_val = *(reinterpret_cast<TW *>(weights_ptr + w * run_info.weights_stride_y) + x);
+
+ acc += input_val * weights_val;
+ in_sum += input_val;
+ we_sum += weights_val;
+
+ offs += dilation.x() * run_info.input_stride_y;
+ }
+
+ weights_ptr += run_info.weights_stride_z;
+ input_offset += dilation.y() * run_info.input_stride_z;
+ }
+
+ T out_vals{ 0 };
+
+ acc -= in_sum * weights_qoffset;
+ acc -= we_sum * input_qoffset;
+ acc += k_offset;
if(has_biases)
{
- acc.at(i) += *reinterpret_cast<int32_t *>(biases_it.ptr() + i * sizeof(int32_t));
+ acc += *(reinterpret_cast<int32_t *>(biases_it.ptr()) + x);
}
- const int out_mul = output_multiplier.at(id.x() + i);
- const int out_shift = output_shift.at(id.x() + i);
+ const int32_t out_mul = output_multiplier.at(x);
+ const int32_t out_shift = output_shift.at(x);
+
if(out_shift < 0)
{
- acc.at(i) = saturating_doubling_high_mul(acc.at(i) * (1 << (-out_shift)), out_mul) + output_qoffset;
+ acc = saturating_doubling_high_mul(acc * (1 << (-out_shift)), out_mul) + output_qoffset;
}
else
{
- acc.at(i) = rounding_divide_by_exp2(saturating_doubling_high_mul(acc.at(i), out_mul), out_shift) + output_qoffset;
+ acc = rounding_divide_by_exp2(saturating_doubling_high_mul(acc, out_mul), out_shift) + output_qoffset;
}
- out_vals[i] = static_cast<T>(utility::clamp<int32_t, T>(acc.at(i)));
- }
- wrapper::vstore(reinterpret_cast<T *>(output_it.ptr()), out_vals);
+ out_vals = static_cast<T>(utility::clamp<AccType, T>(acc));
+ *(reinterpret_cast<T *>(output_it.ptr()) + x) = out_vals;
+ }
},
input_it, weights_it, biases_it, output_it);
}
@@ -312,36 +480,36 @@
void depthwise_loop_generic_quantized(const ITensor *input, const ITensor *weights, const ITensor *biases, ITensor *output, const PadStrideInfo &conv_info,
const Size2D &dilation, unsigned int depth_multiplier, std::vector<int> output_multiplier, std::vector<int> output_shift, const Window &window, bool has_biases)
{
- const size_t input_stride_y = input->info()->strides_in_bytes().y();
- const size_t input_stride_z = input->info()->strides_in_bytes().z();
- const size_t input_max_offset = input->info()->strides_in_bytes().z() * input->info()->dimension(2) - (input->info()->padding().bottom + input->info()->padding().top) *
- input->info()->strides_in_bytes().y();
- const size_t weights_width = weights->info()->dimension(1);
- const size_t weights_height = weights->info()->dimension(2);
- const size_t weights_stride_y = weights->info()->strides_in_bytes().y();
- const size_t weights_stride_z = weights->info()->strides_in_bytes().z();
- const size_t conv_stride_x = conv_info.stride().first;
- const size_t conv_stride_y = conv_info.stride().second;
- const size_t conv_pad_left = conv_info.pad_left();
- const size_t conv_pad_top = conv_info.pad_top();
+ using AccType = int32_t;
+
+ const auto run_info = DepthwiseConvolutionRunInfo(*input->info(), *weights->info(), conv_info, window, depth_multiplier);
+
+ const auto out_of_bound_value = PixelValue(static_cast<uint64_t>(0), input->info()->data_type(), input->info()->quantization_info()).get<T>();
const int32_t input_qoffset = input->info()->quantization_info().uniform().offset;
const int32_t weights_qoffset = weights->info()->quantization_info().uniform().offset;
const int32_t output_qoffset = output->info()->quantization_info().uniform().offset;
- const int32_t k_offset = weights_width * weights_height * input_qoffset * weights_qoffset;
+ const int32_t k_offset = run_info.weights_width * run_info.weights_height * input_qoffset * weights_qoffset;
- Window win_input = window;
- win_input.set(Window::DimY, Window::Dimension(0, 0, 0));
- win_input.set(Window::DimZ, Window::Dimension(0, 0, 0));
+ Window execution_window = window;
+ execution_window.set(Window::DimX, Window::Dimension(0, run_info.input_depth, 1));
- Window win_weights = win_input;
- win_weights.set(3, Window::Dimension(0, 0, 0));
+ Window win_input = execution_window;
+ win_input.set(Window::DimY, dim_manual_loop);
+ win_input.set(Window::DimZ, dim_manual_loop);
- win_input.set_dimension_step(Window::DimX, 1);
+ Window win_weights = window;
+ win_weights.set_dimension_step(Window::DimX, run_info.x_step);
+ win_weights.set(Window::DimY, dim_manual_loop);
+ win_weights.set(Window::DimZ, dim_manual_loop);
+ win_weights.set(Window::DimW, dim_manual_loop);
+
+ Window win_output = window;
+ win_output.set_dimension_step(Window::DimX, run_info.x_step);
Iterator input_it(input, win_input);
Iterator weights_it(weights, win_weights);
- Iterator output_it(output, window);
+ Iterator output_it(output, win_output);
Iterator biases_it{};
if(has_biases)
@@ -349,38 +517,39 @@
biases_it = Iterator(biases, win_weights);
}
- execute_window_loop(window, [&](const Coordinates & id)
+ execute_window_loop(execution_window, [&](const Coordinates & id)
{
- std::vector<int32_t> acc(depth_multiplier, 0);
- std::vector<int32_t> we_sum(depth_multiplier, 0);
- int32_t in_sum = 0;
+ std::vector<AccType> acc(depth_multiplier, 0);
+ std::vector<AccType> we_sum(depth_multiplier, 0);
+ AccType in_sum = 0;
- const int input_y = id.y() * conv_stride_x - conv_pad_left;
- const int input_z = id.z() * conv_stride_y - conv_pad_top;
- int input_offset = input_y * input_stride_y + input_z * input_stride_z;
+ const int32_t input_y = id.y() * run_info.conv_stride_x - run_info.conv_pad_left;
+ const int32_t input_z = id.z() * run_info.conv_stride_y - run_info.conv_pad_top;
+ int64_t input_offset = input_y * run_info.input_stride_y + input_z * run_info.input_stride_z;
auto weights_ptr = weights_it.ptr();
- for(size_t h = 0; h < weights_height; ++h)
+ for(size_t h = 0; h < run_info.weights_height; ++h)
{
int offs = input_offset;
- for(size_t w = 0; w < weights_width; ++w)
+ for(size_t w = 0; w < run_info.weights_width; ++w)
{
- const auto input_val = *(reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), input_max_offset)));
+ const bool is_valid_region = is_valid_input_region(input_y, input_z, w, h, run_info, dilation);
+ const auto input_val = is_valid_region ? *(reinterpret_cast<T *>(input_it.ptr() + std::min(static_cast<size_t>(offs), run_info.input_max_offset))) : out_of_bound_value;
for(size_t m = 0; m < depth_multiplier; ++m)
{
- const auto weights_val = *(reinterpret_cast<TW *>(weights_ptr + m * sizeof(T) + w * weights_stride_y));
+ const auto weights_val = *(reinterpret_cast<TW *>(weights_ptr + m * sizeof(T) + w * run_info.weights_stride_y));
acc.at(m) += input_val * weights_val;
we_sum.at(m) += weights_val;
}
- offs += dilation.x() * input_stride_y;
+ offs += dilation.x() * run_info.input_stride_y;
in_sum += input_val;
}
- weights_ptr += weights_stride_z;
- input_offset += dilation.y() * input_stride_z;
+ weights_ptr += run_info.weights_stride_z;
+ input_offset += dilation.y() * run_info.input_stride_z;
}
for(size_t m = 0; m < depth_multiplier; ++m)
@@ -394,8 +563,8 @@
acc.at(m) += *(reinterpret_cast<int32_t *>(biases_it.ptr() + m * sizeof(int32_t)));
}
- const int out_mul = output_multiplier.at(id.x() + m);
- const int out_shift = output_shift.at(id.x() + m);
+ const int32_t out_mul = output_multiplier.at(id.x() * depth_multiplier + m);
+ const int32_t out_shift = output_shift.at(id.x() * depth_multiplier + m);
if(out_shift < 0)
{
acc.at(m) = saturating_doubling_high_mul(acc.at(m) * (1 << (-out_shift)), out_mul) + output_qoffset;
@@ -404,7 +573,7 @@
{
acc.at(m) = rounding_divide_by_exp2(saturating_doubling_high_mul(acc.at(m), out_mul), out_shift) + output_qoffset;
}
- *(reinterpret_cast<T *>(output_it.ptr() + m * sizeof(T))) = static_cast<T>(utility::clamp<int32_t, T>(acc.at(m)));
+ *(reinterpret_cast<T *>(output_it.ptr() + m * sizeof(T))) = static_cast<T>(utility::clamp<AccType, T>(acc.at(m)));
}
},
input_it, weights_it, biases_it, output_it);
@@ -458,54 +627,13 @@
return Status{};
}
-
-std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITensorInfo *weights, ITensorInfo *biases,
- ITensorInfo *output, const PadStrideInfo &conv_info,
- unsigned int depth_multiplier, const Size2D &dilation)
-{
- // Get convolved dimensions
- const TensorShape output_shape = misc::shape_calculator::compute_depthwise_convolution_shape(*input, *weights, conv_info, depth_multiplier, dilation);
-
- // Output auto inizialitation if not yet initialized
- auto_init_if_empty(*output, input->clone()->set_is_resizable(true).reset_padding().set_tensor_shape(output_shape).set_quantization_info(output->quantization_info()));
-
- // Configure kernel window (generic)
- const unsigned int num_elems_read_per_iteration = (depth_multiplier == 1) ? 8 / element_size_from_data_type(input->data_type()) : 1;
- const unsigned int num_elems_written_per_iteration = num_elems_read_per_iteration * depth_multiplier;
-
- // Configure kernel window
- Window win = calculate_max_window(*output, Steps(num_elems_written_per_iteration));
-
- AccessWindowStatic input_access(input, 0, -conv_info.pad_left(), ceil_to_multiple(num_elems_read_per_iteration, input->dimension(0)),
- input->dimension(1) + std::max(std::max(conv_info.pad_right(), conv_info.pad_bottom()), conv_info.pad_top()));
- AccessWindowHorizontal weights_access(weights, 0, num_elems_written_per_iteration);
- AccessWindowHorizontal output_access(output, 0, num_elems_written_per_iteration);
-
- bool window_changed = update_window_and_padding(win, input_access, weights_access, output_access);
-
- if(biases != nullptr)
- {
- AccessWindowHorizontal biases_access(biases, 0, num_elems_written_per_iteration);
- window_changed |= update_window_and_padding(win, biases_access);
- }
-
- output_access.set_valid_region(win, ValidRegion(Coordinates(), output->tensor_shape()));
-
- Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
- return std::make_pair(err, win);
-}
} // namespace
NEDepthwiseConvolutionLayerNativeKernel::NEDepthwiseConvolutionLayerNativeKernel()
- : _func(), _border_size(0), _input(), _weights(), _biases(), _output(), _conv_info(), _depth_multiplier(1), _dilation(), _output_multiplier(), _output_shift(), _has_biases()
+ : _func(), _input(), _weights(), _biases(), _output(), _conv_info(), _depth_multiplier(1), _dilation(), _output_multiplier(), _output_shift(), _has_biases()
{
}
-BorderSize NEDepthwiseConvolutionLayerNativeKernel::border_size() const
-{
- return _border_size;
-}
-
void NEDepthwiseConvolutionLayerNativeKernel::configure(const ITensor *input, const ITensor *weights, const ITensor *biases, ITensor *output,
const PadStrideInfo &conv_info, unsigned int depth_multiplier, const Size2D &dilation)
{
@@ -518,7 +646,6 @@
_output = output;
_conv_info = conv_info;
_depth_multiplier = depth_multiplier;
- _border_size = BorderSize(_conv_info.pad_left(), 0, std::max(std::max(conv_info.pad_right(), conv_info.pad_bottom()), conv_info.pad_top()), 0);
_dilation = dilation;
_has_biases = (biases != nullptr);
@@ -530,17 +657,17 @@
auto weights_scale = weights->info()->quantization_info().scale();
if(!is_data_type_quantized_per_channel(_weights->info()->data_type()))
{
- for(size_t i = 1; i < _weights->info()->dimension(0); ++i)
+ for(size_t i = 1; i < _weights->info()->dimension(channel_idx); ++i)
{
weights_scale.push_back(weights_scale.front());
}
}
- for(size_t i = 0; i < weights_scale.size(); ++i)
+ for(const auto &s : weights_scale)
{
int32_t out_mult = 0;
int32_t out_shift = 0;
- const float multiplier = input_scale * weights_scale.at(i) / output_scale;
+ const float multiplier = input_scale * s / output_scale;
arm_compute::quantization::calculate_quantized_multiplier(multiplier, &out_mult, &out_shift);
_output_multiplier.push_back(out_mult);
@@ -551,42 +678,42 @@
switch(_weights->info()->data_type())
{
case DataType::QASYMM8:
- _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<uint8_t, uint8_t, 8>;
- pad_vectors(_output_multiplier, _output_shift, 8);
+ _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<uint8_t, uint8_t>;
break;
case DataType::QASYMM8_SIGNED:
- _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<int8_t, int8_t, 8>;
- pad_vectors(_output_multiplier, _output_shift, 8);
+ _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<int8_t, int8_t>;
break;
case DataType::QSYMM8_PER_CHANNEL:
if(_input->info()->data_type() == DataType::QASYMM8)
{
- _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<uint8_t, int8_t, 8>;
+ _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<uint8_t, int8_t>;
}
else
{
- _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<int8_t, int8_t, 8>;
+ _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<int8_t, int8_t>;
}
- pad_vectors(_output_multiplier, _output_shift, 8);
break;
#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
case DataType::F16:
- _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<float16_t, float16_t, 4>;
- pad_vectors(_output_multiplier, _output_shift, 4);
+ _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<float16_t, float16_t>;
break;
#endif // __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
case DataType::F32:
- _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<float, float, 2>;
- pad_vectors(_output_multiplier, _output_shift, 2);
+ _func = &NEDepthwiseConvolutionLayerNativeKernel::run_depthwise<float, float>;
break;
default:
ARM_COMPUTE_ERROR("Data type not supported");
break;
}
- auto win_config = validate_and_configure_window(_input->info(), _weights->info(), (biases != nullptr) ? biases->info() : nullptr, _output->info(), _conv_info, _depth_multiplier, dilation);
- ARM_COMPUTE_ERROR_THROW_ON(win_config.first);
- INEKernel::configure(win_config.second);
+ const TensorShape output_shape = misc::shape_calculator::compute_depthwise_convolution_shape(*input->info(), *weights->info(), conv_info, depth_multiplier, dilation);
+ auto_init_if_empty(*output->info(), input->info()->clone()->set_is_resizable(true).reset_padding().set_tensor_shape(output_shape).set_quantization_info(output->info()->quantization_info()));
+
+ Window win = calculate_max_window(*output->info(), Steps());
+ Coordinates coord;
+ coord.set_num_dimensions(output->info()->num_dimensions());
+ output->info()->set_valid_region(ValidRegion(coord, output->info()->tensor_shape()));
+ INEKernel::configure(win);
}
Status NEDepthwiseConvolutionLayerNativeKernel::validate(const ITensorInfo *input, const ITensorInfo *weights, const ITensorInfo *biases, const ITensorInfo *output, const PadStrideInfo &conv_info,
@@ -594,9 +721,6 @@
const Size2D &dilation)
{
ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, weights, biases, output, conv_info, depth_multiplier, dilation));
- ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), weights->clone().get(), (biases != nullptr) ? biases->clone().get() : nullptr, output->clone().get(), conv_info,
- depth_multiplier, dilation)
- .first);
return Status{};
}
@@ -609,12 +733,7 @@
(this->*_func)(window, _has_biases);
}
-template < typename T, typename TW, int S, typename std::enable_if < std::is_same<T, float>::value
-#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
- || std::is_same<T, float16_t>::value
-#endif // __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
- ,
- int >::type >
+template <typename T, typename TW, NEDepthwiseConvolutionLayerNativeKernel::FloatEnalber<T>>
void NEDepthwiseConvolutionLayerNativeKernel::run_depthwise(const Window &window, bool has_biases)
{
ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
@@ -622,7 +741,7 @@
if(_depth_multiplier == 1)
{
- depthwise_loop_multiplier1_fp<T, S>(_input, _weights, _biases, _output, _conv_info, _dilation, window, has_biases);
+ depthwise_loop_multiplier1_fp<T>(_input, _weights, _biases, _output, _conv_info, _dilation, window, has_biases);
}
else
{
@@ -630,7 +749,7 @@
}
}
-template <typename T, typename TW, int S, typename>
+template <typename T, typename TW, NEDepthwiseConvolutionLayerNativeKernel::Quantized8bitEnalber<T>>
void NEDepthwiseConvolutionLayerNativeKernel::run_depthwise(const Window &window, bool has_biases)
{
ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
@@ -638,7 +757,7 @@
if(_depth_multiplier == 1)
{
- depthwise_loop_multiplier1_quantized<T, TW, S>(_input, _weights, _biases, _output, _conv_info, _dilation, _output_multiplier, _output_shift, window, has_biases);
+ depthwise_loop_multiplier1_quantized<T, TW>(_input, _weights, _biases, _output, _conv_info, _dilation, _output_multiplier, _output_shift, window, has_biases);
}
else
{
diff --git a/src/runtime/NEON/functions/NEDepthwiseConvolutionLayer.cpp b/src/runtime/NEON/functions/NEDepthwiseConvolutionLayer.cpp
index 915a283..6c22523 100644
--- a/src/runtime/NEON/functions/NEDepthwiseConvolutionLayer.cpp
+++ b/src/runtime/NEON/functions/NEDepthwiseConvolutionLayer.cpp
@@ -210,8 +210,8 @@
}
NEDepthwiseConvolutionLayer::NEDepthwiseConvolutionLayerGeneric::NEDepthwiseConvolutionLayerGeneric()
- : _depthwise_conv_kernel(), _fill_border(), _permute_input(), _permute_weights(), _permute_output(), _activationlayer_function(), _permuted_input(), _permuted_weights(), _permuted_output(),
- _is_prepared(false), _is_nchw(false), _is_activationlayer_enabled(false), _original_weights(nullptr)
+ : _depthwise_conv_kernel(), _permute_input(), _permute_weights(), _permute_output(), _activationlayer_function(), _permuted_input(), _permuted_weights(), _permuted_output(), _is_prepared(false),
+ _is_nchw(false), _is_activationlayer_enabled(false), _original_weights(nullptr)
{
}
@@ -244,7 +244,6 @@
_original_weights = weights_to_use;
_depthwise_conv_kernel.configure(input_to_use, weights_to_use, biases, output_to_use, conv_info, depth_multiplier, dilation);
- _fill_border.configure(input_to_use, _depthwise_conv_kernel.border_size(), BorderMode::CONSTANT, PixelValue(static_cast<uint64_t>(0), input->info()->data_type(), input->info()->quantization_info()));
if(_is_nchw)
{
@@ -310,7 +309,6 @@
_permute_input.run();
}
- NEScheduler::get().schedule(&_fill_border, Window::DimX);
NEScheduler::get().schedule(&_depthwise_conv_kernel, Window::DimY);
if(_is_nchw)
diff --git a/tests/validation/NEON/DepthwiseConvolutionLayer.cpp b/tests/validation/NEON/DepthwiseConvolutionLayer.cpp
index 407ebe3..e255fc7 100644
--- a/tests/validation/NEON/DepthwiseConvolutionLayer.cpp
+++ b/tests/validation/NEON/DepthwiseConvolutionLayer.cpp
@@ -177,8 +177,6 @@
TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F32), // Invalid output size
TensorInfo(TensorShape(27U, 13U, 8U), 1, DataType::F32), // Patch size bigger than input width
TensorInfo(TensorShape(27U, 13U, 8U), 1, DataType::F32), // Dilation < 1
- TensorInfo(TensorShape(27U, 13U, 8U), 1, DataType::F32), // Window shrinking
- TensorInfo(TensorShape(32U, 13U, 8U), 1, DataType::QASYMM8), // Window shrinking
}),
framework::dataset::make("WeightsInfo", { TensorInfo(TensorShape(3U, 3U, 2U), 1, DataType::F16),
TensorInfo(TensorShape(3U, 3U, 2U), 1, DataType::F32),
@@ -188,8 +186,6 @@
TensorInfo(TensorShape(3U, 3U, 2U), 1, DataType::F32),
TensorInfo(TensorShape(3U, 3U, 16U), 1, DataType::F32),
TensorInfo(TensorShape(3U, 3U, 16U), 1, DataType::F32),
- TensorInfo(TensorShape(3U, 3U, 16U), 1, DataType::F32),
- TensorInfo(TensorShape(3U, 3U, 24U), 1, DataType::QASYMM8),
})),
framework::dataset::make("BiasesInfo", { TensorInfo(TensorShape(2U), 1, DataType::F32),
TensorInfo(TensorShape(2U), 1, DataType::F32),
@@ -199,8 +195,6 @@
TensorInfo(TensorShape(2U), 1, DataType::F32),
TensorInfo(TensorShape(16U), 1, DataType::F32),
TensorInfo(TensorShape(16U), 1, DataType::F32),
- TensorInfo(TensorShape(16U), 1, DataType::F32),
- TensorInfo(TensorShape(24U), 1, DataType::S32),
})),
framework::dataset::make("OutputInfo", { TensorInfo(TensorShape(25U, 11U, 2U), 1, DataType::F32),
TensorInfo(TensorShape(25U, 11U, 2U), 1, DataType::F32),
@@ -210,8 +204,6 @@
TensorInfo(TensorShape(27U, 13U, 2U), 1, DataType::F32),
TensorInfo(TensorShape(25U, 11U, 16U), 1, DataType::F32),
TensorInfo(TensorShape(25U, 11U, 16U), 1, DataType::F32),
- TensorInfo(TensorShape(25U, 11U, 16U), 1, DataType::F32),
- TensorInfo(TensorShape(32U, 11U, 24U), 1, DataType::QASYMM8),
})),
framework::dataset::make("ConvInfo", { PadStrideInfo(1, 1, 0, 0),
PadStrideInfo(1, 1, 0, 0),
@@ -221,8 +213,6 @@
PadStrideInfo(1, 1, 0, 0),
PadStrideInfo(1, 1, 0, 0),
PadStrideInfo(1, 1, 0, 0),
- PadStrideInfo(1, 1, 0, 0),
- PadStrideInfo(1, 1, 1, 0),
})),
framework::dataset::make("DepthMultiplier", { 1,
1,
@@ -232,8 +222,6 @@
1,
2,
2,
- 2,
- 3,
})),
framework::dataset::make("Dilation", { Size2D(1U, 1U),
Size2D(1U, 1U),
@@ -243,10 +231,8 @@
Size2D(1U, 1U),
Size2D(25U, 1U),
Size2D(0U, 1U),
- Size2D(1U, 1U),
- Size2D(1U, 1U),
})),
- framework::dataset::make("Expected", { false, false, false, false, false, false,false, false, false, false })),
+ framework::dataset::make("Expected", { false, false, false, false, false, false, false, false})),
input_info, weights_info, biases_info, output_info, conv_info, depth_multiplier,dilation, expected)
{
bool is_valid = bool(NEDepthwiseConvolutionLayer::validate(&input_info.clone()->set_is_resizable(false), &weights_info.clone()->set_is_resizable(false), &biases_info.clone()->set_is_resizable(false), &output_info.clone()->set_is_resizable(false), conv_info, depth_multiplier, ActivationLayerInfo(), dilation));
diff --git a/tests/validation/NEON/DepthwiseConvolutionLayerNative.cpp b/tests/validation/NEON/DepthwiseConvolutionLayerNative.cpp
index be0eee3..4755135 100644
--- a/tests/validation/NEON/DepthwiseConvolutionLayerNative.cpp
+++ b/tests/validation/NEON/DepthwiseConvolutionLayerNative.cpp
@@ -102,6 +102,37 @@
TEST_SUITE(NEON)
TEST_SUITE(DepthwiseConvolutionLayerNative)
+
+TEST_CASE(ValidateNoPadding, framework::DatasetMode::ALL)
+{
+ // this test case will ensure that the kernel is not adding implicit padding
+ constexpr uint32_t vector_size = 8; // Asummed vector size of the current native kernel
+ constexpr auto depth = vector_size * 2 + 1; // mis-aligned depth to force padding if exists.
+ constexpr auto data_layout = DataLayout::NHWC;
+ constexpr auto data_type = DataType::F32;
+
+ const auto input_size = Size2D{ 100, 100 }; // random plane size of the input
+ const auto kernel_size = Size2D{ 4, 4 }; // random plane size of the kernel
+ const auto pad_stride_info = PadStrideInfo(3, 3); // random convolution information to
+
+ TensorShape src_shape{ depth, input_size.x(), input_size.y() };
+ TensorShape weights_shape{ depth, kernel_size.x(), kernel_size.y() };
+ TensorShape bias_shape{ depth };
+
+ auto src = create_tensor<Tensor>(src_shape, data_type, 1, QuantizationInfo(), data_layout);
+ auto weights = create_tensor<Tensor>(weights_shape, data_type, 1, QuantizationInfo(), data_layout);
+ auto biases = create_tensor<Tensor>(bias_shape, data_type, 1, QuantizationInfo(), data_layout);
+ auto dst = create_tensor<Tensor>(TensorShape(), data_type, 1, QuantizationInfo(), data_layout);
+
+ NEDepthwiseConvolutionLayerNativeKernel dwc;
+ dwc.configure(&src, &weights, &biases, &dst, pad_stride_info);
+
+ ARM_COMPUTE_EXPECT(src.info()->padding().empty(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(weights.info()->padding().empty(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(biases.info()->padding().empty(), framework::LogLevel::ERRORS);
+ ARM_COMPUTE_EXPECT(dst.info()->padding().empty(), framework::LogLevel::ERRORS);
+}
+
TEST_SUITE(Float)
TEST_SUITE(FP32)
FIXTURE_DATA_TEST_CASE(RunSmall, NEDepthwiseConvolutionLayerNativeFixture<float>, framework::DatasetMode::ALL,