Remove padding from ClScaleKernel
- Merge quantized kernels with fp for bilinear interpolation (both NCHW and NHWC)
- Pass dimensions at compile time rather than at run time
- Use tile-based approach to rework the NCHW kernels
- Remove unused functions/files
Resolve COMPMID-4723
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Change-Id: Ifcdf02beb9daa9f318395751b3c85eb2fe874082
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/6138
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
diff --git a/src/core/CL/cl_kernels/nchw/scale.cl b/src/core/CL/cl_kernels/nchw/scale.cl
index 63a53cc..2b4d6be 100644
--- a/src/core/CL/cl_kernels/nchw/scale.cl
+++ b/src/core/CL/cl_kernels/nchw/scale.cl
@@ -22,7 +22,7 @@
* SOFTWARE.
*/
#include "helpers.h"
-#include "warp_helpers.h"
+#include "tile_helpers.h"
/** Transforms four 2D coordinates. This is used to map the output coordinates to the input coordinates.
*
@@ -87,28 +87,55 @@
* @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_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_nchw(
IMAGE_DECLARATION(in),
- IMAGE_DECLARATION(out),
- const float input_width,
- const float input_height,
- const float scale_x,
- const float scale_y)
+ IMAGE_DECLARATION(out))
{
- 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);
+ const int x = get_global_id(0);
+ const int y = get_global_id(1);
+
+ float8 transformed = transform_nearest((float2)(x * VEC_SIZE, y), (float2)(SCALE_X, SCALE_Y));
#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);
+
+ TILE(SELECT_DATA_TYPE(DATA_TYPE), 1, 4, cond);
+ cond[0].v = CONVERT(((transformed.even < 0) || (transformed.even >= (int)SRC_WIDTH)) || ((transformed.odd < 0) || (transformed.odd >= (int)SRC_HEIGHT)), SELECT_VEC_DATA_TYPE(DATA_TYPE, 4));
+
+ TILE(int, 1, 4, in_x);
+ TILE(int, 1, 4, in_y);
+ in_x[0].v = convert_int4(clamp(transformed.even, 0.f, SRC_WIDTH - 1.f));
+ in_y[0].v = convert_int4(clamp(transformed.odd, 0.f, SRC_HEIGHT - 1.f));
+
+ TILE(DATA_TYPE, 1, VEC_SIZE, out_vals);
+ LOOP_UNROLLING(int, i, 0, 1, VEC_SIZE,
+ {
+ out_vals[0].s[i] = select(*((__global DATA_TYPE *)(in_ptr + in_offset_first_element_in_bytes + in_x[0].s[i] * sizeof(DATA_TYPE) + in_y[0].s[i] * in_stride_y)), (DATA_TYPE)CONSTANT_VALUE, cond[0].s[i]);
+ })
+
+ __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x * out_step_x + y * out_stride_y;
+
+ if(x == get_global_size(0) - 1)
+ {
+#if VEC_SIZE == 1
+ VSTORE_PARTIAL(VEC_SIZE, VEC_SIZE_LEFTOVER)
+ (out_vals[0].s[0], 0, (__global DATA_TYPE *)out_addr);
+#else // VEC_SIZE == 1
+ VSTORE_PARTIAL(VEC_SIZE, VEC_SIZE_LEFTOVER)
+ (out_vals[0].v, 0, (__global DATA_TYPE *)out_addr);
+#endif // VEC_SIZE == 1
+ }
+ else
+ {
+#if VEC_SIZE == 1
+ VSTORE(VEC_SIZE)
+ (out_vals[0].s[0], 0, (__global DATA_TYPE *)out_addr);
+#else // VEC_SIZE == 1
+ VSTORE(VEC_SIZE)
+ (out_vals[0].v, 0, (__global DATA_TYPE *)out_addr);
+#endif // VEC_SIZE == 1
+ }
}
/** Performs an affine transformation on an image interpolating with the BILINEAR method.
@@ -127,22 +154,118 @@
* @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_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_nchw(
IMAGE_DECLARATION(in),
- IMAGE_DECLARATION(out),
- const float input_width,
- const float input_height,
- const float scale_x,
- const float scale_y)
+ IMAGE_DECLARATION(out))
{
- 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 = 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);
+ const int x = get_global_id(0);
+ const int y = get_global_id(1);
+
+ TILE(float, 1, 8, trans_coords);
+ TILE(float, 1, 8, floor_coords);
+ TILE(int, 1, 16, in_x);
+ TILE(int, 1, 16, in_y);
+
+ trans_coords[0].v = transform_bilinear((float2)(x * VEC_SIZE, y), (float2)(SCALE_X, SCALE_Y));
+ floor_coords[0].v = floor(trans_coords[0].v);
+
+ LOOP_UNROLLING(int, i, 0, 1, 4,
+ {
+ LOOP_UNROLLING(int, j, 0, 1, 4,
+ {
+ in_x[0].s[i * 4 + j] = floor_coords[0].s[i * 2 + 0] + (j % 2);
+ in_y[0].s[i * 4 + j] = floor_coords[0].s[i * 2 + 1] + (j > 1);
+ })
+ })
+
+#if defined(BORDER_MODE_CONSTANT)
+ TILE(SELECT_DATA_TYPE(DATA_TYPE), 1, 16, cond);
+ cond[0].v = CONVERT(((in_x[0].v < 0) || (in_x[0].v >= (int)SRC_WIDTH)) || ((in_y[0].v < 0) || (in_y[0].v >= (int)SRC_HEIGHT)), SELECT_VEC_DATA_TYPE(DATA_TYPE, 16));
+#endif // defined(BORDER_MODE_CONSTANT)
+
+ in_x[0].v = clamp(in_x[0].v, 0, (int16)((int)SRC_WIDTH - 1));
+ in_y[0].v = clamp(in_y[0].v, 0, (int16)((int)SRC_HEIGHT - 1));
+
+ TILE(DATA_TYPE, 1, 16, in_vals);
+
+ // Loads the values from the input image
+#if defined(BORDER_MODE_CONSTANT)
+ LOOP_UNROLLING(int, i, 0, 1, 16,
+ {
+ in_vals[0].s[i] = select(*((__global DATA_TYPE *)(in_ptr + in_offset_first_element_in_bytes + in_x[0].s[i] * sizeof(DATA_TYPE) + in_y[0].s[i] * (int)in_stride_y)), (DATA_TYPE)CONSTANT_VALUE, cond[0].s[i]);
+ })
+#else // defined(BORDER_MODE_CONSTANT)
+ LOOP_UNROLLING(int, i, 0, 1, 16,
+ {
+ in_vals[0].s[i] = *((__global DATA_TYPE *)(in_ptr + in_offset_first_element_in_bytes + in_x[0].s[i] * sizeof(DATA_TYPE) + in_y[0].s[i] * (int)in_stride_y));
+ })
+#endif // defined(BORDER_MODE_CONSTANT)
+
+ TILE(float, 1, 8, a);
+ TILE(float, 1, 8, b);
+
+ a[0].v = trans_coords[0].v - floor_coords[0].v;
+ b[0].v = ((float8)(1.f)) - a[0].v;
+
+#if defined(OFFSET) && defined(SCALE)
+ TILE(float, 1, 16, in_vals_f32);
+ TILE(float, 1, 4, out_vals_f32);
+
+ in_vals_f32[0].v = convert_float16(convert_int16(in_vals[0].v) - (int16)OFFSET) * (float16)SCALE;
+
+ // Bilinear interpolation: (in0 * b0 * b1) + (in1 * a0 * b1) + (in2 * b0 * a1) + (in3 * a0 * a1)
+ // (in4 * b2 * b3) + (in5 * a2 * b3) + (in6 * b2 * a3) + (in7 * a2 * a3)
+ // (in8 * b4 * b5) + (in9 * a4 * b5) + (in10 * b4 * a5) + (in11 * a4 * a5)
+ // (in12 * b6 * b7) + (in13 * a6 * b7) + (in14 * b6 * a7) + (in15 * a6 * a7)
+ LOOP_UNROLLING(int, i, 0, 1, 4,
+ {
+ out_vals_f32[0].s[i] = (in_vals_f32[0].s[i * 4 + 0] * b[0].s[i * 2] * b[0].s[i * 2 + 1]) + (in_vals_f32[0].s[i * 4 + 1] * a[0].s[i * 2] * b[0].s[i * 2 + 1]) + (in_vals_f32[0].s[i * 4 + 2] * b[0].s[i * 2] * a[0].s[i * 2 + 1]) + (in_vals_f32[0].s[i * 4 + 3] * a[0].s[i * 2] * a[0].s[i * 2 + 1]);
+ })
+
+ TILE(DATA_TYPE, 1, 4, out_vals_4);
+ TILE(DATA_TYPE, 1, VEC_SIZE, out_vals);
+
+ out_vals_4[0].v = CONVERT_SAT(convert_int4_sat_rtp(out_vals_f32[0].v / (float)SCALE) + OFFSET, VEC_DATA_TYPE(DATA_TYPE, 4));
+
+ LOOP_UNROLLING(int, i, 0, 1, VEC_SIZE,
+ {
+ out_vals[0].s[i] = out_vals_4[0].s[i];
+ })
+#else // defined(OFFSET) && defined(SCALE)
+
+ TILE(DATA_TYPE, 1, VEC_SIZE, out_vals);
+
+ // Bilinear interpolation: (in0 * b0 * b1) + (in1 * a0 * b1) + (in2 * b0 * a1) + (in3 * a0 * a1)
+ // (in4 * b2 * b3) + (in5 * a2 * b3) + (in6 * b2 * a3) + (in7 * a2 * a3)
+ // (in8 * b4 * b5) + (in9 * a4 * b5) + (in10 * b4 * a5) + (in11 * a4 * a5)
+ // (in12 * b6 * b7) + (in13 * a6 * b7) + (in14 * b6 * a7) + (in15 * a6 * a7)
+ LOOP_UNROLLING(int, i, 0, 1, VEC_SIZE,
+ {
+ out_vals[0].s[i] = (in_vals[0].s[i * 4 + 0] * b[0].s[i * 2] * b[0].s[i * 2 + 1]) + (in_vals[0].s[i * 4 + 1] * a[0].s[i * 2] * b[0].s[i * 2 + 1]) + (in_vals[0].s[i * 4 + 2] * b[0].s[i * 2] * a[0].s[i * 2 + 1]) + (in_vals[0].s[i * 4 + 3] * a[0].s[i * 2] * a[0].s[i * 2 + 1]);
+ })
+#endif // defined(OFFSET) && defined(SCALE)
+
+ __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x * out_step_x + y * out_stride_y;
+
+ if(x == get_global_size(0) - 1)
+ {
+#if VEC_SIZE == 1
+ VSTORE_PARTIAL(VEC_SIZE, VEC_SIZE_LEFTOVER)
+ (out_vals[0].s[0], 0, (__global DATA_TYPE *)out_addr);
+#else // VEC_SIZE == 1
+ VSTORE_PARTIAL(VEC_SIZE, VEC_SIZE_LEFTOVER)
+ (out_vals[0].v, 0, (__global DATA_TYPE *)out_addr);
+#endif // VEC_SIZE == 1
+ }
+ else
+ {
+#if VEC_SIZE == 1
+ VSTORE(VEC_SIZE)
+ (out_vals[0].s[0], 0, (__global DATA_TYPE *)out_addr);
+#else // VEC_SIZE == 1
+ VSTORE(VEC_SIZE)
+ (out_vals[0].v, 0, (__global DATA_TYPE *)out_addr);
+#endif // VEC_SIZE == 1
+ }
}
\ No newline at end of file