Skip to content

Commit 7d39426

Browse files
kelvinchoi-intelkelvinchoi-intel
committed
[GPU] Add fs_b_yx_fsv32 format for pooling int8 ref kernel
1 parent 3e3dc0e commit 7d39426

File tree

4 files changed

+25
-10
lines changed

4 files changed

+25
-10
lines changed

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

+1-1
Original file line numberDiff line numberDiff line change
@@ -74,7 +74,7 @@ KERNEL(pooling_gpu_int8_ref)(
7474
const uint f = bf / INPUT0_BATCH_NUM;
7575
const uint b = bf % INPUT0_BATCH_NUM;
7676
const uint z = 0;
77-
#elif OUTPUT_LAYOUT_B_FS_YX_FSV16 || OUTPUT_LAYOUT_BS_FS_YX_BSV32_FSV32 || OUTPUT_LAYOUT_BS_FS_YX_BSV16_FSV32
77+
#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
7878
const uint x = get_global_id(1);
7979
const uint y = get_global_id(2);
8080
const uint bf = (uint)get_global_id(0);

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

+2
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,7 @@ ParamsKey PoolingKernelGPUInt8Ref::GetSupportedKey() const {
1818
k.EnableInputLayout(DataLayout::bfzyx);
1919
k.EnableInputLayout(DataLayout::yxfb);
2020
k.EnableInputLayout(DataLayout::byxf);
21+
k.EnableInputLayout(DataLayout::fs_b_yx_fsv32);
2122
k.EnableInputLayout(DataLayout::b_fs_yx_fsv4);
2223
k.EnableInputLayout(DataLayout::b_fs_yx_fsv32);
2324
k.EnableInputLayout(DataLayout::b_fs_zyx_fsv32);
@@ -30,6 +31,7 @@ ParamsKey PoolingKernelGPUInt8Ref::GetSupportedKey() const {
3031
k.EnableOutputLayout(DataLayout::bfzyx);
3132
k.EnableOutputLayout(DataLayout::yxfb);
3233
k.EnableOutputLayout(DataLayout::byxf);
34+
k.EnableOutputLayout(DataLayout::fs_b_yx_fsv32);
3335
k.EnableOutputLayout(DataLayout::b_fs_yx_fsv4);
3436
k.EnableOutputLayout(DataLayout::b_fs_yx_fsv32);
3537
k.EnableOutputLayout(DataLayout::b_fs_zyx_fsv32);

‎src/plugins/intel_gpu/src/kernel_selector/kernels/pooling/pooling_kernel_gpu_ref.cpp

-2
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,6 @@ ParamsKey PoolingKernelGPURef::GetSupportedKey() const {
99
ParamsKey k;
1010
k.EnableInputDataType(Datatype::F16);
1111
k.EnableInputDataType(Datatype::F32);
12-
k.EnableInputDataType(Datatype::UINT8);
13-
k.EnableInputDataType(Datatype::INT8);
1412
k.EnableOutputDataType(Datatype::F16);
1513
k.EnableOutputDataType(Datatype::F32);
1614
k.EnableOutputDataType(Datatype::UINT8);

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

+22-7
Original file line numberDiff line numberDiff line change
@@ -1521,20 +1521,30 @@ TEST(pooling_forward_gpu, fs_b_yx_fsv32_int8)
15211521
{
15221522
auto& engine = get_test_engine();
15231523

1524-
auto input_prim = engine.allocate_memory({ data_types::f32, format::yxfb, { 1, 32, 1, 2 } });
1524+
auto input_prim = engine.allocate_memory({ data_types::f32, format::yxfb, { 2, 32, 1, 2 } });
15251525

15261526
topology topology;
15271527
topology.add(input_layout("input", input_prim->get_layout()));
1528-
topology.add(reorder("reorder_input", input_info("input"), layout(data_types::i8, format::fs_b_yx_fsv32, { 1, 32, 1, 2 })));
1528+
topology.add(reorder("reorder_input", input_info("input"), layout(data_types::i8, format::fs_b_yx_fsv32, { 2, 32, 1, 2 })));
15291529
topology.add(pooling("avg_pooling", input_info("reorder_input"), pooling_mode::average, { 7, 7 }, { 2, 2 }));
1530-
topology.add(reorder("reorder_after_pooling", input_info("avg_pooling"), layout(data_types::f32, format::fs_b_yx_fsv32, { 1, 32, 1, 1 })));
1530+
topology.add(reorder("reorder_after_pooling", input_info("avg_pooling"), layout(data_types::f32, format::fs_b_yx_fsv32, { 2, 32, 1, 1 })));
15311531

15321532
network network(engine, topology, get_test_default_config(engine));
15331533
set_values(input_prim, { 41.f, 42.f, 43.f, 44.f, 45.f, 46.f, 47.f, 48.f,
15341534
11.f, 12.f, 13.f, 14.f, 15.f, 16.f, 17.f, 18.f,
15351535
21.f, 22.f, 23.f, 24.f, 25.f, 26.f, 27.f, 28.f,
15361536
31.f, 32.f, 33.f, 34.f, 35.f, 36.f, 37.f, 38.f,
15371537

1538+
31.f, 32.f, 33.f, 34.f, 35.f, 36.f, 37.f, 38.f,
1539+
21.f, 22.f, 23.f, 24.f, 25.f, 26.f, 27.f, 28.f,
1540+
51.f, 52.f, 53.f, 54.f, 55.f, 56.f, 57.f, 58.f,
1541+
11.f, 12.f, 13.f, 14.f, 15.f, 16.f, 17.f, 18.f,
1542+
1543+
41.f, 42.f, 43.f, 44.f, 45.f, 46.f, 47.f, 48.f,
1544+
11.f, 12.f, 13.f, 14.f, 15.f, 16.f, 17.f, 18.f,
1545+
21.f, 22.f, 23.f, 24.f, 25.f, 26.f, 27.f, 28.f,
1546+
31.f, 32.f, 33.f, 34.f, 35.f, 36.f, 37.f, 38.f,
1547+
15381548
31.f, 32.f, 33.f, 34.f, 35.f, 36.f, 37.f, 38.f,
15391549
21.f, 22.f, 23.f, 24.f, 25.f, 26.f, 27.f, 28.f,
15401550
51.f, 52.f, 53.f, 54.f, 55.f, 56.f, 57.f, 58.f,
@@ -1549,10 +1559,15 @@ TEST(pooling_forward_gpu, fs_b_yx_fsv32_int8)
15491559

15501560
cldnn::mem_lock<float> output_ptr(output_prim, get_test_stream());
15511561

1552-
std::vector<float> ref_data = { 36.f, 37.f, 38.f, 39.f, 40.f, 41.f, 42.f, 43.f,
1553-
16.f, 17.f, 18.f, 19.f, 20.f, 21.f, 22.f, 23.f,
1554-
36.f, 37.f, 38.f, 39.f, 40.f, 41.f, 42.f, 43.f,
1555-
21.f, 22.f, 23.f, 24.f, 25.f, 26.f, 27.f, 28.f };
1562+
std::vector<float> ref_data = { 41.f, 43.f, 45.f, 47.f, 11.f, 13.f, 15.f, 17.f,
1563+
21.f, 23.f, 25.f, 27.f, 31.f, 33.f, 35.f, 37.f,
1564+
31.f, 33.f, 35.f, 37.f, 21.f, 23.f, 25.f, 27.f,
1565+
51.f, 53.f, 55.f, 57.f, 11.f, 13.f, 15.f, 17.f,
1566+
1567+
42.f, 44.f, 46.f, 48.f, 12.f, 14.f, 16.f, 18.f,
1568+
22.f, 24.f, 26.f, 28.f, 32.f, 34.f, 36.f, 38.f,
1569+
32.f, 34.f, 36.f, 38.f, 22.f, 24.f, 26.f, 28.f,
1570+
52.f, 54.f, 56.f, 58.f, 12.f, 14.f, 16.f, 18.f };
15561571

15571572
for (size_t i = 0; i < ref_data.size(); i++) {
15581573
ASSERT_EQ(ref_data[i], float(output_ptr[i]));

0 commit comments

Comments
 (0)