1 /*
2 * Copyright (c) 2016-2020 Arm Limited.
3 *
4 * SPDX-License-Identifier: MIT
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a copy
7 * of this software and associated documentation files (the "Software"), to
8 * deal in the Software without restriction, including without limitation the
9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10 * sell copies of the Software, and to permit persons to whom the Software is
11 * furnished to do so, subject to the following conditions:
12 *
13 * The above copyright notice and this permission notice shall be included in all
14 * copies or substantial portions of the Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22 * SOFTWARE.
23 */
24 #ifndef ARM_COMPUTE_HELPER_H
25 #define ARM_COMPUTE_HELPER_H
26
27 #include "load_store_utility.h"
28
29 #if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
30 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
31 #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
32
33 #if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8)
34 #pragma OPENCL EXTENSION cl_arm_integer_dot_product_int8 : enable
35 #endif // defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8)
36
37 #if defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8)
38 #pragma OPENCL EXTENSION cl_arm_integer_dot_product_accumulate_int8 : enable
39 #endif // defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8)
40
41 #if defined(ARM_COMPUTE_DEBUG_ENABLED) && defined(cl_arm_printf)
42 #pragma OPENCL EXTENSION cl_arm_printf : enable
43 #endif // defined(ARM_COMPUTE_DEBUG_ENABLED) && defined(cl_arm_printf)
44
45 #define GPU_ARCH_MIDGARD 0x100
46 #define GPU_ARCH_BIFROST 0x200
47
48 /** Concatenate two inputs.
49 *
50 * @param[in] a The first input to be concatenated
51 * @param[in] b The second input to be concatenated
52 *
53 * @return The concatenated output
54 */
55 #define CONCAT(a, b) a##b
56
57 /** Expand the given vector
58 *
59 * @param[in] x The vector to be expanded
60 *
61 * @return The expanded output
62 */
63 #define EXPAND(x) x
64
65 /** Clamp the given value between an upper and lower bound.
66 *
67 * @param[in] x The value to be clamped
68 * @param[in] min_val The lower bound
69 * @param[in] max_val The upper bound
70 *
71 * @return The clamped value.
72 */
73 #define CLAMP(x, min_val, max_val) min(max(x, min_val), max_val)
74
75 /** REVn reverses the given vector whose size is n.
76 * @name REVn
77 *
78 * @param[in] x The vector to be reversed
79 *
80 * @return The reversed vector
81 * @{
82 */
83 #define REV1(x) ((x))
84 #define REV2(x) ((x).s10)
85 #define REV3(x) ((x).s210)
86 #define REV4(x) ((x).s3210)
87 #define REV8(x) ((x).s76543210)
88 #define REV16(x) ((x).sFEDCBA9876543210)
89 /** @} */ // end of group REVn
90
91 /** Reverse the given vector.
92 * @name REVERSE
93 *
94 * @param[in] x The vector to be reversed
95 * @param[in] s The size of the vector
96 *
97 * @return The reversed vector
98 * @{
99 */
100 #define REVERSE_STR(x, s) REV##s((x))
101 #define REVERSE(x, s) REVERSE_STR(x, s)
102 /** @} */ // end of group REVERSE
103
104 /** Circular-right-shift (rotate-right) the vector of size s by the amount of n.
105 * @name ROTs_n
106 *
107 * @param[in] x The vector to be shifted
108 *
109 * @return The shifted vector
110 * @{
111 */
112 #define ROT1_0(x) ((x))
113
114 #define ROT2_0(x) ((x))
115 #define ROT2_1(x) ((x).s10)
116
117 #define ROT3_0(x) ((x))
118 #define ROT3_1(x) ((x).s201)
119 #define ROT3_2(x) ((x).s120)
120
121 #define ROT4_0(x) ((x))
122 #define ROT4_1(x) ((x).s3012)
123 #define ROT4_2(x) ((x).s2301)
124 #define ROT4_3(x) ((x).s1230)
125
126 #define ROT8_0(x) ((x))
127 #define ROT8_1(x) ((x).s70123456)
128 #define ROT8_2(x) ((x).s67012345)
129 #define ROT8_3(x) ((x).s56701234)
130 #define ROT8_4(x) ((x).s45670123)
131 #define ROT8_5(x) ((x).s34567012)
132 #define ROT8_6(x) ((x).s23456701)
133 #define ROT8_7(x) ((x).s12345670)
134
135 #define ROT16_0(x) ((x))
136 #define ROT16_1(x) ((x).sF0123456789ABCDE)
137 #define ROT16_2(x) ((x).sEF0123456789ABCD)
138 #define ROT16_3(x) ((x).sDEF0123456789ABC)
139 #define ROT16_4(x) ((x).sCDEF0123456789AB)
140 #define ROT16_5(x) ((x).sBCDEF0123456789A)
141 #define ROT16_6(x) ((x).sABCDEF0123456789)
142 #define ROT16_7(x) ((x).s9ABCDEF012345678)
143 #define ROT16_8(x) ((x).s89ABCDEF01234567)
144 #define ROT16_9(x) ((x).s789ABCDEF0123456)
145 #define ROT16_10(x) ((x).s6789ABCDEF012345)
146 #define ROT16_11(x) ((x).s56789ABCDEF01234)
147 #define ROT16_12(x) ((x).s456789ABCDEF0123)
148 #define ROT16_13(x) ((x).s3456789ABCDEF012)
149 #define ROT16_14(x) ((x).s23456789ABCDEF01)
150 #define ROT16_15(x) ((x).s123456789ABCDEF0)
151 /** @} */ // end of group ROTs_n
152
153 /** Circular-right-shift (rotate-right) the given vector by the given amount.
154 * @name ROTATE
155 *
156 * @param[in] x The vector to be shifted
157 * @param[in] s The size of the vector
158 * @param[in] n The amount to be shifted
159 *
160 * @return The shifted vector
161 * @{
162 */
163 #define ROTATE_STR(x, s, n) ROT##s##_##n(x)
164 #define ROTATE(x, s, n) ROTATE_STR(x, s, n)
165 /** @} */ // end of group ROTATE
166
167 /** Creates a vector of size n filled with offset values corresponding to the location of each element.
168 * @name V_OFFSn
169 *
170 * @param[in] dt The data type of the output vector
171 *
172 * @return The vector filled with offset values
173 * @{
174 */
175 #define V_OFFS1(dt) (dt##1)(0)
176 #define V_OFFS2(dt) (dt##2)(0, 1)
177 #define V_OFFS3(dt) (dt##3)(0, 1, 2)
178 #define V_OFFS4(dt) (dt##4)(0, 1, 2, 3)
179 #define V_OFFS8(dt) (dt##8)(0, 1, 2, 3, 4, 5, 6, 7)
180 #define V_OFFS16(dt) (dt##16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)
181 /** @} */ // end of group V_OFFSn
182
183 /** Create a vector filled with offset values corresponding to the location of each element.
184 * @name VEC_OFFS
185 *
186 * @param[in] dt The data type of the output vector
187 * @param[in] s The size of the output vector
188 *
189 * @return The vector filled with offset values
190 * @{
191 */
192 #define VEC_OFFS_STR(dt, s) V_OFFS##s(dt)
193 #define VEC_OFFS(dt, s) VEC_OFFS_STR(dt, s)
194 /** @} */ // end of group VEC_OFFS
195
196 #define VLOAD_STR(size) vload##size
197 #define VLOAD(size) VLOAD_STR(size)
198
199 #define PIXEL_UNIT4 1
200 #define PIXEL_UNIT8 2
201 #define PIXEL_UNIT16 4
202
203 /** Utility macro to convert a vector size in pixel unit.
204 *
205 * @name CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT
206 *
207 * @param[in] vec_size Vector size. Only 4,8 and 16 is supported
208 *
209 * @return The pixel unit (number of pixels)
210 * @{
211 */
212 #define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size) PIXEL_UNIT##vec_size
213 #define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(vec_size) CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size)
214 /** @} */ // end of group CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT
215
216 #define read_image2d_floatx1(img, x_coord, y_coord) (float4)(read_imagef(img, (int2)(x_coord, y_coord)));
217 #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)));
218 #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)));
219
220 #if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
221 #define read_image2d_halfx1(img, x_coord, y_coord) (half4)(read_imageh(img, (int2)(x_coord, y_coord)));
222 #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)));
223 #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)));
224 #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
225
226 /** Utility macro to read a 2D OpenCL image object.
227 *
228 * @note Coordinates are not normalized
229 *
230 * @param[in] data_type Data type
231 * @param[in] n0 Number of pixel to read. Only 1,2 and 4 is supported
232 * @param[in] img OpenCL image object
233 * @param[in] x_coord The x coordinate for the top-left pixel
234 * @param[in] y_coord The y coordinate for the top-left pixel
235 *
236 * @return Pixels from the 2D OpenCL image object
237 * @{
238 */
239 #define READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord) read_image2d_##data_type##x##n0(img, x_coord, y_coord)
240 #define READ_IMAGE2D(data_type, n0, img, x_coord, y_coord) READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord)
241
242 #define VSTORE_STR(size) vstore##size
243 #define VSTORE(size) VSTORE_STR(size)
244
245 #define float1 float
246 #define half1 half
247 #define char1 char
248 #define uchar1 uchar
249 #define short1 short
250 #define ushort1 ushort
251 #define int1 int
252 #define uint1 uint
253 #define long1 long
254 #define ulong1 ulong
255 #define double1 double
256
257 #define vload1(OFFSET, PTR) *(OFFSET + PTR)
258 #define vstore1(DATA, OFFSET, PTR) *(OFFSET + PTR) = DATA
259
260 /** Extended partial vstore that correctly handles scalar values as well.
261 * Store the **lower** 0 to (n-1)th elements of the given vector while minimising the amount of vstore ops
262 * @name VSTORE_PARTIAL
263 *
264 * @note With this macro, the passed data can be both a vector and a scalar
265 * @note @p store_size needs to be <= @p size
266 * eg 1: Valid
267 * VSTORE_PARTIAL(16, 15) ...;
268 * eg 2: Invalid
269 * VSTORE_PARTIAL(4, 7) ...;
270 *
271 * @param[in] size The width of @p DATA. Supported values: 1(scalar), 2, 3, 4, 8, 16
272 * @param[in] store_size The number of lower elements to store. Supported values: 1-16, but has to be <= @p size
273 * @{
274 */
275 #define VSTORE_PARTIAL_STR(size, store_size) vstore_partial_##size##_##store_size
276 #define VSTORE_PARTIAL(size, store_size) VSTORE_PARTIAL_STR(size, store_size)
277
278 #define NO_STORE(data, offs, ptr) \
279 { \
280 }
281
282 // Size == 1 (scalar)
283 #define vstore_partial_1_0 NO_STORE
284 #define vstore_partial_1_1 vstore1
285 #define vstore_partial_1_2 NO_STORE
286 #define vstore_partial_1_3 NO_STORE
287 #define vstore_partial_1_4 NO_STORE
288 #define vstore_partial_1_5 NO_STORE
289 #define vstore_partial_1_6 NO_STORE
290 #define vstore_partial_1_7 NO_STORE
291 #define vstore_partial_1_8 NO_STORE
292 #define vstore_partial_1_9 NO_STORE
293 #define vstore_partial_1_10 NO_STORE
294 #define vstore_partial_1_11 NO_STORE
295 #define vstore_partial_1_12 NO_STORE
296 #define vstore_partial_1_13 NO_STORE
297 #define vstore_partial_1_14 NO_STORE
298 #define vstore_partial_1_15 NO_STORE
299 #define vstore_partial_1_16 NO_STORE
300 // Size == 2
301 #define vstore_partial_2_0 NO_STORE
302 #define vstore_partial_2_1 vstore_partial_1
303 #define vstore_partial_2_2 vstore_partial_2
304 #define vstore_partial_2_3 NO_STORE
305 #define vstore_partial_2_4 NO_STORE
306 #define vstore_partial_2_5 NO_STORE
307 #define vstore_partial_2_6 NO_STORE
308 #define vstore_partial_2_7 NO_STORE
309 #define vstore_partial_2_8 NO_STORE
310 #define vstore_partial_2_9 NO_STORE
311 #define vstore_partial_2_10 NO_STORE
312 #define vstore_partial_2_11 NO_STORE
313 #define vstore_partial_2_12 NO_STORE
314 #define vstore_partial_2_13 NO_STORE
315 #define vstore_partial_2_14 NO_STORE
316 #define vstore_partial_2_15 NO_STORE
317 #define vstore_partial_2_16 NO_STORE
318 // Size == 3
319 #define vstore_partial_3_0 NO_STORE
320 #define vstore_partial_3_1 vstore_partial_1
321 #define vstore_partial_3_2 vstore_partial_2
322 #define vstore_partial_3_3 vstore_partial_3
323 #define vstore_partial_3_4 NO_STORE
324 #define vstore_partial_3_5 NO_STORE
325 #define vstore_partial_3_6 NO_STORE
326 #define vstore_partial_3_7 NO_STORE
327 #define vstore_partial_3_8 NO_STORE
328 #define vstore_partial_3_9 NO_STORE
329 #define vstore_partial_3_10 NO_STORE
330 #define vstore_partial_3_11 NO_STORE
331 #define vstore_partial_3_12 NO_STORE
332 #define vstore_partial_3_13 NO_STORE
333 #define vstore_partial_3_14 NO_STORE
334 #define vstore_partial_3_15 NO_STORE
335 #define vstore_partial_3_16 NO_STORE
336 // Size == 4
337 #define vstore_partial_4_0 NO_STORE
338 #define vstore_partial_4_1 vstore_partial_1
339 #define vstore_partial_4_2 vstore_partial_2
340 #define vstore_partial_4_3 vstore_partial_3
341 #define vstore_partial_4_4 vstore_partial_4
342 #define vstore_partial_4_5 NO_STORE
343 #define vstore_partial_4_6 NO_STORE
344 #define vstore_partial_4_7 NO_STORE
345 #define vstore_partial_4_8 NO_STORE
346 #define vstore_partial_4_9 NO_STORE
347 #define vstore_partial_4_10 NO_STORE
348 #define vstore_partial_4_11 NO_STORE
349 #define vstore_partial_4_12 NO_STORE
350 #define vstore_partial_4_13 NO_STORE
351 #define vstore_partial_4_14 NO_STORE
352 #define vstore_partial_4_15 NO_STORE
353 #define vstore_partial_4_16 NO_STORE
354 // Size == 8
355 #define vstore_partial_8_0 NO_STORE
356 #define vstore_partial_8_1 vstore_partial_1
357 #define vstore_partial_8_2 vstore_partial_2
358 #define vstore_partial_8_3 vstore_partial_3
359 #define vstore_partial_8_4 vstore_partial_4
360 #define vstore_partial_8_5 vstore_partial_5
361 #define vstore_partial_8_6 vstore_partial_6
362 #define vstore_partial_8_7 vstore_partial_7
363 #define vstore_partial_8_8 vstore_partial_8
364 #define vstore_partial_8_9 NO_STORE
365 #define vstore_partial_8_10 NO_STORE
366 #define vstore_partial_8_11 NO_STORE
367 #define vstore_partial_8_12 NO_STORE
368 #define vstore_partial_8_13 NO_STORE
369 #define vstore_partial_8_14 NO_STORE
370 #define vstore_partial_8_15 NO_STORE
371 #define vstore_partial_8_16 NO_STORE
372 // Size == 16
373 #define vstore_partial_16_0 NO_STORE
374 #define vstore_partial_16_1 vstore_partial_1
375 #define vstore_partial_16_2 vstore_partial_2
376 #define vstore_partial_16_3 vstore_partial_3
377 #define vstore_partial_16_4 vstore_partial_4
378 #define vstore_partial_16_5 vstore_partial_5
379 #define vstore_partial_16_6 vstore_partial_6
380 #define vstore_partial_16_7 vstore_partial_7
381 #define vstore_partial_16_8 vstore_partial_8
382 #define vstore_partial_16_9 vstore_partial_9
383 #define vstore_partial_16_10 vstore_partial_10
384 #define vstore_partial_16_11 vstore_partial_11
385 #define vstore_partial_16_12 vstore_partial_12
386 #define vstore_partial_16_13 vstore_partial_13
387 #define vstore_partial_16_14 vstore_partial_14
388 #define vstore_partial_16_15 vstore_partial_15
389 #define vstore_partial_16_16 vstore_partial_16
390
391 /** Partial vstore. Store the **lower** 0 to (n-1)th elements of the given vector while minimising the amount of vstore ops
392 * @name vstore_partial_n
393 *
394 * @note @p DATA needs to be a vector not a scalar
395 * @note n needs to be <= the vector width of the input variable @p DATA
396 * eg 1: Valid
397 * vstore_partial_15(var:float16, 0, 0xabcd);
398 * eg 2: Invalid
399 * vstore_partial_7(var:float4, 0, 0xabcd);
400 *
401 * @note in cases n == 1, 2, 3, 4, 8, 16, no extra vstore is invoked, thus there's no performance penalty.
402 *
403 * @param[in] DATA The name of the variable
404 * @param[in] OFFSET Offset in n
405 * @param[in] PTR The base pointer
406 * @{
407 */
408 #define vstore_partial_1(DATA, OFFSET, PTR) \
409 vstore1(DATA.s0, OFFSET, PTR);
410
411 #define vstore_partial_2(DATA, OFFSET, PTR) \
412 vstore2(DATA.s01, OFFSET, PTR);
413
414 #define vstore_partial_3(DATA, OFFSET, PTR) \
415 vstore3(DATA.s012, OFFSET, PTR);
416
417 #define vstore_partial_4(DATA, OFFSET, PTR) \
418 vstore4(DATA.s0123, OFFSET, PTR);
419
420 #define vstore_partial_5(DATA, OFFSET, PTR) \
421 vstore_partial_4(DATA.s0123, OFFSET, PTR); \
422 vstore1(DATA.s4, OFFSET, PTR + 4);
423
424 #define vstore_partial_6(DATA, OFFSET, PTR) \
425 vstore_partial_4(DATA.s0123, OFFSET, PTR); \
426 vstore_partial_2(DATA.s45, OFFSET, PTR + 4);
427
428 #define vstore_partial_7(DATA, OFFSET, PTR) \
429 vstore_partial_4(DATA.s0123, OFFSET, PTR); \
430 vstore_partial_3(DATA.s456, OFFSET, PTR + 4);
431
432 #define vstore_partial_8(DATA, OFFSET, PTR) \
433 vstore8(DATA.s01234567, OFFSET, PTR);
434
435 #define vstore_partial_9(DATA, OFFSET, PTR) \
436 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
437 vstore1(DATA.s8, OFFSET, PTR + 8);
438
439 #define vstore_partial_10(DATA, OFFSET, PTR) \
440 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
441 vstore_partial_2(DATA.s89, OFFSET, PTR + 8);
442
443 #define vstore_partial_11(DATA, OFFSET, PTR) \
444 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
445 vstore_partial_3(DATA.s89a, OFFSET, PTR + 8);
446
447 #define vstore_partial_12(DATA, OFFSET, PTR) \
448 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
449 vstore_partial_4(DATA.s89ab, OFFSET, PTR + 8);
450
451 #define vstore_partial_13(DATA, OFFSET, PTR) \
452 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
453 vstore_partial_5(DATA.s89abcdef, OFFSET, PTR + 8);
454
455 #define vstore_partial_14(DATA, OFFSET, PTR) \
456 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
457 vstore_partial_6(DATA.s89abcdef, OFFSET, PTR + 8);
458
459 #define vstore_partial_15(DATA, OFFSET, PTR) \
460 vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
461 vstore_partial_7(DATA.s89abcdef, OFFSET, PTR + 8);
462
463 #define vstore_partial_16(DATA, OFFSET, PTR) \
464 vstore16(DATA, OFFSET, PTR);
465 /** @} */ // end of groupd vstore_partial_n
466 /** @} */ // end of groupd VSTORE_PARTIAL
467
468 // Convert built-in functions with _sat modifier are not supported in floating point so we create defines
469 // without _sat to overcome this issue
470 #define convert_float_sat convert_float
471 #define convert_float1_sat convert_float
472 #define convert_float2_sat convert_float2
473 #define convert_float3_sat convert_float3
474 #define convert_float4_sat convert_float4
475 #define convert_float8_sat convert_float8
476 #define convert_float16_sat convert_float16
477 #define convert_half_sat convert_float
478 #define convert_half1_sat convert_half
479 #define convert_half2_sat convert_half2
480 #define convert_half3_sat convert_half3
481 #define convert_half4_sat convert_half4
482 #define convert_half8_sat convert_half8
483 #define convert_half16_sat convert_half16
484
485 #define convert_float1 convert_float
486 #define convert_half1 convert_half
487 #define convert_char1 convert_char
488 #define convert_uchar1 convert_uchar
489 #define convert_short1 convert_short
490 #define convert_ushort1 convert_ushort
491 #define convert_int1 convert_int
492 #define convert_uint1 convert_uint
493 #define convert_long1 convert_long
494 #define convert_ulong1 convert_ulong
495 #define convert_double1 convert_double
496
497 #define convert_char1_sat convert_char_sat
498 #define convert_uchar1_sat convert_uchar_sat
499 #define convert_short1_sat convert_short_sat
500 #define convert_ushort1_sat convert_ushort_sat
501 #define convert_int1_sat convert_int_sat
502 #define convert_uint1_sat convert_uint_sat
503 #define convert_long1_sat convert_long_sat
504 #define convert_ulong1_sat convert_ulong_sat
505 #define convert_double1_sat convert_double_sat
506
507 #define VEC_DATA_TYPE_STR(type, size) type##size
508 #define VEC_DATA_TYPE(type, size) VEC_DATA_TYPE_STR(type, size)
509
510 #define CONVERT_STR(x, type) (convert_##type((x)))
511 #define CONVERT(x, type) CONVERT_STR(x, type)
512
513 #define CONVERT_SAT_STR(x, type) (convert_##type##_sat((x)))
514 #define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type)
515
516 #define CONVERT_SAT_ROUND_STR(x, type, round) (convert_##type##_sat_##round((x)))
517 #define CONVERT_SAT_ROUND(x, type, round) CONVERT_SAT_ROUND_STR(x, type, round)
518
519 #define select_vec_dt_uchar(size) uchar##size
520 #define select_vec_dt_char(size) char##size
521 #define select_vec_dt_ushort(size) ushort##size
522 #define select_vec_dt_short(size) short##size
523 #define select_vec_dt_half(size) short##size
524 #define select_vec_dt_uint(size) uint##size
525 #define select_vec_dt_int(size) int##size
526 #define select_vec_dt_float(size) int##size
527 #define select_vec_dt_ulong(size) ulong##size
528 #define select_vec_dt_long(size) long##size
529
530 #define SELECT_VEC_DATA_TYPE_STR(type, size) select_vec_dt_##type(size)
531 #define SELECT_VEC_DATA_TYPE(type, size) SELECT_VEC_DATA_TYPE_STR(type, size)
532 #define SELECT_DATA_TYPE(type) SELECT_VEC_DATA_TYPE_STR(type, 1)
533
534 #define sum_reduce_1(x) (x)
535 #define sum_reduce_2(x) ((x).s0) + ((x).s1)
536 #define sum_reduce_3(x) sum_reduce_2((x).s01) + ((x).s2)
537 #define sum_reduce_4(x) sum_reduce_2((x).s01) + sum_reduce_2((x).s23)
538 #define sum_reduce_8(x) sum_reduce_4((x).s0123) + sum_reduce_4((x).s4567)
539 #define sum_reduce_16(x) sum_reduce_8((x).s01234567) + sum_reduce_8((x).s89ABCDEF)
540
541 #define SUM_REDUCE_STR(x, size) sum_reduce_##size(x)
542 #define SUM_REDUCE(x, size) SUM_REDUCE_STR(x, size)
543
544 #define max_reduce_1(x) (x)
545 #define max_reduce_2(x) max(((x).s0), ((x).s1))
546 #define max_reduce_3(x) max(max_reduce_2((x).s01), ((x).s2))
547 #define max_reduce_4(x) max(max_reduce_2((x).s01), max_reduce_2((x).s23))
548 #define max_reduce_8(x) max(max_reduce_4((x).s0123), max_reduce_4((x).s4567))
549 #define max_reduce_16(x) max(max_reduce_8((x).s01234567), max_reduce_8((x).s89ABCDEF))
550
551 #define MAX_REDUCE_STR(x, size) max_reduce_##size(x)
552 #define MAX_REDUCE(x, size) MAX_REDUCE_STR(x, size)
553
554 #define VECTOR_DECLARATION(name) \
555 __global uchar *name##_ptr, \
556 uint name##_stride_x, \
557 uint name##_step_x, \
558 uint name##_offset_first_element_in_bytes
559
560 #define IMAGE_DECLARATION(name) \
561 __global uchar *name##_ptr, \
562 uint name##_stride_x, \
563 uint name##_step_x, \
564 uint name##_stride_y, \
565 uint name##_step_y, \
566 uint name##_offset_first_element_in_bytes
567
568 #define TENSOR3D_DECLARATION(name) \
569 __global uchar *name##_ptr, \
570 uint name##_stride_x, \
571 uint name##_step_x, \
572 uint name##_stride_y, \
573 uint name##_step_y, \
574 uint name##_stride_z, \
575 uint name##_step_z, \
576 uint name##_offset_first_element_in_bytes
577
578 #define TENSOR4D_DECLARATION(name) \
579 __global uchar *name##_ptr, \
580 uint name##_stride_x, \
581 uint name##_step_x, \
582 uint name##_stride_y, \
583 uint name##_step_y, \
584 uint name##_stride_z, \
585 uint name##_step_z, \
586 uint name##_stride_w, \
587 uint name##_step_w, \
588 uint name##_offset_first_element_in_bytes
589
590 #define CONVERT_TO_VECTOR_STRUCT(name) \
591 update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x)
592
593 #define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name) \
594 update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0)
595
596 #define CONVERT_TO_IMAGE_STRUCT(name) \
597 update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y)
598
599 #define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name) \
600 update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0)
601
602 #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \
603 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)
604
605 #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name) \
606 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)
607
608 #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \
609 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)
610
611 #define CONVERT_TO_TENSOR3D_STRUCT(name) \
612 update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \
613 name##_stride_z, name##_step_z)
614
615 #define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name) \
616 update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0)
617
618 #define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size) \
619 update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \
620 name##_stride_z, name##_step_z, name##_stride_w, name##_step_w, mod_size)
621
622 #define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size) \
623 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)
624
625 #define CONVERT_TO_TENSOR3D_STRUCT_NO_UPDATE_PTR(name) \
626 tensor3D_ptr_no_update(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \
627 name##_stride_z, name##_step_z)
628
629 /** Structure to hold Vector information */
630 typedef struct Vector
631 {
632 __global uchar *ptr; /**< Pointer to the starting postion of the buffer */
633 int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
634 int stride_x; /**< Stride of the image in X dimension (in bytes) */
635 } Vector;
636
637 /** Structure to hold Image information */
638 typedef struct Image
639 {
640 __global uchar *ptr; /**< Pointer to the starting postion of the buffer */
641 int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
642 int stride_x; /**< Stride of the image in X dimension (in bytes) */
643 int stride_y; /**< Stride of the image in Y dimension (in bytes) */
644 } Image;
645
646 /** Structure to hold 3D tensor information */
647 typedef struct Tensor3D
648 {
649 __global uchar *ptr; /**< Pointer to the starting postion of the buffer */
650 int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
651 int stride_x; /**< Stride of the image in X dimension (in bytes) */
652 int stride_y; /**< Stride of the image in Y dimension (in bytes) */
653 int stride_z; /**< Stride of the image in Z dimension (in bytes) */
654 } Tensor3D;
655
656 /** Structure to hold 4D tensor information */
657 typedef struct Tensor4D
658 {
659 __global uchar *ptr; /**< Pointer to the starting postion of the buffer */
660 int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
661 int stride_x; /**< Stride of the image in X dimension (in bytes) */
662 int stride_y; /**< Stride of the image in Y dimension (in bytes) */
663 int stride_z; /**< Stride of the image in Z dimension (in bytes) */
664 int stride_w; /**< Stride of the image in W dimension (in bytes) */
665 } Tensor4D;
666
667 /** Wrap vector information into an Vector structure, and make the pointer point at this workitem's data.
668 *
669 * @param[in] ptr Pointer to the starting postion of the buffer
670 * @param[in] offset_first_element_in_bytes The offset of the first element in the source vector
671 * @param[in] stride_x Stride of the vector in X dimension (in bytes)
672 * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
673 *
674 * @return An image object
675 */
update_vector_workitem_ptr(__global uchar * ptr,uint offset_first_element_in_bytes,uint stride_x,uint step_x)676 inline Vector update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x)
677 {
678 Vector vector =
679 {
680 .ptr = ptr,
681 .offset_first_element_in_bytes = offset_first_element_in_bytes,
682 .stride_x = stride_x,
683 };
684 vector.ptr += vector.offset_first_element_in_bytes + get_global_id(0) * step_x;
685 return vector;
686 }
687
688 /** Wrap image information into an Image structure, and make the pointer point at this workitem's data.
689 *
690 * @param[in] ptr Pointer to the starting postion of the buffer
691 * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
692 * @param[in] stride_x Stride of the image in X dimension (in bytes)
693 * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
694 * @param[in] stride_y Stride of the image in Y dimension (in bytes)
695 * @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
696 *
697 * @return An image object
698 */
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)699 inline 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)
700 {
701 Image img =
702 {
703 .ptr = ptr,
704 .offset_first_element_in_bytes = offset_first_element_in_bytes,
705 .stride_x = stride_x,
706 .stride_y = stride_y
707 };
708 img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y;
709 return img;
710 }
711
712 /** Wrap 3D tensor information into an image structure, and make the pointer point at this workitem's data.
713 *
714 * @param[in] ptr Pointer to the starting postion of the buffer
715 * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
716 * @param[in] stride_x Stride of the image in X dimension (in bytes)
717 * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
718 * @param[in] stride_y Stride of the image in Y dimension (in bytes)
719 * @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
720 * @param[in] stride_z Stride of the image in Z dimension (in bytes)
721 * @param[in] step_z stride_z * number of elements along Z processed per workitem(in bytes)
722 *
723 * @return A 3D tensor object
724 */
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)725 inline 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)
726 {
727 Image img =
728 {
729 .ptr = ptr,
730 .offset_first_element_in_bytes = offset_first_element_in_bytes,
731 .stride_x = stride_x,
732 .stride_y = stride_y
733 };
734 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;
735 return img;
736 }
737
738 /** Wrap 3D tensor information into an tensor structure, and make the pointer point at this workitem's data.
739 *
740 * @param[in] ptr Pointer to the starting postion of the buffer
741 * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
742 * @param[in] stride_x Stride of the image in X dimension (in bytes)
743 * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
744 * @param[in] stride_y Stride of the image in Y dimension (in bytes)
745 * @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
746 * @param[in] stride_z Stride of the image in Z dimension (in bytes)
747 * @param[in] step_z stride_z * number of elements along Z processed per workitem(in bytes)
748 *
749 * @return A 3D tensor object
750 */
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)751 inline 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)
752 {
753 Tensor3D tensor =
754 {
755 .ptr = ptr,
756 .offset_first_element_in_bytes = offset_first_element_in_bytes,
757 .stride_x = stride_x,
758 .stride_y = stride_y,
759 .stride_z = stride_z
760 };
761 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;
762 return tensor;
763 }
764
765 /** Wrap 3D tensor information into an tensor structure.
766 *
767 * @param[in] ptr Pointer to the starting postion of the buffer
768 * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
769 * @param[in] stride_x Stride of the image in X dimension (in bytes)
770 * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
771 * @param[in] stride_y Stride of the image in Y dimension (in bytes)
772 * @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
773 * @param[in] stride_z Stride of the image in Z dimension (in bytes)
774 * @param[in] step_z stride_z * number of elements along Z processed per workitem(in bytes)
775 *
776 * @return A 3D tensor object
777 */
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)778 inline 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)
779 {
780 Tensor3D tensor =
781 {
782 .ptr = ptr,
783 .offset_first_element_in_bytes = offset_first_element_in_bytes,
784 .stride_x = stride_x,
785 .stride_y = stride_y,
786 .stride_z = stride_z
787 };
788 return tensor;
789 }
790
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,uint step_w,uint mod_size)791 inline 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,
792 uint step_w,
793 uint mod_size)
794 {
795 Tensor4D tensor =
796 {
797 .ptr = ptr,
798 .offset_first_element_in_bytes = offset_first_element_in_bytes,
799 .stride_x = stride_x,
800 .stride_y = stride_y,
801 .stride_z = stride_z,
802 .stride_w = stride_w
803 };
804
805 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;
806 return tensor;
807 }
808
809 /** Get the pointer position of a Vector
810 *
811 * @param[in] vec Pointer to the starting position of the buffer
812 * @param[in] x Relative X position
813 */
vector_offset(const Vector * vec,int x)814 inline __global const uchar *vector_offset(const Vector *vec, int x)
815 {
816 return vec->ptr + x * vec->stride_x;
817 }
818
819 /** Get the pointer position of a Image
820 *
821 * @param[in] img Pointer to the starting position of the buffer
822 * @param[in] x Relative X position
823 * @param[in] y Relative Y position
824 */
offset(const Image * img,int x,int y)825 inline __global uchar *offset(const Image *img, int x, int y)
826 {
827 return img->ptr + x * img->stride_x + y * img->stride_y;
828 }
829
830 /** Get the pointer position of a Tensor3D
831 *
832 * @param[in] tensor Pointer to the starting position of the buffer
833 * @param[in] x Relative X position
834 * @param[in] y Relative Y position
835 * @param[in] z Relative Z position
836 */
tensor3D_offset(const Tensor3D * tensor,int x,int y,int z)837 inline __global const uchar *tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
838 {
839 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z;
840 }
841
842 /** Get the pointer position of a Tensor4D
843 *
844 * @param[in] tensor Pointer to the starting position of the buffer
845 * @param[in] x Relative X position
846 * @param[in] y Relative Y position
847 * @param[in] z Relative Z position
848 * @param[in] w Relative W position
849 */
tensor4D_offset(const Tensor4D * tensor,int x,int y,int z,int w)850 inline __global const uchar *tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
851 {
852 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + w * tensor->stride_w;
853 }
854
855 /** Get the offset for a given linear index of a Tensor3D
856 *
857 * @param[in] tensor Pointer to the starting position of the buffer
858 * @param[in] width Width of the input tensor
859 * @param[in] height Height of the input tensor
860 * @param[in] depth Depth of the input tensor
861 * @param[in] index Linear index
862 */
tensor3D_index2ptr(const Tensor3D * tensor,uint width,uint height,uint depth,uint index)863 inline __global const uchar *tensor3D_index2ptr(const Tensor3D *tensor, uint width, uint height, uint depth, uint index)
864 {
865 uint num_elements = width * height;
866
867 const uint z = index / num_elements;
868
869 index %= num_elements;
870
871 const uint y = index / width;
872
873 index %= width;
874
875 const uint x = index;
876
877 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + tensor->offset_first_element_in_bytes;
878 }
879
880 #endif // _HELPER_H
881