From d304e80112e5d246b4bb884d1211eef9381bf034 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Fri, 6 Jul 2018 10:17:33 +0100 Subject: COMPMID-1349: Add support for QASYMM8 LOGISTIC activation in CLActivationLayer Change-Id: Ibabce61cf5427de80078a6468023bed05f5e7c2c Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/139006 Tested-by: Jenkins Reviewed-by: Anthony Barbier Reviewed-by: Georgios Pinitas --- src/core/CL/cl_kernels/activation_layer_qa8.cl | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) (limited to 'src/core/CL/cl_kernels') diff --git a/src/core/CL/cl_kernels/activation_layer_qa8.cl b/src/core/CL/cl_kernels/activation_layer_qa8.cl index 66e54ed6ad..8f6a807613 100644 --- a/src/core/CL/cl_kernels/activation_layer_qa8.cl +++ b/src/core/CL/cl_kernels/activation_layer_qa8.cl @@ -24,7 +24,18 @@ #include "helpers.h" #define TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) +#define VEC_FLOAT VEC_DATA_TYPE(float, VEC_SIZE) +// Logistic Activation +inline TYPE logistic_op(TYPE x) +{ + VEC_FLOAT x_flt = CONVERT(x, VEC_FLOAT); + x_flt = round(x_flt - (float)O1_VAL) * ((float)S1_VAL); + x_flt = 1.f / (1.f + exp(-x_flt)); + + const TYPE x_u8 = CONVERT_SAT(round(x_flt / ((float)S1_VAL)) + (float)O1_VAL, TYPE); + return x_u8; +} // RELU Activation inline TYPE relu_op(TYPE x) { @@ -119,4 +130,4 @@ __kernel void activation_layer_qa8( (data, 0, (__global DATA_TYPE *)output.ptr); } -#endif /* defined(ACT) */ \ No newline at end of file +#endif /* defined(ACT) */ -- cgit v1.2.1