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

[GPU] Support dynamic BatchToSpace and SpaceToBatch (#18630)

Browse Source
This commit is contained in:
Kelvin Choi
2023-08-09 14:56:55 +09:00
committed by GitHub
parent 5db7579bea
commit e60da26894
20 changed files with 1054 additions and 73 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
src/plugins/intel_gpu/include/intel_gpu/primitives/batch_to_space.hpp
View File

@@ -61,18 +61,32 @@ struct batch_to_space : public primitive_base<batch_to_space> {
block_shape(block_shape),
crops_begin(crops_begin),
crops_end(crops_end),
out_size(out_size) {}
out_size(out_size),
shape_constant(1) {}
batch_to_space(const primitive_id& id,
const std::vector<input_info>& inputs,
const tensor& out_size,
const padding& output_padding = padding())
: primitive_base(id, inputs, {output_padding}),
block_shape(tensor()),
crops_begin(tensor()),
crops_end(tensor()),
out_size(out_size),
shape_constant(0) {}
tensor block_shape;
tensor crops_begin;
tensor crops_end;
tensor out_size;
int64_t shape_constant;
size_t hash() const override {
size_t seed = primitive::hash();
seed = hash_combine(seed, block_shape.hash());
seed = hash_combine(seed, crops_begin.hash());
seed = hash_combine(seed, crops_end.hash());
seed = hash_combine(seed, shape_constant);
return seed;
}
@@ -93,6 +107,7 @@ struct batch_to_space : public primitive_base<batch_to_space> {
ob << crops_begin;
ob << crops_end;
ob << out_size;
ob << shape_constant;
}
void load(BinaryInputBuffer& ib) override {
@@ -101,6 +116,7 @@ struct batch_to_space : public primitive_base<batch_to_space> {
ib >> crops_begin;
ib >> crops_end;
ib >> out_size;
ib >> shape_constant;
}
};
} // namespace cldnn

21
src/plugins/intel_gpu/include/intel_gpu/primitives/space_to_batch.hpp
View File

@@ -58,18 +58,32 @@ struct space_to_batch : public primitive_base<space_to_batch> {
block_shape(block_shape),
pads_begin(pads_begin),
pads_end(pads_end),
out_size(out_size) {}
out_size(out_size),
shape_constant(1) {}
space_to_batch(const primitive_id& id,
const std::vector<input_info>& inputs,
const tensor& out_size,
const padding& output_padding = padding())
: primitive_base(id, inputs, {output_padding}),
block_shape(tensor()),
pads_begin(tensor()),
pads_end(tensor()),
out_size(out_size),
shape_constant(0) {}
tensor block_shape;
tensor pads_begin;
tensor pads_end;
tensor out_size;
int64_t shape_constant;
size_t hash() const override {
size_t seed = primitive::hash();
seed = hash_combine(seed, block_shape.hash());
seed = hash_combine(seed, pads_begin.hash());
seed = hash_combine(seed, pads_end.hash());
seed = hash_combine(seed, shape_constant);
return seed;
}
@@ -81,7 +95,8 @@ struct space_to_batch : public primitive_base<space_to_batch> {
return block_shape == rhs_casted.block_shape &&
pads_begin == rhs_casted.pads_begin &&
pads_end == rhs_casted.pads_end;
pads_end == rhs_casted.pads_end &&
shape_constant == rhs_casted.shape_constant;
}
void save(BinaryOutputBuffer& ob) const override {
@@ -90,6 +105,7 @@ struct space_to_batch : public primitive_base<space_to_batch> {
ob << pads_begin;
ob << pads_end;
ob << out_size;
ob << shape_constant;
}
void load(BinaryInputBuffer& ib) override {
@@ -98,6 +114,7 @@ struct space_to_batch : public primitive_base<space_to_batch> {
ib >> pads_begin;
ib >> pads_end;
ib >> out_size;
ib >> shape_constant;
}
};
} // namespace cldnn

90
src/plugins/intel_gpu/src/graph/batch_to_space.cpp
View File

@@ -10,6 +10,8 @@
#include <string>
#include <vector>
#include "batch_to_space_shape_inference.hpp"
namespace cldnn {
GPU_DEFINE_PRIMITIVE_TYPE_ID(batch_to_space)
@@ -66,6 +68,94 @@ layout batch_to_space_inst::calc_output_layout(batch_to_space_node const& node,
return layout{output_type, input_format, desc->out_size};
}
static std::vector<int32_t> tensor_to_vec(const tensor& t, const format f) {
std::vector<int32_t> vec(cldnn::format::dimension(f));
for (size_t i = 0; i < vec.size(); ++i) {
vec[i] = t.sizes()[i];
}
std::reverse(vec.begin() + 2, vec.end());
return vec;
}
template<typename ShapeType>
std::vector<layout> batch_to_space_inst::calc_output_layouts(batch_to_space_node const& /*node*/, const kernel_impl_params& impl_param) {
auto desc = impl_param.typed_desc<batch_to_space>();
auto input0_layout = impl_param.get_input_layout(0);
auto input0_shape = input0_layout.get<ShapeType>();
auto input0_size = input0_shape.size();
auto input0_format = input0_layout.format;
auto& constant_mem = impl_param.memory_deps;
auto block_data = desc->block_shape;
auto begin_data = desc->crops_begin;
auto end_data = desc->crops_end;
auto output_type = desc->output_data_types[0].value_or(input0_layout.data_type);
if (impl_param.has_fused_primitives())
output_type = impl_param.get_fused_output_layout().data_type;
if (desc->shape_constant == 0 && (!constant_mem.count(1) || !constant_mem.count(2) || !constant_mem.count(3))) {
auto out_shape = ov::PartialShape::dynamic(input0_size);
return { layout{out_shape, output_type, input0_format } };
}
ShapeType block_shape = desc->shape_constant == 0 ? impl_param.get_input_layout(1).get<ShapeType>() : ov::Shape{ input0_size };
ShapeType begin_shape = desc->shape_constant == 0 ? impl_param.get_input_layout(2).get<ShapeType>() : ov::Shape{ input0_size };
ShapeType end_shape = desc->shape_constant == 0 ? impl_param.get_input_layout(3).get<ShapeType>() : ov::Shape{ input0_size };
ov::op::v1::BatchToSpace op;
std::vector<ShapeType> output_shapes = {ShapeType{}};
std::vector<ShapeType> input_shapes = {
input0_shape,
block_shape,
begin_shape,
end_shape
};
std::map<size_t, ngraph::HostTensorPtr> const_data;
if (desc->shape_constant) {
auto block_sizes = tensor_to_vec(block_data, input0_format);
auto begin_sizes = tensor_to_vec(begin_data, input0_format);
auto end_sizes = tensor_to_vec(end_data, input0_format);
auto block_values = static_cast<void*>(block_sizes.data());
auto begin_values = static_cast<void*>(begin_sizes.data());
auto end_values = static_cast<void*>(end_sizes.data());
auto block_tensor = make_host_tensor({ block_shape, data_types::i32, input0_format }, block_values);
auto begin_tensor = make_host_tensor({ begin_shape, data_types::i32, input0_format }, begin_values);
auto end_tensor = make_host_tensor({ end_shape, data_types::i32, input0_format }, end_values);
const_data.emplace(1, block_tensor);
const_data.emplace(2, begin_tensor);
const_data.emplace(3, end_tensor);
output_shapes = ov::op::v1::shape_infer(&op, input_shapes, ov::make_tensor_accessor(const_data));
} else {
auto block_mem = constant_mem.at(1);
auto begin_mem = constant_mem.at(2);
auto end_mem = constant_mem.at(3);
cldnn::mem_lock<uint8_t, mem_lock_type::read> lock1(block_mem, impl_param.get_stream());
cldnn::mem_lock<uint8_t, mem_lock_type::read> lock2(begin_mem, impl_param.get_stream());
cldnn::mem_lock<uint8_t, mem_lock_type::read> lock3(end_mem, impl_param.get_stream());
auto block_tensor = make_host_tensor(block_mem->get_layout(), lock1.data());
auto begin_tensor = make_host_tensor(begin_mem->get_layout(), lock2.data());
auto end_tensor = make_host_tensor(end_mem->get_layout(), lock3.data());
const_data.emplace(1, block_tensor);
const_data.emplace(2, begin_tensor);
const_data.emplace(3, end_tensor);
output_shapes = ov::op::v1::shape_infer(&op, input_shapes, ov::make_tensor_accessor(const_data));
}
return { layout{output_shapes[0], output_type, input0_format} };
}
template std::vector<layout> batch_to_space_inst::calc_output_layouts<ov::PartialShape>(batch_to_space_node const& node, const kernel_impl_params& impl_param);
std::string batch_to_space_inst::to_string(batch_to_space_node const& node) {
auto desc = node.get_primitive();
auto node_info = node.desc_to_json();

31
src/plugins/intel_gpu/src/graph/impls/ocl/batch_to_space.cpp
View File

@@ -26,9 +26,34 @@ struct batch_to_space_impl : typed_primitive_impl_ocl<batch_to_space> {
auto params = get_default_params<kernel_selector::batch_to_space_params>(impl_param);
auto optional_params = get_default_optional_params<kernel_selector::batch_to_space_optional_params>(impl_param.get_program());
params.block_shape = convert_dim_vector(primitive->block_shape);
params.crops_begin = convert_dim_vector(primitive->crops_begin);
params.crops_end = convert_dim_vector(primitive->crops_end);
if (primitive->shape_constant) {
params.block_type = kernel_selector::base_params::ArgType::Constant;
params.block_shape = convert_dim_vector(primitive->block_shape);
params.begin_type = kernel_selector::base_params::ArgType::Constant;
params.crops_begin = convert_dim_vector(primitive->crops_begin);
params.end_type = kernel_selector::base_params::ArgType::Constant;
params.crops_end = convert_dim_vector(primitive->crops_end);
} else {
params.block_input_index = 1;
params.block_type = kernel_selector::base_params::ArgType::Input;
auto block_layout = impl_param.get_input_layout(params.block_input_index);
params.inputs.push_back(convert_data_tensor(block_layout));
params.block_dims = block_layout.count();
params.begin_input_index = 2;
params.begin_type = kernel_selector::base_params::ArgType::Input;
auto begin_layout = impl_param.get_input_layout(params.begin_input_index);
params.inputs.push_back(convert_data_tensor(begin_layout));
params.begin_dims = begin_layout.count();
params.end_input_index = 3;
params.end_type = kernel_selector::base_params::ArgType::Input;
auto end_layout = impl_param.get_input_layout(params.end_input_index);
params.inputs.push_back(convert_data_tensor(end_layout));
params.end_dims = end_layout.count();
}
return {params, optional_params};
}

38
src/plugins/intel_gpu/src/graph/impls/ocl/space_to_batch.cpp
View File

@@ -22,17 +22,47 @@ struct space_to_batch_impl : typed_primitive_impl_ocl<space_to_batch> {
return make_unique<space_to_batch_impl>(*this);
}
static kernel_params_t get_kernel_params(const kernel_impl_params& impl_param) {
static kernel_params_t get_kernel_params(const kernel_impl_params& impl_param, bool is_shape_agnostic = false) {
const auto& primitive = impl_param.typed_desc<space_to_batch>();
auto params = get_default_params<kernel_selector::space_to_batch_params>(impl_param);
auto optional_params = get_default_optional_params<kernel_selector::space_to_batch_optional_params>(impl_param.get_program());
params.block_shape = convert_dim_vector(primitive->block_shape);
params.pads_begin = convert_dim_vector(primitive->pads_begin);
params.pads_end = convert_dim_vector(primitive->pads_end);
if (primitive->shape_constant) {
params.block_type = kernel_selector::base_params::ArgType::Constant;
params.block_shape = convert_dim_vector(primitive->block_shape);
params.begin_type = kernel_selector::base_params::ArgType::Constant;
params.pads_begin = convert_dim_vector(primitive->pads_begin);
params.end_type = kernel_selector::base_params::ArgType::Constant;
params.pads_end = convert_dim_vector(primitive->pads_end);
} else {
params.block_input_index = 1;
params.block_type = kernel_selector::base_params::ArgType::Input;
auto block_layout = impl_param.get_input_layout(params.block_input_index);
params.inputs.push_back(convert_data_tensor(block_layout));
params.block_dims = block_layout.count();
params.begin_input_index = 2;
params.begin_type = kernel_selector::base_params::ArgType::Input;
auto begin_layout = impl_param.get_input_layout(params.begin_input_index);
params.inputs.push_back(convert_data_tensor(begin_layout));
params.begin_dims = begin_layout.count();
params.end_input_index = 3;
params.end_type = kernel_selector::base_params::ArgType::Input;
auto end_layout = impl_param.get_input_layout(params.end_input_index);
params.inputs.push_back(convert_data_tensor(end_layout));
params.end_dims = end_layout.count();
}
return {params, optional_params};
}
void update_dispatch_data(const kernel_impl_params& impl_param) override {
auto kernel_params = get_kernel_params(impl_param, true);
(_kernel_data.update_dispatch_data_func)(kernel_params.first, _kernel_data);
}
};
namespace detail {

22
src/plugins/intel_gpu/src/graph/include/batch_to_space_inst.h
View File

@@ -10,6 +10,16 @@
namespace cldnn {
template <>
struct typed_program_node<batch_to_space> : public typed_program_node_base<batch_to_space> {
using parent = typed_program_node_base<batch_to_space>;
public:
using parent::parent;
program_node& input(size_t index = 0) const { return get_dependency(index); }
std::vector<size_t> get_shape_infer_dependencies() const override { return {1, 2, 3}; }
};
using batch_to_space_node = typed_program_node<batch_to_space>;
template <>
@@ -18,8 +28,20 @@ class typed_primitive_inst<batch_to_space> : public typed_primitive_inst_base<ba
using parent::parent;
public:
template<typename ShapeType>
static std::vector<layout> calc_output_layouts(batch_to_space_node const& /*node*/, const kernel_impl_params& impl_param);
static layout calc_output_layout(batch_to_space_node const& node, kernel_impl_params const& impl_param);
static std::string to_string(batch_to_space_node const& node);
bool need_reset_output_memory() const override {
const auto desc = _impl_params->typed_desc<batch_to_space>();
if (desc->shape_constant) {
return true;
}
return false;
}
typed_primitive_inst(network& network, batch_to_space_node const& desc);
};

22
src/plugins/intel_gpu/src/graph/include/space_to_batch_inst.h
View File

@@ -11,6 +11,17 @@
namespace cldnn {
template <>
struct typed_program_node<space_to_batch> : public typed_program_node_base<space_to_batch> {
using parent = typed_program_node_base<space_to_batch>;
public:
using parent::parent;
program_node& input(size_t index = 0) const { return get_dependency(index); }
std::vector<size_t> get_shape_infer_dependencies() const override { return {1, 2, 3}; }
};
using space_to_batch_node = typed_program_node<space_to_batch>;
template <>
@@ -19,9 +30,20 @@ class typed_primitive_inst<space_to_batch> : public typed_primitive_inst_base<sp
using parent::parent;
public:
template<typename ShapeType>
static std::vector<layout> calc_output_layouts(space_to_batch_node const& /*node*/, const kernel_impl_params& impl_param);
static layout calc_output_layout(space_to_batch_node const& node, kernel_impl_params const& impl_param);
static std::string to_string(space_to_batch_node const& node);
bool need_reset_output_memory() const override {
const auto desc = _impl_params->typed_desc<space_to_batch>();
if (!desc->shape_constant) {
return true;
}
return false;
}
typed_primitive_inst(network& network, space_to_batch_node const& desc);
};

91
src/plugins/intel_gpu/src/graph/space_to_batch.cpp
View File

@@ -11,6 +11,8 @@
#include <string>
#include <vector>
#include "space_to_batch_shape_inference.hpp"
namespace cldnn {
GPU_DEFINE_PRIMITIVE_TYPE_ID(space_to_batch)
@@ -58,6 +60,95 @@ layout space_to_batch_inst::calc_output_layout(space_to_batch_node const& node,
return layout{output_type, input_format, desc->out_size};
}
static std::vector<int32_t> tensor_to_vec(const tensor& t, const format f) {
std::vector<int32_t> vec(cldnn::format::dimension(f));
for (size_t i = 0; i < vec.size(); ++i) {
vec[i] = t.sizes()[i];
}
std::reverse(vec.begin() + 2, vec.end());
return vec;
}
template<typename ShapeType>
std::vector<layout> space_to_batch_inst::calc_output_layouts(space_to_batch_node const& /*node*/, const kernel_impl_params& impl_param) {
auto desc = impl_param.typed_desc<space_to_batch>();
auto input0_layout = impl_param.get_input_layout(0);
auto input0_shape = input0_layout.get<ShapeType>();
auto input0_size = input0_shape.size();
auto input0_format = input0_layout.format;
auto& constant_mem = impl_param.memory_deps;
auto block_data = desc->block_shape;
auto begin_data = desc->pads_begin;
auto end_data = desc->pads_end;
if (desc->shape_constant == 0 && (!constant_mem.count(1) || !constant_mem.count(2) || !constant_mem.count(3))) {
auto out_shape = ov::PartialShape::dynamic(input0_size);
return { layout{out_shape, input0_layout.data_type, input0_format } };
}
ShapeType block_shape = desc->shape_constant == 0 ? impl_param.get_input_layout(1).get<ShapeType>() : ov::Shape{ input0_size };
ShapeType begin_shape = desc->shape_constant == 0 ? impl_param.get_input_layout(2).get<ShapeType>() : ov::Shape{ input0_size };
ShapeType end_shape = desc->shape_constant == 0 ? impl_param.get_input_layout(3).get<ShapeType>() : ov::Shape{ input0_size };
ov::op::v1::SpaceToBatch op;
std::vector<ShapeType> output_shapes = {ShapeType{}};
std::vector<ShapeType> input_shapes = {
input0_shape,
block_shape,
begin_shape,
end_shape
};
std::map<size_t, ngraph::HostTensorPtr> const_data;
if (desc->shape_constant) {
auto block_sizes = tensor_to_vec(block_data, input0_format);
auto begin_sizes = tensor_to_vec(begin_data, input0_format);
auto end_sizes = tensor_to_vec(end_data, input0_format);
auto block_values = static_cast<void*>(block_sizes.data());
auto begin_values = static_cast<void*>(begin_sizes.data());
auto end_values = static_cast<void*>(end_sizes.data());
auto block_tensor = make_host_tensor({ block_shape, data_types::i32, input0_format }, block_values);
auto begin_tensor = make_host_tensor({ begin_shape, data_types::i32, input0_format }, begin_values);
auto end_tensor = make_host_tensor({ end_shape, data_types::i32, input0_format }, end_values);
const_data.emplace(1, block_tensor);
const_data.emplace(2, begin_tensor);
const_data.emplace(3, end_tensor);
output_shapes = ov::op::v1::shape_infer(&op, input_shapes, ov::make_tensor_accessor(const_data));
} else {
auto block_mem = constant_mem.at(1);
auto begin_mem = constant_mem.at(2);
auto end_mem = constant_mem.at(3);
cldnn::mem_lock<uint8_t, mem_lock_type::read> lock1(block_mem, impl_param.get_stream());
cldnn::mem_lock<uint8_t, mem_lock_type::read> lock2(begin_mem, impl_param.get_stream());
cldnn::mem_lock<uint8_t, mem_lock_type::read> lock3(end_mem, impl_param.get_stream());
auto block_tensor = make_host_tensor(block_mem->get_layout(), lock1.data());
auto begin_tensor = make_host_tensor(begin_mem->get_layout(), lock2.data());
auto end_tensor = make_host_tensor(end_mem->get_layout(), lock3.data());
const_data.emplace(1, block_tensor);
const_data.emplace(2, begin_tensor);
const_data.emplace(3, end_tensor);
output_shapes = ov::op::v1::shape_infer(&op, input_shapes, ov::make_tensor_accessor(const_data));
}
auto output_type = desc->output_data_types[0].value_or(input0_layout.data_type);
if (impl_param.has_fused_primitives())
output_type = impl_param.get_fused_output_layout().data_type;
return { layout{output_shapes[0], output_type, input0_layout.format} };
}
template std::vector<layout> space_to_batch_inst::calc_output_layouts<ov::PartialShape>(space_to_batch_node const& node, const kernel_impl_params& impl_param);
std::string space_to_batch_inst::to_string(space_to_batch_node const& node) {
auto desc = node.get_primitive();
auto node_info = node.desc_to_json();

83
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/batch_to_space_ref.cl
View File

@@ -5,6 +5,15 @@
#include "include/batch_headers/fetch_data.cl"
KERNEL(batch_to_space_ref)(const __global INPUT0_TYPE* input,
#ifdef BLOCK_TYPE
const __global BLOCK_TYPE* block,
#endif
#ifdef BEGIN_TYPE
const __global BEGIN_TYPE* begin,
#endif
#ifdef END_TYPE
const __global END_TYPE* end,
#endif
__global OUTPUT_TYPE* output
#if HAS_FUSED_OPS_DECLS
, FUSED_OPS_DECLS
@@ -14,6 +23,48 @@ KERNEL(batch_to_space_ref)(const __global INPUT0_TYPE* input,
const uint batch = get_global_id(0);
const uint feature = get_global_id(1);
#ifdef BLOCK_TYPE
const uint block_f = block[1];
#if BLOCK_DIMS == 3
const uint block_x = 1;
const uint block_y = block[BLOCK_DIMS-1];
const uint block_z = 1;
const uint block_w = 1;
#else
const uint block_x = BLOCK_DIMS > 2 ? block[BLOCK_DIMS-1] : 1;
const uint block_y = BLOCK_DIMS > 3 ? block[BLOCK_DIMS-2] : 1;
const uint block_z = BLOCK_DIMS > 4 ? block[BLOCK_DIMS-3] : 1;
const uint block_w = BLOCK_DIMS > 5 ? block[BLOCK_DIMS-4] : 1;
#endif
#else
const uint block_f = BLOCK_SHAPE_FEATURE;
const uint block_x = BLOCK_SHAPE_X;
const uint block_y = BLOCK_SHAPE_Y;
const uint block_z = BLOCK_SHAPE_Z;
const uint block_w = BLOCK_SHAPE_W;
#endif
#ifdef BEGIN_TYPE
const uint begin_f = begin[1];
#if BEGIN_DIMS == 3
const uint begin_x = 0;
const uint begin_y = begin[BEGIN_DIMS-1];
const uint begin_z = 0;
const uint begin_w = 0;
#else
const uint begin_x = BEGIN_DIMS > 2 ? begin[BEGIN_DIMS-1] : 0;
const uint begin_y = BEGIN_DIMS > 3 ? begin[BEGIN_DIMS-2] : 0;
const uint begin_z = BEGIN_DIMS > 4 ? begin[BEGIN_DIMS-3] : 0;
const uint begin_w = BEGIN_DIMS > 5 ? begin[BEGIN_DIMS-4] : 0;
#endif
#else
const uint begin_f = CROPS_BEGIN_FEATURE;
const uint begin_x = CROPS_BEGIN_X;
const uint begin_y = CROPS_BEGIN_Y;
const uint begin_z = CROPS_BEGIN_Z;
const uint begin_w = CROPS_BEGIN_W;
#endif
#if OUTPUT_LAYOUT_BFYX || OUTPUT_LAYOUT_B_FS_YX_FSV16
const uint w = 0;
const uint z = 0;
@@ -30,9 +81,9 @@ KERNEL(batch_to_space_ref)(const __global INPUT0_TYPE* input,
const uint y = yx / OUTPUT_SIZE_X;
const uint x = yx % OUTPUT_SIZE_X;
const uint input_w = 0;
const uint input_z = (z + CROPS_BEGIN_Z) / BLOCK_SHAPE_Z;
const uint input_z = (z + begin_z) / block_z;
const uint offset_w = 0;
const uint offset_z = (z + CROPS_BEGIN_Z) % BLOCK_SHAPE_Z;
const uint offset_z = (z + begin_z) % block_z;
#elif OUTPUT_LAYOUT_BFWZYX
const uint w = (uint)get_global_id(2) / (OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z);
const uint zyx = (uint)get_global_id(2) % (OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z);
@@ -40,25 +91,25 @@ KERNEL(batch_to_space_ref)(const __global INPUT0_TYPE* input,
const uint z = zyx / (OUTPUT_SIZE_X * OUTPUT_SIZE_Y);
const uint y = yx / OUTPUT_SIZE_X;
const uint x = yx % OUTPUT_SIZE_X;
const uint input_w = (w + CROPS_BEGIN_W) / BLOCK_SHAPE_W;
const uint input_z = (z + CROPS_BEGIN_Z) / BLOCK_SHAPE_Z;
const uint offset_w = (w + CROPS_BEGIN_W) % BLOCK_SHAPE_W;
const uint offset_z = (z + CROPS_BEGIN_Z) % BLOCK_SHAPE_Z;
const uint input_w = (w + begin_w) / block_w;
const uint input_z = (z + begin_z) / block_z;
const uint offset_w = (w + begin_w) % block_w;
const uint offset_z = (z + begin_z) % block_z;
#endif
const uint input_feature = (feature + CROPS_BEGIN_FEATURE) / BLOCK_SHAPE_FEATURE;
const uint offset_feature = (feature + CROPS_BEGIN_FEATURE) % BLOCK_SHAPE_FEATURE;
const uint input_feature = (feature + begin_f) / block_f;
const uint offset_feature = (feature + begin_f) % block_f;
const uint input_y = (y + CROPS_BEGIN_Y) / BLOCK_SHAPE_Y;
const uint offset_y = (y + CROPS_BEGIN_Y) % BLOCK_SHAPE_Y;
const uint input_y = (y + begin_y) / block_y;
const uint offset_y = (y + begin_y) % block_y;
const uint input_x = (x + CROPS_BEGIN_X) / BLOCK_SHAPE_X;
const uint offset_x = (x + CROPS_BEGIN_X) % BLOCK_SHAPE_X;
const uint input_x = (x + begin_x) / block_x;
const uint offset_x = (x + begin_x) % block_x;
const uint offset_batch = ((offset_feature * BLOCK_SHAPE_W * BLOCK_SHAPE_Z * BLOCK_SHAPE_Y +
offset_w * BLOCK_SHAPE_Z * BLOCK_SHAPE_Y +
offset_z * BLOCK_SHAPE_Y +
offset_y) * BLOCK_SHAPE_X +
const uint offset_batch = ((offset_feature * block_w * block_z * block_y +
offset_w * block_z * block_y +
offset_z * block_y +
offset_y) * block_x +
offset_x) * OUTPUT_BATCH_NUM;
const uint input_batch = batch + offset_batch;

73
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/space_to_batch_ref.cl
View File

@@ -5,6 +5,15 @@
#include "include/batch_headers/fetch_data.cl"
KERNEL(space_to_batch_ref)(const __global INPUT0_TYPE* input,
#ifdef BLOCK_TYPE
const __global BLOCK_TYPE* block,
#endif
#ifdef BEGIN_TYPE
const __global BEGIN_TYPE* begin,
#endif
#ifdef END_TYPE
const __global END_TYPE* end,
#endif
__global OUTPUT_TYPE* output
#if HAS_FUSED_OPS_DECLS
, FUSED_OPS_DECLS
@@ -37,20 +46,64 @@ KERNEL(space_to_batch_ref)(const __global INPUT0_TYPE* input,
const uint input_batch = batch % INPUT0_BATCH_NUM;
const uint offset_batch = batch / INPUT0_BATCH_NUM;
const int input_feature = feature * BLOCK_SHAPE_FEATURE - PADS_BEGIN_FEATURE +
offset_batch / (BLOCK_SHAPE_W * BLOCK_SHAPE_Z * BLOCK_SHAPE_Y * BLOCK_SHAPE_X);
const uint offset_feature = offset_batch % (BLOCK_SHAPE_W * BLOCK_SHAPE_Z * BLOCK_SHAPE_Y * BLOCK_SHAPE_X);
const int input_w = w * BLOCK_SHAPE_W - PADS_BEGIN_W + offset_feature / (BLOCK_SHAPE_Z * BLOCK_SHAPE_Y * BLOCK_SHAPE_X);
const uint offset_w = offset_feature % (BLOCK_SHAPE_Z * BLOCK_SHAPE_Y * BLOCK_SHAPE_X);
#ifdef BLOCK_TYPE
const uint block_f = block[1];
#if BLOCK_DIMS == 3
const uint block_x = 1;
const uint block_y = block[BLOCK_DIMS-1];
const uint block_z = 1;
const uint block_w = 1;
#else
const uint block_x = BLOCK_DIMS > 2 ? block[BLOCK_DIMS-1] : 1;
const uint block_y = BLOCK_DIMS > 3 ? block[BLOCK_DIMS-2] : 1;
const uint block_z = BLOCK_DIMS > 4 ? block[BLOCK_DIMS-3] : 1;
const uint block_w = BLOCK_DIMS > 5 ? block[BLOCK_DIMS-4] : 1;
#endif
#else
const uint block_f = BLOCK_SHAPE_FEATURE;
const uint block_x = BLOCK_SHAPE_X;
const uint block_y = BLOCK_SHAPE_Y;
const uint block_z = BLOCK_SHAPE_Z;
const uint block_w = BLOCK_SHAPE_W;
#endif
const int input_z = z * BLOCK_SHAPE_Z - PADS_BEGIN_Z + offset_w / (BLOCK_SHAPE_Y * BLOCK_SHAPE_X);
const uint offset_z = offset_w % (BLOCK_SHAPE_Y * BLOCK_SHAPE_X);
const int input_y = y * BLOCK_SHAPE_Y - PADS_BEGIN_Y + offset_z / BLOCK_SHAPE_X;
const uint offset_y = offset_z % BLOCK_SHAPE_X;
#ifdef BEGIN_TYPE
const uint begin_f = begin[1];
#if BEGIN_DIMS == 3
const uint begin_x = 0;
const uint begin_y = begin[BEGIN_DIMS-1];
const uint begin_z = 0;
const uint begin_w = 0;
#else
const uint begin_x = BEGIN_DIMS > 2 ? begin[BEGIN_DIMS-1] : 0;
const uint begin_y = BEGIN_DIMS > 3 ? begin[BEGIN_DIMS-2] : 0;
const uint begin_z = BEGIN_DIMS > 4 ? begin[BEGIN_DIMS-3] : 0;
const uint begin_w = BEGIN_DIMS > 5 ? begin[BEGIN_DIMS-4] : 0;
#endif
#else
const uint begin_f = PADS_BEGIN_FEATURE;
const uint begin_x = PADS_BEGIN_X;
const uint begin_y = PADS_BEGIN_Y;
const uint begin_z = PADS_BEGIN_Z;
const uint begin_w = PADS_BEGIN_W;
#endif
const int input_x = x * BLOCK_SHAPE_X - PADS_BEGIN_X + offset_y;
const int input_feature = feature * block_f - begin_f +
offset_batch / (block_w * block_z * block_y * block_x);
const uint offset_feature = offset_batch % (block_w * block_z * block_y * block_x);
const int input_w = w * block_w - begin_w + offset_feature / (block_z * block_y * block_x);
const uint offset_w = offset_feature % (block_z * block_y * block_x);
const int input_z = z * block_z - begin_z + offset_w / (block_y * block_x);
const uint offset_z = offset_w % (block_y * block_x);
const int input_y = y * block_y - begin_y + offset_z / block_x;
const uint offset_y = offset_z % block_x;
const int input_x = x * block_x - begin_x + offset_y;
#if OUTPUT_DIMS == 4
const int input_index = INPUT0_GET_INDEX(input_batch, input_feature, input_y, input_x);

30
src/plugins/intel_gpu/src/kernel_selector/kernels/batch_to_space/batch_to_space_kernel_base.cpp
View File

@@ -48,6 +48,10 @@ CommonDispatchData BatchToSpaceKernelBase::SetDefault(const batch_to_space_param
return dispatchData;
}
inline std::string GetInputTypeStr(size_t idx) {
return "INPUT" + std::to_string(idx) + "_TYPE";
}
JitConstants BatchToSpaceKernelBase::GetJitConstants(const batch_to_space_params& params) const {
JitConstants jit = MakeBaseParamsJitConstants(params);
@@ -70,9 +74,26 @@ JitConstants BatchToSpaceKernelBase::GetJitConstants(const batch_to_space_params
}
};
makeJitConstForParam(jit, "BLOCK_SHAPE", params.block_shape, 1);
makeJitConstForParam(jit, "CROPS_BEGIN", params.crops_begin, 0);
makeJitConstForParam(jit, "CROPS_END", params.crops_end, 0);
if (params.block_type == base_params::ArgType::Input) {
jit.AddConstant(MakeJitConstant("BLOCK_TYPE", GetInputTypeStr(params.block_input_index)));
jit.AddConstant(MakeJitConstant("BLOCK_DIMS", params.block_dims));
} else {
makeJitConstForParam(jit, "BLOCK_SHAPE", params.block_shape, 1);
}
if (params.begin_type == base_params::ArgType::Input) {
jit.AddConstant(MakeJitConstant("BEGIN_TYPE", GetInputTypeStr(params.begin_input_index)));
jit.AddConstant(MakeJitConstant("BEGIN_DIMS", params.begin_dims));
} else {
makeJitConstForParam(jit, "CROPS_BEGIN", params.crops_begin, 0);
}
if (params.end_type == base_params::ArgType::Input) {
jit.AddConstant(MakeJitConstant("END_TYPE", GetInputTypeStr(params.end_input_index)));
jit.AddConstant(MakeJitConstant("END_DIMS", params.end_dims));
} else {
makeJitConstForParam(jit, "CROPS_END", params.crops_end, 0);
}
return jit;
}
@@ -93,7 +114,8 @@ KernelsData BatchToSpaceKernelBase::GetCommonKernelsData(const Params& params, c
auto& kernel = kd.kernels[0];
FillCLKernelData(kernel, dispatchData, params.engineInfo, kernelName, jit, entry_point,
"", false, false, 1, GetFusedPrimitiveInputsCount(params));
"", false, false, static_cast<int>(newParams.inputs.size()),
GetFusedPrimitiveInputsCount(params), 1, newParams.has_dynamic_tensors());
return { kd };
}

12
src/plugins/intel_gpu/src/kernel_selector/kernels/batch_to_space/batch_to_space_kernel_base.h
View File

@@ -17,6 +17,18 @@ struct batch_to_space_params : public base_params {
DimTensor<uint32_t> block_shape;
DimTensor<uint32_t> crops_begin;
DimTensor<uint32_t> crops_end;
base_params::ArgType block_type = base_params::ArgType::Input;
base_params::ArgType begin_type = base_params::ArgType::Input;
base_params::ArgType end_type = base_params::ArgType::Input;
size_t block_dims = 0;
size_t begin_dims = 0;
size_t end_dims = 0;
size_t block_input_index = 0;
size_t begin_input_index = 0;
size_t end_input_index = 0;
};
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

2
src/plugins/intel_gpu/src/kernel_selector/kernels/batch_to_space/batch_to_space_kernel_ref.cpp
View File

@@ -14,11 +14,13 @@ ParamsKey BatchToSpaceKernelRef::GetSupportedKey() const {
k.EnableInputDataType(Datatype::F32);
k.EnableInputDataType(Datatype::UINT8);
k.EnableInputDataType(Datatype::INT8);
k.EnableInputDataType(Datatype::INT32);
k.EnableOutputDataType(Datatype::F16);
k.EnableOutputDataType(Datatype::F32);
k.EnableOutputDataType(Datatype::UINT8);
k.EnableOutputDataType(Datatype::INT8);
k.EnableOutputDataType(Datatype::INT32);
k.EnableInputLayout(DataLayout::bfyx);
k.EnableInputLayout(DataLayout::bfzyx);

30
src/plugins/intel_gpu/src/kernel_selector/kernels/space_to_batch/space_to_batch_kernel_base.cpp
View File

@@ -48,6 +48,10 @@ CommonDispatchData SpaceToBatchKernelBase::SetDefault(const space_to_batch_param
return dispatchData;
}
inline std::string GetInputTypeStr(size_t idx) {
return "INPUT" + std::to_string(idx) + "_TYPE";
}
JitConstants SpaceToBatchKernelBase::GetJitConstants(const space_to_batch_params& params) const {
JitConstants jit = MakeBaseParamsJitConstants(params);
@@ -70,9 +74,26 @@ JitConstants SpaceToBatchKernelBase::GetJitConstants(const space_to_batch_params
}
};
makeJitConstForParam(jit, "BLOCK_SHAPE", params.block_shape, 1);
makeJitConstForParam(jit, "PADS_BEGIN", params.pads_begin, 0);
makeJitConstForParam(jit, "PADS_END", params.pads_end, 0);
if (params.block_type == base_params::ArgType::Input) {
jit.AddConstant(MakeJitConstant("BLOCK_TYPE", GetInputTypeStr(params.block_input_index)));
jit.AddConstant(MakeJitConstant("BLOCK_DIMS", params.block_dims));
} else {
makeJitConstForParam(jit, "BLOCK_SHAPE", params.block_shape, 1);
}
if (params.begin_type == base_params::ArgType::Input) {
jit.AddConstant(MakeJitConstant("BEGIN_TYPE", GetInputTypeStr(params.begin_input_index)));
jit.AddConstant(MakeJitConstant("BEGIN_DIMS", params.begin_dims));
} else {
makeJitConstForParam(jit, "PADS_BEGIN", params.pads_begin, 0);
}
if (params.end_type == base_params::ArgType::Input) {
jit.AddConstant(MakeJitConstant("END_TYPE", GetInputTypeStr(params.end_input_index)));
jit.AddConstant(MakeJitConstant("END_DIMS", params.end_dims));
} else {
makeJitConstForParam(jit, "PADS_END", params.pads_end, 0);
}
return jit;
}
@@ -93,7 +114,8 @@ KernelsData SpaceToBatchKernelBase::GetCommonKernelsData(const Params& params, c
auto& kernel = kd.kernels[0];
FillCLKernelData(kernel, dispatchData, params.engineInfo, kernelName, jit, entry_point,
"", false, false, 1, GetFusedPrimitiveInputsCount(params));
"", false, false, static_cast<int>(newParams.inputs.size()),
GetFusedPrimitiveInputsCount(params), 1, newParams.has_dynamic_tensors());
return { kd };
}

12
src/plugins/intel_gpu/src/kernel_selector/kernels/space_to_batch/space_to_batch_kernel_base.h
View File

@@ -17,6 +17,18 @@ struct space_to_batch_params : public base_params {
DimTensor<uint32_t> block_shape;
DimTensor<uint32_t> pads_begin;
DimTensor<uint32_t> pads_end;
base_params::ArgType block_type = base_params::ArgType::Input;
base_params::ArgType begin_type = base_params::ArgType::Input;
base_params::ArgType end_type = base_params::ArgType::Input;
size_t block_dims = 0;
size_t begin_dims = 0;
size_t end_dims = 0;
size_t block_input_index = 0;
size_t begin_input_index = 0;
size_t end_input_index = 0;
};
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

2
src/plugins/intel_gpu/src/kernel_selector/kernels/space_to_batch/space_to_batch_kernel_ref.cpp
View File

@@ -14,11 +14,13 @@ ParamsKey SpaceToBatchKernelRef::GetSupportedKey() const {
k.EnableInputDataType(Datatype::F32);
k.EnableInputDataType(Datatype::UINT8);
k.EnableInputDataType(Datatype::INT8);
k.EnableInputDataType(Datatype::INT32);
k.EnableOutputDataType(Datatype::F16);
k.EnableOutputDataType(Datatype::F32);
k.EnableOutputDataType(Datatype::UINT8);
k.EnableOutputDataType(Datatype::INT8);
k.EnableOutputDataType(Datatype::INT32);
k.EnableInputLayout(DataLayout::bfyx);
k.EnableInputLayout(DataLayout::bfzyx);

49
src/plugins/intel_gpu/src/plugin/ops/batch_to_space.cpp
View File

@@ -18,33 +18,52 @@ static void CreateBatchToSpaceOp(Program& p, const std::shared_ptr<ngraph::op::v
auto inputs = p.GetInputInfo(op);
std::string layerName = layer_type_name_ID(op);
auto rank = op->get_input_shape(0).size();
auto rank = op->get_input_partial_shape(0).size();
auto format = cldnn::format::get_default_format(rank);
std::vector<cldnn::tensor> tensor_inputs;
tensor_inputs.reserve(3);
bool non_constant_input = false;
for (size_t i = 1; i < 4; ++i) {
auto inConst = std::dynamic_pointer_cast<ngraph::op::Constant>(op->get_input_node_shared_ptr(i));
OPENVINO_ASSERT(inConst != nullptr, "[GPU] Unsupported parameter nodes type in ", op->get_friendly_name(), " (", op->get_type_name(), ")");
std::vector<int32_t> sizes = inConst->cast_vector<int32_t>();
int32_t default_size = i == 1 ? 1 : 0;
for (size_t s = sizes.size(); s < format.dimension(); s++) {
sizes.push_back(default_size);
bool is_const_input = (inConst != nullptr);
OPENVINO_ASSERT((i == 1) || (i >= 2 && non_constant_input != is_const_input),
"[GPU] Unsupported mixed node with constant and parameter in ", op->get_friendly_name(), " (", op->get_type_name(), ")");
if (!inConst) {
non_constant_input = true;
}
tensor_inputs.emplace_back(format, sizes, default_size);
}
auto out_size = tensor_from_dims(op->get_output_shape(0));
auto batchToSpacePrim = cldnn::batch_to_space(layerName,
inputs[0], // input
tensor_inputs[0], // block_shape
tensor_inputs[1], // crops_begin
tensor_inputs[2], // crops_end
out_size);
auto output_pshape = op->get_output_partial_shape(0);
auto out_size = output_pshape.is_static() ? tensor_from_dims(output_pshape.to_shape()) : cldnn::tensor();
p.add_primitive(*op, batchToSpacePrim);
if (non_constant_input) {
auto batchToSpacePrim = cldnn::batch_to_space(layerName, inputs, out_size);
p.add_primitive(*op, batchToSpacePrim);
} else {
for (size_t i = 1; i < 4; ++i) {
auto inConst = std::dynamic_pointer_cast<ngraph::op::Constant>(op->get_input_node_shared_ptr(i));
std::vector<int32_t> sizes = inConst->cast_vector<int32_t>();
int32_t default_size = i == 1 ? 1 : 0;
for (size_t s = sizes.size(); s < format.dimension(); s++) {
sizes.push_back(default_size);
}
tensor_inputs.emplace_back(format, sizes, default_size);
}
auto batchToSpacePrim = cldnn::batch_to_space(layerName,
inputs[0], // input
tensor_inputs[0], // block_shape
tensor_inputs[1], // crops_begin
tensor_inputs[2], // crops_end
out_size);
p.add_primitive(*op, batchToSpacePrim);
}
}
REGISTER_FACTORY_IMPL(v1, BatchToSpace);

46
src/plugins/intel_gpu/src/plugin/ops/space_to_batch.cpp
View File

@@ -24,30 +24,48 @@ static void CreateSpaceToBatchOp(Program& p, const std::shared_ptr<ngraph::op::v
std::vector<cldnn::tensor> tensor_inputs;
tensor_inputs.reserve(3);
bool non_constant_input = false;
for (size_t i = 1; i < 4; ++i) {
auto inConst = std::dynamic_pointer_cast<ngraph::op::Constant>(op->get_input_node_shared_ptr(i));
OPENVINO_ASSERT(inConst != nullptr, "[GPU] Unsupported parameter nodes type in ", op->get_friendly_name(), " (", op->get_type_name(), ")");
std::vector<int32_t> sizes = inConst->cast_vector<int32_t>();
int32_t default_size = i == 1 ? 1 : 0;
for (size_t s = sizes.size(); s < format.dimension(); s++) {
sizes.push_back(default_size);
bool is_const_input = (inConst != nullptr);
OPENVINO_ASSERT((i == 1) || (i >= 2 && non_constant_input != is_const_input),
"[GPU] Unsupported mixed node with constant and parameter in ", op->get_friendly_name(), " (", op->get_type_name(), ")");
if (!inConst) {
non_constant_input = true;
}
tensor_inputs.emplace_back(format, sizes, default_size);
}
auto output_pshape = op->get_output_partial_shape(0);
// In case of dynamic shapes pass dummy shape value to space_to_batch primitive
// To be removed once we enable internal shape infer for all operations
auto output_pshape = op->get_output_partial_shape(0);
auto out_size = output_pshape.is_static() ? tensor_from_dims(output_pshape.to_shape()) : cldnn::tensor();
auto spaceToBatchPrim = cldnn::space_to_batch(layerName,
inputs[0], // input
tensor_inputs[0], // block_shape
tensor_inputs[1], // crops_begin
tensor_inputs[2], // crops_end
out_size);
if (non_constant_input) {
auto spaceToBatchPrim = cldnn::space_to_batch(layerName, inputs, out_size);
p.add_primitive(*op, spaceToBatchPrim);
} else {
for (size_t i = 1; i < 4; ++i) {
auto inConst = std::dynamic_pointer_cast<ngraph::op::Constant>(op->get_input_node_shared_ptr(i));
p.add_primitive(*op, spaceToBatchPrim);
std::vector<int32_t> sizes = inConst->cast_vector<int32_t>();
int32_t default_size = i == 1 ? 1 : 0;
for (size_t s = sizes.size(); s < format.dimension(); s++) {
sizes.push_back(default_size);
}
tensor_inputs.emplace_back(format, sizes, default_size);
}
auto spaceToBatchPrim = cldnn::space_to_batch(layerName,
inputs[0], // input data
tensor_inputs[0], // block_shape
tensor_inputs[1], // crops_begin
tensor_inputs[2], // crops_end
out_size);
p.add_primitive(*op, spaceToBatchPrim);
}
}
REGISTER_FACTORY_IMPL(v1, SpaceToBatch);

227
src/plugins/intel_gpu/tests/functional/single_layer_tests/dynamic/batch_to_space.cpp Normal file
View File

@@ -0,0 +1,227 @@
// Copyright (C) 2023 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "shared_test_classes/single_layer/batch_to_space.hpp"
#include "shared_test_classes/base/ov_subgraph.hpp"
#include "ngraph_functions/builders.hpp"
#include "common_test_utils/test_constants.hpp"
#include "common_test_utils/ov_tensor_utils.hpp"
using namespace InferenceEngine;
using namespace ov::test;
namespace GPULayerTestsDefinitions {
struct BatchToSpaceParams {
std::vector<int64_t> block;
std::vector<int64_t> begin;
std::vector<int64_t> end;
};
typedef std::tuple<
InputShape, // Input shapes
BatchToSpaceParams,
ElementType, // Element type
ngraph::helpers::InputLayerType, // block/begin/end input type
std::map<std::string, std::string> // Additional network configuration
> BatchToSpaceParamsLayerParamSet;
class BatchToSpaceLayerGPUTest : public testing::WithParamInterface<BatchToSpaceParamsLayerParamSet>,
virtual public SubgraphBaseTest {
public:
static std::string getTestCaseName(const testing::TestParamInfo<BatchToSpaceParamsLayerParamSet>& obj) {
InputShape shapes;
BatchToSpaceParams params;
ElementType elementType;
ngraph::helpers::InputLayerType restInputType;
TargetDevice targetDevice;
std::map<std::string, std::string> additionalConfig;
std::tie(shapes, params, elementType, restInputType, additionalConfig) = obj.param;
std::ostringstream results;
results << "IS=" << ov::test::utils::partialShape2str({shapes.first}) << "_";
results << "TS=";
for (const auto& item : shapes.second) {
results << ov::test::utils::vec2str(item) << "_";
}
results << "netPRC=" << elementType << "_";
results << "block=" << ov::test::utils::vec2str(params.block) << "_";
results << "begin=" << ov::test::utils::vec2str(params.begin) << "_";
results << "end=" << ov::test::utils::vec2str(params.end) << "_";
results << "restInputType=" << restInputType << "_";
results << "config=(";
for (const auto& configEntry : additionalConfig) {
results << configEntry.first << ", " << configEntry.second << ":";
}
results << ")";
return results.str();
}
void generate_inputs(const std::vector<ngraph::Shape>& targetInputStaticShapes) override {
inputs.clear();
const auto& funcInputs = function->inputs();
for (size_t i = 0; i < funcInputs.size(); ++i) {
const auto& funcInput = funcInputs[i];
ov::Tensor tensor;
if (i == 1) {
tensor = ov::Tensor(funcInput.get_element_type(), targetInputStaticShapes[i]);
auto *dataPtr = tensor.data<float>();
for (size_t i = 0; i < block.size(); i++) {
dataPtr[i] = static_cast<float>(block[i]);
}
} else if (i == 2) {
tensor = ov::Tensor(funcInput.get_element_type(), targetInputStaticShapes[i]);
auto *dataPtr = tensor.data<float>();
for (size_t i = 0; i < begin.size(); i++) {
dataPtr[i] = static_cast<float>(begin[i]);
}
} else if (i == 3) {
tensor = ov::Tensor(funcInput.get_element_type(), targetInputStaticShapes[i]);
auto *dataPtr = tensor.data<float>();
for (size_t i = 0; i < end.size(); i++) {
dataPtr[i] = static_cast<float>(end[i]);
}
} else {
tensor = ov::test::utils::create_and_fill_tensor(funcInput.get_element_type(), targetInputStaticShapes[i]);
}
inputs.insert({funcInput.get_node_shared_ptr(), tensor});
}
inferRequestNum++;
}
protected:
std::vector<int64_t> block;
std::vector<int64_t> begin;
std::vector<int64_t> end;
size_t inferRequestNum = 0;
void SetUp() override {
InputShape shapes;
BatchToSpaceParams ssParams;
ngraph::helpers::InputLayerType restInputType;
std::map<std::string, std::string> additionalConfig;
std::tie(shapes, ssParams, inType, restInputType, additionalConfig) = this->GetParam();
block = ssParams.block;
begin = ssParams.begin;
end = ssParams.end;
targetDevice = ov::test::utils::DEVICE_GPU;
std::vector<InputShape> inputShapes;
inputShapes.push_back(shapes);
if (restInputType == ngraph::helpers::InputLayerType::PARAMETER) {
inputShapes.push_back(InputShape({static_cast<int64_t>(block.size())}, std::vector<ov::Shape>(shapes.second.size(), {block.size()})));
inputShapes.push_back(InputShape({static_cast<int64_t>(begin.size())}, std::vector<ov::Shape>(shapes.second.size(), {begin.size()})));
inputShapes.push_back(InputShape({static_cast<int64_t>(end.size())}, std::vector<ov::Shape>(shapes.second.size(), {end.size()})));
}
init_input_shapes(inputShapes);
auto params = ngraph::builder::makeDynamicParams(inType, {inputDynamicShapes.front()});
std::shared_ptr<ov::Node> blockInput, beginInput, endInput;
if (restInputType == ngraph::helpers::InputLayerType::PARAMETER) {
auto blockNode = std::make_shared<ngraph::opset1::Parameter>(ngraph::element::Type_t::i64, ov::Shape{block.size()});
auto beginNode = std::make_shared<ngraph::opset1::Parameter>(ngraph::element::Type_t::i64, ov::Shape{begin.size()});
auto endNode = std::make_shared<ngraph::opset1::Parameter>(ngraph::element::Type_t::i64, ov::Shape{end.size()});
params.push_back(blockNode);
params.push_back(beginNode);
params.push_back(endNode);
blockInput = blockNode;
beginInput = beginNode;
endInput = endNode;
} else {
blockInput = std::make_shared<ngraph::opset1::Constant>(ngraph::element::Type_t::i64, ov::Shape{block.size()}, block);
beginInput = std::make_shared<ngraph::opset1::Constant>(ngraph::element::Type_t::i64, ov::Shape{begin.size()}, begin);
endInput = std::make_shared<ngraph::opset1::Constant>(ngraph::element::Type_t::i64, ov::Shape{end.size()}, end);
}
auto ss = std::make_shared<ngraph::op::v1::BatchToSpace>(params[0], blockInput, beginInput, endInput);
ngraph::ResultVector results;
for (size_t i = 0; i < ss->get_output_size(); i++) {
results.push_back(std::make_shared<ngraph::opset1::Result>(ss->output(i)));
}
function = std::make_shared<ngraph::Function>(results, params, "BatchToSpaceFuncTest");
}
};
TEST_P(BatchToSpaceLayerGPUTest, CompareWithRefs) {
SKIP_IF_CURRENT_TEST_IS_DISABLED()
run();
}
namespace {
std::map<std::string, std::string> emptyAdditionalConfig;
const std::vector<ElementType> inputPrecisions = {
ElementType::f32
};
const std::vector<ngraph::helpers::InputLayerType> restInputTypes = {
ngraph::helpers::InputLayerType::CONSTANT,
ngraph::helpers::InputLayerType::PARAMETER
};
const std::vector<InputShape> inputShapesDynamic3D = {
{{-1, -1, -1}, {{48, 3, 3}, {24, 4, 5}}},
};
const std::vector<BatchToSpaceParams> paramsPlain3D = {
BatchToSpaceParams{ { 1, 2, 4 }, { 0, 0, 1 }, { 0, 0, 1 } },
BatchToSpaceParams{ { 1, 3, 2 }, { 0, 1, 0 }, { 0, 2, 1 } },
};
INSTANTIATE_TEST_SUITE_P(smoke_CompareWithRefs_Plain_Dynamic_3D, BatchToSpaceLayerGPUTest,
::testing::Combine(
::testing::ValuesIn(inputShapesDynamic3D),
::testing::ValuesIn(paramsPlain3D),
::testing::ValuesIn(inputPrecisions),
::testing::ValuesIn(restInputTypes),
::testing::Values(emptyAdditionalConfig)),
BatchToSpaceLayerGPUTest::getTestCaseName);
const std::vector<InputShape> inputShapesDynamic4D = {
{{-1, -1, -1, -1}, {{48, 3, 3, 1}, {24, 4, 5, 6}}},
};
const std::vector<BatchToSpaceParams> paramsPlain4D = {
BatchToSpaceParams{ { 1, 2, 4, 3 }, { 0, 0, 1, 0 }, { 0, 0, 1, 0 } },
BatchToSpaceParams{ { 1, 3, 2, 4 }, { 0, 1, 0, 1 }, { 0, 2, 1, 3 } },
};
INSTANTIATE_TEST_SUITE_P(smoke_CompareWithRefs_Plain_Dynamic_4D, BatchToSpaceLayerGPUTest,
::testing::Combine(
::testing::ValuesIn(inputShapesDynamic4D),
::testing::ValuesIn(paramsPlain4D),
::testing::ValuesIn(inputPrecisions),
::testing::ValuesIn(restInputTypes),
::testing::Values(emptyAdditionalConfig)),
BatchToSpaceLayerGPUTest::getTestCaseName);
const std::vector<InputShape> inputShapesDynamic5D = {
{{-1, -1, -1, -1, -1}, {{48, 3, 3, 1, 5}, {96, 4, 5, 6, 7}}},
};
const std::vector<BatchToSpaceParams> paramsPlain5D = {
BatchToSpaceParams{ { 1, 2, 4, 3, 2 }, { 0, 0, 1, 0, 2 }, { 0, 0, 1, 0, 3 } },
BatchToSpaceParams{ { 1, 3, 2, 4, 2 }, { 0, 1, 0, 1, 3 }, { 0, 2, 1, 3, 2 } },
};
INSTANTIATE_TEST_SUITE_P(smoke_CompareWithRefs_Plain_Dynamic_5D, BatchToSpaceLayerGPUTest,
::testing::Combine(
::testing::ValuesIn(inputShapesDynamic5D),
::testing::ValuesIn(paramsPlain5D),
::testing::ValuesIn(inputPrecisions),
::testing::ValuesIn(restInputTypes),
::testing::Values(emptyAdditionalConfig)),
BatchToSpaceLayerGPUTest::getTestCaseName);
} // namespace
} // namespace GPULayerTestsDefinitions

228
src/plugins/intel_gpu/tests/functional/single_layer_tests/dynamic/space_to_batch.cpp Normal file
View File

@@ -0,0 +1,228 @@
// Copyright (C) 2023 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#include "shared_test_classes/single_layer/space_to_batch.hpp"
#include "shared_test_classes/base/ov_subgraph.hpp"
#include "ngraph_functions/builders.hpp"
#include "common_test_utils/test_constants.hpp"
#include "common_test_utils/ov_tensor_utils.hpp"
using namespace InferenceEngine;
using namespace ov::test;
namespace GPULayerTestsDefinitions {
struct SpaceToBatchParams {
std::vector<int64_t> block;
std::vector<int64_t> begin;
std::vector<int64_t> end;
};
typedef std::tuple<
InputShape, // Input shapes
SpaceToBatchParams,
ElementType, // Element type
ngraph::helpers::InputLayerType, // block/begin/end input type
std::map<std::string, std::string> // Additional network configuration
> SpaceToBatchParamsLayerParamSet;
class SpaceToBatchLayerGPUTest : public testing::WithParamInterface<SpaceToBatchParamsLayerParamSet>,
virtual public SubgraphBaseTest {
public:
static std::string getTestCaseName(const testing::TestParamInfo<SpaceToBatchParamsLayerParamSet>& obj) {
InputShape shapes;
SpaceToBatchParams params;
ElementType elementType;
ngraph::helpers::InputLayerType restInputType;
TargetDevice targetDevice;
std::map<std::string, std::string> additionalConfig;
std::tie(shapes, params, elementType, restInputType, additionalConfig) = obj.param;
std::ostringstream results;
results << "IS=" << ov::test::utils::partialShape2str({shapes.first}) << "_";
results << "TS=";
for (const auto& item : shapes.second) {
results << ov::test::utils::vec2str(item) << "_";
}
results << "netPRC=" << elementType << "_";
results << "block=" << ov::test::utils::vec2str(params.block) << "_";
results << "begin=" << ov::test::utils::vec2str(params.begin) << "_";
results << "end=" << ov::test::utils::vec2str(params.end) << "_";
results << "restInputType=" << restInputType << "_";
results << "config=(";
for (const auto& configEntry : additionalConfig) {
results << configEntry.first << ", " << configEntry.second << ":";
}
results << ")";
return results.str();
}
void generate_inputs(const std::vector<ngraph::Shape>& targetInputStaticShapes) override {
inputs.clear();
const auto& funcInputs = function->inputs();
for (size_t i = 0; i < funcInputs.size(); ++i) {
const auto& funcInput = funcInputs[i];
ov::Tensor tensor;
if (i == 1) {
tensor = ov::Tensor(funcInput.get_element_type(), targetInputStaticShapes[i]);
auto *dataPtr = tensor.data<float>();
for (size_t i = 0; i < block.size(); i++) {
dataPtr[i] = static_cast<float>(block[i]);
}
} else if (i == 2) {
tensor = ov::Tensor(funcInput.get_element_type(), targetInputStaticShapes[i]);
auto *dataPtr = tensor.data<float>();
for (size_t i = 0; i < begin.size(); i++) {
dataPtr[i] = static_cast<float>(begin[i]);
}
} else if (i == 3) {
tensor = ov::Tensor(funcInput.get_element_type(), targetInputStaticShapes[i]);
auto *dataPtr = tensor.data<float>();
for (size_t i = 0; i < end.size(); i++) {
dataPtr[i] = static_cast<float>(end[i]);
}
} else {
tensor = ov::test::utils::create_and_fill_tensor(funcInput.get_element_type(), targetInputStaticShapes[i]);
}
inputs.insert({funcInput.get_node_shared_ptr(), tensor});
}
inferRequestNum++;
}
protected:
std::vector<int64_t> block;
std::vector<int64_t> begin;
std::vector<int64_t> end;
size_t inferRequestNum = 0;
void SetUp() override {
InputShape shapes;
SpaceToBatchParams ssParams;
ngraph::helpers::InputLayerType restInputType;
std::map<std::string, std::string> additionalConfig;
std::tie(shapes, ssParams, inType, restInputType, additionalConfig) = this->GetParam();
block = ssParams.block;
begin = ssParams.begin;
end = ssParams.end;
targetDevice = ov::test::utils::DEVICE_GPU;
std::vector<InputShape> inputShapes;
inputShapes.push_back(shapes);
if (restInputType == ngraph::helpers::InputLayerType::PARAMETER) {
inputShapes.push_back(InputShape({static_cast<int64_t>(block.size())}, std::vector<ov::Shape>(shapes.second.size(), {block.size()})));
inputShapes.push_back(InputShape({static_cast<int64_t>(begin.size())}, std::vector<ov::Shape>(shapes.second.size(), {begin.size()})));
inputShapes.push_back(InputShape({static_cast<int64_t>(end.size())}, std::vector<ov::Shape>(shapes.second.size(), {end.size()})));
}
init_input_shapes(inputShapes);
auto params = ngraph::builder::makeDynamicParams(inType, {inputDynamicShapes.front()});
std::shared_ptr<ov::Node> blockInput, beginInput, endInput;
if (restInputType == ngraph::helpers::InputLayerType::PARAMETER) {
auto blockNode = std::make_shared<ngraph::opset1::Parameter>(ngraph::element::Type_t::i64, ov::Shape{block.size()});
auto beginNode = std::make_shared<ngraph::opset1::Parameter>(ngraph::element::Type_t::i64, ov::Shape{begin.size()});
auto endNode = std::make_shared<ngraph::opset1::Parameter>(ngraph::element::Type_t::i64, ov::Shape{end.size()});
params.push_back(blockNode);
params.push_back(beginNode);
params.push_back(endNode);
blockInput = blockNode;
beginInput = beginNode;
endInput = endNode;
} else {
blockInput = std::make_shared<ngraph::opset1::Constant>(ngraph::element::Type_t::i64, ov::Shape{block.size()}, block);
beginInput = std::make_shared<ngraph::opset1::Constant>(ngraph::element::Type_t::i64, ov::Shape{begin.size()}, begin);
endInput = std::make_shared<ngraph::opset1::Constant>(ngraph::element::Type_t::i64, ov::Shape{end.size()}, end);
}
auto ss = std::make_shared<ngraph::op::v1::SpaceToBatch>(params[0], blockInput, beginInput, endInput);
ngraph::ResultVector results;
for (size_t i = 0; i < ss->get_output_size(); i++) {
results.push_back(std::make_shared<ngraph::opset1::Result>(ss->output(i)));
}
function = std::make_shared<ngraph::Function>(results, params, "SpaceToBatchFuncTest");
}
};
TEST_P(SpaceToBatchLayerGPUTest, CompareWithRefs) {
SKIP_IF_CURRENT_TEST_IS_DISABLED()
run();
}
namespace {
std::map<std::string, std::string> emptyAdditionalConfig;
const std::vector<ElementType> inputPrecisions = {
ElementType::f32
};
const std::vector<ngraph::helpers::InputLayerType> restInputTypes = {
ngraph::helpers::InputLayerType::CONSTANT,
ngraph::helpers::InputLayerType::PARAMETER
};
const std::vector<InputShape> inputShapesDynamic3D = {
{{-1, -1, -1}, {{2, 3, 6}}},
};
const std::vector<SpaceToBatchParams> paramsPlain3D = {
SpaceToBatchParams{ { 1, 2, 3 }, { 0, 2, 2 }, { 0, 3, 1 } },
SpaceToBatchParams{ { 1, 4, 5 }, { 0, 4, 5 }, { 0, 9, 4 } },
};
INSTANTIATE_TEST_SUITE_P(smoke_CompareWithRefs_Dynamic3D, SpaceToBatchLayerGPUTest,
::testing::Combine(
::testing::ValuesIn(inputShapesDynamic3D),
::testing::ValuesIn(paramsPlain3D),
::testing::ValuesIn(inputPrecisions),
::testing::ValuesIn(restInputTypes),
::testing::Values(emptyAdditionalConfig)),
SpaceToBatchLayerGPUTest::getTestCaseName);
const std::vector<InputShape> inputShapesDynamic4D = {
{{-1, -1, -1, -1}, {{2, 3, 6, 5}}},
};
const std::vector<SpaceToBatchParams> paramsPlain4D = {
SpaceToBatchParams{ { 1, 1, 2, 3 }, { 0, 2, 2, 2 }, { 0, 3, 4, 5 } },
SpaceToBatchParams{ { 1, 1, 4, 5 }, { 0, 2, 4, 5 }, { 0, 3, 2, 5 } },
};
INSTANTIATE_TEST_SUITE_P(smoke_CompareWithRefs_Dynamic4D, SpaceToBatchLayerGPUTest,
::testing::Combine(
::testing::ValuesIn(inputShapesDynamic4D),
::testing::ValuesIn(paramsPlain4D),
::testing::ValuesIn(inputPrecisions),
::testing::ValuesIn(restInputTypes),
::testing::Values(emptyAdditionalConfig)),
SpaceToBatchLayerGPUTest::getTestCaseName);
const std::vector<InputShape> inputShapesDynamic5D = {
{{-1, -1, -1, -1, -1}, {{2, 3, 6, 5, 7}}},
};
const std::vector<SpaceToBatchParams> paramsPlain5D = {
SpaceToBatchParams{ { 1, 1, 2, 3, 7 }, { 0, 2, 2, 2, 4 }, { 0, 3, 4, 5, 3 } },
SpaceToBatchParams{ { 1, 1, 4, 5, 8 }, { 0, 2, 4, 5, 5 }, { 0, 3, 2, 5, 4 } },
};
INSTANTIATE_TEST_SUITE_P(smoke_CompareWithRefs_Dynamic5D, SpaceToBatchLayerGPUTest,
::testing::Combine(
::testing::ValuesIn(inputShapesDynamic5D),
::testing::ValuesIn(paramsPlain5D),
::testing::ValuesIn(inputPrecisions),
::testing::ValuesIn(restInputTypes),
::testing::Values(emptyAdditionalConfig)),
SpaceToBatchLayerGPUTest::getTestCaseName);
} // namespace
} // namespace GPULayerTestsDefinitions