diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2018-02-07 15:38:12 +0000 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:47:18 +0000 |
| commit | 1167487ea8e54a76d0a3625e0aa84e2ad9ffd317 (patch) | |
| tree | 287dbc45e895c6b637fecc692c04bd4ae59580ae /src/core/CL/cl_kernels/batchnormalization_layer.cl | |
| parent | 4e1e7dcd581adecd5ad9c0f9503fc3c43f8222ef (diff) | |
| download | ComputeLibrary-1167487ea8e54a76d0a3625e0aa84e2ad9ffd317.tar.gz | |
COMPMID-897 Merge batch normalization with bounded relu
Change-Id: I9a607fe620f795cdea1a99fdd3f5f8c2fc76f980
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/119234
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/batchnormalization_layer.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/batchnormalization_layer.cl | 21 |
1 files changed, 20 insertions, 1 deletions
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 |