Skip to content

Commit 8e0d8dd

Browse files
okhovanokhovan
authoredSep 12, 2023
[GPU] Pad-12 (openvinotoolkit#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
1 parent 016c7de commit 8e0d8dd

File tree

15 files changed

+429
-104
lines changed

15 files changed

+429
-104
lines changed
 

‎src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp

+3
Original file line numberDiff line numberDiff line change
@@ -256,6 +256,9 @@ REGISTER_FACTORY(v10, Unique);
256256
REGISTER_FACTORY(v11, Interpolate);
257257
REGISTER_FACTORY(v11, TopK);
258258

259+
// ------------------------------ Supported v12 ops ----------------------------- //
260+
REGISTER_FACTORY(v12, Pad);
261+
259262
// --------------------------- Supported internal ops --------------------------- //
260263
REGISTER_FACTORY(internal, NonMaxSuppressionIEInternal);
261264
REGISTER_FACTORY(internal, GenerateProposalsIEInternal);

‎src/plugins/intel_gpu/include/intel_gpu/primitives/border.hpp

+16-7
Original file line numberDiff line numberDiff line change
@@ -10,17 +10,17 @@ namespace cldnn {
1010

1111
/// @brief Adds border around input.
1212
///
13-
/// @details Applies border of specified type around input data. The size of output data is increased
13+
/// @details Applies border of specified type around input data. The size of output data is increased or decreased
1414
/// by @c pads_begin and by @c pads_end.
1515
/// @n
1616
/// @n@b Requirements:
17-
/// @n - @c pads_begin and @c pads_end must be non-negative on all dimensions and compatible
17+
/// @n - @c pads_begin and @c pads_end must be compatible
1818
/// with size of input (describe the same dimensions).
1919
/// @n - For @c PadMode equal to @c SYMMETRIC, @c pads_begin and @c pads_end
2020
/// must be lower than or equal to size of input on corresponding dimension (for all dimensions)
2121
/// @n - For @c PadMode equal to @c REFLECT, @c pads_begin and @c pads_end
2222
/// must be lower than size of input on corresponding dimension (for all dimensions)
23-
/// @n Breaking any of this conditions will cause exeption throw.
23+
/// @n Breaking any of this conditions will cause exception throw.
2424
struct border : public primitive_base<border> {
2525
CLDNN_DECLARE_PRIMITIVE(border)
2626

@@ -40,12 +40,13 @@ struct border : public primitive_base<border> {
4040
/// @param id An identifier of new primitive.
4141
/// @param inputs An identifier list of primitives which are not constant input.
4242
/// @param non_constant_input_mask Bit mask whether inputs are non-constant or not
43-
/// @param pads_begin Sizes of border that needs to be added from left
43+
/// @param pads_begin Sizes of border that needs to be added (or removed) from left
4444
/// (in X dimension) and from top (in Y dimension).
45-
/// @param pads_end Sizes of border that needs to be added from right
45+
/// @param pads_end Sizes of border that needs to be added (or removed) from right
4646
/// (in X dimension) and from bottom (in Y dimension).
4747
/// @param pad_mode Value of elements which is used for paddings
4848
/// @param pad_value Pad's value in case of PadMode::CONSTANT
49+
/// @param allow_negative_pad Allow negative values in pads_begin and pad_end to remove borders
4950
/// @param output_padding Optional padding for output from primitive.
5051
border(const primitive_id& id,
5152
const std::vector<input_info>& inputs,
@@ -54,12 +55,14 @@ struct border : public primitive_base<border> {
5455
const ov::CoordinateDiff& pads_end = {},
5556
const ov::op::PadMode pad_mode = ov::op::PadMode::CONSTANT,
5657
const float pad_value = 0.0f,
58+
const bool allow_negative_pad = false,
5759
const padding& output_padding = padding())
5860
: primitive_base(id, inputs, {output_padding}),
5961
pads_begin(pads_begin),
6062
pads_end(pads_end),
6163
pad_mode(pad_mode),
6264
pad_value(pad_value),
65+
allow_negative_pad(allow_negative_pad),
6366
non_constant_input_mask(non_constant_input_mask) {}
6467

6568
/// @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<border> {
6972
/// @brief Type of border that needs to be added to the input.
7073
ov::op::PadMode pad_mode = ov::op::PadMode::CONSTANT;
7174
/// @brief Border value that is used in constant mode.
72-
float pad_value = 0.0f;
75+
float pad_value{0.0};
76+
/// @brief Allow negative values in pads_begin and pad_end.
77+
bool allow_negative_pad{false};
7378
/// @brief Bit mask whether input is non-constant or not. Position is defined at PAD_NON_CONST_INPUT.
7479
int32_t non_constant_input_mask = 0;
7580

@@ -79,6 +84,7 @@ struct border : public primitive_base<border> {
7984
seed = hash_range(seed, pads_end.begin(), pads_end.end());
8085
seed = hash_combine(seed, pad_mode);
8186
seed = hash_combine(seed, pad_value);
87+
seed = hash_combine(seed, allow_negative_pad);
8288
seed = hash_combine(seed, non_constant_input_mask);
8389
return seed;
8490
}
@@ -92,7 +98,8 @@ struct border : public primitive_base<border> {
9298
return pads_begin == rhs_casted.pads_begin &&
9399
pads_end == rhs_casted.pads_end &&
94100
pad_mode == rhs_casted.pad_mode &&
95-
pad_value == rhs_casted.pad_value;
101+
pad_value == rhs_casted.pad_value &&
102+
allow_negative_pad == rhs_casted.allow_negative_pad;
96103
}
97104

98105
void save(BinaryOutputBuffer& ob) const override {
@@ -102,6 +109,7 @@ struct border : public primitive_base<border> {
102109
ob << make_data(&pad_mode, sizeof(ov::op::PadMode));
103110
ob << pad_value;
104111
ob << non_constant_input_mask;
112+
ob << allow_negative_pad;
105113
}
106114

107115
void load(BinaryInputBuffer& ib) override {
@@ -111,6 +119,7 @@ struct border : public primitive_base<border> {
111119
ib >> make_data(&pad_mode, sizeof(ov::op::PadMode));
112120
ib >> pad_value;
113121
ib >> non_constant_input_mask;
122+
ib >> allow_negative_pad;
114123
}
115124
};
116125
} // namespace cldnn

‎src/plugins/intel_gpu/src/graph/border.cpp

+19-17
Original file line numberDiff line numberDiff line change
@@ -107,6 +107,7 @@ std::string border_inst::to_string(border_node const& node) {
107107
border_info.add("pads_end", desc->pads_end);
108108
border_info.add("pad mode", desc->pad_mode);
109109
border_info.add("pad value", std::to_string(desc->pad_value));
110+
border_info.add("negative_pad", std::to_string(desc->allow_negative_pad));
110111

111112
node_info->add("border info", border_info);
112113

@@ -122,23 +123,24 @@ border_inst::typed_primitive_inst(network& network, border_node const& node) : p
122123
}
123124

124125
const auto& input_sizes = input_layout.get_dims();
125-
auto pad_mode = argument->pad_mode;
126-
127-
// Check if sizes of border are in proper range.
128-
CLDNN_ERROR_BOOL(node.id(),
129-
"pads_begin border sizes",
130-
std::any_of(argument->pads_begin.begin(), argument->pads_begin.end(),
131-
[](std::ptrdiff_t pad) {
132-
return pad < 0;
133-
}),
134-
"Invalid border size: negative value");
135-
CLDNN_ERROR_BOOL(node.id(),
136-
"pads_end border sizes",
137-
std::any_of(argument->pads_end.begin(), argument->pads_end.end(),
138-
[](std::ptrdiff_t pad) {
139-
return pad < 0;
140-
}),
141-
"Invalid border size: negative value");
126+
const auto pad_mode = argument->pad_mode;
127+
const bool allow_negative_pad = argument->allow_negative_pad;
128+
129+
const auto check_negative_pad = [](std::ptrdiff_t pad) {
130+
return pad < 0;
131+
};
132+
133+
if (!allow_negative_pad) {
134+
// Check if sizes of border are in proper range.
135+
CLDNN_ERROR_BOOL(node.id(),
136+
"pads_begin border sizes",
137+
std::any_of(argument->pads_begin.begin(), argument->pads_begin.end(), check_negative_pad),
138+
"Invalid border size: negative value");
139+
CLDNN_ERROR_BOOL(node.id(),
140+
"pads_end border sizes",
141+
std::any_of(argument->pads_end.begin(), argument->pads_end.end(), check_negative_pad),
142+
"Invalid border size: negative value");
143+
}
142144

143145
if (pad_mode == ov::op::PadMode::SYMMETRIC) {
144146
bool valid_pads = true;

‎src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp

+4-2
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,7 @@ struct border_impl : typed_primitive_impl_ocl<border> {
4646
begin_vec.insert(begin_vec.end(), zeros_to_add, 0);
4747
}
4848
std::vector<tensor::value_type> pads_begin(begin_vec.begin(), begin_vec.end());
49-
params.lt_sizes = convert_dim_vector(tensor(pads_format, pads_begin, 0));
49+
params.lt_sizes = convert_dim_vector<int32_t>(tensor(pads_format, pads_begin, 0));
5050
} else {
5151
params.begin_type = kernel_selector::base_params::ArgType::Input;
5252

@@ -65,7 +65,7 @@ struct border_impl : typed_primitive_impl_ocl<border> {
6565
end_vec.insert(end_vec.end(), zeros_to_add, 0);
6666
}
6767
std::vector<tensor::value_type> pads_end(end_vec.begin(), end_vec.end());
68-
params.rb_sizes = convert_dim_vector(tensor(pads_format, pads_end, 0));
68+
params.rb_sizes = convert_dim_vector<int32_t>(tensor(pads_format, pads_end, 0));
6969
} else {
7070
params.end_type = kernel_selector::base_params::ArgType::Input;
7171

@@ -100,6 +100,8 @@ struct border_impl : typed_primitive_impl_ocl<border> {
100100
OPENVINO_ASSERT(false, "[GPU] Encountered unhandled enum case: PadMode during translation to kernel selector enumeration.");
101101
}
102102

103+
params.allow_negative_pad = primitive->allow_negative_pad;
104+
103105
return {params, optional_params};
104106
}
105107

‎src/plugins/intel_gpu/src/kernel_selector/cl_kernels/border_gpu_ref.cl

+58-58
Original file line numberDiff line numberDiff line change
@@ -19,19 +19,19 @@ KERNEL(border_gpu_ref)(
1919
__global OUTPUT_TYPE* output)
2020
{
2121
#ifdef BEGIN_TYPE
22-
const uint begin_b = begin[0];
23-
const uint begin_f = begin[1];
22+
const int begin_b = begin[0];
23+
const int begin_f = begin[1];
2424
uint begin_offset = 2;
2525
#if INPUT0_DIMS == 6
26-
const uint begin_w = begin[begin_offset];
26+
const int begin_w = begin[begin_offset];
2727
begin_offset += 1;
2828
#endif
2929
#if INPUT0_DIMS >= 5
30-
const uint begin_z = begin[begin_offset];
30+
const int begin_z = begin[begin_offset];
3131
begin_offset += 1;
3232
#endif
33-
const uint begin_y = begin[begin_offset];
34-
const uint begin_x = begin[begin_offset + 1];
33+
const int begin_y = begin[begin_offset];
34+
const int begin_x = begin[begin_offset + 1];
3535
#else
3636
const uint begin_b = LT_SIZES_BATCH_NUM;
3737
const uint begin_f = LT_SIZES_FEATURE_NUM;
@@ -46,19 +46,19 @@ KERNEL(border_gpu_ref)(
4646
#endif
4747

4848
#ifdef END_TYPE
49-
const uint end_b = end[0];
50-
const uint end_f = end[1];
49+
const int end_b = end[0];
50+
const int end_f = end[1];
5151
uint end_offset = 2;
5252
#if INPUT0_DIMS == 6
53-
const uint end_w = end[end_offset];
53+
const int end_w = end[end_offset];
5454
end_offset += 1;
5555
#endif
5656
#if INPUT0_DIMS >= 5
57-
const uint end_z = end[end_offset];
57+
const int end_z = end[end_offset];
5858
end_offset += 1;
5959
#endif
60-
const uint end_y = end[end_offset];
61-
const uint end_x = end[end_offset + 1];
60+
const int end_y = end[end_offset];
61+
const int end_x = end[end_offset + 1];
6262
#else
6363
const uint end_b = RB_SIZES_BATCH_NUM;
6464
const uint end_f = RB_SIZES_FEATURE_NUM;
@@ -74,65 +74,65 @@ KERNEL(border_gpu_ref)(
7474

7575
// [CONSTEXPR]
7676
// Border sizes(left-top):
77-
const uint blt_sb = begin_b;
78-
const uint blt_sf = begin_f;
79-
const uint blt_sy = begin_y;
80-
const uint blt_sx = begin_x;
77+
const int blt_sb = begin_b;
78+
const int blt_sf = begin_f;
79+
const int blt_sy = begin_y;
80+
const int blt_sx = begin_x;
8181
#if INPUT0_DIMS == 6
82-
const uint blt_sw = begin_w;
82+
const int blt_sw = begin_w;
8383
#else
84-
const uint blt_sw = 0;
84+
const int blt_sw = 0;
8585
#endif
8686
#if INPUT0_DIMS >= 5
87-
const uint blt_sz = begin_z;
87+
const int blt_sz = begin_z;
8888
#else
89-
const uint blt_sz = 0;
89+
const int blt_sz = 0;
9090
#endif
9191

9292
// Border sizes(right-bottom):
93-
const uint brb_sb = end_b;
94-
const uint brb_sf = end_f;
95-
const uint brb_sy = end_y;
96-
const uint brb_sx = end_x;
93+
const int brb_sb = end_b;
94+
const int brb_sf = end_f;
95+
const int brb_sy = end_y;
96+
const int brb_sx = end_x;
9797
#if INPUT0_DIMS == 6
98-
const uint brb_sw = end_w;
98+
const int brb_sw = end_w;
9999
#else
100-
const uint brb_sw = 0;
100+
const int brb_sw = 0;
101101
#endif
102102
#if INPUT0_DIMS >= 5
103-
const uint brb_sz = end_z;
103+
const int brb_sz = end_z;
104104
#else
105-
const uint brb_sz = 0;
105+
const int brb_sz = 0;
106106
#endif
107107

108108
// Input sizes:
109-
const uint in_sx = INPUT0_SIZE_X;
110-
const uint in_sy = INPUT0_SIZE_Y;
111-
const uint in_sz = INPUT0_SIZE_Z;
112-
const uint in_sw = INPUT0_SIZE_W;
113-
const uint in_sf = INPUT0_FEATURE_NUM;
114-
const uint in_sb = INPUT0_BATCH_NUM;
109+
const int in_sx = INPUT0_SIZE_X;
110+
const int in_sy = INPUT0_SIZE_Y;
111+
const int in_sz = INPUT0_SIZE_Z;
112+
const int in_sw = INPUT0_SIZE_W;
113+
const int in_sf = INPUT0_FEATURE_NUM;
114+
const int in_sb = INPUT0_BATCH_NUM;
115115

116116
// Input limits (exclusive; tested on output position):
117-
const uint in_lx = in_sx + blt_sx;
118-
const uint in_ly = in_sy + blt_sy;
119-
const uint in_lz = in_sz + blt_sz;
120-
const uint in_lw = in_sw + blt_sw;
121-
const uint in_lf = in_sf + blt_sf;
122-
const uint in_lb = in_sb + blt_sb;
117+
const int in_lx = in_sx + blt_sx;
118+
const int in_ly = in_sy + blt_sy;
119+
const int in_lz = in_sz + blt_sz;
120+
const int in_lw = in_sw + blt_sw;
121+
const int in_lf = in_sf + blt_sf;
122+
const int in_lb = in_sb + blt_sb;
123123

124-
const uint out_xz = (uint) get_global_id(0);
125-
const uint out_yw = (uint) get_global_id(1);
126-
const uint out_fb = (uint) get_global_id(2);
124+
const int out_xz = get_global_id(0);
125+
const int out_yw = get_global_id(1);
126+
const int out_fb = get_global_id(2);
127127

128-
const uint out_f = out_fb % OUTPUT_FEATURE_NUM;
129-
const uint out_b = out_fb / OUTPUT_FEATURE_NUM;
128+
const int out_f = out_fb % OUTPUT_FEATURE_NUM;
129+
const int out_b = out_fb / OUTPUT_FEATURE_NUM;
130130

131-
const uint out_x = out_xz % OUTPUT_SIZE_X;
132-
const uint out_z = out_xz / OUTPUT_SIZE_X;
131+
const int out_x = out_xz % OUTPUT_SIZE_X;
132+
const int out_z = out_xz / OUTPUT_SIZE_X;
133133

134-
const uint out_y = out_yw % OUTPUT_SIZE_Y;
135-
const uint out_w = out_yw / OUTPUT_SIZE_Y;
134+
const int out_y = out_yw % OUTPUT_SIZE_Y;
135+
const int out_w = out_yw / OUTPUT_SIZE_Y;
136136

137137
#ifdef BORDER_TYPE_CONSTANT
138138
#ifdef BORDER_VALUE_TYPE
@@ -148,14 +148,14 @@ KERNEL(border_gpu_ref)(
148148
out_f >= blt_sf & out_f < in_lf &
149149
out_b >= blt_sb & out_b < in_lb)
150150
{
151-
const uint in_x = out_x - blt_sx;
152-
const uint in_y = out_y - blt_sy;
153-
const uint in_z = out_z - blt_sz;
154-
const uint in_w = out_w - blt_sw;
155-
const uint in_f = out_f - blt_sf;
156-
const uint in_b = out_b - blt_sb;
157-
158-
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);
151+
const int in_x = out_x - blt_sx;
152+
const int in_y = out_y - blt_sy;
153+
const int in_z = out_z - blt_sz;
154+
const int in_w = out_w - blt_sw;
155+
const int in_f = out_f - blt_sf;
156+
const int in_b = out_b - blt_sb;
157+
158+
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);
159159
in_val = input[in_pos];
160160
}
161161
#elif defined BORDER_TYPE_EDGE
@@ -192,6 +192,6 @@ KERNEL(border_gpu_ref)(
192192
#error Unsupported border type.
193193
#endif
194194

195-
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);
195+
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);
196196
output[out_pos] = in_val;
197197
}

‎src/plugins/intel_gpu/src/kernel_selector/kernels/border/border_kernel_base.h

+6-3
Original file line numberDiff line numberDiff line change
@@ -12,16 +12,19 @@ namespace kernel_selector {
1212
// border_params
1313
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
1414
struct border_params : public base_params {
15-
DimTensor<> lt_sizes;
16-
DimTensor<> rb_sizes;
15+
DimTensor<int32_t> lt_sizes;
16+
DimTensor<int32_t> rb_sizes;
1717
BorderType b_type;
1818
float border_value;
19+
bool allow_negative_pad;
1920

2021
ArgType begin_type;
2122
ArgType end_type;
2223
ArgType pad_value_type;
2324

24-
border_params() : base_params(KernelType::BORDER), b_type(BorderType::CONSTANT), border_value(0.0f),
25+
26+
border_params() : base_params(KernelType::BORDER), b_type(BorderType::CONSTANT),
27+
border_value(0.0f), allow_negative_pad(false),
2528
begin_type(ArgType::Constant), end_type(ArgType::Constant), pad_value_type(ArgType::Constant) {}
2629

2730
ParamsKey GetParamsKey() const override {

0 commit comments

Comments
 (0)
Please sign in to comment.