26 #define ADD_OP(a, b) ((a) + (b))
27 #define SUB_OP(a, b) ((a) - (b))
28 #define MUL_OP(a, b) ((a) * (b))
29 #define INVSQRT_OP(a) rsqrt((a))
30 #define SQCVT_SAT(a) (a)
32 #if defined(VEC_SIZE) && defined(DATA_TYPE) && defined(ACTIVATION_TYPE)
80 #ifndef USE_DEFAULT_BETA
83 #ifndef USE_DEFAULT_GAMMA
96 #ifndef USE_DEFAULT_BETA
99 #ifndef USE_DEFAULT_GAMMA
114 const int current_slice = get_global_id(2);
117 denominator = *((__global DATA_TYPE *)(var.
ptr + current_slice * var.
stride_x));
121 numerator = *((__global DATA_TYPE *)(mean.
ptr + current_slice * mean.
stride_x));
122 numerator =
SUB_OP(data, numerator);
123 x_bar =
MUL_OP(numerator, denominator);
125 #ifndef USE_DEFAULT_GAMMA
127 gamma_vec = *((__global DATA_TYPE *)(gamma.
ptr + current_slice * gamma.
stride_x));
129 res =
MUL_OP(gamma_vec, x_bar);
135 #ifndef USE_DEFAULT_BETA
137 beta_vec = *((__global DATA_TYPE *)(beta.
ptr + current_slice * beta.
stride_x));
139 res =
ADD_OP(res, beta_vec);
145 (res, 0, (__global DATA_TYPE *)out.
ptr);