commit 4d33630096c769dd43716dd5607f151e3d5abef7
author Michele Di Giorgio <michele.digiorgio@arm.com> Fri Mar 02 09:43:54 2018 +0000
committer Anthony Barbier <anthony.barbier@arm.com> Fri Nov 02 16:49:37 2018 +0000
tree 762897c2acac9553c0dad688d0c21842c8edff16
parent 1cd41495153c4e89d6195b42f870967339c1a13b

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>

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 0b61b56..29b62d3 100644
--- a/src/core/CL/cl_kernels/batchnormalization_layer.cl
+++ b/src/core/CL/cl_kernels/batchnormalization_layer.cl
@@ -93,8 +93,12 @@
 #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 @@
 #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 @@
     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 @@
     numerator = SUB_OP(data, numerator);
     x_bar     = MUL_OP(numerator, denominator);
 
-    gamma_vec = *((__global DATA_TYPE *)(gamma.ptr + current_slice * gamma.stride_x));
-    beta_vec  = *((__global DATA_TYPE *)(beta.ptr + current_slice * beta.stride_x));
-
+#ifndef USE_DEFAULT_GAMMA
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
-    res = ADD_OP(MUL_OP(gamma_vec, x_bar), beta_vec);
+    gamma_vec = *((__global DATA_TYPE *)(gamma.ptr + current_slice * gamma.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)
+    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 @@
     (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) */