aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/scale.cl
diff options
context:
space:
mode:
authorsteniu01 <steven.niu@arm.com>2017-09-11 15:29:12 +0100
committerAnthony Barbier <anthony.barbier@arm.com>2018-11-02 16:35:24 +0000
commitf81652dc970a8071f41c4138508c39684ed9111a (patch)
tree6bbda484b46784f1e9ddf132849c15d16ee61230 /src/core/CL/cl_kernels/scale.cl
parent9fe414430c3c989b1cdc79d41e031495aed2cb7c (diff)
downloadComputeLibrary-f81652dc970a8071f41c4138508c39684ed9111a.tar.gz
COMPMID-516 Increase tolerance rate of Scale, Conv, fully connected and GEMM
This patch also fix the scale kernel issue where it was calcuated the scale factor inside the gpu but now in the CPU. The GPU and CPU gave different result for simple float division operation Change-Id: Ib6709cb6c41dcf4fc0fa4eb79e481430695bf40e Reviewed-on: http://mpd-gerrit.cambridge.arm.com/87266 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Moritz Pflanzer <moritz.pflanzer@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/scale.cl')
-rw-r--r--src/core/CL/cl_kernels/scale.cl20
1 files changed, 10 insertions, 10 deletions
diff --git a/src/core/CL/cl_kernels/scale.cl b/src/core/CL/cl_kernels/scale.cl
index b3398bd11c..0106ce095c 100644
--- a/src/core/CL/cl_kernels/scale.cl
+++ b/src/core/CL/cl_kernels/scale.cl
@@ -70,20 +70,20 @@ inline const float8 transform_bilinear(const float2 coord, const float2 scale)
* @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] output_width Output image width
- * @param[in] output_height Output 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(
IMAGE_DECLARATION(in),
IMAGE_DECLARATION(out),
const float input_width,
const float input_height,
- const float output_width,
- const float output_height)
+ 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)(input_width / output_width, input_height / output_height);
+ const float2 r = (float2)(scale_x, scale_y);
const float8 tc = clamp_to_border(transform_nearest(get_current_coords(), r), input_width, input_height);
vstore4(read_texels4(&in, convert_int8(tc)), 0, (__global DATA_TYPE *)out.ptr);
}
@@ -104,20 +104,20 @@ __kernel void scale_nearest_neighbour(
* @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] output_width Output image width
- * @param[in] output_height Output 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(
IMAGE_DECLARATION(in),
IMAGE_DECLARATION(out),
const float input_width,
const float input_height,
- const float output_width,
- const float output_height)
+ 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)(input_width / output_width, input_height / output_height);
+ 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);
}