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

[GPU] Enable Select 5d (#24544) 

### Details:
 - *Enable select 5d in cl kernel*
 - *Added basic unit test for select 5d*

### Tickets:
 - *140122*
This commit is contained in:
Steve Yoo 2024-05-28 10:11:23 +09:00 committed by GitHub
parent 4cf2ae0192
commit e6496dbbb8
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
5 changed files with 56 additions and 12 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
src/plugins/intel_gpu/src/graph/impls/ocl/select.cpp
View File

@ -80,6 +80,7 @@ attach_select_impl::attach_select_impl() {
format::bfyx,
format::byxf,
format::yxfb,
format::bfzyx,
};
implementation_map<select>::add(impl_types::ocl,
@ -89,7 +90,8 @@ attach_select_impl::attach_select_impl() {
static_formats);
auto dyn_formats = {
format::bfyx
format::bfyx,
format::bfzyx,
};
implementation_map<select>::add(impl_types::ocl,

33
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/select_gpu_ref.cl
View File

@ -4,23 +4,46 @@
#include "include/batch_headers/fetch_data.cl"
#define INPUT_0 input0[INPUT0_GET_INDEX_SAFE(b, f, y, x)]
#define INPUT_1 input1[INPUT1_GET_INDEX_SAFE(b, f, y, x)]
#define INPUT_2 input2[INPUT2_GET_INDEX_SAFE(b, f, y, x)]
#if OUTPUT_DIMS == 5
#define INPUT_0 input0[INPUT0_GET_INDEX_SAFE(b, f, z, y, x)]
#if INPUT1_DIMS == 4
#define INPUT_1 input1[INPUT1_GET_INDEX_SAFE(b, f, y, x)]
#else
#define INPUT_1 input1[INPUT1_GET_INDEX_SAFE(b, f, z, y, x)]
#endif
#if INPUT2_DIMS == 4
#define INPUT_2 input2[INPUT2_GET_INDEX_SAFE(b, f, y, x)]
#else
#define INPUT_2 input2[INPUT2_GET_INDEX_SAFE(b, f, z, y, x)]
#endif
#elif OUTPUT_DIMS == 4
#define INPUT_0 input0[INPUT0_GET_INDEX_SAFE(b, f, y, x)]
#define INPUT_1 input1[INPUT1_GET_INDEX_SAFE(b, f, y, x)]
#define INPUT_2 input2[INPUT2_GET_INDEX_SAFE(b, f, y, x)]
#endif
KERNEL(select)(
OPTIONAL_SHAPE_INFO_ARG
INPUTS_DECLS
__global OUTPUT_TYPE* output)
{
const uint x = (uint)get_global_id(0);
const uint x = (uint)get_global_id(0);
#if OUTPUT_DIMS == 5
const uint yz = (uint)get_global_id(1);
const uint y = yz % OUTPUT_SIZE_Y;
const uint z = yz / OUTPUT_SIZE_Y;
#elif OUTPUT_DIMS == 4
const uint y = (uint)get_global_id(1);
#endif
const uint bf = (uint)get_global_id(2);
const uint b = bf % OUTPUT_BATCH_NUM;
const uint f = bf / OUTPUT_BATCH_NUM;
#if OUTPUT_DIMS == 5
uint output_offset = OUTPUT_GET_INDEX(b, f, z, y, x);
#elif OUTPUT_DIMS == 4
uint output_offset = OUTPUT_GET_INDEX(b, f, y, x);
#endif
const OUTPUT_TYPE res = select(INPUT_2, INPUT_1, MASK);

22
src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_base.cpp
View File

@ -97,12 +97,26 @@ SelectKernelBase::DispatchData SelectKernelBase::SetDefault(const select_params&
DispatchData dispatchData;
const auto& out = params.outputs[0];
const auto& in = params.inputs[0];
std::vector<std::vector<Tensor::DataChannelName>> dims_by_gws;
dispatchData.gws = { out.X().v, out.Y().v, out.Feature().v * out.Batch().v };
switch (out.Dimentions()) {
case 4:
dispatchData.gws = { out.X().v, out.Y().v, out.Feature().v * out.Batch().v };
std::vector<std::vector<Tensor::DataChannelName>> dims_by_gws = {{ Tensor::DataChannelName::X },
{ Tensor::DataChannelName::Y },
{ Tensor::DataChannelName::FEATURE, Tensor::DataChannelName::BATCH }};
dims_by_gws = {{ Tensor::DataChannelName::X },
{ Tensor::DataChannelName::Y },
{ Tensor::DataChannelName::FEATURE, Tensor::DataChannelName::BATCH }};
break;
case 5:
dispatchData.gws = { out.X().v, out.Y().v * out.Z().v, out.Feature().v * out.Batch().v };
dims_by_gws = {{ Tensor::DataChannelName::X },
{ Tensor::DataChannelName::Y, Tensor::DataChannelName::Z },
{ Tensor::DataChannelName::FEATURE, Tensor::DataChannelName::BATCH }};
break;
default:
throw std::invalid_argument("Unsupported data layout for select primitive");
}
dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo, in.GetLayout(), out.GetLayout(), dims_by_gws);

2
src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_ref.cpp
View File

@ -27,10 +27,12 @@ ParamsKey SelectKernelRef::GetSupportedKey() const {
k.EnableInputLayout(DataLayout::bfyx);
k.EnableInputLayout(DataLayout::yxfb);
k.EnableInputLayout(DataLayout::byxf);
k.EnableInputLayout(DataLayout::bfzyx);
k.EnableOutputLayout(DataLayout::bfyx);
k.EnableOutputLayout(DataLayout::yxfb);
k.EnableOutputLayout(DataLayout::byxf);
k.EnableOutputLayout(DataLayout::bfzyx);
k.EnableBatching();
k.EnableTensorPitches();

7
src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/select.cpp
View File

@ -19,7 +19,8 @@ const std::vector<std::vector<ov::Shape>> noneShapes = {
{{8}, {8}, {8}},
{{4, 5}, {4, 5}, {4, 5}},
{{3, 4, 5}, {3, 4, 5}, {3, 4, 5}},
{{2, 3, 4, 5}, {2, 3, 4, 5}, {2, 3, 4, 5}}
{{2, 3, 4, 5}, {2, 3, 4, 5}, {2, 3, 4, 5}},
{{2, 2, 2, 2, 2}, {2, 2, 2, 2, 2}, {2, 2, 2, 2, 2}},
};
const std::vector<std::vector<ov::Shape>> numpyShapes = {
@ -44,7 +45,9 @@ const std::vector<std::vector<ov::Shape>> numpyShapes = {
{{5, 1, 8}, {2, 1, 9, 8}, {2, 5, 9, 8}},
{{6, 1, 1, 8}, {6, 7, 1, 8}, {2, 1}},
{{5, 1, 1, 1}, {5, 7, 8, 6}, {1, 8, 6}},
{{1, 1, 3}, {1, 3, 1}, {3, 1, 1}}
{{1, 1, 3}, {1, 3, 1}, {3, 1, 1}},
{{2, 2, 2, 2, 2}, {2, 2, 2, 2}, {2, 2, 2, 2, 2}},
{{2, 2, 2, 2, 2}, {2, 2, 2, 2, 2}, {2, 2, 2, 2}}
};
INSTANTIATE_TEST_SUITE_P(smoke_CLDNN_TestsSelect_none,