diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2020-02-11 17:21:31 +0000 |
|---|---|---|
| committer | SiCong Li <sicong.li@arm.com> | 2020-04-03 08:51:12 +0000 |
| commit | 8b2a7d3aa119e7f1d6a03690d05eb27c5d178b9f (patch) | |
| tree | 9fb4f4f328f7a17de13bef109834e8ad8a21d2ee /src/core/CL/cl_kernels/elementwise_operation.cl | |
| parent | 15e4d876643c37e1db36ee1190ec52319479ffaf (diff) | |
| download | ComputeLibrary-8b2a7d3aa119e7f1d6a03690d05eb27c5d178b9f.tar.gz | |
COMPMID-3101 Fuse activation with floating point elementwise operation layers in CL
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Change-Id: I1693f8664ba7c0dc8c076bbe7365cef1e667bd25
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2718
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/elementwise_operation.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/elementwise_operation.cl | 12 |
1 files changed, 11 insertions, 1 deletions
diff --git a/src/core/CL/cl_kernels/elementwise_operation.cl b/src/core/CL/cl_kernels/elementwise_operation.cl index 42d6d33e03..9b87b526f7 100644 --- a/src/core/CL/cl_kernels/elementwise_operation.cl +++ b/src/core/CL/cl_kernels/elementwise_operation.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -44,6 +44,11 @@ #define OP_FUN_NAME(op) OP_FUN_NAME_STR(op) #if defined(OP) && defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_OUT) && defined(VEC_SIZE) + +#if defined(ACTIVATION_TYPE) +#include "activation_float_helpers.h" +#endif // defined(ACTIVATION_TYPE) + /** This function executes an element-wise operation among two tensors. * * @attention The input and output data_types need to be passed at compile time using -DDATA_TYPE_IN1, -DDATA_TYPE_IN2 and -DDATA_TYPE_OUT: @@ -94,7 +99,12 @@ __kernel void OP_FUN_NAME(OP)( in_b = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)); // Calculate and store result +#if defined(ACTIVATION_TYPE) + VSTORE(VEC_SIZE) + (ACTIVATION(ACTIVATION_TYPE, DATA_TYPE_OUT, CONVERT(OP(in_a, in_b), VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)), A_VAL, B_VAL), 0, (__global DATA_TYPE_OUT *)out.ptr); +#else // defined(ACTIVATION_TYPE) VSTORE(VEC_SIZE) (OP(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr); +#endif // defined(ACTIVATION_TYPE) } #endif /* defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_OUT) && defined(VEC_SIZE) */ |