From 8e0d8dd36b52f5d1cd8ed94a1c3550694d8f35e4 Mon Sep 17 00:00:00 2001 From: Oleksii Khovan Date: Tue, 12 Sep 2023 08:18:04 +0200 Subject: [PATCH] [GPU] Pad-12 (#19083) * GPU primitive and kernel changes to support Pad-12 * Exclude Pad-12 from GPU transformations pipeline * add unit tests * add single-layet test for Pad-12 --- .../intel_gpu/plugin/primitives_list.hpp | 3 + .../include/intel_gpu/primitives/border.hpp | 23 +- src/plugins/intel_gpu/src/graph/border.cpp | 34 +-- .../intel_gpu/src/graph/impls/ocl/border.cpp | 6 +- .../cl_kernels/border_gpu_ref.cl | 114 +++++----- .../kernels/border/border_kernel_base.h | 9 +- src/plugins/intel_gpu/src/plugin/ops/pad.cpp | 17 +- .../src/plugin/transformations_pipeline.cpp | 2 + .../single_layer_tests/pad.cpp | 63 ++++++ .../tests/unit/test_cases/border_gpu_test.cpp | 197 +++++++++++++++++- .../shared/include/single_layer_tests/pad.hpp | 4 + .../shared_test_classes/single_layer/pad.hpp | 19 ++ .../src/single_layer/pad.cpp | 2 +- .../include/ngraph_functions/builders.hpp | 6 +- .../ngraph_functions/src/pad.cpp | 30 ++- 15 files changed, 427 insertions(+), 102 deletions(-) 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 c36997481aa..7dd03dc685e 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 @@ -256,6 +256,9 @@ REGISTER_FACTORY(v10, Unique); REGISTER_FACTORY(v11, Interpolate); REGISTER_FACTORY(v11, TopK); +// ------------------------------ Supported v12 ops ----------------------------- // +REGISTER_FACTORY(v12, Pad); + // --------------------------- Supported internal ops --------------------------- // REGISTER_FACTORY(internal, NonMaxSuppressionIEInternal); REGISTER_FACTORY(internal, GenerateProposalsIEInternal); diff --git a/src/plugins/intel_gpu/include/intel_gpu/primitives/border.hpp b/src/plugins/intel_gpu/include/intel_gpu/primitives/border.hpp index def9b956fff..4d0768c2f13 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/primitives/border.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/primitives/border.hpp @@ -10,17 +10,17 @@ namespace cldnn { /// @brief Adds border around input. /// -/// @details Applies border of specified type around input data. The size of output data is increased +/// @details Applies border of specified type around input data. The size of output data is increased or decreased /// by @c pads_begin and by @c pads_end. /// @n /// @n@b Requirements: -/// @n - @c pads_begin and @c pads_end must be non-negative on all dimensions and compatible +/// @n - @c pads_begin and @c pads_end must be compatible /// with size of input (describe the same dimensions). /// @n - For @c PadMode equal to @c SYMMETRIC, @c pads_begin and @c pads_end /// must be lower than or equal to size of input on corresponding dimension (for all dimensions) /// @n - For @c PadMode equal to @c REFLECT, @c pads_begin and @c pads_end /// must be lower than size of input on corresponding dimension (for all dimensions) -/// @n Breaking any of this conditions will cause exeption throw. +/// @n Breaking any of this conditions will cause exception throw. struct border : public primitive_base { CLDNN_DECLARE_PRIMITIVE(border) @@ -40,12 +40,13 @@ struct border : public primitive_base { /// @param id An identifier of new primitive. /// @param inputs An identifier list of primitives which are not constant input. /// @param non_constant_input_mask Bit mask whether inputs are non-constant or not - /// @param pads_begin Sizes of border that needs to be added from left + /// @param pads_begin Sizes of border that needs to be added (or removed) from left /// (in X dimension) and from top (in Y dimension). - /// @param pads_end Sizes of border that needs to be added from right + /// @param pads_end Sizes of border that needs to be added (or removed) from right /// (in X dimension) and from bottom (in Y dimension). /// @param pad_mode Value of elements which is used for paddings /// @param pad_value Pad's value in case of PadMode::CONSTANT + /// @param allow_negative_pad Allow negative values in pads_begin and pad_end to remove borders /// @param output_padding Optional padding for output from primitive. border(const primitive_id& id, const std::vector& inputs, @@ -54,12 +55,14 @@ struct border : public primitive_base { const ov::CoordinateDiff& pads_end = {}, const ov::op::PadMode pad_mode = ov::op::PadMode::CONSTANT, const float pad_value = 0.0f, + const bool allow_negative_pad = false, const padding& output_padding = padding()) : primitive_base(id, inputs, {output_padding}), pads_begin(pads_begin), pads_end(pads_end), pad_mode(pad_mode), pad_value(pad_value), + allow_negative_pad(allow_negative_pad), non_constant_input_mask(non_constant_input_mask) {} /// @brief Sizes of border that needs to be added from left (in X dimension) and from top (in Y dimension). @@ -69,7 +72,9 @@ struct border : public primitive_base { /// @brief Type of border that needs to be added to the input. ov::op::PadMode pad_mode = ov::op::PadMode::CONSTANT; /// @brief Border value that is used in constant mode. - float pad_value = 0.0f; + float pad_value{0.0}; + /// @brief Allow negative values in pads_begin and pad_end. + bool allow_negative_pad{false}; /// @brief Bit mask whether input is non-constant or not. Position is defined at PAD_NON_CONST_INPUT. int32_t non_constant_input_mask = 0; @@ -79,6 +84,7 @@ struct border : public primitive_base { seed = hash_range(seed, pads_end.begin(), pads_end.end()); seed = hash_combine(seed, pad_mode); seed = hash_combine(seed, pad_value); + seed = hash_combine(seed, allow_negative_pad); seed = hash_combine(seed, non_constant_input_mask); return seed; } @@ -92,7 +98,8 @@ struct border : public primitive_base { return pads_begin == rhs_casted.pads_begin && pads_end == rhs_casted.pads_end && pad_mode == rhs_casted.pad_mode && - pad_value == rhs_casted.pad_value; + pad_value == rhs_casted.pad_value && + allow_negative_pad == rhs_casted.allow_negative_pad; } void save(BinaryOutputBuffer& ob) const override { @@ -102,6 +109,7 @@ struct border : public primitive_base { ob << make_data(&pad_mode, sizeof(ov::op::PadMode)); ob << pad_value; ob << non_constant_input_mask; + ob << allow_negative_pad; } void load(BinaryInputBuffer& ib) override { @@ -111,6 +119,7 @@ struct border : public primitive_base { ib >> make_data(&pad_mode, sizeof(ov::op::PadMode)); ib >> pad_value; ib >> non_constant_input_mask; + ib >> allow_negative_pad; } }; } // namespace cldnn diff --git a/src/plugins/intel_gpu/src/graph/border.cpp b/src/plugins/intel_gpu/src/graph/border.cpp index 2e662781b01..97d8bbc30a6 100644 --- a/src/plugins/intel_gpu/src/graph/border.cpp +++ b/src/plugins/intel_gpu/src/graph/border.cpp @@ -107,6 +107,7 @@ std::string border_inst::to_string(border_node const& node) { border_info.add("pads_end", desc->pads_end); border_info.add("pad mode", desc->pad_mode); border_info.add("pad value", std::to_string(desc->pad_value)); + border_info.add("negative_pad", std::to_string(desc->allow_negative_pad)); node_info->add("border info", border_info); @@ -122,23 +123,24 @@ border_inst::typed_primitive_inst(network& network, border_node const& node) : p } const auto& input_sizes = input_layout.get_dims(); - auto pad_mode = argument->pad_mode; + const auto pad_mode = argument->pad_mode; + const bool allow_negative_pad = argument->allow_negative_pad; - // Check if sizes of border are in proper range. - CLDNN_ERROR_BOOL(node.id(), - "pads_begin border sizes", - std::any_of(argument->pads_begin.begin(), argument->pads_begin.end(), - [](std::ptrdiff_t pad) { - return pad < 0; - }), - "Invalid border size: negative value"); - CLDNN_ERROR_BOOL(node.id(), - "pads_end border sizes", - std::any_of(argument->pads_end.begin(), argument->pads_end.end(), - [](std::ptrdiff_t pad) { - return pad < 0; - }), - "Invalid border size: negative value"); + const auto check_negative_pad = [](std::ptrdiff_t pad) { + return pad < 0; + }; + + if (!allow_negative_pad) { + // Check if sizes of border are in proper range. + CLDNN_ERROR_BOOL(node.id(), + "pads_begin border sizes", + std::any_of(argument->pads_begin.begin(), argument->pads_begin.end(), check_negative_pad), + "Invalid border size: negative value"); + CLDNN_ERROR_BOOL(node.id(), + "pads_end border sizes", + std::any_of(argument->pads_end.begin(), argument->pads_end.end(), check_negative_pad), + "Invalid border size: negative value"); + } if (pad_mode == ov::op::PadMode::SYMMETRIC) { bool valid_pads = true; diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp index 9265f345e25..97096e72720 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp @@ -46,7 +46,7 @@ struct border_impl : typed_primitive_impl_ocl { begin_vec.insert(begin_vec.end(), zeros_to_add, 0); } std::vector pads_begin(begin_vec.begin(), begin_vec.end()); - params.lt_sizes = convert_dim_vector(tensor(pads_format, pads_begin, 0)); + params.lt_sizes = convert_dim_vector(tensor(pads_format, pads_begin, 0)); } else { params.begin_type = kernel_selector::base_params::ArgType::Input; @@ -65,7 +65,7 @@ struct border_impl : typed_primitive_impl_ocl { end_vec.insert(end_vec.end(), zeros_to_add, 0); } std::vector pads_end(end_vec.begin(), end_vec.end()); - params.rb_sizes = convert_dim_vector(tensor(pads_format, pads_end, 0)); + params.rb_sizes = convert_dim_vector(tensor(pads_format, pads_end, 0)); } else { params.end_type = kernel_selector::base_params::ArgType::Input; @@ -100,6 +100,8 @@ struct border_impl : typed_primitive_impl_ocl { OPENVINO_ASSERT(false, "[GPU] Encountered unhandled enum case: PadMode during translation to kernel selector enumeration."); } + params.allow_negative_pad = primitive->allow_negative_pad; + return {params, optional_params}; } diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/border_gpu_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/border_gpu_ref.cl index 6eccd3e3546..97298456773 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/border_gpu_ref.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/border_gpu_ref.cl @@ -19,19 +19,19 @@ KERNEL(border_gpu_ref)( __global OUTPUT_TYPE* output) { #ifdef BEGIN_TYPE - const uint begin_b = begin[0]; - const uint begin_f = begin[1]; + const int begin_b = begin[0]; + const int begin_f = begin[1]; uint begin_offset = 2; #if INPUT0_DIMS == 6 - const uint begin_w = begin[begin_offset]; + const int begin_w = begin[begin_offset]; begin_offset += 1; #endif #if INPUT0_DIMS >= 5 - const uint begin_z = begin[begin_offset]; + const int begin_z = begin[begin_offset]; begin_offset += 1; #endif - const uint begin_y = begin[begin_offset]; - const uint begin_x = begin[begin_offset + 1]; + const int begin_y = begin[begin_offset]; + const int begin_x = begin[begin_offset + 1]; #else const uint begin_b = LT_SIZES_BATCH_NUM; const uint begin_f = LT_SIZES_FEATURE_NUM; @@ -46,19 +46,19 @@ KERNEL(border_gpu_ref)( #endif #ifdef END_TYPE - const uint end_b = end[0]; - const uint end_f = end[1]; + const int end_b = end[0]; + const int end_f = end[1]; uint end_offset = 2; #if INPUT0_DIMS == 6 - const uint end_w = end[end_offset]; + const int end_w = end[end_offset]; end_offset += 1; #endif #if INPUT0_DIMS >= 5 - const uint end_z = end[end_offset]; + const int end_z = end[end_offset]; end_offset += 1; #endif - const uint end_y = end[end_offset]; - const uint end_x = end[end_offset + 1]; + const int end_y = end[end_offset]; + const int end_x = end[end_offset + 1]; #else const uint end_b = RB_SIZES_BATCH_NUM; const uint end_f = RB_SIZES_FEATURE_NUM; @@ -74,65 +74,65 @@ KERNEL(border_gpu_ref)( // [CONSTEXPR] // Border sizes(left-top): - const uint blt_sb = begin_b; - const uint blt_sf = begin_f; - const uint blt_sy = begin_y; - const uint blt_sx = begin_x; + const int blt_sb = begin_b; + const int blt_sf = begin_f; + const int blt_sy = begin_y; + const int blt_sx = begin_x; #if INPUT0_DIMS == 6 - const uint blt_sw = begin_w; + const int blt_sw = begin_w; #else - const uint blt_sw = 0; + const int blt_sw = 0; #endif #if INPUT0_DIMS >= 5 - const uint blt_sz = begin_z; + const int blt_sz = begin_z; #else - const uint blt_sz = 0; + const int blt_sz = 0; #endif // Border sizes(right-bottom): - const uint brb_sb = end_b; - const uint brb_sf = end_f; - const uint brb_sy = end_y; - const uint brb_sx = end_x; + const int brb_sb = end_b; + const int brb_sf = end_f; + const int brb_sy = end_y; + const int brb_sx = end_x; #if INPUT0_DIMS == 6 - const uint brb_sw = end_w; + const int brb_sw = end_w; #else - const uint brb_sw = 0; + const int brb_sw = 0; #endif #if INPUT0_DIMS >= 5 - const uint brb_sz = end_z; + const int brb_sz = end_z; #else - const uint brb_sz = 0; + const int brb_sz = 0; #endif // Input sizes: - const uint in_sx = INPUT0_SIZE_X; - const uint in_sy = INPUT0_SIZE_Y; - const uint in_sz = INPUT0_SIZE_Z; - const uint in_sw = INPUT0_SIZE_W; - const uint in_sf = INPUT0_FEATURE_NUM; - const uint in_sb = INPUT0_BATCH_NUM; + const int in_sx = INPUT0_SIZE_X; + const int in_sy = INPUT0_SIZE_Y; + const int in_sz = INPUT0_SIZE_Z; + const int in_sw = INPUT0_SIZE_W; + const int in_sf = INPUT0_FEATURE_NUM; + const int in_sb = INPUT0_BATCH_NUM; // Input limits (exclusive; tested on output position): - const uint in_lx = in_sx + blt_sx; - const uint in_ly = in_sy + blt_sy; - const uint in_lz = in_sz + blt_sz; - const uint in_lw = in_sw + blt_sw; - const uint in_lf = in_sf + blt_sf; - const uint in_lb = in_sb + blt_sb; + const int in_lx = in_sx + blt_sx; + const int in_ly = in_sy + blt_sy; + const int in_lz = in_sz + blt_sz; + const int in_lw = in_sw + blt_sw; + const int in_lf = in_sf + blt_sf; + const int in_lb = in_sb + blt_sb; - const uint out_xz = (uint) get_global_id(0); - const uint out_yw = (uint) get_global_id(1); - const uint out_fb = (uint) get_global_id(2); + const int out_xz = get_global_id(0); + const int out_yw = get_global_id(1); + const int out_fb = get_global_id(2); - const uint out_f = out_fb % OUTPUT_FEATURE_NUM; - const uint out_b = out_fb / OUTPUT_FEATURE_NUM; + const int out_f = out_fb % OUTPUT_FEATURE_NUM; + const int out_b = out_fb / OUTPUT_FEATURE_NUM; - const uint out_x = out_xz % OUTPUT_SIZE_X; - const uint out_z = out_xz / OUTPUT_SIZE_X; + const int out_x = out_xz % OUTPUT_SIZE_X; + const int out_z = out_xz / OUTPUT_SIZE_X; - const uint out_y = out_yw % OUTPUT_SIZE_Y; - const uint out_w = out_yw / OUTPUT_SIZE_Y; + const int out_y = out_yw % OUTPUT_SIZE_Y; + const int out_w = out_yw / OUTPUT_SIZE_Y; #ifdef BORDER_TYPE_CONSTANT #ifdef BORDER_VALUE_TYPE @@ -148,14 +148,14 @@ KERNEL(border_gpu_ref)( out_f >= blt_sf & out_f < in_lf & out_b >= blt_sb & out_b < in_lb) { - const uint in_x = out_x - blt_sx; - const uint in_y = out_y - blt_sy; - const uint in_z = out_z - blt_sz; - const uint in_w = out_w - blt_sw; - const uint in_f = out_f - blt_sf; - const uint in_b = out_b - blt_sb; + const int in_x = out_x - blt_sx; + const int in_y = out_y - blt_sy; + const int in_z = out_z - blt_sz; + const int in_w = out_w - blt_sw; + const int in_f = out_f - blt_sf; + const int in_b = out_b - blt_sb; - const uint in_pos = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR in_b, in_f, in_w, in_z, in_y, in_x); + const int in_pos = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR in_b, in_f, in_w, in_z, in_y, in_x); in_val = input[in_pos]; } #elif defined BORDER_TYPE_EDGE @@ -192,6 +192,6 @@ KERNEL(border_gpu_ref)( #error Unsupported border type. #endif - const uint out_pos = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR out_b, out_f, out_w, out_z, out_y, out_x); + const int out_pos = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR out_b, out_f, out_w, out_z, out_y, out_x); output[out_pos] = in_val; } diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/border/border_kernel_base.h b/src/plugins/intel_gpu/src/kernel_selector/kernels/border/border_kernel_base.h index aa295a7c35b..0fb98d9bc23 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/border/border_kernel_base.h +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/border/border_kernel_base.h @@ -12,16 +12,19 @@ namespace kernel_selector { // border_params //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// struct border_params : public base_params { - DimTensor<> lt_sizes; - DimTensor<> rb_sizes; + DimTensor lt_sizes; + DimTensor rb_sizes; BorderType b_type; float border_value; + bool allow_negative_pad; ArgType begin_type; ArgType end_type; ArgType pad_value_type; - border_params() : base_params(KernelType::BORDER), b_type(BorderType::CONSTANT), border_value(0.0f), + + border_params() : base_params(KernelType::BORDER), b_type(BorderType::CONSTANT), + border_value(0.0f), allow_negative_pad(false), begin_type(ArgType::Constant), end_type(ArgType::Constant), pad_value_type(ArgType::Constant) {} ParamsKey GetParamsKey() const override { diff --git a/src/plugins/intel_gpu/src/plugin/ops/pad.cpp b/src/plugins/intel_gpu/src/plugin/ops/pad.cpp index 18a5225dac6..af894693a02 100644 --- a/src/plugins/intel_gpu/src/plugin/ops/pad.cpp +++ b/src/plugins/intel_gpu/src/plugin/ops/pad.cpp @@ -13,7 +13,7 @@ namespace ov { namespace intel_gpu { -static void CreatePadOp(ProgramBuilder& p, const std::shared_ptr& op) { +static void CreatePadOpInternal(ProgramBuilder& p, const std::shared_ptr& op, bool allow_negative_pad) { validate_inputs_count(op, {3, 4}); auto inputs = p.GetInputInfo(op); std::string layerName = layer_type_name_ID(op); @@ -56,18 +56,27 @@ static void CreatePadOp(ProgramBuilder& p, const std::shared_ptrget_pad_mode(), - pad_value); + pad_value, + allow_negative_pad); + p.add_primitive(*op, borderPrim); +} - p.add_primitive(*op, tilePrim); +static void CreatePadOp(ProgramBuilder& p, const std::shared_ptr& op) { + CreatePadOpInternal(p, op, false); +} + +static void CreatePadOp(ProgramBuilder& p, const std::shared_ptr& op) { + CreatePadOpInternal(p, op, true); } REGISTER_FACTORY_IMPL(v1, Pad); +REGISTER_FACTORY_IMPL(v12, Pad); } // namespace intel_gpu } // namespace ov diff --git a/src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp b/src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp index d10eb959395..a569404fb1c 100644 --- a/src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp +++ b/src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp @@ -102,6 +102,7 @@ #include "transformations/op_conversions/convert_shapeof3.hpp" #include "transformations/op_conversions/convert_topk11_downgrade.hpp" #include "transformations/op_conversions/eye_decomposition.hpp" +#include "transformations/op_conversions/convert_pad12_downgrade.hpp" #include "transformations/convert_precision.hpp" #include "transformations/init_node_info.hpp" #include "transformations/rt_info/fused_names_attribute.hpp" @@ -269,6 +270,7 @@ void TransformationsPipeline::apply(std::shared_ptr func) { manager.register_pass(); manager.register_pass(); manager.register_pass(); + manager.register_pass(); precisions_map int_convert_precision_map { {ov::element::i64, ov::element::i32}, diff --git a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/pad.cpp b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/pad.cpp index dc469ca8cf4..e207e8911e8 100644 --- a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/pad.cpp +++ b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/pad.cpp @@ -21,6 +21,7 @@ const std::vector argPadValue = {0.f, 1.f, 2.f, -1.f}; const std::vector padMode = { ngraph::helpers::PadMode::EDGE, ngraph::helpers::PadMode::REFLECT, + ngraph::helpers::PadMode::SYMMETRIC }; INSTANTIATE_TEST_SUITE_P(smoke_Pad2DConst, @@ -82,4 +83,66 @@ INSTANTIATE_TEST_SUITE_P(smoke_Pad4D, testing::Values(ov::test::utils::DEVICE_GPU)), PadLayerTest::getTestCaseName); +const std::vector> padsBegin2DMixed = {{0, 0}, {1, 1}, {-2, 0}, {0, 3}, {2, -2}}; +const std::vector> padsEnd2DMixed = {{0, 0}, {1, 1}, {0, 1}, {-3, -2}, {2, -1}}; + +INSTANTIATE_TEST_SUITE_P(smoke_Pad2DConst, + PadLayerTest12, + testing::Combine(testing::ValuesIn(padsEnd2DMixed), + testing::ValuesIn(padsEnd2D), + testing::ValuesIn(argPadValue), + testing::Values(ngraph::helpers::PadMode::CONSTANT), + testing::ValuesIn(netPrecisions), + testing::Values(InferenceEngine::Precision::UNSPECIFIED), + testing::Values(InferenceEngine::Precision::UNSPECIFIED), + testing::Values(InferenceEngine::Layout::ANY), + testing::Values(std::vector{13, 5}), + testing::Values(ov::test::utils::DEVICE_GPU)), + PadLayerTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(smoke_Pad2D, + PadLayerTest12, + testing::Combine(testing::ValuesIn(padsBegin2DMixed), + testing::ValuesIn(padsEnd2DMixed), + testing::Values(-333), + testing::ValuesIn(padMode), + testing::ValuesIn(netPrecisions), + testing::Values(InferenceEngine::Precision::UNSPECIFIED), + testing::Values(InferenceEngine::Precision::UNSPECIFIED), + testing::Values(InferenceEngine::Layout::ANY), + testing::Values(std::vector{13, 5}), + testing::Values(ov::test::utils::DEVICE_GPU)), + PadLayerTest::getTestCaseName); + +const std::vector> padsBegin4DMixed = {{0, 0, 0, 0}, {0, 3, 0, 0}, {0, 0, 0, 1}, {0, 0, -1, 1}, {2, 0, 0, 0}, {0, 3, 0, -1}}; +const std::vector> padsEnd4DMixed = {{0, 0, 0, 0}, {0, 3, 0, 0}, {1, 0, 0, 0}, {0, 0, 0, 2}, {1, -3, 0, 0}, {0, 3, 0, -1}}; + +INSTANTIATE_TEST_SUITE_P(smoke_Pad4DConst, + PadLayerTest12, + testing::Combine(testing::ValuesIn(padsBegin4DMixed), + testing::ValuesIn(padsEnd4DMixed), + testing::ValuesIn(argPadValue), + testing::Values(ngraph::helpers::PadMode::CONSTANT), + testing::ValuesIn(netPrecisions), + testing::Values(InferenceEngine::Precision::UNSPECIFIED), + testing::Values(InferenceEngine::Precision::UNSPECIFIED), + testing::Values(InferenceEngine::Layout::ANY), + testing::Values(std::vector{3, 5, 10, 11}), + testing::Values(ov::test::utils::DEVICE_GPU)), + PadLayerTest::getTestCaseName); + +INSTANTIATE_TEST_SUITE_P(smoke_Pad4D, + PadLayerTest12, + testing::Combine(testing::ValuesIn(padsBegin4DMixed), + testing::ValuesIn(padsEnd4DMixed), + testing::Values(-333), + testing::ValuesIn(padMode), + testing::ValuesIn(netPrecisions), + testing::Values(InferenceEngine::Precision::UNSPECIFIED), + testing::Values(InferenceEngine::Precision::UNSPECIFIED), + testing::Values(InferenceEngine::Layout::ANY), + testing::Values(std::vector{3, 5, 10, 11}), + testing::Values(ov::test::utils::DEVICE_GPU)), + PadLayerTest::getTestCaseName); + } // namespace diff --git a/src/plugins/intel_gpu/tests/unit/test_cases/border_gpu_test.cpp b/src/plugins/intel_gpu/tests/unit/test_cases/border_gpu_test.cpp index 2d5a1b631e4..a8f30a0da42 100644 --- a/src/plugins/intel_gpu/tests/unit/test_cases/border_gpu_test.cpp +++ b/src/plugins/intel_gpu/tests/unit/test_cases/border_gpu_test.cpp @@ -16,6 +16,7 @@ using namespace cldnn; using namespace ::tests; +namespace { template static std::vector generate_rnd_real_input( const std::vector sizes, @@ -55,6 +56,7 @@ using border_test_param = std::tuple, // shape in std::array, // coord diff lt std::array, // coord diff rb + bool, // allow negative pads bool>; // is_caching_test template @@ -65,11 +67,12 @@ public: T pad_value; format::type fmt; std::array sh_in, cd_lt, cd_rb, sh_out; + bool allow_negative_pads; bool is_caching_test; void SetUp() override { ::testing::TestWithParam>::SetUp(); rg.set_seed(GET_SUITE_NAME); - std::tie(pad_mode, pad_value, fmt, sh_in, cd_lt, cd_rb, is_caching_test) = this->GetParam(); + std::tie(pad_mode, pad_value, fmt, sh_in, cd_lt, cd_rb, allow_negative_pads, is_caching_test) = this->GetParam(); sh_out = {sh_in[0] + cd_lt[0] + cd_rb[0], sh_in[1] + cd_lt[1] + cd_rb[1], sh_in[2] + cd_lt[2] + cd_rb[2], @@ -88,7 +91,8 @@ public: ov::CoordinateDiff(cd_lt.begin(), cd_lt.end()), ov::CoordinateDiff(cd_rb.begin(), cd_rb.end()), pad_mode, - pad_value), + pad_value, + allow_negative_pads), reorder("output", input_info("border"), cldnn::format::bfyx, T_dt)); cldnn::network::ptr target_network = get_network(engine, target_topology, get_test_default_config(engine), get_test_stream_ptr(), is_caching_test); target_network->set_input_data("input", input); @@ -103,7 +107,8 @@ public: ov::CoordinateDiff(cd_lt.begin(), cd_lt.end()), ov::CoordinateDiff(cd_rb.begin(), cd_rb.end()), pad_mode, - pad_value)); + pad_value, + allow_negative_pads)); cldnn::network base_network(engine, base_topology, get_test_default_config(engine)); base_network.set_input_data("input", input); @@ -123,6 +128,7 @@ INSTANTIATE_TEST_SUITE_P(border_test_i8, testing::Values(std::array{2, 3, 4, 5}), testing::Values(std::array{1, 2, 3, 4}), testing::Values(std::array{1, 1, 1, 1}), + testing::Values(false), testing::Values(false))); using border_test_u8 = border_test; TEST_P(border_test_u8, border_test_u8) {} @@ -134,6 +140,7 @@ INSTANTIATE_TEST_SUITE_P(border_test_u8, testing::Values(std::array{2, 3, 4, 5}), testing::Values(std::array{1, 2, 3, 4}), testing::Values(std::array{1, 1, 1, 1}), + testing::Values(false), testing::Values(false))); using border_test_i32 = border_test; TEST_P(border_test_i32, border_test_i32) {} @@ -145,7 +152,19 @@ INSTANTIATE_TEST_SUITE_P(border_test_i32, testing::Values(std::array{2, 3, 4, 5}), testing::Values(std::array{1, 2, 3, 4}), testing::Values(std::array{1, 1, 1, 1}), + testing::Values(false), testing::Values(false))); +INSTANTIATE_TEST_SUITE_P(negative_pads, + border_test_i32, + testing::Combine(testing::Values(PAD_MODES), + testing::Values(-333), + testing::Values(format::type::b_fs_yx_fsv16), + testing::Values(std::array{6, 8, 7, 11}), + testing::ValuesIn({std::array{-1, -2, -2, -3}, std::array{-1, 3, 4, -3}}), + testing::ValuesIn({std::array{-1, -2, -2, -1}, std::array{2, -3, 3, -2}}), + testing::Values(true), + testing::Values(false))); + using border_test_f16 = border_test; TEST_P(border_test_f16, border_test_f16) {} INSTANTIATE_TEST_SUITE_P(border_test_f16, @@ -156,6 +175,7 @@ INSTANTIATE_TEST_SUITE_P(border_test_f16, testing::Values(std::array{2, 3, 4, 5}), testing::Values(std::array{1, 2, 3, 4}), testing::Values(std::array{1, 1, 1, 1}), + testing::Values(false), testing::Values(false))); INSTANTIATE_TEST_SUITE_P(export_import, border_test_f16, @@ -165,6 +185,7 @@ INSTANTIATE_TEST_SUITE_P(export_import, testing::Values(std::array{2, 3, 4, 5}), testing::Values(std::array{1, 2, 3, 4}), testing::Values(std::array{1, 1, 1, 1}), + testing::Values(false), testing::Values(true))); using border_test_f32 = border_test; TEST_P(border_test_f32, border_test_f32) {} @@ -176,6 +197,7 @@ INSTANTIATE_TEST_SUITE_P(border_test_f32, testing::Values(std::array{2, 3, 4, 5}), testing::Values(std::array{1, 2, 3, 4}), testing::Values(std::array{1, 1, 1, 1}), + testing::Values(false), testing::Values(false))); INSTANTIATE_TEST_SUITE_P(bsv16fsv16_reorder, @@ -186,6 +208,7 @@ INSTANTIATE_TEST_SUITE_P(bsv16fsv16_reorder, testing::Values(std::array{2, 3, 4, 5}), testing::Values(std::array{1, 2, 3, 4}), testing::Values(std::array{1, 1, 1, 1}), + testing::Values(false), testing::Values(false))); TEST(border_gpu, bsv16fsv16_without_reorder) { @@ -1636,3 +1659,171 @@ TEST(border_gpu, basic_bfyx_2x1x2x3_1x2x3x4_border_constant_dynamic) { } } } + +struct border_dynamic_test_param { + ov::op::PadMode mode; + std::array in_shape; + std::array lt; + std::array rb; +}; + +class border_dynamic_test : public ::testing::TestWithParam { +public: + void SetUp() override { + ::testing::TestWithParam::SetUp(); + + const border_dynamic_test_param p = this->GetParam(); + + mode = p.mode; + in_size_b = p.in_shape[0]; + in_size_f = p.in_shape[1]; + in_size_y = p.in_shape[2]; + in_size_x = p.in_shape[3]; + + blt_size_b = p.lt[0]; + blt_size_f = p.lt[1]; + blt_size_y = p.lt[2]; + blt_size_x = p.lt[3]; + + brb_size_b = p.rb[0]; + brb_size_f = p.rb[1]; + brb_size_y = p.rb[2]; + brb_size_x = p.rb[3]; + + out_size_b = in_size_b + blt_size_b + brb_size_b; + out_size_f = in_size_f + blt_size_f + brb_size_f; + out_size_y = in_size_y + blt_size_y + brb_size_y; + out_size_x = in_size_x + blt_size_x + brb_size_x; + + auto& engine = get_test_engine(); + + const auto input_layout_dynamic = layout{ov::PartialShape::dynamic(4), data_types::f32, format::bfyx}; + const auto input_layout_static = layout{ov::PartialShape{in_size_b, in_size_f, in_size_y, in_size_x}, data_types::f32, format::bfyx}; + const auto input = engine.allocate_memory(input_layout_static); + const auto pads_begin = engine.allocate_memory({{4}, data_types::i32, format::bfyx}); + const auto pads_end = engine.allocate_memory({{4}, data_types::i32, format::bfyx}); + + set_values(pads_begin, {blt_size_b, blt_size_f, blt_size_y, blt_size_x}); + set_values(pads_end, {brb_size_b, brb_size_f, brb_size_y, brb_size_x}); + + constexpr auto pad_value = -333.0f; + + topology topology; + topology.add(input_layout("input", input_layout_dynamic)); + topology.add(data("pads_begin", pads_begin)); + topology.add(data("pads_end", pads_end)); + topology.add(border("output", + {input_info("input"), input_info("pads_begin"), input_info("pads_end")}, + cldnn::border::PAD_NON_CONST_INPUT::BEGIN | + cldnn::border::PAD_NON_CONST_INPUT::END, + std::vector{}, + std::vector{}, + mode, + pad_value, + true)); + + const std::vector sizes{ static_cast(in_size_b), static_cast(in_size_f), + static_cast(in_size_y), static_cast(in_size_x) }; + const std::vector input_data = generate_rnd_real_input(sizes, -8.0f, 8.0f); + set_values(input, input_data); + + ExecutionConfig config = get_test_default_config(engine); + config.set_property(ov::intel_gpu::allow_new_shape_infer(true)); + network network(engine, topology, config); + network.set_input_data("input", input); + + const auto inst = network.get_primitive("output"); + const auto impl = inst->get_impl(); + ASSERT_TRUE(impl != nullptr); + ASSERT_TRUE(impl->is_dynamic()); + + const auto outputs = network.execute(); + ASSERT_EQ(outputs.size(), size_t(1)); + ASSERT_EQ(outputs.begin()->first, "output"); + + const auto output = outputs.at("output").get_memory(); + const cldnn::mem_lock output_ptr(output, get_test_stream()); + + const auto expected_size = out_size_b * out_size_f * out_size_y * out_size_x; + ASSERT_EQ(output_ptr.size(), expected_size); + + for (auto b = 0; b < out_size_b; ++b) { + for (auto f = 0; f < out_size_f; ++f) { + for (auto y = 0; y < out_size_y; ++y) { + for (auto x = 0; x < out_size_x; ++x) { + const auto output_off = ((b * out_size_f + f) * out_size_y + y) * out_size_x + x; + ASSERT_GE(output_off, 0); + + if (mode == ov::op::PadMode::CONSTANT) { + if (b < blt_size_b || b >= out_size_b - brb_size_b || + f < blt_size_f || f >= out_size_f - brb_size_f || + y < blt_size_y || y >= out_size_y - brb_size_y || + x < blt_size_x || x >= out_size_x - brb_size_x) { + ASSERT_EQ(output_ptr[output_off], pad_value); + } else { + const auto input_off = (((b - blt_size_b) * in_size_f + f - blt_size_f) * in_size_y + y - blt_size_y) * in_size_x + x - blt_size_x; // BFYX + ASSERT_GE(input_off, 0); + ASSERT_EQ(output_ptr[output_off], input_data[input_off]); + } + } else { + int in_b, in_f, in_y, in_x; + CalcInIndices(b, f, y, x, in_b, in_f, in_y, in_x); + const auto input_off = ((in_b * in_size_f + in_f) * in_size_y + in_y) * in_size_x + in_x; + ASSERT_GE(input_off, 0); + ASSERT_EQ(output_ptr[output_off], input_data[input_off]); + } + } + } + } + } + } + +private: + void CalcInIndices(const int b, const int f, const int y, const int x, int& in_b, int& in_f, int& in_y, int& in_x) { + switch (mode) { + case ov::op::PadMode::REFLECT: { + in_b = (b >= blt_size_b && b < out_size_b - brb_size_b) ? b - blt_size_b : (b < blt_size_b ? blt_size_b - b : in_size_b + out_size_b - brb_size_b - 2 - b); + in_f = (f >= blt_size_f && f < out_size_f - brb_size_f) ? f - blt_size_f : (f < blt_size_f ? blt_size_f - f : in_size_f + out_size_f - brb_size_f - 2 - f); + in_y = (y >= blt_size_y && y < out_size_y - brb_size_y) ? y - blt_size_y : (y < blt_size_y ? blt_size_y - y : in_size_y + out_size_y - brb_size_y - 2 - y); + in_x = (x >= blt_size_x && x < out_size_x - brb_size_x) ? x - blt_size_x : (x < blt_size_x ? blt_size_x - x : in_size_x + out_size_x - brb_size_x - 2 - x); + break; + } + case ov::op::PadMode::SYMMETRIC: { + in_b = (b >= blt_size_b && b < out_size_b - brb_size_b) ? b - blt_size_b : (b < blt_size_b ? blt_size_b - 1 - b : in_size_b + out_size_b - brb_size_b - 1 - b); + in_f = (f >= blt_size_f && f < out_size_f - brb_size_f) ? f - blt_size_f : (f < blt_size_f ? blt_size_f - 1 - f : in_size_f + out_size_f - brb_size_f - 1 - f); + in_y = (y >= blt_size_y && y < out_size_y - brb_size_y) ? y - blt_size_y : (y < blt_size_y ? blt_size_y - 1 - y : in_size_y + out_size_y - brb_size_y - 1 - y); + in_x = (x >= blt_size_x && x < out_size_x - brb_size_x) ? x - blt_size_x : (x < blt_size_x ? blt_size_x - 1 - x : in_size_x + out_size_x - brb_size_x - 1 - x); + break; + } + case ov::op::PadMode::EDGE: { + in_b = (b >= blt_size_b && b < out_size_b - brb_size_b) ? b - blt_size_b : (b < blt_size_b ? 0 : in_size_b - 1); + in_f = (f >= blt_size_f && f < out_size_f - brb_size_f) ? f - blt_size_f : (f < blt_size_f ? 0 : in_size_f - 1); + in_y = (y >= blt_size_y && y < out_size_y - brb_size_y) ? y - blt_size_y : (y < blt_size_y ? 0 : in_size_y - 1); + in_x = (x >= blt_size_x && x < out_size_x - brb_size_x) ? x - blt_size_x : (x < blt_size_x ? 0 : in_size_x - 1); + break; + } + default: { + throw std::runtime_error("Invalid PadMode"); + } + } + } + + ov::op::PadMode mode; + int in_size_b, in_size_f, in_size_y, in_size_x; + int blt_size_b, blt_size_f, blt_size_y, blt_size_x; + int brb_size_b, brb_size_f, brb_size_y, brb_size_x; + int out_size_b, out_size_f, out_size_y, out_size_x; +}; + +const std::vector dynamic_params { + {ov::op::PadMode::CONSTANT, {2, 3, 5, 4}, {-1, 2, -2, 3}, {2, -1, 3, -2}}, + {ov::op::PadMode::EDGE, {3, 4, 6, 5}, {-1, 1, -3, 2}, {3, -1, 1, -3}}, + {ov::op::PadMode::REFLECT, {3, 4, 6, 5}, {-1, 1, -3, 2}, {2, -1, 2, -3}}, + {ov::op::PadMode::SYMMETRIC, {2, 3, 5, 4}, {-1, 2, -2, 3}, {2, -1, 3, -2}} + }; +TEST_P(border_dynamic_test, border_dynamic_test) {} +INSTANTIATE_TEST_SUITE_P(border_dynamic_test, + border_dynamic_test, + ::testing::ValuesIn(dynamic_params)); +}; // namespace + diff --git a/src/tests/functional/plugin/shared/include/single_layer_tests/pad.hpp b/src/tests/functional/plugin/shared/include/single_layer_tests/pad.hpp index a58f05ca27a..8bdf9a3d2c2 100644 --- a/src/tests/functional/plugin/shared/include/single_layer_tests/pad.hpp +++ b/src/tests/functional/plugin/shared/include/single_layer_tests/pad.hpp @@ -12,4 +12,8 @@ TEST_P(PadLayerTest, CompareWithRefs) { Run(); } +TEST_P(PadLayerTest12, CompareWithRefs) { + Run(); +} + } // namespace LayerTestsDefinitions diff --git a/src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/pad.hpp b/src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/pad.hpp index ff9a9f2712b..613a0659b24 100644 --- a/src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/pad.hpp +++ b/src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/pad.hpp @@ -33,6 +33,25 @@ public: protected: void SetUp() override; + virtual std::shared_ptr CreatePadOp(const ngraph::Output& data, + const std::vector& padsBegin, + const std::vector& padsEnd, + float argPadValue, + ngraph::helpers::PadMode padMode) const { + const auto pad = ngraph::builder::makePad(data, padsBegin, padsEnd, argPadValue, padMode, false); + return pad; + } }; +class PadLayerTest12 : public PadLayerTest { +protected: + std::shared_ptr CreatePadOp(const ngraph::Output& data, + const std::vector& padsBegin, + const std::vector& padsEnd, + float argPadValue, + ngraph::helpers::PadMode padMode) const override { + const auto pad = ngraph::builder::makePad(data, padsBegin, padsEnd, argPadValue, padMode, true); + return pad; + } +}; } // namespace LayerTestsDefinitions diff --git a/src/tests/functional/shared_test_classes/src/single_layer/pad.cpp b/src/tests/functional/shared_test_classes/src/single_layer/pad.cpp index 6f796483b2b..2c92716bed8 100644 --- a/src/tests/functional/shared_test_classes/src/single_layer/pad.cpp +++ b/src/tests/functional/shared_test_classes/src/single_layer/pad.cpp @@ -46,7 +46,7 @@ void PadLayerTest::SetUp() { ov::ParameterVector params{std::make_shared(ngPrc, ov::Shape(inputShape))}; auto paramOuts = ngraph::helpers::convert2OutputVector( ngraph::helpers::castOps2Nodes(params)); - auto pad = ngraph::builder::makePad(paramOuts[0], padsBegin, padsEnd, argPadValue, padMode); + auto pad = CreatePadOp(paramOuts[0], padsBegin, padsEnd, argPadValue, padMode); ngraph::ResultVector results{std::make_shared(pad)}; function = std::make_shared(results, params, "pad"); } diff --git a/src/tests/ngraph_helpers/ngraph_functions/include/ngraph_functions/builders.hpp b/src/tests/ngraph_helpers/ngraph_functions/include/ngraph_functions/builders.hpp index d3f4b516589..b4ce38a5921 100644 --- a/src/tests/ngraph_helpers/ngraph_functions/include/ngraph_functions/builders.hpp +++ b/src/tests/ngraph_helpers/ngraph_functions/include/ngraph_functions/builders.hpp @@ -521,13 +521,15 @@ std::shared_ptr makePad(const ngraph::Output& data, const std::vector& padsBegin, const std::vector& padsEnd, float argPadValue, - ngraph::helpers::PadMode padMode); + ngraph::helpers::PadMode padMode, + const bool allow_negative_pad = false); std::shared_ptr makePad(const ov::Output& in, const ov::Output& beginNode, const ov::Output& endNode, const ov::Output& valueNode, - ngraph::helpers::PadMode padMode); + ngraph::helpers::PadMode padMode, + const bool allow_negative_pad = false); std::shared_ptr makeBatchNormInference(const ngraph::Output& data, double epsilon); diff --git a/src/tests/ngraph_helpers/ngraph_functions/src/pad.cpp b/src/tests/ngraph_helpers/ngraph_functions/src/pad.cpp index 962353afebb..8b53059a023 100644 --- a/src/tests/ngraph_helpers/ngraph_functions/src/pad.cpp +++ b/src/tests/ngraph_helpers/ngraph_functions/src/pad.cpp @@ -13,7 +13,8 @@ std::shared_ptr makePad(const ngraph::Output& data, const std::vector& padsBegin, const std::vector& padsEnd, float argPadValue, - ngraph::helpers::PadMode padMode) { + ngraph::helpers::PadMode padMode, + const bool allow_negative_pad) { ngraph::op::PadMode pad_mode; switch (padMode) { case ngraph::helpers::PadMode::CONSTANT: @@ -37,14 +38,20 @@ std::shared_ptr makePad(const ngraph::Output& data, auto pads_end = std::make_shared(ngraph::element::i64, ngraph::Shape{padsEnd.size()}, padsEnd.data()); auto arg_pad_value = std::make_shared(data.get_element_type(), ngraph::Shape{}, &argPadValue); - return std::make_shared(data, pads_begin, pads_end, arg_pad_value, pad_mode); + + if (allow_negative_pad) { + return std::make_shared(data, pads_begin, pads_end, arg_pad_value, pad_mode); + } else { + return std::make_shared(data, pads_begin, pads_end, arg_pad_value, pad_mode); + } } std::shared_ptr makePad(const ov::Output& in, const ov::Output& beginNode, const ov::Output& endNode, const ov::Output& valueNode, - ngraph::helpers::PadMode padMode) { + ngraph::helpers::PadMode padMode, + const bool allow_negative_pad) { ngraph::op::PadMode pad_mode; switch (padMode) { case ngraph::helpers::PadMode::CONSTANT: @@ -62,10 +69,19 @@ std::shared_ptr makePad(const ov::Output& in, default: throw std::runtime_error("Can't create layer for this pad mode"); } - if (valueNode.get_node_shared_ptr() == nullptr) - return std::make_shared(in, beginNode, endNode, pad_mode); - else - return std::make_shared(in, beginNode, endNode, valueNode, pad_mode); + if (valueNode.get_node_shared_ptr() == nullptr) { + if (allow_negative_pad) { + return std::make_shared(in, beginNode, endNode, pad_mode); + } else { + return std::make_shared(in, beginNode, endNode, pad_mode); + } + } else { + if (allow_negative_pad) { + return std::make_shared(in, beginNode, endNode, valueNode, pad_mode); + } else { + return std::make_shared(in, beginNode, endNode, valueNode, pad_mode); + } + } } } // namespace builder