From c8097129aca9e410ec1fe0ff8f568db7a89b7cd2 Mon Sep 17 00:00:00 2001 From: Rob Hughes Date: Thu, 23 Nov 2017 14:28:54 +0000 Subject: COMPMID-556 Add saturation to 8-bit activation. This prevents undefined overflow from occurring when an output value cannot be expressed in the output space Change-Id: I4045a44ee257a829f6cfe708dcbad37bb91c0ee4 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/110393 Reviewed-by: Georgios Pinitas Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com --- src/core/CL/cl_kernels/activation_layer_qa8.cl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'src/core') diff --git a/src/core/CL/cl_kernels/activation_layer_qa8.cl b/src/core/CL/cl_kernels/activation_layer_qa8.cl index 4d9bf0efad..910a93fdc1 100644 --- a/src/core/CL/cl_kernels/activation_layer_qa8.cl +++ b/src/core/CL/cl_kernels/activation_layer_qa8.cl @@ -91,8 +91,8 @@ __kernel void activation_layer_qa8( // requantize to output space float16 fdata = convert_float16(data); - fdata = round((fdata - O1_VAL) * (S1_VAL / S2_VAL) + O2_VAL); - uchar16 qdata = convert_uchar16(fdata); + fdata = round((fdata - (float)O1_VAL) * ((float)S1_VAL / (float)S2_VAL) + (float)O2_VAL); + uchar16 qdata = convert_uchar16_sat(fdata); // Store result VSTORE(VEC_SIZE) -- cgit v1.2.1