diff options
| author | Rob Hughes <robert.hughes@arm.com> | 2017-11-23 14:28:54 +0000 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:41:04 +0000 |
| commit | c8097129aca9e410ec1fe0ff8f568db7a89b7cd2 (patch) | |
| tree | c3afb6fbbc4604c52f34fc5913975ba7e33b4244 | |
| parent | b7b31538eb9137e4d3e8de6d381dcbe9fc58df94 (diff) | |
| download | ComputeLibrary-c8097129aca9e410ec1fe0ff8f568db7a89b7cd2.tar.gz | |
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 <georgios.pinitas@arm.com>
Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/activation_layer_qa8.cl | 4 |
1 files changed, 2 insertions, 2 deletions
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) |