diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2018-02-16 15:17:23 +0000 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:49:16 +0000 |
| commit | 99ac60bca77e9977c844cc1293751d63ddc3065c (patch) | |
| tree | 30867264809a61c817e5c543a6ac997620f4beaa /src/core/CL/cl_kernels/activation_layer_qa8.cl | |
| parent | 60954c671ffdc3422bbdb728fc022eb6896c1e17 (diff) | |
| download | ComputeLibrary-99ac60bca77e9977c844cc1293751d63ddc3065c.tar.gz | |
COMPMID-853 Fuse CL DepthwiseConvolution with Activation for QASYM8
Change-Id: I287908f76af458ad4b4d865d353dc37e33877250
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/120839
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 | 34 |
1 files changed, 23 insertions, 11 deletions
diff --git a/src/core/CL/cl_kernels/activation_layer_qa8.cl b/src/core/CL/cl_kernels/activation_layer_qa8.cl index cb31e99efb..66e54ed6ad 100644 --- a/src/core/CL/cl_kernels/activation_layer_qa8.cl +++ b/src/core/CL/cl_kernels/activation_layer_qa8.cl @@ -44,6 +44,26 @@ inline TYPE lu_brelu_op(TYPE x) #define ACTIVATION_OP2(op, x) op##_op(x) #define ACTIVATION_OP(op, x) ACTIVATION_OP2(op, x) +#if defined(O1_VAL) && defined(O2_VAL) && defined(S1_VAL) && defined(S2_VAL) +#define PERFORM_ACTIVATION_QA8(act, data) \ + ({ \ + data = ACTIVATION_OP(act, data); \ + \ + VEC_DATA_TYPE(float, VEC_SIZE) \ + fdata = CONVERT(data, VEC_DATA_TYPE(float, VEC_SIZE)); \ + \ + fdata = round((fdata - (float)O1_VAL) * ((float)S1_VAL / (float)S2_VAL) + (float)O2_VAL); \ + data = CONVERT_SAT(fdata, VEC_DATA_TYPE(uchar, VEC_SIZE)); \ + }) +#else /* defined(O1_VAL) && defined(O2_VAL) && defined(S1_VAL) && defined(S2_VAL) */ +#define PERFORM_ACTIVATION_QA8(act, data) \ + ({ \ + data = ACTIVATION_OP(act, data); \ + }) +#endif /* defined(O1_VAL) && defined(O2_VAL) && defined(S1_VAL) && defined(S2_VAL) */ + +#if defined(ACT) + /** This performs an activation function on QASYMM8 inputs. * * @note In order to perform the activation function "in-place", the pre-processor -DIN_PLACE must be passed at compile time @@ -92,19 +112,11 @@ __kernel void activation_layer_qa8( // Load data TYPE data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input.ptr); - // Perform activation - data = ACTIVATION_OP(ACT, data); - -#if defined(O1_VAL) && defined(O2_VAL) && defined(S1_VAL) && defined(S2_VAL) - // requantize to output space - VEC_DATA_TYPE(float, VEC_SIZE) - fdata = CONVERT(data, VEC_DATA_TYPE(float, VEC_SIZE)); - - fdata = round((fdata - (float)O1_VAL) * ((float)S1_VAL / (float)S2_VAL) + (float)O2_VAL); - data = CONVERT_SAT(fdata, VEC_DATA_TYPE(uchar, VEC_SIZE)); -#endif // defined(O1_VAL) && defined(O2_VAL) && defined(S1_VAL) && defined(S2_VAL) + data = PERFORM_ACTIVATION_QA8(ACT, data); // Store result VSTORE(VEC_SIZE) (data, 0, (__global DATA_TYPE *)output.ptr); } + +#endif /* defined(ACT) */
\ No newline at end of file |