diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2019-11-12 14:45:14 +0000 |
|---|---|---|
| committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2019-11-12 21:33:29 +0000 |
| commit | af724e762b622b677d79e4ac63b22b4c5581f05d (patch) | |
| tree | e87d37e07fa3d0fbb9efa6b4a04e39ad5963c40b | |
| parent | 7b12bfb369d11eeb3e51ce178955dd0186fa8b52 (diff) | |
| download | ComputeLibrary-af724e762b622b677d79e4ac63b22b4c5581f05d.tar.gz | |
COMPMID-2897 CL LogSoftmax QASYMM8 pre-commit tests are intermittently failing
Change-Id: I5cac5d40cf345cf91cfbea1e60ef78ede048cc1e
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-on: https://review.mlplatform.org/c/2267
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/softmax_layer_quantized.cl | 13 |
1 files 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)); } |