From e2428a013014feac81af60f65f9b16cc244327aa Mon Sep 17 00:00:00 2001 From: Usama Arif Date: Thu, 9 May 2019 11:03:17 +0100 Subject: COMPMID-2194: Refactor activation function macro in OpenCL. New Macros for activation. This commit contains the new macros for activation. Only the activation_layer utilizes the new macros in this commit. Change-Id: I2fa8567cc876e8cb67a1e876652bc348b7ed23ea Signed-off-by: Usama Arif Reviewed-on: https://review.mlplatform.org/c/1104 Comments-Addressed: Arm Jenkins Reviewed-by: Gian Marco Iodice Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/activation_layer.cl | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) (limited to 'src/core/CL/cl_kernels/activation_layer.cl') diff --git a/src/core/CL/cl_kernels/activation_layer.cl b/src/core/CL/cl_kernels/activation_layer.cl index cf1f434972..d820753b57 100644 --- a/src/core/CL/cl_kernels/activation_layer.cl +++ b/src/core/CL/cl_kernels/activation_layer.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -21,12 +21,11 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ -#define TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) -#define SELECT_TYPE VEC_DATA_TYPE(SELECT_DATA_TYPE, VEC_SIZE) +#if defined(ACT) && defined(DATA_TYPE) && defined(VEC_SIZE) -#include "activation_helpers.h" +#define TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) -#if defined(ACT) +#include "activation_float_helpers.h" /** This performs an activation function floating point inputs. * @@ -74,11 +73,11 @@ __kernel void activation_layer( TYPE data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input.ptr); // Perform activation - data = ACTIVATION_OP(ACT, data); + data = ACTIVATION(ACT, DATA_TYPE, data, A_VAL, B_VAL); // Store result VSTORE(VEC_SIZE) (data, 0, (__global DATA_TYPE *)output.ptr); } -#endif /* defined(ACT) */ \ No newline at end of file +#endif /* defined(ACT) */ -- cgit v1.2.1