Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

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

Merged
merged 4 commits into from
Dec 4, 2024
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
0b0f177
[GPU] Add int8 input type for pooling ref kernel to cover fs_b_yx_fsv…
kelvinchoi-intel Nov 1, 2024
c13f570
[GPU] Add fs_b_yx_fsv32 format for pooling int8 ref kernel
kelvinchoi-intel Nov 27, 2024
bac20a6
Add more unittests
kelvinchoi-intel Nov 29, 2024
247a48a
Update unittest to reduce the number of test
kelvinchoi-intel Dec 4, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion src/plugins/intel_gpu/src/kernel_selector/cl_kernels/pooling_gpu_int8_ref.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 2 additions & 0 deletions src/plugins/intel_gpu/src/kernel_selector/kernels/pooling/pooling_kernel_gpu_int8_ref.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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);
Expand Down
34 changes: 34 additions & 0 deletions src/plugins/intel_gpu/tests/unit/test_cases/pooling_gpu_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down
Loading