diff options
Diffstat (limited to 'src/core/CL/cl_kernels/scale.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/scale.cl | 127 |
1 files changed, 124 insertions, 3 deletions
diff --git a/src/core/CL/cl_kernels/scale.cl b/src/core/CL/cl_kernels/scale.cl index a2ae8c4dd6..744f28a918 100644 --- a/src/core/CL/cl_kernels/scale.cl +++ b/src/core/CL/cl_kernels/scale.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -83,7 +83,7 @@ inline const float8 transform_bilinear(const float2 coord, const float2 scale) * @param[in] scale_x The scale factor along x dimension * @param[in] scale_y The scale factor along y dimension */ -__kernel void scale_nearest_neighbour( +__kernel void scale_nearest_neighbour_nchw( IMAGE_DECLARATION(in), IMAGE_DECLARATION(out), const float input_width, @@ -119,7 +119,7 @@ __kernel void scale_nearest_neighbour( * @param[in] scale_x The scale factor along x dimension * @param[in] scale_y The scale factor along y dimension */ -__kernel void scale_bilinear( +__kernel void scale_bilinear_nchw( IMAGE_DECLARATION(in), IMAGE_DECLARATION(out), const float input_width, @@ -133,3 +133,124 @@ __kernel void scale_bilinear( const float8 tc = transform_bilinear(get_current_coords(), r); vstore4(bilinear_interpolate_with_border(&in, tc, input_width, input_height, BORDER_SIZE), 0, (__global DATA_TYPE *)out.ptr); } + +/** Performs scale on an image interpolating with the NEAREAST NEIGHBOUR method. Input and output are single channel F32. (NHWC) + * + * @note Sampling policy to used is passed as -DSAMPLING_POLICY_(TYPE) e.g. -DSAMPLING_POLICY_TOP_LEFT + * + * @param[in] in_ptr Pointer to the source image. Supported data types: U8/S16/F16/F32. + * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] in_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] in_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] in_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] in_stride_z Stride of the source image in Z dimension (in bytes) + * @param[in] in_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] in_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] out_ptr Pointer to the destination image. Supported data types: same as @p in_ptr + * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] out_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] out_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] out_stride_z Stride of the destination image in Z dimension (in bytes) + * @param[in] out_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[in] input_width Input image width + * @param[in] input_height Input image height + * @param[in] scale_x The scale factor along x dimension + * @param[in] scale_y The scale factor along y dimension + */ +__kernel void scale_nearest_neighbour_nhwc( + TENSOR3D_DECLARATION(in), + TENSOR3D_DECLARATION(out), + const float input_width, + const float input_height, + const float scale_x, + const float scale_y) +{ + Tensor3D in = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(in); + Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out); + + const float new_x = (get_global_id(1) + 0.5f) * scale_x; + const float new_y = (get_global_id(2) + 0.5f) * scale_y; + const float clamped_x = clamp(new_x, 0.0f, input_width - 1); + const float clamped_y = clamp(new_y, 0.0f, input_height - 1); + + *((__global DATA_TYPE *)out.ptr) = *((__global DATA_TYPE *)tensor3D_offset(&in, get_global_id(0), convert_int(clamped_x), convert_int(clamped_y))); +} + +/** Performs scale on an image interpolating with the BILINEAR method. (NHWC) + * + * @note Sampling policy to be used is passed as -DSAMPLING_POLICY_(TYPE) e.g. -DSAMPLING_POLICY_TOP_LEFT + * @note If border mode replicate is used, is should be passed as -DBORDER_MODE_REPLICATE + * + * @param[in] in_ptr Pointer to the source image. Supported data types: U8/S16/F16/F32. + * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] in_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] in_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] in_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] in_stride_z Stride of the source image in Z dimension (in bytes) + * @param[in] in_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] in_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] out_ptr Pointer to the destination image. Supported data types: same as @p in_ptr + * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] out_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] out_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] out_stride_z Stride of the destination image in Z dimension (in bytes) + * @param[in] out_step_z dst_stride_y * number of elements along Z processed per workitem(in bytes) + * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[in] input_width Input image width + * @param[in] input_height Input image height + * @param[in] scale_x The scale factor along x dimension + * @param[in] scale_y The scale factor along y dimension + */ +__kernel void scale_bilinear_nhwc( + TENSOR3D_DECLARATION(in), + TENSOR3D_DECLARATION(out), + const float input_width, + const float input_height, + const float scale_x, + const float scale_y) +{ + Tensor3D in = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(in); + Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out); + +#ifdef SAMPLING_POLICY_TOP_LEFT + const float new_x = get_global_id(1) * scale_x; + const float new_y = get_global_id(2) * scale_y; +#elif SAMPLING_POLICY_CENTER + const float new_x = (get_global_id(1) + 0.5f) * scale_x - 0.5f; + const float new_y = (get_global_id(2) + 0.5f) * scale_y - 0.5f; +#else /* SAMPLING_POLICY */ +#error("Unsupported sampling policy"); +#endif /* SAMPLING_POLICY */ + + const float new_xf = floor(new_x); + const float new_yf = floor(new_y); + float clamped_x = clamp(new_xf, 0.0f, input_width - 1); + float clamped_x1 = clamp(new_xf + 1, 0.0f, input_width - 1); + float clamped_x_ = clamped_x; + float clamped_x1_ = clamped_x1; + const float clamped_y = clamp(new_yf, 0.0f, input_height - 1); + const float clamped_y1 = clamp(new_yf + 1, 0.0f, input_height - 1); + +#ifndef BORDER_MODE_REPLICATE + clamped_x1 = select(clamped_x1, 0.0f - BORDER_SIZE, new_yf + 1 < 0.f || new_yf + 1 > input_height - 1 || new_xf + 1 < 0.f || new_xf + 1 > input_width - 1); + clamped_x_ = select(clamped_x_, 0.0f - BORDER_SIZE, new_yf + 1 > input_height - 1 || new_xf < 0.f || new_xf > input_width - 1); + clamped_x = select(clamped_x, 0.0f - BORDER_SIZE, new_yf < 0.f || new_yf > input_height - 1 || new_xf < 0.f || new_xf > input_width - 1); + clamped_x1_ = select(clamped_x1_, 0.0f - BORDER_SIZE, new_xf + 1 < 0.f || new_xf + 1 > input_width - 1 || new_yf < 0.f || new_yf > input_height - 1); +#endif /* BORDER_MODE_REPLICATE */ + + float4 ins = (float4)(*((__global DATA_TYPE *)tensor3D_offset(&in, get_global_id(0), convert_int(clamped_x), convert_int(clamped_y))), + *((__global DATA_TYPE *)tensor3D_offset(&in, get_global_id(0), convert_int(clamped_x1_), convert_int(clamped_y))), + *((__global DATA_TYPE *)tensor3D_offset(&in, get_global_id(0), convert_int(clamped_x_), convert_int(clamped_y1))), + *((__global DATA_TYPE *)tensor3D_offset(&in, get_global_id(0), convert_int(clamped_x1), convert_int(clamped_y1)))); + + const float a = new_x - new_xf; + const float b = 1.f - a; + const float a1 = new_y - new_yf; + const float b1 = 1.f - a1; + const float fr = ((ins.s0 * b * b1) + (ins.s1 * a * b1) + (ins.s2 * b * a1) + (ins.s3 * a * a1)); + + *((__global DATA_TYPE *)out.ptr) = CONVERT(fr, DATA_TYPE); +} |