From 1167487ea8e54a76d0a3625e0aa84e2ad9ffd317 Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Wed, 7 Feb 2018 15:38:12 +0000 Subject: COMPMID-897 Merge batch normalization with bounded relu Change-Id: I9a607fe620f795cdea1a99fdd3f5f8c2fc76f980 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/119234 Tested-by: Jenkins Reviewed-by: Gian Marco Iodice Reviewed-by: Georgios Pinitas --- src/core/CL/cl_kernels/batchnormalization_layer.cl | 21 ++++++++++++++++++++- 1 file changed, 20 insertions(+), 1 deletion(-) (limited to 'src/core/CL/cl_kernels/batchnormalization_layer.cl') diff --git a/src/core/CL/cl_kernels/batchnormalization_layer.cl b/src/core/CL/cl_kernels/batchnormalization_layer.cl index fbffefb3c0..5ddeb1a6a1 100644 --- a/src/core/CL/cl_kernels/batchnormalization_layer.cl +++ b/src/core/CL/cl_kernels/batchnormalization_layer.cl @@ -23,6 +23,8 @@ */ #include "helpers.h" +#if defined(VEC_SIZE) && defined(DATA_TYPE) + #if defined(FIXED_POINT_POSITION) #include "fixed_point.h" @@ -42,6 +44,16 @@ #endif /* FIXED_POINT_POSITION */ +#if defined(LU_BRELU) +#define ACTIVATION_FUNC(x) CLAMP(x, (DATA_TYPE)B_VAL, (DATA_TYPE)A_VAL) +#elif defined(BRELU) +#define ACTIVATION_FUNC(x) CLAMP(x, (DATA_TYPE)0, (DATA_TYPE)A_VAL) +#elif defined(RELU) +#define ACTIVATION_FUNC(x) max(x, (DATA_TYPE)0) +#else /* FUSED_ACT */ +#define ACTIVATION_FUNC(x) (x) +#endif /* FUSED_ACT */ + /** Apply batch normalization. * * @param[in] input_ptr Pointer to the first source tensor. Supported data types: QS8/QS16/F16/F32 @@ -126,6 +138,13 @@ __kernel void batchnormalization_layer(TENSOR3D_DECLARATION(input), gamma_vec = *((__global DATA_TYPE *)(gamma.ptr + current_slice * gamma.stride_x)); beta_vec = *((__global DATA_TYPE *)(beta.ptr + current_slice * beta.stride_x)); + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + res = ADD_OP(MUL_OP(gamma_vec, x_bar), beta_vec); + + res = ACTIVATION_FUNC(res); + VSTORE(VEC_SIZE) - (ADD_OP(MUL_OP(gamma_vec, x_bar), beta_vec), 0, (__global DATA_TYPE *)out.ptr); + (res, 0, (__global DATA_TYPE *)out.ptr); } + +#endif /* defined(VEC_SIZE) && defined(DATA_TYPE) */ \ No newline at end of file -- cgit v1.2.1