diff options
| author | Pablo Tello <pablo.tello@arm.com> | 2021-06-07 15:06:17 +0100 |
|---|---|---|
| committer | Sheri Zhang <sheri.zhang@arm.com> | 2021-06-07 16:09:00 +0000 |
| commit | 18834933ac69330190a974c42ce0a22264a76431 (patch) | |
| tree | a983c8b04385b117f9ebadeae1b2a7a8952ddbbf /src/core/CL/cl_kernels/pooling_layer.cl | |
| parent | bdcdc39d89b6a6556f5c0483af5379f75eae0c55 (diff) | |
| download | ComputeLibrary-18834933ac69330190a974c42ce0a22264a76431.tar.gz | |
Revert "Add optimization for global pooling in pooling_layer.cl"
* Resolves COMPMID-4544
This reverts commit 50929ef951880469b9d579323d2f9c9f5025327d.
Signed-off-by: Pablo Tello <pablo.tello@arm.com>
Change-Id: Ibc40054cecc2fe99bff48154f11fe5b602b4a64a
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/c/VisualCompute/ComputeLibrary/+/330380
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Sheri Zhang <sheri.zhang@arm.com>
Comments-Addressed: bsgcomp <bsgcomp@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5774
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/pooling_layer.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/pooling_layer.cl | 32 |
1 files changed, 1 insertions, 31 deletions
diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl index ba61674c7b..8944c9b1ac 100644 --- a/src/core/CL/cl_kernels/pooling_layer.cl +++ b/src/core/CL/cl_kernels/pooling_layer.cl @@ -723,34 +723,6 @@ __kernel void pooling_layer_MxN_nhwc( VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE) res0 = INITIAL_VALUE; -#if POOL_SIZE_X == SRC_WIDTH && POOL_SIZE_Y == SRC_HEIGHT && PAD_X == 0 && PAD_Y == 0 - // Global pooling path - - int filter_size = POOL_SIZE_X * POOL_SIZE_Y; - -#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 && PAD_X == 0 && PAD_Y == 0 - int idx_in_w = idx_out_w * STRIDE_X - PAD_X; int idx_in_h = idx_out_h * STRIDE_Y - PAD_Y; @@ -785,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 && PAD_X == 0 && PAD_Y == 0 - + #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) |