commit adaae7e453cc4cc07905daca68fa7b938555d581
author Georgios Pinitas <georgios.pinitas@arm.com> Mon Oct 30 15:56:32 2017 +0000
committer Anthony Barbier <anthony.barbier@arm.com> Fri Nov 02 16:35:24 2018 +0000
tree 712003aa080cf577b2b2af318f219367897dc0e4
parent cf3935ffd4c67d9396c2435a3a28d3a159753105

COMPMID-647: Exclude padding pixels from averaging factor.

Adds support for excluding the padding pixels from the average scaling
factor calculation.

Change-Id: Ia13fbfeae235aff564db74191613921848231a01
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/93715
Reviewed-by: Robert Hughes <robert.hughes@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>

src/core/CL/cl_kernels/pooling_layer.cl

diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl
index 6379c28..635c44a 100644
--- a/src/core/CL/cl_kernels/pooling_layer.cl
+++ b/src/core/CL/cl_kernels/pooling_layer.cl
@@ -186,10 +186,14 @@
 DATA_TYPE calculate_avg_scale(const int pool_size, const int upper_bound_w, const int upper_bound_h,
                               const int pad_x, const int pad_y, const int stride_x, const int stride_y)
 {
-    const int start_x = get_global_id(0) * stride_x - pad_x;
-    const int start_y = get_global_id(1) * stride_y - pad_y;
+    int       start_x = get_global_id(0) * stride_x - pad_x;
+    int       start_y = get_global_id(1) * stride_y - pad_y;
     const int end_x   = min(start_x + pool_size, upper_bound_w);
     const int end_y   = min(start_y + pool_size, upper_bound_h);
+#if defined(EXCLUDE_PADDING)
+    start_x = max(0, start_x);
+    start_y = max(0, start_y);
+#endif /* defined(EXCLUDE_PADDING) */
     return ((end_y - start_y) * (end_x - start_x));
 }
 
@@ -334,10 +338,14 @@
 calculate_avg_scale4(const int pool_size, const int upper_bound_w, const int upper_bound_h,
                      const int pad_x, const int pad_y, const int stride_x, const int stride_y)
 {
-    const int4 start_x = ((int4)get_global_id(0) * 4 + (int4)(0, 1, 2, 3)) * (int4)stride_x - (int4)pad_x;
-    const int  start_y = get_global_id(1) * stride_y - pad_y;
+    int4       start_x = ((int4)get_global_id(0) * 4 + (int4)(0, 1, 2, 3)) * (int4)stride_x - (int4)pad_x;
+    int        start_y = get_global_id(1) * stride_y - pad_y;
     const int4 end_x   = min(start_x + (int4)pool_size, (int4)upper_bound_w);
     const int  end_y   = min(start_y + pool_size, upper_bound_h);
+#if defined(EXCLUDE_PADDING)
+    start_x = max((int4)0, start_x);
+    start_y = max(0, start_y);
+#endif /* defined(EXCLUDE_PADDING) */
     return (VEC_DATA_TYPE(DATA_TYPE, 4))(1.f) / CONVERT_VECTOR4(DATA_TYPE)(((int4)(end_y - start_y)) * (end_x - start_x));
 }