1/* 2 * Copyright (c) 2016-2019 Arm Limited. 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.h" 25 26#ifndef DATA_TYPE 27#define DATA_TYPE short 28#endif /* DATA_TYPE */ 29 30#ifndef COMPUTE_TYPE 31#define COMPUTE_TYPE int 32#endif /* COMPUTE_TYPE */ 33 34#ifndef DATA_TYPE_OUT 35#define DATA_TYPE_OUT uchar 36#endif /* DATA_TYPE_OUT */ 37 38/** Compute a 1D horizontal convolution of size 9 for 8 bytes assuming the input is made of 1 channel of 1 byte (i.e 8 pixels). 39 * 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 third left pixel 44 * @param[in] left4_coeff Weight of the left pixel 45 * @param[in] middle_coeff Weight of the middle pixel 46 * @param[in] right1_coeff Weight of the right pixel 47 * @param[in] right2_coeff Weight of the second right pixel 48 * @param[in] right3_coeff Weight of the third right pixel 49 * @param[in] right4_coeff Weight of the most right pixel 50 * 51 * @return a short8 containing 8 convoluted values. 52 */ 53VEC_DATA_TYPE(DATA_TYPE, 8) 54convolution1x9( 55 __global const uchar *left_pixel, 56 const short left1_coeff, 57 const short left2_coeff, 58 const short left3_coeff, 59 const short left4_coeff, 60 const short middle_coeff, 61 const short right1_coeff, 62 const short right2_coeff, 63 const short right3_coeff, 64 const short right4_coeff) 65{ 66 uchar16 temp = vload16(0, left_pixel); 67 68 VEC_DATA_TYPE(DATA_TYPE, 8) 69 left1 = CONVERT(temp.s01234567, VEC_DATA_TYPE(DATA_TYPE, 8)); 70 VEC_DATA_TYPE(DATA_TYPE, 8) 71 left2 = CONVERT(temp.s12345678, VEC_DATA_TYPE(DATA_TYPE, 8)); 72 VEC_DATA_TYPE(DATA_TYPE, 8) 73 left3 = CONVERT(temp.s23456789, VEC_DATA_TYPE(DATA_TYPE, 8)); 74 VEC_DATA_TYPE(DATA_TYPE, 8) 75 left4 = CONVERT(temp.s3456789a, VEC_DATA_TYPE(DATA_TYPE, 8)); 76 VEC_DATA_TYPE(DATA_TYPE, 8) 77 middle = CONVERT(temp.s456789ab, VEC_DATA_TYPE(DATA_TYPE, 8)); 78 VEC_DATA_TYPE(DATA_TYPE, 8) 79 right1 = CONVERT(temp.s56789abc, VEC_DATA_TYPE(DATA_TYPE, 8)); 80 VEC_DATA_TYPE(DATA_TYPE, 8) 81 right2 = CONVERT(temp.s6789abcd, VEC_DATA_TYPE(DATA_TYPE, 8)); 82 VEC_DATA_TYPE(DATA_TYPE, 8) 83 right3 = CONVERT(temp.s789abcde, VEC_DATA_TYPE(DATA_TYPE, 8)); 84 VEC_DATA_TYPE(DATA_TYPE, 8) 85 right4 = CONVERT(temp.s89abcdef, VEC_DATA_TYPE(DATA_TYPE, 8)); 86 87 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 + left4 * (VEC_DATA_TYPE(DATA_TYPE, 88 8))left4_coeff + middle * (VEC_DATA_TYPE(DATA_TYPE, 8))middle_coeff + right1 * (VEC_DATA_TYPE(DATA_TYPE, 8))right1_coeff + right2 * (VEC_DATA_TYPE(DATA_TYPE, 89 8))right2_coeff + right3 * (VEC_DATA_TYPE(DATA_TYPE, 8))right3_coeff + right4 * (VEC_DATA_TYPE(DATA_TYPE, 8))right4_coeff; 90} 91 92/** Compute a 1D vertical convolution of size 9 for 8 bytes assuming the input is made of 1 channel of 1 byte (i.e 8 pixels). 93 * 94 * @param[in] src Pointer to source image. 95 * @param[in] up1_coeff Weight of the most up pixel 96 * @param[in] up2_coeff Weight of the second up pixel 97 * @param[in] up3_coeff Weight of the third up pixel 98 * @param[in] up4_coeff Weight of the up pixel 99 * @param[in] middle_coeff Weight of the middle pixel 100 * @param[in] down1_coeff Weight of the down pixel 101 * @param[in] down2_coeff Weight of the second down pixel 102 * @param[in] down3_coeff Weight of the third down pixel 103 * @param[in] down4_coeff Weight of the most down pixel 104 * 105 * @return a short8 containing 8 convoluted values. 106 */ 107VEC_DATA_TYPE(COMPUTE_TYPE, 8) 108convolution9x1( 109 Image *src, 110 const short up1_coeff, 111 const short up2_coeff, 112 const short up3_coeff, 113 const short up4_coeff, 114 const short middle_coeff, 115 const short down1_coeff, 116 const short down2_coeff, 117 const short down3_coeff, 118 const short down4_coeff) 119{ 120 VEC_DATA_TYPE(COMPUTE_TYPE, 8) 121 val; 122 VEC_DATA_TYPE(COMPUTE_TYPE, 8) 123 out = (VEC_DATA_TYPE(COMPUTE_TYPE, 8))0; 124 125 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, -4)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 126 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))up1_coeff; 127 128 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, -3)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 129 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))up2_coeff; 130 131 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, -2)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 132 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))up3_coeff; 133 134 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, -1)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 135 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))up4_coeff; 136 137 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 0)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 138 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))middle_coeff; 139 140 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 1)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 141 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))down1_coeff; 142 143 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 2)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 144 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))down2_coeff; 145 146 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 3)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 147 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))down3_coeff; 148 149 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 4)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 150 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))down4_coeff; 151 152 return out; 153} 154 155/** Apply a 9x9 convolution matrix to a single channel U8 input image and return the result. 156 * 157 * Convolution matrix layout:\n 158 * [ mat0, mat1, mat2, mat3 , mat4, mat5, mat6, mat7, mat8 ]\n 159 * [ mat9, mat10, mat11, mat12, mat13, mat14, mat15, mat16, mat17 ]\n 160 * [ mat18, mat19, mat20, mat21, mat22, mat23, mat24, mat25, mat26 ]\n 161 * [ mat27, mat28, mat29, mat30, mat31, mat32, mat33, mat34, mat35 ]\n 162 * [ mat36, mat37, mat38, mat39, mat40, mat41, mat42, mat43, mat44 ]\n 163 * [ mat45, mat46, mat47, mat48, mat49, mat50, mat51, mat52, mat53 ]\n 164 * [ mat54, mat55, mat56, mat57, mat58, mat59, mat60, mat61, mat62 ] 165 * [ mat63, mat64, mat65, mat66, mat67, mat68, mat69, mat70, mat71 ] 166 * [ mat72, mat73, mat74, mat75, mat76, mat77, mat78, mat79, mat80 ] 167 * 168 * @param[in] src A pointer to source Image structure. 169 * @param[in] mat0 Coefficient from the convolution matrix 170 * @param[in] mat1 Coefficient from the convolution matrix 171 * @param[in] mat2 Coefficient from the convolution matrix 172 * @param[in] mat3 Coefficient from the convolution matrix 173 * @param[in] mat4 Coefficient from the convolution matrix 174 * @param[in] mat5 Coefficient from the convolution matrix 175 * @param[in] mat6 Coefficient from the convolution matrix 176 * @param[in] mat7 Coefficient from the convolution matrix 177 * @param[in] mat8 Coefficient from the convolution matrix 178 * @param[in] mat9 Coefficient from the convolution matrix 179 * @param[in] mat10 Coefficient from the convolution matrix 180 * @param[in] mat11 Coefficient from the convolution matrix 181 * @param[in] mat12 Coefficient from the convolution matrix 182 * @param[in] mat13 Coefficient from the convolution matrix 183 * @param[in] mat14 Coefficient from the convolution matrix 184 * @param[in] mat15 Coefficient from the convolution matrix 185 * @param[in] mat16 Coefficient from the convolution matrix 186 * @param[in] mat17 Coefficient from the convolution matrix 187 * @param[in] mat18 Coefficient from the convolution matrix 188 * @param[in] mat19 Coefficient from the convolution matrix 189 * @param[in] mat20 Coefficient from the convolution matrix 190 * @param[in] mat21 Coefficient from the convolution matrix 191 * @param[in] mat22 Coefficient from the convolution matrix 192 * @param[in] mat23 Coefficient from the convolution matrix 193 * @param[in] mat24 Coefficient from the convolution matrix 194 * @param[in] mat25 Coefficient from the convolution matrix 195 * @param[in] mat26 Coefficient from the convolution matrix 196 * @param[in] mat27 Coefficient from the convolution matrix 197 * @param[in] mat28 Coefficient from the convolution matrix 198 * @param[in] mat29 Coefficient from the convolution matrix 199 * @param[in] mat30 Coefficient from the convolution matrix 200 * @param[in] mat31 Coefficient from the convolution matrix 201 * @param[in] mat32 Coefficient from the convolution matrix 202 * @param[in] mat33 Coefficient from the convolution matrix 203 * @param[in] mat34 Coefficient from the convolution matrix 204 * @param[in] mat35 Coefficient from the convolution matrix 205 * @param[in] mat36 Coefficient from the convolution matrix 206 * @param[in] mat37 Coefficient from the convolution matrix 207 * @param[in] mat38 Coefficient from the convolution matrix 208 * @param[in] mat39 Coefficient from the convolution matrix 209 * @param[in] mat40 Coefficient from the convolution matrix 210 * @param[in] mat41 Coefficient from the convolution matrix 211 * @param[in] mat42 Coefficient from the convolution matrix 212 * @param[in] mat43 Coefficient from the convolution matrix 213 * @param[in] mat44 Coefficient from the convolution matrix 214 * @param[in] mat45 Coefficient from the convolution matrix 215 * @param[in] mat46 Coefficient from the convolution matrix 216 * @param[in] mat47 Coefficient from the convolution matrix 217 * @param[in] mat48 Coefficient from the convolution matrix 218 * @param[in] mat49 Coefficient from the convolution matrix 219 * @param[in] mat50 Coefficient from the convolution matrix 220 * @param[in] mat51 Coefficient from the convolution matrix 221 * @param[in] mat52 Coefficient from the convolution matrix 222 * @param[in] mat53 Coefficient from the convolution matrix 223 * @param[in] mat54 Coefficient from the convolution matrix 224 * @param[in] mat55 Coefficient from the convolution matrix 225 * @param[in] mat56 Coefficient from the convolution matrix 226 * @param[in] mat57 Coefficient from the convolution matrix 227 * @param[in] mat58 Coefficient from the convolution matrix 228 * @param[in] mat59 Coefficient from the convolution matrix 229 * @param[in] mat60 Coefficient from the convolution matrix 230 * @param[in] mat61 Coefficient from the convolution matrix 231 * @param[in] mat62 Coefficient from the convolution matrix 232 * @param[in] mat63 Coefficient from the convolution matrix 233 * @param[in] mat64 Coefficient from the convolution matrix 234 * @param[in] mat65 Coefficient from the convolution matrix 235 * @param[in] mat66 Coefficient from the convolution matrix 236 * @param[in] mat67 Coefficient from the convolution matrix 237 * @param[in] mat68 Coefficient from the convolution matrix 238 * @param[in] mat69 Coefficient from the convolution matrix 239 * @param[in] mat70 Coefficient from the convolution matrix 240 * @param[in] mat71 Coefficient from the convolution matrix 241 * @param[in] mat72 Coefficient from the convolution matrix 242 * @param[in] mat73 Coefficient from the convolution matrix 243 * @param[in] mat74 Coefficient from the convolution matrix 244 * @param[in] mat75 Coefficient from the convolution matrix 245 * @param[in] mat76 Coefficient from the convolution matrix 246 * @param[in] mat77 Coefficient from the convolution matrix 247 * @param[in] mat78 Coefficient from the convolution matrix 248 * @param[in] mat79 Coefficient from the convolution matrix 249 * @param[in] mat80 Coefficient from the convolution matrix 250 * @param[in] scale Convolution matrix scale (Sum of the coefficients, or 1 if the sum is 0) 251 * 252 */ 253short8 convolution9x9( 254 Image *src, 255 const short mat0, const short mat1, const short mat2, const short mat3, const short mat4, 256 const short mat5, const short mat6, const short mat7, const short mat8, const short mat9, 257 const short mat10, const short mat11, const short mat12, const short mat13, const short mat14, 258 const short mat15, const short mat16, const short mat17, const short mat18, const short mat19, 259 const short mat20, const short mat21, const short mat22, const short mat23, const short mat24, 260 const short mat25, const short mat26, const short mat27, const short mat28, const short mat29, 261 const short mat30, const short mat31, const short mat32, const short mat33, const short mat34, 262 const short mat35, const short mat36, const short mat37, const short mat38, const short mat39, 263 const short mat40, const short mat41, const short mat42, const short mat43, const short mat44, 264 const short mat45, const short mat46, const short mat47, const short mat48, const short mat49, 265 const short mat50, const short mat51, const short mat52, const short mat53, const short mat54, 266 const short mat55, const short mat56, const short mat57, const short mat58, const short mat59, 267 const short mat60, const short mat61, const short mat62, const short mat63, const short mat64, 268 const short mat65, const short mat66, const short mat67, const short mat68, const short mat69, 269 const short mat70, const short mat71, const short mat72, const short mat73, const short mat74, 270 const short mat75, const short mat76, const short mat77, const short mat78, const short mat79, 271 const short mat80, uint scale) 272{ 273 VEC_DATA_TYPE(DATA_TYPE, 8) 274 pixels; 275 276 pixels = convolution1x9(offset(src, -4, -4), mat0, mat1, mat2, mat3, mat4, mat5, mat6, mat7, mat8); 277 pixels += convolution1x9(offset(src, -4, -3), mat9, mat10, mat11, mat12, mat13, mat14, mat15, mat16, mat17); 278 pixels += convolution1x9(offset(src, -4, -2), mat18, mat19, mat20, mat21, mat22, mat23, mat24, mat25, mat26); 279 pixels += convolution1x9(offset(src, -4, -1), mat27, mat28, mat29, mat30, mat31, mat32, mat33, mat34, mat35); 280 pixels += convolution1x9(offset(src, -4, 0), mat36, mat37, mat38, mat39, mat40, mat41, mat42, mat43, mat44); 281 pixels += convolution1x9(offset(src, -4, 1), mat45, mat46, mat47, mat48, mat49, mat50, mat51, mat52, mat53); 282 pixels += convolution1x9(offset(src, -4, 2), mat54, mat55, mat56, mat57, mat58, mat59, mat60, mat61, mat62); 283 pixels += convolution1x9(offset(src, -4, 3), mat63, mat64, mat65, mat66, mat67, mat68, mat69, mat70, mat71); 284 pixels += convolution1x9(offset(src, -4, 4), mat72, mat73, mat74, mat75, mat76, mat77, mat78, mat79, mat80); 285 286 if(scale > 0) 287 { 288 pixels /= (VEC_DATA_TYPE(DATA_TYPE, 8))scale; 289 } 290 291 return convert_short8_sat(pixels); 292} 293 294#ifndef DYNAMIC_MATRIX_CONVOLUTION 295 296/** Apply a 1x9 static convolution matrix to a single channel U8 input image and output a single temporary channel image. 297 * 298 * @attention The matrix coefficients (MAT0, MAT1, MAT2, MAT3, MAT4, MAT5, MAT6, MAT7, MAT8) and DATA_TYPE need to be passed at compile time:\n 299 * e.g. -DMAT0=7 -DMAT1=8, ... -DMAT8=8, -DCOMPUTE_TYPE=int 300 * 301 * @param[in] src_ptr Pointer to the source image. Supported data types: U8 302 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes) 303 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) 304 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes) 305 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) 306 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image 307 * @param[out] dst_ptr Pointer to the destination image. Supported data types: U16, S16, S32 308 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes) 309 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) 310 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes) 311 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) 312 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image 313 */ 314__kernel void convolution_separable1x9_static( 315 IMAGE_DECLARATION(src), 316 IMAGE_DECLARATION(dst)) 317{ 318 Image src = CONVERT_TO_IMAGE_STRUCT(src); 319 Image dst = CONVERT_TO_IMAGE_STRUCT(dst); 320 321 // Output pixels 322 VEC_DATA_TYPE(DATA_TYPE, 8) 323 pixels = convolution1x9(offset(&src, -4, 0), MAT0, MAT1, MAT2, MAT3, MAT4, MAT5, MAT6, MAT7, MAT8); 324 325 // Store result in dst 326 vstore8(pixels, 0, (__global DATA_TYPE *)dst.ptr); 327} 328 329/** Apply a 9x1 static convolution matrix to a single channel U8 input image and output a single channel image. 330 * 331 * @attention The matrix coefficients (MAT9, MAT10, ... MAT17, SCALE), COMPUTE_TYPE and DATA_TYPE_OUT need to be passed at compile time:\n 332 * e.g. -DMAT9=9 -DMAT10=10, ... -DMAT17=17, -DSCALE=6, -DCOMPUTE_TYPE=int, -DDATA_TYPE_OUT=int 333 * 334 * @param[in] src_ptr Pointer to the source image. Supported data types: U16, S16, S32 335 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes) 336 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) 337 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes) 338 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) 339 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image 340 * @param[out] dst_ptr Pointer to the destination image. Supported data types: U8, S16 341 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes) 342 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) 343 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes) 344 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) 345 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image 346 */ 347__kernel void convolution_separable9x1_static( 348 IMAGE_DECLARATION(src), 349 IMAGE_DECLARATION(dst)) 350{ 351 Image src = CONVERT_TO_IMAGE_STRUCT(src); 352 Image dst = CONVERT_TO_IMAGE_STRUCT(dst); 353 354 // Output pixels 355 VEC_DATA_TYPE(COMPUTE_TYPE, 8) 356 pixels = convolution9x1(&src, MAT9, MAT10, MAT11, MAT12, MAT13, MAT14, MAT15, MAT16, MAT17); 357 358 // Divide by the scale 359 pixels = pixels / (VEC_DATA_TYPE(COMPUTE_TYPE, 8))SCALE; 360 361 // Store result in dst 362 vstore8(CONVERT_SAT(pixels, VEC_DATA_TYPE(DATA_TYPE_OUT, 8)), 0, (__global DATA_TYPE_OUT *)dst.ptr); 363} 364 365/** Apply a static 9x9 convolution matrix to a single channel U8 input image and output a single channel image including borders 366 * 367 * @attention The matrix coefficients(MAT0, MAT1, ... MAT80, SCALE), DATA_TYPE_OUT need to be passed at compile time:\n 368 * e.g. -DMAT0=0 -DMAT1=1, ... -DMAT80=80, -DSCALE=6, -DDATA_TYPE_OUT=int 369 * 370 * @param[in] src_ptr Pointer to the source image. Supported data types: U8 371 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes) 372 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) 373 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes) 374 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) 375 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image 376 * @param[out] dst_ptr Pointer to the destination image. Supported data types: U8, S16 377 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes) 378 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) 379 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes) 380 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) 381 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image 382 */ 383__kernel void convolution9x9_static( 384 IMAGE_DECLARATION(src), 385 IMAGE_DECLARATION(dst)) 386{ 387 Image src = CONVERT_TO_IMAGE_STRUCT(src); 388 Image dst = CONVERT_TO_IMAGE_STRUCT(dst); 389 390 short8 pixels = convolution9x9(&src, 391 MAT0, MAT1, MAT2, MAT3, MAT4, MAT5, MAT6, MAT7, MAT8, MAT9, MAT10, MAT11, MAT12, MAT13, 392 MAT14, MAT15, MAT16, MAT17, MAT18, MAT19, MAT20, MAT21, MAT22, MAT23, MAT24, MAT25, 393 MAT26, MAT27, MAT28, MAT29, MAT30, MAT31, MAT32, MAT33, MAT34, MAT35, MAT36, MAT37, 394 MAT38, MAT39, MAT40, MAT41, MAT42, MAT43, MAT44, MAT45, MAT46, MAT47, MAT48, MAT49, 395 MAT50, MAT51, MAT52, MAT53, MAT54, MAT55, MAT56, MAT57, MAT58, MAT59, MAT60, MAT61, 396 MAT62, MAT63, MAT64, MAT65, MAT66, MAT67, MAT68, MAT69, MAT70, MAT71, MAT72, MAT73, 397 MAT74, MAT75, MAT76, MAT77, MAT78, MAT79, MAT80, SCALE); 398 399 // Store the result as is in dst 400 vstore8(CONVERT_SAT(pixels, VEC_DATA_TYPE(DATA_TYPE_OUT, 8)), 0, (__global DATA_TYPE_OUT *)dst.ptr); 401} 402 403#endif // DYNAMIC_MATRIX_CONVOLUTION 404