diff options
| author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2018-07-06 10:17:33 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:54:10 +0000 |
| commit | d304e80112e5d246b4bb884d1211eef9381bf034 (patch) | |
| tree | 456943f81602b2df7ae9909d51a942dc3f15cc97 /src/core/CL/cl_kernels/activation_layer_qa8.cl | |
| parent | d93991e290618a685b67506c78090350e6aee43f (diff) | |
| download | ComputeLibrary-d304e80112e5d246b4bb884d1211eef9381bf034.tar.gz | |
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 <bsgcomp@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/activation_layer_qa8.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/activation_layer_qa8.cl | 13 |
1 files changed, 12 insertions, 1 deletions
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) */ |