From e6496dbbb8e8548cb1e3e1e7ddc0cbac29527c8e Mon Sep 17 00:00:00 2001 From: Steve Yoo Date: Tue, 28 May 2024 10:11:23 +0900 Subject: [PATCH] [GPU] Enable Select 5d (#24544) ### Details: - *Enable select 5d in cl kernel* - *Added basic unit test for select 5d* ### Tickets: - *140122* --- .../intel_gpu/src/graph/impls/ocl/select.cpp | 4 ++- .../cl_kernels/select_gpu_ref.cl | 33 ++++++++++++++++--- .../kernels/select/select_kernel_base.cpp | 22 ++++++++++--- .../kernels/select/select_kernel_ref.cpp | 2 ++ .../single_layer_tests/select.cpp | 7 ++-- 5 files changed, 56 insertions(+), 12 deletions(-) diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/select.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/select.cpp index b96b50008f7..90a36900671 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/select.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/select.cpp @@ -80,6 +80,7 @@ attach_select_impl::attach_select_impl() { format::bfyx, format::byxf, format::yxfb, + format::bfzyx, }; implementation_map::add(impl_types::ocl, diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/select_gpu_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/select_gpu_ref.cl index 68171f61bc0..d56b5299ac2 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/select_gpu_ref.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/select_gpu_ref.cl @@ -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); diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_base.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_base.cpp index cd4281d8059..30ff171c025 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_base.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_base.cpp @@ -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> 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> 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); diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_ref.cpp index f5526e37bed..eca8964f980 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_ref.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/select/select_kernel_ref.cpp @@ -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(); diff --git a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/select.cpp b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/select.cpp index 48451217b1e..ecd58b840bf 100644 --- a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/select.cpp +++ b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/select.cpp @@ -19,7 +19,8 @@ const std::vector> 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> numpyShapes = { @@ -44,7 +45,9 @@ const std::vector> 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,