diff options
| author | Sang-Hoon Park <sang-hoon.park@arm.com> | 2020-07-07 09:36:09 +0100 |
|---|---|---|
| committer | Sang-Hoon Park <sang-hoon.park@arm.com> | 2020-07-08 10:00:17 +0000 |
| commit | a0205b987509d239b1635024fe8f334a4534f56e (patch) | |
| tree | b1751b74d14073d2c711de486d3f0efb1a4d2c6c /src/core/CL/cl_kernels/softmax_layer_quantized.cl | |
| parent | f9b595adbdc3f6f51ffa2c1f2aa70d0262d0db2d (diff) | |
| download | ComputeLibrary-a0205b987509d239b1635024fe8f334a4534f56e.tar.gz | |
COMPMID-3574: add logarithm to LogSoftmaxLayer
Missed logarithm for the summation is added to NEON, CL and reference
backends. To avoid complex changes, log softmax layer on CL backend
doesn't support quantized data types. Tests and doxygen comments
are modified accordingly.
Change-Id: Iafd29291be8b81345cb4999b2668dbc3ae0c3345
Signed-off-by: Sang-Hoon Park <sang-hoon.park@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3517
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Reviewed-by: SiCong Li <sicong.li@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/softmax_layer_quantized.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/softmax_layer_quantized.cl | 13 |
1 files changed, 3 insertions, 10 deletions
diff --git a/src/core/CL/cl_kernels/softmax_layer_quantized.cl b/src/core/CL/cl_kernels/softmax_layer_quantized.cl index 5d35e50b1f..81e7b896d5 100644 --- a/src/core/CL/cl_kernels/softmax_layer_quantized.cl +++ b/src/core/CL/cl_kernels/softmax_layer_quantized.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -570,14 +570,12 @@ __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 -#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)); int16 shifted_sum_minus_one = shifted_sum_minus_one_1; int16 shifted_scale = ASYMM_ONE_OVER_ONE_PLUS_X_FOR_X_IN_0_1(shifted_sum_minus_one, 16); -#endif /* LOG_SOFTMAX */ // It was already calculated in prev layer, should be stored into tmp output and reused int16 data_diff = vload16(0, (__global int *)offset(&src, 0, 0)); @@ -589,18 +587,13 @@ __kernel void softmax_layer_norm_quantized( } #endif /* defined(INPUT_BETA_MULTIPLIER) && defined(INPUT_BETA_LEFT_SHIFT) */ -#ifdef LOG_SOFTMAX - 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 */ 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); #ifdef QASYMM8_SIGNED - data = ADD_OP(data, (int16)(MIN_VALUE), int, 16); + data = ADD_OP(data, (int16)(MIN_VALUE), int, 16); #endif /* QASYMM8_SIGNED */ - data = select(MIN_VALUE, data, data_diff >= (int16)(DIFF_MIN)); -#endif /* LOG_SOFTMAX */ + data = select(MIN_VALUE, data, data_diff >= (int16)(DIFF_MIN)); vstore16(CONVERT_SAT(data, VEC_DATA_TYPE(DATA_TYPE, 16)), 0, (__global DATA_TYPE *)offset(&dst, 0, 0)); } |