diff options
| author | Georgios Pinitas <georgios.pinitas@arm.com> | 2017-10-30 15:56:32 +0000 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:35:24 +0000 |
| commit | adaae7e453cc4cc07905daca68fa7b938555d581 (patch) | |
| tree | 712003aa080cf577b2b2af318f219367897dc0e4 /src/core/CL/cl_kernels/pooling_layer.cl | |
| parent | cf3935ffd4c67d9396c2435a3a28d3a159753105 (diff) | |
| download | ComputeLibrary-adaae7e453cc4cc07905daca68fa7b938555d581.tar.gz | |
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>
Diffstat (limited to 'src/core/CL/cl_kernels/pooling_layer.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/pooling_layer.cl | 16 |
1 files changed, 12 insertions, 4 deletions
diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl index 6379c288b3..635c44a849 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 @@ VEC_DATA_TYPE(DATA_TYPE, 4) 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)); } |