diff options
| author | Georgios Pinitas <georgios.pinitas@arm.com> | 2018-10-22 16:17:20 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:55:45 +0000 |
| commit | 60e98253f1e3df1723e7b8f4c996b544aa7c7205 (patch) | |
| tree | 45ca11d6fb0a16974fc8681bc7161a6ad2b1af2e /src/core/CL/cl_kernels/depthwise_convolution_quantized.cl | |
| parent | c04a0e8f93c620d05444251e1ae55dcf8c660a1b (diff) | |
| download | ComputeLibrary-60e98253f1e3df1723e7b8f4c996b544aa7c7205.tar.gz | |
COMPMID-1451: Fuse activation in DepthwiseConvolution.
Change-Id: Id964d9068e18aaa13ab8adcbf7a9375b034ea6c3
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/154651
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/depthwise_convolution_quantized.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/depthwise_convolution_quantized.cl | 18 |
1 files changed, 9 insertions, 9 deletions
diff --git a/src/core/CL/cl_kernels/depthwise_convolution_quantized.cl b/src/core/CL/cl_kernels/depthwise_convolution_quantized.cl index 7cd48790c6..3239885abc 100644 --- a/src/core/CL/cl_kernels/depthwise_convolution_quantized.cl +++ b/src/core/CL/cl_kernels/depthwise_convolution_quantized.cl @@ -720,7 +720,7 @@ __kernel void depthwise_convolution_3x3_quantized_nhwc( Image dst = CONVERT_TENSOR3D_TO_IMAGE_STRUCT(dst); VSTORE(VEC_SIZE) - (res, 0, dst.ptr); + (ACTIVATION_FUNC(res), 0, dst.ptr); } #endif // defined(CONV_STRIDE_X) && defined(CONV_STRIDE_Y) @@ -953,18 +953,18 @@ __kernel void depthwise_convolution_3x3_quantized_nhwc_stride1( __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x * dst_step_x + y * dst_step_y + (z * NUM_PLANES_PROCESSED) * dst_step_z; VSTORE(VEC_SIZE) - (res0, 0, dst_addr + 0 * dst_stride_y); + (ACTIVATION_FUNC(res0), 0, dst_addr + 0 * dst_stride_y); VSTORE(VEC_SIZE) - (res1, 0, dst_addr + 1 * dst_stride_y); + (ACTIVATION_FUNC(res1), 0, dst_addr + 1 * dst_stride_y); #if((DST_DIM_2 % NUM_PLANES_PROCESSED) != 0) if((z * NUM_PLANES_PROCESSED + 1) < DST_DIM_2) #endif // ((DST_DIM_2 % NUM_PLANES_PROCESSED) != 0) { VSTORE(VEC_SIZE) - (res2, 0, dst_addr + 0 * dst_stride_y + 1 * dst_stride_z); + (ACTIVATION_FUNC(res2), 0, dst_addr + 0 * dst_stride_y + 1 * dst_stride_z); VSTORE(VEC_SIZE) - (res3, 0, dst_addr + 1 * dst_stride_y + 1 * dst_stride_z); + (ACTIVATION_FUNC(res3), 0, dst_addr + 1 * dst_stride_y + 1 * dst_stride_z); } } @@ -1159,18 +1159,18 @@ __kernel void depthwise_convolution_3x3_quantized_dot8_nhwc_stride1( __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x * dst_step_x + y * dst_step_y + (z * NUM_PLANES_PROCESSED) * dst_step_z; VSTORE(VEC_SIZE) - (res0, 0, dst_addr + 0 * dst_stride_y); + (ACTIVATION_FUNC(res0), 0, dst_addr + 0 * dst_stride_y); VSTORE(VEC_SIZE) - (res1, 0, dst_addr + 1 * dst_stride_y); + (ACTIVATION_FUNC(res1), 0, dst_addr + 1 * dst_stride_y); #if((DST_DIM_2 % NUM_PLANES_PROCESSED) != 0) if((z * NUM_PLANES_PROCESSED + 1) < DST_DIM_2) #endif // ((DST_DIM_2 % NUM_PLANES_PROCESSED) != 0) { VSTORE(VEC_SIZE) - (res2, 0, dst_addr + 0 * dst_stride_y + 1 * dst_stride_z); + (ACTIVATION_FUNC(res2), 0, dst_addr + 0 * dst_stride_y + 1 * dst_stride_z); VSTORE(VEC_SIZE) - (res3, 0, dst_addr + 1 * dst_stride_y + 1 * dst_stride_z); + (ACTIVATION_FUNC(res3), 0, dst_addr + 1 * dst_stride_y + 1 * dst_stride_z); } } #endif // defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) |