From b33c94205244ee633380bc77d8856dc95f972912 Mon Sep 17 00:00:00 2001 From: Andrei Gorbachev Date: Thu, 4 Apr 2024 12:40:56 +0100 Subject: [PATCH] [GPU] Strideslice revert (#23845) ### Details: - This PR reverts: https://github.com/openvinotoolkit/openvino/pull/23659 Tickets: 137330, 137713 --- .../cl_kernels/strided_slice_ref.cl | 59 +++++-------------- .../skip_tests_config.cpp | 2 + 2 files changed, 17 insertions(+), 44 deletions(-) diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/strided_slice_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/strided_slice_ref.cl index a768aa426eb918..eff062e449f097 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/strided_slice_ref.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/strided_slice_ref.cl @@ -289,62 +289,33 @@ KERNEL(strided_slice_ref)(OPTIONAL_SHAPE_INFO_ARG #if NEW_AXIS_MODE // If NEW_AXIS_MODE that just copy input to output -#ifdef INPUT0_LAYOUT_BFYX - const uint index_in_batch = (feature * (uint)get_global_size(2) + (uint)get_global_id(2)) % (OUTPUT_SIZE_X * OUTPUT_SIZE_Y); - const uint input_feature_id = (feature * (uint)get_global_size(2) + (uint)get_global_id(2)) / (OUTPUT_SIZE_X * OUTPUT_SIZE_Y); +#ifdef OUTPUT_LAYOUT_BFYX const uint w_input = 0; const uint z_input = 0; - const uint y_input = index_in_batch / OUTPUT_SIZE_X; - const uint x_input = index_in_batch % OUTPUT_SIZE_X; -#elif INPUT0_LAYOUT_BFZYX - const uint index_in_batch = (feature * (uint)get_global_size(2) + (uint)get_global_id(2)) % (OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z); - const uint input_feature_id = (feature * (uint)get_global_size(2) + (uint)get_global_id(2)) / (OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z); + const uint y_input = (uint)get_global_id(2) / INPUT0_SIZE_X; + const uint x_input = (uint)get_global_id(2) % INPUT0_SIZE_X; +#elif OUTPUT_LAYOUT_BFZYX const uint w_input = 0; - const uint yx_input = index_in_batch % (INPUT0_SIZE_X * INPUT0_SIZE_Y); - const uint z_input = index_in_batch / (INPUT0_SIZE_X * INPUT0_SIZE_Y); + const uint yx_input = (uint)get_global_id(2) % (INPUT0_SIZE_X * INPUT0_SIZE_Y); + const uint z_input = (uint)get_global_id(2) / (INPUT0_SIZE_X * INPUT0_SIZE_Y); const uint y_input = yx_input / INPUT0_SIZE_X; const uint x_input = yx_input % INPUT0_SIZE_X; -#elif INPUT0_LAYOUT_BFWZYX - const uint index_in_batch = (feature * (uint)get_global_size(2) + (uint)get_global_id(2)) % (OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z * OUTPUT_SIZE_W); - const uint input_feature_id = (feature * (uint)get_global_size(2) + (uint)get_global_id(2)) / (OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z * OUTPUT_SIZE_W); - const uint zyx_input = index_in_batch % (INPUT0_SIZE_X * INPUT0_SIZE_Y * INPUT0_SIZE_Z); - const uint w_input = index_in_batch / (INPUT0_SIZE_X * INPUT0_SIZE_Y * INPUT0_SIZE_Z); +#elif OUTPUT_LAYOUT_BFWZYX + const uint zyx_input = (uint)get_global_id(2) % (INPUT0_SIZE_X * INPUT0_SIZE_Y * INPUT0_SIZE_Z); + const uint w_input = (uint)get_global_id(2) / (INPUT0_SIZE_X * INPUT0_SIZE_Y * INPUT0_SIZE_Z); const uint z_input = zyx_input / (INPUT0_SIZE_X * INPUT0_SIZE_Y); const uint yx_input = zyx_input % (INPUT0_SIZE_X * INPUT0_SIZE_Y); const uint y_input = yx_input / INPUT0_SIZE_X; const uint x_input = yx_input % INPUT0_SIZE_X; #endif - const uint input_index = INPUT0_OFFSET + batch * INPUT0_BATCH_PITCH + - input_feature_id * INPUT0_FEATURE_PITCH + - w_input * OUTPUT_W_PITCH + - z_input * OUTPUT_Z_PITCH + - y_input * OUTPUT_Y_PITCH + - x_input * OUTPUT_X_PITCH; - -#ifdef OUTPUT_LAYOUT_BFYX - const uint y = (uint)get_global_id(2) / OUTPUT_SIZE_X; - const uint x = (uint)get_global_id(2) % OUTPUT_SIZE_X; - const uint output_index = OUTPUT_GET_INDEX(batch, feature, y, x); -#elif OUTPUT_LAYOUT_BFZYX - const uint yx = (uint)get_global_id(2) % (OUTPUT_SIZE_X * OUTPUT_SIZE_Y); - const uint z = (uint)get_global_id(2) / (OUTPUT_SIZE_X * OUTPUT_SIZE_Y); - const uint y = yx / OUTPUT_SIZE_X; - const uint x = yx % OUTPUT_SIZE_X; - const uint output_index = OUTPUT_GET_INDEX(batch, feature, z, y, x); -#elif OUTPUT_LAYOUT_BFWZYX - const uint zyx = (uint)get_global_id(2) % (OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z); - const uint w = (uint)get_global_id(2) / (OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z); - const uint z = zyx / (OUTPUT_SIZE_X * OUTPUT_SIZE_Y); - const uint yx = zyx % (OUTPUT_SIZE_X * OUTPUT_SIZE_Y); - const uint y = yx / OUTPUT_SIZE_X; - const uint x = yx % OUTPUT_SIZE_X; - const uint output_index = OUTPUT_GET_INDEX(batch, feature, w, z, y, x); -#endif - - output[output_index] = input[input_index]; - + feature * INPUT0_FEATURE_PITCH + + w_input * INPUT0_W_PITCH + + z_input * INPUT0_Z_PITCH + + y_input * INPUT0_Y_PITCH + + x_input * INPUT0_X_PITCH; + output[input_index] = input[input_index]; #else // NEW_AXIS_MODE #ifdef OUTPUT_LAYOUT_BFYX const uint w = 0; diff --git a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/skip_tests_config.cpp b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/skip_tests_config.cpp index d705eb52414b13..516e1a54b6792e 100644 --- a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/skip_tests_config.cpp +++ b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/skip_tests_config.cpp @@ -205,6 +205,8 @@ std::vector disabledTestPatterns() { R"(.*smoke_ConvolutionBackpropData3D_AutoPadding_OutputPaddingDefined/ConvolutionBackpropDataLayerTest.Inference/.*_TS=\{\(1.16.5.5.5\)\}_.*_netPRC=f16_.*)", R"(.*smoke_PSROIPooling_average/PSROIPoolingLayerTest.Inference/IS=\(3.8.16.16\)_coord_shape=\(10.5\)_out_dim=2_group_size=2_scale=(0.625|1)_bins_x=1_bins_y=1_mode=average_modelType=f16.*)", R"(.*smoke_RDFT_5d_last_axis/RDFTLayerTest.Inference/IS=\(10.4.8.2.5\)_modelType=f32_Axes=\(0.1.2.3.4\)_SignalSize=\(\).*)", + // Issue: 129991 + R"(.*StridedSliceLayerTest.*TS=.*2.2.4.1*.*)", // Issue: 136862 R"(.*smoke_ConditionGPUTest_static/StaticConditionLayerGPUTest.CompareWithRefs/IS=\(3.6\)_netPRC=i8_ifCond=PARAM_targetDevice=GPU_.*)", #if defined(_WIN32)