diff options
| -rw-r--r-- | arm_compute/core/NEON/kernels/NEPoolingLayerKernel.h | 12 | ||||
| -rw-r--r-- | arm_compute/core/Types.h | 15 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/pooling_layer.cl | 16 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLPoolingLayerKernel.cpp | 9 | ||||
| -rw-r--r-- | src/core/NEON/kernels/NEPoolingLayerKernel.cpp | 96 | ||||
| -rw-r--r-- | tests/validation/CL/PoolingLayer.cpp | 11 | ||||
| -rw-r--r-- | tests/validation/CPP/PoolingLayer.cpp | 36 | ||||
| -rw-r--r-- | tests/validation/NEON/PoolingLayer.cpp | 10 | ||||
| -rw-r--r-- | tests/validation/fixtures/PoolingLayerFixture.h | 10 |
9 files changed, 131 insertions, 84 deletions
diff --git a/arm_compute/core/NEON/kernels/NEPoolingLayerKernel.h b/arm_compute/core/NEON/kernels/NEPoolingLayerKernel.h index 9d7c75179a..0a57a26f17 100644 --- a/arm_compute/core/NEON/kernels/NEPoolingLayerKernel.h +++ b/arm_compute/core/NEON/kernels/NEPoolingLayerKernel.h @@ -66,14 +66,14 @@ private: * @param[in] window_input Input region on which to execute the kernel. * @param[in] window Output region on which to execute the kernel. */ - template <PoolingType pooling_type> + template <PoolingType pooling_type, bool exclude_padding = false> void pooling2_f32(const Window &window_input, const Window &window); /** Function to perform 2x2 pooling for float16_t. * * @param[in] window_input Input region on which to execute the kernel. * @param[in] window Output region on which to execute the kernel. */ - template <PoolingType pooling_type> + template <PoolingType pooling_type, bool exclude_padding = false> void pooling2_f16(const Window &window_input, const Window &window); /** Function to perform 2x2 pooling for 8bit fixed point. @@ -95,14 +95,14 @@ private: * @param[in] window_input Input region on which to execute the kernel. * @param[in] window Output region on which to execute the kernel. */ - template <PoolingType pooling_type> + template <PoolingType pooling_type, bool exclude_padding = false> void pooling3_f32(const Window &window_input, const Window &window); /** Function to perform 3x3 pooling. * * @param[in] window_input Input region on which to execute the kernel. * @param[in] window Output region on which to execute the kernel. */ - template <PoolingType pooling_type> + template <PoolingType pooling_type, bool exclude_padding = false> void pooling3_f16(const Window &window_input, const Window &window); /** Function to perform 3x3 pooling for 8bit fixed point. * @@ -123,14 +123,14 @@ private: * @param[in] window_input Input region on which to execute the kernel. * @param[in] window Output region on which to execute the kernel. */ - template <PoolingType pooling_type> + template <PoolingType pooling_type, bool exclude_padding = false> void pooling7_f32(const Window &window_input, const Window &window); /** Function to perform NxN pooling. * * @param[in] window_input Input region on which to execute the kernel. * @param[in] window Output region on which to execute the kernel. */ - template <PoolingType pooling_type> + template <PoolingType pooling_type, bool exclude_padding = false> void poolingN_f32(const Window &window_input, const Window &window); /** Common signature for all the specialised Pooling functions * diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h index e567bac860..e8be6127a8 100644 --- a/arm_compute/core/Types.h +++ b/arm_compute/core/Types.h @@ -545,9 +545,15 @@ public: * @param[in] pool_type Pooling type @ref PoolingType. Defaults to @ref PoolingType::MAX * @param[in] pool_size (Optional) Pooling size, in elements, across x and y. Defaults to 2. * @param[in] pad_stride_info (Optional) Padding and stride information @ref PadStrideInfo + * @param[in] exclude_padding (Optional) Strategy when accounting padding in calculations. + * True will exclude padding while false will not (Used in AVG/L2 pooling to determine the pooling area). + * Defaults to false; */ - PoolingLayerInfo(PoolingType pool_type = PoolingType::MAX, unsigned int pool_size = 2, PadStrideInfo pad_stride_info = PadStrideInfo()) - : _pool_type(pool_type), _pool_size(pool_size), _pad_stride_info(pad_stride_info) + PoolingLayerInfo(PoolingType pool_type = PoolingType::MAX, + unsigned int pool_size = 2, + PadStrideInfo pad_stride_info = PadStrideInfo(), + bool exclude_padding = false) + : _pool_type(pool_type), _pool_size(pool_size), _pad_stride_info(pad_stride_info), _exclude_padding(exclude_padding) { } PoolingType pool_type() const @@ -562,11 +568,16 @@ public: { return _pad_stride_info; } + bool exclude_padding() const + { + return _exclude_padding; + } private: PoolingType _pool_type; unsigned int _pool_size; PadStrideInfo _pad_stride_info; + bool _exclude_padding; }; /** ROI Pooling Layer Information class */ diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl index 6379c288b3..635c44a849 100644 --- a/src/core/CL/cl_kernels/pooling_layer.cl +++ b/src/core/CL/cl_kernels/pooling_layer.cl @@ -186,10 +186,14 @@ DATA_TYPE calculate_avg_scale(const int pool_size, const int upper_bound_w, const int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y) { - const int start_x = get_global_id(0) * stride_x - pad_x; - const int start_y = get_global_id(1) * stride_y - pad_y; + int start_x = get_global_id(0) * stride_x - pad_x; + int start_y = get_global_id(1) * stride_y - pad_y; const int end_x = min(start_x + pool_size, upper_bound_w); const int end_y = min(start_y + pool_size, upper_bound_h); +#if defined(EXCLUDE_PADDING) + start_x = max(0, start_x); + start_y = max(0, start_y); +#endif /* defined(EXCLUDE_PADDING) */ return ((end_y - start_y) * (end_x - start_x)); } @@ -334,10 +338,14 @@ VEC_DATA_TYPE(DATA_TYPE, 4) calculate_avg_scale4(const int pool_size, const int upper_bound_w, const int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y) { - const int4 start_x = ((int4)get_global_id(0) * 4 + (int4)(0, 1, 2, 3)) * (int4)stride_x - (int4)pad_x; - const int start_y = get_global_id(1) * stride_y - pad_y; + int4 start_x = ((int4)get_global_id(0) * 4 + (int4)(0, 1, 2, 3)) * (int4)stride_x - (int4)pad_x; + int start_y = get_global_id(1) * stride_y - pad_y; const int4 end_x = min(start_x + (int4)pool_size, (int4)upper_bound_w); const int end_y = min(start_y + pool_size, upper_bound_h); +#if defined(EXCLUDE_PADDING) + start_x = max((int4)0, start_x); + start_y = max(0, start_y); +#endif /* defined(EXCLUDE_PADDING) */ return (VEC_DATA_TYPE(DATA_TYPE, 4))(1.f) / CONVERT_VECTOR4(DATA_TYPE)(((int4)(end_y - start_y)) * (end_x - start_x)); } diff --git a/src/core/CL/kernels/CLPoolingLayerKernel.cpp b/src/core/CL/kernels/CLPoolingLayerKernel.cpp index 542d5dcf9f..8b8f61e621 100644 --- a/src/core/CL/kernels/CLPoolingLayerKernel.cpp +++ b/src/core/CL/kernels/CLPoolingLayerKernel.cpp @@ -61,6 +61,7 @@ void CLPoolingLayerKernel::configure(const ICLTensor *input, ICLTensor *output, const PoolingType pool_type = pool_info.pool_type(); const int pool_size = pool_info.pool_size(); const PadStrideInfo pad_stride_info = pool_info.pad_stride_info(); + bool exclude_padding = pool_info.exclude_padding(); std::tie(pool_pad_x, pool_pad_y) = pad_stride_info.pad(); std::tie(pool_stride_x, pool_stride_y) = pad_stride_info.stride(); @@ -109,8 +110,12 @@ void CLPoolingLayerKernel::configure(const ICLTensor *input, ICLTensor *output, build_opts.emplace(("-DSTRIDE_X=" + support::cpp11::to_string(pool_stride_x))); if(pool_type != PoolingType::MAX) { - build_opts.emplace(("-DMAX_WIDTH=" + support::cpp11::to_string(input->info()->dimension(0) + pool_pad_x))); - build_opts.emplace(("-DMAX_HEIGHT=" + support::cpp11::to_string(input->info()->dimension(1) + pool_pad_y))); + if(exclude_padding) + { + build_opts.emplace("-DEXCLUDE_PADDING"); + } + build_opts.emplace(("-DMAX_WIDTH=" + support::cpp11::to_string(input->info()->dimension(0) + (exclude_padding ? 0 : pool_pad_x)))); + build_opts.emplace(("-DMAX_HEIGHT=" + support::cpp11::to_string(input->info()->dimension(1) + (exclude_padding ? 0 : pool_pad_y)))); build_opts.emplace(("-DSTRIDE_Y=" + support::cpp11::to_string(pool_stride_y))); build_opts.emplace(("-DPAD_X=" + support::cpp11::to_string(pool_pad_x))); build_opts.emplace(("-DPAD_Y=" + support::cpp11::to_string(pool_pad_y))); diff --git a/src/core/NEON/kernels/NEPoolingLayerKernel.cpp b/src/core/NEON/kernels/NEPoolingLayerKernel.cpp index 0024e33723..122540b07e 100644 --- a/src/core/NEON/kernels/NEPoolingLayerKernel.cpp +++ b/src/core/NEON/kernels/NEPoolingLayerKernel.cpp @@ -47,13 +47,19 @@ using namespace arm_compute; namespace { +template <bool exclude_padding> inline float calculate_avg_scale(const Coordinates &id, const int pool_size, const int upper_bound_w, const int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y) { - const int start_x = id.x() * stride_x - pad_x; - const int start_y = id.y() * stride_y - pad_y; + int start_x = id.x() * stride_x - pad_x; + int start_y = id.y() * stride_y - pad_y; const int end_x = std::min(start_x + pool_size, upper_bound_w); const int end_y = std::min(start_y + pool_size, upper_bound_h); + if(exclude_padding) + { + start_x = std::max(0, start_x); + start_y = std::max(0, start_y); + } return 1.f / ((end_y - start_y) * (end_x - start_x)); } @@ -105,6 +111,7 @@ void NEPoolingLayerKernel::configure(const ITensor *input, ITensor *output, cons PoolingType pool_type = pool_info.pool_type(); int pool_size = pool_info.pool_size(); const PadStrideInfo pad_stride_info = pool_info.pad_stride_info(); + bool exclude_padding = pool_info.exclude_padding(); std::tie(pool_pad_x, pool_pad_y) = pad_stride_info.pad(); std::tie(pool_stride_x, pool_stride_y) = pad_stride_info.stride(); @@ -117,6 +124,7 @@ void NEPoolingLayerKernel::configure(const ITensor *input, ITensor *output, cons ARM_COMPUTE_ERROR_ON((supported_pool_sizes.find(pool_size) == supported_pool_sizes.end()) && (input->info()->data_type() != DataType::F32)); ARM_COMPUTE_ERROR_ON(pool_pad_x >= pool_size || pool_pad_y >= pool_size); ARM_COMPUTE_ERROR_ON(is_data_type_fixed_point(input->info()->data_type()) && pool_stride_x > 2); + ARM_COMPUTE_ERROR_ON(exclude_padding && is_data_type_fixed_point(input->info()->data_type())); // Check output dimensions std::tie(pooled_w, pooled_h) = scaled_dimensions(input->info()->dimension(0), input->info()->dimension(1), @@ -268,13 +276,13 @@ void NEPoolingLayerKernel::configure(const ITensor *input, ITensor *output, cons switch(pool_type) { case PoolingType::AVG: - _func = &NEPoolingLayerKernel::pooling2_f16<PoolingType::AVG>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling2_f16<PoolingType::AVG, true> : &NEPoolingLayerKernel::pooling2_f16<PoolingType::AVG, false>; break; case PoolingType::L2: - _func = &NEPoolingLayerKernel::pooling2_f16<PoolingType::L2>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling2_f16<PoolingType::L2, true> : &NEPoolingLayerKernel::pooling2_f16<PoolingType::L2, false>; break; case PoolingType::MAX: - _func = &NEPoolingLayerKernel::pooling2_f16<PoolingType::MAX>; + _func = &NEPoolingLayerKernel::pooling2_f16<PoolingType::MAX, false>; break; default: ARM_COMPUTE_ERROR("Unsupported pooling type!"); @@ -285,13 +293,13 @@ void NEPoolingLayerKernel::configure(const ITensor *input, ITensor *output, cons switch(pool_type) { case PoolingType::AVG: - _func = &NEPoolingLayerKernel::pooling2_f32<PoolingType::AVG>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling2_f32<PoolingType::AVG, true> : &NEPoolingLayerKernel::pooling2_f32<PoolingType::AVG, false>; break; case PoolingType::L2: - _func = &NEPoolingLayerKernel::pooling2_f32<PoolingType::L2>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling2_f32<PoolingType::L2, true> : &NEPoolingLayerKernel::pooling2_f32<PoolingType::L2, false>; break; case PoolingType::MAX: - _func = &NEPoolingLayerKernel::pooling2_f32<PoolingType::MAX>; + _func = &NEPoolingLayerKernel::pooling2_f32<PoolingType::MAX, false>; break; default: ARM_COMPUTE_ERROR("Unsupported pooling type!"); @@ -332,13 +340,13 @@ void NEPoolingLayerKernel::configure(const ITensor *input, ITensor *output, cons switch(pool_type) { case PoolingType::AVG: - _func = &NEPoolingLayerKernel::pooling3_f16<PoolingType::AVG>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling3_f16<PoolingType::AVG, true> : &NEPoolingLayerKernel::pooling3_f16<PoolingType::AVG, false>; break; case PoolingType::L2: - _func = &NEPoolingLayerKernel::pooling3_f16<PoolingType::L2>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling3_f16<PoolingType::L2, true> : &NEPoolingLayerKernel::pooling3_f16<PoolingType::L2, false>; break; case PoolingType::MAX: - _func = &NEPoolingLayerKernel::pooling3_f16<PoolingType::MAX>; + _func = &NEPoolingLayerKernel::pooling3_f16<PoolingType::MAX, false>; break; default: ARM_COMPUTE_ERROR("Unsupported pooling type!"); @@ -349,13 +357,13 @@ void NEPoolingLayerKernel::configure(const ITensor *input, ITensor *output, cons switch(pool_type) { case PoolingType::AVG: - _func = &NEPoolingLayerKernel::pooling3_f32<PoolingType::AVG>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling3_f32<PoolingType::AVG, true> : &NEPoolingLayerKernel::pooling3_f32<PoolingType::AVG, false>; break; case PoolingType::L2: - _func = &NEPoolingLayerKernel::pooling3_f32<PoolingType::L2>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling3_f32<PoolingType::L2, true> : &NEPoolingLayerKernel::pooling3_f32<PoolingType::L2, false>; break; case PoolingType::MAX: - _func = &NEPoolingLayerKernel::pooling3_f32<PoolingType::MAX>; + _func = &NEPoolingLayerKernel::pooling3_f32<PoolingType::MAX, false>; break; default: ARM_COMPUTE_ERROR("Unsupported pooling type!"); @@ -366,13 +374,13 @@ void NEPoolingLayerKernel::configure(const ITensor *input, ITensor *output, cons switch(pool_type) { case PoolingType::AVG: - _func = &NEPoolingLayerKernel::pooling7_f32<PoolingType::AVG>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling7_f32<PoolingType::AVG, true> : &NEPoolingLayerKernel::pooling7_f32<PoolingType::AVG, false>; break; case PoolingType::L2: - _func = &NEPoolingLayerKernel::pooling7_f32<PoolingType::L2>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::pooling7_f32<PoolingType::L2, true> : &NEPoolingLayerKernel::pooling7_f32<PoolingType::L2, false>; break; case PoolingType::MAX: - _func = &NEPoolingLayerKernel::pooling7_f32<PoolingType::MAX>; + _func = &NEPoolingLayerKernel::pooling7_f32<PoolingType::MAX, false>; break; default: ARM_COMPUTE_ERROR("Unsupported pooling type!"); @@ -382,13 +390,13 @@ void NEPoolingLayerKernel::configure(const ITensor *input, ITensor *output, cons switch(pool_type) { case PoolingType::AVG: - _func = &NEPoolingLayerKernel::poolingN_f32<PoolingType::AVG>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::poolingN_f32<PoolingType::AVG, true> : &NEPoolingLayerKernel::poolingN_f32<PoolingType::AVG, false>; break; case PoolingType::L2: - _func = &NEPoolingLayerKernel::poolingN_f32<PoolingType::L2>; + _func = (exclude_padding) ? &NEPoolingLayerKernel::poolingN_f32<PoolingType::L2, true> : &NEPoolingLayerKernel::poolingN_f32<PoolingType::L2, false>; break; case PoolingType::MAX: - _func = &NEPoolingLayerKernel::poolingN_f32<PoolingType::MAX>; + _func = &NEPoolingLayerKernel::poolingN_f32<PoolingType::MAX, false>; break; default: ARM_COMPUTE_ERROR("Unsupported pooling type!"); @@ -533,7 +541,7 @@ void NEPoolingLayerKernel::pooling2_q16(const Window &window_input, const Window input, output); } -template <PoolingType pooling_type> +template <PoolingType pooling_type, bool exclude_padding> void NEPoolingLayerKernel::pooling3_f16(const Window &window_input, const Window &window) { #ifdef ARM_COMPUTE_AARCH64_V8_2 @@ -547,8 +555,8 @@ void NEPoolingLayerKernel::pooling3_f16(const Window &window_input, const Window int pool_stride_y = 0; std::tie(pool_pad_x, pool_pad_y) = _pool_info.pad_stride_info().pad(); std::tie(pool_stride_x, pool_stride_y) = _pool_info.pad_stride_info().stride(); - const int upper_bound_w = _input->info()->dimension(0) + pool_pad_x; - const int upper_bound_h = _input->info()->dimension(1) + pool_pad_y; + const int upper_bound_w = _input->info()->dimension(0) + (exclude_padding ? 0 : pool_pad_x); + const int upper_bound_h = _input->info()->dimension(1) + (exclude_padding ? 0 : pool_pad_y); const unsigned char *const input_top_ptr = _input->ptr_to_element(Coordinates(-static_cast<int>(pool_pad_x), -static_cast<int>(pool_pad_y))); const unsigned char *const input_middle_ptr = _input->ptr_to_element(Coordinates(-static_cast<int>(pool_pad_x), -static_cast<int>(pool_pad_y) + 1)); @@ -572,7 +580,7 @@ void NEPoolingLayerKernel::pooling3_f16(const Window &window_input, const Window if(pooling_type != PoolingType::MAX) { // Calculate scale - const float scale = calculate_avg_scale(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); + const float scale = calculate_avg_scale<exclude_padding>(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); const float16x4_t scale_v = vdup_n_f16(scale); // Perform pooling const float16x4_t sum_data = vadd_f16(vadd_f16(top_data, bottom_data), middle_data); @@ -602,7 +610,7 @@ void NEPoolingLayerKernel::pooling3_f16(const Window &window_input, const Window #endif /* ARM_COMPUTE_AARCH64_V8_2 */ } -template <PoolingType pooling_type> +template <PoolingType pooling_type, bool exclude_padding> void NEPoolingLayerKernel::pooling2_f16(const Window &window_input, const Window &window) { #ifdef ARM_COMPUTE_AARCH64_V8_2 @@ -612,8 +620,8 @@ void NEPoolingLayerKernel::pooling2_f16(const Window &window_input, const Window int pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y = 0; std::tie(pool_pad_x, pool_pad_y) = _pool_info.pad_stride_info().pad(); std::tie(pool_stride_x, pool_stride_y) = _pool_info.pad_stride_info().stride(); - const int upper_bound_w = _input->info()->dimension(0) + pool_pad_x; - const int upper_bound_h = _input->info()->dimension(1) + pool_pad_y; + const int upper_bound_w = _input->info()->dimension(0) + (exclude_padding ? 0 : pool_pad_x); + const int upper_bound_h = _input->info()->dimension(1) + (exclude_padding ? 0 : pool_pad_y); const unsigned char *const input_top_ptr = _input->ptr_to_element(Coordinates(-static_cast<int>(pool_pad_x), -static_cast<int>(pool_pad_y))); const unsigned char *const input_bottom_ptr = _input->ptr_to_element(Coordinates(-static_cast<int>(pool_pad_x), -static_cast<int>(pool_pad_y) + 1)); @@ -635,7 +643,7 @@ void NEPoolingLayerKernel::pooling2_f16(const Window &window_input, const Window if(pooling_type != PoolingType::MAX) { - const float scale = calculate_avg_scale(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); + const float scale = calculate_avg_scale<exclude_padding>(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); const float16x8_t scale_v = vdupq_n_f16(scale); res = vmulq_f16(scale_v, vaddq_f16(bottom_data.val[1], vaddq_f16(bottom_data.val[0], vaddq_f16(top_data.val[0], top_data.val[1])))); } @@ -661,7 +669,7 @@ void NEPoolingLayerKernel::pooling2_f16(const Window &window_input, const Window #endif /* ARM_COMPUTE_AARCH64_V8_2 */ } -template <PoolingType pooling_type> +template <PoolingType pooling_type, bool exclude_padding> void NEPoolingLayerKernel::pooling2_f32(const Window &window_input, const Window &window) { Iterator input(_input, window_input); @@ -674,8 +682,8 @@ void NEPoolingLayerKernel::pooling2_f32(const Window &window_input, const Window int pool_stride_y = 0; std::tie(pool_pad_x, pool_pad_y) = _pool_info.pad_stride_info().pad(); std::tie(pool_stride_x, pool_stride_y) = _pool_info.pad_stride_info().stride(); - const int upper_bound_w = _input->info()->dimension(0) + pool_pad_x; - const int upper_bound_h = _input->info()->dimension(1) + pool_pad_y; + const int upper_bound_w = _input->info()->dimension(0) + (exclude_padding ? 0 : pool_pad_x); + const int upper_bound_h = _input->info()->dimension(1) + (exclude_padding ? 0 : pool_pad_y); const uint8_t *const input_top_ptr = _input->ptr_to_element(Coordinates(-static_cast<int>(pool_pad_x), -static_cast<int>(pool_pad_y))); const uint8_t *const input_bottom_ptr = _input->ptr_to_element(Coordinates(-static_cast<int>(pool_pad_x), -static_cast<int>(pool_pad_y) + 1)); @@ -697,7 +705,7 @@ void NEPoolingLayerKernel::pooling2_f32(const Window &window_input, const Window if(pooling_type != PoolingType::MAX) { // Calculate scale - float scale = calculate_avg_scale(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); + float scale = calculate_avg_scale<exclude_padding>(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); const float32x2_t scale_v = vdup_n_f32(scale); // Perform pooling @@ -868,7 +876,7 @@ void NEPoolingLayerKernel::pooling3_q16(const Window &window_input, const Window input, output); } -template <PoolingType pooling_type> +template <PoolingType pooling_type, bool exclude_padding> void NEPoolingLayerKernel::pooling3_f32(const Window &window_input, const Window &window) { Iterator input(_input, window_input); @@ -881,8 +889,8 @@ void NEPoolingLayerKernel::pooling3_f32(const Window &window_input, const Window int pool_stride_y = 0; std::tie(pool_pad_x, pool_pad_y) = _pool_info.pad_stride_info().pad(); std::tie(pool_stride_x, pool_stride_y) = _pool_info.pad_stride_info().stride(); - const int upper_bound_w = _input->info()->dimension(0) + pool_pad_x; - const int upper_bound_h = _input->info()->dimension(1) + pool_pad_y; + const int upper_bound_w = _input->info()->dimension(0) + (exclude_padding ? 0 : pool_pad_x); + const int upper_bound_h = _input->info()->dimension(1) + (exclude_padding ? 0 : pool_pad_y); const uint8_t *const input_top_ptr = _input->ptr_to_element(Coordinates(-static_cast<int>(pool_pad_x), -static_cast<int>(pool_pad_y))); const uint8_t *const input_middle_ptr = _input->ptr_to_element(Coordinates(-static_cast<int>(pool_pad_x), -static_cast<int>(pool_pad_y) + 1)); @@ -907,7 +915,7 @@ void NEPoolingLayerKernel::pooling3_f32(const Window &window_input, const Window if(pooling_type != PoolingType::MAX) { // Calculate scale - float scale = calculate_avg_scale(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); + float scale = calculate_avg_scale<exclude_padding>(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); const float32x2_t scale_v = vdup_n_f32(scale); // Perform pooling @@ -935,7 +943,7 @@ void NEPoolingLayerKernel::pooling3_f32(const Window &window_input, const Window input, output); } -template <PoolingType pooling_type> +template <PoolingType pooling_type, bool exclude_padding> void NEPoolingLayerKernel::pooling7_f32(const Window &window_input, const Window &window) { Iterator input(_input, window_input); @@ -948,8 +956,8 @@ void NEPoolingLayerKernel::pooling7_f32(const Window &window_input, const Window int pool_stride_y = 0; std::tie(pool_pad_x, pool_pad_y) = _pool_info.pad_stride_info().pad(); std::tie(pool_stride_x, pool_stride_y) = _pool_info.pad_stride_info().stride(); - const int upper_bound_w = _input->info()->dimension(0) + pool_pad_x; - const int upper_bound_h = _input->info()->dimension(1) + pool_pad_y; + const int upper_bound_w = _input->info()->dimension(0) + (exclude_padding ? 0 : pool_pad_x); + const int upper_bound_h = _input->info()->dimension(1) + (exclude_padding ? 0 : pool_pad_y); std::array<const uint8_t *, pool_size> input_ptrs{ {} }; for(int i = 0; i < pool_size; ++i) @@ -964,7 +972,7 @@ void NEPoolingLayerKernel::pooling7_f32(const Window &window_input, const Window if(pooling_type != PoolingType::MAX) { // Calculate scale - float scale = calculate_avg_scale(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); + float scale = calculate_avg_scale<exclude_padding>(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); const float32x2_t scale_v = vdup_n_f32(scale); // Perform pooling @@ -1017,7 +1025,7 @@ void NEPoolingLayerKernel::pooling7_f32(const Window &window_input, const Window input, output); } -template <PoolingType pooling_type> +template <PoolingType pooling_type, bool exclude_padding> void NEPoolingLayerKernel::poolingN_f32(const Window &window_input, const Window &window) { Iterator input(_input, window_input); @@ -1030,8 +1038,8 @@ void NEPoolingLayerKernel::poolingN_f32(const Window &window_input, const Window int pool_stride_y = 0; std::tie(pool_pad_x, pool_pad_y) = _pool_info.pad_stride_info().pad(); std::tie(pool_stride_x, pool_stride_y) = _pool_info.pad_stride_info().stride(); - const int upper_bound_w = _input->info()->dimension(0) + pool_pad_x; - const int upper_bound_h = _input->info()->dimension(1) + pool_pad_y; + const int upper_bound_w = _input->info()->dimension(0) + (exclude_padding ? 0 : pool_pad_x); + const int upper_bound_h = _input->info()->dimension(1) + (exclude_padding ? 0 : pool_pad_y); execute_window_loop(window, [&](const Coordinates & id) { @@ -1040,7 +1048,7 @@ void NEPoolingLayerKernel::poolingN_f32(const Window &window_input, const Window if(pooling_type != PoolingType::MAX) { // Calculate scale - const float scale = calculate_avg_scale(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); + const float scale = calculate_avg_scale<exclude_padding>(id, pool_size, upper_bound_w, upper_bound_h, pool_pad_x, pool_pad_y, pool_stride_x, pool_stride_y); // Perform pooling float32x4_t vres = vdupq_n_f32(0.0f); diff --git a/tests/validation/CL/PoolingLayer.cpp b/tests/validation/CL/PoolingLayer.cpp index 809c80f28c..87b86fedf2 100644 --- a/tests/validation/CL/PoolingLayer.cpp +++ b/tests/validation/CL/PoolingLayer.cpp @@ -44,13 +44,14 @@ namespace validation namespace { /** Input data set for float data types */ -const auto PoolingLayerDatasetFP = combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { 2, 3, 4, 7, 9 })), - framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) })); +const auto PoolingLayerDatasetFP = combine(combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { 2, 3, 4, 7, 9 })), + framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) })), + framework::dataset::make("ExcludePadding", { true, false })); /** Input data set for quantized data types */ -const auto PoolingLayerDatasetQS = combine(combine(framework::dataset::make("PoolingType", { PoolingType::MAX, PoolingType::AVG }), framework::dataset::make("PoolingSize", { 2, 3 })), - framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) })); - +const auto PoolingLayerDatasetQS = combine(combine(combine(framework::dataset::make("PoolingType", { PoolingType::MAX, PoolingType::AVG }), framework::dataset::make("PoolingSize", { 2, 3 })), + framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) })), + framework::dataset::make("ExcludePadding", { true, false })); constexpr AbsoluteTolerance<float> tolerance_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for float types */ constexpr AbsoluteTolerance<float> tolerance_f16(0.01f); /**< Tolerance value for comparing reference's output against implementation's output for float types */ constexpr AbsoluteTolerance<float> tolerance_qs8(3); /**< Tolerance value for comparing reference's output against implementation's output for quantized input */ diff --git a/tests/validation/CPP/PoolingLayer.cpp b/tests/validation/CPP/PoolingLayer.cpp index 85a8343d87..4f755ce2c4 100644 --- a/tests/validation/CPP/PoolingLayer.cpp +++ b/tests/validation/CPP/PoolingLayer.cpp @@ -54,12 +54,13 @@ TensorShape calculate_output_shape(TensorShape shape, PoolingLayerInfo info) template <typename T, typename std::enable_if<is_floating_point<T>::value, int>::type> SimpleTensor<T> pooling_layer(const SimpleTensor<T> &src, PoolingLayerInfo info) { - const int pool_size = info.pool_size(); - PoolingType type = info.pool_type(); - int pool_stride_x = info.pad_stride_info().stride().first; - int pool_stride_y = info.pad_stride_info().stride().second; - int pad_x = info.pad_stride_info().pad().first; - int pad_y = info.pad_stride_info().pad().second; + const int pool_size = info.pool_size(); + PoolingType type = info.pool_type(); + int pool_stride_x = info.pad_stride_info().stride().first; + int pool_stride_y = info.pad_stride_info().stride().second; + int pad_x = info.pad_stride_info().pad().first; + int pad_y = info.pad_stride_info().pad().second; + bool exclude_padding = info.exclude_padding(); const auto w_src = static_cast<int>(src.shape()[0]); const auto h_src = static_cast<int>(src.shape()[1]); @@ -122,6 +123,11 @@ SimpleTensor<T> pooling_layer(const SimpleTensor<T> &src, PoolingLayerInfo info) hstart = std::max(hstart, 0); wend = std::min(wend, w_src); hend = std::min(hend, h_src); + // Exclude padding pixels from the average + if(exclude_padding) + { + pool = (hend - hstart) * (wend - wstart); + } if(type == PoolingType::AVG) { @@ -157,12 +163,13 @@ SimpleTensor<T> pooling_layer(const SimpleTensor<T> &src, PoolingLayerInfo info) template <typename T, typename std::enable_if<std::is_integral<T>::value, int>::type> SimpleTensor<T> pooling_layer(const SimpleTensor<T> &src, PoolingLayerInfo info) { - const int pool_size = info.pool_size(); - PoolingType type = info.pool_type(); - int pool_stride_x = info.pad_stride_info().stride().first; - int pool_stride_y = info.pad_stride_info().stride().second; - int pad_x = info.pad_stride_info().pad().first; - int pad_y = info.pad_stride_info().pad().second; + const int pool_size = info.pool_size(); + PoolingType type = info.pool_type(); + int pool_stride_x = info.pad_stride_info().stride().first; + int pool_stride_y = info.pad_stride_info().stride().second; + int pad_x = info.pad_stride_info().pad().first; + int pad_y = info.pad_stride_info().pad().second; + bool exclude_padding = info.exclude_padding(); const auto w_src = static_cast<int>(src.shape()[0]); const auto h_src = static_cast<int>(src.shape()[1]); @@ -224,6 +231,11 @@ SimpleTensor<T> pooling_layer(const SimpleTensor<T> &src, PoolingLayerInfo info) hstart = std::max(hstart, 0); wend = std::min(wend, w_src); hend = std::min(hend, h_src); + // Exclude padding pixels from the average + if(exclude_padding) + { + pool = (hend - hstart) * (wend - wstart); + } using namespace fixed_point_arithmetic; diff --git a/tests/validation/NEON/PoolingLayer.cpp b/tests/validation/NEON/PoolingLayer.cpp index a721fb9d15..13384620bd 100644 --- a/tests/validation/NEON/PoolingLayer.cpp +++ b/tests/validation/NEON/PoolingLayer.cpp @@ -44,12 +44,14 @@ namespace validation namespace { /** Input data set for float data types */ -const auto PoolingLayerDatasetFP = combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { 2, 3, 7, 9 })), - framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) })); +const auto PoolingLayerDatasetFP = combine(combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { 2, 3, 7, 9 })), + framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) })), + framework::dataset::make("ExcludePadding", { true, false })); /** Input data set for quantized data types */ -const auto PoolingLayerDatasetQS = combine(combine(framework::dataset::make("PoolingType", { PoolingType::MAX, PoolingType::AVG }), framework::dataset::make("PoolingSize", { 2, 3 })), - framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) })); +const auto PoolingLayerDatasetQS = combine(combine(combine(framework::dataset::make("PoolingType", { PoolingType::MAX, PoolingType::AVG }), framework::dataset::make("PoolingSize", { 2, 3 })), + framework::dataset::make("PadStride", { PadStrideInfo(1, 1, 0, 0), PadStrideInfo(2, 1, 0, 0), PadStrideInfo(1, 2, 1, 1), PadStrideInfo(2, 2, 1, 0) })), + framework::dataset::make("ExcludePadding", { false })); constexpr AbsoluteTolerance<float> tolerance_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for float types */ #ifdef ARM_COMPUTE_AARCH64_V8_2 diff --git a/tests/validation/fixtures/PoolingLayerFixture.h b/tests/validation/fixtures/PoolingLayerFixture.h index 775c4125fc..09b9e0ef1a 100644 --- a/tests/validation/fixtures/PoolingLayerFixture.h +++ b/tests/validation/fixtures/PoolingLayerFixture.h @@ -47,10 +47,10 @@ class PoolingLayerValidationFixedPointFixture : public framework::Fixture { public: template <typename...> - void setup(TensorShape shape, PoolingType pool_type, int pool_size, PadStrideInfo pad_stride_info, DataType data_type, int fractional_bits) + void setup(TensorShape shape, PoolingType pool_type, int pool_size, PadStrideInfo pad_stride_info, bool exclude_padding, DataType data_type, int fractional_bits) { _fractional_bits = fractional_bits; - PoolingLayerInfo info(pool_type, pool_size, pad_stride_info); + PoolingLayerInfo info(pool_type, pool_size, pad_stride_info, exclude_padding); _target = compute_target(shape, info, data_type, fractional_bits); _reference = compute_reference(shape, info, data_type, fractional_bits); @@ -123,9 +123,9 @@ class PoolingLayerValidationFixture : public PoolingLayerValidationFixedPointFix { public: template <typename...> - void setup(TensorShape shape, PoolingType pool_type, int pool_size, PadStrideInfo pad_stride_info, DataType data_type) + void setup(TensorShape shape, PoolingType pool_type, int pool_size, PadStrideInfo pad_stride_info, bool exclude_padding, DataType data_type) { - PoolingLayerValidationFixedPointFixture<TensorType, AccessorType, FunctionType, T>::setup(shape, pool_type, pool_size, pad_stride_info, data_type, 0); + PoolingLayerValidationFixedPointFixture<TensorType, AccessorType, FunctionType, T>::setup(shape, pool_type, pool_size, pad_stride_info, exclude_padding, data_type, 0); } }; @@ -136,7 +136,7 @@ public: template <typename...> void setup(TensorShape shape, PoolingType pool_type, DataType data_type) { - PoolingLayerValidationFixedPointFixture<TensorType, AccessorType, FunctionType, T>::setup(shape, pool_type, shape.x(), PadStrideInfo(1, 1, 0, 0), data_type, 0); + PoolingLayerValidationFixedPointFixture<TensorType, AccessorType, FunctionType, T>::setup(shape, pool_type, shape.x(), PadStrideInfo(1, 1, 0, 0), true, data_type, 0); } }; } // namespace validation |