[GPU] Support fs_b_yx_fsv32 and int8 case for pooling (openvinotoolkit#27371)

### Details:
 - *Support fs_b_yx_fsv32 and int8 case for pooling*

### Tickets:
 - *157507*

Author: kelvinchoi-intel

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

@@ -74,7 +74,7 @@ KERNEL(pooling_gpu_int8_ref)(
     const uint f    = bf / INPUT0_BATCH_NUM;
     const uint b    = bf % INPUT0_BATCH_NUM;
     const uint z    = 0;
-#elif OUTPUT_LAYOUT_B_FS_YX_FSV16 || OUTPUT_LAYOUT_BS_FS_YX_BSV32_FSV32 || OUTPUT_LAYOUT_BS_FS_YX_BSV16_FSV32
+#elif OUTPUT_LAYOUT_B_FS_YX_FSV16 || OUTPUT_LAYOUT_BS_FS_YX_BSV32_FSV32 || OUTPUT_LAYOUT_BS_FS_YX_BSV16_FSV32 || OUTPUT_LAYOUT_FS_B_YX_FSV32
     const uint x = get_global_id(1);
     const uint y = get_global_id(2);
     const uint bf = (uint)get_global_id(0);

src/plugins/intel_gpu/src/kernel_selector/kernels/pooling/pooling_kernel_gpu_int8_ref.cpp

@@ -18,6 +18,7 @@ ParamsKey PoolingKernelGPUInt8Ref::GetSupportedKey() const {
     k.EnableInputLayout(DataLayout::bfzyx);
     k.EnableInputLayout(DataLayout::yxfb);
     k.EnableInputLayout(DataLayout::byxf);
+    k.EnableInputLayout(DataLayout::fs_b_yx_fsv32);
     k.EnableInputLayout(DataLayout::b_fs_yx_fsv4);
     k.EnableInputLayout(DataLayout::b_fs_yx_fsv32);
     k.EnableInputLayout(DataLayout::b_fs_zyx_fsv32);
@@ -30,6 +31,7 @@ ParamsKey PoolingKernelGPUInt8Ref::GetSupportedKey() const {
     k.EnableOutputLayout(DataLayout::bfzyx);
     k.EnableOutputLayout(DataLayout::yxfb);
     k.EnableOutputLayout(DataLayout::byxf);
+    k.EnableOutputLayout(DataLayout::fs_b_yx_fsv32);
     k.EnableOutputLayout(DataLayout::b_fs_yx_fsv4);
     k.EnableOutputLayout(DataLayout::b_fs_yx_fsv32);
     k.EnableOutputLayout(DataLayout::b_fs_zyx_fsv32);

src/plugins/intel_gpu/tests/unit/test_cases/pooling_gpu_test.cpp

@@ -2276,6 +2276,40 @@ class pooling_scale_random_test_base : public pooling_random_test_base<InputT, M
     VF<output_t> _shift;
 };
 
+using pooling_random_test_int8_uint8 = pooling_random_test;
+
+TEST_P(pooling_random_test_int8_uint8, avg_int8) {
+    auto test_case = pooling_random_test_base<int8_t, pooling_mode::average>();
+    ASSERT_NO_FATAL_FAILURE(test_case.run_random(GetParam(), false));
+}
+
+TEST_P(pooling_random_test_int8_uint8, max_int8) {
+    auto test_case = pooling_random_test_base<int8_t, pooling_mode::max>();
+    ASSERT_NO_FATAL_FAILURE(test_case.run_random(GetParam(), false));
+}
+
+TEST_P(pooling_random_test_int8_uint8, avg_uint8) {
+    auto test_case = pooling_random_test_base<uint8_t, pooling_mode::average>();
+    ASSERT_NO_FATAL_FAILURE(test_case.run_random(GetParam(), false));
+}
+
+TEST_P(pooling_random_test_int8_uint8, max_uint8) {
+    auto test_case = pooling_random_test_base<uint8_t, pooling_mode::max>();
+    ASSERT_NO_FATAL_FAILURE(test_case.run_random(GetParam(), false));
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    smoke_low_precision,
+    pooling_random_test_int8_uint8,
+    testing::Combine(testing::Values(1, 2),
+                     testing::Values(3, 8),
+                     testing::Values(std::tuple<size_t, size_t, size_t>(12, 12, 1)),
+                     testing::Values(std::tuple<size_t, size_t, size_t>(4, 4, 1)),
+                     testing::Values(std::tuple<int, int, int>(2, 2, 1)),
+                     testing::Values(std::tuple<int, int, int>(0, 0, 0)),
+                     testing::Values(format::fs_b_yx_fsv32)),
+    testing::internal::DefaultParamName<pooling_random_test_params>);
+
 using pooling_random_test_fp16_fp32 = pooling_random_test;
 
 TEST_P(pooling_random_test_fp16_fp32, avg_fp16) {