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 5 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 left pixel 43 * @param[in] middle_coeff Weight of the middle pixel 44 * @param[in] right1_coeff Weight of the right pixel 45 * @param[in] right2_coeff Weight of the most right pixel 46 * 47 * @return a short8 containing 8 convoluted values. 48 */ 49VEC_DATA_TYPE(DATA_TYPE, 8) 50convolution1x5( 51 __global const uchar *left_pixel, 52 const short left1_coeff, 53 const short left2_coeff, 54 const short middle_coeff, 55 const short right1_coeff, 56 const short right2_coeff) 57{ 58 uchar16 temp = vload16(0, left_pixel); 59 60 VEC_DATA_TYPE(DATA_TYPE, 8) 61 left1 = CONVERT(temp.s01234567, VEC_DATA_TYPE(DATA_TYPE, 8)); 62 VEC_DATA_TYPE(DATA_TYPE, 8) 63 left2 = CONVERT(temp.s12345678, VEC_DATA_TYPE(DATA_TYPE, 8)); 64 VEC_DATA_TYPE(DATA_TYPE, 8) 65 middle = CONVERT(temp.s23456789, VEC_DATA_TYPE(DATA_TYPE, 8)); 66 VEC_DATA_TYPE(DATA_TYPE, 8) 67 right1 = CONVERT(temp.s3456789a, VEC_DATA_TYPE(DATA_TYPE, 8)); 68 VEC_DATA_TYPE(DATA_TYPE, 8) 69 right2 = CONVERT(temp.s456789ab, VEC_DATA_TYPE(DATA_TYPE, 8)); 70 71 return left1 * (VEC_DATA_TYPE(DATA_TYPE, 8))left1_coeff + left2 * (VEC_DATA_TYPE(DATA_TYPE, 8))left2_coeff 72 + middle * (VEC_DATA_TYPE(DATA_TYPE, 8))middle_coeff + right1 * (VEC_DATA_TYPE(DATA_TYPE, 8))right1_coeff + right2 * (VEC_DATA_TYPE(DATA_TYPE, 8))right2_coeff; 73} 74 75/** Compute a 1D vertical convolution of size 5 for 8 bytes assuming the input is made of 1 channel of 1 byte (i.e 8 pixels). 76 * 77 * @param[in] src Pointer to source image. 78 * @param[in] up1_coeff Weight of the most up pixel 79 * @param[in] up2_coeff Weight of the up pixel 80 * @param[in] middle_coeff Weight of the middle pixel 81 * @param[in] down1_coeff Weight of the down pixel 82 * @param[in] down2_coeff Weight of the most down pixel 83 * 84 * @return a short8 containing 8 convoluted values. 85 */ 86VEC_DATA_TYPE(COMPUTE_TYPE, 8) 87convolution5x1( 88 Image *src, 89 const short up1_coeff, 90 const short up2_coeff, 91 const short middle_coeff, 92 const short down1_coeff, 93 const short down2_coeff) 94{ 95 VEC_DATA_TYPE(COMPUTE_TYPE, 8) 96 val; 97 VEC_DATA_TYPE(COMPUTE_TYPE, 8) 98 out = (VEC_DATA_TYPE(COMPUTE_TYPE, 8))0; 99 100 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, -2)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 101 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))up1_coeff; 102 103 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, -1)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 104 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))up2_coeff; 105 106 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 0)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 107 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))middle_coeff; 108 109 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 1)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 110 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))down1_coeff; 111 112 val = CONVERT(vload8(0, (__global DATA_TYPE *)offset(src, 0, 2)), VEC_DATA_TYPE(COMPUTE_TYPE, 8)); 113 out += val * (VEC_DATA_TYPE(COMPUTE_TYPE, 8))down2_coeff; 114 115 return out; 116} 117 118/** Apply a 5x5 convolution matrix to a single channel U8 input image and return the result. 119 * 120 * Convolution matrix layout:\n 121 * [ mat0, mat1, mat2, mat3 , mat4 ]\n 122 * [ mat5, mat6, mat7, mat8, mat9 ]\n 123 * [ mat10, mat11, mat12, mat13, mat14 ]\n 124 * [ mat15, mat16, mat17, mat18, mat19 ]\n 125 * [ mat20, mat21, mat22, mat23, mat24 ] 126 * 127 * @param[in] src A pointer to source Image structure. 128 * @param[in] mat0 Coefficient from the convolution matrix 129 * @param[in] mat1 Coefficient from the convolution matrix 130 * @param[in] mat2 Coefficient from the convolution matrix 131 * @param[in] mat3 Coefficient from the convolution matrix 132 * @param[in] mat4 Coefficient from the convolution matrix 133 * @param[in] mat5 Coefficient from the convolution matrix 134 * @param[in] mat6 Coefficient from the convolution matrix 135 * @param[in] mat7 Coefficient from the convolution matrix 136 * @param[in] mat8 Coefficient from the convolution matrix 137 * @param[in] mat9 Coefficient from the convolution matrix 138 * @param[in] mat10 Coefficient from the convolution matrix 139 * @param[in] mat11 Coefficient from the convolution matrix 140 * @param[in] mat12 Coefficient from the convolution matrix 141 * @param[in] mat13 Coefficient from the convolution matrix 142 * @param[in] mat14 Coefficient from the convolution matrix 143 * @param[in] mat15 Coefficient from the convolution matrix 144 * @param[in] mat16 Coefficient from the convolution matrix 145 * @param[in] mat17 Coefficient from the convolution matrix 146 * @param[in] mat18 Coefficient from the convolution matrix 147 * @param[in] mat19 Coefficient from the convolution matrix 148 * @param[in] mat20 Coefficient from the convolution matrix 149 * @param[in] mat21 Coefficient from the convolution matrix 150 * @param[in] mat22 Coefficient from the convolution matrix 151 * @param[in] mat23 Coefficient from the convolution matrix 152 * @param[in] mat24 Coefficient from the convolution matrix 153 * @param[in] scale Convolution matrix scale (Sum of the coefficients, or 1 if the sum is 0) 154 * 155 * @return a short8 containing 8 convoluted and scaled values. 156 */ 157short8 convolution5x5( 158 Image *src, 159 const short mat0, const short mat1, const short mat2, const short mat3, const short mat4, 160 const short mat5, const short mat6, const short mat7, const short mat8, const short mat9, 161 const short mat10, const short mat11, const short mat12, const short mat13, const short mat14, 162 const short mat15, const short mat16, const short mat17, const short mat18, const short mat19, 163 const short mat20, const short mat21, const short mat22, const short mat23, const short mat24, 164 uint scale) 165{ 166 VEC_DATA_TYPE(DATA_TYPE, 8) 167 pixels; 168 169 pixels = convolution1x5(offset(src, -2, -2), mat0, mat1, mat2, mat3, mat4); 170 pixels += convolution1x5(offset(src, -2, -1), mat5, mat6, mat7, mat8, mat9); 171 pixels += convolution1x5(offset(src, -2, 0), mat10, mat11, mat12, mat13, mat14); 172 pixels += convolution1x5(offset(src, -2, 1), mat15, mat16, mat17, mat18, mat19); 173 pixels += convolution1x5(offset(src, -2, 2), mat20, mat21, mat22, mat23, mat24); 174 175 if(scale > 0) 176 { 177 pixels /= (VEC_DATA_TYPE(DATA_TYPE, 8))scale; 178 } 179 180 return convert_short8_sat(pixels); 181} 182 183#ifndef DYNAMIC_MATRIX_CONVOLUTION 184 185/** Apply a 1x5 static convolution matrix to a single channel U8 input image and output a single temporary channel image(Support U16, S16, S32). 186 * 187 * @attention The matrix coefficients (MAT0, MAT1, MAT2, MAT3, MAT4) and DATA_TYPE need to be passed at compile time:\n 188 * e.g. -DMAT0=1 -DMAT2=2, -DMAT3=3, -DMAT4=4, -DDATA_TYPE=int 189 * 190 * @param[in] src_ptr Pointer to the source image. Supported data types: U8 191 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes) 192 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) 193 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes) 194 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) 195 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image 196 * @param[out] dst_ptr Pointer to the destination image. Supported data types: U16, S16, S32 197 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes) 198 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) 199 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes) 200 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) 201 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image 202 */ 203__kernel void convolution_separable1x5_static( 204 IMAGE_DECLARATION(src), 205 IMAGE_DECLARATION(dst)) 206{ 207 Image src = CONVERT_TO_IMAGE_STRUCT(src); 208 Image dst = CONVERT_TO_IMAGE_STRUCT(dst); 209 210 // Output pixels 211 VEC_DATA_TYPE(DATA_TYPE, 8) 212 pixels = convolution1x5(offset(&src, -2, 0), MAT0, MAT1, MAT2, MAT3, MAT4); 213 214 // Store result in dst 215 vstore8(pixels, 0, (__global DATA_TYPE *)dst.ptr); 216} 217 218/** Apply a 5x1 static convolution matrix to a single channel U8 input image and output a single channel image. 219 * 220 * @attention The matrix coefficients (MAT5, MAT6, MAT7, MAT8, MAT9, SCALE), COMPUTE_TYPE and DATA_TYPE_OUT need to be passed at compile time:\n 221 * e.g. -DMAT5=1 -DMAT6=2, -DMAT7=3, -DMAT8=4, -DMAT9=5, -DSCALE=6, -DCOMPUTE_TYPE=int, -DDATA_TYPE_OUT=int 222 * 223 * @param[in] src_ptr Pointer to the source image. Supported data types: U16, S16, S32 224 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes) 225 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) 226 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes) 227 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) 228 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image 229 * @param[out] dst_ptr Pointer to the destination image. Supported data types: U8, S16 230 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes) 231 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) 232 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes) 233 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) 234 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image 235 */ 236__kernel void convolution_separable5x1_static( 237 IMAGE_DECLARATION(src), 238 IMAGE_DECLARATION(dst)) 239{ 240 Image src = CONVERT_TO_IMAGE_STRUCT(src); 241 Image dst = CONVERT_TO_IMAGE_STRUCT(dst); 242 243 // Output pixels 244 VEC_DATA_TYPE(COMPUTE_TYPE, 8) 245 pixels = convolution5x1(&src, MAT5, MAT6, MAT7, MAT8, MAT9); 246 247 // Divide by the scale 248 pixels /= (VEC_DATA_TYPE(COMPUTE_TYPE, 8))SCALE; 249 250 // Store result in dst 251 vstore8(CONVERT_SAT(pixels, VEC_DATA_TYPE(DATA_TYPE_OUT, 8)), 0, (__global DATA_TYPE_OUT *)dst.ptr); 252} 253 254/** Apply a static 5x5 convolution matrix to a single channel U8 input image and output a single channel image including borders 255 * 256 * @attention The matrix coefficients(MAT0, MAT1, ... MAT24, SCALE), DATA_TYPE_OUT need to be passed at compile time:\n 257 * e.g. -DMAT0=1 -DMAT1=2, ... -DMAT24=24, -DSCALE=6, -DDATA_TYPE_OUT=int 258 * 259 * @param[in] src_ptr Pointer to the source image. Supported data types: U8 260 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes) 261 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) 262 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes) 263 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) 264 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image 265 * @param[out] dst_ptr Pointer to the destination image. Supported data types: U8, S16 266 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes) 267 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) 268 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes) 269 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) 270 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image 271 */ 272__kernel void convolution5x5_static( 273 IMAGE_DECLARATION(src), 274 IMAGE_DECLARATION(dst)) 275{ 276 Image src = CONVERT_TO_IMAGE_STRUCT(src); 277 Image dst = CONVERT_TO_IMAGE_STRUCT(dst); 278 279 short8 pixels = convolution5x5(&src, 280 MAT0, MAT1, MAT2, MAT3, MAT4, MAT5, MAT6, MAT7, MAT8, MAT9, MAT10, MAT11, MAT12, MAT13, 281 MAT14, MAT15, MAT16, MAT17, MAT18, MAT19, MAT20, MAT21, MAT22, MAT23, MAT24, SCALE); 282 283 // Store the result as is in dst 284 vstore8(CONVERT_SAT(pixels, VEC_DATA_TYPE(DATA_TYPE_OUT, 8)), 0, (__global DATA_TYPE_OUT *)dst.ptr); 285} 286 287#endif // DYNAMIC_MATRIX_CONVOLUTION 288