From 9fc3be6250964a2da74cb7a05cf8e352a896ac80 Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Sat, 29 May 2021 04:01:51 +0100 Subject: Fuse activation in ClDirectConv2dKernel for float types Resolves: COMPMID-4430 Signed-off-by: Georgios Pinitas Change-Id: I9a40033e09223d601460a7e52cc297c58c9a2737 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5757 Reviewed-by: Michele Di Giorgio Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/direct_convolution.cl | 4 ++++ 1 file changed, 4 insertions(+) (limited to 'src/core/CL') 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); -- cgit v1.2.1