From af724e762b622b677d79e4ac63b22b4c5581f05d Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Tue, 12 Nov 2019 14:45:14 +0000 Subject: COMPMID-2897 CL LogSoftmax QASYMM8 pre-commit tests are intermittently failing Change-Id: I5cac5d40cf345cf91cfbea1e60ef78ede048cc1e Signed-off-by: Giorgio Arena Reviewed-on: https://review.mlplatform.org/c/2267 Reviewed-by: Michele Di Giorgio Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/softmax_layer_quantized.cl | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/src/core/CL/cl_kernels/softmax_layer_quantized.cl b/src/core/CL/cl_kernels/softmax_layer_quantized.cl index 8ccc5d3dd5..ce3bd7bc43 100644 --- a/src/core/CL/cl_kernels/softmax_layer_quantized.cl +++ b/src/core/CL/cl_kernels/softmax_layer_quantized.cl @@ -560,8 +560,8 @@ __kernel void softmax_layer_norm_quantized( int sum_val = *((__global int *)offset(&sum, 0, get_global_id(1))); // It will be better to calculate this in prev layer and pass here as parameter - uint sum_val_u = convert_uint(sum_val); #ifndef LOG_SOFTMAX + uint sum_val_u = convert_uint(sum_val); int headroom_plus_one = clz(sum_val_u); int num_bits_over_unit = EXP_ACCUMULATION_INT_BITS - headroom_plus_one; int shifted_sum_minus_one_1 = convert_int((sum_val_u << headroom_plus_one) - (1u << 31)); @@ -578,15 +578,16 @@ __kernel void softmax_layer_norm_quantized( data_diff_mult = ASYMM_MULT(data_diff * (1 << INPUT_BETA_LEFT_SHIFT), INPUT_BETA_MULTIPLIER, 16); } #endif /* defined(INPUT_BETA_MULTIPLIER) && defined(INPUT_BETA_LEFT_SHIFT) */ - int16 data = ASYMM_EXP_ON_NEGATIVE_VALUES(data_diff_mult, SCALED_DIFF_INT_BITS, 16); #ifdef LOG_SOFTMAX - data = SUB_OP(data_diff_mult, (int16)sum_val_u, int, 16); + long16 data = SUB_OP(convert_long16(data_diff_mult), (long16)(sum_val), long, 16); + data = select(0L, data, convert_long16(data_diff) >= (long16)(DIFF_MIN)); #else /* LOG_SOFTMAX */ - data = ASYMM_MULT(shifted_scale, data, 16); - data = ASYMM_ROUNDING_DIVIDE_BY_POW2(data, num_bits_over_unit + 31 - 8, 16); + int16 data = ASYMM_EXP_ON_NEGATIVE_VALUES(data_diff_mult, SCALED_DIFF_INT_BITS, 16); + data = ASYMM_MULT(shifted_scale, data, 16); + data = ASYMM_ROUNDING_DIVIDE_BY_POW2(data, num_bits_over_unit + 31 - 8, 16); + data = select(0, data, data_diff >= (int16)(DIFF_MIN)); #endif /* LOG_SOFTMAX */ - data = select(0, data, data_diff >= (int16)(DIFF_MIN)); vstore16(convert_uchar16_sat(data), 0, (__global uchar *)offset(&dst, 0, 0)); } -- cgit v1.2.1