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.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.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/softmax_layer.cl | 3 |
1 files changed, 2 insertions, 1 deletions
diff --git a/src/core/CL/cl_kernels/softmax_layer.cl b/src/core/CL/cl_kernels/softmax_layer.cl index 767cf4c4f7..0c9f8c1c66 100644 --- a/src/core/CL/cl_kernels/softmax_layer.cl +++ b/src/core/CL/cl_kernels/softmax_layer.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -112,6 +112,7 @@ __kernel void softmax_layer_norm( VEC_DATA_TYPE(DATA_TYPE, 16) data = vload16(0, (__global DATA_TYPE *)offset(&src, 0, 0)); #ifdef LOG_SOFTMAX + sum_val = log(sum_val); vstore16(SUB_OP(data, sum_val, DATA_TYPE, 16), 0, (__global DATA_TYPE *)offset(&dst, 0, 0)); #else /* LOG_SOFTMAX */ vstore16(DIV_OP(data, sum_val, DATA_TYPE, 16), 0, (__global DATA_TYPE *)offset(&dst, 0, 0)); |