Gitiles
Code Review Sign In
review.mlplatform.org / ml / ComputeLibrary / e5d76e1574103de405df625e48e5294ea106060c / . / src / core / CL / cl_kernels / direct_convolution_quantized.cl
blob: b80d4f587e3e4a716aceb77d1e7586264e71ccf5 [file] [log] [blame]
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]1/*
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]2 * Copyright (c) 2017-2021 Arm Limited.
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]3 *
4 * SPDX-License-Identifier: MIT
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a copy
7 * of this software and associated documentation files (the "Software"), to
8 * deal in the Software without restriction, including without limitation the
9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10 * sell copies of the Software, and to permit persons to whom the Software is
11 * furnished to do so, subject to the following conditions:
12 *
13 * The above copyright notice and this permission notice shall be included in all
14 * copies or substantial portions of the Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22 * SOFTWARE.
23 */
24#include "helpers_asymm.h"
25
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]26#undef CONVERT_SAT_STR
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]27#undef CONVERT_SAT
28
Michele Di Giorgio14cbfb22019-10-23 10:53:10 +0100[diff] [blame]29#if defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) && defined(OUTPUT_MULTIPLIER) && defined(OUTPUT_SHIFT)
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]30
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]31#define CONVERT_SAT_STR(x, type) (convert_##type##8_sat((x)))
32#define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type)
33
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100[diff] [blame]34#if KERNEL_SIZE == 9
35
36#if STRIDE_X == 1
37#define CONVOLUTION1x9(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x9_STRIDE1(acc, src_row_ptr, weights_row_ptr)
38#elif STRIDE_X == 2
39#define CONVOLUTION1x9(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x9_STRIDE2(acc, src_row_ptr, weights_row_ptr)
40#else /* STRIDE_X not equals 1 or 2 */
41#error "STRIDE_X larger than 2 is not supported"
42#endif /* STRIDE_X */
43
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]44#define CONVOLUTION1x9_STRIDE1(acc, src_row_ptr, weights_row_ptr) \
45 ({ \
46 int8 weights_values0 = convert_int8(vload8(0, weights_row_ptr)); \
47 int weights_value1 = convert_int(*(weights_row_ptr + 8)); \
48 int16 src0 = convert_int16(vload16(0, src_row_ptr)); \
49 acc += (src0.lo + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \
50 acc += ((int8)(src0.s1234, src0.s5678) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \
51 acc += ((int8)(src0.s2345, src0.s6789) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \
52 acc += ((int8)(src0.s3456, src0.s789A) + INPUT_OFFSET) * ((int8)weights_values0.s3 + WEIGHTS_OFFSET); \
53 acc += ((int8)(src0.s4567, src0.s89AB) + INPUT_OFFSET) * ((int8)weights_values0.s4 + WEIGHTS_OFFSET); \
54 acc += ((int8)(src0.s5678, src0.s9ABC) + INPUT_OFFSET) * ((int8)weights_values0.s5 + WEIGHTS_OFFSET); \
55 acc += ((int8)(src0.s6789, src0.sABCD) + INPUT_OFFSET) * ((int8)weights_values0.s6 + WEIGHTS_OFFSET); \
56 acc += ((int8)(src0.s789A, src0.sBCDE) + INPUT_OFFSET) * ((int8)weights_values0.s7 + WEIGHTS_OFFSET); \
57 acc += ((int8)(src0.s89AB, src0.sCDEF) + INPUT_OFFSET) * ((int8)weights_value1 + WEIGHTS_OFFSET); \
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100[diff] [blame]58 })
59
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]60#define CONVOLUTION1x9_STRIDE2(acc, src_row_ptr, weights_row_ptr) \
61 ({ \
62 int8 weights_values0 = convert_int8(vload8(0, weights_row_ptr)); \
63 int weights_value1 = convert_int(*(weights_row_ptr + 8)); \
64 int16 src0 = convert_int16(vload16(0, src_row_ptr)); \
65 int8 src1 = convert_int8(vload8(0, src_row_ptr + 16)); \
66 acc += (src0.even + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \
67 acc += ((int8)(src0.s1357, src0.s9BDF) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \
68 acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \
69 acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + INPUT_OFFSET) * ((int8)weights_values0.s3 + WEIGHTS_OFFSET); \
70 acc += ((int8)(src0.s468A, src0.sCE, src1.s02) + INPUT_OFFSET) * ((int8)weights_values0.s4 + WEIGHTS_OFFSET); \
71 acc += ((int8)(src0.s579B, src0.sDF, src1.s13) + INPUT_OFFSET) * ((int8)weights_values0.s5 + WEIGHTS_OFFSET); \
72 acc += ((int8)(src0.s68AC, src0.sE, src1.s024) + INPUT_OFFSET) * ((int8)weights_values0.s6 + WEIGHTS_OFFSET); \
73 acc += ((int8)(src0.s79BD, src0.sF, src1.s135) + INPUT_OFFSET) * ((int8)weights_values0.s7 + WEIGHTS_OFFSET); \
74 acc += ((int8)(src0.s8ACE, src1.s0246) + INPUT_OFFSET) * ((int8)weights_value1 + WEIGHTS_OFFSET); \
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100[diff] [blame]75 })
76
77#elif KERNEL_SIZE == 5
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]78
79#if STRIDE_X == 1
80#define CONVOLUTION1x5(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr)
81#elif STRIDE_X == 2
82#define CONVOLUTION1x5(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr)
83#else /* STRIDE_X not equals 1 or 2 */
84#error "STRIDE_X larger than 2 is not supported"
85#endif /* STRIDE_X */
86
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]87#define CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr) \
88 ({ \
89 int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \
90 int weights_value1 = convert_int(*(weights_row_ptr + 4)); \
91 int8 src0 = convert_int8(vload8(0, src_row_ptr)); \
92 int4 src1 = convert_int4(vload4(0, src_row_ptr + 8)); \
93 acc += (src0 + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \
94 acc += ((int8)(src0.s1234, src0.s567, src1.s0) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \
95 acc += ((int8)(src0.s234, src0.s567, src1.s01) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \
96 acc += ((int8)(src0.s345, src0.s67, src1.s012) + INPUT_OFFSET) * ((int8)weights_values0.s3 + WEIGHTS_OFFSET); \
97 acc += ((int8)(src0.s45, src0.s67, src1.s0123) + INPUT_OFFSET) * ((int8)weights_value1 + WEIGHTS_OFFSET); \
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]98 })
99
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]100#define CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr) \
101 ({ \
102 int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \
103 int weights_value1 = convert_int(*(weights_row_ptr + 4)); \
104 int16 src0 = convert_int16(vload16(0, src_row_ptr)); \
105 int4 src1 = convert_int4(vload4(0, src_row_ptr + 16)); \
106 acc += (src0.even + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \
107 acc += ((int8)(src0.s1357, src0.s9BDF) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \
108 acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \
109 acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + INPUT_OFFSET) * ((int8)weights_values0.s3 + WEIGHTS_OFFSET); \
110 acc += ((int8)(src0.s468a, src0.sCE, src1.s02) + INPUT_OFFSET) * ((int8)weights_value1 + WEIGHTS_OFFSET); \
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]111 })
112
113#elif KERNEL_SIZE == 3
114
115#if STRIDE_X == 1
116#define CONVOLUTION1x3(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr)
117#elif STRIDE_X == 2
118#define CONVOLUTION1x3(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr)
119#else /* STRIDE_X not equals 1 or 2 */
120#error "STRIDE_X larger than 2 is not supported"
121#endif /* STRIDE_X */
122
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]123#define CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr) \
124 ({ \
125 int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \
126 int8 src0 = convert_int8(vload8(0, src_row_ptr)); \
127 int2 src1 = convert_int2(vload2(0, src_row_ptr + 8)); \
128 acc += (src0 + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \
129 acc += ((int8)(src0.s1234, src0.s567, src1.s0) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \
130 acc += ((int8)(src0.s234, src0.s567, src1.s01) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]131 })
132
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]133#define CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr) \
134 ({ \
135 int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \
136 int16 src0 = convert_int16(vload16(0, src_row_ptr)); \
137 int src1 = convert_int(*(src_row_ptr + 16)); \
138 acc += (src0.even + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \
139 acc += ((int8)(src0.s1357, src0.s9BDF) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \
140 acc += ((int8)(src0.s2468, src0.sACE, src1) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]141 })
142
143#elif KERNEL_SIZE == 1
144
145#if STRIDE_X == 3
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]146#define INPUT_VALUE extract_input_stride3
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]147#elif STRIDE_X == 2
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]148#define INPUT_VALUE extract_input_stride2
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]149#elif STRIDE_X == 1
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]150#define INPUT_VALUE extract_input_stride1
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]151
152#else /* STRIDE_X not equals 1, 2 or 3 */
153#error "Only support strides 1, 2 and 3"
154#endif /* STRIDE_X */
155
156/** Extracts a 1D horizontal vector from the input tensor with stride as 1.
157 *
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]158 * @param[in] input_value Pointer to the first value.
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]159 *
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]160 * @return extracted input values.
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]161 */
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]162inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride1(__global const DATA_TYPE *input_value)
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]163{
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]164 return vload8(0, input_value);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]165}
166
167/** Extracts a 1D horizontal vector from the input tensor with stride as 2.
168 *
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]169 * @param[in] input_value Pointer to the first value.
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]170 *
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]171 * @return extracted input values.
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]172 */
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]173inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride2(__global const DATA_TYPE *input_value)
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]174{
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]175 VEC_DATA_TYPE(DATA_TYPE, 16)
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]176 temp = vload16(0, input_value);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]177 return temp.s02468ace;
178}
179
180/** Extracts a 1D horizontal vector from the input tensor with stride as 3 and 8-bit data size.
181 *
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]182 * @param[in] input_value Pointer to the first value.
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]183 *
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]184 * @return extracted input values.
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]185 */
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]186inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3(__global const DATA_TYPE *input_value)
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]187{
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]188 VEC_DATA_TYPE(DATA_TYPE, 16)
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]189 temp1 = vload16(0, input_value);
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]190 VEC_DATA_TYPE(DATA_TYPE, 16)
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]191 temp2 = vload16(0, input_value + 12);
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]192 return (VEC_DATA_TYPE(DATA_TYPE, 8))(temp1.s0369, temp2.s0369);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]193}
194
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100[diff] [blame]195#else /* KERNEL_SIZE not equals 1, 3 , 5, 9 */
196#error "Only kernel sizes 1, 3, 5 and 9 are supported"
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]197#endif /* KERNEL_SIZE */
198
199/** This kernel performs a direct convolution to convolve the low three dimensions.
200 *
201 * @note The convolution stride x must be passed at compile time using -DSTRIDE_X e.g. -DSTRIDE_X=1
202 * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH
203 * @note If biases are used then -DHAS_BIAS has to be passed at compile time
Michele Di Giorgio14cbfb22019-10-23 10:53:10 +0100[diff] [blame]204 * @note The output quantization multiplier must be passed at compile time using -DOUTPUT_MULTIPLIER e.g. -DOUTPUT_MULTIPLIER=1234
205 * @note The output quantization shift must be passed at compile time using -DOUTPUT_SHIFT e.g. -DOUTPUT_SHIFT=4
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]206 * @note The input offset quantization parameter must be passed at compile time using -DINPUT_OFFSET e.g. -DINPUT_OFFSET=3
207 * @note The weights offset quantization parameter must be passed at compile time using -DWEIGHTS_OFFSET e.g. -DWEIGHTS_OFFSET=3
208 * @note The destination offset quantization parameter must be passed at compile time using -DOUTPUT_OFFSET e.g. -DOUTPUT_OFFSET=3
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]209 *
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]210 * @param[in] src_ptr Pointer to the source tensor. Supported data types: QASYMM8/QASYMM8_SIGNED
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]211 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes)
212 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
213 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes)
214 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
215 * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes)
216 * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes)
217 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor
218 * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr
219 * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes)
220 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
221 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes)
222 * @param[in] dst_step_y dst_stride_y * number of elements along Z processed per workitem(in bytes)
223 * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes)
224 * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes)
225 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]226 * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]227 * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes)
228 * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes)
229 * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes)
230 * @param[in] weights_step_y weights_stride_y * number of elements along y processed per workitem(in bytes)
231 * @param[in] weights_stride_z Stride of the weights tensor in Z dimension (in bytes)
232 * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes)
233 * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor
Georgios Pinitas540d0082017-11-17 10:55:00 +0000[diff] [blame]234 * @param[in] biases_ptr Pointer to the biases tensor. Supported data types: S32
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]235 * @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes)
236 * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes)
237 * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor
238 * @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]239 */
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100[diff] [blame]240__kernel void direct_convolution_quantized(
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]241 TENSOR3D_DECLARATION(src),
242 TENSOR3D_DECLARATION(dst),
243 TENSOR3D_DECLARATION(weights),
244#ifdef HAS_BIAS
245 VECTOR_DECLARATION(biases),
246#endif /* defined(HAS_BIAS) */
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]247 unsigned int weights_stride_w)
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]248{
249 Image src = CONVERT_TO_IMAGE_STRUCT(src);
250 Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights);
251 Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);
252
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]253 int8 values0 = 0;
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]254
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]255 __global DATA_TYPE *weights_addr = (__global DATA_TYPE *)tensor3D_offset(&weights, 0, 0, 0);
256 __global DATA_TYPE *src_addr = (__global DATA_TYPE *)offset(&src, 0, 0);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]257
258 const int kernel_index = get_global_id(2);
259 weights_addr += kernel_index * weights_stride_w;
260
261 for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d)
262 {
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100[diff] [blame]263#if KERNEL_SIZE == 9
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]264 CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_y));
265 CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y));
266 CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_y));
267 CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_y));
268 CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_y));
269 CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 5 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 5 * weights_stride_y));
270 CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 6 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 6 * weights_stride_y));
271 CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 7 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 7 * weights_stride_y));
272 CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 8 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 8 * weights_stride_y));
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100[diff] [blame]273#elif KERNEL_SIZE == 5
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]274 CONVOLUTION1x5(values0, (__global DATA_TYPE *)src_addr, (__global DATA_TYPE *)weights_addr);
275 CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y));
276 CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_y));
277 CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_y));
278 CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_y));
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]279#elif KERNEL_SIZE == 3
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]280 CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_y));
281 CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y));
282 CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_y));
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]283#elif KERNEL_SIZE == 1
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]284 int weight = convert_int(*(__global DATA_TYPE *)weights_addr);
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]285 int8 input_value = convert_int8(INPUT_VALUE((__global DATA_TYPE *)src_addr));
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]286 values0 += (input_value + INPUT_OFFSET) * ((int8)weight + WEIGHTS_OFFSET);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]287#endif /* (KERNEL_SIZE == 1) || (KERNEL_SIZE == 3) || (KERNEL_SIZE == 5) */
288
289 src_addr += src_stride_z;
290 weights_addr += weights_stride_z;
291 }
292
293#ifdef HAS_BIAS
Georgios Pinitas540d0082017-11-17 10:55:00 +0000[diff] [blame]294 Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases);
295 __global int *bias_addr = ((__global int *)(vector_offset(&biases, kernel_index)));
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]296 values0 += (int8)(*bias_addr);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]297#endif /* defined(HAS_BIAS) */
298
Michele Di Giorgio14cbfb22019-10-23 10:53:10 +0100[diff] [blame]299#if OUTPUT_SHIFT < 0
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]300 values0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_GREATER_THAN_ONE(values0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8);
Michele Di Giorgio14cbfb22019-10-23 10:53:10 +0100[diff] [blame]301#else // OUTPUT_SHIFT < 0
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]302 values0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(values0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8);
Michele Di Giorgio14cbfb22019-10-23 10:53:10 +0100[diff] [blame]303#endif // OUTPUT_SHIFT < 0
Gian Marco Iodiceff1fe3e2021-01-02 09:58:51 +0000[diff] [blame]304 values0 = values0 + OUTPUT_OFFSET;
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]305
Giorgio Arenae620a832020-02-17 16:33:20 +0000[diff] [blame]306 vstore8(CONVERT_SAT(values0, DATA_TYPE), 0, (__global DATA_TYPE *)dst.ptr);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]307}
Michele Di Giorgio14cbfb22019-10-23 10:53:10 +0100[diff] [blame]308#endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) && defined(OUTPUT_MULTIPLIER) && defined(OUTPUT_SHIFT)