From 22ee17fda6ba44c05c537d309923557eb1febd1c Mon Sep 17 00:00:00 2001 From: Tetiana Gubanova Date: Mon, 23 May 2022 18:48:55 +0300 Subject: [PATCH] [GPU] AdaptiveMaxPool and AdaptiveAvgPool gpu implementations (#11556) * Add kernel for AdaptivePooling * Add GPU primitive for AdaptivePooling * Add single-layer tests for GPU * Add adaptive pooling unit tests --- .../intel_gpu/plugin/primitives_list.hpp | 2 + .../intel_gpu/primitives/adaptive_pooling.hpp | 63 ++++++ .../intel_gpu/src/graph/adaptive_pooling.cpp | 39 ++++ .../src/graph/impls/ocl/adaptive_pooling.cpp | 91 ++++++++ .../src/graph/impls/ocl/register.cpp | 1 + .../src/graph/impls/ocl/register.hpp | 1 + .../src/graph/include/adaptive_pooling_inst.h | 48 +++++ .../src/kernel_selector/common/common_types.h | 3 +- .../adaptive_pooling_kernel_ref.cpp | 113 ++++++++++ .../adaptive_pooling_kernel_ref.h | 34 +++ .../adaptive_pooling_kernel_selector.cpp | 21 ++ .../adaptive_pooling_kernel_selector.h | 18 ++ .../cl_kernels/adaptive_pooling_gpu_ref.cl | 107 ++++++++++ .../src/plugin/ops/adaptive_pooling.cpp | 78 +++++++ .../adaptive_avg_pooling_gpu_test.cpp | 169 +++++++++++++++ .../adaptive_max_pooling_gpu_test.cpp | 198 ++++++++++++++++++ .../single_layer_tests/adaptive_pooling.cpp | 77 +++++++ 17 files changed, 1062 insertions(+), 1 deletion(-) create mode 100644 src/plugins/intel_gpu/include/intel_gpu/primitives/adaptive_pooling.hpp create mode 100644 src/plugins/intel_gpu/src/graph/adaptive_pooling.cpp create mode 100644 src/plugins/intel_gpu/src/graph/impls/ocl/adaptive_pooling.cpp create mode 100644 src/plugins/intel_gpu/src/graph/include/adaptive_pooling_inst.h create mode 100644 src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.cpp create mode 100644 src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.h create mode 100644 src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.cpp create mode 100644 src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.h create mode 100644 src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/adaptive_pooling_gpu_ref.cl create mode 100644 src/plugins/intel_gpu/src/plugin/ops/adaptive_pooling.cpp create mode 100644 src/plugins/intel_gpu/tests/test_cases/adaptive_avg_pooling_gpu_test.cpp create mode 100644 src/plugins/intel_gpu/tests/test_cases/adaptive_max_pooling_gpu_test.cpp create mode 100644 src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/adaptive_pooling.cpp diff --git a/src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp b/src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp index ded3227964f..706fa6d37e2 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp @@ -225,6 +225,8 @@ REGISTER_FACTORY(v8, I420toRGB); REGISTER_FACTORY(v8, I420toBGR); REGISTER_FACTORY(v8, RandomUniform) REGISTER_FACTORY(v8, MaxPool); +REGISTER_FACTORY(v8, AdaptiveAvgPool); +REGISTER_FACTORY(v8, AdaptiveMaxPool); // --------------------------- Supported internal ops --------------------------- // REGISTER_FACTORY(internal, NonMaxSuppressionIEInternal); diff --git a/src/plugins/intel_gpu/include/intel_gpu/primitives/adaptive_pooling.hpp b/src/plugins/intel_gpu/include/intel_gpu/primitives/adaptive_pooling.hpp new file mode 100644 index 00000000000..4705c75619d --- /dev/null +++ b/src/plugins/intel_gpu/include/intel_gpu/primitives/adaptive_pooling.hpp @@ -0,0 +1,63 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once +#include "primitive.hpp" +#include + +namespace cldnn { + +enum class adaptive_pooling_mode : int32_t { + max, + average +}; + +struct adaptive_pooling : public primitive_base { + CLDNN_DECLARE_PRIMITIVE(adaptive_pooling) + + /// @brief Constructs AdaptiveAvgPooling primitive. + /// @param id This primitive id. + /// @param input Input primitive id. + /// @param output_size Output data size of the primitive + adaptive_pooling(const primitive_id &id, + const primitive_id &input, + tensor output_size, + const primitive_id &ext_prim_id = "") + : primitive_base(id, {input}, ext_prim_id), + mode{adaptive_pooling_mode::average}, + output_size{output_size} {} + + /// @brief Constructs AdaptiveMaxPooling primitive. + /// @param id This primitive id. + /// @param input Input primitive id. + /// @param output_shape Output shape primitive id. + /// @param output_size Output data size of the primitive + /// @param indices_output Indices output primitive id. + /// @param index_element_type Data type of indices output. + adaptive_pooling(const primitive_id &id, + const primitive_id &input, + tensor output_size, + const primitive_id &indices_output, + data_types index_element_type, + const primitive_id &ext_prim_id = "") + : primitive_base(id, {input, indices_output}, ext_prim_id), + mode{adaptive_pooling_mode::max}, + output_size{output_size}, + indices_output{indices_output}, + index_element_type{index_element_type} {} + + adaptive_pooling_mode mode; + tensor output_size; + primitive_id indices_output; + data_types index_element_type{data_types::i64}; + +protected: + std::vector> get_dependencies() const override { + std::vector> ret; + if (!indices_output.empty()) + ret.push_back(indices_output); + return ret; + } +}; +} // namespace cldnn \ No newline at end of file diff --git a/src/plugins/intel_gpu/src/graph/adaptive_pooling.cpp b/src/plugins/intel_gpu/src/graph/adaptive_pooling.cpp new file mode 100644 index 00000000000..199f16bc794 --- /dev/null +++ b/src/plugins/intel_gpu/src/graph/adaptive_pooling.cpp @@ -0,0 +1,39 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "adaptive_pooling_inst.h" +#include "primitive_type_base.h" +#include "intel_gpu/runtime/error_handler.hpp" +#include "json_object.h" +#include + +namespace cldnn { +primitive_type_id adaptive_pooling::type_id() { + static primitive_type_base instance; + return &instance; +} + +layout adaptive_pooling_inst::calc_output_layout(const adaptive_pooling_node& node) { + const auto data_layout = node.input().get_output_layout(); + const auto prim = node.get_primitive(); + return {data_layout.data_type, data_layout.format, prim->output_size}; +} + +std::string adaptive_pooling_inst::to_string(const adaptive_pooling_node& node) { + const auto prim = node.get_primitive(); + + std::stringstream primitive_description; + + json_composite info; + const auto mode = prim->mode == adaptive_pooling_mode::max ? "max" : "average"; + info.add("mode", mode); + info.add("output_size", prim->output_size); + + auto node_info = node.desc_to_json(); + node_info->add("adaptive_pooling_info", info); + node_info->dump(primitive_description); + + return primitive_description.str(); +} +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/adaptive_pooling.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/adaptive_pooling.cpp new file mode 100644 index 00000000000..84a9fde6f0f --- /dev/null +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/adaptive_pooling.cpp @@ -0,0 +1,91 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "intel_gpu/primitives/adaptive_pooling.hpp" +#include "adaptive_pooling_inst.h" +#include "primitive_base.hpp" +#include "impls/implementation_map.hpp" +#include "kernel_selector_helper.h" + +#include "adaptive_pooling/adaptive_pooling_kernel_selector.h" +#include "adaptive_pooling/adaptive_pooling_kernel_ref.h" + + +namespace cldnn { +namespace ocl { +struct adaptive_pooling_impl : public typed_primitive_impl_ocl { + using parent = typed_primitive_impl_ocl; + using parent::parent; + + std::unique_ptr clone() const override { + return make_unique(*this); + } + +protected: + kernel_arguments_data get_arguments(typed_primitive_inst& instance, int32_t) const override { + kernel_arguments_data args; + const auto num_inputs = instance.inputs_memory_count(); + for (size_t i = 0; i < num_inputs; ++i) { + args.inputs.push_back(instance.input_memory_ptr(i)); + } + + args.outputs = {instance.output_memory_ptr()}; + return args; + } + +public: + static primitive_impl* create(const adaptive_pooling_node& arg) { + auto params = get_default_params(arg); + auto optional_params = get_default_optional_params(arg.get_program()); + + const auto& primitive = arg.get_primitive(); + if (primitive->mode == adaptive_pooling_mode::average) { + params.mode = kernel_selector::PoolType::AVG; + } else { + params.mode = kernel_selector::PoolType::MAX; + + switch (primitive->index_element_type) { + case cldnn::data_types::i32: { + params.poolIndexElementType = kernel_selector::Datatype::INT32; + break; + } + case cldnn::data_types::i64: { + params.poolIndexElementType = kernel_selector::Datatype::INT64; + break; + } + default: + throw std::runtime_error{"Not supported index element type"}; + } + + params.inputs.push_back(convert_data_tensor(arg.output_indices().get_output_layout())); + } + + const auto& kernel_selector = kernel_selector::adaptive_pooling_kernel_selector::Instance(); + const auto best_kernels = kernel_selector.GetBestKernels(params, optional_params); + + CLDNN_ERROR_BOOL(arg.id(), + "best_kernels.empty()", + best_kernels.empty(), + "Cannot find a proper kernel with this arguments"); + + return new adaptive_pooling_impl(arg, best_kernels[0]); + } +}; + +namespace detail { +attach_adaptive_pooling_impl::attach_adaptive_pooling_impl() { + implementation_map::add(impl_types::ocl, adaptive_pooling_impl::create, { + std::make_tuple(data_types::f16, format::bfyx), + std::make_tuple(data_types::f16, format::bfzyx), + std::make_tuple(data_types::f32, format::bfyx), + std::make_tuple(data_types::f32, format::bfzyx), + std::make_tuple(data_types::i32, format::bfyx), + std::make_tuple(data_types::i32, format::bfzyx), + std::make_tuple(data_types::i64, format::bfyx), + std::make_tuple(data_types::i64, format::bfzyx), + }); +} +} // namespace detail +} // namespace ocl +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp index 7417cc6ada2..20c34cfcb34 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp @@ -13,6 +13,7 @@ namespace ocl { void register_implementations() { REGISTER_OCL(activation); + REGISTER_OCL(adaptive_pooling); REGISTER_OCL(arg_max_min); REGISTER_OCL(average_unpooling); REGISTER_OCL(binary_convolution); diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp b/src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp index 3b5c4b80304..2f7d4ea1b6b 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp @@ -84,6 +84,7 @@ namespace detail { } REGISTER_OCL(activation); +REGISTER_OCL(adaptive_pooling); REGISTER_OCL(arg_max_min); REGISTER_OCL(average_unpooling); REGISTER_OCL(batch_to_space); diff --git a/src/plugins/intel_gpu/src/graph/include/adaptive_pooling_inst.h b/src/plugins/intel_gpu/src/graph/include/adaptive_pooling_inst.h new file mode 100644 index 00000000000..16c20331ced --- /dev/null +++ b/src/plugins/intel_gpu/src/graph/include/adaptive_pooling_inst.h @@ -0,0 +1,48 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +/////////////////////////////////////////////////////////////////////////////////////////////////// + +#pragma once +#include "intel_gpu/primitives/adaptive_pooling.hpp" +#include "primitive_inst.h" +#include + +namespace cldnn { +template <> +struct typed_program_node : public typed_program_node_base { + using parent = typed_program_node_base; + +public: + using parent::parent; + + program_node& input() const { return get_dependency(0); } + program_node& output_shape() const { return get_dependency(1); } + program_node& output_indices() const { + CLDNN_ERROR_BOOL(id(), + "poolingMode != max", + get_primitive()->mode != adaptive_pooling_mode::max, + "output indices are relevant only for AdaptiveMaxPooling"); + return get_dependency(2); + } +}; + +using adaptive_pooling_node = typed_program_node; + +template <> +class typed_primitive_inst : public typed_primitive_inst_base { + using parent = typed_primitive_inst_base; + +public: + static layout calc_output_layout(const adaptive_pooling_node& node); + static std::string to_string(const adaptive_pooling_node& node); + + typed_primitive_inst(network& network, const adaptive_pooling_node& node) + : parent(network, node) + {} +}; + +using adaptive_pooling_inst = typed_primitive_inst; + +} // namespace cldnn diff --git a/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h b/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h index 6d39413c169..c5cd0090162 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h +++ b/src/plugins/intel_gpu/src/kernel_selector/common/common_types.h @@ -80,7 +80,8 @@ enum class KernelType { EXPERIMENTAL_DETECTRON_ROI_FEATURE_EXTRACTOR, EXPERIMENTAL_DETECTRON_TOPK_ROIS, CONVERT_COLOR, - RANDOM_UNIFORM + RANDOM_UNIFORM, + ADAPTIVE_POOLING }; //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.cpp new file mode 100644 index 00000000000..2909d0a041b --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.cpp @@ -0,0 +1,113 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "adaptive_pooling_kernel_ref.h" +#include "kernel_selector_utils.h" +#include +#include + +namespace kernel_selector { + +ParamsKey AdaptivePoolingRef::GetSupportedKey() const { + ParamsKey k; + k.EnableInputDataType(Datatype::INT32); + k.EnableInputDataType(Datatype::INT64); + k.EnableInputDataType(Datatype::F32); + k.EnableInputDataType(Datatype::F16); + k.EnableOutputDataType(Datatype::INT32); + k.EnableOutputDataType(Datatype::INT64); + k.EnableOutputDataType(Datatype::F32); + k.EnableOutputDataType(Datatype::F16); + k.EnableDifferentTypes(); + k.EnableAllInputLayout(); + k.EnableAllOutputLayout(); + k.EnableBatching(); + return k; +} + +KernelsPriority AdaptivePoolingRef::GetKernelsPriority(const Params&, const optional_params&) const { + return DONT_USE_IF_HAVE_SOMETHING_ELSE; +} + +bool AdaptivePoolingRef::Validate(const Params& p, const optional_params& o) const { + if (p.GetType() != KernelType::ADAPTIVE_POOLING + || o.GetType() != KernelType::ADAPTIVE_POOLING) { + return false; + } + + const auto& params = dynamic_cast(p); + const auto& inputs = params.inputs; + + if (!((params.mode == PoolType::MAX && inputs.size() == 2) || + (params.mode == PoolType::AVG && inputs.size() == 1))) { + return false; + } + + const auto input_dims = inputs[0].Dimentions(); + if (input_dims < 2 || input_dims > 5) { + return false; + } + + return true; +} + +namespace { +AdaptivePoolingRef::DispatchData SetDefault(const adaptive_pooling_params& params) { + AdaptivePoolingRef::DispatchData dispatch_data; + const auto& output = params.outputs[0]; + + dispatch_data.gws[0] = output.X().v; + dispatch_data.gws[1] = output.Y().v * output.Z().v; + dispatch_data.gws[2] = output.Batch().v * output.Feature().v; + + dispatch_data.lws[0] = 1; + dispatch_data.lws[1] = 1; + dispatch_data.lws[2] = 1; + + return dispatch_data; +} +} // namespace + +KernelsData AdaptivePoolingRef::GetKernelsData(const Params& params, const optional_params& options) const { + if (!Validate(params, options)) { + return {}; + } + + KernelData kd = KernelData::Default(params); + const adaptive_pooling_params& new_params = static_cast(params); + + const auto dispatchData = SetDefault(new_params); + const auto entry_point = GetEntryPoint(kernelName, new_params.layerID, params, options); + + auto cldnn_jit = MakeBaseParamsJitConstants(new_params); + + cldnn_jit.AddConstant(MakeJitConstant(toString(new_params.mode) + "_POOLING", 1)); + + if (new_params.mode == PoolType::MAX) { + cldnn_jit.Merge(MakeTypeJitConstants(new_params.poolIndexElementType, "INDICES")); + } + + const auto accumulator_type = new_params.inputs[0].GetDType(); + cldnn_jit.Merge(MakeTypeJitConstants(accumulator_type, "ACCUMULATOR")); + + const auto jit = CreateJit(kernelName, cldnn_jit, entry_point); + + auto& kernel = kd.kernels[0]; + KernelBase::CheckDispatchData(kernelName, dispatchData, params.engineInfo.maxWorkGroupSize); + kernel.params.workGroups.global = dispatchData.gws; + kernel.params.workGroups.local = dispatchData.lws; + kernel.code.kernelString = GetKernelString(kernelName, jit, entry_point, params.engineInfo); + + auto& arguments = kernel.params.arguments; + arguments.push_back({ArgumentDescriptor::Types::INPUT, 0}); // input data + arguments.push_back({ArgumentDescriptor::Types::OUTPUT, 0}); // output + if (new_params.mode == PoolType::MAX) { + arguments.push_back({ArgumentDescriptor::Types::INPUT, 1}); // indices + } + + KernelsData kernelsData; + kernelsData.push_back(std::move(kd)); + return kernelsData; +} +} // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.h b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.h new file mode 100644 index 00000000000..0ba5176da76 --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.h @@ -0,0 +1,34 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "kernel_base_opencl.h" + +namespace kernel_selector { +struct adaptive_pooling_params : public base_params { + adaptive_pooling_params() : base_params(KernelType::ADAPTIVE_POOLING) {} + + PoolType mode{PoolType::MAX}; + Datatype poolIndexElementType = Datatype::INT64; +}; + +struct adaptive_pooling_optional_params : public optional_params { + adaptive_pooling_optional_params() : optional_params(KernelType::ADAPTIVE_POOLING) {} +}; + +class AdaptivePoolingRef : public KernelBaseOpenCL { +public: + AdaptivePoolingRef() : KernelBaseOpenCL("adaptive_pooling_gpu_ref") {} + ~AdaptivePoolingRef() = default; + + using DispatchData = CommonDispatchData; + + KernelsData GetKernelsData(const Params& params, const optional_params& options) const override; + KernelsPriority GetKernelsPriority(const Params& params, const optional_params& options) const override; + ParamsKey GetSupportedKey() const override; +protected: + bool Validate(const Params& p, const optional_params& o) const override; +}; +} // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.cpp b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.cpp new file mode 100644 index 00000000000..6db85c0be30 --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.cpp @@ -0,0 +1,21 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "adaptive_pooling_kernel_selector.h" +#include "adaptive_pooling_kernel_ref.h" + +namespace kernel_selector { +adaptive_pooling_kernel_selector::adaptive_pooling_kernel_selector() { + Attach(); +} + +adaptive_pooling_kernel_selector& adaptive_pooling_kernel_selector::Instance() { + static adaptive_pooling_kernel_selector instance_; + return instance_; +} + +KernelsData adaptive_pooling_kernel_selector::GetBestKernels(const Params& params, const optional_params& options) const { + return GetNaiveBestKernel(params, options, KernelType::ADAPTIVE_POOLING); +} +} // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.h b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.h new file mode 100644 index 00000000000..fc19b5ad348 --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.h @@ -0,0 +1,18 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "kernel_selector.h" + +namespace kernel_selector { +class adaptive_pooling_kernel_selector : public kernel_selector_base { +public: + static adaptive_pooling_kernel_selector& Instance(); + + adaptive_pooling_kernel_selector(); + + KernelsData GetBestKernels(const Params& params, const optional_params& options) const override; +}; +} // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/adaptive_pooling_gpu_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/adaptive_pooling_gpu_ref.cl new file mode 100644 index 00000000000..76dc629bbb5 --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/adaptive_pooling_gpu_ref.cl @@ -0,0 +1,107 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" + +#if MAX_POOLING + #define INIT_VAL ACCUMULATOR_VAL_MIN +#elif AVG_POOLING + #define INIT_VAL ACCUMULATOR_VAL_ZERO +#else + #error +#endif + +KERNEL(adaptive_pooling_gpu)( + const __global INPUT0_TYPE* input, + __global OUTPUT_TYPE* output +#if MAX_POOLING + , __global INDICES_TYPE* indices +#endif +) +{ + const uint bf = (uint)get_global_id(2); + const uint f = bf % INPUT0_FEATURE_NUM; + const uint b = bf / INPUT0_FEATURE_NUM; + + const uint x = (uint)get_global_id(0); +#if OUTPUT_DIMS == 5 + const uint y = (uint)get_global_id(1) % OUTPUT_SIZE_Y; + const uint z = (uint)get_global_id(1) / OUTPUT_SIZE_Y; + + const uint batch_and_feature_offset = INPUT0_GET_INDEX(b, f, 0, 0, 0); +#else + const uint y = (uint)get_global_id(1); + const uint z = 0; + + const uint batch_and_feature_offset = INPUT0_GET_INDEX(b, f, 0, 0); +#endif + + ACCUMULATOR_TYPE result = INIT_VAL; + +#if MAX_POOLING + INDICES_TYPE result_idx = 0; +#elif AVG_POOLING + uint num_elements = 0; +#else +#error +#endif + +#if OUTPUT_DIMS == 5 + uint z_start = z * INPUT0_SIZE_Z / OUTPUT_SIZE_Z; + uint z_end = ceil((float)((z + 1) * INPUT0_SIZE_Z) / OUTPUT_SIZE_Z); +#endif + uint y_start = y * INPUT0_SIZE_Y / OUTPUT_SIZE_Y; + uint y_end = ceil((float)((y + 1) * INPUT0_SIZE_Y) / OUTPUT_SIZE_Y); + uint x_start = x * INPUT0_SIZE_X / OUTPUT_SIZE_X; + uint x_end = ceil((float)((x + 1) * INPUT0_SIZE_X) / OUTPUT_SIZE_X); + + +#if OUTPUT_DIMS == 5 + for (uint k = z_start; k < z_end; ++k) { + const uint z_offset = k * INPUT0_SIZE_Y * INPUT0_SIZE_X; +#else + const uint z_offset = 0; +#endif + for (uint j = y_start; j < y_end; ++j) { + uint y_offset = z_offset + j * INPUT0_SIZE_X; + + for (uint i = x_start; i < x_end; ++i) { + uint idx_within_feature = y_offset + i; + + const current_input_value = TO_ACCUMULATOR_TYPE(input[batch_and_feature_offset + idx_within_feature]); +#if MAX_POOLING + if (current_input_value > result) { + result = current_input_value; + result_idx = idx_within_feature; + } +#elif AVG_POOLING + result += TO_ACCUMULATOR_TYPE(current_input_value); + ++num_elements; +#else +#error +#endif + } + } +#if OUTPUT_DIMS == 5 + } +#endif + +#if OUTPUT_DIMS == 5 + const uint output_pos = OUTPUT_GET_INDEX(b, f, z, y, x); +#else + const uint output_pos = OUTPUT_GET_INDEX(b, f, y, x); +#endif + +#if MAX_POOLING + output[output_pos] = result; + indices[output_pos] = result_idx; +#elif AVG_POOLING + output[output_pos] = result / TO_ACCUMULATOR_TYPE(max(num_elements, (uint)1)); +#else +#error +#endif +} + +#undef INIT_VAL diff --git a/src/plugins/intel_gpu/src/plugin/ops/adaptive_pooling.cpp b/src/plugins/intel_gpu/src/plugin/ops/adaptive_pooling.cpp new file mode 100644 index 00000000000..f47553d722a --- /dev/null +++ b/src/plugins/intel_gpu/src/plugin/ops/adaptive_pooling.cpp @@ -0,0 +1,78 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "intel_gpu/plugin/program.hpp" +#include "intel_gpu/plugin/common_utils.hpp" + +#include "ngraph/op/adaptive_max_pool.hpp" + +#include "intel_gpu/primitives/mutable_data.hpp" +#include "intel_gpu/primitives/adaptive_pooling.hpp" + +namespace ov { +namespace runtime { +namespace intel_gpu { + +static void CreateAdaptiveAvgPoolOp(Program& p, const std::shared_ptr& op) { + p.ValidateInputs(op, {2}); + + const auto input_primitives = p.GetInputPrimitiveIDs(op); + const auto layer_name = layer_type_name_ID(op); + const auto op_friendly_name = op->get_friendly_name(); + + const cldnn::adaptive_pooling poolPrim{layer_name, + input_primitives[0], + tensor_from_dims(op->get_output_shape(0)), + op_friendly_name}; + p.AddPrimitive(poolPrim); + p.AddPrimitiveToProfiler(poolPrim, op); +} + +static void CreateAdaptiveMaxPoolOp(Program& p, const std::shared_ptr& op) { + p.ValidateInputs(op, {2}); + if (op->get_output_size() != 2) { + IE_THROW() << "AdaptiveMaxPool requires 2 outputs"; + } + + auto input_primitives = p.GetInputPrimitiveIDs(op); + const auto layer_type_name = layer_type_name_ID(op); + const auto layer_name = layer_type_name + ".0"; + const auto op_friendly_name = op->get_friendly_name(); + + const auto indices_precision = op->get_output_element_type(1); + const auto indices_shape = op->get_output_shape(1); + const cldnn::layout indices_layout{DataTypeFromPrecision(indices_precision), + DefaultFormatForDims(indices_shape.size()), + tensor_from_dims(indices_shape)}; + const auto indices_memory = p.GetEngine().allocate_memory(indices_layout); + + const cldnn::primitive_id indices_id_w = layer_type_name + "_md_write"; + const cldnn::mutable_data indices_mutable_prim_w{indices_id_w, indices_memory, op_friendly_name}; + p.primitiveIDs[indices_id_w] = indices_id_w; + p.AddPrimitive(indices_mutable_prim_w); + + input_primitives.push_back(indices_id_w); + + const cldnn::adaptive_pooling poolPrim{layer_name, + input_primitives[0], + tensor_from_dims(op->get_output_shape(0)), + input_primitives.back(), + DataTypeFromPrecision(op->get_index_element_type()), + op_friendly_name}; + p.AddPrimitive(poolPrim); + + const cldnn::primitive_id indices_id_r = layer_type_name + ".1"; + const cldnn::mutable_data indices_mutable_prim_r{indices_id_r, {layer_name}, indices_memory, op_friendly_name}; + p.primitiveIDs[indices_id_r] = indices_id_r; + p.AddPrimitive(indices_mutable_prim_r); + + p.AddPrimitiveToProfiler(poolPrim, op); +} + +REGISTER_FACTORY_IMPL(v8, AdaptiveAvgPool); +REGISTER_FACTORY_IMPL(v8, AdaptiveMaxPool); + +} // namespace intel_gpu +} // namespace runtime +} // namespace ov diff --git a/src/plugins/intel_gpu/tests/test_cases/adaptive_avg_pooling_gpu_test.cpp b/src/plugins/intel_gpu/tests/test_cases/adaptive_avg_pooling_gpu_test.cpp new file mode 100644 index 00000000000..954f8aeb570 --- /dev/null +++ b/src/plugins/intel_gpu/tests/test_cases/adaptive_avg_pooling_gpu_test.cpp @@ -0,0 +1,169 @@ +// Copyright (C) 2021-2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "test_utils.h" +#include +#include +#include + +#include +#include + +using namespace cldnn; +using namespace ::tests; + +template +struct AdaptiveAvgPoolingParams { + format inputFormat; + tensor inputTensor; + std::vector inputs; + tensor outputTensor; + std::vector outputs; +}; + +template +struct adaptive_avg_pooling_test + : public ::testing::TestWithParam > { +public: + void test() { + auto data_type = type_to_data_type::value; + AdaptiveAvgPoolingParams params = testing::TestWithParam >::GetParam(); + auto &engine = get_test_engine(); + + auto input = engine.allocate_memory({data_type, params.inputFormat, params.inputTensor}); + + set_values(input, params.inputs); + + const std::string input_id = "adaptive_avg_input_id"; + const std::string adaptive_avg_pooling_id = "adaptive_avg_pooling_id"; + topology topology; + topology.add(input_layout(input_id, input->get_layout())); + + topology.add(adaptive_pooling(adaptive_avg_pooling_id, input_id, params.outputTensor)); + + network network(engine, topology); + + network.set_input_data(input_id, input); + + auto result = network.execute(); + + auto out_mem = result.at(adaptive_avg_pooling_id).get_memory(); + cldnn::mem_lock out_ptr(out_mem, get_test_stream()); + + ASSERT_EQ(params.outputTensor.count(), out_ptr.size()); + for (size_t i = 0; i < params.outputs.size(); ++i) { + EXPECT_NEAR(params.outputs[i], out_ptr[i], 0.005) << "at i = " << i; + } + } +}; + +template +std::vector getValues(const std::vector &values) { + std::vector result(values.begin(), values.end()); + return result; +} + +template +std::vector> generateAdaptiveAvgPoolingParams() { + std::vector> result = { + {format::bfyx, + tensor(1, 2, 7, 3), + getValues({0, 4, 1, 3, -2, -5, -2, -2, 1, -3, 1, -3, -4, 0, -2, 1, -1, -2, 3, -1, -3, + -1, -2, 3, 4, -3, -4, 1, 2, 0, -4, -5, -2, -2, -3, 2, 3, 1, -5, 2, -4, -2}), + tensor(1, 2, 3, 3), + getValues({1.66666663, + 0.66666669, + -3., + -1.33333337, + -1.66666663, + -2.33333325, + -0.66666669, + 0., + -0.33333334, + + 0., + 1.33333337, + -2., + -0.66666669, + -3.66666675, + -2.33333325, + 2., + -0.66666669, + -1.33333337}) + }, + {format::bfyx, + tensor(1, 3, 10, 7), + getValues({-2, -3, -4, 3, -5, 4, 0, -4, -2, -4, -5, 0, -3, 0, -2, 0, 0, -5, -4, -1, 3, -1, 0, -1, + 0, -2, 0, 4, 1, 4, 0, -1, -4, 2, -2, -5, -1, -1, -2, 1, 2, -2, -1, 2, 0, -1, 0, -5, + 4, 4, 3, 0, -4, -4, -4, -2, 0, 1, -2, -1, 4, -2, -4, 1, -1, -3, -4, -1, 1, -4, + + -2, -4, -5, 0, -4, 3, 4, -5, -4, -2, 0, 2, -4, -3, 3, -1, 1, -4, -5, 4, 2, -5, 2, -3, + 0, 4, 3, 3, 1, 2, -1, -4, 1, -3, -3, -2, 3, 4, -2, -5, 1, 4, 4, -2, 2, 1, -5, -2, + -5, 1, 1, -2, -3, -3, -1, -5, 1, -3, -5, -3, -4, -1, 4, -3, 4, -1, 4, 3, 1, 4, + + -2, -4, -4, 4, -3, 4, 2, -3, -2, 4, -3, 0, 1, -4, 4, 4, 0, 3, -1, 3, 3, -5, 0, 3, + -3, 1, -2, 4, -5, -5, 1, 0, -1, 0, -3, -2, 0, -3, 3, -2, -2, 0, -3, 4, -1, 2, -2, 2, + -3, -1, -4, -2, 0, 2, 0, 2, 0, -3, 4, 3, -5, -3, -5, 1, -5, -3, -5, 4, -3, 3}), + tensor(1, 3, 3, 3), + getValues({-1.08333337, -0.25000000, -0.91666669, -0.08333334, -0.66666669, + 0.75000000, -0.41666666, -1.33333337, -0.58333331, + + -1.66666663, 0.58333331, -0.16666667, -0.33333334, -0.41666666, + -0.16666667, -0.33333334, -0.66666669, -0.75000000, + + -0.91666669, 0.83333331, -0.16666667, 0., -0.25000000, + -1.16666663, -1.41666663, -0.41666666, -0.08333334}) + }, + {format::bfzyx, + tensor(2, 2, 3, 3, 3), + getValues( + {-5, 1, -3, -4, 4, -4, 3, -3, -1, 0, 0, -2, -4, 2, 0, -4, -5, -2, -4, -4, 0, -2, 3, -3, 4, + -1, -4, + -1, -1, -5, 4, -1, -2, -3, 0, 4, -1, -5, -4, 1, 1, 4, -5, -5, -5, 4, -3, -3, -3, 4, 0, -3, + -5, 1, + 4, 2, 1, -5, -5, 1, 0, -4, -1, 2, -4, -2, 4, 3, 1, -3, -3, -2, -4, -3, -3, 3, -1, 1, 2, 2, + -4, + -5, -4, 1, 3, -4, -1, 2, 4, -5, 0, 1, -2, 0, 0, -2, 3, -2, -5, -3, -5, -2, -1, 3, -2, 4, 3, + -3}), + tensor(2, 2, 2, 2, 2), + getValues({-0.750, -0.250, -1.375, -1.125, -1.125, -0.500, -0.875, -1.250, + -0.375, -1.625, -1., -0.500, -0.250, -0.750, -1.875, -0.625, + 0.125, -0.375, -1.625, -1.250, 0., -1., 0.875, -0.375, + -1.125, -1.375, 0.750, -1.875, -0.625, -1.125, 1.250, -1.}), + }, + }; + return result; +} + +struct PrintToStringParamName { + template + std::string operator()(const testing::TestParamInfo > ¶m) { + std::stringstream buf; + buf << " input tensor " << param.param.inputTensor.to_string() + << " output tensor " << param.param.outputTensor.to_string(); + return buf.str(); + } +}; + +using adaptive_avg_pooling_test_f32 = adaptive_avg_pooling_test; +using adaptive_avg_pooling_test_f16 = adaptive_avg_pooling_test; + +TEST_P(adaptive_avg_pooling_test_f32, adaptive_avg_pooling_test_f32) { + ASSERT_NO_FATAL_FAILURE(test()); +} + +TEST_P(adaptive_avg_pooling_test_f16, adaptive_avg_pooling_test_f16) { + ASSERT_NO_FATAL_FAILURE(test()); +} + + +INSTANTIATE_TEST_SUITE_P(smoke_adaptive_avg_pooling_test_f32, + adaptive_avg_pooling_test_f32, + ::testing::ValuesIn(generateAdaptiveAvgPoolingParams()), + PrintToStringParamName()); + +INSTANTIATE_TEST_SUITE_P(smoke_adaptive_avg_pooling_test_f16, + adaptive_avg_pooling_test_f16, + ::testing::ValuesIn(generateAdaptiveAvgPoolingParams()), + PrintToStringParamName()); diff --git a/src/plugins/intel_gpu/tests/test_cases/adaptive_max_pooling_gpu_test.cpp b/src/plugins/intel_gpu/tests/test_cases/adaptive_max_pooling_gpu_test.cpp new file mode 100644 index 00000000000..fc19264bc2d --- /dev/null +++ b/src/plugins/intel_gpu/tests/test_cases/adaptive_max_pooling_gpu_test.cpp @@ -0,0 +1,198 @@ +// Copyright (C) 2021-2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "test_utils.h" +#include +#include +#include +#include + +#include +#include + +using namespace cldnn; +using namespace ::tests; + +template +struct AdaptiveMaxPoolingParams { + format inputFormat; + tensor inputTensor; + std::vector inputs; + tensor outputTensor; + std::vector outputs; + std::vector output_indices; +}; + +template +struct adaptive_max_pooling_test + : public ::testing::TestWithParam > { +public: + void test() { + auto data_type = type_to_data_type::value; + AdaptiveMaxPoolingParams params = testing::TestWithParam >::GetParam(); + auto &engine = get_test_engine(); + + auto input = engine.allocate_memory({data_type, params.inputFormat, params.inputTensor}); + auto indices_output = engine.allocate_memory({data_types::i32, params.inputFormat, params.outputTensor}); + + set_values(input, params.inputs); + + const std::string input_id = "adaptive_max_input_id"; + const std::string adaptive_max_pooling_id = "adaptive_max_pooling_id"; + const std::string output_indices_id = "output_indices_id"; + topology topology; + topology.add(input_layout(input_id, input->get_layout())); + topology.add(mutable_data(output_indices_id, indices_output)); + + topology.add(adaptive_pooling(adaptive_max_pooling_id, input_id, params.outputTensor, output_indices_id, + data_types::i32)); + + network network(engine, topology); + + network.set_input_data(input_id, input); + + auto result = network.execute(); + + auto out_mem = result.at(adaptive_max_pooling_id).get_memory(); + cldnn::mem_lock out_ptr(out_mem, get_test_stream()); + cldnn::mem_lock out_indices(indices_output, get_test_stream()); + + ASSERT_EQ(params.outputTensor.count(), out_ptr.size()); + ASSERT_EQ(params.outputTensor.count(), out_indices.size()); + for (size_t i = 0; i < params.outputs.size(); ++i) { + EXPECT_NEAR(params.outputs[i], out_ptr[i], 0.005) << "at i = " << i; + EXPECT_EQ(params.output_indices[i], out_indices[i]); + } + } +}; + +template +std::vector getValues(const std::vector &values) { + std::vector result(values.begin(), values.end()); + return result; +} + +template +std::vector> generateAdaptiveMaxPoolingParams() { + std::vector> result = { + {format::bfyx, + tensor(2, 3, 1, 7), + getValues({0, 4, 1, 3, -2, -5, -2, -2, 1, -3, 1, -3, -4, 0, -2, 1, -1, -2, 3, -1, -3, + -1, -2, 3, 4, -3, -4, 1, 2, 0, -4, -5, -2, -2, -3, 2, 3, 1, -5, 2, -4, -2}), + tensor(2, 3, 1, 3), + getValues({4, + 3, + -2, + 1, + 1, + 0, + 1, + 3, + 3, + 3, + 4, + 1, + 2, + -2, + -2, + 3, + 2, + 2}), + std::vector{1, + 3, + 4, + 1, + 3, + 6, + 1, + 4, + 4, + 2, + 3, + 6, + 0, + 4, + 4, + 1, + 4, + 4} + }, + {format::bfyx, + tensor(1, 3, 10, 7), + getValues({0, -2, -5, -5, 2, 3, 2, -3, 1, -2, -4, -1, -1, -1, 2, -4, 3, -5, -1, -1, 1, 2, 4, -2, + -3, -2, 0, -5, 2, -4, -1, -4, 4, 2, 1, -2, 2, -3, 0, 1, -3, 3, -1, 4, 0, 2, 0, 3, + 4, -4, 1, 4, -1, -5, -2, 4, -3, 3, 2, 1, 0, 4, 2, -5, 2, -5, -2, -1, 4, 2, + + 0, 4, -2, 0, -5, -3, 4, -4, -2, -2, 2, 1, 4, 3, 2, -5, -4, -4, 0, 1, 4, -4, -3, 3, + 3, 4, -2, -3, -4, -2, 0, 1, -1, 3, -2, 2, 0, -3, -1, -1, 0, 0, 2, 2, -2, 1, -3, 1, + 2, 4, 3, -5, -4, 1, -4, 2, 0, -2, -5, 2, -3, -2, -3, -4, 2, -2, -4, 2, -4, -3, + + 1, -5, -1, -5, 2, 1, 3, 4, 3, 0, -5, 4, -3, -4, -1, 2, -4, 2, 0, -5, -3, 0, 2, -3, + -5, 3, -2, -1, -5, -4, -5, 0, -5, -1, -3, 3, 3, -4, -3, -4, -5, 4, -1, 1, -1, -4, 1, + -3, + -4, -1, -2, -3, -5, 2, 2, -5, 1, 1, -5, -4, 0, 2, 4, 2, 0, 2, 4, 0, -5, 2}), + tensor(1, 3, 3, 3), + getValues({4, 3, 3, 4, 4, 4, 4, 4, 4, + 4, 4, 4, 4, 4, 4, 3, 2, 4, + 4, 3, 4, 4, 3, 3, 4, 4, 4}), + std::vector{22, 5, 16, 22, 43, 48, 43, 43, 48, + 1, 6, 6, 20, 25, 49, 50, 43, 49, + 11, 6, 7, 41, 25, 36, 41, 66, 66} + }, + {format::bfzyx, + tensor(2, 2, 3, 3, 3), + getValues( + {-5, 1, -3, -4, 4, -4, 3, -3, -1, 0, 0, -2, -4, 2, 0, -4, -5, -2, -4, -4, 0, -2, 3, -3, 4, + -1, -4, + -1, -1, -5, 4, -1, -2, -3, 0, 4, -1, -5, -4, 1, 1, 4, -5, -5, -5, 4, -3, -3, -3, 4, 0, -3, + -5, 1, + 4, 2, 1, -5, -5, 1, 0, -4, -1, 2, -4, -2, 4, 3, 1, -3, -3, -2, -4, -3, -3, 3, -1, 1, 2, 2, + -4, + -5, -4, 1, 3, -4, -1, 2, 4, -5, 0, 1, -2, 0, 0, -2, 3, -2, -5, -3, -5, -2, -1, 3, -2, 4, 3, + -3}), + tensor(2, 2, 2, 2, 2), + getValues({4, 4, 4, 4, 3, 3, 4, 3, + 4, 4, 4, 4, 4, 4, 4, 4, + 4, 3, 4, 3, 4, 3, 4, 3, + 3, 1, 4, 4, 3, 3, 4, 3}), + std::vector{4, 4, 4, 4, 22, 22, 24, 22, + 3, 14, 3, 8, 18, 14, 22, 14, + 0, 13, 12, 13, 12, 13, 12, 13, + 3, 2, 7, 7, 22, 22, 24, 22} + }, + }; + return result; +} + +struct PrintToStringParamName { + template + std::string operator()(const testing::TestParamInfo > ¶m) { + std::stringstream buf; + buf << " input tensor " << param.param.inputTensor.to_string() + << " output tensor " << param.param.outputTensor.to_string(); + return buf.str(); + } +}; + +using adaptive_max_pooling_test_f32 = adaptive_max_pooling_test; +using adaptive_max_pooling_test_f16 = adaptive_max_pooling_test; + +TEST_P(adaptive_max_pooling_test_f32, adaptive_max_pooling_test_f32) { + ASSERT_NO_FATAL_FAILURE(test()); +} + +TEST_P(adaptive_max_pooling_test_f16, adaptive_max_pooling_test_f16) { + ASSERT_NO_FATAL_FAILURE(test()); +} + + +INSTANTIATE_TEST_SUITE_P(smoke_adaptive_max_pooling_test_f32, + adaptive_max_pooling_test_f32, + ::testing::ValuesIn(generateAdaptiveMaxPoolingParams()), + PrintToStringParamName()); + +INSTANTIATE_TEST_SUITE_P(smoke_adaptive_max_pooling_test_f16, + adaptive_max_pooling_test_f16, + ::testing::ValuesIn(generateAdaptiveMaxPoolingParams()), + PrintToStringParamName()); diff --git a/src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/adaptive_pooling.cpp b/src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/adaptive_pooling.cpp new file mode 100644 index 00000000000..f19829c7be4 --- /dev/null +++ b/src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/adaptive_pooling.cpp @@ -0,0 +1,77 @@ +// Copyright (C) 2022 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include +#include + +#include "single_layer_tests/adaptive_pooling.hpp" +#include "common_test_utils/test_constants.hpp" + +using namespace ngraph::helpers; +using namespace LayerTestsDefinitions; +using namespace ngraph::element; + +namespace { +const std::vector poolingModes = {"max", "avg"}; + +const std::vector netPrecisions = { + InferenceEngine::Precision::FP32, + InferenceEngine::Precision::FP16, +}; + +const std::vector> inputShapes1D = { + {1, 3, 5}, + {1, 1, 17}, +}; +const std::vector> outputShapes1D = { + {2}, + {5}, +}; + +INSTANTIATE_TEST_SUITE_P(smoke_AdaptivePooling1D, AdaPoolLayerTest, + ::testing::Combine( + ::testing::ValuesIn(inputShapes1D), + ::testing::ValuesIn(outputShapes1D), + ::testing::ValuesIn(poolingModes), + ::testing::ValuesIn(netPrecisions), + ::testing::Values(CommonTestUtils::DEVICE_GPU)), + AdaPoolLayerTest::getTestCaseName); + +const std::vector> inputShapes2D = { + {1, 3, 4, 6}, + {1, 1, 17, 5}, +}; +const std::vector> outputShapes2D = { + {2, 4}, + {4, 5}, +}; + +INSTANTIATE_TEST_SUITE_P(smoke_AdaptivePooling2D, AdaPoolLayerTest, + ::testing::Combine( + ::testing::ValuesIn(inputShapes2D), + ::testing::ValuesIn(outputShapes2D), + ::testing::ValuesIn(poolingModes), + ::testing::ValuesIn(netPrecisions), + ::testing::Values(CommonTestUtils::DEVICE_GPU)), + AdaPoolLayerTest::getTestCaseName); + +const std::vector> inputShapes3D = { + {1, 1, 3, 3, 3}, + {1, 3, 5, 7, 11}, +}; +const std::vector> outputShapes3D = { + {2, 2, 2}, + {4, 5, 3}, +}; + +INSTANTIATE_TEST_SUITE_P(smoke_AdaptivePooling3D, AdaPoolLayerTest, + ::testing::Combine( + ::testing::ValuesIn(inputShapes3D), + ::testing::ValuesIn(outputShapes3D), + ::testing::ValuesIn(poolingModes), + ::testing::ValuesIn(netPrecisions), + ::testing::Values(CommonTestUtils::DEVICE_GPU)), + AdaPoolLayerTest::getTestCaseName); + +} // namespace