diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/pooling_gpu_int8_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/pooling_gpu_int8_ref.cl index d2f374215a1d4d..c2c47394b0b147 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/pooling_gpu_int8_ref.cl +++ b/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); diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/pooling/pooling_kernel_gpu_int8_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/pooling/pooling_kernel_gpu_int8_ref.cpp index 0c3ccf6d3604f1..9eee25f57ef600 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/pooling/pooling_kernel_gpu_int8_ref.cpp +++ b/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); diff --git a/src/plugins/intel_gpu/tests/unit/test_cases/pooling_gpu_test.cpp b/src/plugins/intel_gpu/tests/unit/test_cases/pooling_gpu_test.cpp index 8f76297493315b..461474335e903a 100644 --- a/src/plugins/intel_gpu/tests/unit/test_cases/pooling_gpu_test.cpp +++ b/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 _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(); + 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(); + 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(); + 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(); + 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(12, 12, 1)), + testing::Values(std::tuple(4, 4, 1)), + testing::Values(std::tuple(2, 2, 1)), + testing::Values(std::tuple(0, 0, 0)), + testing::Values(format::fs_b_yx_fsv32)), + testing::internal::DefaultParamName); + using pooling_random_test_fp16_fp32 = pooling_random_test; TEST_P(pooling_random_test_fp16_fp32, avg_fp16) {