Skip to content

Commit

Permalink
[GPU] Strideslice revert (openvinotoolkit#23845)
Browse files Browse the repository at this point in the history 
### Details:
- This PR reverts:
openvinotoolkit#23659

Tickets:
137330, 137713
  • Loading branch information
@andrei-cv @bbielawx
andrei-cv authored and bbielawx committed Apr 12, 2024
1 parent 11f42c3 commit b33c942
Show file tree
Hide file tree
Showing 2 changed files with 17 additions and 44 deletions.
59 changes: 15 additions & 44 deletions src/plugins/intel_gpu/src/kernel_selector/cl_kernels/strided_slice_ref.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 2 additions & 0 deletions src/plugins/intel_gpu/tests/functional/shared_tests_instances/skip_tests_config.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -205,6 +205,8 @@ std::vector<std::string> 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)
Expand Down

0 comments on commit b33c942

Please sign in to comment.