diff options
| author | SiCong Li <sicong.li@arm.com> | 2017-07-27 17:59:20 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-09-17 14:16:42 +0100 |
| commit | 02dfb2cd0bfab1bcff141db4598c23055e67311b (patch) | |
| tree | 033efc87dc1127d9d6cbf07a0e490392ef439637 /src/core/CL/cl_kernels | |
| parent | a36ccf1b5c66a5f25bdde1ce69b60e49ed67bb41 (diff) | |
| download | ComputeLibrary-02dfb2cd0bfab1bcff141db4598c23055e67311b.tar.gz | |
COMPMID-457 Fix F16 NormalizationLayer CL kernel
Change-Id: I307c7ef6a49c852615c4425dc8dc0b1066a6974f
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/81895
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels')
| -rw-r--r-- | src/core/CL/cl_kernels/normalization_layer.cl | 2 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/pixelwise_mul_float.cl | 2 |
2 files changed, 2 insertions, 2 deletions
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); |