diff options
| -rw-r--r-- | src/core/CL/cl_kernels/instance_normalization.cl | 14 | ||||
| -rw-r--r-- | tests/validation/reference/InstanceNormalizationLayer.cpp | 19 |
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) |