IntenseWebs/openvino
3
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

[GPU] AdaptiveMaxPool and AdaptiveAvgPool gpu implementations (#11556)

Browse Source 
* Add kernel for AdaptivePooling

* Add GPU primitive for AdaptivePooling

* Add single-layer tests for GPU

* Add adaptive pooling unit tests
This commit is contained in:
Tetiana Gubanova
2022-05-23 18:48:55 +03:00
committed by GitHub
parent ff6ea62ce0
commit 22ee17fda6
17 changed files with 1062 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp
View File

@@ -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);

63
src/plugins/intel_gpu/include/intel_gpu/primitives/adaptive_pooling.hpp Normal file
View File

@@ -0,0 +1,63 @@
// Copyright (C) 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#pragma once
#include "primitive.hpp"
#include <vector>
namespace cldnn {
enum class adaptive_pooling_mode : int32_t {
max,
average
};
struct adaptive_pooling : public primitive_base<adaptive_pooling> {
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<std::reference_wrapper<const primitive_id>> get_dependencies() const override {
std::vector<std::reference_wrapper<const primitive_id>> ret;
if (!indices_output.empty())
ret.push_back(indices_output);
return ret;
}
};
} // namespace cldnn

39
src/plugins/intel_gpu/src/graph/adaptive_pooling.cpp Normal file
View File

@@ -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 <string>
namespace cldnn {
primitive_type_id adaptive_pooling::type_id() {
static primitive_type_base<adaptive_pooling> 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

91
src/plugins/intel_gpu/src/graph/impls/ocl/adaptive_pooling.cpp Normal file
View File

@@ -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<adaptive_pooling> {
using parent = typed_primitive_impl_ocl<adaptive_pooling>;
using parent::parent;
std::unique_ptr<primitive_impl> clone() const override {
return make_unique<adaptive_pooling_impl>(*this);
}
protected:
kernel_arguments_data get_arguments(typed_primitive_inst<adaptive_pooling>& 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<kernel_selector::adaptive_pooling_params>(arg);
auto optional_params = get_default_optional_params<kernel_selector::adaptive_pooling_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<adaptive_pooling>::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

1
src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp
View File

@@ -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);

1
src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp
View File

@@ -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);

48
src/plugins/intel_gpu/src/graph/include/adaptive_pooling_inst.h Normal file
View File

@@ -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 <string>
namespace cldnn {
template <>
struct typed_program_node<adaptive_pooling> : public typed_program_node_base<adaptive_pooling> {
using parent = typed_program_node_base<adaptive_pooling>;
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<adaptive_pooling>;
template <>
class typed_primitive_inst<adaptive_pooling> : public typed_primitive_inst_base<adaptive_pooling> {
using parent = typed_primitive_inst_base<adaptive_pooling>;
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<adaptive_pooling>;
} // namespace cldnn

3
src/plugins/intel_gpu/src/kernel_selector/common/common_types.h
View File

@@ -80,7 +80,8 @@ enum class KernelType {
EXPERIMENTAL_DETECTRON_ROI_FEATURE_EXTRACTOR,
EXPERIMENTAL_DETECTRON_TOPK_ROIS,
CONVERT_COLOR,
RANDOM_UNIFORM
RANDOM_UNIFORM,
ADAPTIVE_POOLING
};
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

113
src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.cpp Normal file
View File

@@ -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 <algorithm>
#include <string>
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<const adaptive_pooling_params&>(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<adaptive_pooling_params>(params);
const adaptive_pooling_params& new_params = static_cast<const adaptive_pooling_params&>(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

34
src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_ref.h Normal file
View File

@@ -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

21
src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.cpp Normal file
View File

@@ -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<AdaptivePoolingRef>();
}
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

18
src/plugins/intel_gpu/src/kernel_selector/core/actual_kernels/adaptive_pooling/adaptive_pooling_kernel_selector.h Normal file
View File

@@ -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

107
src/plugins/intel_gpu/src/kernel_selector/core/cl_kernels/adaptive_pooling_gpu_ref.cl Normal file
View File

@@ -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

78
src/plugins/intel_gpu/src/plugin/ops/adaptive_pooling.cpp Normal file
View File

@@ -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<ngraph::op::v8::AdaptiveAvgPool>& 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<ngraph::op::v8::AdaptiveMaxPool>& 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

169
src/plugins/intel_gpu/tests/test_cases/adaptive_avg_pooling_gpu_test.cpp Normal file
View File

@@ -0,0 +1,169 @@
// Copyright (C) 2021-2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "test_utils.h"
#include <intel_gpu/primitives/input_layout.hpp>
#include <intel_gpu/primitives/activation.hpp>
#include <intel_gpu/primitives/adaptive_pooling.hpp>
#include <cstddef>
#include <string>
using namespace cldnn;
using namespace ::tests;
template<typename T>
struct AdaptiveAvgPoolingParams {
format inputFormat;
tensor inputTensor;
std::vector<T> inputs;
tensor outputTensor;
std::vector<T> outputs;
};
template<typename T>
struct adaptive_avg_pooling_test
: public ::testing::TestWithParam<AdaptiveAvgPoolingParams<T> > {
public:
void test() {
auto data_type = type_to_data_type<T>::value;
AdaptiveAvgPoolingParams<T> params = testing::TestWithParam<AdaptiveAvgPoolingParams<T> >::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<T> 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<typename T>
std::vector<T> getValues(const std::vector<float> &values) {
std::vector<T> result(values.begin(), values.end());
return result;
}
template<typename T>
std::vector<AdaptiveAvgPoolingParams<T>> generateAdaptiveAvgPoolingParams() {
std::vector<AdaptiveAvgPoolingParams<T>> result = {
{format::bfyx,
tensor(1, 2, 7, 3),
getValues<T>({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<T>({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<T>({-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<T>({-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<T>(
{-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<T>({-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<class T>
std::string operator()(const testing::TestParamInfo<AdaptiveAvgPoolingParams<T> > &param) {
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<float>;
using adaptive_avg_pooling_test_f16 = adaptive_avg_pooling_test<half_t>;
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<float>()),
PrintToStringParamName());
INSTANTIATE_TEST_SUITE_P(smoke_adaptive_avg_pooling_test_f16,
adaptive_avg_pooling_test_f16,
::testing::ValuesIn(generateAdaptiveAvgPoolingParams<half_t>()),
PrintToStringParamName());

198
src/plugins/intel_gpu/tests/test_cases/adaptive_max_pooling_gpu_test.cpp Normal file
View File

@@ -0,0 +1,198 @@
// Copyright (C) 2021-2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "test_utils.h"
#include <intel_gpu/primitives/input_layout.hpp>
#include <intel_gpu/primitives/activation.hpp>
#include <intel_gpu/primitives/adaptive_pooling.hpp>
#include <intel_gpu/primitives/mutable_data.hpp>
#include <cstddef>
#include <string>
using namespace cldnn;
using namespace ::tests;
template<typename T>
struct AdaptiveMaxPoolingParams {
format inputFormat;
tensor inputTensor;
std::vector<T> inputs;
tensor outputTensor;
std::vector<T> outputs;
std::vector<int32_t> output_indices;
};
template<typename T>
struct adaptive_max_pooling_test
: public ::testing::TestWithParam<AdaptiveMaxPoolingParams<T> > {
public:
void test() {
auto data_type = type_to_data_type<T>::value;
AdaptiveMaxPoolingParams<T> params = testing::TestWithParam<AdaptiveMaxPoolingParams<T> >::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<T> out_ptr(out_mem, get_test_stream());
cldnn::mem_lock<int32_t> 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<typename T>
std::vector<T> getValues(const std::vector<float> &values) {
std::vector<T> result(values.begin(), values.end());
return result;
}
template<typename T>
std::vector<AdaptiveMaxPoolingParams<T>> generateAdaptiveMaxPoolingParams() {
std::vector<AdaptiveMaxPoolingParams<T>> result = {
{format::bfyx,
tensor(2, 3, 1, 7),
getValues<T>({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<T>({4,
3,
-2,
1,
1,
0,
1,
3,
3,
3,
4,
1,
2,
-2,
-2,
3,
2,
2}),
std::vector<int32_t>{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<T>({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<T>({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<int32_t>{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<T>(
{-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<T>({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<int32_t>{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<class T>
std::string operator()(const testing::TestParamInfo<AdaptiveMaxPoolingParams<T> > &param) {
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<float>;
using adaptive_max_pooling_test_f16 = adaptive_max_pooling_test<half_t>;
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<float>()),
PrintToStringParamName());
INSTANTIATE_TEST_SUITE_P(smoke_adaptive_max_pooling_test_f16,
adaptive_max_pooling_test_f16,
::testing::ValuesIn(generateAdaptiveMaxPoolingParams<half_t>()),
PrintToStringParamName());

77
src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/adaptive_pooling.cpp Normal file
View File

@@ -0,0 +1,77 @@
// Copyright (C) 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include <string>
#include <vector>
#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<std::string> poolingModes = {"max", "avg"};
const std::vector<InferenceEngine::Precision> netPrecisions = {
InferenceEngine::Precision::FP32,
InferenceEngine::Precision::FP16,
};
const std::vector<std::vector<size_t>> inputShapes1D = {
{1, 3, 5},
{1, 1, 17},
};
const std::vector<std::vector<int>> 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<std::vector<size_t>> inputShapes2D = {
{1, 3, 4, 6},
{1, 1, 17, 5},
};
const std::vector<std::vector<int>> 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<std::vector<size_t>> inputShapes3D = {
{1, 1, 3, 3, 3},
{1, 3, 5, 7, 11},
};
const std::vector<std::vector<int>> 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