From bf3c6626e98b9e1be435fce9fdabc9d21f3b5b3a Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Thu, 8 Mar 2018 11:52:27 +0000 Subject: COMPMID-803: Add NHWC data format support for CL batch normalisation Change-Id: Ie37588f60b9cfc7b1d09b1e8628fcfb4b17e0717 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/123834 Tested-by: Jenkins Reviewed-by: Anthony Barbier --- src/core/CL/cl_kernels/batchnormalization_layer.cl | 125 +++++++++++++++++++-- 1 file changed, 117 insertions(+), 8 deletions(-) (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 29b62d3d92..9c980da62a 100644 --- a/src/core/CL/cl_kernels/batchnormalization_layer.cl +++ b/src/core/CL/cl_kernels/batchnormalization_layer.cl @@ -87,19 +87,19 @@ * @param[in] gamma_offset_first_element_in_bytes The offset of the first element in the gamma source tensor * @param[in] epsilon Epsilon parameter in the batch normalization equation */ -__kernel void batchnormalization_layer(TENSOR3D_DECLARATION(input), +__kernel void batchnormalization_layer_nchw(TENSOR3D_DECLARATION(input), #ifndef IN_PLACE - TENSOR3D_DECLARATION(output), + TENSOR3D_DECLARATION(output), #endif /* not IN_PLACE */ - VECTOR_DECLARATION(mean), - VECTOR_DECLARATION(var), + VECTOR_DECLARATION(mean), + VECTOR_DECLARATION(var), #ifndef USE_DEFAULT_BETA - VECTOR_DECLARATION(beta), + VECTOR_DECLARATION(beta), #endif /* USE_DEFAULT_BETA */ #ifndef USE_DEFAULT_GAMMA - VECTOR_DECLARATION(gamma), + VECTOR_DECLARATION(gamma), #endif /* USE_DEFAULT_GAMMA */ - float epsilon) + float epsilon) { Tensor3D in = CONVERT_TO_TENSOR3D_STRUCT(input); #ifdef IN_PLACE @@ -145,7 +145,7 @@ __kernel void batchnormalization_layer(TENSOR3D_DECLARATION(input), res = MUL_OP(gamma_vec, x_bar); #else /* USE_DEFAULT_GAMMA */ // gamma is equal to 1, no need to perform multiplications - res = x_bar; + res = x_bar; #endif /* USE_DEFAULT_GAMMA */ #ifndef USE_DEFAULT_BETA @@ -161,4 +161,113 @@ __kernel void batchnormalization_layer(TENSOR3D_DECLARATION(input), (res, 0, (__global DATA_TYPE *)out.ptr); } +/** Apply batch normalization on tensors with NHWC format. + * + * @param[in] input_ptr Pointer to the first source tensor. Supported data types: QS8/QS16/F16/F32 + * @param[in] input_stride_x Stride of the first source tensor in X dimension (in bytes) + * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input_stride_y Stride of the first source tensor in Y dimension (in bytes) + * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input_stride_z Stride of the first source tensor in Z dimension (in bytes) + * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] input_offset_first_element_in_bytes The offset of the first element in the first source tensor + * @param[out] output_ptr Pointer to the destination tensor. Supported data types: same as @p input_ptr + * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] mean_ptr Pointer to the mean source tensor. Supported data types: same as @p input_ptr + * @param[in] mean_stride_x Stride of the mean source tensor in X dimension (in bytes) + * @param[in] mean_step_x mean_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] mean_offset_first_element_in_bytes The offset of the first element in the mean source tensor + * @param[in] var_ptr Pointer to the var tensor. Supported data types: same as @p input_ptr + * @param[in] var_stride_x Stride of the var tensor in X dimension (in bytes) + * @param[in] var_step_x var_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] var_offset_first_element_in_bytes The offset of the first element in the var source tensor + * @param[in] beta_ptr Pointer to the beta source tensor. Supported data types: same as @p input_ptr + * @param[in] beta_stride_x Stride of the beta source tensor in X dimension (in bytes) + * @param[in] beta_step_x beta_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] beta_offset_first_element_in_bytes The offset of the first element in the beta source tensor + * @param[in] gamma_ptr Pointer to the gamma source tensor. Supported data types: same as @p input_ptr + * @param[in] gamma_stride_x Stride of the gamma source tensor in X dimension (in bytes) + * @param[in] gamma_step_x gamma_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] gamma_offset_first_element_in_bytes The offset of the first element in the gamma source tensor + * @param[in] epsilon Epsilon parameter in the batch normalization equation + */ +__kernel void batchnormalization_layer_nhwc(TENSOR3D_DECLARATION(input), +#ifndef IN_PLACE + TENSOR3D_DECLARATION(output), +#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); +#ifdef IN_PLACE + Tensor3D out = in; +#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); +#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; + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + denominator = 0; + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + numerator = 0; + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + x_bar = 0; + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + res = 0; + + const int current_slice = get_global_id(0); + + data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in.ptr); + denominator = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(var.ptr + current_slice * VEC_SIZE * var.stride_x)); + denominator = INVSQRT_OP(ADD_OP(denominator, ((VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))SQCVT_SAT(epsilon)))); + + // Calculate x bar and store results + numerator = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(mean.ptr + current_slice * VEC_SIZE * mean.stride_x)); + numerator = SUB_OP(data, numerator); + x_bar = MUL_OP(numerator, denominator); + +#ifndef USE_DEFAULT_GAMMA + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + gamma_vec = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(gamma.ptr + current_slice * VEC_SIZE * 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 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(beta.ptr + current_slice * VEC_SIZE * 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); + + VSTORE(VEC_SIZE) + (res, 0, (__global DATA_TYPE *)out.ptr); +} #endif /* defined(VEC_SIZE) && defined(DATA_TYPE) */ -- cgit v1.2.1