From 7a49c7993a8b2c670f7caa90e3fbe8a5f1c03078 Mon Sep 17 00:00:00 2001 From: Daniil Efremov Date: Tue, 14 Nov 2017 21:25:34 +0700 Subject: COMPMID-661: issue# 23 Scale border fix (#26) Changes in CL and reference in terms of border handling. Change-Id: I5bed95b1f4c308629d7113455dc8a55d74500bcd Reviewed-on: http://mpd-gerrit.cambridge.arm.com/95742 Reviewed-by: Anthony Barbier Tested-by: Kaizen --- src/core/CL/cl_kernels/scale.cl | 4 +-- src/core/CL/cl_kernels/warp_helpers.h | 53 +++++++++++++++++++++++++++-------- src/core/CL/kernels/CLScaleKernel.cpp | 27 ++++++++++++------ 3 files changed, 61 insertions(+), 23 deletions(-) (limited to 'src/core') diff --git a/src/core/CL/cl_kernels/scale.cl b/src/core/CL/cl_kernels/scale.cl index 0106ce095c..d533d970c7 100644 --- a/src/core/CL/cl_kernels/scale.cl +++ b/src/core/CL/cl_kernels/scale.cl @@ -84,7 +84,7 @@ __kernel void scale_nearest_neighbour( 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(transform_nearest(get_current_coords(), r), input_width, input_height); + const float8 tc = clamp_to_border_with_size(transform_nearest(get_current_coords(), r), input_width, input_height, BORDER_SIZE); vstore4(read_texels4(&in, convert_int8(tc)), 0, (__global DATA_TYPE *)out.ptr); } @@ -119,5 +119,5 @@ __kernel void scale_bilinear( Image out = CONVERT_TO_IMAGE_STRUCT(out); const float2 r = (float2)(scale_x, scale_y); const float8 tc = transform_bilinear(get_current_coords(), r); - vstore4(bilinear_interpolate(&in, tc, input_width, input_height), 0, (__global DATA_TYPE *)out.ptr); + vstore4(bilinear_interpolate_with_border(&in, tc, input_width, input_height, BORDER_SIZE), 0, (__global DATA_TYPE *)out.ptr); } diff --git a/src/core/CL/cl_kernels/warp_helpers.h b/src/core/CL/cl_kernels/warp_helpers.h index 26a8b859a4..76f0a4a104 100644 --- a/src/core/CL/cl_kernels/warp_helpers.h +++ b/src/core/CL/cl_kernels/warp_helpers.h @@ -23,6 +23,22 @@ */ #include "helpers.h" +/** Clamps the given coordinates to the borders according to the border size. + * + * @param[in] coords Vector of 2D coordinates to clamp. Even positions are X coords, odd positions are Y coords. + * @param[in] width Width of the image + * @param[in] height Height of the image + * @param[in] border_size Border size of the image + * + */ +inline const float8 clamp_to_border_with_size(float8 coords, const float width, const float height, const float border_size) +{ + const float4 clamped_x = clamp(coords.even, 0.0f - border_size, width - 1 + border_size); + const float4 clamped_y = clamp(coords.odd, 0.0f - border_size, height - 1 + border_size); + return (float8)(clamped_x.s0, clamped_y.s0, clamped_x.s1, clamped_y.s1, clamped_x.s2, clamped_y.s2, clamped_x.s3, clamped_y.s3); +} + +/* FIXME(COMPMID-682): Clamp border properly in UNDEFINED border mode in Warp, Scale, Remap */ /** Clamps the given coordinates to the borders. * * @param[in] coords Vector of 2D coordinates to clamp. Even positions are X coords, odd positions are Y coords. @@ -32,9 +48,7 @@ */ inline const float8 clamp_to_border(float8 coords, const float width, const float height) { - const float4 clamped_x = clamp(coords.even, -1.0f, width); - const float4 clamped_y = clamp(coords.odd, -1.0f, height); - return (float8)(clamped_x.s0, clamped_y.s0, clamped_x.s1, clamped_y.s1, clamped_x.s2, clamped_y.s2, clamped_x.s3, clamped_y.s3); + return clamp_to_border_with_size(coords, width, height, 1); } /** Reads four texels from the input image. The coords vector is used to determine which texels to be read. @@ -72,23 +86,25 @@ inline const float8 get_neighbour_coords(const float2 coord) /** Computes the bilinear interpolation for each set of coordinates in the vector coords and returns the values * - * @param[in] in Pointer to the source image. - * @param[in] coords Vector of four 2D coordinates. Even pos is x and odd y. - * @param[in] width Width of the image - * @param[in] height Height of the image + * @param[in] in Pointer to the source image. + * @param[in] coords Vector of four 2D coordinates. Even pos is x and odd y. + * @param[in] width Width of the image + * @param[in] height Height of the image + * @param[in] border_size Border size */ -inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate(const Image *in, const float8 coords, const float width, const float height) +inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate_with_border(const Image *in, const float8 coords, const float width, const float height, const float border_size) { // If any of the 4 texels is out of the image's boundaries we use the border value (REPLICATE or CONSTANT) for any texel out of the image. // Sets the 4x4 coordinates for each of the four input texels const float8 fc = floor(coords); const float16 c1 = (float16)( - clamp_to_border(get_neighbour_coords((float2)(fc.s0, fc.s1)), width, height), - clamp_to_border(get_neighbour_coords((float2)(fc.s2, fc.s3)), width, height)); + clamp_to_border_with_size(get_neighbour_coords((float2)(fc.s0, fc.s1)), width, height, border_size), + clamp_to_border_with_size(get_neighbour_coords((float2)(fc.s2, fc.s3)), width, height, border_size)); const float16 c2 = (float16)( - clamp_to_border(get_neighbour_coords((float2)(fc.s4, fc.s5)), width, height), - clamp_to_border(get_neighbour_coords((float2)(fc.s6, fc.s7)), width, height)); + clamp_to_border_with_size(get_neighbour_coords((float2)(fc.s4, fc.s5)), width, height, border_size), + clamp_to_border_with_size(get_neighbour_coords((float2)(fc.s6, fc.s7)), width, height, border_size)); + // Loads the values from the input image const float16 t = (float16)( /* tl, tr, bl, br */ @@ -109,3 +125,16 @@ inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate(const Image *in, c ((t.sc * b.s6 * b.s7) + (t.sd * a.s6 * b.s7) + (t.se * b.s6 * a.s7) + (t.sf * a.s6 * a.s7))); return CONVERT(fr, VEC_DATA_TYPE(DATA_TYPE, 4)); } + +/* FIXME(COMPMID-682): Clamp border properly in UNDEFINED border mode in Warp, Scale, Remap */ +/** Computes the bilinear interpolation for each set of coordinates in the vector coords and returns the values + * + * @param[in] in Pointer to the source image. + * @param[in] coords Vector of four 2D coordinates. Even pos is x and odd y. + * @param[in] width Width of the image + * @param[in] height Height of the image +*/ +inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate(const Image *in, const float8 coords, const float width, const float height) +{ + return bilinear_interpolate_with_border(in, coords, width, height, 1); +} diff --git a/src/core/CL/kernels/CLScaleKernel.cpp b/src/core/CL/kernels/CLScaleKernel.cpp index 82ebe644ea..6a5d24c943 100644 --- a/src/core/CL/kernels/CLScaleKernel.cpp +++ b/src/core/CL/kernels/CLScaleKernel.cpp @@ -54,11 +54,14 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo _input = input; _output = output; - /* Compute the ratio between source width/height and destination width/height */ + // Compute the ratio between source width/height and destination width/height const auto wr = static_cast(input->info()->dimension(0)) / static_cast(output->info()->dimension(0)); const auto hr = static_cast(input->info()->dimension(1)) / static_cast(output->info()->dimension(1)); - /* Area interpolation behaves as Nearest Neighbour in case of up-sampling */ + // Compute actual border size + BorderSize border = border_undefined ? BorderSize(0) : border_size(); + + // Area interpolation behaves as Nearest Neighbour in case of up-sampling if(policy == InterpolationPolicy::AREA && wr <= 1.f && hr <= 1.f) { policy = InterpolationPolicy::NEAREST_NEIGHBOR; @@ -69,11 +72,14 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo } // Create kernel - std::set build_opts = { ("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())) }; - std::string interpolation_name = string_from_interpolation_policy(policy); + CLBuildOptions build_opts; + build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); + build_opts.add_option("-DBORDER_SIZE=" + support::cpp11::to_string(border.right)); + + std::string interpolation_name = string_from_interpolation_policy(policy); std::transform(interpolation_name.begin(), interpolation_name.end(), interpolation_name.begin(), ::tolower); std::string kernel_name = "scale_" + interpolation_name; - _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts)); + _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); // Configure kernel window constexpr unsigned int num_elems_processed_per_iteration = 4; @@ -84,15 +90,18 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo // Reads can occur within the valid region of the input AccessWindowStatic input_access(input->info(), - input_valid_region.anchor[0] - border_size().left, input_valid_region.anchor[1] - border_size().top, - input_valid_region.anchor[0] + input_valid_region.shape[0] + border_size().right, - input_valid_region.anchor[1] + input_valid_region.shape[1] + border_size().bottom); + input_valid_region.anchor[0] - border.left, input_valid_region.anchor[1] - border.top, + input_valid_region.anchor[0] + input_valid_region.shape[0] + border.right, + input_valid_region.anchor[1] + input_valid_region.shape[1] + border.bottom); AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); update_window_and_padding(win, input_access, output_access); - output_access.set_valid_region(win, calculate_valid_region_scale(*(input->info()), output->info()->tensor_shape(), policy, border_size(), + output_access.set_valid_region(win, calculate_valid_region_scale(*(input->info()), + output->info()->tensor_shape(), + policy, + border, border_undefined)); ICLKernel::configure(win); -- cgit v1.2.1