diff options
| author | Sang-Hoon Park <sang-hoon.park@arm.com> | 2020-07-02 10:49:39 +0100 |
|---|---|---|
| committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2020-07-03 17:39:20 +0000 |
| commit | 94d5051bc56a59d857f8a09560e2da5c0d7894b0 (patch) | |
| tree | 577d9ce10f73608a71c033b495b54457e6c5b1ab /src/core/CL/cl_kernels | |
| parent | 2aad21a900a21f467b3ec6b37420f892f0d80221 (diff) | |
| download | ComputeLibrary-94d5051bc56a59d857f8a09560e2da5c0d7894b0.tar.gz | |
COMPMID-3539: Change indexing for nearest neighbor with aligned corners
For nearest neighbor interpolation policy with aligned corners
all of NEON, CL and reference use round() rather than float to
find the nearest integer.
Change-Id: If0360da870e983303bf0424ca1100084084c1efc
Signed-off-by: Sang-Hoon Park <sang-hoon.park@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3495
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels')
| -rw-r--r-- | src/core/CL/cl_kernels/scale.cl | 26 |
1 files changed, 17 insertions, 9 deletions
diff --git a/src/core/CL/cl_kernels/scale.cl b/src/core/CL/cl_kernels/scale.cl index 499f9ea53f..971087021f 100644 --- a/src/core/CL/cl_kernels/scale.cl +++ b/src/core/CL/cl_kernels/scale.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -100,10 +100,14 @@ __kernel void scale_nearest_neighbour_nchw( const float scale_x, const float scale_y) { - Image in = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in); - Image out = CONVERT_TO_IMAGE_STRUCT(out); - const float2 r = (float2)(scale_x, scale_y); - const float8 tc = clamp_to_border_with_size(transform_nearest(get_current_coords(), r), input_width, input_height, BORDER_SIZE); + Image in = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in); + Image out = CONVERT_TO_IMAGE_STRUCT(out); + const float2 r = (float2)(scale_x, scale_y); + float8 transformed = transform_nearest(get_current_coords(), r); +#ifdef ALIGN_CORNERS + transformed = round(transformed); +#endif // ALIGN_CORNERS + const float8 tc = clamp_to_border_with_size(transformed, input_width, input_height, BORDER_SIZE); vstore4(read_texels4(&in, convert_int8(tc)), 0, (__global DATA_TYPE *)out.ptr); } @@ -182,14 +186,18 @@ __kernel void scale_nearest_neighbour_nhwc( Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT(out, DEPTH_OUT); #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; + float new_x = get_global_id(1) * scale_x; + 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; + float new_x = (get_global_id(1) + 0.5f) * scale_x; + float new_y = ((get_global_id(2) % DEPTH_OUT) + 0.5f) * scale_y; #else /* SAMPLING_POLICY */ #error("Unsupported sampling policy"); #endif /* SAMPLING_POLICY */ +#ifdef ALIGN_CORNERS + new_x = round(new_x); + new_y = round(new_y); +#endif /* ALIGN_CORNERS */ const float clamped_x = clamp(new_x, 0.0f, input_width - 1); const float clamped_y = clamp(new_y, 0.0f, input_height - 1); |