Gitiles
Code Review Sign In
review.mlplatform.org / ml / ComputeLibrary / 7d7c4206eb5fce3a1e93bad7a91841ff6d904f23 / . / src / core / CL / cl_kernels / direct_convolution_quantized.cl
blob: 0a8c5faecf6d10c2ead3c39613eb3368871d87ba [file] [log] [blame]
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]1/*
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]2 * Copyright (c) 2017-2020 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
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 + weight_offset); \
50 acc += ((int8)(src0.s1234, src0.s5678) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \
51 acc += ((int8)(src0.s2345, src0.s6789) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \
52 acc += ((int8)(src0.s3456, src0.s789A) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \
53 acc += ((int8)(src0.s4567, src0.s89AB) + input_offset) * ((int8)weights_values0.s4 + weight_offset); \
54 acc += ((int8)(src0.s5678, src0.s9ABC) + input_offset) * ((int8)weights_values0.s5 + weight_offset); \
55 acc += ((int8)(src0.s6789, src0.sABCD) + input_offset) * ((int8)weights_values0.s6 + weight_offset); \
56 acc += ((int8)(src0.s789A, src0.sBCDE) + input_offset) * ((int8)weights_values0.s7 + weight_offset); \
57 acc += ((int8)(src0.s89AB, src0.sCDEF) + input_offset) * ((int8)weights_value1 + weight_offset); \
58 })
59
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 + weight_offset); \
67 acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \
68 acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \
69 acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \
70 acc += ((int8)(src0.s468A, src0.sCE, src1.s02) + input_offset) * ((int8)weights_values0.s4 + weight_offset); \
71 acc += ((int8)(src0.s579B, src0.sDF, src1.s13) + input_offset) * ((int8)weights_values0.s5 + weight_offset); \
72 acc += ((int8)(src0.s68AC, src0.sE, src1.s024) + input_offset) * ((int8)weights_values0.s6 + weight_offset); \
73 acc += ((int8)(src0.s79BD, src0.sF, src1.s135) + input_offset) * ((int8)weights_values0.s7 + weight_offset); \
74 acc += ((int8)(src0.s8ACE, src1.s0246) + input_offset) * ((int8)weights_value1 + weight_offset); \
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
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 + weight_offset); \
94 acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \
95 acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \
96 acc += ((int8)(src0.s345, src0.s67, src1.s012) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \
97 acc += ((int8)(src0.s45, src0.s67, src1.s0123) + input_offset) * ((int8)weights_value1 + weight_offset); \
98 })
99
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 + weight_offset); \
107 acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \
108 acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \
109 acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \
110 acc += ((int8)(src0.s468a, src0.sCE, src1.s02) + input_offset) * ((int8)weights_value1 + weight_offset); \
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
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 + weight_offset); \
129 acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \
130 acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \
131 })
132
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 + weight_offset); \
139 acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \
140 acc += ((int8)(src0.s2468, src0.sACE, src1) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \
141 })
142
143#elif KERNEL_SIZE == 1
144
145#if STRIDE_X == 3
146#define INPUT_PIXEL extract_input_stride3
147#elif STRIDE_X == 2
148#define INPUT_PIXEL extract_input_stride2
149#elif STRIDE_X == 1
150#define INPUT_PIXEL extract_input_stride1
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 *
158 * @param[in] input_pixel Pointer to the first pixel.
159 *
160 * @return extracted input pixels.
161 */
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]162inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride1(__global const DATA_TYPE *input_pixel)
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]163{
164 return vload8(0, input_pixel);
165}
166
167/** Extracts a 1D horizontal vector from the input tensor with stride as 2.
168 *
169 * @param[in] input_pixel Pointer to the first pixel.
170 *
171 * @return extracted input pixels.
172 */
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]173inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride2(__global const DATA_TYPE *input_pixel)
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)
176 temp = vload16(0, input_pixel);
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 *
182 * @param[in] input_pixel Pointer to the first pixel.
183 *
184 * @return extracted input pixels.
185 */
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]186inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3(__global const DATA_TYPE *input_pixel)
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)
189 temp1 = vload16(0, input_pixel);
190 VEC_DATA_TYPE(DATA_TYPE, 16)
191 temp2 = vload16(0, input_pixel + 12);
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
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]206 *
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]207 * @param[in] src_ptr Pointer to the source tensor. Supported data types: QASYMM8/QASYMM8_SIGNED
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]208 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes)
209 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
210 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes)
211 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
212 * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes)
213 * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes)
214 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor
215 * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr
216 * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes)
217 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
218 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes)
219 * @param[in] dst_step_y dst_stride_y * number of elements along Z processed per workitem(in bytes)
220 * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes)
221 * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes)
222 * @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]223 * @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]224 * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes)
225 * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes)
226 * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes)
227 * @param[in] weights_step_y weights_stride_y * number of elements along y processed per workitem(in bytes)
228 * @param[in] weights_stride_z Stride of the weights tensor in Z dimension (in bytes)
229 * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes)
230 * @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]231 * @param[in] biases_ptr Pointer to the biases tensor. Supported data types: S32
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]232 * @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes)
233 * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes)
234 * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor
235 * @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension
236 * @param[in] input_offset Input offset quantization parameter
237 * @param[in] weight_offset Weights offset quantization parameter
238 * @param[in] output_offset Output offset quantization parameter
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) */
247 unsigned int weights_stride_w,
248 int input_offset,
249 int weight_offset,
Michele Di Giorgio14cbfb22019-10-23 10:53:10 +0100[diff] [blame]250 int output_offset)
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]251{
252 Image src = CONVERT_TO_IMAGE_STRUCT(src);
253 Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights);
254 Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);
255
256 int8 pixels0 = 0;
257
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]258 __global DATA_TYPE *weights_addr = (__global DATA_TYPE *)tensor3D_offset(&weights, 0, 0, 0);
259 __global DATA_TYPE *src_addr = (__global DATA_TYPE *)offset(&src, 0, 0);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]260
261 const int kernel_index = get_global_id(2);
262 weights_addr += kernel_index * weights_stride_w;
263
264 for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d)
265 {
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100[diff] [blame]266#if KERNEL_SIZE == 9
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]267 CONVOLUTION1x9(pixels0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_y));
268 CONVOLUTION1x9(pixels0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y));
269 CONVOLUTION1x9(pixels0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_y));
270 CONVOLUTION1x9(pixels0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_y));
271 CONVOLUTION1x9(pixels0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_y));
272 CONVOLUTION1x9(pixels0, (__global DATA_TYPE *)(src_addr + 5 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 5 * weights_stride_y));
273 CONVOLUTION1x9(pixels0, (__global DATA_TYPE *)(src_addr + 6 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 6 * weights_stride_y));
274 CONVOLUTION1x9(pixels0, (__global DATA_TYPE *)(src_addr + 7 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 7 * weights_stride_y));
275 CONVOLUTION1x9(pixels0, (__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]276#elif KERNEL_SIZE == 5
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]277 CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)src_addr, (__global DATA_TYPE *)weights_addr);
278 CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y));
279 CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_y));
280 CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_y));
281 CONVOLUTION1x5(pixels0, (__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]282#elif KERNEL_SIZE == 3
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]283 CONVOLUTION1x3(pixels0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_y));
284 CONVOLUTION1x3(pixels0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y));
285 CONVOLUTION1x3(pixels0, (__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]286#elif KERNEL_SIZE == 1
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]287 int weight = convert_int(*(__global DATA_TYPE *)weights_addr);
288 int8 input_pixel = convert_int8(INPUT_PIXEL((__global DATA_TYPE *)src_addr));
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]289 pixels0 += (input_pixel + input_offset) * ((int8)weight + weight_offset);
290#endif /* (KERNEL_SIZE == 1) || (KERNEL_SIZE == 3) || (KERNEL_SIZE == 5) */
291
292 src_addr += src_stride_z;
293 weights_addr += weights_stride_z;
294 }
295
296#ifdef HAS_BIAS
Georgios Pinitas540d0082017-11-17 10:55:00 +0000[diff] [blame]297 Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases);
298 __global int *bias_addr = ((__global int *)(vector_offset(&biases, kernel_index)));
299 pixels0 += (int8)(*bias_addr);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]300#endif /* defined(HAS_BIAS) */
301
Michele Di Giorgio14cbfb22019-10-23 10:53:10 +0100[diff] [blame]302#if OUTPUT_SHIFT < 0
303 pixels0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_GREATER_THAN_ONE(pixels0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8);
304#else // OUTPUT_SHIFT < 0
305 pixels0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(pixels0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8);
306#endif // OUTPUT_SHIFT < 0
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]307 pixels0 = pixels0 + output_offset;
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]308
Sheri Zhang681f2d42020-02-20 11:23:08 +0000[diff] [blame]309 vstore8(CONVERT_SAT(pixels0, DATA_TYPE), 0, (__global DATA_TYPE *)dst.ptr);
Chunosovd621bca2017-11-03 17:33:15 +0700[diff] [blame]310}
Michele Di Giorgio14cbfb22019-10-23 10:53:10 +0100[diff] [blame]311#endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) && defined(OUTPUT_MULTIPLIER) && defined(OUTPUT_SHIFT)