diff options
| author | Michalis Spyrou <michalis.spyrou@arm.com> | 2019-07-09 14:21:06 +0100 |
|---|---|---|
| committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2019-07-11 13:31:25 +0000 |
| commit | d473386e4d5e0edcf55e13a2bf3c422a23fac0de (patch) | |
| tree | 64c4132e8a46c809639d719e0426a7e6b9dd0371 /src/core/CL | |
| parent | 9c9b70b9d30482d34f4f9c9dbc6479df163f96a1 (diff) | |
| download | ComputeLibrary-d473386e4d5e0edcf55e13a2bf3c422a23fac0de.tar.gz | |
COMPMID-2447: Align TFlite nearest neighbor NE/CL functions with ACL
Change-Id: Idd7b23247491d6e2e31d19b2a8aa522470ca174c
Signed-off-by: Michalis Spyrou <michalis.spyrou@arm.com>
Reviewed-on: https://review.mlplatform.org/c/1500
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Diffstat (limited to 'src/core/CL')
| -rw-r--r-- | src/core/CL/cl_kernels/scale.cl | 22 |
1 files changed, 19 insertions, 3 deletions
diff --git a/src/core/CL/cl_kernels/scale.cl b/src/core/CL/cl_kernels/scale.cl index 5ac6443c98..499f9ea53f 100644 --- a/src/core/CL/cl_kernels/scale.cl +++ b/src/core/CL/cl_kernels/scale.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -33,10 +33,19 @@ */ inline const float8 transform_nearest(const float2 coord, const float2 scale) { +#ifdef SAMPLING_POLICY_TOP_LEFT + const float4 in_x_coords = (float4)(coord.s0, 1 + coord.s0, 2 + coord.s0, 3 + coord.s0); + const float4 new_x = in_x_coords * (float4)(scale.s0); + const float4 new_y = (float4)(coord.s1 * scale.s1); + return (float8)(new_x.s0, new_y.s0, new_x.s1, new_y.s1, new_x.s2, new_y.s2, new_x.s3, new_y.s3); +#elif SAMPLING_POLICY_CENTER const float4 in_x_coords = (float4)(coord.s0, 1 + coord.s0, 2 + coord.s0, 3 + coord.s0); const float4 new_x = (in_x_coords + ((float4)(0.5f))) * (float4)(scale.s0); const float4 new_y = (float4)((coord.s1 + 0.5f) * scale.s1); return (float8)(new_x.s0, new_y.s0, new_x.s1, new_y.s1, new_x.s2, new_y.s2, new_x.s3, new_y.s3); +#else /* SAMPLING_POLICY */ +#error("Unsupported sampling policy"); +#endif /* SAMPLING_POLICY */ } /** Transforms four 2D coordinates. This is used to map the output coordinates to the input coordinates. @@ -172,8 +181,15 @@ __kernel void scale_nearest_neighbour_nhwc( Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(in, 0); Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT(out, DEPTH_OUT); - const float new_x = (get_global_id(1) + 0.5f) * scale_x; - const float new_y = ((get_global_id(2) % DEPTH_OUT) + 0.5f) * scale_y; +#ifdef SAMPLING_POLICY_TOP_LEFT + const float new_x = get_global_id(1) * scale_x; + const float new_y = (get_global_id(2) % DEPTH_OUT) * scale_y; +#elif SAMPLING_POLICY_CENTER + const float new_x = (get_global_id(1) + 0.5f) * scale_x; + const float new_y = ((get_global_id(2) % DEPTH_OUT) + 0.5f) * scale_y; +#else /* SAMPLING_POLICY */ +#error("Unsupported sampling policy"); +#endif /* SAMPLING_POLICY */ const float clamped_x = clamp(new_x, 0.0f, input_width - 1); const float clamped_y = clamp(new_y, 0.0f, input_height - 1); |