diff options
| author | Michalis Spyrou <michalis.spyrou@arm.com> | 2021-04-29 17:23:02 +0100 |
|---|---|---|
| committer | Sheri Zhang <sheri.zhang@arm.com> | 2021-04-29 18:51:40 +0000 |
| commit | baebdb9a2c5ac1e7a25e4fb718dabcd508fe475a (patch) | |
| tree | cf9fb7b1ee807c3c1258d0701b0a9f6268f8af15 | |
| parent | 3ecf9fefa6f6299a0736599f150d4791cc8345d9 (diff) | |
| download | ComputeLibrary-baebdb9a2c5ac1e7a25e4fb718dabcd508fe475a.tar.gz | |
Remove Global pooling optimization
It doesn't take into consideration the padding.
Change-Id: Ie7a2a0c7aee93fd7007deb6e46f6bec0e7f06066
Signed-off-by: Michalis Spyrou <michalis.spyrou@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5529
Reviewed-by: TeresaARM <teresa.charlinreyes@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/pooling_layer.cl | 31 |
1 files changed, 1 insertions, 30 deletions
diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl index ebf7c5c078..8944c9b1ac 100644 --- a/src/core/CL/cl_kernels/pooling_layer.cl +++ b/src/core/CL/cl_kernels/pooling_layer.cl @@ -712,7 +712,6 @@ __kernel void pooling_layer_MxN_nhwc( int idx_out_n = GET_SPATIAL_IDX(2, 1, 0) / DST_HEIGHT; #else //DST_BATCH_SIZE != 1 int idx_out_h = GET_SPATIAL_IDX(2, 1, 0); - ; int idx_out_n = 0; #endif // DST_BATCH_SIZE != 1 @@ -738,32 +737,6 @@ __kernel void pooling_layer_MxN_nhwc( int filter_size = POOL_SIZE_X * POOL_SIZE_Y; #endif // defined(EXCLUDE_PADDING) -#if POOL_SIZE_X == SRC_WIDTH && POOL_SIZE_Y == SRC_HEIGHT - // Global pooling path - -#pragma unroll 8 - for(int y = 0; y < POOL_SIZE_X * POOL_SIZE_Y; ++y) - { - VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE) - data0; -#if defined(FP_MIXED_PRECISION) - // In case of FP_MIXED_PRECISION, ACC_DATA_TYPE is != DATA_TYPE - data0 = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(in_base_ptr), VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE)); -#else // defined(FP_MIXED_PRECISION) - data0 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(in_base_ptr)); -#endif // defined(FP_MIXED_PRECISION) - -#if defined(POOL_L2) - // Raise to power of 2 for L2 Pooling - data0 *= data0; -#endif // defined(POOL_L2) - - res0 = POOL_OP(res0, data0); - - in_base_ptr += input_stride_y; - } -#else // POOL_SIZE_X == SRC_WIDTH && POOL_SIZE_Y == SRC_HEIGHT - for(int y = pool_y_s; y < pool_y_e; ++y) { for(int x = pool_x_s; x < pool_x_e; ++x) @@ -784,9 +757,7 @@ __kernel void pooling_layer_MxN_nhwc( res0 = POOL_OP(res0, data0); } } - -#endif // POOL_SIZE_X == SRC_WIDTH && POOL_SIZE_Y == SRC_HEIGHT - + #if defined(POOL_AVG) || defined(POOL_L2) res0 /= (VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))filter_size; #endif // defined(POOL_AVG) || defined(POOL_L2) |