Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[GPU] Pad-12 #19083

Merged
merged 4 commits into from
Sep 12, 2023
Merged

[GPU] Pad-12 #19083

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
bea7f97
GPU primitive and kernel changes to support Pad-12
okhovan Aug 8, 2023
d37e254
Exclude Pad-12 from GPU transformations pipeline
okhovan Aug 8, 2023
f4e60f5
add unit tests
okhovan Aug 8, 2023
4aae63c
add single-layet test for Pad-12
okhovan Aug 8, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
23 changes: 16 additions & 7 deletions src/plugins/intel_gpu/include/intel_gpu/primitives/border.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<border> {
CLDNN_DECLARE_PRIMITIVE(border)

Expand All @@ -40,12 +40,13 @@ struct border : public primitive_base<border> {
/// @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<input_info>& inputs,
Expand All @@ -54,12 +55,14 @@ struct border : public primitive_base<border> {
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).
Expand All @@ -69,7 +72,9 @@ struct border : public primitive_base<border> {
/// @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;

Expand All @@ -79,6 +84,7 @@ struct border : public primitive_base<border> {
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;
}
Expand All @@ -92,7 +98,8 @@ struct border : public primitive_base<border> {
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 {
Expand All @@ -102,6 +109,7 @@ struct border : public primitive_base<border> {
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 {
Expand All @@ -111,6 +119,7 @@ struct border : public primitive_base<border> {
ib >> make_data(&pad_mode, sizeof(ov::op::PadMode));
ib >> pad_value;
ib >> non_constant_input_mask;
ib >> allow_negative_pad;
}
};
} // namespace cldnn
36 changes: 19 additions & 17 deletions src/plugins/intel_gpu/src/graph/border.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand All @@ -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;

// 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 pad_mode = argument->pad_mode;
const bool allow_negative_pad = argument->allow_negative_pad;

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;
Expand Down
6 changes: 4 additions & 2 deletions src/plugins/intel_gpu/src/graph/impls/ocl/border.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,7 +46,7 @@ struct border_impl : typed_primitive_impl_ocl<border> {
begin_vec.insert(begin_vec.end(), zeros_to_add, 0);
}
std::vector<tensor::value_type> 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<int32_t>(tensor(pads_format, pads_begin, 0));
} else {
params.begin_type = kernel_selector::base_params::ArgType::Input;

Expand All @@ -65,7 +65,7 @@ struct border_impl : typed_primitive_impl_ocl<border> {
end_vec.insert(end_vec.end(), zeros_to_add, 0);
}
std::vector<tensor::value_type> 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<int32_t>(tensor(pads_format, pads_end, 0));
} else {
params.end_type = kernel_selector::base_params::ArgType::Input;

Expand Down Expand Up @@ -100,6 +100,8 @@ struct border_impl : typed_primitive_impl_ocl<border> {
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};
}

Expand Down
116 changes: 58 additions & 58 deletions src/plugins/intel_gpu/src/kernel_selector/cl_kernels/border_gpu_ref.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -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;
Expand All @@ -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
Expand All @@ -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 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_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 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
Expand Down Expand Up @@ -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;
}
9 changes: 6 additions & 3 deletions src/plugins/intel_gpu/src/kernel_selector/kernels/border/border_kernel_base.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,16 +12,19 @@ namespace kernel_selector {
// border_params
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
struct border_params : public base_params {
DimTensor<> lt_sizes;
DimTensor<> rb_sizes;
DimTensor<int32_t> lt_sizes;
DimTensor<int32_t> 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 {
Expand Down
Loading