From e1bd2a7b55c1a6aef01703baf000acc41c576ec0 Mon Sep 17 00:00:00 2001 From: Wilson Seok Date: Wed, 27 Dec 2023 23:31:41 +0900 Subject: [PATCH] Revert "[GPU] Fix accuracy problems brought by permute_f_y kernel (#20241)" (#21830) (#21876) This reverts commit 6719133af29d55477b5a1a0a271d177fecc5934f. Co-authored-by: Pavel Durandin --- .../cl_kernels/permute_f_y_axes.cl | 43 ++++++------ .../permute/permute_kernel_f_y_axes.cpp | 69 ++++++++----------- .../kernels/permute/permute_kernel_f_y_axes.h | 6 +- .../permute/permute_kernel_selector.cpp | 2 +- .../single_layer_tests/transpose.cpp | 16 ----- .../unit/fusions/permute_fusion_test.cpp | 8 +-- 6 files changed, 52 insertions(+), 92 deletions(-) diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/permute_f_y_axes.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/permute_f_y_axes.cl index a76f021ae0e..3f3e0a3f8c9 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/permute_f_y_axes.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/permute_f_y_axes.cl @@ -25,22 +25,21 @@ KERNEL (permute_f_y_axes)( #endif ) { - const int f_idx = get_global_id(2); - const int b_idx = get_global_id(1);; - const int y_idx = get_global_id(0); + const int bf = get_global_id(2); + const int f_idx = bf % INPUT0_FEATURE_NUM; + const int b_idx = bf / INPUT0_FEATURE_NUM; + const int x_start = get_global_id(0) * BLOCK_SIZE; + const int y_idx = get_global_id(1); - for (int x_idx = 0; x_idx < J_TIMES * VEC_SIZE; x_idx+=VEC_SIZE) { - IN_VEC_TYPE resv = READ_VEC(0, &input[INPUT0_GET_INDEX(b_idx, f_idx, y_idx, x_idx)]); + __attribute__((opencl_unroll_hint(J_TIMES))) + for (int j = 0; j < J_TIMES; ++j) { + const int x_idx = x_start + j * VEC_SIZE; + IN_VEC_TYPE res = READ_VEC(0, &input[INPUT0_GET_INDEX(b_idx, f_idx, y_idx, x_idx)]); #if HAS_FUSED_OPS - OUT_VEC_TYPE result; - for (int i = 0; i < VEC_SIZE; i++) - { - INPUT0_TYPE res = resv[i]; - FUSED_OPS; - result[i] = FUSED_OPS_RESULT; - } + FUSED_OPS_VEC; + OUT_VEC_TYPE result = FUSED_OPS_RESULT_VEC; #else - OUT_VEC_TYPE result = ACTIVATION(resv, ACTIVATION_PARAMS); + OUT_VEC_TYPE result = ACTIVATION(res, ACTIVATION_PARAMS); #endif const int output_idx = OUTPUT_GET_INDEX(b_idx, y_idx, f_idx, x_idx); WRITE_VEC(result, 0, &output[output_idx]); @@ -133,16 +132,18 @@ KERNEL (permute_f_y_axes)( for (int j = 0; j < J_TIMES; ++j) { const int j_vec = j * VEC_SIZE; const int y_idx = y_begin + j_vec; - IN_VEC_TYPE resv = READ_VEC(0, &input[INPUT0_GET_INDEX(b_idx, f_idx, y_idx, x_idx)]); - for (int k = 0; k < VEC_SIZE; ++k) { + IN_VEC_TYPE res = READ_VEC(0, &input[INPUT0_GET_INDEX(b_idx, f_idx, y_idx, x_idx)]); #if HAS_FUSED_OPS - INPUT0_TYPE res = resv[k]; - FUSED_OPS - transpose_buf[j_vec + k][bf_local] = FUSED_OPS_RESULT; -#else - transpose_buf[j_vec + k][bf_local] = ACTIVATION(resv[k], ACTIVATION_PARAMS); -#endif + FUSED_OPS_VEC; + OUT_VEC_TYPE result = FUSED_OPS_RESULT_VEC; + for (int k = 0; k < VEC_SIZE; ++k) { + transpose_buf[j_vec + k][bf_local] = result[k]; } +#else + for (int k = 0; k < VEC_SIZE; ++k) { + transpose_buf[j_vec + k][bf_local] = ACTIVATION(res[k], ACTIVATION_PARAMS); + } +#endif } __attribute__((opencl_unroll_hint(J_TIMES))) diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_f_y_axes.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_f_y_axes.cpp index 08d6ed89a78..e6fc63ac5d7 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_f_y_axes.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_f_y_axes.cpp @@ -128,46 +128,29 @@ JitConstants PermuteKernel_f_y_axes::GetJitConstants(const permute_params& param } } + const size_t tile_width = GetTileWidth(params); + const size_t vector_size = std::min(tile_width, static_cast(4)); + const size_t tile_size = GetTileSize(params); + const size_t j_times = tile_size / vector_size; + const size_t feature_block_size = GetFeatureBlockSize(params); + jit.AddConstant(MakeJitConstant("BLOCK_SIZE", tile_width)); + jit.AddConstant(MakeJitConstant("VEC_SIZE", vector_size)); + jit.AddConstant(MakeJitConstant("J_TIMES", j_times)); + jit.AddConstant(MakeJitConstant("TILE_SIZE", tile_size)); + jit.AddConstant(MakeJitConstant("FEATURE_BLOCK_SIZE", feature_block_size)); + const auto layout = params.inputs.front().GetLayout(); - std::size_t vector_size; - if (IsSimpleMemCopyOperation(params)) { - vector_size = GetDivisor(params.inputs[0].X().v); - auto j_times = params.inputs[0].X().v / vector_size; - jit.AddConstant(MakeJitConstant("VEC_SIZE", vector_size)); - jit.AddConstant(MakeJitConstant("J_TIMES", j_times)); - } else { - const size_t tile_width = GetTileWidth(params); - vector_size = std::min(tile_width, static_cast(4)); - const size_t tile_size = GetTileSize(params); - const size_t j_times = tile_size / vector_size; - const size_t feature_block_size = GetFeatureBlockSize(params); - jit.AddConstant(MakeJitConstant("VEC_SIZE", vector_size)); - jit.AddConstant(MakeJitConstant("J_TIMES", j_times)); - jit.AddConstant(MakeJitConstant("TILE_SIZE", tile_size)); - jit.AddConstant(MakeJitConstant("FEATURE_BLOCK_SIZE", feature_block_size)); - if (!SimpleLayout(layout)) { - const auto subgroup_size = Is3DTranspose(params) ? feature_block_size : tile_size; - jit.AddConstant(MakeJitConstant("SUB_GROUP_SIZE", subgroup_size)); - } + if (!SimpleLayout(layout)) { + const auto subgroup_size = Is3DTranspose(params) ? feature_block_size : tile_size; + jit.AddConstant(MakeJitConstant("SUB_GROUP_SIZE", subgroup_size)); } if (!params.fused_ops.empty()) { - if (IsSimpleMemCopyOperation(params)) { - const std::vector original_output_order = {"b_idx", "y_idx", "f_idx", "(x_idx+i)"}; - const FusedOpsConfiguration conf_scalar = {"", original_output_order, "res", params.inputs[0].GetDType(), 1}; - jit.Merge(MakeFusedOpsJitConstants(params, {conf_scalar})); - } else if (SimpleLayout(layout)) { - const std::vector original_output_order = {"b_idx", "(y_idx+k)", "f_idx", "x_idx"}; - const FusedOpsConfiguration conf_scalar = {"", original_output_order, "res", params.inputs[0].GetDType(), 1}; - jit.Merge(MakeFusedOpsJitConstants(params, {conf_scalar})); - } else { - const std::vector original_output_order = {"b_idx", "y_idx", "f_idx", "x_idx"}; - const FusedOpsConfiguration conf_scalar = {"", original_output_order, "res", params.inputs[0].GetDType(), 1}; - const FusedOpsConfiguration conf_vec = {"_VEC", original_output_order, "res", params.inputs[0].GetDType(), vector_size}; - jit.Merge(MakeFusedOpsJitConstants(params, {conf_scalar, conf_vec})); - } + const std::vector original_output_order = {"b_idx", "f_idx", "y_idx", "x_idx"}; + const FusedOpsConfiguration conf_scalar = {"", original_output_order, "res", params.inputs[0].GetDType(), 1}; + const FusedOpsConfiguration conf_vec = {"_VEC", original_output_order, "res", params.inputs[0].GetDType(), vector_size}; + jit.Merge(MakeFusedOpsJitConstants(params, {conf_scalar, conf_vec})); } - return jit; } @@ -175,12 +158,14 @@ static inline std::vector GetGWS(const permute_params& params) { const auto& in = params.inputs[0]; std::vector gws; auto block_size = IsSimpleMemCopyOperation(params) ? GetTileWidth(params) : GetTileSize(params); - if (IsSimpleMemCopyOperation(params)) { - return {in.Y().v, in.Batch().v, in.Feature().v}; - } else if (Is3DTranspose(params)) { - gws = {in.X().v / block_size, in.Y().v / GetFeatureBlockSize(params), (in.Batch().v * in.Feature().v)}; - } else { // the case where x is 1 + if (params.inputs[0].X().v == 1) { gws = {in.X().v, in.Y().v / block_size, (in.Batch().v * in.Feature().v)}; + } else { + if (Is3DTranspose(params)) { + gws = {in.X().v / block_size, in.Y().v / GetFeatureBlockSize(params), (in.Batch().v * in.Feature().v)}; + } else { + gws = {in.X().v / block_size, in.Y().v, (in.Batch().v * in.Feature().v)}; + } } return gws; } @@ -192,9 +177,9 @@ CommonDispatchData PermuteKernel_f_y_axes::SetDefault(const permute_params& para auto in_layout = params.inputs[0].GetLayout(); auto out_layout = params.outputs[0].GetLayout(); const std::vector> dims_by_gws = { + {Tensor::DataChannelName::X}, {Tensor::DataChannelName::Y}, - {Tensor::DataChannelName::BATCH}, - {Tensor::DataChannelName::FEATURE}}; + {Tensor::DataChannelName::FEATURE, Tensor::DataChannelName::BATCH}}; dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo, in_layout, out_layout, dims_by_gws); } else if (Is3DTranspose(params)) { diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_f_y_axes.h b/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_f_y_axes.h index 615b1ce0d34..eeeb27fd321 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_f_y_axes.h +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_f_y_axes.h @@ -23,11 +23,7 @@ protected: JitConstants GetJitConstants(const permute_params& params, const CommonDispatchData& dispatchData) const override; CommonDispatchData SetDefault(const permute_params& params) const override; std::vector GetSupportedFusedOps() const override { - return { - FusedOpType::ACTIVATION, - FusedOpType::QUANTIZE, - FusedOpType::ELTWISE - }; + return {FusedOpType::ACTIVATION, FusedOpType::QUANTIZE, FusedOpType::ELTWISE}; } }; } // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_selector.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_selector.cpp index 64e825b23b6..1e37d5f22ae 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_selector.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_selector.cpp @@ -16,7 +16,7 @@ permute_kernel_selector::permute_kernel_selector() { Attach(); Attach(); Attach(); - Attach(); + // Attach(); } KernelsData permute_kernel_selector::GetBestKernels(const Params& params, const optional_params& options) const { diff --git a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/transpose.cpp b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/transpose.cpp index fd344effe2c..981e77d2f5f 100644 --- a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/transpose.cpp +++ b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/transpose.cpp @@ -40,22 +40,6 @@ INSTANTIATE_TEST_SUITE_P(smoke_Transpose, testing::Values(ov::test::utils::DEVICE_GPU)), TransposeLayerTest::getTestCaseName); -const std::vector> levitShapes = { - {{1, 4, 196, 32}}, - {{1, 196, 4, 32}}, - {{1, 196, 4, 16}}, - {{1, 196, 8, 64}}, - {{1, 24, 196, 1}}, -}; - -INSTANTIATE_TEST_SUITE_P(smoke_Transpose_Levit, - TransposeLayerTest, - testing::Combine(testing::Values(std::vector{0, 2, 1, 3}), - testing::Values(ov::element::f16), - testing::ValuesIn(ov::test::static_shapes_to_test_representation(levitShapes)), - testing::Values(ov::test::utils::DEVICE_GPU)), - TransposeLayerTest::getTestCaseName); - /** * 5D permute tests */ diff --git a/src/plugins/intel_gpu/tests/unit/fusions/permute_fusion_test.cpp b/src/plugins/intel_gpu/tests/unit/fusions/permute_fusion_test.cpp index 32306ec3670..aa7919eb7b6 100644 --- a/src/plugins/intel_gpu/tests/unit/fusions/permute_fusion_test.cpp +++ b/src/plugins/intel_gpu/tests/unit/fusions/permute_fusion_test.cpp @@ -171,10 +171,7 @@ public: #define CASE_PERMUTE_TILE_BFZYX_TO_BFYXZ_0 { 1, 8, 8, 2, 2 }, { 1, 8, 2, 8, 2 }, { 0, 1, 3, 4, 2 }, tensor{ 0 }, data_types::f32, format::bfzyx, data_types::f32, format::bfzyx // permute_f_y_axes -#define CASE_PERMUTE_TILE_BFYX_TO_BYFX_0 { 1, 8, 4, 16 }, { 1, 16, 4, 8 }, { 0, 2, 1, 3 }, tensor{ 0 }, data_types::f32, format::bfyx, data_types::f32, format::bfyx -#define CASE_PERMUTE_TILE_BFYX_TO_BYFX_1 { 1, 8, 1, 16 }, { 1, 16, 1, 8 }, { 0, 2, 1, 3 }, tensor{ 0 }, data_types::f32, format::bfyx, data_types::f32, format::bfyx -#define CASE_PERMUTE_TILE_BFYX_TO_BYFX_2 { 1, 8, 4, 16 }, { 1, 16, 4, 8 }, { 0, 2, 1, 3 }, tensor{ 0 }, data_types::f32, format::b_fs_yx_fsv16, data_types::f32, format::b_fs_yx_fsv16 -#define CASE_PERMUTE_TILE_BFYX_TO_BYFX_3 { 1, 8, 1, 16 }, { 1, 16, 1, 8 }, { 0, 2, 1, 3 }, tensor{ 0 }, data_types::f32, format::b_fs_yx_fsv16, data_types::f32, format::b_fs_yx_fsv16 +#define CASE_PERMUTE_TILE_BFYX_TO_BYFX_0 { 1, 8, 4, 2 }, { 1, 2, 4, 8 }, { 0, 2, 1, 3 }, tensor{ 0 }, data_types::f32, format::bfyx, data_types::f32, format::bfyx class permute_activation_scale_eltwise: public PermuteFusingTest {}; TEST_P(permute_activation_scale_eltwise, basic) { @@ -274,9 +271,6 @@ INSTANTIATE_TEST_SUITE_P(fusings_gpu, permute_activation_scale_eltwise, ::testin // Fusing tests for permute_f_y_axes permute_params{ CASE_PERMUTE_TILE_BFYX_TO_BYFX_0, 2, 5 }, - permute_params{ CASE_PERMUTE_TILE_BFYX_TO_BYFX_1, 2, 5 }, - permute_params{ CASE_PERMUTE_TILE_BFYX_TO_BYFX_2, 2, 5 }, - permute_params{ CASE_PERMUTE_TILE_BFYX_TO_BYFX_3, 2, 5 }, })); class permute_quant_u8: public PermuteFusingTest {};