aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/pooling_layer.cl
diff options
context:
space:
mode:
Diffstat (limited to 'src/core/CL/cl_kernels/pooling_layer.cl')
-rw-r--r--src/core/CL/cl_kernels/pooling_layer.cl31
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)