aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorManuel Bottini <manuel.bottini@arm.com>2019-10-03 12:03:19 +0100
committerManuel Bottini <manuel.bottini@arm.com>2019-10-07 12:52:40 +0000
commit6665f82b38ce904aa588230546d66e65d38f20d6 (patch)
tree360cddcf65dceed0c5408f7242bb58e9e307d7aa
parentcce2ea6194f3e781cc9eb86c6d2bd50be0c97b1e (diff)
downloadComputeLibrary-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>
-rw-r--r--src/core/CL/cl_kernels/instance_normalization.cl14
-rw-r--r--tests/validation/reference/InstanceNormalizationLayer.cpp19
2 files changed, 14 insertions, 19 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;
diff --git a/tests/validation/reference/InstanceNormalizationLayer.cpp b/tests/validation/reference/InstanceNormalizationLayer.cpp
index 0e5c02aa99..ad0ac1be68 100644
--- a/tests/validation/reference/InstanceNormalizationLayer.cpp
+++ b/tests/validation/reference/InstanceNormalizationLayer.cpp
@@ -52,26 +52,21 @@ SimpleTensor<T> instance_normalization(const SimpleTensor<T> &src, float gamma,
for(size_t c_i = 0; c_i < c_size; ++c_i)
{
float sum_h_w = 0;
- //Compute mean
+ float sum_sq_h_w = 0;
+
for(size_t h_i = 0; h_i < h_size; ++h_i)
{
for(size_t w_i = 0; w_i < w_size; ++w_i)
{
- sum_h_w += src[coord2index(src.shape(), Coordinates(w_i, h_i, c_i, n_i))];
+ float val = src[coord2index(src.shape(), Coordinates(w_i, h_i, c_i, n_i))];
+ sum_h_w += val;
+ sum_sq_h_w += val*val;
}
}
+ //Compute mean
const float mean_h_w = sum_h_w / (h_size * w_size);
-
//Compute variance
- float partial_var_h_w = 0;
- for(size_t h_i = 0; h_i < h_size; ++h_i)
- {
- for(size_t w_i = 0; w_i < w_size; ++w_i)
- {
- partial_var_h_w += std::pow(src[coord2index(src.shape(), Coordinates(w_i, h_i, c_i, n_i))] - mean_h_w, 2);
- }
- }
- const float var_h_w = partial_var_h_w / (h_size * w_size);
+ const float var_h_w = sum_sq_h_w / (h_size * w_size) - mean_h_w * mean_h_w;;
//Apply mean
for(size_t h_i = 0; h_i < h_size; ++h_i)