2 * Copyright (c) 2016, 2017 ARM Limited.
4 * SPDX-License-Identifier: MIT
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:
13 * The above copyright notice and this permission notice shall be included in all
14 * copies or substantial portions of the Software.
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
27 #define DATA_TYPE short
28 #endif /* DATA_TYPE */
31 #define COMPUTE_TYPE int
32 #endif /* COMPUTE_TYPE */
35 #define DATA_TYPE_OUT uchar
36 #endif /* DATA_TYPE_OUT */
38 /** Compute a 1D horizontal convolution of size 7 for 8 bytes assuming the input is made of 1 channel of 1 byte (i.e 8 pixels).
40 * @param[in] left_pixel Pointer to the left pixel
41 * @param[in] left1_coeff Weight of the most left pixel
42 * @param[in] left2_coeff Weight of the second left pixel
43 * @param[in] left3_coeff Weight of the left pixel
44 * @param[in] middle_coeff Weight of the middle pixel
45 * @param[in] right1_coeff Weight of the right pixel
46 * @param[in] right2_coeff Weight of the second right pixel
47 * @param[in] right3_coeff Weight of the most right pixel
49 * @return a short8 containing 8 convoluted values.
51 VEC_DATA_TYPE(DATA_TYPE, 8)
53 __global const uchar *left_pixel,
54 const short left1_coeff,
55 const short left2_coeff,
56 const short left3_coeff,
57 const short middle_coeff,
58 const short right1_coeff,
59 const short right2_coeff,
60 const short right3_coeff)
62 uchar16 temp = vload16(0, left_pixel);
64 VEC_DATA_TYPE(DATA_TYPE, 8)
65 left1 = CONVERT(temp.s01234567, VEC_DATA_TYPE(DATA_TYPE, 8));
66 VEC_DATA_TYPE(DATA_TYPE, 8)
67 left2 = CONVERT(temp.s12345678, VEC_DATA_TYPE(DATA_TYPE, 8));
68 VEC_DATA_TYPE(DATA_TYPE, 8)
69 left3 = CONVERT(temp.s23456789, VEC_DATA_TYPE(DATA_TYPE, 8));
70 VEC_DATA_TYPE(DATA_TYPE, 8)
71 middle = CONVERT(temp.s3456789a, VEC_DATA_TYPE(DATA_TYPE, 8));
72 VEC_DATA_TYPE(DATA_TYPE, 8)
73 right1 = CONVERT(temp.s456789ab, VEC_DATA_TYPE(DATA_TYPE, 8));
74 VEC_DATA_TYPE(DATA_TYPE, 8)
75 right2 = CONVERT(temp.s56789abc, VEC_DATA_TYPE(DATA_TYPE, 8));
76 VEC_DATA_TYPE(DATA_TYPE, 8)
77 right3 = CONVERT(temp.s6789abcd, VEC_DATA_TYPE(DATA_TYPE, 8));
79 return left1 * (VEC_DATA_TYPE(DATA_TYPE, 8))left1_coeff + left2 * (VEC_DATA_TYPE(DATA_TYPE, 8))left2_coeff + left3 * (VEC_DATA_TYPE(DATA_TYPE, 8))left3_coeff + middle * (VEC_DATA_TYPE(DATA_TYPE,
80 8))middle_coeff + right1 * (VEC_DATA_TYPE(DATA_TYPE, 8))right1_coeff + right2 * (VEC_DATA_TYPE(DATA_TYPE, 8))right2_coeff + right3 * (VEC_DATA_TYPE(DATA_TYPE, 8))right3_coeff;
83 /** Compute a 1D vertical convolution of size 7 for 8 bytes assuming the input is made of 1 channel of 1 byte (i.e 8 pixels).
85 * @param[in] src Pointer to source image.
86 * @param[in] up1_coeff Weight of the most up pixel
87 * @param[in] up2_coeff Weight of the second up pixel
88 * @param[in] up3_coeff Weight of the up pixel
89 * @param[in] middle_coeff Weight of the middle pixel
90 * @param[in] down1_coeff Weight of the down pixel
91 * @param[in] down2_coeff Weight of the second down pixel
92 * @param[in] down3_coeff Weight of the third down pixel
94 * @return a short8 containing 8 convoluted values.
96 VEC_DATA_TYPE(COMPUTE_TYPE, 8)
99 const short up1_coeff,
100 const short up2_coeff,
101 const short up3_coeff,
102 const short middle_coeff,
103 const short down1_coeff,
104 const short down2_coeff,
105 const short down3_coeff)
107 VEC_DATA_TYPE(COMPUTE_TYPE, 8)
109 VEC_DATA_TYPE(COMPUTE_TYPE, 8)
110 out = (VEC_DATA_TYPE(COMPUTE_TYPE, 8))0;
112 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, -3)), VEC_DATA_TYPE(COMPUTE_TYPE, 8));
113 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))up1_coeff;
115 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, -2)), VEC_DATA_TYPE(COMPUTE_TYPE, 8));
116 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))up2_coeff;
118 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, -1)), VEC_DATA_TYPE(COMPUTE_TYPE, 8));
119 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))up3_coeff;
121 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 0)), VEC_DATA_TYPE(COMPUTE_TYPE, 8));
122 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))middle_coeff;
124 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 1)), VEC_DATA_TYPE(COMPUTE_TYPE, 8));
125 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))down1_coeff;
127 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 2)), VEC_DATA_TYPE(COMPUTE_TYPE, 8));
128 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))down2_coeff;
130 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 3)), VEC_DATA_TYPE(COMPUTE_TYPE, 8));
131 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))down3_coeff;
136 /** Apply a 7x7 convolution matrix to a single channel U8 input image and return the result.
138 * Convolution matrix layout:\n
139 * [ mat0, mat1, mat2, mat3 , mat4, mat5, mat6 ]\n
140 * [ mat7, mat8, mat9, mat10, mat11, mat12, mat13 ]\n
141 * [ mat14, mat15, mat16, mat17, mat18, mat19, mat20 ]\n
142 * [ mat21, mat22, mat23, mat24, mat25, mat26, mat27 ]\n
143 * [ mat28, mat29, mat30, mat31, mat32, mat33, mat34 ]\n
144 * [ mat35, mat36, mat37, mat38, mat39, mat40, mat41 ]\n
145 * [ mat42, mat43, mat44, mat45, mat46, mat47, mat48 ]
147 * @param[in] src A pointer to source Image structure.
148 * @param[in] mat0 Coefficient from the convolution matrix
149 * @param[in] mat1 Coefficient from the convolution matrix
150 * @param[in] mat2 Coefficient from the convolution matrix
151 * @param[in] mat3 Coefficient from the convolution matrix
152 * @param[in] mat4 Coefficient from the convolution matrix
153 * @param[in] mat5 Coefficient from the convolution matrix
154 * @param[in] mat6 Coefficient from the convolution matrix
155 * @param[in] mat0 Coefficient from the convolution matrix
156 * @param[in] mat7 Coefficient from the convolution matrix
157 * @param[in] mat8 Coefficient from the convolution matrix
158 * @param[in] mat9 Coefficient from the convolution matrix
159 * @param[in] mat10 Coefficient from the convolution matrix
160 * @param[in] mat11 Coefficient from the convolution matrix
161 * @param[in] mat12 Coefficient from the convolution matrix
162 * @param[in] mat13 Coefficient from the convolution matrix
163 * @param[in] mat14 Coefficient from the convolution matrix
164 * @param[in] mat15 Coefficient from the convolution matrix
165 * @param[in] mat16 Coefficient from the convolution matrix
166 * @param[in] mat10 Coefficient from the convolution matrix
167 * @param[in] mat17 Coefficient from the convolution matrix
168 * @param[in] mat18 Coefficient from the convolution matrix
169 * @param[in] mat19 Coefficient from the convolution matrix
170 * @param[in] mat20 Coefficient from the convolution matrix
171 * @param[in] mat21 Coefficient from the convolution matrix
172 * @param[in] mat22 Coefficient from the convolution matrix
173 * @param[in] mat23 Coefficient from the convolution matrix
174 * @param[in] mat24 Coefficient from the convolution matrix
175 * @param[in] mat25 Coefficient from the convolution matrix
176 * @param[in] mat26 Coefficient from the convolution matrix
177 * @param[in] mat27 Coefficient from the convolution matrix
178 * @param[in] mat28 Coefficient from the convolution matrix
179 * @param[in] mat29 Coefficient from the convolution matrix
180 * @param[in] mat30 Coefficient from the convolution matrix
181 * @param[in] mat31 Coefficient from the convolution matrix
182 * @param[in] mat32 Coefficient from the convolution matrix
183 * @param[in] mat33 Coefficient from the convolution matrix
184 * @param[in] mat34 Coefficient from the convolution matrix
185 * @param[in] mat35 Coefficient from the convolution matrix
186 * @param[in] mat36 Coefficient from the convolution matrix
187 * @param[in] mat37 Coefficient from the convolution matrix
188 * @param[in] mat38 Coefficient from the convolution matrix
189 * @param[in] mat39 Coefficient from the convolution matrix
190 * @param[in] mat40 Coefficient from the convolution matrix
191 * @param[in] mat41 Coefficient from the convolution matrix
192 * @param[in] mat42 Coefficient from the convolution matrix
193 * @param[in] mat43 Coefficient from the convolution matrix
194 * @param[in] mat44 Coefficient from the convolution matrix
195 * @param[in] mat45 Coefficient from the convolution matrix
196 * @param[in] mat46 Coefficient from the convolution matrix
197 * @param[in] mat47 Coefficient from the convolution matrix
198 * @param[in] mat48 Coefficient from the convolution matrix
199 * @param[in] scale Convolution matrix scale (Sum of the coefficients, or 1 if the sum is 0)
202 short8 convolution7x7(
204 const short mat0, const short mat1, const short mat2, const short mat3, const short mat4,
205 const short mat5, const short mat6, const short mat7, const short mat8, const short mat9,
206 const short mat10, const short mat11, const short mat12, const short mat13, const short mat14,
207 const short mat15, const short mat16, const short mat17, const short mat18, const short mat19,
208 const short mat20, const short mat21, const short mat22, const short mat23, const short mat24,
209 const short mat25, const short mat26, const short mat27, const short mat28, const short mat29,
210 const short mat30, const short mat31, const short mat32, const short mat33, const short mat34,
211 const short mat35, const short mat36, const short mat37, const short mat38, const short mat39,
212 const short mat40, const short mat41, const short mat42, const short mat43, const short mat44,
213 const short mat45, const short mat46, const short mat47, const short mat48, uint scale)
215 VEC_DATA_TYPE(DATA_TYPE, 8)
218 pixels = convolution1x7(offset(src, -3, -3), mat0, mat1, mat2, mat3, mat4, mat5, mat6);
219 pixels += convolution1x7(offset(src, -3, -2), mat7, mat8, mat9, mat10, mat11, mat12, mat13);
220 pixels += convolution1x7(offset(src, -3, -1), mat14, mat15, mat16, mat17, mat18, mat19, mat20);
221 pixels += convolution1x7(offset(src, -3, 0), mat21, mat22, mat23, mat24, mat25, mat26, mat27);
222 pixels += convolution1x7(offset(src, -3, 1), mat28, mat29, mat30, mat31, mat32, mat33, mat34);
223 pixels += convolution1x7(offset(src, -3, 2), mat35, mat36, mat37, mat38, mat39, mat40, mat41);
224 pixels += convolution1x7(offset(src, -3, 3), mat42, mat43, mat44, mat45, mat46, mat47, mat48);
228 pixels /= (VEC_DATA_TYPE(DATA_TYPE, 8))scale;
231 return convert_short8_sat(pixels);
234 #ifndef DYNAMIC_MATRIX_CONVOLUTION
236 /** Apply a 1x7 static convolution matrix to a single channel U8 input image and output a single temporary channel image.
238 * @attention The matrix coefficients (MAT0, MAT1, MAT2, MAT3, MAT4, MAT5, MAT6) and DATA_TYPE need to be passed at compile time:\n
239 * e.g. -DMAT0=1 -DMAT1=2, ... -DMAT6=6, -DDATA_TYPE=int
241 * @param[in] src_ptr Pointer to the source image. Supported data types: U8
242 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes)
243 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
244 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)
245 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
246 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image
247 * @param[out] dst_ptr Pointer to the destination image. Supported data types: U16, S16, S32
248 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes)
249 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
250 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes)
251 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
252 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image
254 __kernel void convolution_separable1x7_static(
255 IMAGE_DECLARATION(src),
256 IMAGE_DECLARATION(dst))
258 Image src = CONVERT_TO_IMAGE_STRUCT(src);
259 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
262 VEC_DATA_TYPE(DATA_TYPE, 8)
263 pixels = convolution1x7(offset(&src, -3, 0), MAT0, MAT1, MAT2, MAT3, MAT4, MAT5, MAT6);
265 // Store result in dst
266 vstore8(pixels, 0, (__global DATA_TYPE *)dst.ptr);
269 /** Apply a 7x1 static convolution matrix to a single channel U8 input image and output a single channel image.
271 * @attention The matrix coefficients (MAT7, MAT8, MAT9, MAT10, MAT11, MAT12, MAT13, SCALE), COMPUTE_TYPE and DATA_TYPE_OUT need to be passed at compile time:\n
272 * e.g. -DMAT0=7 -DMAT1=8, ... -DMAT24=13, -DSCALE=6, -DCOMPUTE_TYPE=int, -DDATA_TYPE_OUT=int
274 * @param[in] src_ptr Pointer to the source image. Supported data types: U16, S16, S32
275 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes)
276 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
277 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)
278 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
279 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image
280 * @param[out] dst_ptr Pointer to the destination image. Supported data types: U8, S16
281 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes)
282 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
283 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes)
284 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
285 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image
287 __kernel void convolution_separable7x1_static(
288 IMAGE_DECLARATION(src),
289 IMAGE_DECLARATION(dst))
291 Image src = CONVERT_TO_IMAGE_STRUCT(src);
292 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
295 VEC_DATA_TYPE(COMPUTE_TYPE, 8)
296 pixels = convolution7x1(&src, MAT7, MAT8, MAT9, MAT10, MAT11, MAT12, MAT13);
298 // Divide by the scale
299 pixels /= (VEC_DATA_TYPE(COMPUTE_TYPE, 8))SCALE;
301 // Store result in dst
302 vstore8(CONVERT_SAT(pixels, VEC_DATA_TYPE(DATA_TYPE_OUT, 8)), 0, (__global DATA_TYPE_OUT *)dst.ptr);
305 /** Apply a static 7x7 convolution matrix to a single channel U8 input image and output a single channel U8 image including the borders.
307 * @attention The matrix coefficients(MAT0, MAT1, ... MAT48, SCALE), DATA_TYPE_OUT need to be passed at compile time:\n
308 * e.g. -DMAT0=7 -DMAT1=8, ... -DMAT48=48, -DSCALE=6, -DDATA_TYPE_OUT=int
310 * @param[in] src_ptr Pointer to the source image. Supported data types: U8
311 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes)
312 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
313 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)
314 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
315 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image
316 * @param[out] dst_ptr Pointer to the destination image. Supported data types: U8, S16
317 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes)
318 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
319 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes)
320 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
321 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image
323 __kernel void convolution7x7_static(
324 IMAGE_DECLARATION(src),
325 IMAGE_DECLARATION(dst))
327 Image src = CONVERT_TO_IMAGE_STRUCT(src);
328 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
330 short8 pixels = convolution7x7(&src,
331 MAT0, MAT1, MAT2, MAT3, MAT4, MAT5, MAT6, MAT7, MAT8, MAT9, MAT10, MAT11, MAT12, MAT13,
332 MAT14, MAT15, MAT16, MAT17, MAT18, MAT19, MAT20, MAT21, MAT22, MAT23, MAT24, MAT25,
333 MAT26, MAT27, MAT28, MAT29, MAT30, MAT31, MAT32, MAT33, MAT34, MAT35, MAT36, MAT37,
334 MAT38, MAT39, MAT40, MAT41, MAT42, MAT43, MAT44, MAT45, MAT46, MAT47, MAT48, SCALE);
336 // Clamp results to [ 0, 255 ] and store them in dst
337 vstore8(CONVERT_SAT(pixels, VEC_DATA_TYPE(DATA_TYPE_OUT, 8)), 0, (__global DATA_TYPE_OUT *)dst.ptr);
340 #endif // DYNAMIC_MATRIX_CONVOLUTION