diff options
| author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2018-03-02 09:43:54 +0000 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:49:37 +0000 |
| commit | 4d33630096c769dd43716dd5607f151e3d5abef7 (patch) | |
| tree | 762897c2acac9553c0dad688d0c21842c8edff16 /src/core/CL/cl_kernels/batchnormalization_layer.cl | |
| parent | 1cd41495153c4e89d6195b42f870967339c1a13b (diff) | |
| download | ComputeLibrary-4d33630096c769dd43716dd5607f151e3d5abef7.tar.gz | |
COMPMID-987: Make beta and gamma optional in BatchNormalization
Currently we have beta and gamma compulsory in Batch normalization. There are
network that might not need one or both of those. Thus these should be optional
with beta(offset) defaulting to zero and gamma(scale) to 1. Will also reduce
some memory requirements.
Change-Id: I15bf1ec14b814be2acebf1be1a4fba9c4fbd3190
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/123237
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/batchnormalization_layer.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/batchnormalization_layer.cl | 35 |
1 files changed, 26 insertions, 9 deletions
diff --git a/src/core/CL/cl_kernels/batchnormalization_layer.cl b/src/core/CL/cl_kernels/batchnormalization_layer.cl index 0b61b5638c..29b62d3d92 100644 --- a/src/core/CL/cl_kernels/batchnormalization_layer.cl +++ b/src/core/CL/cl_kernels/batchnormalization_layer.cl @@ -93,8 +93,12 @@ __kernel void batchnormalization_layer(TENSOR3D_DECLARATION(input), #endif /* not IN_PLACE */ VECTOR_DECLARATION(mean), VECTOR_DECLARATION(var), +#ifndef USE_DEFAULT_BETA VECTOR_DECLARATION(beta), +#endif /* USE_DEFAULT_BETA */ +#ifndef USE_DEFAULT_GAMMA VECTOR_DECLARATION(gamma), +#endif /* USE_DEFAULT_GAMMA */ float epsilon) { Tensor3D in = CONVERT_TO_TENSOR3D_STRUCT(input); @@ -103,10 +107,14 @@ __kernel void batchnormalization_layer(TENSOR3D_DECLARATION(input), #else /* IN_PLACE */ Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output); #endif /* IN_PLACE */ - Vector mean = CONVERT_TO_VECTOR_STRUCT(mean); - Vector var = CONVERT_TO_VECTOR_STRUCT(var); - Vector beta = CONVERT_TO_VECTOR_STRUCT(beta); + Vector mean = CONVERT_TO_VECTOR_STRUCT(mean); + Vector var = CONVERT_TO_VECTOR_STRUCT(var); +#ifndef USE_DEFAULT_BETA + Vector beta = CONVERT_TO_VECTOR_STRUCT(beta); +#endif /* USE_DEFAULT_BETA */ +#ifndef USE_DEFAULT_GAMMA Vector gamma = CONVERT_TO_VECTOR_STRUCT(gamma); +#endif /* USE_DEFAULT_GAMMA */ VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) data = 0; @@ -117,9 +125,7 @@ __kernel void batchnormalization_layer(TENSOR3D_DECLARATION(input), VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) x_bar = 0; VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - gamma_vec = 0; - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - beta_vec = 0; + res = 0; const int current_slice = get_global_id(2); @@ -132,11 +138,22 @@ __kernel void batchnormalization_layer(TENSOR3D_DECLARATION(input), numerator = SUB_OP(data, numerator); x_bar = MUL_OP(numerator, denominator); +#ifndef USE_DEFAULT_GAMMA + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) gamma_vec = *((__global DATA_TYPE *)(gamma.ptr + current_slice * gamma.stride_x)); - beta_vec = *((__global DATA_TYPE *)(beta.ptr + current_slice * beta.stride_x)); + res = MUL_OP(gamma_vec, x_bar); +#else /* USE_DEFAULT_GAMMA */ + // gamma is equal to 1, no need to perform multiplications + res = x_bar; +#endif /* USE_DEFAULT_GAMMA */ + +#ifndef USE_DEFAULT_BETA VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - res = ADD_OP(MUL_OP(gamma_vec, x_bar), beta_vec); + beta_vec = *((__global DATA_TYPE *)(beta.ptr + current_slice * beta.stride_x)); + // beta is not zero, hence we need to perform the addition + res = ADD_OP(res, beta_vec); +#endif /* USE_DEFAULT_BETA */ res = ACTIVATION_FUNC(res); @@ -144,4 +161,4 @@ __kernel void batchnormalization_layer(TENSOR3D_DECLARATION(input), (res, 0, (__global DATA_TYPE *)out.ptr); } -#endif /* defined(VEC_SIZE) && defined(DATA_TYPE) */
\ No newline at end of file +#endif /* defined(VEC_SIZE) && defined(DATA_TYPE) */ |