[GPU] Bugfix reorder for byfx format (#25782)

+ Reorder returns OOR error while handling byfx from a fused permute
parent

### Details:
 - *item1*
 - *...*

### Tickets:
 - CVS-147330

---------

Signed-off-by: Min, Byung-il <byungil.min@intel.com>

Author: byungilm

src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data.cl

@@ -27,8 +27,15 @@ KERNEL (reorder_data)(
 #endif
     )
 {
+#if INPUT0_LAYOUT_BYFX
+    // GWS_FEATURE takes Y for byfx format
+    const uint b = get_global_id(GWS_BATCH);
+    const uint y = get_global_id(GWS_FEATURE);
+#else
     const uint b = get_global_id(GWS_BATCH);
     const uint f = get_global_id(GWS_FEATURE);
+#endif
+
 #if   INPUT0_DIMS == 2
     const uint y = 0;
     const uint x = 0;
@@ -37,8 +44,14 @@ KERNEL (reorder_data)(
     const uint u = 0;
     const uint v = 0;
 #elif INPUT0_DIMS == 4
-    const uint y = ((uint)(get_global_id(GWS_YX))) / INPUT0_SIZE_X;
-    const uint x = ((uint)(get_global_id(GWS_YX))) % INPUT0_SIZE_X;
+    #if INPUT0_LAYOUT_BYFX
+        // GWS_YX takes (F and X) axes for byfx format
+        const uint f = ((uint)(get_global_id(GWS_YX))) / INPUT0_SIZE_X;
+        const uint x = ((uint)(get_global_id(GWS_YX))) % INPUT0_SIZE_X;
+    #else
+        const uint y = ((uint)(get_global_id(GWS_YX))) / INPUT0_SIZE_X;
+        const uint x = ((uint)(get_global_id(GWS_YX))) % INPUT0_SIZE_X;
+    #endif
     const uint z = 0;
     const uint w = 0;
     const uint u = 0;

src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_kernel_base.cpp

@@ -189,6 +189,12 @@ ReorderKernelBase::DispatchData ReorderKernelBase::SetDefault(const reorder_para
         dispatchData.lws[0] = 1;
         dispatchData.lws[1] = 16;
         dispatchData.lws[2] = 1;
+    } else if (input_l == DataLayout::byfx) {
+        auto first_primary_axis_size = dispatchData.gws[0];  // X axis
+        auto second_primiary_axis_size =  dispatchData.gws[1];  // YF axes
+        dispatchData.gws[0] = first_primary_axis_size * input.Feature().v;  // takes XF axes
+        dispatchData.gws[1] = second_primiary_axis_size / input.Feature().v;  // takes Y axis
+        dispatchData.lws = {1, 1, 1};
     }
 
     return dispatchData;

src/plugins/intel_gpu/tests/unit/fusions/gemm_fusion_test.cpp

@@ -142,6 +142,8 @@ class GemmFusingTest : public ::BaseFusingTest<gemm_test_params> {
 #define CASE_GEMM_PERMUTES_FUSION_FP16_3 { { 17, 11, 2, 18 }, { 17, 11, 18, 4 } }, { 17, 11, 2, 4 }, data_types::f16, data_types::f16, data_types::f16, format::bfyx, data_types::f16, format::bfyx
 #define CASE_GEMM_PERMUTES_FUSION_FP16_4 { { 3, 2, 10, 12 }, { 3, 2, 12, 20 } }, { 3, 2, 10, 20 }, data_types::f16, data_types::f16, data_types::f16, format::bfyx, data_types::f16, format::bfyx
 #define CASE_GEMM_PERMUTES_FUSION_FP16_5 { { 3, 2, 16, 32 }, { 3, 2, 32, 16} }, { 3, 2, 16, 16 }, data_types::f16, data_types::f16, data_types::f16, format::bfyx, data_types::f16, format::bfyx
+#define CASE_GEMM_PERMUTES_FUSION_FP16_6 { { 3, 2, 16, 32 },  { 3, 16, 2, 32} }, { 3, 2, 2, 32 }, data_types::f16, data_types::f16, data_types::f16, format::bfyx, data_types::f16, format::bfyx
+
 class gemm_3in_quantize_i8 : public GemmFusingTest {};
 TEST_P(gemm_3in_quantize_i8, basic) {
     // TODO: Fix me, refer PR(#15873)
@@ -757,4 +759,40 @@ INSTANTIATE_TEST_SUITE_P(
         gemm_test_params{CASE_GEMM_PERMUTES_FUSION_FP16_3, 3, 6, "", broadcast_kinds::feature/*dummy*/, eltwise_mode::sum/*dummy*/, {{0, 2, 1, 3} /*byfx*/, {1, 2, 3, 0} /*xbfy*/, {0, 2, 1, 3} /*byfx*/}},
     }));
 
+class permute_gemm_reorder : public GemmFusingTestOneDNN {};
+TEST_P(permute_gemm_reorder, fused_permute_gemm_with_reorder) {
+    auto p = GetParam();
+    auto in_lay0 = get_input_layout(p, 0);
+    auto in_lay1 = get_input_layout(p, 1);
+    auto permute_in_lay0 = get_permute_input_shape(in_lay0.get_shape(), p.permute_orders[0]);
+    auto permute_in_lay1 = get_permute_input_shape(in_lay1.get_shape(), p.permute_orders[1]);
+    in_lay0.set_partial_shape(permute_in_lay0);
+    in_lay1.set_partial_shape(permute_in_lay1);
+    create_topologies(
+        input_layout("input0", in_lay0),
+        input_layout("input1", in_lay1),
+        permute("permute0", input_info("input0"), p.permute_orders[0]),
+        reorder("reorder_permute", input_info("permute0"), p.default_format, data_types::f32),
+        permute("permute1", input_info("input1"), p.permute_orders[1]),
+        gemm("gemm_prim", { input_info("permute0"), input_info("permute1") }, data_types::f16),
+        reorder("reorder_bfyx", input_info("gemm_prim"), p.default_format, data_types::f32),
+        eltwise("eltwise", { input_info("reorder_permute"), input_info("reorder_bfyx") }, eltwise_mode::sum, data_types::f32)
+    );
+
+    tolerance = default_tolerance(data_types::f16);
+    execute(p, false);
+}
+
+#define CASE_PERMUTES_GEMM_FUSION_FP16_1 { { 1, 12, 20, 64 }, { 1, 12, 64, 64 } }, { 1, 12, 20, 64 }, data_types::f16, data_types::f16, data_types::f16, format::bfyx, data_types::f16, format::bfyx
+#define CASE_PERMUTES_GEMM_FUSION_FP16_2 {  { 3, 2, 10, 12 },   { 3, 2, 12, 1 } },   { 3, 2, 10, 1 }, data_types::f16, data_types::f16, data_types::f16, format::bfyx, data_types::f16, format::bfyx
+
+INSTANTIATE_TEST_SUITE_P(
+    fusings_gpu, permute_gemm_reorder, ::testing::ValuesIn(std::vector<gemm_test_params>{
+        gemm_test_params{CASE_PERMUTES_GEMM_FUSION_FP16_1, 4, 6, "", broadcast_kinds::feature/*dummy*/, eltwise_mode::sum/*dummy*/, {{0, 2, 1, 3} /*byfx*/, {0, 2, 1, 3} /*byfx*/}},
+        gemm_test_params{CASE_PERMUTES_GEMM_FUSION_FP16_1, 4, 6, "", broadcast_kinds::feature/*dummy*/, eltwise_mode::sum/*dummy*/, {{0, 2, 1, 3} /*byfx*/, {1, 2, 3, 0} /*xbfy*/}},
+        gemm_test_params{CASE_PERMUTES_GEMM_FUSION_FP16_2, 4, 6, "", broadcast_kinds::feature/*dummy*/, eltwise_mode::sum/*dummy*/, {{0, 2, 1, 3} /*byfx*/, {0, 2, 1, 3} /*byfx*/}},
+        gemm_test_params{CASE_PERMUTES_GEMM_FUSION_FP16_2, 4, 6, "", broadcast_kinds::feature/*dummy*/, eltwise_mode::sum/*dummy*/, {{0, 2, 1, 3} /*byfx*/, {1, 2, 3, 0} /*xbfy*/}},
+    }));
+
+
 #endif // ENABLE_ONEDNN_FOR_GPU