1R"( 2 3 4 5#ifndef ARM_COMPUTE_HELPER_H 6#define ARM_COMPUTE_HELPER_H 7 8 9 10 11#define STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 12 VSTORE(N0) \ 13 (BASENAME##0, 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0)); 14 15#define STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 16 STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 17 VSTORE(N0) \ 18 (BASENAME##1, 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1)); 19 20#define STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 21 STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 22 VSTORE(N0) \ 23 (BASENAME##2, 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2)); 24 25#define STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 26 STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 27 VSTORE(N0) \ 28 (BASENAME##3, 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3)); 29 30#define STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 31 STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 32 VSTORE(N0) \ 33 (BASENAME##4, 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4)); 34 35#define STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 36 STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 37 VSTORE(N0) \ 38 (BASENAME##5, 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5)); 39 40#define STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 41 STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 42 VSTORE(N0) \ 43 (BASENAME##6, 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6)); 44 45#define STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 46 STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 47 VSTORE(N0) \ 48 (BASENAME##7, 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7)); 49 50#define STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 51 STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 52 VSTORE(N0) \ 53 (BASENAME##8, 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8)); 54 55#define STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 56 STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 57 VSTORE(N0) \ 58 (BASENAME##9, 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9)); 59 60#define STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 61 STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 62 VSTORE(N0) \ 63 (BASENAME##A, 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A)); 64 65#define STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 66 STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 67 VSTORE(N0) \ 68 (BASENAME##B, 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B)); 69 70#define STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 71 STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 72 VSTORE(N0) \ 73 (BASENAME##C, 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C)); 74 75#define STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 76 STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 77 VSTORE(N0) \ 78 (BASENAME##D, 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D)); 79 80#define STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 81 STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 82 VSTORE(N0) \ 83 (BASENAME##E, 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E)); 84 85#define STORE_ROW_16(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 86 STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 87 VSTORE(N0) \ 88 (BASENAME##F, 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F)); 89 90 91 92#define CONVERT_STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 93 VSTORE(N0) \ 94 (CONVERT_SAT((BASENAME##0), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0)); 95 96#define CONVERT_STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 97 CONVERT_STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 98 VSTORE(N0) \ 99 (CONVERT_SAT((BASENAME##1), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1)); 100 101#define CONVERT_STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 102 CONVERT_STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 103 VSTORE(N0) \ 104 (CONVERT_SAT((BASENAME##2), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2)); 105 106#define CONVERT_STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 107 CONVERT_STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 108 VSTORE(N0) \ 109 (CONVERT_SAT((BASENAME##3), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3)); 110 111#define CONVERT_STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 112 CONVERT_STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 113 VSTORE(N0) \ 114 (CONVERT_SAT((BASENAME##4), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4)); 115 116#define CONVERT_STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 117 CONVERT_STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 118 VSTORE(N0) \ 119 (CONVERT_SAT((BASENAME##5), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5)); 120 121#define CONVERT_STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 122 CONVERT_STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 123 VSTORE(N0) \ 124 (CONVERT_SAT((BASENAME##6), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6)); 125 126#define CONVERT_STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 127 CONVERT_STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 128 VSTORE(N0) \ 129 (CONVERT_SAT((BASENAME##7), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7)); 130 131#define CONVERT_STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 132 CONVERT_STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 133 VSTORE(N0) \ 134 (CONVERT_SAT((BASENAME##8), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8)); 135 136#define CONVERT_STORE_ROW_10(N0, DATA, BASENAME, PTR, STRIDE_Y, Z) \ 137 CONVERT_STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 138 VSTORE(N0) \ 139 (CONVERT_SAT((BASENAME##9), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9)); 140 141#define CONVERT_STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 142 CONVERT_STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 143 VSTORE(N0) \ 144 (CONVERT_SAT((BASENAME##A), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A)); 145 146#define CONVERT_STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 147 CONVERT_STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 148 VSTORE(N0) \ 149 (CONVERT_SAT((BASENAME##B), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B)); 150 151#define CONVERT_STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 152 CONVERT_STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 153 VSTORE(N0) \ 154 (CONVERT_SAT((BASENAME##C), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C)); 155 156#define CONVERT_STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 157 CONVERT_STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 158 VSTORE(N0) \ 159 (CONVERT_SAT((BASENAME##D), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D)); 160 161#define CONVERT_STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 162 CONVERT_STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 163 VSTORE(N0) \ 164 (CONVERT_SAT((BASENAME##E), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E)); 165 166#define CONVERT_STORE_ROW_16(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 167 CONVERT_STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 168 VSTORE(N0) \ 169 (CONVERT_SAT((BASENAME##F), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F)); 170 171 172 173 174#define STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_ROW_##M0(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 175#define STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 176 177 178 179#define CONVERT_STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) CONVERT_STORE_ROW_##M0(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 180#define CONVERT_STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) CONVERT_STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 181 182 183 184#define STORE_ROW_PARTIAL_1(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 185 VSTORE_PARTIAL(N0, STORE_N0) \ 186 (BASENAME##0, 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0)); 187 188#define STORE_ROW_PARTIAL_2(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 189 STORE_ROW_PARTIAL_1(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 190 VSTORE_PARTIAL(N0, STORE_N0) \ 191 (BASENAME##1, 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1)); 192 193#define STORE_ROW_PARTIAL_3(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 194 STORE_ROW_PARTIAL_2(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 195 VSTORE_PARTIAL(N0, STORE_N0) \ 196 (BASENAME##2, 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2)); 197 198#define STORE_ROW_PARTIAL_4(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 199 STORE_ROW_PARTIAL_3(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 200 VSTORE_PARTIAL(N0, STORE_N0) \ 201 (BASENAME##3, 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3)); 202 203#define STORE_ROW_PARTIAL_5(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 204 STORE_ROW_PARTIAL_4(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 205 VSTORE_PARTIAL(N0, STORE_N0) \ 206 (BASENAME##4, 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4)); 207 208#define STORE_ROW_PARTIAL_6(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 209 STORE_ROW_PARTIAL_5(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 210 VSTORE_PARTIAL(N0, STORE_N0) \ 211 (BASENAME##5, 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5)); 212 213#define STORE_ROW_PARTIAL_7(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 214 STORE_ROW_PARTIAL_6(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 215 VSTORE_PARTIAL(N0, STORE_N0) \ 216 (BASENAME##6, 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6)); 217 218#define STORE_ROW_PARTIAL_8(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 219 STORE_ROW_PARTIAL_7(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 220 VSTORE_PARTIAL(N0, STORE_N0) \ 221 (BASENAME##7, 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7)); 222 223#define STORE_ROW_PARTIAL_9(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 224 STORE_ROW_PARTIAL_8(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 225 VSTORE_PARTIAL(N0, STORE_N0) \ 226 (BASENAME##8, 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8)); 227 228#define STORE_ROW_PARTIAL_10(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 229 STORE_ROW_PARTIAL_9(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 230 VSTORE_PARTIAL(N0, STORE_N0) \ 231 (BASENAME##9, 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9)); 232 233#define STORE_ROW_PARTIAL_11(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 234 STORE_ROW_PARTIAL_10(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 235 VSTORE_PARTIAL(N0, STORE_N0) \ 236 (BASENAME##A, 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A)); 237 238#define STORE_ROW_PARTIAL_12(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 239 STORE_ROW_PARTIAL_11(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 240 VSTORE_PARTIAL(N0, STORE_N0) \ 241 (BASENAME##B, 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B)); 242 243#define STORE_ROW_PARTIAL_13(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 244 STORE_ROW_PARTIAL_12(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 245 VSTORE_PARTIAL(N0, STORE_N0) \ 246 (BASENAME##C, 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C)); 247 248#define STORE_ROW_PARTIAL_14(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 249 STORE_ROW_PARTIAL_13(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 250 VSTORE_PARTIAL(N0, STORE_N0) \ 251 (BASENAME##D, 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D)); 252 253#define STORE_ROW_PARTIAL_15(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 254 STORE_ROW_PARTIAL_14(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 255 VSTORE_PARTIAL(N0, STORE_N0) \ 256 (BASENAME##E, 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E)); 257 258#define STORE_ROW_PARTIAL_16(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 259 STORE_ROW_PARTIAL_15(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 260 VSTORE_PARTIAL(N0, STORE_N0) \ 261 (BASENAME##F, 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F)); 262 263 264 265#define STORE_BLOCK_PARTIAL_STR(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_ROW_PARTIAL_##STORE_M0(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 266#define STORE_BLOCK_PARTIAL(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_BLOCK_PARTIAL_STR(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 267 268#define STORE_BLOCK_PARTIAL_IN_X_AND_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 269 if(!(PARTIAL_COND_X) && !(PARTIAL_COND_Y)) \ 270 { \ 271 STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 272 } \ 273 else if((PARTIAL_COND_Y) && !(PARTIAL_COND_X)) \ 274 { \ 275 STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 276 } \ 277 else if(!(PARTIAL_COND_Y) && (PARTIAL_COND_X)) \ 278 { \ 279 STORE_BLOCK_PARTIAL(M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 280 } \ 281 else \ 282 { \ 283 STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 284 } 285 286#define STORE_BLOCK_PARTIAL_IN_X(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_N0, PARTIAL_COND_X) \ 287 if(!(PARTIAL_COND_X)) \ 288 { \ 289 STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 290 } \ 291 else \ 292 { \ 293 STORE_BLOCK_PARTIAL(M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 294 } 295 296#define STORE_BLOCK_PARTIAL_IN_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_COND_Y) \ 297 if(!(PARTIAL_COND_Y)) \ 298 { \ 299 STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 300 } \ 301 else \ 302 { \ 303 STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 304 } 305 306 307#if defined(PARTIAL_STORE_M0) && defined(PARTIAL_STORE_N0) 308 309 310#if PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 == 0 311 312#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 313 STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 314 315#elif PARTIAL_STORE_M0 > 0 && PARTIAL_STORE_N0 == 0 316 317#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 318 STORE_BLOCK_PARTIAL_IN_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_COND_Y) 319 320#elif PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 > 0 321 322#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 323 STORE_BLOCK_PARTIAL_IN_X(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_N0, PARTIAL_COND_X) 324 325#else 326 327#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 328 STORE_BLOCK_PARTIAL_IN_X_AND_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) 329 330#endif 331 332#endif 333 334 335#if defined(PARTIAL_STORE_M0) 336 337#define COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) \ 338 ((uint)(max(0, (int)(y * M0) - (int)((M0 - PARTIAL_STORE_M0) % M0)))) 339#else 340#define COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) \ 341 ((uint)(y * M0)) 342#endif 343 344 345 346#define STORE_VECTOR_SELECT(basename, data_type, ptr, vec_size, leftover, cond) \ 347 STORE_BLOCK_PARTIAL_IN_X(1, vec_size, data_type, basename, ptr, 0, 0, leftover, cond) 348 349 350#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16) 351#pragma OPENCL EXTENSION cl_khr_fp16 : enable 352#endif 353 354#if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) 355#pragma OPENCL EXTENSION cl_arm_integer_dot_product_int8 : enable 356#endif 357 358#if defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8) 359#pragma OPENCL EXTENSION cl_arm_integer_dot_product_accumulate_int8 : enable 360#endif 361 362#if defined(ARM_COMPUTE_DEBUG_ENABLED) && defined(cl_arm_printf) 363#pragma OPENCL EXTENSION cl_arm_printf : enable 364#endif 365 366#define GPU_ARCH_MIDGARD 0x100 367#define GPU_ARCH_BIFROST 0x200 368#define GPU_ARCH_VALHALL 0x300 369 370 371#define CONCAT(a, b) a##b 372 373 374#define EXPAND(x) x 375 376 377#define CLAMP(x, min_val, max_val) min(max(x, min_val), max_val) 378 379 380#define REV1(x) ((x)) 381#define REV2(x) ((x).s10) 382#define REV3(x) ((x).s210) 383#define REV4(x) ((x).s3210) 384#define REV8(x) ((x).s76543210) 385#define REV16(x) ((x).sFEDCBA9876543210) 386 387 388 389#define REVERSE_STR(x, s) REV##s((x)) 390#define REVERSE(x, s) REVERSE_STR(x, s) 391 392 393 394#define ROT1_0(x) ((x)) 395#define ROT1_1(x) ((x)) 396 397#define ROT2_0(x) ((x)) 398#define ROT2_1(x) ((x).s10) 399#define ROT2_2(x) ((x)) 400 401#define ROT3_0(x) ((x)) 402#define ROT3_1(x) ((x).s201) 403#define ROT3_2(x) ((x).s120) 404#define ROT3_3(x) ((x)) 405 406#define ROT4_0(x) ((x)) 407#define ROT4_1(x) ((x).s3012) 408#define ROT4_2(x) ((x).s2301) 409#define ROT4_3(x) ((x).s1230) 410#define ROT4_4(x) ((x)) 411 412#define ROT8_0(x) ((x)) 413#define ROT8_1(x) ((x).s70123456) 414#define ROT8_2(x) ((x).s67012345) 415#define ROT8_3(x) ((x).s56701234) 416#define ROT8_4(x) ((x).s45670123) 417#define ROT8_5(x) ((x).s34567012) 418#define ROT8_6(x) ((x).s23456701) 419#define ROT8_7(x) ((x).s12345670) 420#define ROT8_8(x) ((x)) 421 422#define ROT16_0(x) ((x)) 423#define ROT16_1(x) ((x).sF0123456789ABCDE) 424#define ROT16_2(x) ((x).sEF0123456789ABCD) 425#define ROT16_3(x) ((x).sDEF0123456789ABC) 426#define ROT16_4(x) ((x).sCDEF0123456789AB) 427#define ROT16_5(x) ((x).sBCDEF0123456789A) 428#define ROT16_6(x) ((x).sABCDEF0123456789) 429#define ROT16_7(x) ((x).s9ABCDEF012345678) 430#define ROT16_8(x) ((x).s89ABCDEF01234567) 431#define ROT16_9(x) ((x).s789ABCDEF0123456) 432#define ROT16_10(x) ((x).s6789ABCDEF012345) 433#define ROT16_11(x) ((x).s56789ABCDEF01234) 434#define ROT16_12(x) ((x).s456789ABCDEF0123) 435#define ROT16_13(x) ((x).s3456789ABCDEF012) 436#define ROT16_14(x) ((x).s23456789ABCDEF01) 437#define ROT16_15(x) ((x).s123456789ABCDEF0) 438#define ROT16_16(x) ((x)) 439 440 441 442#define ROTATE_STR(x, s, n) ROT##s##_##n(x) 443#define ROTATE(x, s, n) ROTATE_STR(x, s, n) 444 445 446 447#define V_OFFS1(dt) (dt##1)(0) 448#define V_OFFS2(dt) (dt##2)(0, 1) 449#define V_OFFS3(dt) (dt##3)(0, 1, 2) 450#define V_OFFS4(dt) (dt##4)(0, 1, 2, 3) 451#define V_OFFS8(dt) (dt##8)(0, 1, 2, 3, 4, 5, 6, 7) 452#define V_OFFS16(dt) (dt##16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15) 453 454 455 456#define VEC_OFFS_STR(dt, s) V_OFFS##s(dt) 457#define VEC_OFFS(dt, s) VEC_OFFS_STR(dt, s) 458 459 460#define VLOAD_STR(size) vload##size 461#define VLOAD(size) VLOAD_STR(size) 462 463 464#define VLOAD_PARTIAL_STR(size, load_size) vload_partial_##size##_##load_size 465#define VLOAD_PARTIAL(size, load_size) VLOAD_PARTIAL_STR(size, load_size) 466 467#define NO_LOAD(data, offs, ptr) \ 468 { \ 469 } 470 471 472#define vload_partial_1_0 NO_LOAD 473#define vload_partial_1_1 vload1 474#define vload_partial_1_2 NO_LOAD 475#define vload_partial_1_3 NO_LOAD 476#define vload_partial_1_4 NO_LOAD 477#define vload_partial_1_5 NO_LOAD 478#define vload_partial_1_6 NO_LOAD 479#define vload_partial_1_7 NO_LOAD 480#define vload_partial_1_8 NO_LOAD 481#define vload_partial_1_9 NO_LOAD 482#define vload_partial_1_10 NO_LOAD 483#define vload_partial_1_11 NO_LOAD 484#define vload_partial_1_12 NO_LOAD 485#define vload_partial_1_13 NO_LOAD 486#define vload_partial_1_14 NO_LOAD 487#define vload_partial_1_15 NO_LOAD 488#define vload_partial_1_16 NO_LOAD 489 490#define vload_partial_2_0 NO_LOAD 491#define vload_partial_2_1 vload_partial_1 492#define vload_partial_2_2 vload_partial_2 493#define vload_partial_2_3 NO_LOAD 494#define vload_partial_2_4 NO_LOAD 495#define vload_partial_2_5 NO_LOAD 496#define vload_partial_2_6 NO_LOAD 497#define vload_partial_2_7 NO_LOAD 498#define vload_partial_2_8 NO_LOAD 499#define vload_partial_2_9 NO_LOAD 500#define vload_partial_2_10 NO_LOAD 501#define vload_partial_2_11 NO_LOAD 502#define vload_partial_2_12 NO_LOAD 503#define vload_partial_2_13 NO_LOAD 504#define vload_partial_2_14 NO_LOAD 505#define vload_partial_2_15 NO_LOAD 506#define vload_partial_2_16 NO_LOAD 507 508#define vload_partial_3_0 NO_LOAD 509#define vload_partial_3_1 vload_partial_1 510#define vload_partial_3_2 vload_partial_2 511#define vload_partial_3_3 vload_partial_3 512#define vload_partial_3_4 NO_LOAD 513#define vload_partial_3_5 NO_LOAD 514#define vload_partial_3_6 NO_LOAD 515#define vload_partial_3_7 NO_LOAD 516#define vload_partial_3_8 NO_LOAD 517#define vload_partial_3_9 NO_LOAD 518#define vload_partial_3_10 NO_LOAD 519#define vload_partial_3_11 NO_LOAD 520#define vload_partial_3_12 NO_LOAD 521#define vload_partial_3_13 NO_LOAD 522#define vload_partial_3_14 NO_LOAD 523#define vload_partial_3_15 NO_LOAD 524#define vload_partial_3_16 NO_LOAD 525 526#define vload_partial_4_0 NO_LOAD 527#define vload_partial_4_1 vload_partial_1 528#define vload_partial_4_2 vload_partial_2 529#define vload_partial_4_3 vload_partial_3 530#define vload_partial_4_4 vload_partial_4 531#define vload_partial_4_5 NO_LOAD 532#define vload_partial_4_6 NO_LOAD 533#define vload_partial_4_7 NO_LOAD 534#define vload_partial_4_8 NO_LOAD 535#define vload_partial_4_9 NO_LOAD 536#define vload_partial_4_10 NO_LOAD 537#define vload_partial_4_11 NO_LOAD 538#define vload_partial_4_12 NO_LOAD 539#define vload_partial_4_13 NO_LOAD 540#define vload_partial_4_14 NO_LOAD 541#define vload_partial_4_15 NO_LOAD 542#define vload_partial_4_16 NO_LOAD 543 544#define vload_partial_8_0 NO_LOAD 545#define vload_partial_8_1 vload_partial_1 546#define vload_partial_8_2 vload_partial_2 547#define vload_partial_8_3 vload_partial_3 548#define vload_partial_8_4 vload_partial_4 549#define vload_partial_8_5 vload_partial_5 550#define vload_partial_8_6 vload_partial_6 551#define vload_partial_8_7 vload_partial_7 552#define vload_partial_8_8 vload_partial_8 553#define vload_partial_8_9 NO_LOAD 554#define vload_partial_8_10 NO_LOAD 555#define vload_partial_8_11 NO_LOAD 556#define vload_partial_8_12 NO_LOAD 557#define vload_partial_8_13 NO_LOAD 558#define vload_partial_8_14 NO_LOAD 559#define vload_partial_8_15 NO_LOAD 560#define vload_partial_8_16 NO_LOAD 561 562#define vload_partial_16_0 NO_LOAD 563#define vload_partial_16_1 vload_partial_1 564#define vload_partial_16_2 vload_partial_2 565#define vload_partial_16_3 vload_partial_3 566#define vload_partial_16_4 vload_partial_4 567#define vload_partial_16_5 vload_partial_5 568#define vload_partial_16_6 vload_partial_6 569#define vload_partial_16_7 vload_partial_7 570#define vload_partial_16_8 vload_partial_8 571#define vload_partial_16_9 vload_partial_9 572#define vload_partial_16_10 vload_partial_10 573#define vload_partial_16_11 vload_partial_11 574#define vload_partial_16_12 vload_partial_12 575#define vload_partial_16_13 vload_partial_13 576#define vload_partial_16_14 vload_partial_14 577#define vload_partial_16_15 vload_partial_15 578#define vload_partial_16_16 vload_partial_16 579 580 581#define vload_partial_1(DATA, OFFSET, PTR) \ 582 DATA.s0 = vload1(OFFSET, PTR); 583 584#define vload_partial_2(DATA, OFFSET, PTR) \ 585 DATA.s01 = vload2(OFFSET, PTR); 586 587#define vload_partial_3(DATA, OFFSET, PTR) \ 588 DATA.s012 = vload3(OFFSET, PTR); 589 590#define vload_partial_4(DATA, OFFSET, PTR) \ 591 DATA.s0123 = vload4(OFFSET, PTR); 592 593#define vload_partial_5(DATA, OFFSET, PTR) \ 594 vload_partial_4(DATA.s0123, OFFSET, PTR); \ 595 DATA.s4 = vload1(OFFSET, PTR + 4); 596 597#define vload_partial_6(DATA, OFFSET, PTR) \ 598 vload_partial_4(DATA.s0123, OFFSET, PTR); \ 599 vload_partial_2(DATA.s45, OFFSET, PTR + 4); 600 601#define vload_partial_7(DATA, OFFSET, PTR) \ 602 vload_partial_4(DATA.s0123, OFFSET, PTR); \ 603 vload_partial_3(DATA.s456, OFFSET, PTR + 4); 604 605#define vload_partial_8(DATA, OFFSET, PTR) \ 606 DATA.s01234567 = vload8(OFFSET, PTR); 607 608#define vload_partial_9(DATA, OFFSET, PTR) \ 609 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 610 DATA.s8 = vload1(OFFSET, PTR + 8); 611 612#define vload_partial_10(DATA, OFFSET, PTR) \ 613 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 614 vload_partial_2(DATA.s89, OFFSET, PTR + 8); 615 616#define vload_partial_11(DATA, OFFSET, PTR) \ 617 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 618 vload_partial_3(DATA.s89A, OFFSET, PTR + 8); 619 620#define vload_partial_12(DATA, OFFSET, PTR) \ 621 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 622 vload_partial_4(DATA.s89AB, OFFSET, PTR + 8); 623 624#define vload_partial_13(DATA, OFFSET, PTR) \ 625 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 626 vload_partial_5(DATA.s89ABCDEF, OFFSET, PTR + 8); 627 628#define vload_partial_14(DATA, OFFSET, PTR) \ 629 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 630 vload_partial_6(DATA.s89ABCDEF, OFFSET, PTR + 8); 631 632#define vload_partial_15(DATA, OFFSET, PTR) \ 633 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 634 vload_partial_7(DATA.s89ABCDEF, OFFSET, PTR + 8); 635 636#define vload_partial_16(DATA, OFFSET, PTR) \ 637 DATA = vload16(OFFSET, PTR); 638 639 640 641#define PIXEL_UNIT4 1 642#define PIXEL_UNIT8 2 643#define PIXEL_UNIT16 4 644 645 646#define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size) PIXEL_UNIT##vec_size 647#define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(vec_size) CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size) 648 649 650#define read_image2d_floatx1(img, x_coord, y_coord) (float4)(read_imagef(img, (int2)(x_coord, y_coord))); 651#define read_image2d_floatx2(img, x_coord, y_coord) (float8)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord))); 652#define read_image2d_floatx4(img, x_coord, y_coord) (float16)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord)), read_imagef(img, (int2)(x_coord + 2, y_coord)), read_imagef(img, (int2)(x_coord + 3, y_coord))); 653 654#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16) 655#define read_image2d_halfx1(img, x_coord, y_coord) (half4)(read_imageh(img, (int2)(x_coord, y_coord))); 656#define read_image2d_halfx2(img, x_coord, y_coord) (half8)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord))); 657#define read_image2d_halfx4(img, x_coord, y_coord) (half16)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord)), read_imageh(img, (int2)(x_coord + 2, y_coord)), read_imageh(img, (int2)(x_coord + 3, y_coord))); 658#endif 659 660#define write_image2d_floatx1(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values)); 661#define write_image2d_floatx2(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values.s0123), write_imagef(img, (int2)(x_coord + 1, y_coord), values.s4567)); 662#define write_image2d_floatx4(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values.s0123), write_imagef(img, (int2)(x_coord + 1, y_coord), values.s4567), write_imagef(img, (int2)(x_coord + 2, y_coord), values.s89AB), write_imagef(img, (int2)(x_coord + 3, y_coord), values.sCDEF)); 663 664#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16) 665#define write_image2d_halfx1(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values)); 666#define write_image2d_halfx2(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values.s0123), write_imageh(img, (int2)(x_coord + 1, y_coord), values.s4567)); 667#define write_image2d_halfx4(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values.s0123), write_imageh(img, (int2)(x_coord + 1, y_coord), values.s4567), write_imageh(img, (int2)(x_coord + 2, y_coord), values.s89AB), write_imageh(img, (int2)(x_coord + 3, y_coord), values.sCDEF)); 668#endif 669 670 671#define READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord) read_image2d_##data_type##x##n0(img, x_coord, y_coord) 672#define READ_IMAGE2D(data_type, n0, img, x_coord, y_coord) READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord) 673 674 675#define WRITE_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord, values) write_image2d_##data_type##x##n0(img, x_coord, y_coord, values) 676#define WRITE_IMAGE2D(data_type, n0, img, x_coord, y_coord, values) WRITE_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord, values) 677 678#define VSTORE_STR(size) vstore##size 679#define VSTORE(size) VSTORE_STR(size) 680 681#define float1 float 682#define half1 half 683#define char1 char 684#define uchar1 uchar 685#define short1 short 686#define ushort1 ushort 687#define int1 int 688#define uint1 uint 689#define long1 long 690#define ulong1 ulong 691#define double1 double 692 693#define vload1(OFFSET, PTR) *(OFFSET + PTR) 694#define vstore1(DATA, OFFSET, PTR) *(OFFSET + PTR) = DATA 695 696 697#define VSTORE_PARTIAL_STR(size, store_size) vstore_partial_##size##_##store_size 698#define VSTORE_PARTIAL(size, store_size) VSTORE_PARTIAL_STR(size, store_size) 699 700#define NO_STORE(data, offs, ptr) \ 701 { \ 702 } 703 704 705#define vstore_partial_1_0 NO_STORE 706#define vstore_partial_1_1 vstore1 707#define vstore_partial_1_2 NO_STORE 708#define vstore_partial_1_3 NO_STORE 709#define vstore_partial_1_4 NO_STORE 710#define vstore_partial_1_5 NO_STORE 711#define vstore_partial_1_6 NO_STORE 712#define vstore_partial_1_7 NO_STORE 713#define vstore_partial_1_8 NO_STORE 714#define vstore_partial_1_9 NO_STORE 715#define vstore_partial_1_10 NO_STORE 716#define vstore_partial_1_11 NO_STORE 717#define vstore_partial_1_12 NO_STORE 718#define vstore_partial_1_13 NO_STORE 719#define vstore_partial_1_14 NO_STORE 720#define vstore_partial_1_15 NO_STORE 721#define vstore_partial_1_16 NO_STORE 722 723#define vstore_partial_2_0 NO_STORE 724#define vstore_partial_2_1 vstore_partial_1 725#define vstore_partial_2_2 vstore_partial_2 726#define vstore_partial_2_3 NO_STORE 727#define vstore_partial_2_4 NO_STORE 728#define vstore_partial_2_5 NO_STORE 729#define vstore_partial_2_6 NO_STORE 730#define vstore_partial_2_7 NO_STORE 731#define vstore_partial_2_8 NO_STORE 732#define vstore_partial_2_9 NO_STORE 733#define vstore_partial_2_10 NO_STORE 734#define vstore_partial_2_11 NO_STORE 735#define vstore_partial_2_12 NO_STORE 736#define vstore_partial_2_13 NO_STORE 737#define vstore_partial_2_14 NO_STORE 738#define vstore_partial_2_15 NO_STORE 739#define vstore_partial_2_16 NO_STORE 740 741#define vstore_partial_3_0 NO_STORE 742#define vstore_partial_3_1 vstore_partial_1 743#define vstore_partial_3_2 vstore_partial_2 744#define vstore_partial_3_3 vstore_partial_3 745#define vstore_partial_3_4 NO_STORE 746#define vstore_partial_3_5 NO_STORE 747#define vstore_partial_3_6 NO_STORE 748#define vstore_partial_3_7 NO_STORE 749#define vstore_partial_3_8 NO_STORE 750#define vstore_partial_3_9 NO_STORE 751#define vstore_partial_3_10 NO_STORE 752#define vstore_partial_3_11 NO_STORE 753#define vstore_partial_3_12 NO_STORE 754#define vstore_partial_3_13 NO_STORE 755#define vstore_partial_3_14 NO_STORE 756#define vstore_partial_3_15 NO_STORE 757#define vstore_partial_3_16 NO_STORE 758 759#define vstore_partial_4_0 NO_STORE 760#define vstore_partial_4_1 vstore_partial_1 761#define vstore_partial_4_2 vstore_partial_2 762#define vstore_partial_4_3 vstore_partial_3 763#define vstore_partial_4_4 vstore_partial_4 764#define vstore_partial_4_5 NO_STORE 765#define vstore_partial_4_6 NO_STORE 766#define vstore_partial_4_7 NO_STORE 767#define vstore_partial_4_8 NO_STORE 768#define vstore_partial_4_9 NO_STORE 769#define vstore_partial_4_10 NO_STORE 770#define vstore_partial_4_11 NO_STORE 771#define vstore_partial_4_12 NO_STORE 772#define vstore_partial_4_13 NO_STORE 773#define vstore_partial_4_14 NO_STORE 774#define vstore_partial_4_15 NO_STORE 775#define vstore_partial_4_16 NO_STORE 776 777#define vstore_partial_8_0 NO_STORE 778#define vstore_partial_8_1 vstore_partial_1 779#define vstore_partial_8_2 vstore_partial_2 780#define vstore_partial_8_3 vstore_partial_3 781#define vstore_partial_8_4 vstore_partial_4 782#define vstore_partial_8_5 vstore_partial_5 783#define vstore_partial_8_6 vstore_partial_6 784#define vstore_partial_8_7 vstore_partial_7 785#define vstore_partial_8_8 vstore_partial_8 786#define vstore_partial_8_9 NO_STORE 787#define vstore_partial_8_10 NO_STORE 788#define vstore_partial_8_11 NO_STORE 789#define vstore_partial_8_12 NO_STORE 790#define vstore_partial_8_13 NO_STORE 791#define vstore_partial_8_14 NO_STORE 792#define vstore_partial_8_15 NO_STORE 793#define vstore_partial_8_16 NO_STORE 794 795#define vstore_partial_16_0 NO_STORE 796#define vstore_partial_16_1 vstore_partial_1 797#define vstore_partial_16_2 vstore_partial_2 798#define vstore_partial_16_3 vstore_partial_3 799#define vstore_partial_16_4 vstore_partial_4 800#define vstore_partial_16_5 vstore_partial_5 801#define vstore_partial_16_6 vstore_partial_6 802#define vstore_partial_16_7 vstore_partial_7 803#define vstore_partial_16_8 vstore_partial_8 804#define vstore_partial_16_9 vstore_partial_9 805#define vstore_partial_16_10 vstore_partial_10 806#define vstore_partial_16_11 vstore_partial_11 807#define vstore_partial_16_12 vstore_partial_12 808#define vstore_partial_16_13 vstore_partial_13 809#define vstore_partial_16_14 vstore_partial_14 810#define vstore_partial_16_15 vstore_partial_15 811#define vstore_partial_16_16 vstore_partial_16 812 813 814#define vstore_partial_1(DATA, OFFSET, PTR) \ 815 vstore1(DATA.s0, OFFSET, PTR); 816 817#define vstore_partial_2(DATA, OFFSET, PTR) \ 818 vstore2(DATA.s01, OFFSET, PTR); 819 820#define vstore_partial_3(DATA, OFFSET, PTR) \ 821 vstore3(DATA.s012, OFFSET, PTR); 822 823#define vstore_partial_4(DATA, OFFSET, PTR) \ 824 vstore4(DATA.s0123, OFFSET, PTR); 825 826#define vstore_partial_5(DATA, OFFSET, PTR) \ 827 vstore_partial_4(DATA.s0123, OFFSET, PTR); \ 828 vstore1(DATA.s4, OFFSET, PTR + 4); 829 830#define vstore_partial_6(DATA, OFFSET, PTR) \ 831 vstore_partial_4(DATA.s0123, OFFSET, PTR); \ 832 vstore_partial_2(DATA.s45, OFFSET, PTR + 4); 833 834#define vstore_partial_7(DATA, OFFSET, PTR) \ 835 vstore_partial_4(DATA.s0123, OFFSET, PTR); \ 836 vstore_partial_3(DATA.s456, OFFSET, PTR + 4); 837 838#define vstore_partial_8(DATA, OFFSET, PTR) \ 839 vstore8(DATA.s01234567, OFFSET, PTR); 840 841#define vstore_partial_9(DATA, OFFSET, PTR) \ 842 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 843 vstore1(DATA.s8, OFFSET, PTR + 8); 844 845#define vstore_partial_10(DATA, OFFSET, PTR) \ 846 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 847 vstore_partial_2(DATA.s89, OFFSET, PTR + 8); 848 849#define vstore_partial_11(DATA, OFFSET, PTR) \ 850 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 851 vstore_partial_3(DATA.s89a, OFFSET, PTR + 8); 852 853#define vstore_partial_12(DATA, OFFSET, PTR) \ 854 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 855 vstore_partial_4(DATA.s89ab, OFFSET, PTR + 8); 856 857#define vstore_partial_13(DATA, OFFSET, PTR) \ 858 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 859 vstore_partial_5(DATA.s89abcdef, OFFSET, PTR + 8); 860 861#define vstore_partial_14(DATA, OFFSET, PTR) \ 862 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 863 vstore_partial_6(DATA.s89abcdef, OFFSET, PTR + 8); 864 865#define vstore_partial_15(DATA, OFFSET, PTR) \ 866 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 867 vstore_partial_7(DATA.s89abcdef, OFFSET, PTR + 8); 868 869#define vstore_partial_16(DATA, OFFSET, PTR) \ 870 vstore16(DATA, OFFSET, PTR); 871 872 873 874 875 876#define convert_float_sat convert_float 877#define convert_float1_sat convert_float 878#define convert_float2_sat convert_float2 879#define convert_float3_sat convert_float3 880#define convert_float4_sat convert_float4 881#define convert_float8_sat convert_float8 882#define convert_float16_sat convert_float16 883#define convert_half_sat convert_float 884#define convert_half1_sat convert_half 885#define convert_half2_sat convert_half2 886#define convert_half3_sat convert_half3 887#define convert_half4_sat convert_half4 888#define convert_half8_sat convert_half8 889#define convert_half16_sat convert_half16 890 891#define convert_float1 convert_float 892#define convert_half1 convert_half 893#define convert_char1 convert_char 894#define convert_uchar1 convert_uchar 895#define convert_short1 convert_short 896#define convert_ushort1 convert_ushort 897#define convert_int1 convert_int 898#define convert_uint1 convert_uint 899#define convert_long1 convert_long 900#define convert_ulong1 convert_ulong 901#define convert_double1 convert_double 902 903#define convert_char1_sat convert_char_sat 904#define convert_uchar1_sat convert_uchar_sat 905#define convert_uchar2_sat convert_uchar2_sat 906#define convert_uchar3_sat convert_uchar3_sat 907#define convert_uchar4_sat convert_uchar4_sat 908#define convert_uchar8_sat convert_uchar8_sat 909#define convert_uchar16_sat convert_uchar16_sat 910#define convert_short1_sat convert_short_sat 911#define convert_ushort1_sat convert_ushort_sat 912#define convert_int1_sat convert_int_sat 913#define convert_uint1_sat convert_uint_sat 914#define convert_long1_sat convert_long_sat 915#define convert_ulong1_sat convert_ulong_sat 916#define convert_double1_sat convert_double_sat 917 918#define VEC_DATA_TYPE_STR(type, size) type##size 919#define VEC_DATA_TYPE(type, size) VEC_DATA_TYPE_STR(type, size) 920 921#define CONVERT_STR(x, type) (convert_##type((x))) 922#define CONVERT(x, type) CONVERT_STR(x, type) 923 924#define CONVERT_SAT_STR(x, type) (convert_##type##_sat((x))) 925#define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type) 926 927#define CONVERT_SAT_ROUND_STR(x, type, round) (convert_##type##_sat_##round((x))) 928#define CONVERT_SAT_ROUND(x, type, round) CONVERT_SAT_ROUND_STR(x, type, round) 929 930#define select_vec_dt_uchar(size) uchar##size 931#define select_vec_dt_char(size) char##size 932#define select_vec_dt_ushort(size) ushort##size 933#define select_vec_dt_short(size) short##size 934#define select_vec_dt_half(size) short##size 935#define select_vec_dt_uint(size) uint##size 936#define select_vec_dt_int(size) int##size 937#define select_vec_dt_float(size) int##size 938#define select_vec_dt_ulong(size) ulong##size 939#define select_vec_dt_long(size) long##size 940 941#define SELECT_VEC_DATA_TYPE_STR(type, size) select_vec_dt_##type(size) 942#define SELECT_VEC_DATA_TYPE(type, size) SELECT_VEC_DATA_TYPE_STR(type, size) 943#define SELECT_DATA_TYPE(type) SELECT_VEC_DATA_TYPE_STR(type, 1) 944 945#define signed_int_vec_dt_uchar(size) char##size 946#define signed_int_vec_dt_char(size) char##size 947#define signed_int_vec_dt_ushort(size) short##size 948#define signed_int_vec_dt_short(size) short##size 949#define signed_int_vec_dt_half(size) short##size 950#define signed_int_vec_dt_uint(size) int##size 951#define signed_int_vec_dt_int(size) int##size 952#define signed_int_vec_dt_float(size) int##size 953#define signed_int_vec_dt_ulong(size) long##size 954#define signed_int_vec_dt_long(size) long##size 955 956#define SIGNED_INT_VEC_DATA_TYPE_STR(type, size) signed_int_vec_dt_##type(size) 957#define SIGNED_INT_VEC_DATA_TYPE(type, size) SIGNED_INT_VEC_DATA_TYPE_STR(type, size) 958#define SIGNED_INT_DATA_TYPE(type) SIGNED_INT_VEC_DATA_TYPE_STR(type, 1) 959 960#define sum_reduce_1(x) (x) 961#define sum_reduce_2(x) ((x).s0) + ((x).s1) 962#define sum_reduce_3(x) sum_reduce_2((x).s01) + ((x).s2) 963#define sum_reduce_4(x) sum_reduce_2((x).s01) + sum_reduce_2((x).s23) 964#define sum_reduce_8(x) sum_reduce_4((x).s0123) + sum_reduce_4((x).s4567) 965#define sum_reduce_16(x) sum_reduce_8((x).s01234567) + sum_reduce_8((x).s89ABCDEF) 966 967#define SUM_REDUCE_STR(x, size) sum_reduce_##size(x) 968#define SUM_REDUCE(x, size) SUM_REDUCE_STR(x, size) 969 970#define prod_reduce_1(x) (x) 971#define prod_reduce_2(x) ((x).s0) * ((x).s1) 972#define prod_reduce_3(x) prod_reduce_2((x).s01) * ((x).s2) 973#define prod_reduce_4(x) prod_reduce_2((x).s01) * prod_reduce_2((x).s23) 974#define prod_reduce_8(x) prod_reduce_4((x).s0123) * prod_reduce_4((x).s4567) 975#define prod_reduce_16(x) prod_reduce_8((x).s01234567) * prod_reduce_8((x).s89ABCDEF) 976 977#define PROD_REDUCE_STR(x, size) prod_reduce_##size(x) 978#define PROD_REDUCE(x, size) PROD_REDUCE_STR(x, size) 979 980#define max_reduce_1(x) (x) 981#define max_reduce_2(x) max(((x).s0), ((x).s1)) 982#define max_reduce_3(x) max(max_reduce_2((x).s01), ((x).s2)) 983#define max_reduce_4(x) max(max_reduce_2((x).s01), max_reduce_2((x).s23)) 984#define max_reduce_8(x) max(max_reduce_4((x).s0123), max_reduce_4((x).s4567)) 985#define max_reduce_16(x) max(max_reduce_8((x).s01234567), max_reduce_8((x).s89ABCDEF)) 986 987#define MAX_REDUCE_STR(x, size) max_reduce_##size(x) 988#define MAX_REDUCE(x, size) MAX_REDUCE_STR(x, size) 989 990#define VECTOR_DECLARATION(name) \ 991 __global uchar *name##_ptr, \ 992 uint name##_stride_x, \ 993 uint name##_step_x, \ 994 uint name##_offset_first_element_in_bytes 995 996#define IMAGE_DECLARATION(name) \ 997 __global uchar *name##_ptr, \ 998 uint name##_stride_x, \ 999 uint name##_step_x, \ 1000 uint name##_stride_y, \ 1001 uint name##_step_y, \ 1002 uint name##_offset_first_element_in_bytes 1003 1004#define TENSOR3D_DECLARATION(name) \ 1005 __global uchar *name##_ptr, \ 1006 uint name##_stride_x, \ 1007 uint name##_step_x, \ 1008 uint name##_stride_y, \ 1009 uint name##_step_y, \ 1010 uint name##_stride_z, \ 1011 uint name##_step_z, \ 1012 uint name##_offset_first_element_in_bytes 1013 1014#define TENSOR4D_DECLARATION(name) \ 1015 __global uchar *name##_ptr, \ 1016 uint name##_stride_x, \ 1017 uint name##_step_x, \ 1018 uint name##_stride_y, \ 1019 uint name##_step_y, \ 1020 uint name##_stride_z, \ 1021 uint name##_step_z, \ 1022 uint name##_stride_w, \ 1023 uint name##_step_w, \ 1024 uint name##_offset_first_element_in_bytes 1025 1026#define TENSOR5D_DECLARATION(name) \ 1027 __global uchar *name##_ptr, \ 1028 uint name##_stride_x, \ 1029 uint name##_step_x, \ 1030 uint name##_stride_y, \ 1031 uint name##_step_y, \ 1032 uint name##_stride_z, \ 1033 uint name##_step_z, \ 1034 uint name##_stride_w, \ 1035 uint name##_step_w, \ 1036 uint name##_stride_v, \ 1037 uint name##_step_v, \ 1038 uint name##_offset_first_element_in_bytes 1039 1040#define CONVERT_TO_VECTOR_STRUCT(name) \ 1041 update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x) 1042 1043#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name) \ 1044 update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0) 1045 1046#define CONVERT_TO_IMAGE_STRUCT(name) \ 1047 update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y) 1048 1049#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name) \ 1050 update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0) 1051 1052#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \ 1053 update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, name##_step_z) 1054 1055#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name) \ 1056 update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, name##_step_z) 1057 1058#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \ 1059 update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, name##_step_z) 1060 1061#define CONVERT_TO_TENSOR3D_STRUCT(name) \ 1062 update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \ 1063 name##_stride_z, name##_step_z) 1064 1065#define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name) \ 1066 update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0) 1067 1068#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size) \ 1069 update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \ 1070 name##_stride_z, name##_step_z, name##_stride_w, name##_step_w, mod_size) 1071 1072#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size) \ 1073 update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0, name##_stride_w, 0, mod_size) 1074 1075#define CONVERT_TO_TENSOR3D_STRUCT_NO_UPDATE_PTR(name) \ 1076 tensor3D_ptr_no_update(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \ 1077 name##_stride_z, name##_step_z) 1078 1079 1080typedef struct Vector 1081{ 1082 __global uchar *ptr; 1083 int offset_first_element_in_bytes; 1084 int stride_x; 1085} Vector; 1086 1087 1088typedef struct Image 1089{ 1090 __global uchar *ptr; 1091 int offset_first_element_in_bytes; 1092 int stride_x; 1093 int stride_y; 1094} Image; 1095 1096 1097typedef struct Tensor3D 1098{ 1099 __global uchar *ptr; 1100 int offset_first_element_in_bytes; 1101 int stride_x; 1102 int stride_y; 1103 int stride_z; 1104} Tensor3D; 1105 1106 1107typedef struct Tensor4D 1108{ 1109 __global uchar *ptr; 1110 int offset_first_element_in_bytes; 1111 int stride_x; 1112 int stride_y; 1113 int stride_z; 1114 int stride_w; 1115} Tensor4D; 1116 1117 1118inline Vector update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x) 1119{ 1120 Vector vector = 1121 { 1122 .ptr = ptr, 1123 .offset_first_element_in_bytes = offset_first_element_in_bytes, 1124 .stride_x = stride_x, 1125 }; 1126 vector.ptr += vector.offset_first_element_in_bytes + get_global_id(0) * step_x; 1127 return vector; 1128} 1129 1130 1131inline Image update_image_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y) 1132{ 1133 Image img = 1134 { 1135 .ptr = ptr, 1136 .offset_first_element_in_bytes = offset_first_element_in_bytes, 1137 .stride_x = stride_x, 1138 .stride_y = stride_y 1139 }; 1140 img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y; 1141 return img; 1142} 1143 1144 1145inline Image update_image_from_tensor3D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z) 1146{ 1147 Image img = 1148 { 1149 .ptr = ptr, 1150 .offset_first_element_in_bytes = offset_first_element_in_bytes, 1151 .stride_x = stride_x, 1152 .stride_y = stride_y 1153 }; 1154 img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z; 1155 return img; 1156} 1157 1158 1159inline Tensor3D update_tensor3D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z) 1160{ 1161 Tensor3D tensor = 1162 { 1163 .ptr = ptr, 1164 .offset_first_element_in_bytes = offset_first_element_in_bytes, 1165 .stride_x = stride_x, 1166 .stride_y = stride_y, 1167 .stride_z = stride_z 1168 }; 1169 tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z; 1170 return tensor; 1171} 1172 1173 1174inline Tensor3D tensor3D_ptr_no_update(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z) 1175{ 1176 Tensor3D tensor = 1177 { 1178 .ptr = ptr, 1179 .offset_first_element_in_bytes = offset_first_element_in_bytes, 1180 .stride_x = stride_x, 1181 .stride_y = stride_y, 1182 .stride_z = stride_z 1183 }; 1184 return tensor; 1185} 1186 1187inline Tensor4D update_tensor4D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z, uint stride_w, 1188 uint step_w, 1189 uint mod_size) 1190{ 1191 Tensor4D tensor = 1192 { 1193 .ptr = ptr, 1194 .offset_first_element_in_bytes = offset_first_element_in_bytes, 1195 .stride_x = stride_x, 1196 .stride_y = stride_y, 1197 .stride_z = stride_z, 1198 .stride_w = stride_w 1199 }; 1200 1201 tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + (get_global_id(2) % mod_size) * step_z + (get_global_id(2) / mod_size) * step_w; 1202 return tensor; 1203} 1204 1205 1206inline __global const uchar *vector_offset(const Vector *vec, int x) 1207{ 1208 return vec->ptr + x * vec->stride_x; 1209} 1210 1211 1212inline __global uchar *offset(const Image *img, int x, int y) 1213{ 1214 return img->ptr + x * img->stride_x + y * img->stride_y; 1215} 1216 1217 1218inline __global const uchar *tensor3D_offset(const Tensor3D *tensor, int x, int y, int z) 1219{ 1220 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z; 1221} 1222 1223 1224inline __global const uchar *tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w) 1225{ 1226 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + w * tensor->stride_w; 1227} 1228 1229 1230inline __global const uchar *tensor3D_index2ptr(const Tensor3D *tensor, uint width, uint height, uint depth, uint index) 1231{ 1232 uint num_elements = width * height; 1233 1234 const uint z = index / num_elements; 1235 1236 index %= num_elements; 1237 1238 const uint y = index / width; 1239 1240 index %= width; 1241 1242 const uint x = index; 1243 1244 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + tensor->offset_first_element_in_bytes; 1245} 1246 1247#endif 1248 1249#if GPU_ARCH == GPU_ARCH_BIFROST 1250#define MLA(a, b, c) (fma(c, b, a)) 1251#else 1252#define MLA(a, b, c) ((b) * (c) + (a)) 1253#endif 1254 1255 1256#define hard_swish_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (x * ((min(max((x + (DATA_TYPE)3.0), (DATA_TYPE)0.0), (DATA_TYPE)6.0)) * (DATA_TYPE)0.166666667)) 1257 1258 1259#define logistic_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) ((DATA_TYPE)1.0 / ((DATA_TYPE)1.0 + exp(-x))) 1260 1261 1262#define tanh_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) ((DATA_TYPE)A_VAL * tanh((DATA_TYPE)B_VAL * x)) 1263 1264 1265#define relu_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (max((DATA_TYPE)0.0, x)) 1266 1267 1268#define brelu_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (min((DATA_TYPE)A_VAL, max((DATA_TYPE)0.0, x))) 1269 1270 1271#define lu_brelu_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (min(max(x, (DATA_TYPE)B_VAL), (DATA_TYPE)A_VAL)) 1272 1273 1274#define lrelu_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) ((min(x, (DATA_TYPE)0.0) * (DATA_TYPE)A_VAL) + max(x, (DATA_TYPE)0.0)) 1275 1276 1277#define srelu_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (log((DATA_TYPE)1.0 + exp(x))) 1278 1279 1280#define elu_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (select(((DATA_TYPE)A_VAL * (exp(x) - (DATA_TYPE)1.0)), x, (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))isgreaterequal(x, (DATA_TYPE)0.0))) 1281 1282 1283#define abs_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (fabs(x)) 1284 1285 1286#define square_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (x * x) 1287 1288 1289#define sqrt_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (sqrt(x)) 1290 1291 1292#define linear_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (MLA((DATA_TYPE)B_VAL, (DATA_TYPE)A_VAL, x)) 1293 1294 1295#define gelu_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (x * (DATA_TYPE)0.5 * ((DATA_TYPE)1.0 + erf(x / (DATA_TYPE)1.41421356237))) 1296 1297 1298#define identity_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) (x) 1299 1300#define ACT_OP(op, DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) op##_op(DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) 1301 1302#define ACTIVATION(op, DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) ACT_OP(op, DATA_TYPE, VEC_SIZE, x, A_VAL, B_VAL) 1303 1304#ifndef ARM_COMPUTE_HELPER_H 1305#define ARM_COMPUTE_HELPER_H 1306 1307 1308 1309 1310#define STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1311 VSTORE(N0) \ 1312 (BASENAME##0, 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0)); 1313 1314#define STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1315 STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1316 VSTORE(N0) \ 1317 (BASENAME##1, 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1)); 1318 1319#define STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1320 STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1321 VSTORE(N0) \ 1322 (BASENAME##2, 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2)); 1323 1324#define STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1325 STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1326 VSTORE(N0) \ 1327 (BASENAME##3, 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3)); 1328 1329#define STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1330 STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1331 VSTORE(N0) \ 1332 (BASENAME##4, 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4)); 1333 1334#define STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1335 STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1336 VSTORE(N0) \ 1337 (BASENAME##5, 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5)); 1338 1339#define STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1340 STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1341 VSTORE(N0) \ 1342 (BASENAME##6, 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6)); 1343 1344#define STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1345 STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1346 VSTORE(N0) \ 1347 (BASENAME##7, 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7)); 1348 1349#define STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1350 STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1351 VSTORE(N0) \ 1352 (BASENAME##8, 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8)); 1353 1354#define STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1355 STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1356 VSTORE(N0) \ 1357 (BASENAME##9, 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9)); 1358 1359#define STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1360 STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1361 VSTORE(N0) \ 1362 (BASENAME##A, 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A)); 1363 1364#define STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1365 STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1366 VSTORE(N0) \ 1367 (BASENAME##B, 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B)); 1368 1369#define STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1370 STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1371 VSTORE(N0) \ 1372 (BASENAME##C, 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C)); 1373 1374#define STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1375 STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1376 VSTORE(N0) \ 1377 (BASENAME##D, 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D)); 1378 1379#define STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1380 STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1381 VSTORE(N0) \ 1382 (BASENAME##E, 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E)); 1383 1384#define STORE_ROW_16(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1385 STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1386 VSTORE(N0) \ 1387 (BASENAME##F, 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F)); 1388 1389 1390 1391#define CONVERT_STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1392 VSTORE(N0) \ 1393 (CONVERT_SAT((BASENAME##0), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0)); 1394 1395#define CONVERT_STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1396 CONVERT_STORE_ROW_1(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1397 VSTORE(N0) \ 1398 (CONVERT_SAT((BASENAME##1), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1)); 1399 1400#define CONVERT_STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1401 CONVERT_STORE_ROW_2(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1402 VSTORE(N0) \ 1403 (CONVERT_SAT((BASENAME##2), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2)); 1404 1405#define CONVERT_STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1406 CONVERT_STORE_ROW_3(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1407 VSTORE(N0) \ 1408 (CONVERT_SAT((BASENAME##3), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3)); 1409 1410#define CONVERT_STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1411 CONVERT_STORE_ROW_4(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1412 VSTORE(N0) \ 1413 (CONVERT_SAT((BASENAME##4), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4)); 1414 1415#define CONVERT_STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1416 CONVERT_STORE_ROW_5(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1417 VSTORE(N0) \ 1418 (CONVERT_SAT((BASENAME##5), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5)); 1419 1420#define CONVERT_STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1421 CONVERT_STORE_ROW_6(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1422 VSTORE(N0) \ 1423 (CONVERT_SAT((BASENAME##6), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6)); 1424 1425#define CONVERT_STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1426 CONVERT_STORE_ROW_7(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1427 VSTORE(N0) \ 1428 (CONVERT_SAT((BASENAME##7), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7)); 1429 1430#define CONVERT_STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1431 CONVERT_STORE_ROW_8(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1432 VSTORE(N0) \ 1433 (CONVERT_SAT((BASENAME##8), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8)); 1434 1435#define CONVERT_STORE_ROW_10(N0, DATA, BASENAME, PTR, STRIDE_Y, Z) \ 1436 CONVERT_STORE_ROW_9(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1437 VSTORE(N0) \ 1438 (CONVERT_SAT((BASENAME##9), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9)); 1439 1440#define CONVERT_STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1441 CONVERT_STORE_ROW_10(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1442 VSTORE(N0) \ 1443 (CONVERT_SAT((BASENAME##A), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A)); 1444 1445#define CONVERT_STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1446 CONVERT_STORE_ROW_11(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1447 VSTORE(N0) \ 1448 (CONVERT_SAT((BASENAME##B), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B)); 1449 1450#define CONVERT_STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1451 CONVERT_STORE_ROW_12(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1452 VSTORE(N0) \ 1453 (CONVERT_SAT((BASENAME##C), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C)); 1454 1455#define CONVERT_STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1456 CONVERT_STORE_ROW_13(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1457 VSTORE(N0) \ 1458 (CONVERT_SAT((BASENAME##D), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D)); 1459 1460#define CONVERT_STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1461 CONVERT_STORE_ROW_14(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1462 VSTORE(N0) \ 1463 (CONVERT_SAT((BASENAME##E), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E)); 1464 1465#define CONVERT_STORE_ROW_16(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1466 CONVERT_STORE_ROW_15(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1467 VSTORE(N0) \ 1468 (CONVERT_SAT((BASENAME##F), VEC_DATA_TYPE(DATA_TYPE, N0)), 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F)); 1469 1470 1471 1472 1473#define STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_ROW_##M0(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 1474#define STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 1475 1476 1477 1478#define CONVERT_STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) CONVERT_STORE_ROW_##M0(N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 1479#define CONVERT_STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) CONVERT_STORE_BLOCK_STR(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 1480 1481 1482 1483#define STORE_ROW_PARTIAL_1(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1484 VSTORE_PARTIAL(N0, STORE_N0) \ 1485 (BASENAME##0, 0, (__global DATA_TYPE *)(PTR + 0 * STRIDE_Y + Z##0)); 1486 1487#define STORE_ROW_PARTIAL_2(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1488 STORE_ROW_PARTIAL_1(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1489 VSTORE_PARTIAL(N0, STORE_N0) \ 1490 (BASENAME##1, 0, (__global DATA_TYPE *)(PTR + 1 * STRIDE_Y + Z##1)); 1491 1492#define STORE_ROW_PARTIAL_3(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1493 STORE_ROW_PARTIAL_2(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1494 VSTORE_PARTIAL(N0, STORE_N0) \ 1495 (BASENAME##2, 0, (__global DATA_TYPE *)(PTR + 2 * STRIDE_Y + Z##2)); 1496 1497#define STORE_ROW_PARTIAL_4(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1498 STORE_ROW_PARTIAL_3(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1499 VSTORE_PARTIAL(N0, STORE_N0) \ 1500 (BASENAME##3, 0, (__global DATA_TYPE *)(PTR + 3 * STRIDE_Y + Z##3)); 1501 1502#define STORE_ROW_PARTIAL_5(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1503 STORE_ROW_PARTIAL_4(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1504 VSTORE_PARTIAL(N0, STORE_N0) \ 1505 (BASENAME##4, 0, (__global DATA_TYPE *)(PTR + 4 * STRIDE_Y + Z##4)); 1506 1507#define STORE_ROW_PARTIAL_6(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1508 STORE_ROW_PARTIAL_5(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1509 VSTORE_PARTIAL(N0, STORE_N0) \ 1510 (BASENAME##5, 0, (__global DATA_TYPE *)(PTR + 5 * STRIDE_Y + Z##5)); 1511 1512#define STORE_ROW_PARTIAL_7(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1513 STORE_ROW_PARTIAL_6(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1514 VSTORE_PARTIAL(N0, STORE_N0) \ 1515 (BASENAME##6, 0, (__global DATA_TYPE *)(PTR + 6 * STRIDE_Y + Z##6)); 1516 1517#define STORE_ROW_PARTIAL_8(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1518 STORE_ROW_PARTIAL_7(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1519 VSTORE_PARTIAL(N0, STORE_N0) \ 1520 (BASENAME##7, 0, (__global DATA_TYPE *)(PTR + 7 * STRIDE_Y + Z##7)); 1521 1522#define STORE_ROW_PARTIAL_9(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1523 STORE_ROW_PARTIAL_8(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1524 VSTORE_PARTIAL(N0, STORE_N0) \ 1525 (BASENAME##8, 0, (__global DATA_TYPE *)(PTR + 8 * STRIDE_Y + Z##8)); 1526 1527#define STORE_ROW_PARTIAL_10(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1528 STORE_ROW_PARTIAL_9(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1529 VSTORE_PARTIAL(N0, STORE_N0) \ 1530 (BASENAME##9, 0, (__global DATA_TYPE *)(PTR + 9 * STRIDE_Y + Z##9)); 1531 1532#define STORE_ROW_PARTIAL_11(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1533 STORE_ROW_PARTIAL_10(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1534 VSTORE_PARTIAL(N0, STORE_N0) \ 1535 (BASENAME##A, 0, (__global DATA_TYPE *)(PTR + 10 * STRIDE_Y + Z##A)); 1536 1537#define STORE_ROW_PARTIAL_12(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1538 STORE_ROW_PARTIAL_11(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1539 VSTORE_PARTIAL(N0, STORE_N0) \ 1540 (BASENAME##B, 0, (__global DATA_TYPE *)(PTR + 11 * STRIDE_Y + Z##B)); 1541 1542#define STORE_ROW_PARTIAL_13(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1543 STORE_ROW_PARTIAL_12(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1544 VSTORE_PARTIAL(N0, STORE_N0) \ 1545 (BASENAME##C, 0, (__global DATA_TYPE *)(PTR + 12 * STRIDE_Y + Z##C)); 1546 1547#define STORE_ROW_PARTIAL_14(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1548 STORE_ROW_PARTIAL_13(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1549 VSTORE_PARTIAL(N0, STORE_N0) \ 1550 (BASENAME##D, 0, (__global DATA_TYPE *)(PTR + 13 * STRIDE_Y + Z##D)); 1551 1552#define STORE_ROW_PARTIAL_15(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1553 STORE_ROW_PARTIAL_14(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1554 VSTORE_PARTIAL(N0, STORE_N0) \ 1555 (BASENAME##E, 0, (__global DATA_TYPE *)(PTR + 14 * STRIDE_Y + Z##E)); 1556 1557#define STORE_ROW_PARTIAL_16(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1558 STORE_ROW_PARTIAL_15(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) \ 1559 VSTORE_PARTIAL(N0, STORE_N0) \ 1560 (BASENAME##F, 0, (__global DATA_TYPE *)(PTR + 15 * STRIDE_Y + Z##F)); 1561 1562 1563 1564#define STORE_BLOCK_PARTIAL_STR(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_ROW_PARTIAL_##STORE_M0(N0, STORE_N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 1565#define STORE_BLOCK_PARTIAL(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) STORE_BLOCK_PARTIAL_STR(STORE_M0, STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 1566 1567#define STORE_BLOCK_PARTIAL_IN_X_AND_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 1568 if(!(PARTIAL_COND_X) && !(PARTIAL_COND_Y)) \ 1569 { \ 1570 STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 1571 } \ 1572 else if((PARTIAL_COND_Y) && !(PARTIAL_COND_X)) \ 1573 { \ 1574 STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 1575 } \ 1576 else if(!(PARTIAL_COND_Y) && (PARTIAL_COND_X)) \ 1577 { \ 1578 STORE_BLOCK_PARTIAL(M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 1579 } \ 1580 else \ 1581 { \ 1582 STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 1583 } 1584 1585#define STORE_BLOCK_PARTIAL_IN_X(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_N0, PARTIAL_COND_X) \ 1586 if(!(PARTIAL_COND_X)) \ 1587 { \ 1588 STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 1589 } \ 1590 else \ 1591 { \ 1592 STORE_BLOCK_PARTIAL(M0, PARTIAL_STORE_N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 1593 } 1594 1595#define STORE_BLOCK_PARTIAL_IN_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_COND_Y) \ 1596 if(!(PARTIAL_COND_Y)) \ 1597 { \ 1598 STORE_BLOCK_PARTIAL(M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 1599 } \ 1600 else \ 1601 { \ 1602 STORE_BLOCK_PARTIAL(PARTIAL_STORE_M0, N0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z); \ 1603 } 1604 1605 1606#if defined(PARTIAL_STORE_M0) && defined(PARTIAL_STORE_N0) 1607 1608 1609#if PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 == 0 1610 1611#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 1612 STORE_BLOCK(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z) 1613 1614#elif PARTIAL_STORE_M0 > 0 && PARTIAL_STORE_N0 == 0 1615 1616#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 1617 STORE_BLOCK_PARTIAL_IN_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_COND_Y) 1618 1619#elif PARTIAL_STORE_M0 == 0 && PARTIAL_STORE_N0 > 0 1620 1621#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 1622 STORE_BLOCK_PARTIAL_IN_X(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_N0, PARTIAL_COND_X) 1623 1624#else 1625 1626#define STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ 1627 STORE_BLOCK_PARTIAL_IN_X_AND_Y(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X) 1628 1629#endif 1630 1631#endif 1632 1633 1634#if defined(PARTIAL_STORE_M0) 1635 1636#define COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) \ 1637 ((uint)(max(0, (int)(y * M0) - (int)((M0 - PARTIAL_STORE_M0) % M0)))) 1638#else 1639#define COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) \ 1640 ((uint)(y * M0)) 1641#endif 1642 1643 1644 1645#define STORE_VECTOR_SELECT(basename, data_type, ptr, vec_size, leftover, cond) \ 1646 STORE_BLOCK_PARTIAL_IN_X(1, vec_size, data_type, basename, ptr, 0, 0, leftover, cond) 1647 1648 1649#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16) 1650#pragma OPENCL EXTENSION cl_khr_fp16 : enable 1651#endif 1652 1653#if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) 1654#pragma OPENCL EXTENSION cl_arm_integer_dot_product_int8 : enable 1655#endif 1656 1657#if defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8) 1658#pragma OPENCL EXTENSION cl_arm_integer_dot_product_accumulate_int8 : enable 1659#endif 1660 1661#if defined(ARM_COMPUTE_DEBUG_ENABLED) && defined(cl_arm_printf) 1662#pragma OPENCL EXTENSION cl_arm_printf : enable 1663#endif 1664 1665#define GPU_ARCH_MIDGARD 0x100 1666#define GPU_ARCH_BIFROST 0x200 1667#define GPU_ARCH_VALHALL 0x300 1668 1669 1670#define CONCAT(a, b) a##b 1671 1672 1673#define EXPAND(x) x 1674 1675 1676#define CLAMP(x, min_val, max_val) min(max(x, min_val), max_val) 1677 1678 1679#define REV1(x) ((x)) 1680#define REV2(x) ((x).s10) 1681#define REV3(x) ((x).s210) 1682#define REV4(x) ((x).s3210) 1683#define REV8(x) ((x).s76543210) 1684#define REV16(x) ((x).sFEDCBA9876543210) 1685 1686 1687 1688#define REVERSE_STR(x, s) REV##s((x)) 1689#define REVERSE(x, s) REVERSE_STR(x, s) 1690 1691 1692 1693#define ROT1_0(x) ((x)) 1694#define ROT1_1(x) ((x)) 1695 1696#define ROT2_0(x) ((x)) 1697#define ROT2_1(x) ((x).s10) 1698#define ROT2_2(x) ((x)) 1699 1700#define ROT3_0(x) ((x)) 1701#define ROT3_1(x) ((x).s201) 1702#define ROT3_2(x) ((x).s120) 1703#define ROT3_3(x) ((x)) 1704 1705#define ROT4_0(x) ((x)) 1706#define ROT4_1(x) ((x).s3012) 1707#define ROT4_2(x) ((x).s2301) 1708#define ROT4_3(x) ((x).s1230) 1709#define ROT4_4(x) ((x)) 1710 1711#define ROT8_0(x) ((x)) 1712#define ROT8_1(x) ((x).s70123456) 1713#define ROT8_2(x) ((x).s67012345) 1714#define ROT8_3(x) ((x).s56701234) 1715#define ROT8_4(x) ((x).s45670123) 1716#define ROT8_5(x) ((x).s34567012) 1717#define ROT8_6(x) ((x).s23456701) 1718#define ROT8_7(x) ((x).s12345670) 1719#define ROT8_8(x) ((x)) 1720 1721#define ROT16_0(x) ((x)) 1722#define ROT16_1(x) ((x).sF0123456789ABCDE) 1723#define ROT16_2(x) ((x).sEF0123456789ABCD) 1724#define ROT16_3(x) ((x).sDEF0123456789ABC) 1725#define ROT16_4(x) ((x).sCDEF0123456789AB) 1726#define ROT16_5(x) ((x).sBCDEF0123456789A) 1727#define ROT16_6(x) ((x).sABCDEF0123456789) 1728#define ROT16_7(x) ((x).s9ABCDEF012345678) 1729#define ROT16_8(x) ((x).s89ABCDEF01234567) 1730#define ROT16_9(x) ((x).s789ABCDEF0123456) 1731#define ROT16_10(x) ((x).s6789ABCDEF012345) 1732#define ROT16_11(x) ((x).s56789ABCDEF01234) 1733#define ROT16_12(x) ((x).s456789ABCDEF0123) 1734#define ROT16_13(x) ((x).s3456789ABCDEF012) 1735#define ROT16_14(x) ((x).s23456789ABCDEF01) 1736#define ROT16_15(x) ((x).s123456789ABCDEF0) 1737#define ROT16_16(x) ((x)) 1738 1739 1740 1741#define ROTATE_STR(x, s, n) ROT##s##_##n(x) 1742#define ROTATE(x, s, n) ROTATE_STR(x, s, n) 1743 1744 1745 1746#define V_OFFS1(dt) (dt##1)(0) 1747#define V_OFFS2(dt) (dt##2)(0, 1) 1748#define V_OFFS3(dt) (dt##3)(0, 1, 2) 1749#define V_OFFS4(dt) (dt##4)(0, 1, 2, 3) 1750#define V_OFFS8(dt) (dt##8)(0, 1, 2, 3, 4, 5, 6, 7) 1751#define V_OFFS16(dt) (dt##16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15) 1752 1753 1754 1755#define VEC_OFFS_STR(dt, s) V_OFFS##s(dt) 1756#define VEC_OFFS(dt, s) VEC_OFFS_STR(dt, s) 1757 1758 1759#define VLOAD_STR(size) vload##size 1760#define VLOAD(size) VLOAD_STR(size) 1761 1762 1763#define VLOAD_PARTIAL_STR(size, load_size) vload_partial_##size##_##load_size 1764#define VLOAD_PARTIAL(size, load_size) VLOAD_PARTIAL_STR(size, load_size) 1765 1766#define NO_LOAD(data, offs, ptr) \ 1767 { \ 1768 } 1769 1770 1771#define vload_partial_1_0 NO_LOAD 1772#define vload_partial_1_1 vload1 1773#define vload_partial_1_2 NO_LOAD 1774#define vload_partial_1_3 NO_LOAD 1775#define vload_partial_1_4 NO_LOAD 1776#define vload_partial_1_5 NO_LOAD 1777#define vload_partial_1_6 NO_LOAD 1778#define vload_partial_1_7 NO_LOAD 1779#define vload_partial_1_8 NO_LOAD 1780#define vload_partial_1_9 NO_LOAD 1781#define vload_partial_1_10 NO_LOAD 1782#define vload_partial_1_11 NO_LOAD 1783#define vload_partial_1_12 NO_LOAD 1784#define vload_partial_1_13 NO_LOAD 1785#define vload_partial_1_14 NO_LOAD 1786#define vload_partial_1_15 NO_LOAD 1787#define vload_partial_1_16 NO_LOAD 1788 1789#define vload_partial_2_0 NO_LOAD 1790#define vload_partial_2_1 vload_partial_1 1791#define vload_partial_2_2 vload_partial_2 1792#define vload_partial_2_3 NO_LOAD 1793#define vload_partial_2_4 NO_LOAD 1794#define vload_partial_2_5 NO_LOAD 1795#define vload_partial_2_6 NO_LOAD 1796#define vload_partial_2_7 NO_LOAD 1797#define vload_partial_2_8 NO_LOAD 1798#define vload_partial_2_9 NO_LOAD 1799#define vload_partial_2_10 NO_LOAD 1800#define vload_partial_2_11 NO_LOAD 1801#define vload_partial_2_12 NO_LOAD 1802#define vload_partial_2_13 NO_LOAD 1803#define vload_partial_2_14 NO_LOAD 1804#define vload_partial_2_15 NO_LOAD 1805#define vload_partial_2_16 NO_LOAD 1806 1807#define vload_partial_3_0 NO_LOAD 1808#define vload_partial_3_1 vload_partial_1 1809#define vload_partial_3_2 vload_partial_2 1810#define vload_partial_3_3 vload_partial_3 1811#define vload_partial_3_4 NO_LOAD 1812#define vload_partial_3_5 NO_LOAD 1813#define vload_partial_3_6 NO_LOAD 1814#define vload_partial_3_7 NO_LOAD 1815#define vload_partial_3_8 NO_LOAD 1816#define vload_partial_3_9 NO_LOAD 1817#define vload_partial_3_10 NO_LOAD 1818#define vload_partial_3_11 NO_LOAD 1819#define vload_partial_3_12 NO_LOAD 1820#define vload_partial_3_13 NO_LOAD 1821#define vload_partial_3_14 NO_LOAD 1822#define vload_partial_3_15 NO_LOAD 1823#define vload_partial_3_16 NO_LOAD 1824 1825#define vload_partial_4_0 NO_LOAD 1826#define vload_partial_4_1 vload_partial_1 1827#define vload_partial_4_2 vload_partial_2 1828#define vload_partial_4_3 vload_partial_3 1829#define vload_partial_4_4 vload_partial_4 1830#define vload_partial_4_5 NO_LOAD 1831#define vload_partial_4_6 NO_LOAD 1832#define vload_partial_4_7 NO_LOAD 1833#define vload_partial_4_8 NO_LOAD 1834#define vload_partial_4_9 NO_LOAD 1835#define vload_partial_4_10 NO_LOAD 1836#define vload_partial_4_11 NO_LOAD 1837#define vload_partial_4_12 NO_LOAD 1838#define vload_partial_4_13 NO_LOAD 1839#define vload_partial_4_14 NO_LOAD 1840#define vload_partial_4_15 NO_LOAD 1841#define vload_partial_4_16 NO_LOAD 1842 1843#define vload_partial_8_0 NO_LOAD 1844#define vload_partial_8_1 vload_partial_1 1845#define vload_partial_8_2 vload_partial_2 1846#define vload_partial_8_3 vload_partial_3 1847#define vload_partial_8_4 vload_partial_4 1848#define vload_partial_8_5 vload_partial_5 1849#define vload_partial_8_6 vload_partial_6 1850#define vload_partial_8_7 vload_partial_7 1851#define vload_partial_8_8 vload_partial_8 1852#define vload_partial_8_9 NO_LOAD 1853#define vload_partial_8_10 NO_LOAD 1854#define vload_partial_8_11 NO_LOAD 1855#define vload_partial_8_12 NO_LOAD 1856#define vload_partial_8_13 NO_LOAD 1857#define vload_partial_8_14 NO_LOAD 1858#define vload_partial_8_15 NO_LOAD 1859#define vload_partial_8_16 NO_LOAD 1860 1861#define vload_partial_16_0 NO_LOAD 1862#define vload_partial_16_1 vload_partial_1 1863#define vload_partial_16_2 vload_partial_2 1864#define vload_partial_16_3 vload_partial_3 1865#define vload_partial_16_4 vload_partial_4 1866#define vload_partial_16_5 vload_partial_5 1867#define vload_partial_16_6 vload_partial_6 1868#define vload_partial_16_7 vload_partial_7 1869#define vload_partial_16_8 vload_partial_8 1870#define vload_partial_16_9 vload_partial_9 1871#define vload_partial_16_10 vload_partial_10 1872#define vload_partial_16_11 vload_partial_11 1873#define vload_partial_16_12 vload_partial_12 1874#define vload_partial_16_13 vload_partial_13 1875#define vload_partial_16_14 vload_partial_14 1876#define vload_partial_16_15 vload_partial_15 1877#define vload_partial_16_16 vload_partial_16 1878 1879 1880#define vload_partial_1(DATA, OFFSET, PTR) \ 1881 DATA.s0 = vload1(OFFSET, PTR); 1882 1883#define vload_partial_2(DATA, OFFSET, PTR) \ 1884 DATA.s01 = vload2(OFFSET, PTR); 1885 1886#define vload_partial_3(DATA, OFFSET, PTR) \ 1887 DATA.s012 = vload3(OFFSET, PTR); 1888 1889#define vload_partial_4(DATA, OFFSET, PTR) \ 1890 DATA.s0123 = vload4(OFFSET, PTR); 1891 1892#define vload_partial_5(DATA, OFFSET, PTR) \ 1893 vload_partial_4(DATA.s0123, OFFSET, PTR); \ 1894 DATA.s4 = vload1(OFFSET, PTR + 4); 1895 1896#define vload_partial_6(DATA, OFFSET, PTR) \ 1897 vload_partial_4(DATA.s0123, OFFSET, PTR); \ 1898 vload_partial_2(DATA.s45, OFFSET, PTR + 4); 1899 1900#define vload_partial_7(DATA, OFFSET, PTR) \ 1901 vload_partial_4(DATA.s0123, OFFSET, PTR); \ 1902 vload_partial_3(DATA.s456, OFFSET, PTR + 4); 1903 1904#define vload_partial_8(DATA, OFFSET, PTR) \ 1905 DATA.s01234567 = vload8(OFFSET, PTR); 1906 1907#define vload_partial_9(DATA, OFFSET, PTR) \ 1908 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 1909 DATA.s8 = vload1(OFFSET, PTR + 8); 1910 1911#define vload_partial_10(DATA, OFFSET, PTR) \ 1912 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 1913 vload_partial_2(DATA.s89, OFFSET, PTR + 8); 1914 1915#define vload_partial_11(DATA, OFFSET, PTR) \ 1916 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 1917 vload_partial_3(DATA.s89A, OFFSET, PTR + 8); 1918 1919#define vload_partial_12(DATA, OFFSET, PTR) \ 1920 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 1921 vload_partial_4(DATA.s89AB, OFFSET, PTR + 8); 1922 1923#define vload_partial_13(DATA, OFFSET, PTR) \ 1924 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 1925 vload_partial_5(DATA.s89ABCDEF, OFFSET, PTR + 8); 1926 1927#define vload_partial_14(DATA, OFFSET, PTR) \ 1928 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 1929 vload_partial_6(DATA.s89ABCDEF, OFFSET, PTR + 8); 1930 1931#define vload_partial_15(DATA, OFFSET, PTR) \ 1932 vload_partial_8(DATA.s01234567, OFFSET, PTR); \ 1933 vload_partial_7(DATA.s89ABCDEF, OFFSET, PTR + 8); 1934 1935#define vload_partial_16(DATA, OFFSET, PTR) \ 1936 DATA = vload16(OFFSET, PTR); 1937 1938 1939 1940#define PIXEL_UNIT4 1 1941#define PIXEL_UNIT8 2 1942#define PIXEL_UNIT16 4 1943 1944 1945#define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size) PIXEL_UNIT##vec_size 1946#define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(vec_size) CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size) 1947 1948 1949#define read_image2d_floatx1(img, x_coord, y_coord) (float4)(read_imagef(img, (int2)(x_coord, y_coord))); 1950#define read_image2d_floatx2(img, x_coord, y_coord) (float8)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord))); 1951#define read_image2d_floatx4(img, x_coord, y_coord) (float16)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord)), read_imagef(img, (int2)(x_coord + 2, y_coord)), read_imagef(img, (int2)(x_coord + 3, y_coord))); 1952 1953#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16) 1954#define read_image2d_halfx1(img, x_coord, y_coord) (half4)(read_imageh(img, (int2)(x_coord, y_coord))); 1955#define read_image2d_halfx2(img, x_coord, y_coord) (half8)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord))); 1956#define read_image2d_halfx4(img, x_coord, y_coord) (half16)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord)), read_imageh(img, (int2)(x_coord + 2, y_coord)), read_imageh(img, (int2)(x_coord + 3, y_coord))); 1957#endif 1958 1959#define write_image2d_floatx1(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values)); 1960#define write_image2d_floatx2(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values.s0123), write_imagef(img, (int2)(x_coord + 1, y_coord), values.s4567)); 1961#define write_image2d_floatx4(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values.s0123), write_imagef(img, (int2)(x_coord + 1, y_coord), values.s4567), write_imagef(img, (int2)(x_coord + 2, y_coord), values.s89AB), write_imagef(img, (int2)(x_coord + 3, y_coord), values.sCDEF)); 1962 1963#if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16) 1964#define write_image2d_halfx1(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values)); 1965#define write_image2d_halfx2(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values.s0123), write_imageh(img, (int2)(x_coord + 1, y_coord), values.s4567)); 1966#define write_image2d_halfx4(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values.s0123), write_imageh(img, (int2)(x_coord + 1, y_coord), values.s4567), write_imageh(img, (int2)(x_coord + 2, y_coord), values.s89AB), write_imageh(img, (int2)(x_coord + 3, y_coord), values.sCDEF)); 1967#endif 1968 1969 1970#define READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord) read_image2d_##data_type##x##n0(img, x_coord, y_coord) 1971#define READ_IMAGE2D(data_type, n0, img, x_coord, y_coord) READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord) 1972 1973 1974#define WRITE_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord, values) write_image2d_##data_type##x##n0(img, x_coord, y_coord, values) 1975#define WRITE_IMAGE2D(data_type, n0, img, x_coord, y_coord, values) WRITE_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord, values) 1976 1977#define VSTORE_STR(size) vstore##size 1978#define VSTORE(size) VSTORE_STR(size) 1979 1980#define float1 float 1981#define half1 half 1982#define char1 char 1983#define uchar1 uchar 1984#define short1 short 1985#define ushort1 ushort 1986#define int1 int 1987#define uint1 uint 1988#define long1 long 1989#define ulong1 ulong 1990#define double1 double 1991 1992#define vload1(OFFSET, PTR) *(OFFSET + PTR) 1993#define vstore1(DATA, OFFSET, PTR) *(OFFSET + PTR) = DATA 1994 1995 1996#define VSTORE_PARTIAL_STR(size, store_size) vstore_partial_##size##_##store_size 1997#define VSTORE_PARTIAL(size, store_size) VSTORE_PARTIAL_STR(size, store_size) 1998 1999#define NO_STORE(data, offs, ptr) \ 2000 { \ 2001 } 2002 2003 2004#define vstore_partial_1_0 NO_STORE 2005#define vstore_partial_1_1 vstore1 2006#define vstore_partial_1_2 NO_STORE 2007#define vstore_partial_1_3 NO_STORE 2008#define vstore_partial_1_4 NO_STORE 2009#define vstore_partial_1_5 NO_STORE 2010#define vstore_partial_1_6 NO_STORE 2011#define vstore_partial_1_7 NO_STORE 2012#define vstore_partial_1_8 NO_STORE 2013#define vstore_partial_1_9 NO_STORE 2014#define vstore_partial_1_10 NO_STORE 2015#define vstore_partial_1_11 NO_STORE 2016#define vstore_partial_1_12 NO_STORE 2017#define vstore_partial_1_13 NO_STORE 2018#define vstore_partial_1_14 NO_STORE 2019#define vstore_partial_1_15 NO_STORE 2020#define vstore_partial_1_16 NO_STORE 2021 2022#define vstore_partial_2_0 NO_STORE 2023#define vstore_partial_2_1 vstore_partial_1 2024#define vstore_partial_2_2 vstore_partial_2 2025#define vstore_partial_2_3 NO_STORE 2026#define vstore_partial_2_4 NO_STORE 2027#define vstore_partial_2_5 NO_STORE 2028#define vstore_partial_2_6 NO_STORE 2029#define vstore_partial_2_7 NO_STORE 2030#define vstore_partial_2_8 NO_STORE 2031#define vstore_partial_2_9 NO_STORE 2032#define vstore_partial_2_10 NO_STORE 2033#define vstore_partial_2_11 NO_STORE 2034#define vstore_partial_2_12 NO_STORE 2035#define vstore_partial_2_13 NO_STORE 2036#define vstore_partial_2_14 NO_STORE 2037#define vstore_partial_2_15 NO_STORE 2038#define vstore_partial_2_16 NO_STORE 2039 2040#define vstore_partial_3_0 NO_STORE 2041#define vstore_partial_3_1 vstore_partial_1 2042#define vstore_partial_3_2 vstore_partial_2 2043#define vstore_partial_3_3 vstore_partial_3 2044#define vstore_partial_3_4 NO_STORE 2045#define vstore_partial_3_5 NO_STORE 2046#define vstore_partial_3_6 NO_STORE 2047#define vstore_partial_3_7 NO_STORE 2048#define vstore_partial_3_8 NO_STORE 2049#define vstore_partial_3_9 NO_STORE 2050#define vstore_partial_3_10 NO_STORE 2051#define vstore_partial_3_11 NO_STORE 2052#define vstore_partial_3_12 NO_STORE 2053#define vstore_partial_3_13 NO_STORE 2054#define vstore_partial_3_14 NO_STORE 2055#define vstore_partial_3_15 NO_STORE 2056#define vstore_partial_3_16 NO_STORE 2057 2058#define vstore_partial_4_0 NO_STORE 2059#define vstore_partial_4_1 vstore_partial_1 2060#define vstore_partial_4_2 vstore_partial_2 2061#define vstore_partial_4_3 vstore_partial_3 2062#define vstore_partial_4_4 vstore_partial_4 2063#define vstore_partial_4_5 NO_STORE 2064#define vstore_partial_4_6 NO_STORE 2065#define vstore_partial_4_7 NO_STORE 2066#define vstore_partial_4_8 NO_STORE 2067#define vstore_partial_4_9 NO_STORE 2068#define vstore_partial_4_10 NO_STORE 2069#define vstore_partial_4_11 NO_STORE 2070#define vstore_partial_4_12 NO_STORE 2071#define vstore_partial_4_13 NO_STORE 2072#define vstore_partial_4_14 NO_STORE 2073#define vstore_partial_4_15 NO_STORE 2074#define vstore_partial_4_16 NO_STORE 2075 2076#define vstore_partial_8_0 NO_STORE 2077#define vstore_partial_8_1 vstore_partial_1 2078#define vstore_partial_8_2 vstore_partial_2 2079#define vstore_partial_8_3 vstore_partial_3 2080#define vstore_partial_8_4 vstore_partial_4 2081#define vstore_partial_8_5 vstore_partial_5 2082#define vstore_partial_8_6 vstore_partial_6 2083#define vstore_partial_8_7 vstore_partial_7 2084#define vstore_partial_8_8 vstore_partial_8 2085#define vstore_partial_8_9 NO_STORE 2086#define vstore_partial_8_10 NO_STORE 2087#define vstore_partial_8_11 NO_STORE 2088#define vstore_partial_8_12 NO_STORE 2089#define vstore_partial_8_13 NO_STORE 2090#define vstore_partial_8_14 NO_STORE 2091#define vstore_partial_8_15 NO_STORE 2092#define vstore_partial_8_16 NO_STORE 2093 2094#define vstore_partial_16_0 NO_STORE 2095#define vstore_partial_16_1 vstore_partial_1 2096#define vstore_partial_16_2 vstore_partial_2 2097#define vstore_partial_16_3 vstore_partial_3 2098#define vstore_partial_16_4 vstore_partial_4 2099#define vstore_partial_16_5 vstore_partial_5 2100#define vstore_partial_16_6 vstore_partial_6 2101#define vstore_partial_16_7 vstore_partial_7 2102#define vstore_partial_16_8 vstore_partial_8 2103#define vstore_partial_16_9 vstore_partial_9 2104#define vstore_partial_16_10 vstore_partial_10 2105#define vstore_partial_16_11 vstore_partial_11 2106#define vstore_partial_16_12 vstore_partial_12 2107#define vstore_partial_16_13 vstore_partial_13 2108#define vstore_partial_16_14 vstore_partial_14 2109#define vstore_partial_16_15 vstore_partial_15 2110#define vstore_partial_16_16 vstore_partial_16 2111 2112 2113#define vstore_partial_1(DATA, OFFSET, PTR) \ 2114 vstore1(DATA.s0, OFFSET, PTR); 2115 2116#define vstore_partial_2(DATA, OFFSET, PTR) \ 2117 vstore2(DATA.s01, OFFSET, PTR); 2118 2119#define vstore_partial_3(DATA, OFFSET, PTR) \ 2120 vstore3(DATA.s012, OFFSET, PTR); 2121 2122#define vstore_partial_4(DATA, OFFSET, PTR) \ 2123 vstore4(DATA.s0123, OFFSET, PTR); 2124 2125#define vstore_partial_5(DATA, OFFSET, PTR) \ 2126 vstore_partial_4(DATA.s0123, OFFSET, PTR); \ 2127 vstore1(DATA.s4, OFFSET, PTR + 4); 2128 2129#define vstore_partial_6(DATA, OFFSET, PTR) \ 2130 vstore_partial_4(DATA.s0123, OFFSET, PTR); \ 2131 vstore_partial_2(DATA.s45, OFFSET, PTR + 4); 2132 2133#define vstore_partial_7(DATA, OFFSET, PTR) \ 2134 vstore_partial_4(DATA.s0123, OFFSET, PTR); \ 2135 vstore_partial_3(DATA.s456, OFFSET, PTR + 4); 2136 2137#define vstore_partial_8(DATA, OFFSET, PTR) \ 2138 vstore8(DATA.s01234567, OFFSET, PTR); 2139 2140#define vstore_partial_9(DATA, OFFSET, PTR) \ 2141 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 2142 vstore1(DATA.s8, OFFSET, PTR + 8); 2143 2144#define vstore_partial_10(DATA, OFFSET, PTR) \ 2145 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 2146 vstore_partial_2(DATA.s89, OFFSET, PTR + 8); 2147 2148#define vstore_partial_11(DATA, OFFSET, PTR) \ 2149 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 2150 vstore_partial_3(DATA.s89a, OFFSET, PTR + 8); 2151 2152#define vstore_partial_12(DATA, OFFSET, PTR) \ 2153 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 2154 vstore_partial_4(DATA.s89ab, OFFSET, PTR + 8); 2155 2156#define vstore_partial_13(DATA, OFFSET, PTR) \ 2157 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 2158 vstore_partial_5(DATA.s89abcdef, OFFSET, PTR + 8); 2159 2160#define vstore_partial_14(DATA, OFFSET, PTR) \ 2161 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 2162 vstore_partial_6(DATA.s89abcdef, OFFSET, PTR + 8); 2163 2164#define vstore_partial_15(DATA, OFFSET, PTR) \ 2165 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \ 2166 vstore_partial_7(DATA.s89abcdef, OFFSET, PTR + 8); 2167 2168#define vstore_partial_16(DATA, OFFSET, PTR) \ 2169 vstore16(DATA, OFFSET, PTR); 2170 2171 2172 2173 2174 2175#define convert_float_sat convert_float 2176#define convert_float1_sat convert_float 2177#define convert_float2_sat convert_float2 2178#define convert_float3_sat convert_float3 2179#define convert_float4_sat convert_float4 2180#define convert_float8_sat convert_float8 2181#define convert_float16_sat convert_float16 2182#define convert_half_sat convert_float 2183#define convert_half1_sat convert_half 2184#define convert_half2_sat convert_half2 2185#define convert_half3_sat convert_half3 2186#define convert_half4_sat convert_half4 2187#define convert_half8_sat convert_half8 2188#define convert_half16_sat convert_half16 2189 2190#define convert_float1 convert_float 2191#define convert_half1 convert_half 2192#define convert_char1 convert_char 2193#define convert_uchar1 convert_uchar 2194#define convert_short1 convert_short 2195#define convert_ushort1 convert_ushort 2196#define convert_int1 convert_int 2197#define convert_uint1 convert_uint 2198#define convert_long1 convert_long 2199#define convert_ulong1 convert_ulong 2200#define convert_double1 convert_double 2201 2202#define convert_char1_sat convert_char_sat 2203#define convert_uchar1_sat convert_uchar_sat 2204#define convert_uchar2_sat convert_uchar2_sat 2205#define convert_uchar3_sat convert_uchar3_sat 2206#define convert_uchar4_sat convert_uchar4_sat 2207#define convert_uchar8_sat convert_uchar8_sat 2208#define convert_uchar16_sat convert_uchar16_sat 2209#define convert_short1_sat convert_short_sat 2210#define convert_ushort1_sat convert_ushort_sat 2211#define convert_int1_sat convert_int_sat 2212#define convert_uint1_sat convert_uint_sat 2213#define convert_long1_sat convert_long_sat 2214#define convert_ulong1_sat convert_ulong_sat 2215#define convert_double1_sat convert_double_sat 2216 2217#define VEC_DATA_TYPE_STR(type, size) type##size 2218#define VEC_DATA_TYPE(type, size) VEC_DATA_TYPE_STR(type, size) 2219 2220#define CONVERT_STR(x, type) (convert_##type((x))) 2221#define CONVERT(x, type) CONVERT_STR(x, type) 2222 2223#define CONVERT_SAT_STR(x, type) (convert_##type##_sat((x))) 2224#define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type) 2225 2226#define CONVERT_SAT_ROUND_STR(x, type, round) (convert_##type##_sat_##round((x))) 2227#define CONVERT_SAT_ROUND(x, type, round) CONVERT_SAT_ROUND_STR(x, type, round) 2228 2229#define select_vec_dt_uchar(size) uchar##size 2230#define select_vec_dt_char(size) char##size 2231#define select_vec_dt_ushort(size) ushort##size 2232#define select_vec_dt_short(size) short##size 2233#define select_vec_dt_half(size) short##size 2234#define select_vec_dt_uint(size) uint##size 2235#define select_vec_dt_int(size) int##size 2236#define select_vec_dt_float(size) int##size 2237#define select_vec_dt_ulong(size) ulong##size 2238#define select_vec_dt_long(size) long##size 2239 2240#define SELECT_VEC_DATA_TYPE_STR(type, size) select_vec_dt_##type(size) 2241#define SELECT_VEC_DATA_TYPE(type, size) SELECT_VEC_DATA_TYPE_STR(type, size) 2242#define SELECT_DATA_TYPE(type) SELECT_VEC_DATA_TYPE_STR(type, 1) 2243 2244#define signed_int_vec_dt_uchar(size) char##size 2245#define signed_int_vec_dt_char(size) char##size 2246#define signed_int_vec_dt_ushort(size) short##size 2247#define signed_int_vec_dt_short(size) short##size 2248#define signed_int_vec_dt_half(size) short##size 2249#define signed_int_vec_dt_uint(size) int##size 2250#define signed_int_vec_dt_int(size) int##size 2251#define signed_int_vec_dt_float(size) int##size 2252#define signed_int_vec_dt_ulong(size) long##size 2253#define signed_int_vec_dt_long(size) long##size 2254 2255#define SIGNED_INT_VEC_DATA_TYPE_STR(type, size) signed_int_vec_dt_##type(size) 2256#define SIGNED_INT_VEC_DATA_TYPE(type, size) SIGNED_INT_VEC_DATA_TYPE_STR(type, size) 2257#define SIGNED_INT_DATA_TYPE(type) SIGNED_INT_VEC_DATA_TYPE_STR(type, 1) 2258 2259#define sum_reduce_1(x) (x) 2260#define sum_reduce_2(x) ((x).s0) + ((x).s1) 2261#define sum_reduce_3(x) sum_reduce_2((x).s01) + ((x).s2) 2262#define sum_reduce_4(x) sum_reduce_2((x).s01) + sum_reduce_2((x).s23) 2263#define sum_reduce_8(x) sum_reduce_4((x).s0123) + sum_reduce_4((x).s4567) 2264#define sum_reduce_16(x) sum_reduce_8((x).s01234567) + sum_reduce_8((x).s89ABCDEF) 2265 2266#define SUM_REDUCE_STR(x, size) sum_reduce_##size(x) 2267#define SUM_REDUCE(x, size) SUM_REDUCE_STR(x, size) 2268 2269#define prod_reduce_1(x) (x) 2270#define prod_reduce_2(x) ((x).s0) * ((x).s1) 2271#define prod_reduce_3(x) prod_reduce_2((x).s01) * ((x).s2) 2272#define prod_reduce_4(x) prod_reduce_2((x).s01) * prod_reduce_2((x).s23) 2273#define prod_reduce_8(x) prod_reduce_4((x).s0123) * prod_reduce_4((x).s4567) 2274#define prod_reduce_16(x) prod_reduce_8((x).s01234567) * prod_reduce_8((x).s89ABCDEF) 2275 2276#define PROD_REDUCE_STR(x, size) prod_reduce_##size(x) 2277#define PROD_REDUCE(x, size) PROD_REDUCE_STR(x, size) 2278 2279#define max_reduce_1(x) (x) 2280#define max_reduce_2(x) max(((x).s0), ((x).s1)) 2281#define max_reduce_3(x) max(max_reduce_2((x).s01), ((x).s2)) 2282#define max_reduce_4(x) max(max_reduce_2((x).s01), max_reduce_2((x).s23)) 2283#define max_reduce_8(x) max(max_reduce_4((x).s0123), max_reduce_4((x).s4567)) 2284#define max_reduce_16(x) max(max_reduce_8((x).s01234567), max_reduce_8((x).s89ABCDEF)) 2285 2286#define MAX_REDUCE_STR(x, size) max_reduce_##size(x) 2287#define MAX_REDUCE(x, size) MAX_REDUCE_STR(x, size) 2288 2289#define VECTOR_DECLARATION(name) \ 2290 __global uchar *name##_ptr, \ 2291 uint name##_stride_x, \ 2292 uint name##_step_x, \ 2293 uint name##_offset_first_element_in_bytes 2294 2295#define IMAGE_DECLARATION(name) \ 2296 __global uchar *name##_ptr, \ 2297 uint name##_stride_x, \ 2298 uint name##_step_x, \ 2299 uint name##_stride_y, \ 2300 uint name##_step_y, \ 2301 uint name##_offset_first_element_in_bytes 2302 2303#define TENSOR3D_DECLARATION(name) \ 2304 __global uchar *name##_ptr, \ 2305 uint name##_stride_x, \ 2306 uint name##_step_x, \ 2307 uint name##_stride_y, \ 2308 uint name##_step_y, \ 2309 uint name##_stride_z, \ 2310 uint name##_step_z, \ 2311 uint name##_offset_first_element_in_bytes 2312 2313#define TENSOR4D_DECLARATION(name) \ 2314 __global uchar *name##_ptr, \ 2315 uint name##_stride_x, \ 2316 uint name##_step_x, \ 2317 uint name##_stride_y, \ 2318 uint name##_step_y, \ 2319 uint name##_stride_z, \ 2320 uint name##_step_z, \ 2321 uint name##_stride_w, \ 2322 uint name##_step_w, \ 2323 uint name##_offset_first_element_in_bytes 2324 2325#define TENSOR5D_DECLARATION(name) \ 2326 __global uchar *name##_ptr, \ 2327 uint name##_stride_x, \ 2328 uint name##_step_x, \ 2329 uint name##_stride_y, \ 2330 uint name##_step_y, \ 2331 uint name##_stride_z, \ 2332 uint name##_step_z, \ 2333 uint name##_stride_w, \ 2334 uint name##_step_w, \ 2335 uint name##_stride_v, \ 2336 uint name##_step_v, \ 2337 uint name##_offset_first_element_in_bytes 2338 2339#define CONVERT_TO_VECTOR_STRUCT(name) \ 2340 update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x) 2341 2342#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name) \ 2343 update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0) 2344 2345#define CONVERT_TO_IMAGE_STRUCT(name) \ 2346 update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y) 2347 2348#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name) \ 2349 update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0) 2350 2351#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \ 2352 update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, name##_step_z) 2353 2354#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name) \ 2355 update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, name##_step_z) 2356 2357#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \ 2358 update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, name##_step_z) 2359 2360#define CONVERT_TO_TENSOR3D_STRUCT(name) \ 2361 update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \ 2362 name##_stride_z, name##_step_z) 2363 2364#define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name) \ 2365 update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0) 2366 2367#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size) \ 2368 update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \ 2369 name##_stride_z, name##_step_z, name##_stride_w, name##_step_w, mod_size) 2370 2371#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size) \ 2372 update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0, name##_stride_w, 0, mod_size) 2373 2374#define CONVERT_TO_TENSOR3D_STRUCT_NO_UPDATE_PTR(name) \ 2375 tensor3D_ptr_no_update(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \ 2376 name##_stride_z, name##_step_z) 2377 2378 2379typedef struct Vector 2380{ 2381 __global uchar *ptr; 2382 int offset_first_element_in_bytes; 2383 int stride_x; 2384} Vector; 2385 2386 2387typedef struct Image 2388{ 2389 __global uchar *ptr; 2390 int offset_first_element_in_bytes; 2391 int stride_x; 2392 int stride_y; 2393} Image; 2394 2395 2396typedef struct Tensor3D 2397{ 2398 __global uchar *ptr; 2399 int offset_first_element_in_bytes; 2400 int stride_x; 2401 int stride_y; 2402 int stride_z; 2403} Tensor3D; 2404 2405 2406typedef struct Tensor4D 2407{ 2408 __global uchar *ptr; 2409 int offset_first_element_in_bytes; 2410 int stride_x; 2411 int stride_y; 2412 int stride_z; 2413 int stride_w; 2414} Tensor4D; 2415 2416 2417inline Vector update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x) 2418{ 2419 Vector vector = 2420 { 2421 .ptr = ptr, 2422 .offset_first_element_in_bytes = offset_first_element_in_bytes, 2423 .stride_x = stride_x, 2424 }; 2425 vector.ptr += vector.offset_first_element_in_bytes + get_global_id(0) * step_x; 2426 return vector; 2427} 2428 2429 2430inline Image update_image_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y) 2431{ 2432 Image img = 2433 { 2434 .ptr = ptr, 2435 .offset_first_element_in_bytes = offset_first_element_in_bytes, 2436 .stride_x = stride_x, 2437 .stride_y = stride_y 2438 }; 2439 img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y; 2440 return img; 2441} 2442 2443 2444inline Image update_image_from_tensor3D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z) 2445{ 2446 Image img = 2447 { 2448 .ptr = ptr, 2449 .offset_first_element_in_bytes = offset_first_element_in_bytes, 2450 .stride_x = stride_x, 2451 .stride_y = stride_y 2452 }; 2453 img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z; 2454 return img; 2455} 2456 2457 2458inline Tensor3D update_tensor3D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z) 2459{ 2460 Tensor3D tensor = 2461 { 2462 .ptr = ptr, 2463 .offset_first_element_in_bytes = offset_first_element_in_bytes, 2464 .stride_x = stride_x, 2465 .stride_y = stride_y, 2466 .stride_z = stride_z 2467 }; 2468 tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z; 2469 return tensor; 2470} 2471 2472 2473inline Tensor3D tensor3D_ptr_no_update(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z) 2474{ 2475 Tensor3D tensor = 2476 { 2477 .ptr = ptr, 2478 .offset_first_element_in_bytes = offset_first_element_in_bytes, 2479 .stride_x = stride_x, 2480 .stride_y = stride_y, 2481 .stride_z = stride_z 2482 }; 2483 return tensor; 2484} 2485 2486inline Tensor4D update_tensor4D_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y, uint stride_z, uint step_z, uint stride_w, 2487 uint step_w, 2488 uint mod_size) 2489{ 2490 Tensor4D tensor = 2491 { 2492 .ptr = ptr, 2493 .offset_first_element_in_bytes = offset_first_element_in_bytes, 2494 .stride_x = stride_x, 2495 .stride_y = stride_y, 2496 .stride_z = stride_z, 2497 .stride_w = stride_w 2498 }; 2499 2500 tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + (get_global_id(2) % mod_size) * step_z + (get_global_id(2) / mod_size) * step_w; 2501 return tensor; 2502} 2503 2504 2505inline __global const uchar *vector_offset(const Vector *vec, int x) 2506{ 2507 return vec->ptr + x * vec->stride_x; 2508} 2509 2510 2511inline __global uchar *offset(const Image *img, int x, int y) 2512{ 2513 return img->ptr + x * img->stride_x + y * img->stride_y; 2514} 2515 2516 2517inline __global const uchar *tensor3D_offset(const Tensor3D *tensor, int x, int y, int z) 2518{ 2519 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z; 2520} 2521 2522 2523inline __global const uchar *tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w) 2524{ 2525 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + w * tensor->stride_w; 2526} 2527 2528 2529inline __global const uchar *tensor3D_index2ptr(const Tensor3D *tensor, uint width, uint height, uint depth, uint index) 2530{ 2531 uint num_elements = width * height; 2532 2533 const uint z = index / num_elements; 2534 2535 index %= num_elements; 2536 2537 const uint y = index / width; 2538 2539 index %= width; 2540 2541 const uint x = index; 2542 2543 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + tensor->offset_first_element_in_bytes; 2544} 2545 2546#endif 2547 2548#ifndef SRC_CORE_CL_CL_KERNELS_TILE_HELPERS 2549#define SRC_CORE_CL_CL_KERNELS_TILE_HELPERS 2550 2551 2552 2553 2554#define TILE_VECTOR_SIZE1 1 2555#define TILE_VECTOR_SIZE2 2 2556#define TILE_VECTOR_SIZE3 3 2557#define TILE_VECTOR_SIZE4 4 2558#define TILE_VECTOR_SIZE5 8 2559#define TILE_VECTOR_SIZE6 8 2560#define TILE_VECTOR_SIZE7 8 2561#define TILE_VECTOR_SIZE8 8 2562#define TILE_VECTOR_SIZE9 16 2563#define TILE_VECTOR_SIZE10 16 2564#define TILE_VECTOR_SIZE11 16 2565#define TILE_VECTOR_SIZE12 16 2566#define TILE_VECTOR_SIZE13 16 2567#define TILE_VECTOR_SIZE14 16 2568#define TILE_VECTOR_SIZE15 16 2569#define TILE_VECTOR_SIZE16 16 2570 2571#define TILE_VECTOR_TYPE1(DATA_TYPE) DATA_TYPE##1 2572#define TILE_VECTOR_TYPE2(DATA_TYPE) DATA_TYPE##2 2573#define TILE_VECTOR_TYPE3(DATA_TYPE) DATA_TYPE##3 2574#define TILE_VECTOR_TYPE4(DATA_TYPE) DATA_TYPE##4 2575#define TILE_VECTOR_TYPE5(DATA_TYPE) DATA_TYPE##8 2576#define TILE_VECTOR_TYPE6(DATA_TYPE) DATA_TYPE##8 2577#define TILE_VECTOR_TYPE7(DATA_TYPE) DATA_TYPE##8 2578#define TILE_VECTOR_TYPE8(DATA_TYPE) DATA_TYPE##8 2579#define TILE_VECTOR_TYPE9(DATA_TYPE) DATA_TYPE##16 2580#define TILE_VECTOR_TYPE10(DATA_TYPE) DATA_TYPE##16 2581#define TILE_VECTOR_TYPE11(DATA_TYPE) DATA_TYPE##16 2582#define TILE_VECTOR_TYPE12(DATA_TYPE) DATA_TYPE##16 2583#define TILE_VECTOR_TYPE13(DATA_TYPE) DATA_TYPE##16 2584#define TILE_VECTOR_TYPE14(DATA_TYPE) DATA_TYPE##16 2585#define TILE_VECTOR_TYPE15(DATA_TYPE) DATA_TYPE##16 2586#define TILE_VECTOR_TYPE16(DATA_TYPE) DATA_TYPE##16 2587 2588 2589#define TILE(DATA_TYPE, H, W, BASENAME) TILE_STR(DATA_TYPE, H, W, BASENAME) 2590#define TILE_STR(DATA_TYPE, H, W, BASENAME) \ 2591 union { \ 2592 DATA_TYPE s[TILE_VECTOR_SIZE##W]; \ 2593 TILE_VECTOR_TYPE##W(DATA_TYPE) v; \ 2594 } BASENAME[H] 2595 2596#define TENSOR4D_IMAGE(name) \ 2597 __read_only image2d_t name##_img, \ 2598 __global uchar *name##_ptr, \ 2599 uint name##_stride_x, \ 2600 uint name##_step_x, \ 2601 uint name##_stride_y, \ 2602 uint name##_step_y, \ 2603 uint name##_stride_z, \ 2604 uint name##_step_z, \ 2605 uint name##_stride_w, \ 2606 uint name##_step_w, \ 2607 uint name##_offset_first_element_in_bytes 2608 2609#define TENSOR4D_BUFFER(name) \ 2610 __global uchar *name##_ptr, \ 2611 uint name##_stride_x, \ 2612 uint name##_step_x, \ 2613 uint name##_stride_y, \ 2614 uint name##_step_y, \ 2615 uint name##_stride_z, \ 2616 uint name##_step_z, \ 2617 uint name##_stride_w, \ 2618 uint name##_step_w, \ 2619 uint name##_offset_first_element_in_bytes 2620 2621#define TENSOR4D_STR(name, type) TENSOR4D_##type(name) 2622#define TENSOR4D(name, type) TENSOR4D_STR(name, type) 2623 2624#define TENSOR4D_T_IMAGE(name) \ 2625 __read_only image2d_t name##_img, \ 2626 __global uchar *name##_ptr, \ 2627 uint name##_stride_y, \ 2628 uint name##_stride_z, \ 2629 uint name##_stride_w, \ 2630 uint name##_c, \ 2631 uint name##_w, \ 2632 uint name##_h, \ 2633 uint name##_n, \ 2634 uint name##_offset_first_element_in_bytes 2635 2636#define TENSOR4D_T_BUFFER(name) \ 2637 __global uchar *name##_ptr, \ 2638 uint name##_stride_y, \ 2639 uint name##_stride_z, \ 2640 uint name##_stride_w, \ 2641 uint name##_c, \ 2642 uint name##_w, \ 2643 uint name##_h, \ 2644 uint name##_n, \ 2645 uint name##_offset_first_element_in_bytes 2646 2647#define TENSOR4D_T_STR(name, type) TENSOR4D_T_##type(name) 2648 2649 2650#define TENSOR4D_T(name, type) TENSOR4D_T_STR(name, type) 2651 2652#define TENSOR4D_RO_T_IMAGE(name) \ 2653 __read_only image2d_t name##_img, \ 2654 TENSOR4D_T_BUFFER(name) 2655 2656#define TENSOR4D_RO_T_BUFFER(name) TENSOR4D_T_BUFFER(name) 2657 2658#define TENSOR4D_RO_T_STR(name, type) TENSOR4D_RO_T_##type(name) 2659 2660 2661#define TENSOR4D_RO_T(name, type) TENSOR4D_RO_T_STR(name, type) 2662 2663#define TENSOR4D_WO_T_IMAGE(name) \ 2664 __write_only image2d_t name##_img, \ 2665 TENSOR4D_T_BUFFER(name) 2666 2667#define TENSOR4D_WO_T_BUFFER(name) TENSOR4D_T_BUFFER(name) 2668 2669#define TENSOR4D_WO_T_STR(name, type) TENSOR4D_WO_T_##type(name) 2670 2671 2672#define TENSOR4D_WO_T(name, type) TENSOR4D_WO_T_STR(name, type) 2673 2674#define TENSOR3D_T_IMAGE(name) \ 2675 __read_only image2d_t name##_img, \ 2676 __global uchar *name##_ptr, \ 2677 uint name##_stride_y, \ 2678 uint name##_stride_z, \ 2679 uint name##_w, \ 2680 uint name##_h, \ 2681 uint name##_n, \ 2682 uint name##_offset_first_element_in_bytes 2683 2684#define TENSOR3D_T_BUFFER(name) \ 2685 __global uchar *name##_ptr, \ 2686 uint name##_stride_y, \ 2687 uint name##_stride_z, \ 2688 uint name##_w, \ 2689 uint name##_h, \ 2690 uint name##_n, \ 2691 uint name##_offset_first_element_in_bytes 2692 2693#define TENSOR3D_T_STR(name, type) TENSOR3D_T_##type(name) 2694#define TENSOR3D_T(name, type) TENSOR3D_T_STR(name, type) 2695 2696#if !defined(UNROLL_WITH_PRAGMA) 2697#define UNROLL_INCR(idx, step, macro) idx += (step); (macro) 2698 2699#define LOOP_UNROLLING_1(idx, step, macro) (macro) 2700#define LOOP_UNROLLING_2(idx, step, macro) LOOP_UNROLLING_1(idx, step, macro); UNROLL_INCR(idx, step, macro) 2701#define LOOP_UNROLLING_3(idx, step, macro) LOOP_UNROLLING_2(idx, step, macro); UNROLL_INCR(idx, step, macro) 2702#define LOOP_UNROLLING_4(idx, step, macro) LOOP_UNROLLING_3(idx, step, macro); UNROLL_INCR(idx, step, macro) 2703#define LOOP_UNROLLING_5(idx, step, macro) LOOP_UNROLLING_4(idx, step, macro); UNROLL_INCR(idx, step, macro) 2704#define LOOP_UNROLLING_6(idx, step, macro) LOOP_UNROLLING_5(idx, step, macro); UNROLL_INCR(idx, step, macro) 2705#define LOOP_UNROLLING_7(idx, step, macro) LOOP_UNROLLING_6(idx, step, macro); UNROLL_INCR(idx, step, macro) 2706#define LOOP_UNROLLING_8(idx, step, macro) LOOP_UNROLLING_7(idx, step, macro); UNROLL_INCR(idx, step, macro) 2707#define LOOP_UNROLLING_9(idx, step, macro) LOOP_UNROLLING_8(idx, step, macro); UNROLL_INCR(idx, step, macro) 2708#define LOOP_UNROLLING_10(idx, step, macro) LOOP_UNROLLING_9(idx, step, macro); UNROLL_INCR(idx, step, macro) 2709#define LOOP_UNROLLING_11(idx, step, macro) LOOP_UNROLLING_10(idx, step, macro); UNROLL_INCR(idx, step, macro) 2710#define LOOP_UNROLLING_12(idx, step, macro) LOOP_UNROLLING_11(idx, step, macro); UNROLL_INCR(idx, step, macro) 2711#define LOOP_UNROLLING_13(idx, step, macro) LOOP_UNROLLING_12(idx, step, macro); UNROLL_INCR(idx, step, macro) 2712#define LOOP_UNROLLING_14(idx, step, macro) LOOP_UNROLLING_13(idx, step, macro); UNROLL_INCR(idx, step, macro) 2713#define LOOP_UNROLLING_15(idx, step, macro) LOOP_UNROLLING_14(idx, step, macro); UNROLL_INCR(idx, step, macro) 2714#define LOOP_UNROLLING_16(idx, step, macro) LOOP_UNROLLING_15(idx, step, macro); UNROLL_INCR(idx, step, macro) 2715#define LOOP_UNROLLING_17(idx, step, macro) LOOP_UNROLLING_16(idx, step, macro); UNROLL_INCR(idx, step, macro) 2716#define LOOP_UNROLLING_18(idx, step, macro) LOOP_UNROLLING_17(idx, step, macro); UNROLL_INCR(idx, step, macro) 2717#define LOOP_UNROLLING_19(idx, step, macro) LOOP_UNROLLING_18(idx, step, macro); UNROLL_INCR(idx, step, macro) 2718#define LOOP_UNROLLING_20(idx, step, macro) LOOP_UNROLLING_19(idx, step, macro); UNROLL_INCR(idx, step, macro) 2719#define LOOP_UNROLLING_21(idx, step, macro) LOOP_UNROLLING_20(idx, step, macro); UNROLL_INCR(idx, step, macro) 2720#define LOOP_UNROLLING_22(idx, step, macro) LOOP_UNROLLING_21(idx, step, macro); UNROLL_INCR(idx, step, macro) 2721#define LOOP_UNROLLING_23(idx, step, macro) LOOP_UNROLLING_22(idx, step, macro); UNROLL_INCR(idx, step, macro) 2722#define LOOP_UNROLLING_24(idx, step, macro) LOOP_UNROLLING_23(idx, step, macro); UNROLL_INCR(idx, step, macro) 2723#define LOOP_UNROLLING_25(idx, step, macro) LOOP_UNROLLING_24(idx, step, macro); UNROLL_INCR(idx, step, macro) 2724#define LOOP_UNROLLING_26(idx, step, macro) LOOP_UNROLLING_25(idx, step, macro); UNROLL_INCR(idx, step, macro) 2725#define LOOP_UNROLLING_27(idx, step, macro) LOOP_UNROLLING_26(idx, step, macro); UNROLL_INCR(idx, step, macro) 2726#define LOOP_UNROLLING_28(idx, step, macro) LOOP_UNROLLING_27(idx, step, macro); UNROLL_INCR(idx, step, macro) 2727#define LOOP_UNROLLING_29(idx, step, macro) LOOP_UNROLLING_28(idx, step, macro); UNROLL_INCR(idx, step, macro) 2728#define LOOP_UNROLLING_30(idx, step, macro) LOOP_UNROLLING_29(idx, step, macro); UNROLL_INCR(idx, step, macro) 2729#define LOOP_UNROLLING_31(idx, step, macro) LOOP_UNROLLING_30(idx, step, macro); UNROLL_INCR(idx, step, macro) 2730#define LOOP_UNROLLING_32(idx, step, macro) LOOP_UNROLLING_31(idx, step, macro); UNROLL_INCR(idx, step, macro) 2731#define LOOP_UNROLLING_33(idx, step, macro) LOOP_UNROLLING_32(idx, step, macro); UNROLL_INCR(idx, step, macro) 2732#define LOOP_UNROLLING_34(idx, step, macro) LOOP_UNROLLING_33(idx, step, macro); UNROLL_INCR(idx, step, macro) 2733#define LOOP_UNROLLING_35(idx, step, macro) LOOP_UNROLLING_34(idx, step, macro); UNROLL_INCR(idx, step, macro) 2734#define LOOP_UNROLLING_36(idx, step, macro) LOOP_UNROLLING_35(idx, step, macro); UNROLL_INCR(idx, step, macro) 2735#define LOOP_UNROLLING_37(idx, step, macro) LOOP_UNROLLING_36(idx, step, macro); UNROLL_INCR(idx, step, macro) 2736#define LOOP_UNROLLING_38(idx, step, macro) LOOP_UNROLLING_37(idx, step, macro); UNROLL_INCR(idx, step, macro) 2737#define LOOP_UNROLLING_39(idx, step, macro) LOOP_UNROLLING_38(idx, step, macro); UNROLL_INCR(idx, step, macro) 2738#define LOOP_UNROLLING_40(idx, step, macro) LOOP_UNROLLING_39(idx, step, macro); UNROLL_INCR(idx, step, macro) 2739#define LOOP_UNROLLING_41(idx, step, macro) LOOP_UNROLLING_40(idx, step, macro); UNROLL_INCR(idx, step, macro) 2740#define LOOP_UNROLLING_42(idx, step, macro) LOOP_UNROLLING_41(idx, step, macro); UNROLL_INCR(idx, step, macro) 2741#define LOOP_UNROLLING_43(idx, step, macro) LOOP_UNROLLING_42(idx, step, macro); UNROLL_INCR(idx, step, macro) 2742#define LOOP_UNROLLING_44(idx, step, macro) LOOP_UNROLLING_43(idx, step, macro); UNROLL_INCR(idx, step, macro) 2743#define LOOP_UNROLLING_45(idx, step, macro) LOOP_UNROLLING_44(idx, step, macro); UNROLL_INCR(idx, step, macro) 2744#define LOOP_UNROLLING_46(idx, step, macro) LOOP_UNROLLING_45(idx, step, macro); UNROLL_INCR(idx, step, macro) 2745#define LOOP_UNROLLING_47(idx, step, macro) LOOP_UNROLLING_46(idx, step, macro); UNROLL_INCR(idx, step, macro) 2746#define LOOP_UNROLLING_48(idx, step, macro) LOOP_UNROLLING_47(idx, step, macro); UNROLL_INCR(idx, step, macro) 2747#define LOOP_UNROLLING_49(idx, step, macro) LOOP_UNROLLING_48(idx, step, macro); UNROLL_INCR(idx, step, macro) 2748#define LOOP_UNROLLING_50(idx, step, macro) LOOP_UNROLLING_49(idx, step, macro); UNROLL_INCR(idx, step, macro) 2749#define LOOP_UNROLLING_51(idx, step, macro) LOOP_UNROLLING_50(idx, step, macro); UNROLL_INCR(idx, step, macro) 2750#define LOOP_UNROLLING_52(idx, step, macro) LOOP_UNROLLING_51(idx, step, macro); UNROLL_INCR(idx, step, macro) 2751#define LOOP_UNROLLING_53(idx, step, macro) LOOP_UNROLLING_52(idx, step, macro); UNROLL_INCR(idx, step, macro) 2752#define LOOP_UNROLLING_54(idx, step, macro) LOOP_UNROLLING_53(idx, step, macro); UNROLL_INCR(idx, step, macro) 2753#define LOOP_UNROLLING_55(idx, step, macro) LOOP_UNROLLING_54(idx, step, macro); UNROLL_INCR(idx, step, macro) 2754#define LOOP_UNROLLING_56(idx, step, macro) LOOP_UNROLLING_55(idx, step, macro); UNROLL_INCR(idx, step, macro) 2755#define LOOP_UNROLLING_57(idx, step, macro) LOOP_UNROLLING_56(idx, step, macro); UNROLL_INCR(idx, step, macro) 2756#define LOOP_UNROLLING_58(idx, step, macro) LOOP_UNROLLING_57(idx, step, macro); UNROLL_INCR(idx, step, macro) 2757#define LOOP_UNROLLING_59(idx, step, macro) LOOP_UNROLLING_58(idx, step, macro); UNROLL_INCR(idx, step, macro) 2758#define LOOP_UNROLLING_60(idx, step, macro) LOOP_UNROLLING_59(idx, step, macro); UNROLL_INCR(idx, step, macro) 2759#define LOOP_UNROLLING_61(idx, step, macro) LOOP_UNROLLING_60(idx, step, macro); UNROLL_INCR(idx, step, macro) 2760#define LOOP_UNROLLING_62(idx, step, macro) LOOP_UNROLLING_61(idx, step, macro); UNROLL_INCR(idx, step, macro) 2761#define LOOP_UNROLLING_63(idx, step, macro) LOOP_UNROLLING_62(idx, step, macro); UNROLL_INCR(idx, step, macro) 2762#define LOOP_UNROLLING_64(idx, step, macro) LOOP_UNROLLING_63(idx, step, macro); UNROLL_INCR(idx, step, macro) 2763#define LOOP_UNROLLING_65(idx, step, macro) LOOP_UNROLLING_64(idx, step, macro); UNROLL_INCR(idx, step, macro) 2764#define LOOP_UNROLLING_66(idx, step, macro) LOOP_UNROLLING_65(idx, step, macro); UNROLL_INCR(idx, step, macro) 2765#define LOOP_UNROLLING_67(idx, step, macro) LOOP_UNROLLING_66(idx, step, macro); UNROLL_INCR(idx, step, macro) 2766#define LOOP_UNROLLING_68(idx, step, macro) LOOP_UNROLLING_67(idx, step, macro); UNROLL_INCR(idx, step, macro) 2767#define LOOP_UNROLLING_69(idx, step, macro) LOOP_UNROLLING_68(idx, step, macro); UNROLL_INCR(idx, step, macro) 2768#define LOOP_UNROLLING_70(idx, step, macro) LOOP_UNROLLING_69(idx, step, macro); UNROLL_INCR(idx, step, macro) 2769#define LOOP_UNROLLING_71(idx, step, macro) LOOP_UNROLLING_70(idx, step, macro); UNROLL_INCR(idx, step, macro) 2770#define LOOP_UNROLLING_72(idx, step, macro) LOOP_UNROLLING_71(idx, step, macro); UNROLL_INCR(idx, step, macro) 2771#define LOOP_UNROLLING_73(idx, step, macro) LOOP_UNROLLING_72(idx, step, macro); UNROLL_INCR(idx, step, macro) 2772#define LOOP_UNROLLING_74(idx, step, macro) LOOP_UNROLLING_73(idx, step, macro); UNROLL_INCR(idx, step, macro) 2773#define LOOP_UNROLLING_75(idx, step, macro) LOOP_UNROLLING_74(idx, step, macro); UNROLL_INCR(idx, step, macro) 2774#define LOOP_UNROLLING_76(idx, step, macro) LOOP_UNROLLING_75(idx, step, macro); UNROLL_INCR(idx, step, macro) 2775#define LOOP_UNROLLING_77(idx, step, macro) LOOP_UNROLLING_76(idx, step, macro); UNROLL_INCR(idx, step, macro) 2776#define LOOP_UNROLLING_78(idx, step, macro) LOOP_UNROLLING_77(idx, step, macro); UNROLL_INCR(idx, step, macro) 2777#define LOOP_UNROLLING_79(idx, step, macro) LOOP_UNROLLING_78(idx, step, macro); UNROLL_INCR(idx, step, macro) 2778#define LOOP_UNROLLING_80(idx, step, macro) LOOP_UNROLLING_79(idx, step, macro); UNROLL_INCR(idx, step, macro) 2779#define LOOP_UNROLLING_81(idx, step, macro) LOOP_UNROLLING_80(idx, step, macro); UNROLL_INCR(idx, step, macro) 2780#define LOOP_UNROLLING_82(idx, step, macro) LOOP_UNROLLING_81(idx, step, macro); UNROLL_INCR(idx, step, macro) 2781#define LOOP_UNROLLING_83(idx, step, macro) LOOP_UNROLLING_82(idx, step, macro); UNROLL_INCR(idx, step, macro) 2782#define LOOP_UNROLLING_84(idx, step, macro) LOOP_UNROLLING_83(idx, step, macro); UNROLL_INCR(idx, step, macro) 2783#define LOOP_UNROLLING_85(idx, step, macro) LOOP_UNROLLING_84(idx, step, macro); UNROLL_INCR(idx, step, macro) 2784#define LOOP_UNROLLING_86(idx, step, macro) LOOP_UNROLLING_85(idx, step, macro); UNROLL_INCR(idx, step, macro) 2785#define LOOP_UNROLLING_87(idx, step, macro) LOOP_UNROLLING_86(idx, step, macro); UNROLL_INCR(idx, step, macro) 2786#define LOOP_UNROLLING_88(idx, step, macro) LOOP_UNROLLING_87(idx, step, macro); UNROLL_INCR(idx, step, macro) 2787#define LOOP_UNROLLING_89(idx, step, macro) LOOP_UNROLLING_88(idx, step, macro); UNROLL_INCR(idx, step, macro) 2788#define LOOP_UNROLLING_90(idx, step, macro) LOOP_UNROLLING_89(idx, step, macro); UNROLL_INCR(idx, step, macro) 2789#define LOOP_UNROLLING_91(idx, step, macro) LOOP_UNROLLING_90(idx, step, macro); UNROLL_INCR(idx, step, macro) 2790#define LOOP_UNROLLING_92(idx, step, macro) LOOP_UNROLLING_91(idx, step, macro); UNROLL_INCR(idx, step, macro) 2791#define LOOP_UNROLLING_93(idx, step, macro) LOOP_UNROLLING_92(idx, step, macro); UNROLL_INCR(idx, step, macro) 2792#define LOOP_UNROLLING_94(idx, step, macro) LOOP_UNROLLING_93(idx, step, macro); UNROLL_INCR(idx, step, macro) 2793#define LOOP_UNROLLING_95(idx, step, macro) LOOP_UNROLLING_94(idx, step, macro); UNROLL_INCR(idx, step, macro) 2794#define LOOP_UNROLLING_96(idx, step, macro) LOOP_UNROLLING_95(idx, step, macro); UNROLL_INCR(idx, step, macro) 2795#define LOOP_UNROLLING_97(idx, step, macro) LOOP_UNROLLING_96(idx, step, macro); UNROLL_INCR(idx, step, macro) 2796#define LOOP_UNROLLING_98(idx, step, macro) LOOP_UNROLLING_97(idx, step, macro); UNROLL_INCR(idx, step, macro) 2797#define LOOP_UNROLLING_99(idx, step, macro) LOOP_UNROLLING_98(idx, step, macro); UNROLL_INCR(idx, step, macro) 2798#define LOOP_UNROLLING_100(idx, step, macro) LOOP_UNROLLING_99(idx, step, macro); UNROLL_INCR(idx, step, macro) 2799#define LOOP_UNROLLING_101(idx, step, macro) LOOP_UNROLLING_100(idx, step, macro); UNROLL_INCR(idx, step, macro) 2800#define LOOP_UNROLLING_102(idx, step, macro) LOOP_UNROLLING_101(idx, step, macro); UNROLL_INCR(idx, step, macro) 2801#define LOOP_UNROLLING_103(idx, step, macro) LOOP_UNROLLING_102(idx, step, macro); UNROLL_INCR(idx, step, macro) 2802#define LOOP_UNROLLING_104(idx, step, macro) LOOP_UNROLLING_103(idx, step, macro); UNROLL_INCR(idx, step, macro) 2803#define LOOP_UNROLLING_105(idx, step, macro) LOOP_UNROLLING_104(idx, step, macro); UNROLL_INCR(idx, step, macro) 2804#define LOOP_UNROLLING_106(idx, step, macro) LOOP_UNROLLING_105(idx, step, macro); UNROLL_INCR(idx, step, macro) 2805#define LOOP_UNROLLING_107(idx, step, macro) LOOP_UNROLLING_106(idx, step, macro); UNROLL_INCR(idx, step, macro) 2806#define LOOP_UNROLLING_108(idx, step, macro) LOOP_UNROLLING_107(idx, step, macro); UNROLL_INCR(idx, step, macro) 2807#define LOOP_UNROLLING_109(idx, step, macro) LOOP_UNROLLING_108(idx, step, macro); UNROLL_INCR(idx, step, macro) 2808#define LOOP_UNROLLING_110(idx, step, macro) LOOP_UNROLLING_109(idx, step, macro); UNROLL_INCR(idx, step, macro) 2809#define LOOP_UNROLLING_111(idx, step, macro) LOOP_UNROLLING_110(idx, step, macro); UNROLL_INCR(idx, step, macro) 2810#define LOOP_UNROLLING_112(idx, step, macro) LOOP_UNROLLING_111(idx, step, macro); UNROLL_INCR(idx, step, macro) 2811#define LOOP_UNROLLING_113(idx, step, macro) LOOP_UNROLLING_112(idx, step, macro); UNROLL_INCR(idx, step, macro) 2812#define LOOP_UNROLLING_114(idx, step, macro) LOOP_UNROLLING_113(idx, step, macro); UNROLL_INCR(idx, step, macro) 2813#define LOOP_UNROLLING_115(idx, step, macro) LOOP_UNROLLING_114(idx, step, macro); UNROLL_INCR(idx, step, macro) 2814#define LOOP_UNROLLING_116(idx, step, macro) LOOP_UNROLLING_115(idx, step, macro); UNROLL_INCR(idx, step, macro) 2815#define LOOP_UNROLLING_117(idx, step, macro) LOOP_UNROLLING_116(idx, step, macro); UNROLL_INCR(idx, step, macro) 2816#define LOOP_UNROLLING_118(idx, step, macro) LOOP_UNROLLING_117(idx, step, macro); UNROLL_INCR(idx, step, macro) 2817#define LOOP_UNROLLING_119(idx, step, macro) LOOP_UNROLLING_118(idx, step, macro); UNROLL_INCR(idx, step, macro) 2818#define LOOP_UNROLLING_120(idx, step, macro) LOOP_UNROLLING_119(idx, step, macro); UNROLL_INCR(idx, step, macro) 2819#define LOOP_UNROLLING_121(idx, step, macro) LOOP_UNROLLING_120(idx, step, macro); UNROLL_INCR(idx, step, macro) 2820#define LOOP_UNROLLING_122(idx, step, macro) LOOP_UNROLLING_121(idx, step, macro); UNROLL_INCR(idx, step, macro) 2821#define LOOP_UNROLLING_123(idx, step, macro) LOOP_UNROLLING_122(idx, step, macro); UNROLL_INCR(idx, step, macro) 2822#define LOOP_UNROLLING_124(idx, step, macro) LOOP_UNROLLING_123(idx, step, macro); UNROLL_INCR(idx, step, macro) 2823#define LOOP_UNROLLING_125(idx, step, macro) LOOP_UNROLLING_124(idx, step, macro); UNROLL_INCR(idx, step, macro) 2824#define LOOP_UNROLLING_126(idx, step, macro) LOOP_UNROLLING_125(idx, step, macro); UNROLL_INCR(idx, step, macro) 2825#define LOOP_UNROLLING_127(idx, step, macro) LOOP_UNROLLING_126(idx, step, macro); UNROLL_INCR(idx, step, macro) 2826#define LOOP_UNROLLING_128(idx, step, macro) LOOP_UNROLLING_127(idx, step, macro); UNROLL_INCR(idx, step, macro) 2827 2828#define LOOP_UNROLLING_STR(type, idx, start, step, num, macro) \ 2829 { \ 2830 type idx = start; \ 2831 LOOP_UNROLLING_##num(idx, step, macro); \ 2832 } 2833#else 2834#define LOOP_UNROLLING_STR(type, idx, start, step, num, macro) \ 2835 { \ 2836 _Pragma("unroll") \ 2837 for(type idx = start; idx < (num * step); idx += step) \ 2838 { \ 2839 (macro); \ 2840 } \ 2841 } 2842#endif 2843#define LOOP_UNROLLING(type, idx, start, step, num, macro) LOOP_UNROLLING_STR(type, idx, start, step, num, macro) 2844 2845 2846#define GET_SPATIAL_IDX(IDX, N0, PARTIAL_N0) (max((int)(get_global_id(IDX) * N0 - (N0 - PARTIAL_N0) % N0), 0)) 2847 2848 2849#define DOT_PRODUCT_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, b, c) DOT_PRODUCT_INTEGER8_STR(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, b, c) 2850#define DOT_PRODUCT_INTEGER8_STR(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, b, c) DOT_PRODUCT##K0##_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) 2851#define DOT_PRODUCT1_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2852 ({ \ 2853 c += (C_DATA_TYPE)(a) * (C_DATA_TYPE)(b); \ 2854 }) 2855#if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_khr_integer_dot_product) 2856#define DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += dot((A_DATA_TYPE##4)((a).s01, (A_DATA_TYPE##2)(0)), (B_DATA_TYPE##4)(((b).s01), (B_DATA_TYPE##2)(0))); 2857#define DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += dot((A_DATA_TYPE##4)((a).s012, (A_DATA_TYPE)0), (B_DATA_TYPE##4)(((b).s012), (B_DATA_TYPE)0)); 2858#define DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += dot((a), (b)); 2859#elif defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8) 2860#define DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c = arm_dot_acc((A_DATA_TYPE##4)((a).s01, (A_DATA_TYPE##2)(0)), (B_DATA_TYPE##4)(((b).s01), (B_DATA_TYPE##2)(0)), (c)); 2861#define DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c = arm_dot_acc((A_DATA_TYPE##4)((a).s012, (A_DATA_TYPE)0), (B_DATA_TYPE##4)(((b).s012), (B_DATA_TYPE)0), (c)); 2862#define DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c = arm_dot_acc((a), (b), (c)); 2863#elif defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) 2864#define DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += arm_dot((A_DATA_TYPE##4)((a).s01, (A_DATA_TYPE##2)(0)), (B_DATA_TYPE##4)(((b).s01), (B_DATA_TYPE##2)(0))); 2865#define DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += arm_dot((A_DATA_TYPE##4)((a).s012, (A_DATA_TYPE)0), (B_DATA_TYPE##4)(((b).s012), (B_DATA_TYPE)0)); 2866#define DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) c += arm_dot((a), (b)); 2867#else 2868#define DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2869 ({ \ 2870 c += (C_DATA_TYPE)(a).s0 * (C_DATA_TYPE)(b).s0; \ 2871 c += (C_DATA_TYPE)(a).s1 * (C_DATA_TYPE)(b).s1; \ 2872 }) 2873#define DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2874 ({ \ 2875 DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c); \ 2876 c += (C_DATA_TYPE)(a).s2 * (C_DATA_TYPE)(b).s2; \ 2877 }) 2878#define DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, x, y, val) \ 2879 ({ \ 2880 val += (C_DATA_TYPE)(x).s0 * (C_DATA_TYPE)(y).s0; \ 2881 val += (C_DATA_TYPE)(x).s1 * (C_DATA_TYPE)(y).s1; \ 2882 val += (C_DATA_TYPE)(x).s2 * (C_DATA_TYPE)(y).s2; \ 2883 val += (C_DATA_TYPE)(x).s3 * (C_DATA_TYPE)(y).s3; \ 2884 }) 2885#endif 2886#define DOT_PRODUCT5_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2887 ({ \ 2888 DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s0123), ((b).s0123), c); \ 2889 DOT_PRODUCT1_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s4), ((b).s4), c); \ 2890 }) 2891#define DOT_PRODUCT6_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2892 ({ \ 2893 DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s0123), ((b).s0123), c); \ 2894 DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s45), ((b).s45), c); \ 2895 }) 2896#define DOT_PRODUCT7_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2897 ({ \ 2898 DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s0123), ((b).s0123), c); \ 2899 DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s456), ((b).s456), c); \ 2900 }) 2901#define DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2902 ({ \ 2903 DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).lo), ((b).lo), c); \ 2904 DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).hi), ((b).hi), c); \ 2905 }) 2906#define DOT_PRODUCT9_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2907 ({ \ 2908 DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c); \ 2909 DOT_PRODUCT1_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s8), ((b).s8), c); \ 2910 }) 2911#define DOT_PRODUCT10_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2912 ({ \ 2913 DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c); \ 2914 DOT_PRODUCT2_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89), ((b).s89), c); \ 2915 }) 2916#define DOT_PRODUCT11_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2917 ({ \ 2918 DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c); \ 2919 DOT_PRODUCT3_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89A), ((b).s89A), c); \ 2920 }) 2921#define DOT_PRODUCT12_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2922 ({ \ 2923 DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c); \ 2924 DOT_PRODUCT4_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89AB), ((b).s89AB), c); \ 2925 }) 2926#define DOT_PRODUCT13_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2927 ({ \ 2928 DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c); \ 2929 DOT_PRODUCT5_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89ABC), ((b).s89ABC), c); \ 2930 }) 2931#define DOT_PRODUCT14_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2932 ({ \ 2933 DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c); \ 2934 DOT_PRODUCT6_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89ABCD), ((b).s89ABCD), c); \ 2935 }) 2936#define DOT_PRODUCT15_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2937 ({ \ 2938 DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s01234567), ((b).s01234567), c); \ 2939 DOT_PRODUCT7_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).s89ABCDE), ((b).s89ABCDE), c); \ 2940 }) 2941#define DOT_PRODUCT16_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, a, b, c) \ 2942 ({ \ 2943 DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).lo), ((b).lo), c); \ 2944 DOT_PRODUCT8_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, ((a).hi), ((b).hi), c); \ 2945 }) 2946 2947 2948#define REDUCE_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, c) REDUCE_INTEGER8_STR(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, c) 2949#define REDUCE_INTEGER8_STR(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, c) DOT_PRODUCT_INTEGER8(A_DATA_TYPE, B_DATA_TYPE, C_DATA_TYPE, K0, a, (TILE_VECTOR_TYPE##K0(B_DATA_TYPE))1, c) 2950 2951 2952#define V_LOAD(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y) V_LOAD_STR(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y) 2953#define V_LOAD_STR(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y) V_LOAD_##TENSOR_TYPE(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y) 2954#define V_LOAD_BUFFER(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y) \ 2955 VLOAD(WIDTH) \ 2956 (0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (Y) * (STRIDE_Y))) 2957#define V_LOAD_IMAGE(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y) READ_IMAGE2D(DATA_TYPE, CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(WIDTH), TENSOR##_img, (X) / 4, (Y)) 2958 2959 2960#define V_STORE(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y, VALUES) V_STORE_STR(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y, VALUES) 2961#define V_STORE_STR(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, Y, STRIDE_Y, VALUES) V_STORE_##TENSOR_TYPE(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y, VALUES) 2962#define V_STORE_BUFFER(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y, VALUES) \ 2963 VSTORE(WIDTH) \ 2964 (VALUES, 0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (Y) * (STRIDE_Y))) 2965#define V_STORE_IMAGE(DATA_TYPE, WIDTH, TENSOR, X, Y, STRIDE_Y, VALUES) WRITE_IMAGE2D(DATA_TYPE, CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(WIDTH), TENSOR##_img, (X) / 4, (Y), VALUES) 2966 2967 2968#define T_LOAD(DATA_TYPE, HEIGHT, WIDTH, TENSOR_TYPE, TENSOR, X, Y, YI_MULTIPLIER, STRIDE_Y, dst) \ 2969 ({ \ 2970 LOOP_UNROLLING(int, _i, 0, 1, HEIGHT, \ 2971 { \ 2972 dst[_i].v = V_LOAD(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, ((Y) + _i * (int)(YI_MULTIPLIER)), STRIDE_Y); \ 2973 }) \ 2974 }) 2975 2976 2977#define T_LOAD_INDIRECT(DATA_TYPE, HEIGHT, WIDTH, TENSOR_TYPE, TENSOR, X, STRIDE_Y, indirect_y, dst) \ 2978 ({ \ 2979 LOOP_UNROLLING(int, _i, 0, 1, HEIGHT, \ 2980 { \ 2981 dst[_i].v = V_LOAD(DATA_TYPE, WIDTH, TENSOR_TYPE, TENSOR, X, (indirect_y[_i].v), STRIDE_Y); \ 2982 }) \ 2983 }) 2984 2985 2986#define T_LOAD_INDIRECT_WIDTH_SELECT(DATA_TYPE, HEIGHT, WIDTH0, WIDTH1, TENSOR_TYPE, TENSOR, X, STRIDE_Y, WIDTH1_CONDITION, dst, indirect_y) \ 2987 ({ \ 2988 if(WIDTH1_CONDITION) \ 2989 { \ 2990 LOOP_UNROLLING(int, _i, 0, 1, HEIGHT, \ 2991 { \ 2992 VLOAD_PARTIAL(WIDTH0, WIDTH1) \ 2993 (dst[HEIGHT - 1 - _i].v, 0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (indirect_y[HEIGHT - 1 - _i].v) * STRIDE_Y)); \ 2994 }) \ 2995 } \ 2996 else \ 2997 { \ 2998 LOOP_UNROLLING(int, _i, 0, 1, HEIGHT, \ 2999 { \ 3000 dst[HEIGHT - 1 - _i].v = V_LOAD(DATA_TYPE, WIDTH0, TENSOR_TYPE, TENSOR, X, (indirect_y[HEIGHT - 1 - _i].v), STRIDE_Y); \ 3001 }) \ 3002 } \ 3003 }) 3004 3005#define T_LOAD_NHWC(DATA_TYPE, TILE_HEIGHT, TILE_WIDTH, TILE_CHANNELS, TENSOR_TYPE, TENSOR, B, Y, X, C, TENSOR_WIDTH, TENSOR_HEIGHT, STRIDE_Y, dst) \ 3006 ({ \ 3007 LOOP_UNROLLING(int, _yk, 0, 1, TILE_HEIGHT, \ 3008 { \ 3009 LOOP_UNROLLING(int, _xk, 0, 1, TILE_WIDTH, \ 3010 { \ 3011 int _src_y = (X) + _xk + ((Y) + _yk) * (TENSOR_WIDTH); \ 3012 _src_y += (B) * (int)(TENSOR_WIDTH) * (int)(TENSOR_HEIGHT); \ 3013 int _src_valid_y = (((X) + _xk) >= 0 && ((X) + _xk) < (int)(TENSOR_WIDTH) && ((Y) + _yk) >= 0 && ((Y) + _yk) < (int)(TENSOR_HEIGHT)); \ 3014 if(_src_valid_y != 0) \ 3015 { \ 3016 dst[_xk + _yk * (TILE_WIDTH)].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, _src_y, STRIDE_Y); \ 3017 } \ 3018 }) \ 3019 }) \ 3020 }) 3021 3022 3023#define T_LOAD_NHWC_WITH_DILATION(DATA_TYPE, TILE_HEIGHT, TILE_WIDTH, TILE_CHANNELS, TENSOR_TYPE, TENSOR, B, Y, X, C, TENSOR_WIDTH, TENSOR_HEIGHT, DILATION_X, DILATION_Y, BOUNDARY_CHECK, dst) \ 3024 ({ \ 3025 LOOP_UNROLLING(int, _yk, 0, 1, TILE_HEIGHT, \ 3026 { \ 3027 LOOP_UNROLLING(int, _xk, 0, 1, TILE_WIDTH, \ 3028 { \ 3029 int _src_y = (X) + _xk * (DILATION_X); \ 3030 int _src_z = ((Y) + _yk * (DILATION_Y)); \ 3031 int _src_w = (B); \ 3032 bool _src_valid_y = (((X) + _xk * (DILATION_X)) >= 0) && (((X) + _xk * (DILATION_X)) < (int)(TENSOR_WIDTH)) && (((Y) + _yk * (DILATION_Y)) >= 0) && (((Y) + _yk * (DILATION_Y)) < (int)(TENSOR_HEIGHT)); \ 3033 if(!(BOUNDARY_CHECK)) \ 3034 { \ 3035 dst[_xk + _yk * (TILE_WIDTH)].v = VLOAD(TILE_CHANNELS) \ 3036 (0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (C) * sizeof(DATA_TYPE) + (_src_y) * (TENSOR##_stride_y) + (_src_z) * (TENSOR##_stride_z) + (_src_w) * (TENSOR##_stride_w))); \ 3037 } \ 3038 else \ 3039 { \ 3040 if(_src_valid_y) \ 3041 { \ 3042 dst[_xk + _yk * (TILE_WIDTH)].v = VLOAD(TILE_CHANNELS) \ 3043 (0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (C) * sizeof(DATA_TYPE) + (_src_y) * (TENSOR##_stride_y) + (_src_z) * (TENSOR##_stride_z) + (_src_w) * (TENSOR##_stride_w))); \ 3044 } \ 3045 } \ 3046 }) \ 3047 }) \ 3048 }) 3049 3050 3051#define T_LOAD_NHWC_INDIRECT(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, B, Y, X, C, TENSOR_WIDTH, TENSOR_HEIGHT, STRIDE_Y, xi, yi, dst) \ 3052 ({ \ 3053 LOOP_UNROLLING(int, _i, 0, 1, TILE_AREA, \ 3054 { \ 3055 int _src_y = (X) + xi[_i].v + ((Y) + yi[_i].v) * (TENSOR_WIDTH); \ 3056 _src_y += (B) * (int)(TENSOR_WIDTH) * (int)(TENSOR_HEIGHT); \ 3057 int _src_valid_y = (((X) + xi[_i].v) >= 0 && ((X) + xi[_i].v) < (int)(TENSOR_WIDTH) && ((Y) + yi[_i].v) >= 0 && ((Y) + yi[_i].v) < (int)(TENSOR_HEIGHT)); \ 3058 if(_src_valid_y != 0) \ 3059 { \ 3060 dst[_i].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, _src_y, STRIDE_Y); \ 3061 } \ 3062 }) \ 3063 }) 3064 3065 3066#define T_LOAD2D_INDIRECT(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) T_LOAD2D_INDIRECT_STR(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) 3067#define T_LOAD2D_INDIRECT_STR(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) T_LOAD2D_INDIRECT_##TENSOR_TYPE(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) 3068#define T_LOAD2D_INDIRECT_BUFFER(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) \ 3069 ({ \ 3070 LOOP_UNROLLING(int, _i, 0, 1, TILE_AREA, \ 3071 { \ 3072 if(yi[0].s[_i] >= 0) \ 3073 { \ 3074 dst[_i].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, yi[0].s[_i], STRIDE_Y); \ 3075 } \ 3076 }) \ 3077 }) 3078 3079#define T_LOAD2D_INDIRECT_IMAGE(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, STRIDE_Y, yi, dst) \ 3080 ({ \ 3081 LOOP_UNROLLING(int, _i, 0, 1, TILE_AREA, \ 3082 { \ 3083 dst[_i].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, yi[0].s[_i], STRIDE_Y); \ 3084 }) \ 3085 }) 3086 3087 3088#define T_LOAD_NDHWC_INDIRECT(DATA_TYPE, TILE_AREA, TILE_CHANNELS, TENSOR_TYPE, TENSOR, B, Z, Y, X, C, TENSOR_WIDTH, TENSOR_HEIGHT, TENSOR_DEPTH, STRIDE_Y, xi, yi, zi, dst) \ 3089 ({ \ 3090 LOOP_UNROLLING(int, _i, 0, 1, TILE_AREA, \ 3091 { \ 3092 int _src_y = (X) + xi[_i].v + ((Y) + yi[_i].v) * (TENSOR_WIDTH) + ((Z) + zi[_i].v) * (TENSOR_WIDTH * TENSOR_HEIGHT); \ 3093 _src_y += (B) * (int)(TENSOR_WIDTH) * (int)(TENSOR_HEIGHT) * (int)(TENSOR_DEPTH); \ 3094 int _src_valid_y = (((X) + xi[_i].v) >= 0 && ((X) + xi[_i].v) < (int)(TENSOR_WIDTH) && ((Y) + yi[_i].v) >= 0 && ((Y) + yi[_i].v) < (int)(TENSOR_HEIGHT) \ 3095 && ((Z) + zi[_i].v) >= 0 && ((Z) + zi[_i].v) < (int)(TENSOR_DEPTH)); \ 3096 if(_src_valid_y != 0) \ 3097 { \ 3098 dst[_i].v = V_LOAD(DATA_TYPE, TILE_CHANNELS, TENSOR_TYPE, TENSOR, C, _src_y, STRIDE_Y); \ 3099 } \ 3100 }) \ 3101 }) 3102 3103 3104#define T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, HEIGHT, WIDTH0, WIDTH1, TENSOR_TYPE, TENSOR, X, STRIDE_Y, WIDTH1_CONDITION, src, indirect_y) \ 3105 ({ \ 3106 if(WIDTH1_CONDITION) \ 3107 { \ 3108 LOOP_UNROLLING(int, _i, 0, 1, HEIGHT, \ 3109 { \ 3110 VSTORE_PARTIAL(WIDTH0, WIDTH1) \ 3111 (CONVERT(src[HEIGHT - 1 - _i].v, VEC_DATA_TYPE(DATA_TYPE, WIDTH0)), 0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (indirect_y[HEIGHT - 1 - _i].v) * STRIDE_Y)); \ 3112 }) \ 3113 } \ 3114 else \ 3115 { \ 3116 LOOP_UNROLLING(int, _i, 0, 1, HEIGHT, \ 3117 { \ 3118 VSTORE(WIDTH0) \ 3119 (CONVERT(src[HEIGHT - 1 - _i].v, VEC_DATA_TYPE(DATA_TYPE, WIDTH0)), 0, (__global DATA_TYPE *)(TENSOR##_ptr + TENSOR##_offset_first_element_in_bytes + (X) * sizeof(DATA_TYPE) + (indirect_y[HEIGHT - 1 - _i].v) * STRIDE_Y)); \ 3120 }) \ 3121 } \ 3122 }) 3123 3124 3125#define T_OFFSET_CORRECTION(ACC_DATA_TYPE, M0, N0, K0, SRC_OFFSET, WEI_OFFSET, lhs, rhs, dst) \ 3126 ({ \ 3127 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3128 { \ 3129 ACC_DATA_TYPE _tm = 0; \ 3130 LOOP_UNROLLING(int, _k0, 0, 1, K0, \ 3131 { \ 3132 _tm += ((ACC_DATA_TYPE)lhs[_m0].s[_k0] * (ACC_DATA_TYPE)WEI_OFFSET); \ 3133 }) \ 3134 LOOP_UNROLLING(int, _n0, 0, 1, N0, \ 3135 { \ 3136 dst[_m0].s[_n0] += _tm; \ 3137 LOOP_UNROLLING(int, _k0, 0, 1, K0, \ 3138 { \ 3139 dst[_m0].s[_n0] += ((ACC_DATA_TYPE)rhs[_n0].s[_k0] * (ACC_DATA_TYPE)SRC_OFFSET); \ 3140 }) \ 3141 }) \ 3142 }) \ 3143 }) 3144 3145 3146#define T_QUANTIZE8(SRC_DATA_TYPE, DST_DATA_TYPE, QUANTIZATION_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst) T_QUANTIZE8_STR(SRC_DATA_TYPE, DST_DATA_TYPE, QUANTIZATION_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst) 3147#define T_QUANTIZE8_STR(SRC_DATA_TYPE, DST_DATA_TYPE, QUANTIZATION_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst) T_QUANTIZE8_##QUANTIZATION_TYPE(SRC_DATA_TYPE, DST_DATA_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst) 3148 3149 3150#define T_QUANTIZE8_PER_TENSOR(SRC_DATA_TYPE, DST_DATA_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst) \ 3151 ({ \ 3152 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3153 { \ 3154 LOOP_UNROLLING(int, _n0, 0, 1, N0, \ 3155 { \ 3156 SRC_DATA_TYPE _tmp = 0; \ 3157 SRC_DATA_TYPE _src = src[_m0].s[_n0]; \ 3158 _src *= select((SRC_DATA_TYPE)1, ((SRC_DATA_TYPE)1 << (SRC_DATA_TYPE)(-DST_SHIFT)), ((SRC_DATA_TYPE)DST_SHIFT < (SRC_DATA_TYPE)0)); \ 3159 SRC_DATA_TYPE overflow = _src == DST_MULTIPLIER && _src == INT_MIN; \ 3160 long a_64 = (long)(_src); \ 3161 long b_64 = (long)(DST_MULTIPLIER); \ 3162 long ab_64 = a_64 * b_64; \ 3163 long mask1 = 1 << 30; \ 3164 long mask2 = 1 - (1 << 30); \ 3165 long is_positive_or_zero = ab_64 >= 0; \ 3166 long nudge = select(mask2, mask1, is_positive_or_zero); \ 3167 SRC_DATA_TYPE ab_x2_high32 = CONVERT((ab_64 + nudge) / (long)(1ll << 31), SRC_DATA_TYPE); \ 3168 _tmp = select(ab_x2_high32, (SRC_DATA_TYPE)INT_MAX, overflow); \ 3169 if(DST_SHIFT >= 0) \ 3170 { \ 3171 long mask = ((((int)1) << DST_SHIFT) - (long)1); \ 3172 long threshold = _tmp < (int)0 ? (mask >> 1) + (long)1 : (mask >> 1) + 0; \ 3173 _tmp = (_tmp & mask) > threshold ? (_tmp >> DST_SHIFT) + (int)1 : (_tmp >> DST_SHIFT); \ 3174 } \ 3175 _tmp += DST_OFFSET; \ 3176 dst[_m0].s[_n0] = CONVERT_SAT(_tmp, DST_DATA_TYPE); \ 3177 }) \ 3178 }) \ 3179 }) 3180 3181 3182#define T_QUANTIZE8_PER_CHANNEL(SRC_DATA_TYPE, DST_DATA_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst_multipliers, dst_shifts, dst) \ 3183 ({ \ 3184 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3185 { \ 3186 LOOP_UNROLLING(int, _n0, 0, 1, N0, \ 3187 { \ 3188 SRC_DATA_TYPE _tmp = 0; \ 3189 SRC_DATA_TYPE _tmp2 = 0; \ 3190 SRC_DATA_TYPE _src = src[_m0].s[_n0]; \ 3191 SRC_DATA_TYPE _dst_multiplier = dst_multipliers[0].s[_n0]; \ 3192 SRC_DATA_TYPE _dst_shift = dst_shifts[0].s[_n0]; \ 3193 _src *= select((SRC_DATA_TYPE)1, ((SRC_DATA_TYPE)1 << (SRC_DATA_TYPE)(-_dst_shift)), ((SRC_DATA_TYPE)_dst_shift < (SRC_DATA_TYPE)0)); \ 3194 SRC_DATA_TYPE overflow = _src == _dst_multiplier && _src == INT_MIN; \ 3195 long a_64 = (long)(_src); \ 3196 long b_64 = (long)(_dst_multiplier); \ 3197 long ab_64 = a_64 * b_64; \ 3198 long mask1 = 1 << 30; \ 3199 long mask2 = 1 - (1 << 30); \ 3200 long is_positive_or_zero = ab_64 >= 0; \ 3201 long nudge = select(mask2, mask1, is_positive_or_zero); \ 3202 SRC_DATA_TYPE ab_x2_high32 = CONVERT((ab_64 + nudge) / (long)(1ll << 31), SRC_DATA_TYPE); \ 3203 _tmp = select(ab_x2_high32, (SRC_DATA_TYPE)INT_MAX, overflow); \ 3204 long mask = ((((int)1) << _dst_shift) - (int)1); \ 3205 long threshold = (mask >> 1) + any(_tmp); \ 3206 _tmp2 = _tmp >> _dst_shift; \ 3207 _tmp2 += select(0, 1, (_tmp & mask) > threshold); \ 3208 _tmp = select(_tmp, _tmp2, _dst_shift >= 0); \ 3209 _tmp += DST_OFFSET; \ 3210 dst[_m0].s[_n0] = CONVERT_SAT(_tmp, DST_DATA_TYPE); \ 3211 }) \ 3212 }) \ 3213 }) 3214 3215 3216#define T_QUANTIZE8_ASYMMETRIC(SRC_DATA_TYPE, DST_DATA_TYPE, M0, N0, DST_OFFSET, DST_SHIFT, DST_MULTIPLIER, src, dst) \ 3217 ({ \ 3218 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3219 { \ 3220 LOOP_UNROLLING(int, _n0, 0, 1, N0, \ 3221 { \ 3222 SRC_DATA_TYPE _tmp = 0; \ 3223 SRC_DATA_TYPE _src = src[_m0].s[_n0]; \ 3224 _src *= select((SRC_DATA_TYPE)1, ((SRC_DATA_TYPE)1 << (SRC_DATA_TYPE)(-DST_SHIFT)), ((SRC_DATA_TYPE)DST_SHIFT < (SRC_DATA_TYPE)0)); \ 3225 SRC_DATA_TYPE overflow = _src == DST_MULTIPLIER && _src == INT_MIN; \ 3226 long a_64 = (long)(_src); \ 3227 long b_64 = (long)(DST_MULTIPLIER); \ 3228 long ab_64 = a_64 * b_64; \ 3229 long mask1 = 1 << 30; \ 3230 long mask2 = 1 - (1 << 30); \ 3231 long is_positive_or_zero = ab_64 >= 0; \ 3232 long nudge = select(mask2, mask1, is_positive_or_zero); \ 3233 SRC_DATA_TYPE ab_x2_high32 = CONVERT((ab_64 + nudge) / (long)(1ll << 31), SRC_DATA_TYPE); \ 3234 _tmp = select(ab_x2_high32, (SRC_DATA_TYPE)INT_MAX, overflow); \ 3235 if(DST_SHIFT >= 0) \ 3236 { \ 3237 long mask = ((((int)1) << DST_SHIFT) - (int)1); \ 3238 long threshold = _tmp < (int)0 ? (mask >> 1) + (long)1 : (mask >> 1) + 0; \ 3239 _tmp = (_tmp & mask) > threshold ? (_tmp >> DST_SHIFT) + (int)1 : (_tmp >> DST_SHIFT); \ 3240 } \ 3241 _tmp += DST_OFFSET; \ 3242 dst[_m0].s[_n0] = CONVERT_SAT(_tmp, DST_DATA_TYPE); \ 3243 }) \ 3244 }) \ 3245 }) 3246 3247 3248#define T_ROWSET_MASK(DATA_TYPE, M0, N0, VALUE_TO_SET, a, mask) \ 3249 ({ \ 3250 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3251 { \ 3252 LOOP_UNROLLING(int, _n0, 0, 1, N0, \ 3253 { \ 3254 a[_m0].s[_n0] = select((DATA_TYPE)(a[_m0].s[_n0]), (DATA_TYPE)(VALUE_TO_SET), (SELECT_DATA_TYPE(DATA_TYPE))(mask[_m0].v == (DATA_TYPE)0)); \ 3255 }) \ 3256 }) \ 3257 }) 3258 3259 3260#define T_ACTIVATION(DATA_TYPE, M0, N0, ACTIVATION_TYPE, A_VAL, B_VAL, src, dst) \ 3261 ({ \ 3262 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3263 { \ 3264 dst[_m0].v = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, N0, src[_m0].v, A_VAL, B_VAL); \ 3265 }) \ 3266 }) 3267 3268 3269#define relu_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (max((DATA_TYPE)ZERO_VALUE, x)) 3270 3271#define brelu_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (min((DATA_TYPE)A_VAL, max((DATA_TYPE)ZERO_VALUE, x))) 3272 3273#define lu_brelu_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (min(max(x, (DATA_TYPE)B_VAL), (DATA_TYPE)A_VAL)) 3274 3275#define hard_swish_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (x * ((min(max((DATA_TYPE)(x + (DATA_TYPE)3.f), (DATA_TYPE)0.f), (DATA_TYPE)6.f)) * (DATA_TYPE)0.166666667f)) 3276 3277#define identity_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) (x) 3278 3279#define ACT_OP_QUANTIZED(op, DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) op##_op_quantized(DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) 3280#define ACTIVATION_QUANTIZED(op, DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) ACT_OP_QUANTIZED(op, DATA_TYPE, VEC_SIZE, ZERO_VALUE, A_VAL, B_VAL, x) 3281 3282#define V_ADD(A_VAL, B_VAL) ((A_VAL) + (B_VAL)) 3283#define V_SUB(A_VAL, B_VAL) ((A_VAL) - (B_VAL)) 3284#define V_DIV(A_VAL, B_VAL) ((A_VAL) / (B_VAL)) 3285#define V_MUL(A_VAL, B_VAL) ((A_VAL) * (B_VAL)) 3286 3287 3288#define T_ACTIVATION_QUANTIZED(DATA_TYPE, M0, N0, ACTIVATION_TYPE, ZERO_VALUE, A_VAL, B_VAL, src, dst) \ 3289 ({ \ 3290 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3291 { \ 3292 dst[_m0].v = ACTIVATION_QUANTIZED(ACTIVATION_TYPE, DATA_TYPE, N0, ZERO_VALUE, A_VAL, B_VAL, src[_m0].v); \ 3293 }) \ 3294 }) 3295 3296 3297#define T_ADD(DATA_TYPE, M0, N0, lhs, rhs, dst) \ 3298 ({ \ 3299 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3300 { \ 3301 dst[_m0].v = lhs[_m0].v + rhs[_m0].v; \ 3302 }) \ 3303 }) 3304 3305 3306#define T_ADD_CONSTANT(DATA_TYPE, M0, N0, lhs, rhs_constant, dst) \ 3307 ({ \ 3308 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3309 { \ 3310 dst[_m0].v = lhs[_m0].v + (DATA_TYPE)rhs_constant; \ 3311 }) \ 3312 }) 3313 3314#define T_ELTWISE_BROADCAST_ADD_X(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_ADD, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3315#define T_ELTWISE_BROADCAST_LHS_X_ADD(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_LHS_X(V_ADD, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3316#define T_ELTWISE_BROADCAST_RHS_X_ADD(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_ADD, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3317 3318#define T_ELTWISE_BROADCAST_LHS_X_SUB(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_LHS_X(V_SUB, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3319#define T_ELTWISE_BROADCAST_RHS_X_SUB(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_SUB, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3320 3321#define T_ELTWISE_BROADCAST_DIV_X(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_DIV, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3322 3323#define T_ELTWISE_BROADCAST_LHS_X_MUL(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_LHS_X(V_MUL, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3324#define T_ELTWISE_BROADCAST_RHS_X_MUL(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE_BROADCAST_X(V_MUL, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3325 3326 3327#define T_SCALE_CONSTANT(DATA_TYPE, M0, N0, lhs, rhs_constant, dst) \ 3328 ({ \ 3329 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3330 { \ 3331 dst[_m0].v = lhs[_m0].v * (DATA_TYPE)rhs_constant; \ 3332 }) \ 3333 }) 3334 3335 3336#define T_ELTWISE_BROADCAST_X(T_ELWISE_OP, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) \ 3337 ({ \ 3338 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3339 { \ 3340 dst[_m0].v = T_ELWISE_OP(CONVERT(lhs[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)), CONVERT(rhs[0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0))); \ 3341 }) \ 3342 }) 3343 3344 3345#define T_ELTWISE_BROADCAST_LHS_X(T_ELWISE_OP, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) \ 3346 ({ \ 3347 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3348 { \ 3349 dst[_m0].v = T_ELWISE_OP(CONVERT(lhs[0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)), CONVERT(rhs[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0))); \ 3350 }) \ 3351 }) 3352 3353#define T_ELTWISE_ADD(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE(V_ADD, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3354#define T_ELTWISE_SUB(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE(V_SUB, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3355#define T_ELTWISE_DIV(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE(V_DIV, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3356#define T_ELTWISE_MUL(DST_DATA_TYPE, M0, N0, lhs, rhs, dst) T_ELTWISE(V_MUL, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) 3357 3358 3359#define T_ELTWISE(T_ELWISE_OP, DST_DATA_TYPE, M0, N0, lhs, rhs, dst) \ 3360 ({ \ 3361 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3362 { \ 3363 dst[_m0].v = T_ELWISE_OP(CONVERT(lhs[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0)), CONVERT(rhs[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0))); \ 3364 }) \ 3365 }) 3366 3367 3368#define T_FLOOR(DST_DATA_TYPE, M0, N0, src, dst) \ 3369 ({ \ 3370 LOOP_UNROLLING(int, _m0, 0, 1, M0, \ 3371 { \ 3372 dst[_m0].v = floor(CONVERT(src[_m0].v, VEC_DATA_TYPE(DST_DATA_TYPE, N0))); \ 3373 }) \ 3374 }) 3375 3376 3377#define T_MMUL(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, LHS_LAYOUT, RHS_LAYOUT, lhs, rhs, dst) T_MMUL_##LHS_LAYOUT##_##RHS_LAYOUT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) 3378#define T_MMUL_NT_T(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_##LHS_DATA_TYPE##_##RHS_DATA_TYPE##_##DST_DATA_TYPE(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) 3379#define T_MMUL_NT_T_float_float_float(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_FLOAT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) 3380#define T_MMUL_NT_T_half_half_float(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_FLOAT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) 3381#define T_MMUL_NT_T_half_half_half(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_FLOAT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) 3382#define T_MMUL_NT_T_char_char_int(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) 3383#define T_MMUL_NT_T_uchar_uchar_uint(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) 3384#define T_MMUL_NT_T_uchar_uchar_int(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) T_MMUL_NT_T_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) 3385#define T_MMUL_NT_T_FLOAT(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) \ 3386 { \ 3387 LOOP_UNROLLING(int, _m, 0, 1, M0, \ 3388 { \ 3389 LOOP_UNROLLING(int, _n, 0, 1, N0, \ 3390 { \ 3391 LOOP_UNROLLING(int, _k, 0, 1, K0, \ 3392 { \ 3393 dst[_m].s[_n] = fma((DST_DATA_TYPE)(lhs[_m].s[_k]), (DST_DATA_TYPE)(rhs[_n].s[_k]), dst[_m].s[_n]); \ 3394 }) \ 3395 }) \ 3396 }) \ 3397 } 3398 3399#define T_MMUL_NT_T_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, M0, N0, K0, lhs, rhs, dst) \ 3400 ({ \ 3401 LOOP_UNROLLING(int, _m, 0, 1, M0, \ 3402 { \ 3403 LOOP_UNROLLING(int, _n, 0, 1, N0, \ 3404 { \ 3405 DOT_PRODUCT_INTEGER8(LHS_DATA_TYPE, RHS_DATA_TYPE, DST_DATA_TYPE, K0, (lhs[_m].v), (rhs[_n].v), dst[_m].s[_n]); \ 3406 }) \ 3407 }) \ 3408 }) 3409 3410#endif 3411 3412#if defined(NUM_TILES_X) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H) 3413#if defined(VEC_SIZE) && VEC_SIZE == 2 3414 3415__kernel void winograd_output_transform_2x2_3x3_nchw( 3416 TENSOR4D_DECLARATION(src), 3417 TENSOR4D_DECLARATION(dst) 3418#if defined(HAS_BIAS) 3419 , 3420 VECTOR_DECLARATION(bias) 3421#endif 3422) 3423{ 3424 3425#if defined(SRC_DEPTH) 3426 Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, SRC_DEPTH); 3427 const __global uchar *src_addr = tensor4D_offset(&src, 0, 0, 0, 0); 3428#else 3429 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); 3430 const __global uchar *src_addr = tensor3D_offset(&src, 0, 0, 0); 3431#endif 3432 3433 3434 DATA_TYPE d00 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_z)); 3435 DATA_TYPE d01 = *((__global DATA_TYPE *)(src_addr + 1 * src_stride_z)); 3436 DATA_TYPE d02 = *((__global DATA_TYPE *)(src_addr + 2 * src_stride_z)); 3437 DATA_TYPE d03 = *((__global DATA_TYPE *)(src_addr + 3 * src_stride_z)); 3438 3439#if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) 3440 3441 3442 3443 3444 float out00 = d00 + d01 + d02; 3445 float out01 = d01 - d02 - d03; 3446#else 3447 3448 DATA_TYPE d10 = *((__global DATA_TYPE *)(src_addr + 4 * src_stride_z)); 3449 DATA_TYPE d11 = *((__global DATA_TYPE *)(src_addr + 5 * src_stride_z)); 3450 DATA_TYPE d12 = *((__global DATA_TYPE *)(src_addr + 6 * src_stride_z)); 3451 DATA_TYPE d13 = *((__global DATA_TYPE *)(src_addr + 7 * src_stride_z)); 3452 3453 DATA_TYPE d20 = *((__global DATA_TYPE *)(src_addr + 8 * src_stride_z)); 3454 DATA_TYPE d21 = *((__global DATA_TYPE *)(src_addr + 9 * src_stride_z)); 3455 DATA_TYPE d22 = *((__global DATA_TYPE *)(src_addr + 10 * src_stride_z)); 3456 DATA_TYPE d23 = *((__global DATA_TYPE *)(src_addr + 11 * src_stride_z)); 3457 3458 DATA_TYPE d30 = *((__global DATA_TYPE *)(src_addr + 12 * src_stride_z)); 3459 DATA_TYPE d31 = *((__global DATA_TYPE *)(src_addr + 13 * src_stride_z)); 3460 DATA_TYPE d32 = *((__global DATA_TYPE *)(src_addr + 14 * src_stride_z)); 3461 DATA_TYPE d33 = *((__global DATA_TYPE *)(src_addr + 15 * src_stride_z)); 3462 3463 3464 float k0 = d01 + d11 + d21; 3465 float k1 = d02 + d12 + d22; 3466 float k2 = d11 - d21 - d31; 3467 float k3 = d12 - d22 - d32; 3468 3469 3470 3471 3472 3473 3474 float out00 = d10; 3475 float out01 = -d13; 3476 float out10 = d10; 3477 float out11 = -d13; 3478 3479 out00 += d00 + d20 + k0 + k1; 3480 out01 += k0 - k1 - (d03 + d23); 3481 out10 += -d20 - d30 + k2 + k3; 3482 out11 += k2 - k3 + d23 + d33; 3483#endif 3484 3485 int y_in = get_global_id(1); 3486 int x_out = (y_in % NUM_TILES_X) * OUTPUT_TILE_W; 3487 int y_out = (y_in / NUM_TILES_X) * OUTPUT_TILE_H; 3488 int z_out = get_global_id(0); 3489#if defined(SRC_DEPTH) 3490 int batch = get_global_id(2) / SRC_DEPTH; 3491#endif 3492 3493#if defined(HAS_BIAS) 3494 3495 Vector bias = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bias); 3496 3497 float b = (float) * ((__global DATA_TYPE *)(vector_offset(&bias, z_out))); 3498 3499 out00 += (float)b; 3500 out01 += (float)b; 3501#endif 3502 3503 3504#if defined(SRC_DEPTH) 3505 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w; 3506#else 3507 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z; 3508#endif 3509 3510 3511#if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) 3512 const VEC_DATA_TYPE(DATA_TYPE, 2) 3513 out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, CONVERT((VEC_DATA_TYPE(float, 2))(out00, out01), VEC_DATA_TYPE(DATA_TYPE, 2)), A_VAL, B_VAL); 3514 *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_y)) = out0_dt.s0; 3515 *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_y)) = out0_dt.s1; 3516#else 3517 vstore2(ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, CONVERT((VEC_DATA_TYPE(float, 2))(out00, out01), VEC_DATA_TYPE(DATA_TYPE, 2)), A_VAL, B_VAL), 0, 3518 (__global DATA_TYPE *)(dst_addr + 0 * dst_stride_y)); 3519#endif 3520 3521#if !defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) 3522#if defined(HAS_BIAS) 3523 3524 out10 += (DATA_TYPE)b; 3525 out11 += (DATA_TYPE)b; 3526#endif 3527 vstore2(ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, CONVERT((VEC_DATA_TYPE(float, 2))(out10, out11), VEC_DATA_TYPE(DATA_TYPE, 2)), A_VAL, B_VAL), 0, 3528 (__global DATA_TYPE *)(dst_addr + 1 * dst_stride_y)); 3529#endif 3530} 3531#endif 3532 3533#if defined(VEC_SIZE) && VEC_SIZE == 4 3534 3535__kernel void winograd_output_transform_4x4_3x3_nchw( 3536 TENSOR4D_DECLARATION(src), 3537 TENSOR4D_DECLARATION(dst) 3538#if defined(HAS_BIAS) 3539 , 3540 VECTOR_DECLARATION(bias) 3541#endif 3542) 3543{ 3544 3545#if defined(SRC_DEPTH) 3546 Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, SRC_DEPTH); 3547 const __global uchar *src_addr = tensor4D_offset(&src, 0, 0, 0, 0); 3548#else 3549 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); 3550 const __global uchar *src_addr = tensor3D_offset(&src, 0, 0, 0); 3551#endif 3552 3553 3554 DATA_TYPE d00 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_z)); 3555 DATA_TYPE d01 = *((__global DATA_TYPE *)(src_addr + 1 * src_stride_z)); 3556 DATA_TYPE d02 = *((__global DATA_TYPE *)(src_addr + 2 * src_stride_z)); 3557 DATA_TYPE d03 = *((__global DATA_TYPE *)(src_addr + 3 * src_stride_z)); 3558 DATA_TYPE d04 = *((__global DATA_TYPE *)(src_addr + 4 * src_stride_z)); 3559 DATA_TYPE d05 = *((__global DATA_TYPE *)(src_addr + 5 * src_stride_z)); 3560 3561#if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) 3562 3563 float out00 = d00 + d01 + d02 + d03 + d04; 3564 float out01 = d01 - d02 + 2.0f * d03 - 2.0f * d04; 3565 float out02 = d01 + d02 + 4.0f * d03 + 4.0f * d04; 3566 float out03 = d01 - d02 + 8.0f * d03 - 8.0f * d04 + d05; 3567#else 3568 3569 DATA_TYPE d10 = *((__global DATA_TYPE *)(src_addr + 6 * src_stride_z)); 3570 DATA_TYPE d11 = *((__global DATA_TYPE *)(src_addr + 7 * src_stride_z)); 3571 DATA_TYPE d12 = *((__global DATA_TYPE *)(src_addr + 8 * src_stride_z)); 3572 DATA_TYPE d13 = *((__global DATA_TYPE *)(src_addr + 9 * src_stride_z)); 3573 DATA_TYPE d14 = *((__global DATA_TYPE *)(src_addr + 10 * src_stride_z)); 3574 DATA_TYPE d15 = *((__global DATA_TYPE *)(src_addr + 11 * src_stride_z)); 3575 3576 DATA_TYPE d20 = *((__global DATA_TYPE *)(src_addr + 12 * src_stride_z)); 3577 DATA_TYPE d21 = *((__global DATA_TYPE *)(src_addr + 13 * src_stride_z)); 3578 DATA_TYPE d22 = *((__global DATA_TYPE *)(src_addr + 14 * src_stride_z)); 3579 DATA_TYPE d23 = *((__global DATA_TYPE *)(src_addr + 15 * src_stride_z)); 3580 DATA_TYPE d24 = *((__global DATA_TYPE *)(src_addr + 16 * src_stride_z)); 3581 DATA_TYPE d25 = *((__global DATA_TYPE *)(src_addr + 17 * src_stride_z)); 3582 3583 DATA_TYPE d30 = *((__global DATA_TYPE *)(src_addr + 18 * src_stride_z)); 3584 DATA_TYPE d31 = *((__global DATA_TYPE *)(src_addr + 19 * src_stride_z)); 3585 DATA_TYPE d32 = *((__global DATA_TYPE *)(src_addr + 20 * src_stride_z)); 3586 DATA_TYPE d33 = *((__global DATA_TYPE *)(src_addr + 21 * src_stride_z)); 3587 DATA_TYPE d34 = *((__global DATA_TYPE *)(src_addr + 22 * src_stride_z)); 3588 DATA_TYPE d35 = *((__global DATA_TYPE *)(src_addr + 23 * src_stride_z)); 3589 3590 DATA_TYPE d40 = *((__global DATA_TYPE *)(src_addr + 24 * src_stride_z)); 3591 DATA_TYPE d41 = *((__global DATA_TYPE *)(src_addr + 25 * src_stride_z)); 3592 DATA_TYPE d42 = *((__global DATA_TYPE *)(src_addr + 26 * src_stride_z)); 3593 DATA_TYPE d43 = *((__global DATA_TYPE *)(src_addr + 27 * src_stride_z)); 3594 DATA_TYPE d44 = *((__global DATA_TYPE *)(src_addr + 28 * src_stride_z)); 3595 DATA_TYPE d45 = *((__global DATA_TYPE *)(src_addr + 29 * src_stride_z)); 3596 3597 DATA_TYPE d50 = *((__global DATA_TYPE *)(src_addr + 30 * src_stride_z)); 3598 DATA_TYPE d51 = *((__global DATA_TYPE *)(src_addr + 31 * src_stride_z)); 3599 DATA_TYPE d52 = *((__global DATA_TYPE *)(src_addr + 32 * src_stride_z)); 3600 DATA_TYPE d53 = *((__global DATA_TYPE *)(src_addr + 33 * src_stride_z)); 3601 DATA_TYPE d54 = *((__global DATA_TYPE *)(src_addr + 34 * src_stride_z)); 3602 DATA_TYPE d55 = *((__global DATA_TYPE *)(src_addr + 35 * src_stride_z)); 3603 3604 3605 float out00 = (float)d01 + (float)d21 + (float)d41 + (float)d11 + (float)d31; 3606 float out01 = (float)d01 + (float)d21 + (float)d41 + (float)d11 + (float)d31; 3607 float out02 = (float)d01 + (float)d21 + (float)d41 + (float)d11 + (float)d31; 3608 float out03 = (float)d01 + d21 + (float)d41 + (float)d11 + (float)d31; 3609 3610 float k0 = d03 + d04 + d13 + d14 + d23 + d24 + d33 + d34 + d43 + d44; 3611 float k1 = 2.0f * d03 - 2.0f * d04 + 2.0f * d13 - 2.0f * d14 + 2.0f * d23 - 2.0f * d24 + 2.0f * d33 - 2.0f * d34 + 2.0f * d43 - 2.0f * d44; 3612 3613 out00 += k0 + d00 + d02 + d10 + d12 + d20 + d22 + d30 + d32 + d40 + d42; 3614 out01 += k1 - d02 - d12 - d22 - d32 - d42; 3615 out02 += 4.0f * k0 + d02 + d12 + d22 + d32 + d42; 3616 out03 += 4.0f * k1 - d02 - d12 - d22 - d32 - d42 + d05 + d15 + d25 + d35 + d45; 3617 3618 3619 float out10 = d11 - d21 + 2.0f * d31 - 2.0f * d41; 3620 float out11 = d11 - d21 + 2.0f * d31 - 2.0f * d41; 3621 float out12 = d11 - d21 + 2.0f * d31 - 2.0f * d41; 3622 float out13 = d11 - d21 + 2.0f * d31 - 2.0f * d41; 3623 3624 k0 = d13 + d14 - d23 - d24 + 2.0f * d33 + 2.0f * d34 - 2.0f * d43 - 2.0f * d44; 3625 k1 = 2.0f * d13 - 2.0f * d14 - 2.0f * d23 + 2.0f * d24 + 4.0f * d33 - 4.0f * d34 - 4.0f * d43 + 4.0f * d44; 3626 3627 out10 += k0 + d10 + d12 - d20 - d22 + 2.0f * d30 + 2.0f * d32 - 2.0f * d40 - 2.0f * d42; 3628 out11 += k1 - d12 + d22 - 2.0f * d32 + 2.0f * d42; 3629 out12 += 4.0f * k0 + d12 - d22 + 2.0f * d32 - 2.0f * d42; 3630 out13 += 4.0f * k1 - d12 + d15 + d22 - d25 - 2.0f * d32 + 2.0f * d35 + 2.0f * d42 - 2.0f * d45; 3631 3632 3633 float out20 = d11 + d21 + 4.0f * d31 + 4.0f * d41; 3634 float out21 = d11 + d21 + 4.0f * d31 + 4.0f * d41; 3635 float out22 = d11 + d21 + 4.0f * d31 + 4.0f * d41; 3636 float out23 = d11 + d21 + 4.0f * d31 + 4.0f * d41; 3637 3638 k0 = d13 + d14 + d23 + d24 + 4.0f * d33 + 4.0f * d34 + 4.0f * d43 + 4.0f * d44; 3639 k1 = 2.0f * d13 - 2.0f * d14 + 2.0f * d23 - 2.0f * d24 + 8.0f * d33 - 8.0f * d34 + 8.0f * d43 - 8.0f * d44; 3640 3641 out20 += k0 + d10 + d12 + d20 + d22 + 4.0f * d30 + 4.0f * d32 + 4.0f * d40 + 4.0f * d42; 3642 out21 += k1 - d12 - d22 - 4.0f * d32 - 4.0f * d42; 3643 out22 += 4.0f * k0 + d12 + d22 + 4.0f * d32 + 4.0f * d42; 3644 out23 += 4.0f * k1 - d12 + d15 - d22 + d25 - 4.0f * d32 + 4.0f * d35 - 4.0f * d42 + 4.0f * d45; 3645 3646 3647 float out30 = d11 - d21 + 8.0f * d31 - 8.0f * d41 + d51; 3648 float out31 = d11 - d21 + 8.0f * d31 - 8.0f * d41 + d51; 3649 float out32 = d11 - d21 + 8.0f * d31 - 8.0f * d41 + d51; 3650 float out33 = d11 - d21 + 8.0f * d31 - 8.0f * d41 + d51; 3651 3652 k0 = d13 + d14 - d23 - d24 + 8.0f * d33 + 8.0f * d34 - 8.0f * d43 - 8.0f * d44 + d53 + d54; 3653 k1 = 2.0f * d13 - 2.0f * d14 - 2.0f * d23 + 2.0f * d24 + 16.0f * d33 - 16.0f * d34 - 16.0f * d43 + 16.0f * d44 + 2.0f * d53 - 2.0f * d54; 3654 3655 out30 += k0 + d10 + d12 - d20 - d22 + 8.0f * d30 + 8.0f * d32 - 8.0f * d40 - 8.0f * d42 + d50 + d52; 3656 out31 += k1 - d12 + d22 - 8.0f * d32 + 8.0f * d42 - d52; 3657 out32 += 4.0f * k0 + d12 - d22 + 8.0f * d32 - 8.0f * d42 + d52; 3658 out33 += 4.0f * k1 - d12 + d15 + d22 - d25 - 8.0f * d32 + 8.0f * d35 + 8.0f * d42 - 8.0f * d45 - d52 + d55; 3659#endif 3660 3661 int y_in = get_global_id(1); 3662 int x_out = (y_in % NUM_TILES_X) * OUTPUT_TILE_W; 3663 int y_out = (y_in / NUM_TILES_X) * OUTPUT_TILE_H; 3664 int z_out = get_global_id(0); 3665#if defined(SRC_DEPTH) 3666 int batch = get_global_id(2) / SRC_DEPTH; 3667#endif 3668 3669#if defined(HAS_BIAS) 3670 3671 Vector bias = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bias); 3672 3673 float b = (float) * ((__global DATA_TYPE *)(vector_offset(&bias, z_out))); 3674 3675 out00 += (float)b; 3676 out01 += (float)b; 3677 out02 += (float)b; 3678 out03 += (float)b; 3679#endif 3680 3681 3682#if defined(SRC_DEPTH) 3683 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w; 3684#else 3685 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z; 3686#endif 3687 3688 3689 const VEC_DATA_TYPE(DATA_TYPE, 4) 3690 out0_dt = CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), A_VAL, B_VAL), VEC_DATA_TYPE(DATA_TYPE, 4)); 3691 3692#if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) 3693 *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_y)) = out0_dt.s0; 3694 *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_y)) = out0_dt.s1; 3695 *((__global DATA_TYPE *)(dst_addr + 2 * dst_stride_y)) = out0_dt.s2; 3696 *((__global DATA_TYPE *)(dst_addr + 3 * dst_stride_y)) = out0_dt.s3; 3697#else 3698 vstore4(out0_dt, 0, (__global DATA_TYPE *)(dst_addr + 0 * dst_stride_y)); 3699#endif 3700 3701#if !defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) 3702#if defined(HAS_BIAS) 3703 3704 out10 += (float)b; 3705 out11 += (float)b; 3706 out12 += (float)b; 3707 out13 += (float)b; 3708 3709 out20 += (float)b; 3710 out21 += (float)b; 3711 out22 += (float)b; 3712 out23 += (float)b; 3713 3714 out30 += (float)b; 3715 out31 += (float)b; 3716 out32 += (float)b; 3717 out33 += (float)b; 3718#endif 3719 vstore4(CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out10, out11, out12, out13), A_VAL, B_VAL), VEC_DATA_TYPE(DATA_TYPE, 4)), 0, 3720 (__global DATA_TYPE *)(dst_addr + 1 * dst_stride_y)); 3721 vstore4(CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out20, out21, out22, out23), A_VAL, B_VAL), VEC_DATA_TYPE(DATA_TYPE, 4)), 0, 3722 (__global DATA_TYPE *)(dst_addr + 2 * dst_stride_y)); 3723 vstore4(CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out30, out31, out32, out33), A_VAL, B_VAL), VEC_DATA_TYPE(DATA_TYPE, 4)), 0, 3724 (__global DATA_TYPE *)(dst_addr + 3 * dst_stride_y)); 3725#endif 3726} 3727 3728#define COMPUTE_TMP_COL(col, d0, d1, d2, d3, d4, d5, d6, d7, comm_fact) \ 3729 ({ \ 3730 comm_fact.s0 = d1 + d2; \ 3731 comm_fact.s1 = d3 + d4; \ 3732 comm_fact.s2 = d5 + d6; \ 3733 \ 3734 col.s0 = comm_fact.s0 + comm_fact.s1 + 8.f * comm_fact.s2 + d0; \ 3735 col.s2 = comm_fact.s0 + 4.f * comm_fact.s1 + 2.f * comm_fact.s2; \ 3736 \ 3737 comm_fact.s0 = d1 - d2; \ 3738 comm_fact.s1 = d3 - d4; \ 3739 comm_fact.s2 = d5 - d6; \ 3740 \ 3741 col.s1 = comm_fact.s0 + 2.f * comm_fact.s1 + 4.f * comm_fact.s2; \ 3742 col.s3 = comm_fact.s0 + 8.f * comm_fact.s1 + comm_fact.s2 + d7; \ 3743 }) 3744 3745 3746__kernel void winograd_output_transform_4x4_5x5_nchw( 3747 TENSOR4D_DECLARATION(src), 3748 TENSOR4D_DECLARATION(dst) 3749#if defined(HAS_BIAS) 3750 , 3751 VECTOR_DECLARATION(bias) 3752#endif 3753) 3754{ 3755 3756#if defined(SRC_DEPTH) 3757 Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, SRC_DEPTH); 3758 const __global uchar *src_addr = tensor4D_offset(&src, 0, 0, 0, 0); 3759#else 3760 3761 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); 3762 const __global uchar *src_addr = tensor3D_offset(&src, 0, 0, 0); 3763#endif 3764 3765 3766 int y_in = get_global_id(1); 3767 int x_out = (y_in % NUM_TILES_X) * OUTPUT_TILE_W; 3768 int y_out = (y_in / NUM_TILES_X) * OUTPUT_TILE_H; 3769 int z_out = get_global_id(0); 3770#if defined(SRC_DEPTH) 3771 int batch = get_global_id(2) / SRC_DEPTH; 3772#endif 3773 3774#if defined(SRC_DEPTH) 3775 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z + batch * dst_stride_w; 3776#else 3777 3778 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z; 3779#endif 3780 3781 3782 DATA_TYPE d00 = *((__global DATA_TYPE *)(src_addr + 0 * src_stride_z)); 3783 DATA_TYPE d01 = *((__global DATA_TYPE *)(src_addr + 1 * src_stride_z)); 3784 DATA_TYPE d02 = *((__global DATA_TYPE *)(src_addr + 2 * src_stride_z)); 3785 DATA_TYPE d03 = *((__global DATA_TYPE *)(src_addr + 3 * src_stride_z)); 3786 DATA_TYPE d04 = *((__global DATA_TYPE *)(src_addr + 4 * src_stride_z)); 3787 DATA_TYPE d05 = *((__global DATA_TYPE *)(src_addr + 5 * src_stride_z)); 3788 DATA_TYPE d06 = *((__global DATA_TYPE *)(src_addr + 6 * src_stride_z)); 3789 DATA_TYPE d07 = *((__global DATA_TYPE *)(src_addr + 7 * src_stride_z)); 3790 3791#if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) 3792 3793 float out00 = d00 + d01 + d02 + d03 + d04 + 8.0f * d05 + 8.0f * d06; 3794 float out01 = d01 - d02 + 2.0f * d03 - 2.0f * d04 + 4.0f * d05 - 4.0f * d06; 3795 float out02 = d01 + d02 + 4.0f * d03 + 4.0f * d04 + 2.0f * d05 + 2.0f * d06; 3796 float out03 = d01 - d02 + 8.0f * d03 - 8.0f * d04 + d05 - d06 + d07; 3797 3798#if defined(HAS_BIAS) 3799 3800 Vector bias = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bias); 3801 3802 float b = (float) * ((__global DATA_TYPE *)(vector_offset(&bias, z_out))); 3803 3804 out00 += (DATA_TYPE)b; 3805 out01 += (DATA_TYPE)b; 3806 out02 += (DATA_TYPE)b; 3807 out03 += (DATA_TYPE)b; 3808#endif 3809 3810 3811#if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) 3812 VEC_DATA_TYPE(DATA_TYPE, 4) 3813 out0_dt = CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), A_VAL, 3814 B_VAL), 3815 VEC_DATA_TYPE(DATA_TYPE, 4)); 3816 *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_y)) = out0_dt.s0; 3817 *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_y)) = out0_dt.s1; 3818 *((__global DATA_TYPE *)(dst_addr + 2 * dst_stride_y)) = out0_dt.s2; 3819 *((__global DATA_TYPE *)(dst_addr + 3 * dst_stride_y)) = out0_dt.s3; 3820#else 3821 vstore4(CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), A_VAL, B_VAL), VEC_DATA_TYPE(DATA_TYPE, 4)), 3822 0, (__global DATA_TYPE *)(dst_addr)); 3823#endif 3824 3825#else 3826 3827 DATA_TYPE d10 = *((__global DATA_TYPE *)(src_addr + 8 * src_stride_z)); 3828 DATA_TYPE d11 = *((__global DATA_TYPE *)(src_addr + 9 * src_stride_z)); 3829 DATA_TYPE d12 = *((__global DATA_TYPE *)(src_addr + 10 * src_stride_z)); 3830 DATA_TYPE d13 = *((__global DATA_TYPE *)(src_addr + 11 * src_stride_z)); 3831 DATA_TYPE d14 = *((__global DATA_TYPE *)(src_addr + 12 * src_stride_z)); 3832 DATA_TYPE d15 = *((__global DATA_TYPE *)(src_addr + 13 * src_stride_z)); 3833 DATA_TYPE d16 = *((__global DATA_TYPE *)(src_addr + 14 * src_stride_z)); 3834 DATA_TYPE d17 = *((__global DATA_TYPE *)(src_addr + 15 * src_stride_z)); 3835 3836 DATA_TYPE d20 = *((__global DATA_TYPE *)(src_addr + 16 * src_stride_z)); 3837 DATA_TYPE d21 = *((__global DATA_TYPE *)(src_addr + 17 * src_stride_z)); 3838 DATA_TYPE d22 = *((__global DATA_TYPE *)(src_addr + 18 * src_stride_z)); 3839 DATA_TYPE d23 = *((__global DATA_TYPE *)(src_addr + 19 * src_stride_z)); 3840 DATA_TYPE d24 = *((__global DATA_TYPE *)(src_addr + 20 * src_stride_z)); 3841 DATA_TYPE d25 = *((__global DATA_TYPE *)(src_addr + 21 * src_stride_z)); 3842 DATA_TYPE d26 = *((__global DATA_TYPE *)(src_addr + 22 * src_stride_z)); 3843 DATA_TYPE d27 = *((__global DATA_TYPE *)(src_addr + 23 * src_stride_z)); 3844 3845 DATA_TYPE d30 = *((__global DATA_TYPE *)(src_addr + 24 * src_stride_z)); 3846 DATA_TYPE d31 = *((__global DATA_TYPE *)(src_addr + 25 * src_stride_z)); 3847 DATA_TYPE d32 = *((__global DATA_TYPE *)(src_addr + 26 * src_stride_z)); 3848 DATA_TYPE d33 = *((__global DATA_TYPE *)(src_addr + 27 * src_stride_z)); 3849 DATA_TYPE d34 = *((__global DATA_TYPE *)(src_addr + 28 * src_stride_z)); 3850 DATA_TYPE d35 = *((__global DATA_TYPE *)(src_addr + 29 * src_stride_z)); 3851 DATA_TYPE d36 = *((__global DATA_TYPE *)(src_addr + 30 * src_stride_z)); 3852 DATA_TYPE d37 = *((__global DATA_TYPE *)(src_addr + 31 * src_stride_z)); 3853 3854 DATA_TYPE d40 = *((__global DATA_TYPE *)(src_addr + 32 * src_stride_z)); 3855 DATA_TYPE d41 = *((__global DATA_TYPE *)(src_addr + 33 * src_stride_z)); 3856 DATA_TYPE d42 = *((__global DATA_TYPE *)(src_addr + 34 * src_stride_z)); 3857 DATA_TYPE d43 = *((__global DATA_TYPE *)(src_addr + 35 * src_stride_z)); 3858 DATA_TYPE d44 = *((__global DATA_TYPE *)(src_addr + 36 * src_stride_z)); 3859 DATA_TYPE d45 = *((__global DATA_TYPE *)(src_addr + 37 * src_stride_z)); 3860 DATA_TYPE d46 = *((__global DATA_TYPE *)(src_addr + 38 * src_stride_z)); 3861 DATA_TYPE d47 = *((__global DATA_TYPE *)(src_addr + 39 * src_stride_z)); 3862 3863 DATA_TYPE d50 = *((__global DATA_TYPE *)(src_addr + 40 * src_stride_z)); 3864 DATA_TYPE d51 = *((__global DATA_TYPE *)(src_addr + 41 * src_stride_z)); 3865 DATA_TYPE d52 = *((__global DATA_TYPE *)(src_addr + 42 * src_stride_z)); 3866 DATA_TYPE d53 = *((__global DATA_TYPE *)(src_addr + 43 * src_stride_z)); 3867 DATA_TYPE d54 = *((__global DATA_TYPE *)(src_addr + 44 * src_stride_z)); 3868 DATA_TYPE d55 = *((__global DATA_TYPE *)(src_addr + 45 * src_stride_z)); 3869 DATA_TYPE d56 = *((__global DATA_TYPE *)(src_addr + 46 * src_stride_z)); 3870 DATA_TYPE d57 = *((__global DATA_TYPE *)(src_addr + 47 * src_stride_z)); 3871 3872 DATA_TYPE d60 = *((__global DATA_TYPE *)(src_addr + 48 * src_stride_z)); 3873 DATA_TYPE d61 = *((__global DATA_TYPE *)(src_addr + 49 * src_stride_z)); 3874 DATA_TYPE d62 = *((__global DATA_TYPE *)(src_addr + 50 * src_stride_z)); 3875 DATA_TYPE d63 = *((__global DATA_TYPE *)(src_addr + 51 * src_stride_z)); 3876 DATA_TYPE d64 = *((__global DATA_TYPE *)(src_addr + 52 * src_stride_z)); 3877 DATA_TYPE d65 = *((__global DATA_TYPE *)(src_addr + 53 * src_stride_z)); 3878 DATA_TYPE d66 = *((__global DATA_TYPE *)(src_addr + 54 * src_stride_z)); 3879 DATA_TYPE d67 = *((__global DATA_TYPE *)(src_addr + 55 * src_stride_z)); 3880 3881 DATA_TYPE d70 = *((__global DATA_TYPE *)(src_addr + 56 * src_stride_z)); 3882 DATA_TYPE d71 = *((__global DATA_TYPE *)(src_addr + 57 * src_stride_z)); 3883 DATA_TYPE d72 = *((__global DATA_TYPE *)(src_addr + 58 * src_stride_z)); 3884 DATA_TYPE d73 = *((__global DATA_TYPE *)(src_addr + 59 * src_stride_z)); 3885 DATA_TYPE d74 = *((__global DATA_TYPE *)(src_addr + 60 * src_stride_z)); 3886 DATA_TYPE d75 = *((__global DATA_TYPE *)(src_addr + 61 * src_stride_z)); 3887 DATA_TYPE d76 = *((__global DATA_TYPE *)(src_addr + 62 * src_stride_z)); 3888 DATA_TYPE d77 = *((__global DATA_TYPE *)(src_addr + 63 * src_stride_z)); 3889 3890 3891 VEC_DATA_TYPE(float, 4) 3892 comm_fact0, comm_fact1, comm_fact2; 3893 VEC_DATA_TYPE(float, 4) 3894 tmp_col0, tmp_col1, tmp_col2, tmp_col3, tmp_col4, tmp_col5, tmp_col6, tmp_col7; 3895 3896 COMPUTE_TMP_COL(tmp_col0, d00, d10, d20, d30, d40, d50, d60, d70, comm_fact0); 3897 COMPUTE_TMP_COL(tmp_col1, d01, d11, d21, d31, d41, d51, d61, d71, comm_fact0); 3898 COMPUTE_TMP_COL(tmp_col2, d02, d12, d22, d32, d42, d52, d62, d72, comm_fact0); 3899 COMPUTE_TMP_COL(tmp_col3, d03, d13, d23, d33, d43, d53, d63, d73, comm_fact0); 3900 COMPUTE_TMP_COL(tmp_col4, d04, d14, d24, d34, d44, d54, d64, d74, comm_fact0); 3901 COMPUTE_TMP_COL(tmp_col5, d05, d15, d25, d35, d45, d55, d65, d75, comm_fact0); 3902 COMPUTE_TMP_COL(tmp_col6, d06, d16, d26, d36, d46, d56, d66, d76, comm_fact0); 3903 COMPUTE_TMP_COL(tmp_col7, d07, d17, d27, d37, d47, d57, d67, d77, comm_fact0); 3904 3905 3906 comm_fact0 = tmp_col1 + tmp_col2; 3907 comm_fact1 = tmp_col3 + tmp_col4; 3908 comm_fact2 = tmp_col5 + tmp_col6; 3909 3910 VEC_DATA_TYPE(float, 4) 3911 out_col0 = comm_fact0 + comm_fact1 + (float)8.f * comm_fact2 + tmp_col0; 3912 VEC_DATA_TYPE(float, 4) 3913 out_col2 = comm_fact0 + (float)4.f * comm_fact1 + (float)2.f * comm_fact2; 3914 3915 comm_fact0 = tmp_col1 - tmp_col2; 3916 comm_fact1 = tmp_col3 - tmp_col4; 3917 comm_fact2 = tmp_col5 - tmp_col6; 3918 3919 VEC_DATA_TYPE(float, 4) 3920 out_col1 = comm_fact0 + (float)2.f * comm_fact1 + (float)4.f * comm_fact2; 3921 VEC_DATA_TYPE(float, 4) 3922 out_col3 = comm_fact0 + (float)8.f * comm_fact1 + comm_fact2 + tmp_col7; 3923 3924#if defined(HAS_BIAS) 3925 3926 Vector bias = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bias); 3927 3928 float b = (float) * ((__global DATA_TYPE *)(vector_offset(&bias, z_out))); 3929 3930 out_col0 += (VEC_DATA_TYPE(float, 4))b; 3931 out_col1 += (VEC_DATA_TYPE(float, 4))b; 3932 out_col2 += (VEC_DATA_TYPE(float, 4))b; 3933 out_col3 += (VEC_DATA_TYPE(float, 4))b; 3934#endif 3935 3936 3937 vstore4(CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out_col0.s0, out_col1.s0, out_col2.s0, out_col3.s0), A_VAL, B_VAL), 3938 VEC_DATA_TYPE(DATA_TYPE, 4)), 3939 0, (__global DATA_TYPE *)(dst_addr + 0 * dst_stride_y)); 3940 vstore4(CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out_col0.s1, out_col1.s1, out_col2.s1, out_col3.s1), A_VAL, B_VAL), 3941 VEC_DATA_TYPE(DATA_TYPE, 4)), 3942 0, (__global DATA_TYPE *)(dst_addr + 1 * dst_stride_y)); 3943 vstore4(CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out_col0.s2, out_col1.s2, out_col2.s2, out_col3.s2), A_VAL, B_VAL), 3944 VEC_DATA_TYPE(DATA_TYPE, 4)), 3945 0, (__global DATA_TYPE *)(dst_addr + 2 * dst_stride_y)); 3946 vstore4(CONVERT(ACTIVATION(ACTIVATION_TYPE, float, VEC_SIZE, (VEC_DATA_TYPE(float, 4))(out_col0.s3, out_col1.s3, out_col2.s3, out_col3.s3), A_VAL, B_VAL), 3947 VEC_DATA_TYPE(DATA_TYPE, 4)), 3948 0, (__global DATA_TYPE *)(dst_addr + 3 * dst_stride_y)); 3949#endif 3950} 3951#endif 3952 3953#if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) 3954#if defined(VEC_SIZE) && VEC_SIZE == 2 3955 3956__kernel void winograd_output_transform_2x1_3x1_nchw( 3957 TENSOR4D_DECLARATION(src), 3958 TENSOR4D_DECLARATION(dst) 3959#if defined(HAS_BIAS) 3960 , 3961 VECTOR_DECLARATION(bias) 3962#endif 3963) 3964{ 3965 winograd_output_transform_2x2_3x3_nchw(src_ptr, 3966 src_stride_x, 3967 src_step_x, 3968 src_stride_y, 3969 src_step_y, 3970 src_stride_z, 3971 src_step_z, 3972 src_stride_w, 3973 src_step_w, 3974 src_offset_first_element_in_bytes, 3975 dst_ptr, 3976 dst_stride_x, 3977 dst_step_x, 3978 dst_stride_y, 3979 dst_step_y, 3980 dst_stride_z, 3981 dst_step_z, 3982 dst_stride_w, 3983 dst_step_w, 3984 dst_offset_first_element_in_bytes 3985#if defined(HAS_BIAS) 3986 , 3987 bias_ptr, 3988 bias_stride_x, 3989 bias_step_x, 3990 bias_offset_first_element_in_bytes 3991#endif 3992 ); 3993} 3994 3995#endif 3996 3997#if defined(VEC_SIZE) && VEC_SIZE == 4 3998 3999__kernel void winograd_output_transform_4x1_3x1_nchw( 4000 TENSOR4D_DECLARATION(src), 4001 TENSOR4D_DECLARATION(dst) 4002#if defined(HAS_BIAS) 4003 , 4004 VECTOR_DECLARATION(bias) 4005#endif 4006) 4007{ 4008 winograd_output_transform_4x4_3x3_nchw(src_ptr, 4009 src_stride_x, 4010 src_step_x, 4011 src_stride_y, 4012 src_step_y, 4013 src_stride_z, 4014 src_step_z, 4015 src_stride_w, 4016 src_step_w, 4017 src_offset_first_element_in_bytes, 4018 dst_ptr, 4019 dst_stride_x, 4020 dst_step_x, 4021 dst_stride_y, 4022 dst_step_y, 4023 dst_stride_z, 4024 dst_step_z, 4025 dst_stride_w, 4026 dst_step_w, 4027 dst_offset_first_element_in_bytes 4028#if defined(HAS_BIAS) 4029 , 4030 bias_ptr, 4031 bias_stride_x, 4032 bias_step_x, 4033 bias_offset_first_element_in_bytes 4034#endif 4035 ); 4036} 4037 4038 4039__kernel void winograd_output_transform_4x1_5x1_nchw( 4040 TENSOR4D_DECLARATION(src), 4041 TENSOR4D_DECLARATION(dst) 4042#if defined(HAS_BIAS) 4043 , 4044 VECTOR_DECLARATION(bias) 4045#endif 4046) 4047{ 4048 winograd_output_transform_4x4_5x5_nchw(src_ptr, 4049 src_stride_x, 4050 src_step_x, 4051 src_stride_y, 4052 src_step_y, 4053 src_stride_z, 4054 src_step_z, 4055 src_stride_w, 4056 src_step_w, 4057 src_offset_first_element_in_bytes, 4058 dst_ptr, 4059 dst_stride_x, 4060 dst_step_x, 4061 dst_stride_y, 4062 dst_step_y, 4063 dst_stride_z, 4064 dst_step_z, 4065 dst_stride_w, 4066 dst_step_w, 4067 dst_offset_first_element_in_bytes 4068#if defined(HAS_BIAS) 4069 , 4070 bias_ptr, 4071 bias_stride_x, 4072 bias_step_x, 4073 bias_offset_first_element_in_bytes 4074#endif 4075 ); 4076} 4077 4078#endif 4079#endif 4080 4081#if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) 4082#if defined(VEC_SIZE) && VEC_SIZE == 2 4083 4084__kernel void winograd_output_transform_1x2_1x3_nchw( 4085 TENSOR4D_DECLARATION(src), 4086 TENSOR4D_DECLARATION(dst) 4087#if defined(HAS_BIAS) 4088 , 4089 VECTOR_DECLARATION(bias) 4090#endif 4091) 4092{ 4093 winograd_output_transform_2x2_3x3_nchw(src_ptr, 4094 src_stride_x, 4095 src_step_x, 4096 src_stride_y, 4097 src_step_y, 4098 src_stride_z, 4099 src_step_z, 4100 src_stride_w, 4101 src_step_w, 4102 src_offset_first_element_in_bytes, 4103 dst_ptr, 4104 dst_stride_x, 4105 dst_step_x, 4106 dst_stride_y, 4107 dst_step_y, 4108 dst_stride_z, 4109 dst_step_z, 4110 dst_stride_w, 4111 dst_step_w, 4112 dst_offset_first_element_in_bytes 4113#if defined(HAS_BIAS) 4114 , 4115 bias_ptr, 4116 bias_stride_x, 4117 bias_step_x, 4118 bias_offset_first_element_in_bytes 4119#endif 4120 ); 4121} 4122 4123#endif 4124 4125#if defined(VEC_SIZE) && VEC_SIZE == 4 4126 4127__kernel void winograd_output_transform_1x4_1x3_nchw( 4128 TENSOR4D_DECLARATION(src), 4129 TENSOR4D_DECLARATION(dst) 4130#if defined(HAS_BIAS) 4131 , 4132 VECTOR_DECLARATION(bias) 4133#endif 4134) 4135{ 4136 winograd_output_transform_4x4_3x3_nchw(src_ptr, 4137 src_stride_x, 4138 src_step_x, 4139 src_stride_y, 4140 src_step_y, 4141 src_stride_z, 4142 src_step_z, 4143 src_stride_w, 4144 src_step_w, 4145 src_offset_first_element_in_bytes, 4146 dst_ptr, 4147 dst_stride_x, 4148 dst_step_x, 4149 dst_stride_y, 4150 dst_step_y, 4151 dst_stride_z, 4152 dst_step_z, 4153 dst_stride_w, 4154 dst_step_w, 4155 dst_offset_first_element_in_bytes 4156#if defined(HAS_BIAS) 4157 , 4158 bias_ptr, 4159 bias_stride_x, 4160 bias_step_x, 4161 bias_offset_first_element_in_bytes 4162#endif 4163 ); 4164} 4165 4166 4167__kernel void winograd_output_transform_1x4_1x5_nchw( 4168 TENSOR4D_DECLARATION(src), 4169 TENSOR4D_DECLARATION(dst) 4170#if defined(HAS_BIAS) 4171 , 4172 VECTOR_DECLARATION(bias) 4173#endif 4174) 4175{ 4176 winograd_output_transform_4x4_5x5_nchw(src_ptr, 4177 src_stride_x, 4178 src_step_x, 4179 src_stride_y, 4180 src_step_y, 4181 src_stride_z, 4182 src_step_z, 4183 src_stride_w, 4184 src_step_w, 4185 src_offset_first_element_in_bytes, 4186 dst_ptr, 4187 dst_stride_x, 4188 dst_step_x, 4189 dst_stride_y, 4190 dst_step_y, 4191 dst_stride_z, 4192 dst_step_z, 4193 dst_stride_w, 4194 dst_step_w, 4195 dst_offset_first_element_in_bytes 4196#if defined(HAS_BIAS) 4197 , 4198 bias_ptr, 4199 bias_stride_x, 4200 bias_step_x, 4201 bias_offset_first_element_in_bytes 4202#endif 4203 ); 4204} 4205 4206#endif 4207#endif 4208#endif )"