From 02dfb2cd0bfab1bcff141db4598c23055e67311b Mon Sep 17 00:00:00 2001 From: SiCong Li Date: Thu, 27 Jul 2017 17:59:20 +0100 Subject: COMPMID-457 Fix F16 NormalizationLayer CL kernel Change-Id: I307c7ef6a49c852615c4425dc8dc0b1066a6974f Reviewed-on: http://mpd-gerrit.cambridge.arm.com/81895 Reviewed-by: Anthony Barbier Tested-by: Kaizen --- src/core/CL/cl_kernels/normalization_layer.cl | 2 +- src/core/CL/cl_kernels/pixelwise_mul_float.cl | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) (limited to 'src') diff --git a/src/core/CL/cl_kernels/normalization_layer.cl b/src/core/CL/cl_kernels/normalization_layer.cl index 076b0d8909..2305ae0d15 100644 --- a/src/core/CL/cl_kernels/normalization_layer.cl +++ b/src/core/CL/cl_kernels/normalization_layer.cl @@ -146,7 +146,7 @@ __kernel void normalization_layer_in_map_1D(TENSOR3D_DECLARATION(input), acc_vec += vload4(0, (__global DATA_TYPE *)tensor3D_offset(&squared_in, i - current_pos, 0, 0)); } - const float4 normalized = pow((float4)kappa + coeff * (float4)acc_vec, beta); + const float4 normalized = pow((float4)kappa + coeff * CONVERT(acc_vec, float4), beta); const float4 normalized_pixel = CONVERT(vload4(0, (__global DATA_TYPE *)in.ptr), float4) / normalized; diff --git a/src/core/CL/cl_kernels/pixelwise_mul_float.cl b/src/core/CL/cl_kernels/pixelwise_mul_float.cl index 98127e0311..f4f36a02af 100644 --- a/src/core/CL/cl_kernels/pixelwise_mul_float.cl +++ b/src/core/CL/cl_kernels/pixelwise_mul_float.cl @@ -84,7 +84,7 @@ __kernel void pixelwise_mul_float( // Perform multiplication #ifdef DATA_TYPE_FLOAT VEC_DATA_TYPE(DATA_TYPE_OUT, 16) - res = CONVERT(in1_data * in2_data * scale, VEC_DATA_TYPE(DATA_TYPE_OUT, 16)); + res = CONVERT(in1_data * in2_data * (DATA_TYPE_RES)scale, VEC_DATA_TYPE(DATA_TYPE_OUT, 16)); #else /* DATA_TYPE_FLOAT */ VEC_DATA_TYPE(DATA_TYPE_OUT, 16) res = CONVERT_OP_FLOAT(CONVERT_OP_FLOAT((convert_float16(in1_data * in2_data) * scale), VEC_DATA_TYPE(DATA_TYPE_RES, 16), ROUND), VEC_DATA_TYPE(DATA_TYPE_OUT, 16), ROUND); -- cgit v1.2.1