diff options
| author | Manuel Bottini <manuel.bottini@arm.com> | 2019-10-03 12:03:19 +0100 |
|---|---|---|
| committer | Manuel Bottini <manuel.bottini@arm.com> | 2019-10-07 12:52:40 +0000 |
| commit | 6665f82b38ce904aa588230546d66e65d38f20d6 (patch) | |
| tree | 360cddcf65dceed0c5408f7242bb58e9e307d7aa /src/core/CL/cl_kernels/instance_normalization.cl | |
| parent | cce2ea6194f3e781cc9eb86c6d2bd50be0c97b1e (diff) | |
| download | ComputeLibrary-6665f82b38ce904aa588230546d66e65d38f20d6.tar.gz | |
COMPMID-2452: Mismatches FP16 in CL InstanceNormalization
Change-Id: I8a28557370be50b2ccc9534feb9fc552c6ee5cf0
Signed-off-by: Manuel Bottini <manuel.bottini@arm.com>
Reviewed-on: https://review.mlplatform.org/c/2037
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/instance_normalization.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/instance_normalization.cl | 14 |
1 files changed, 7 insertions, 7 deletions
diff --git a/src/core/CL/cl_kernels/instance_normalization.cl b/src/core/CL/cl_kernels/instance_normalization.cl index 699597e8a8..18afcc53a7 100644 --- a/src/core/CL/cl_kernels/instance_normalization.cl +++ b/src/core/CL/cl_kernels/instance_normalization.cl @@ -58,8 +58,8 @@ __kernel void instance_normalization( #endif /* IN_PLACE */ ) { - DATA_TYPE sum = 0.f; - DATA_TYPE sum_sq = 0.f; + float sum = 0.f; + float sum_sq = 0.f; #if defined(NHWC) @@ -73,7 +73,7 @@ __kernel void instance_normalization( { for(int i_h = 0; i_h < DIM_Z; ++i_h) { - DATA_TYPE data = *((__global DATA_TYPE *)input_ptr + pc + i_w * DIM_X + i_h * elements_x_y + pn * elements_x_y_z); + float data = (float)*((__global DATA_TYPE *)input_ptr + pc + i_w * DIM_X + i_h * elements_x_y + pn * elements_x_y_z); sum += data; sum_sq += data * data; } @@ -120,14 +120,14 @@ __kernel void instance_normalization( part_sum_sq.s0 += data * data; } - sum = part_sum.s0; - sum_sq = part_sum_sq.s0; + sum = (float)part_sum.s0; + sum_sq = (float)part_sum_sq.s0; #endif // defined(NHWC) - const DATA_TYPE mean_float = ((float)sum / elements_plane); + const float mean_float = (sum / elements_plane); const DATA_TYPE mean = (DATA_TYPE)mean_float; - const float var_float = ((float)sum_sq / elements_plane) - (mean_float * mean_float); + const float var_float = (sum_sq / elements_plane) - (mean_float * mean_float); const float multip_float = GAMMA / sqrt(var_float + EPSILON); const DATA_TYPE multip = (DATA_TYPE)multip_float; |