diff options
| author | Georgios Pinitas <georgios.pinitas@arm.com> | 2021-05-29 04:01:51 +0100 |
|---|---|---|
| committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2021-06-01 16:16:45 +0000 |
| commit | 9fc3be6250964a2da74cb7a05cf8e352a896ac80 (patch) | |
| tree | cdfeb80dfb556fb05851a8bb36377dcd3a808256 /src/core/CL | |
| parent | b4bb6a03f717a320b935809fde795b3d6ec5a69f (diff) | |
| download | ComputeLibrary-9fc3be6250964a2da74cb7a05cf8e352a896ac80.tar.gz | |
Fuse activation in ClDirectConv2dKernel for float types
Resolves: COMPMID-4430
Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com>
Change-Id: I9a40033e09223d601460a7e52cc297c58c9a2737
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5757
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL')
| -rw-r--r-- | src/core/CL/cl_kernels/direct_convolution.cl | 4 |
1 files changed, 4 insertions, 0 deletions
diff --git a/src/core/CL/cl_kernels/direct_convolution.cl b/src/core/CL/cl_kernels/direct_convolution.cl index e303d2067d..c5444cd7cc 100644 --- a/src/core/CL/cl_kernels/direct_convolution.cl +++ b/src/core/CL/cl_kernels/direct_convolution.cl @@ -22,6 +22,7 @@ * SOFTWARE. */ +#include "activation_float_helpers.h" #include "helpers.h" #include "helpers_asymm.h" #include "tile_helpers.h" @@ -256,6 +257,9 @@ __kernel void direct_convolution_nhwc( T_QUANTIZE8_ASYMMETRIC(ACC_DATA_TYPE, DST_DATA_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, c, cq); #endif // defined(IS_QUANTIZED) + // Apply activation + T_ACTIVATION(DST_DATA_TYPE, M0, N0, ACTIVATION_TYPE, A_VAL, B_VAL, _IOUTPUT_TILE, _IOUTPUT_TILE); + // _IOUTPUT_TILE: c = fp32/fp16, cq=qasymm8 // Store the tile in reverse order so the invalid values are overwritten with the valid ones T_STORE_INDIRECT_WIDTH_SELECT(DST_DATA_TYPE, M0, N0, PARTIAL_N0, DST_TENSOR_TYPE, dst, cout, dst_stride_y, x_cond, _IOUTPUT_TILE, dst_indirect_y); |