Compute Library
 23.11
helpers.h
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2016-2023 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 #define GPU_ARCH_VALHALL 0x300
48 
49 /** Concatenate two inputs.
50  *
51  * @param[in] a The first input to be concatenated
52  * @param[in] b The second input to be concatenated
53  *
54  * @return The concatenated output
55  */
56 #define CONCAT(a, b) a##b
57 
58 /** Expand the given vector
59  *
60  * @param[in] x The vector to be expanded
61  *
62  * @return The expanded output
63  */
64 #define EXPAND(x) x
65 
66 /** Clamp the given value between an upper and lower bound.
67  *
68  * @param[in] x The value to be clamped
69  * @param[in] min_val The lower bound
70  * @param[in] max_val The upper bound
71  *
72  * @return The clamped value.
73  */
74 #define CLAMP(x, min_val, max_val) min(max(x, min_val), max_val)
75 
76 /** REVn reverses the given vector whose size is n.
77  * @name REVn
78  *
79  * @param[in] x The vector to be reversed
80  *
81  * @return The reversed vector
82  * @{
83  */
84 #define REV1(x) ((x))
85 #define REV2(x) ((x).s10)
86 #define REV3(x) ((x).s210)
87 #define REV4(x) ((x).s3210)
88 #define REV8(x) ((x).s76543210)
89 #define REV16(x) ((x).sFEDCBA9876543210)
90 /** @} */ // end of group REVn
91 
92 /** Reverse the given vector.
93  * @name REVERSE
94  *
95  * @param[in] x The vector to be reversed
96  * @param[in] s The size of the vector
97  *
98  * @return The reversed vector
99  * @{
100  */
101 #define REVERSE_STR(x, s) REV##s((x))
102 #define REVERSE(x, s) REVERSE_STR(x, s)
103 /** @} */ // end of group REVERSE
104 
105 /** Circular-right-shift (rotate-right) the vector of size s by the amount of n.
106  * @name ROTs_n
107  *
108  * @param[in] x The vector to be shifted
109  *
110  * @return The shifted vector
111  * @{
112  */
113 #define ROT1_0(x) ((x))
114 #define ROT1_1(x) ((x))
115 
116 #define ROT2_0(x) ((x))
117 #define ROT2_1(x) ((x).s10)
118 #define ROT2_2(x) ((x))
119 
120 #define ROT3_0(x) ((x))
121 #define ROT3_1(x) ((x).s201)
122 #define ROT3_2(x) ((x).s120)
123 #define ROT3_3(x) ((x))
124 
125 #define ROT4_0(x) ((x))
126 #define ROT4_1(x) ((x).s3012)
127 #define ROT4_2(x) ((x).s2301)
128 #define ROT4_3(x) ((x).s1230)
129 #define ROT4_4(x) ((x))
130 
131 #define ROT8_0(x) ((x))
132 #define ROT8_1(x) ((x).s70123456)
133 #define ROT8_2(x) ((x).s67012345)
134 #define ROT8_3(x) ((x).s56701234)
135 #define ROT8_4(x) ((x).s45670123)
136 #define ROT8_5(x) ((x).s34567012)
137 #define ROT8_6(x) ((x).s23456701)
138 #define ROT8_7(x) ((x).s12345670)
139 #define ROT8_8(x) ((x))
140 
141 #define ROT16_0(x) ((x))
142 #define ROT16_1(x) ((x).sF0123456789ABCDE)
143 #define ROT16_2(x) ((x).sEF0123456789ABCD)
144 #define ROT16_3(x) ((x).sDEF0123456789ABC)
145 #define ROT16_4(x) ((x).sCDEF0123456789AB)
146 #define ROT16_5(x) ((x).sBCDEF0123456789A)
147 #define ROT16_6(x) ((x).sABCDEF0123456789)
148 #define ROT16_7(x) ((x).s9ABCDEF012345678)
149 #define ROT16_8(x) ((x).s89ABCDEF01234567)
150 #define ROT16_9(x) ((x).s789ABCDEF0123456)
151 #define ROT16_10(x) ((x).s6789ABCDEF012345)
152 #define ROT16_11(x) ((x).s56789ABCDEF01234)
153 #define ROT16_12(x) ((x).s456789ABCDEF0123)
154 #define ROT16_13(x) ((x).s3456789ABCDEF012)
155 #define ROT16_14(x) ((x).s23456789ABCDEF01)
156 #define ROT16_15(x) ((x).s123456789ABCDEF0)
157 #define ROT16_16(x) ((x))
158 /** @} */ // end of group ROTs_n
159 
160 /** Circular-right-shift (rotate-right) the given vector by the given amount.
161  * @name ROTATE
162  *
163  * @param[in] x The vector to be shifted
164  * @param[in] s The size of the vector
165  * @param[in] n The amount to be shifted
166  *
167  * @return The shifted vector
168  * @{
169  */
170 #define ROTATE_STR(x, s, n) ROT##s##_##n(x)
171 #define ROTATE(x, s, n) ROTATE_STR(x, s, n)
172 /** @} */ // end of group ROTATE
173 
174 /** Creates a vector of size n filled with offset values corresponding to the location of each element.
175  * @name V_OFFSn
176  *
177  * @param[in] dt The data type of the output vector
178  *
179  * @return The vector filled with offset values
180  * @{
181  */
182 #define V_OFFS1(dt) (dt##1)(0)
183 #define V_OFFS2(dt) (dt##2)(0, 1)
184 #define V_OFFS3(dt) (dt##3)(0, 1, 2)
185 #define V_OFFS4(dt) (dt##4)(0, 1, 2, 3)
186 #define V_OFFS8(dt) (dt##8)(0, 1, 2, 3, 4, 5, 6, 7)
187 #define V_OFFS16(dt) (dt##16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)
188 /** @} */ // end of group V_OFFSn
189 
190 /** Create a vector filled with offset values corresponding to the location of each element.
191  * @name VEC_OFFS
192  *
193  * @param[in] dt The data type of the output vector
194  * @param[in] s The size of the output vector
195  *
196  * @return The vector filled with offset values
197  * @{
198  */
199 #define VEC_OFFS_STR(dt, s) V_OFFS##s(dt)
200 #define VEC_OFFS(dt, s) VEC_OFFS_STR(dt, s)
201 /** @} */ // end of group VEC_OFFS
202 
203 #define VLOAD_STR(size) vload##size
204 #define VLOAD(size) VLOAD_STR(size)
205 
206 /** Extended partial vload that correctly handles scalar values as well.
207  * Load the **lower** 0 to (n-1)th elements of the given vector while minimising the amount of load ops
208  * @name VLOAD_PARTIAL
209  *
210  * @note With this macro, the passed data can be both a vector and a scalar
211  * @note @p load_size needs to be <= @p size
212  * eg 1: Valid
213  * VLOAD_PARTIAL(16, 15) ...;
214  * eg 2: Invalid
215  * VLOAD_PARTIAL(4, 7) ...;
216  *
217  * @param[in] size The width of @p DATA. Supported values: 1(scalar), 2, 3, 4, 8, 16
218  * @param[in] load_size The number of lower elements to load. Supported values: 1-16, but has to be <= @p size
219  * @{
220  */
221 #define VLOAD_PARTIAL_STR(size, load_size) vload_partial_##size##_##load_size
222 #define VLOAD_PARTIAL(size, load_size) VLOAD_PARTIAL_STR(size, load_size)
223 
224 #define NO_LOAD(data, offs, ptr) \
225  { \
226  }
227 
228 // Size == 1 (scalar)
229 #define vload_partial_1_0 NO_LOAD
230 #define vload_partial_1_1 vload1
231 #define vload_partial_1_2 NO_LOAD
232 #define vload_partial_1_3 NO_LOAD
233 #define vload_partial_1_4 NO_LOAD
234 #define vload_partial_1_5 NO_LOAD
235 #define vload_partial_1_6 NO_LOAD
236 #define vload_partial_1_7 NO_LOAD
237 #define vload_partial_1_8 NO_LOAD
238 #define vload_partial_1_9 NO_LOAD
239 #define vload_partial_1_10 NO_LOAD
240 #define vload_partial_1_11 NO_LOAD
241 #define vload_partial_1_12 NO_LOAD
242 #define vload_partial_1_13 NO_LOAD
243 #define vload_partial_1_14 NO_LOAD
244 #define vload_partial_1_15 NO_LOAD
245 #define vload_partial_1_16 NO_LOAD
246 // Size == 2
247 #define vload_partial_2_0 NO_LOAD
248 #define vload_partial_2_1 vload_partial_1
249 #define vload_partial_2_2 vload_partial_2
250 #define vload_partial_2_3 NO_LOAD
251 #define vload_partial_2_4 NO_LOAD
252 #define vload_partial_2_5 NO_LOAD
253 #define vload_partial_2_6 NO_LOAD
254 #define vload_partial_2_7 NO_LOAD
255 #define vload_partial_2_8 NO_LOAD
256 #define vload_partial_2_9 NO_LOAD
257 #define vload_partial_2_10 NO_LOAD
258 #define vload_partial_2_11 NO_LOAD
259 #define vload_partial_2_12 NO_LOAD
260 #define vload_partial_2_13 NO_LOAD
261 #define vload_partial_2_14 NO_LOAD
262 #define vload_partial_2_15 NO_LOAD
263 #define vload_partial_2_16 NO_LOAD
264 // Size == 3
265 #define vload_partial_3_0 NO_LOAD
266 #define vload_partial_3_1 vload_partial_1
267 #define vload_partial_3_2 vload_partial_2
268 #define vload_partial_3_3 vload_partial_3
269 #define vload_partial_3_4 NO_LOAD
270 #define vload_partial_3_5 NO_LOAD
271 #define vload_partial_3_6 NO_LOAD
272 #define vload_partial_3_7 NO_LOAD
273 #define vload_partial_3_8 NO_LOAD
274 #define vload_partial_3_9 NO_LOAD
275 #define vload_partial_3_10 NO_LOAD
276 #define vload_partial_3_11 NO_LOAD
277 #define vload_partial_3_12 NO_LOAD
278 #define vload_partial_3_13 NO_LOAD
279 #define vload_partial_3_14 NO_LOAD
280 #define vload_partial_3_15 NO_LOAD
281 #define vload_partial_3_16 NO_LOAD
282 // Size == 4
283 #define vload_partial_4_0 NO_LOAD
284 #define vload_partial_4_1 vload_partial_1
285 #define vload_partial_4_2 vload_partial_2
286 #define vload_partial_4_3 vload_partial_3
287 #define vload_partial_4_4 vload_partial_4
288 #define vload_partial_4_5 NO_LOAD
289 #define vload_partial_4_6 NO_LOAD
290 #define vload_partial_4_7 NO_LOAD
291 #define vload_partial_4_8 NO_LOAD
292 #define vload_partial_4_9 NO_LOAD
293 #define vload_partial_4_10 NO_LOAD
294 #define vload_partial_4_11 NO_LOAD
295 #define vload_partial_4_12 NO_LOAD
296 #define vload_partial_4_13 NO_LOAD
297 #define vload_partial_4_14 NO_LOAD
298 #define vload_partial_4_15 NO_LOAD
299 #define vload_partial_4_16 NO_LOAD
300 // Size == 8
301 #define vload_partial_8_0 NO_LOAD
302 #define vload_partial_8_1 vload_partial_1
303 #define vload_partial_8_2 vload_partial_2
304 #define vload_partial_8_3 vload_partial_3
305 #define vload_partial_8_4 vload_partial_4
306 #define vload_partial_8_5 vload_partial_5
307 #define vload_partial_8_6 vload_partial_6
308 #define vload_partial_8_7 vload_partial_7
309 #define vload_partial_8_8 vload_partial_8
310 #define vload_partial_8_9 NO_LOAD
311 #define vload_partial_8_10 NO_LOAD
312 #define vload_partial_8_11 NO_LOAD
313 #define vload_partial_8_12 NO_LOAD
314 #define vload_partial_8_13 NO_LOAD
315 #define vload_partial_8_14 NO_LOAD
316 #define vload_partial_8_15 NO_LOAD
317 #define vload_partial_8_16 NO_LOAD
318 // Size == 16
319 #define vload_partial_16_0 NO_LOAD
320 #define vload_partial_16_1 vload_partial_1
321 #define vload_partial_16_2 vload_partial_2
322 #define vload_partial_16_3 vload_partial_3
323 #define vload_partial_16_4 vload_partial_4
324 #define vload_partial_16_5 vload_partial_5
325 #define vload_partial_16_6 vload_partial_6
326 #define vload_partial_16_7 vload_partial_7
327 #define vload_partial_16_8 vload_partial_8
328 #define vload_partial_16_9 vload_partial_9
329 #define vload_partial_16_10 vload_partial_10
330 #define vload_partial_16_11 vload_partial_11
331 #define vload_partial_16_12 vload_partial_12
332 #define vload_partial_16_13 vload_partial_13
333 #define vload_partial_16_14 vload_partial_14
334 #define vload_partial_16_15 vload_partial_15
335 #define vload_partial_16_16 vload_partial_16
336 
337 /** Partial vload. Load the **lower** 0 to (n-1)th elements of the given vector while minimising the amount of vload ops
338  * @name vload_partial_n
339  *
340  * @note @p DATA needs to be a vector not a scalar
341  * @note n needs to be <= the vector width of the input variable @p DATA
342  * eg 1: Valid
343  * vload_partial_15(var:float16, 0, 0xabcd);
344  * eg 2: Invalid
345  * vload_partial_7(var:float4, 0, 0xabcd);
346  *
347  * @note in cases n == 1, 2, 3, 4, 8, 16, no extra vload is invoked, thus there's no performance penalty.
348  *
349  * @param[in] DATA The name of the variable where to load the values
350  * @param[in] OFFSET Offset in n
351  * @param[in] PTR The base pointer
352  * @{
353  */
354 #define vload_partial_1(DATA, OFFSET, PTR) DATA.s0 = vload1(OFFSET, PTR);
355 
356 #define vload_partial_2(DATA, OFFSET, PTR) DATA.s01 = vload2(OFFSET, PTR);
357 
358 #define vload_partial_3(DATA, OFFSET, PTR) DATA.s012 = vload3(OFFSET, PTR);
359 
360 #define vload_partial_4(DATA, OFFSET, PTR) DATA.s0123 = vload4(OFFSET, PTR);
361 
362 #define vload_partial_5(DATA, OFFSET, PTR) \
363  vload_partial_4(DATA.s0123, OFFSET, PTR); \
364  DATA.s4 = vload1(OFFSET, PTR + 4);
365 
366 #define vload_partial_6(DATA, OFFSET, PTR) \
367  vload_partial_4(DATA.s0123, OFFSET, PTR); \
368  vload_partial_2(DATA.s45, OFFSET, PTR + 4);
369 
370 #define vload_partial_7(DATA, OFFSET, PTR) \
371  vload_partial_4(DATA.s0123, OFFSET, PTR); \
372  vload_partial_3(DATA.s456, OFFSET, PTR + 4);
373 
374 #define vload_partial_8(DATA, OFFSET, PTR) DATA.s01234567 = vload8(OFFSET, PTR);
375 
376 #define vload_partial_9(DATA, OFFSET, PTR) \
377  vload_partial_8(DATA.s01234567, OFFSET, PTR); \
378  DATA.s8 = vload1(OFFSET, PTR + 8);
379 
380 #define vload_partial_10(DATA, OFFSET, PTR) \
381  vload_partial_8(DATA.s01234567, OFFSET, PTR); \
382  vload_partial_2(DATA.s89, OFFSET, PTR + 8);
383 
384 #define vload_partial_11(DATA, OFFSET, PTR) \
385  vload_partial_8(DATA.s01234567, OFFSET, PTR); \
386  vload_partial_3(DATA.s89A, OFFSET, PTR + 8);
387 
388 #define vload_partial_12(DATA, OFFSET, PTR) \
389  vload_partial_8(DATA.s01234567, OFFSET, PTR); \
390  vload_partial_4(DATA.s89AB, OFFSET, PTR + 8);
391 // For vload_partial_{13,14,15}, an 8-vector size has been passed, because vectors size of size 5,6,7 are not supported
392 #define vload_partial_13(DATA, OFFSET, PTR) \
393  vload_partial_8(DATA.s01234567, OFFSET, PTR); \
394  vload_partial_5(DATA.s89ABCDEF, OFFSET, PTR + 8);
395 
396 #define vload_partial_14(DATA, OFFSET, PTR) \
397  vload_partial_8(DATA.s01234567, OFFSET, PTR); \
398  vload_partial_6(DATA.s89ABCDEF, OFFSET, PTR + 8);
399 
400 #define vload_partial_15(DATA, OFFSET, PTR) \
401  vload_partial_8(DATA.s01234567, OFFSET, PTR); \
402  vload_partial_7(DATA.s89ABCDEF, OFFSET, PTR + 8);
403 
404 #define vload_partial_16(DATA, OFFSET, PTR) DATA = vload16(OFFSET, PTR);
405 /** @} */ // end of groupd vload_partial_n
406 /** @} */ // end of groupd VLOAD_PARTIAL
407 
408 #define PIXEL_UNIT4 1
409 #define PIXEL_UNIT8 2
410 #define PIXEL_UNIT16 4
411 
412 /** Utility macro to convert a vector size in pixel unit.
413  *
414  * @name CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT
415  *
416  * @param[in] vec_size Vector size. Only 4,8 and 16 is supported
417  *
418  * @return The pixel unit (number of pixels)
419  * @{
420  */
421 #define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size) PIXEL_UNIT##vec_size
422 #define CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(vec_size) CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT_STR(vec_size)
423 /** @} */ // end of group CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT
424 
425 #define read_image2d_floatx1(img, x_coord, y_coord) (float4)(read_imagef(img, (int2)(x_coord, y_coord)));
426 #define read_image2d_floatx2(img, x_coord, y_coord) \
427  (float8)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord)));
428 #define read_image2d_floatx4(img, x_coord, y_coord) \
429  (float16)(read_imagef(img, (int2)(x_coord, y_coord)), read_imagef(img, (int2)(x_coord + 1, y_coord)), \
430  read_imagef(img, (int2)(x_coord + 2, y_coord)), read_imagef(img, (int2)(x_coord + 3, y_coord)));
431 
432 #if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
433 #define read_image2d_halfx1(img, x_coord, y_coord) (half4)(read_imageh(img, (int2)(x_coord, y_coord)));
434 #define read_image2d_halfx2(img, x_coord, y_coord) \
435  (half8)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord)));
436 #define read_image2d_halfx4(img, x_coord, y_coord) \
437  (half16)(read_imageh(img, (int2)(x_coord, y_coord)), read_imageh(img, (int2)(x_coord + 1, y_coord)), \
438  read_imageh(img, (int2)(x_coord + 2, y_coord)), read_imageh(img, (int2)(x_coord + 3, y_coord)));
439 #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
440 
441 #define write_image2d_floatx1(img, x_coord, y_coord, values) (write_imagef(img, (int2)(x_coord, y_coord), values));
442 #define write_image2d_floatx2(img, x_coord, y_coord, values) \
443  (write_imagef(img, (int2)(x_coord, y_coord), values.s0123), \
444  write_imagef(img, (int2)(x_coord + 1, y_coord), values.s4567));
445 #define write_image2d_floatx4(img, x_coord, y_coord, values) \
446  (write_imagef(img, (int2)(x_coord, y_coord), values.s0123), \
447  write_imagef(img, (int2)(x_coord + 1, y_coord), values.s4567), \
448  write_imagef(img, (int2)(x_coord + 2, y_coord), values.s89AB), \
449  write_imagef(img, (int2)(x_coord + 3, y_coord), values.sCDEF));
450 
451 #if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
452 #define write_image2d_halfx1(img, x_coord, y_coord, values) (write_imageh(img, (int2)(x_coord, y_coord), values));
453 #define write_image2d_halfx2(img, x_coord, y_coord, values) \
454  (write_imageh(img, (int2)(x_coord, y_coord), values.s0123), \
455  write_imageh(img, (int2)(x_coord + 1, y_coord), values.s4567));
456 #define write_image2d_halfx4(img, x_coord, y_coord, values) \
457  (write_imageh(img, (int2)(x_coord, y_coord), values.s0123), \
458  write_imageh(img, (int2)(x_coord + 1, y_coord), values.s4567), \
459  write_imageh(img, (int2)(x_coord + 2, y_coord), values.s89AB), \
460  write_imageh(img, (int2)(x_coord + 3, y_coord), values.sCDEF));
461 #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) && defined(cl_khr_fp16)
462 
463 /** Utility macro to read a 2D OpenCL image object.
464  *
465  * @note Coordinates are not normalized
466  *
467  * @param[in] data_type Data type
468  * @param[in] n0 Number of pixel to read. Only 1,2 and 4 is supported
469  * @param[in] img OpenCL image object
470  * @param[in] x_coord The x coordinate for the top-left pixel
471  * @param[in] y_coord The y coordinate for the top-left pixel
472  *
473  * @return Pixels from the 2D OpenCL image object
474  * @{
475  */
476 #define READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord) read_image2d_##data_type##x##n0(img, x_coord, y_coord)
477 #define READ_IMAGE2D(data_type, n0, img, x_coord, y_coord) READ_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord)
478 /** @} */
479 
480 /** Utility macro to write a 2D OpenCL image object.
481  *
482  * @note Coordinates are not normalized
483  *
484  * @param[in] data_type Data type
485  * @param[in] n0 Number of pixel to write. Only 1,2 and 4 is supported
486  * @param[in] img OpenCL image object
487  * @param[in] x_coord The x coordinate for the top-left pixel
488  * @param[in] y_coord The y coordinate for the top-left pixel
489  * @param[in] values Values to write
490  *
491  * @{
492  */
493 #define WRITE_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord, values) \
494  write_image2d_##data_type##x##n0(img, x_coord, y_coord, values)
495 #define WRITE_IMAGE2D(data_type, n0, img, x_coord, y_coord, values) \
496  WRITE_IMAGE2D_STR(data_type, n0, img, x_coord, y_coord, values)
497 /** @} */
498 
499 #define VSTORE_STR(size) vstore##size
500 #define VSTORE(size) VSTORE_STR(size)
501 
502 #define float1 float
503 #define half1 half
504 #define char1 char
505 #define uchar1 uchar
506 #define short1 short
507 #define ushort1 ushort
508 #define int1 int
509 #define uint1 uint
510 #define long1 long
511 #define ulong1 ulong
512 #define double1 double
513 
514 #define vload1(OFFSET, PTR) *(OFFSET + PTR)
515 #define vstore1(DATA, OFFSET, PTR) *(OFFSET + PTR) = DATA
516 
517 /** Extended partial vstore that correctly handles scalar values as well.
518  * Store the **lower** 0 to (n-1)th elements of the given vector while minimising the amount of vstore ops
519  * @name VSTORE_PARTIAL
520  *
521  * @note With this macro, the passed data can be both a vector and a scalar
522  * @note @p store_size needs to be <= @p size
523  * eg 1: Valid
524  * VSTORE_PARTIAL(16, 15) ...;
525  * eg 2: Invalid
526  * VSTORE_PARTIAL(4, 7) ...;
527  *
528  * @param[in] size The width of @p DATA. Supported values: 1(scalar), 2, 3, 4, 8, 16
529  * @param[in] store_size The number of lower elements to store. Supported values: 1-16, but has to be <= @p size
530  * @{
531  */
532 #define VSTORE_PARTIAL_STR(size, store_size) vstore_partial_##size##_##store_size
533 #define VSTORE_PARTIAL(size, store_size) VSTORE_PARTIAL_STR(size, store_size)
534 
535 #define NO_STORE(data, offs, ptr) \
536  { \
537  }
538 
539 // Size == 1 (scalar)
540 #define vstore_partial_1_0 NO_STORE
541 #define vstore_partial_1_1 vstore1
542 #define vstore_partial_1_2 NO_STORE
543 #define vstore_partial_1_3 NO_STORE
544 #define vstore_partial_1_4 NO_STORE
545 #define vstore_partial_1_5 NO_STORE
546 #define vstore_partial_1_6 NO_STORE
547 #define vstore_partial_1_7 NO_STORE
548 #define vstore_partial_1_8 NO_STORE
549 #define vstore_partial_1_9 NO_STORE
550 #define vstore_partial_1_10 NO_STORE
551 #define vstore_partial_1_11 NO_STORE
552 #define vstore_partial_1_12 NO_STORE
553 #define vstore_partial_1_13 NO_STORE
554 #define vstore_partial_1_14 NO_STORE
555 #define vstore_partial_1_15 NO_STORE
556 #define vstore_partial_1_16 NO_STORE
557 // Size == 2
558 #define vstore_partial_2_0 NO_STORE
559 #define vstore_partial_2_1 vstore_partial_1
560 #define vstore_partial_2_2 vstore_partial_2
561 #define vstore_partial_2_3 NO_STORE
562 #define vstore_partial_2_4 NO_STORE
563 #define vstore_partial_2_5 NO_STORE
564 #define vstore_partial_2_6 NO_STORE
565 #define vstore_partial_2_7 NO_STORE
566 #define vstore_partial_2_8 NO_STORE
567 #define vstore_partial_2_9 NO_STORE
568 #define vstore_partial_2_10 NO_STORE
569 #define vstore_partial_2_11 NO_STORE
570 #define vstore_partial_2_12 NO_STORE
571 #define vstore_partial_2_13 NO_STORE
572 #define vstore_partial_2_14 NO_STORE
573 #define vstore_partial_2_15 NO_STORE
574 #define vstore_partial_2_16 NO_STORE
575 // Size == 3
576 #define vstore_partial_3_0 NO_STORE
577 #define vstore_partial_3_1 vstore_partial_1
578 #define vstore_partial_3_2 vstore_partial_2
579 #define vstore_partial_3_3 vstore_partial_3
580 #define vstore_partial_3_4 NO_STORE
581 #define vstore_partial_3_5 NO_STORE
582 #define vstore_partial_3_6 NO_STORE
583 #define vstore_partial_3_7 NO_STORE
584 #define vstore_partial_3_8 NO_STORE
585 #define vstore_partial_3_9 NO_STORE
586 #define vstore_partial_3_10 NO_STORE
587 #define vstore_partial_3_11 NO_STORE
588 #define vstore_partial_3_12 NO_STORE
589 #define vstore_partial_3_13 NO_STORE
590 #define vstore_partial_3_14 NO_STORE
591 #define vstore_partial_3_15 NO_STORE
592 #define vstore_partial_3_16 NO_STORE
593 // Size == 4
594 #define vstore_partial_4_0 NO_STORE
595 #define vstore_partial_4_1 vstore_partial_1
596 #define vstore_partial_4_2 vstore_partial_2
597 #define vstore_partial_4_3 vstore_partial_3
598 #define vstore_partial_4_4 vstore_partial_4
599 #define vstore_partial_4_5 NO_STORE
600 #define vstore_partial_4_6 NO_STORE
601 #define vstore_partial_4_7 NO_STORE
602 #define vstore_partial_4_8 NO_STORE
603 #define vstore_partial_4_9 NO_STORE
604 #define vstore_partial_4_10 NO_STORE
605 #define vstore_partial_4_11 NO_STORE
606 #define vstore_partial_4_12 NO_STORE
607 #define vstore_partial_4_13 NO_STORE
608 #define vstore_partial_4_14 NO_STORE
609 #define vstore_partial_4_15 NO_STORE
610 #define vstore_partial_4_16 NO_STORE
611 // Size == 8
612 #define vstore_partial_8_0 NO_STORE
613 #define vstore_partial_8_1 vstore_partial_1
614 #define vstore_partial_8_2 vstore_partial_2
615 #define vstore_partial_8_3 vstore_partial_3
616 #define vstore_partial_8_4 vstore_partial_4
617 #define vstore_partial_8_5 vstore_partial_5
618 #define vstore_partial_8_6 vstore_partial_6
619 #define vstore_partial_8_7 vstore_partial_7
620 #define vstore_partial_8_8 vstore_partial_8
621 #define vstore_partial_8_9 NO_STORE
622 #define vstore_partial_8_10 NO_STORE
623 #define vstore_partial_8_11 NO_STORE
624 #define vstore_partial_8_12 NO_STORE
625 #define vstore_partial_8_13 NO_STORE
626 #define vstore_partial_8_14 NO_STORE
627 #define vstore_partial_8_15 NO_STORE
628 #define vstore_partial_8_16 NO_STORE
629 // Size == 16
630 #define vstore_partial_16_0 NO_STORE
631 #define vstore_partial_16_1 vstore_partial_1
632 #define vstore_partial_16_2 vstore_partial_2
633 #define vstore_partial_16_3 vstore_partial_3
634 #define vstore_partial_16_4 vstore_partial_4
635 #define vstore_partial_16_5 vstore_partial_5
636 #define vstore_partial_16_6 vstore_partial_6
637 #define vstore_partial_16_7 vstore_partial_7
638 #define vstore_partial_16_8 vstore_partial_8
639 #define vstore_partial_16_9 vstore_partial_9
640 #define vstore_partial_16_10 vstore_partial_10
641 #define vstore_partial_16_11 vstore_partial_11
642 #define vstore_partial_16_12 vstore_partial_12
643 #define vstore_partial_16_13 vstore_partial_13
644 #define vstore_partial_16_14 vstore_partial_14
645 #define vstore_partial_16_15 vstore_partial_15
646 #define vstore_partial_16_16 vstore_partial_16
647 
648 /** Partial vstore. Store the **lower** 0 to (n-1)th elements of the given vector while minimising the amount of vstore ops
649  * @name vstore_partial_n
650  *
651  * @note @p DATA needs to be a vector not a scalar
652  * @note n needs to be <= the vector width of the input variable @p DATA
653  * eg 1: Valid
654  * vstore_partial_15(var:float16, 0, 0xabcd);
655  * eg 2: Invalid
656  * vstore_partial_7(var:float4, 0, 0xabcd);
657  *
658  * @note in cases n == 1, 2, 3, 4, 8, 16, no extra vstore is invoked, thus there's no performance penalty.
659  *
660  * @param[in] DATA The name of the variable
661  * @param[in] OFFSET Offset in n
662  * @param[in] PTR The base pointer
663  * @{
664  */
665 #define vstore_partial_1(DATA, OFFSET, PTR) vstore1(DATA.s0, OFFSET, PTR);
666 
667 #define vstore_partial_2(DATA, OFFSET, PTR) vstore2(DATA.s01, OFFSET, PTR);
668 
669 #define vstore_partial_3(DATA, OFFSET, PTR) vstore3(DATA.s012, OFFSET, PTR);
670 
671 #define vstore_partial_4(DATA, OFFSET, PTR) vstore4(DATA.s0123, OFFSET, PTR);
672 
673 #define vstore_partial_5(DATA, OFFSET, PTR) \
674  vstore_partial_4(DATA.s0123, OFFSET, PTR); \
675  vstore1(DATA.s4, OFFSET, PTR + 4);
676 
677 #define vstore_partial_6(DATA, OFFSET, PTR) \
678  vstore_partial_4(DATA.s0123, OFFSET, PTR); \
679  vstore_partial_2(DATA.s45, OFFSET, PTR + 4);
680 
681 #define vstore_partial_7(DATA, OFFSET, PTR) \
682  vstore_partial_4(DATA.s0123, OFFSET, PTR); \
683  vstore_partial_3(DATA.s456, OFFSET, PTR + 4);
684 
685 #define vstore_partial_8(DATA, OFFSET, PTR) vstore8(DATA.s01234567, OFFSET, PTR);
686 
687 #define vstore_partial_9(DATA, OFFSET, PTR) \
688  vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
689  vstore1(DATA.s8, OFFSET, PTR + 8);
690 
691 #define vstore_partial_10(DATA, OFFSET, PTR) \
692  vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
693  vstore_partial_2(DATA.s89, OFFSET, PTR + 8);
694 
695 #define vstore_partial_11(DATA, OFFSET, PTR) \
696  vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
697  vstore_partial_3(DATA.s89a, OFFSET, PTR + 8);
698 
699 #define vstore_partial_12(DATA, OFFSET, PTR) \
700  vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
701  vstore_partial_4(DATA.s89ab, OFFSET, PTR + 8);
702 
703 #define vstore_partial_13(DATA, OFFSET, PTR) \
704  vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
705  vstore_partial_5(DATA.s89abcdef, OFFSET, PTR + 8);
706 
707 #define vstore_partial_14(DATA, OFFSET, PTR) \
708  vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
709  vstore_partial_6(DATA.s89abcdef, OFFSET, PTR + 8);
710 
711 #define vstore_partial_15(DATA, OFFSET, PTR) \
712  vstore_partial_8(DATA.s01234567, OFFSET, PTR); \
713  vstore_partial_7(DATA.s89abcdef, OFFSET, PTR + 8);
714 
715 #define vstore_partial_16(DATA, OFFSET, PTR) vstore16(DATA, OFFSET, PTR);
716 /** @} */ // end of groupd vstore_partial_n
717 /** @} */ // end of groupd VSTORE_PARTIAL
718 
719 // Convert built-in functions with _sat modifier are not supported in floating point so we create defines
720 // without _sat to overcome this issue
721 #define convert_float_sat convert_float
722 #define convert_float1_sat convert_float
723 #define convert_float2_sat convert_float2
724 #define convert_float3_sat convert_float3
725 #define convert_float4_sat convert_float4
726 #define convert_float8_sat convert_float8
727 #define convert_float16_sat convert_float16
728 #define convert_half_sat convert_float
729 #define convert_half1_sat convert_half
730 #define convert_half2_sat convert_half2
731 #define convert_half3_sat convert_half3
732 #define convert_half4_sat convert_half4
733 #define convert_half8_sat convert_half8
734 #define convert_half16_sat convert_half16
735 
736 #define convert_float1 convert_float
737 #define convert_half1 convert_half
738 #define convert_char1 convert_char
739 #define convert_uchar1 convert_uchar
740 #define convert_short1 convert_short
741 #define convert_ushort1 convert_ushort
742 #define convert_int1 convert_int
743 #define convert_uint1 convert_uint
744 #define convert_long1 convert_long
745 #define convert_ulong1 convert_ulong
746 #define convert_double1 convert_double
747 
748 #define convert_char1_sat convert_char_sat
749 #define convert_uchar1_sat convert_uchar_sat
750 #define convert_uchar2_sat convert_uchar2_sat
751 #define convert_uchar3_sat convert_uchar3_sat
752 #define convert_uchar4_sat convert_uchar4_sat
753 #define convert_uchar8_sat convert_uchar8_sat
754 #define convert_uchar16_sat convert_uchar16_sat
755 #define convert_short1_sat convert_short_sat
756 #define convert_ushort1_sat convert_ushort_sat
757 #define convert_int1_sat convert_int_sat
758 #define convert_uint1_sat convert_uint_sat
759 #define convert_long1_sat convert_long_sat
760 #define convert_ulong1_sat convert_ulong_sat
761 #define convert_double1_sat convert_double_sat
762 
763 #define VEC_DATA_TYPE_STR(type, size) type##size
764 #define VEC_DATA_TYPE(type, size) VEC_DATA_TYPE_STR(type, size)
765 
766 #define CONVERT_STR(x, type) (convert_##type((x)))
767 #define CONVERT(x, type) CONVERT_STR(x, type)
768 
769 #define CONVERT_SAT_STR(x, type) (convert_##type##_sat((x)))
770 #define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type)
771 
772 #define CONVERT_SAT_ROUND_STR(x, type, round) (convert_##type##_sat_##round((x)))
773 #define CONVERT_SAT_ROUND(x, type, round) CONVERT_SAT_ROUND_STR(x, type, round)
774 
775 #define select_vec_dt_uchar(size) uchar##size
776 #define select_vec_dt_char(size) char##size
777 #define select_vec_dt_ushort(size) ushort##size
778 #define select_vec_dt_short(size) short##size
779 #define select_vec_dt_half(size) short##size
780 #define select_vec_dt_uint(size) uint##size
781 #define select_vec_dt_int(size) int##size
782 #define select_vec_dt_float(size) int##size
783 #define select_vec_dt_ulong(size) ulong##size
784 #define select_vec_dt_long(size) long##size
785 
786 #define SELECT_VEC_DATA_TYPE_STR(type, size) select_vec_dt_##type(size)
787 #define SELECT_VEC_DATA_TYPE(type, size) SELECT_VEC_DATA_TYPE_STR(type, size)
788 #define SELECT_DATA_TYPE(type) SELECT_VEC_DATA_TYPE_STR(type, 1)
789 
790 #define signed_int_vec_dt_uchar(size) char##size
791 #define signed_int_vec_dt_char(size) char##size
792 #define signed_int_vec_dt_ushort(size) short##size
793 #define signed_int_vec_dt_short(size) short##size
794 #define signed_int_vec_dt_half(size) short##size
795 #define signed_int_vec_dt_uint(size) int##size
796 #define signed_int_vec_dt_int(size) int##size
797 #define signed_int_vec_dt_float(size) int##size
798 #define signed_int_vec_dt_ulong(size) long##size
799 #define signed_int_vec_dt_long(size) long##size
800 
801 #define SIGNED_INT_VEC_DATA_TYPE_STR(type, size) signed_int_vec_dt_##type(size)
802 #define SIGNED_INT_VEC_DATA_TYPE(type, size) SIGNED_INT_VEC_DATA_TYPE_STR(type, size)
803 #define SIGNED_INT_DATA_TYPE(type) SIGNED_INT_VEC_DATA_TYPE_STR(type, 1)
804 
805 #define sum_reduce_1(x) (x)
806 #define sum_reduce_2(x) ((x).s0) + ((x).s1)
807 #define sum_reduce_3(x) sum_reduce_2((x).s01) + ((x).s2)
808 #define sum_reduce_4(x) sum_reduce_2((x).s01) + sum_reduce_2((x).s23)
809 #define sum_reduce_8(x) sum_reduce_4((x).s0123) + sum_reduce_4((x).s4567)
810 #define sum_reduce_16(x) sum_reduce_8((x).s01234567) + sum_reduce_8((x).s89ABCDEF)
811 
812 #define SUM_REDUCE_STR(x, size) sum_reduce_##size(x)
813 #define SUM_REDUCE(x, size) SUM_REDUCE_STR(x, size)
814 
815 #define prod_reduce_1(x) (x)
816 #define prod_reduce_2(x) ((x).s0) * ((x).s1)
817 #define prod_reduce_3(x) prod_reduce_2((x).s01) * ((x).s2)
818 #define prod_reduce_4(x) prod_reduce_2((x).s01) * prod_reduce_2((x).s23)
819 #define prod_reduce_8(x) prod_reduce_4((x).s0123) * prod_reduce_4((x).s4567)
820 #define prod_reduce_16(x) prod_reduce_8((x).s01234567) * prod_reduce_8((x).s89ABCDEF)
821 
822 #define PROD_REDUCE_STR(x, size) prod_reduce_##size(x)
823 #define PROD_REDUCE(x, size) PROD_REDUCE_STR(x, size)
824 
825 #define max_reduce_1(x) (x)
826 #define max_reduce_2(x) max(((x).s0), ((x).s1))
827 #define max_reduce_3(x) max(max_reduce_2((x).s01), ((x).s2))
828 #define max_reduce_4(x) max(max_reduce_2((x).s01), max_reduce_2((x).s23))
829 #define max_reduce_8(x) max(max_reduce_4((x).s0123), max_reduce_4((x).s4567))
830 #define max_reduce_16(x) max(max_reduce_8((x).s01234567), max_reduce_8((x).s89ABCDEF))
831 
832 #define MAX_REDUCE_STR(x, size) max_reduce_##size(x)
833 #define MAX_REDUCE(x, size) MAX_REDUCE_STR(x, size)
834 
835 #define min_reduce_1(x) (x)
836 #define min_reduce_2(x) min(((x).s0), ((x).s1))
837 #define min_reduce_3(x) min(min_reduce_2((x).s01), ((x).s2))
838 #define min_reduce_4(x) min(min_reduce_2((x).s01), min_reduce_2((x).s23))
839 #define min_reduce_8(x) min(min_reduce_4((x).s0123), min_reduce_4((x).s4567))
840 #define min_reduce_16(x) min(min_reduce_8((x).s01234567), min_reduce_8((x).s89ABCDEF))
841 
842 #define MIN_REDUCE_STR(x, size) min_reduce_##size(x)
843 #define MIN_REDUCE(x, size) MIN_REDUCE_STR(x, size)
844 
845 #define VECTOR_DECLARATION(name) \
846  __global uchar *name##_ptr, uint name##_stride_x, uint name##_step_x, uint name##_offset_first_element_in_bytes
847 
848 #define IMAGE_DECLARATION(name) \
849  __global uchar *name##_ptr, uint name##_stride_x, uint name##_step_x, uint name##_stride_y, uint name##_step_y, \
850  uint name##_offset_first_element_in_bytes
851 
852 #define TENSOR3D_DECLARATION(name) \
853  __global uchar *name##_ptr, uint name##_stride_x, uint name##_step_x, uint name##_stride_y, uint name##_step_y, \
854  uint name##_stride_z, uint name##_step_z, uint name##_offset_first_element_in_bytes
855 
856 #define TENSOR4D_DECLARATION(name) \
857  __global uchar *name##_ptr, uint name##_stride_x, uint name##_step_x, uint name##_stride_y, uint name##_step_y, \
858  uint name##_stride_z, uint name##_step_z, uint name##_stride_w, uint name##_step_w, \
859  uint name##_offset_first_element_in_bytes
860 
861 #define TENSOR5D_DECLARATION(name) \
862  __global uchar *name##_ptr, uint name##_stride_x, uint name##_step_x, uint name##_stride_y, uint name##_step_y, \
863  uint name##_stride_z, uint name##_step_z, uint name##_stride_w, uint name##_step_w, uint name##_stride_v, \
864  uint name##_step_v, uint name##_offset_first_element_in_bytes
865 
866 #define CONVERT_TO_VECTOR_STRUCT(name) \
867  update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x)
868 
869 #define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name) \
870  update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0)
871 
872 #define CONVERT_TO_IMAGE_STRUCT(name) \
873  update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, \
874  name##_stride_y, name##_step_y)
875 
876 #define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name) \
877  update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0)
878 
879 #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \
880  update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, \
881  name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, \
882  name##_step_z)
883 
884 #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name) \
885  update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, \
886  name##_stride_y, 0, name##_stride_z, name##_step_z)
887 
888 #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \
889  update_image_from_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, \
890  name##_step_x, name##_stride_y, name##_step_y, name##_stride_z, \
891  name##_step_z)
892 
893 #define CONVERT_TO_TENSOR3D_STRUCT(name) \
894  update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, \
895  name##_stride_y, name##_step_y, name##_stride_z, name##_step_z)
896 
897 #define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name) \
898  update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, \
899  name##_stride_y, 0, name##_stride_z, 0)
900 
901 #define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size) \
902  update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, \
903  name##_stride_y, name##_step_y, name##_stride_z, name##_step_z, name##_stride_w, \
904  name##_step_w, mod_size)
905 
906 #define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size) \
907  update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, \
908  name##_stride_y, 0, name##_stride_z, 0, name##_stride_w, 0, mod_size)
909 
910 #define CONVERT_TO_TENSOR3D_STRUCT_NO_UPDATE_PTR(name) \
911  tensor3D_ptr_no_update(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, \
912  name##_stride_y, name##_step_y, name##_stride_z, name##_step_z)
913 
914 /** Structure to hold Vector information */
915 typedef struct Vector
916 {
917  __global uchar *ptr; /**< Pointer to the starting postion of the buffer */
918  int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
919  int stride_x; /**< Stride of the image in X dimension (in bytes) */
920 } Vector;
921 
922 /** Structure to hold Image information */
923 typedef struct Image
924 {
925  __global uchar *ptr; /**< Pointer to the starting postion of the buffer */
926  int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
927  int stride_x; /**< Stride of the image in X dimension (in bytes) */
928  int stride_y; /**< Stride of the image in Y dimension (in bytes) */
929 } Image;
930 
931 /** Structure to hold 3D tensor information */
932 typedef struct Tensor3D
933 {
934  __global uchar *ptr; /**< Pointer to the starting postion of the buffer */
935  int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
936  int stride_x; /**< Stride of the image in X dimension (in bytes) */
937  int stride_y; /**< Stride of the image in Y dimension (in bytes) */
938  int stride_z; /**< Stride of the image in Z dimension (in bytes) */
939 } Tensor3D;
940 
941 /** Structure to hold 4D tensor information */
942 typedef struct Tensor4D
943 {
944  __global uchar *ptr; /**< Pointer to the starting postion of the buffer */
945  int offset_first_element_in_bytes; /**< The offset of the first element in the source image */
946  int stride_x; /**< Stride of the image in X dimension (in bytes) */
947  int stride_y; /**< Stride of the image in Y dimension (in bytes) */
948  int stride_z; /**< Stride of the image in Z dimension (in bytes) */
949  int stride_w; /**< Stride of the image in W dimension (in bytes) */
950 } Tensor4D;
951 
952 /** Wrap vector information into an Vector structure, and make the pointer point at this workitem's data.
953  *
954  * @param[in] ptr Pointer to the starting postion of the buffer
955  * @param[in] offset_first_element_in_bytes The offset of the first element in the source vector
956  * @param[in] stride_x Stride of the vector in X dimension (in bytes)
957  * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
958  *
959  * @return An image object
960  */
961 inline Vector
962 update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x)
963 {
964  Vector vector = {
965  .ptr = ptr,
966  .offset_first_element_in_bytes = offset_first_element_in_bytes,
967  .stride_x = stride_x,
968  };
969  vector.ptr += vector.offset_first_element_in_bytes + get_global_id(0) * step_x;
970  return vector;
971 }
972 
973 /** Wrap image information into an Image structure, and make the pointer point at this workitem's data.
974  *
975  * @param[in] ptr Pointer to the starting postion of the buffer
976  * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
977  * @param[in] stride_x Stride of the image in X dimension (in bytes)
978  * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
979  * @param[in] stride_y Stride of the image in Y dimension (in bytes)
980  * @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
981  *
982  * @return An image object
983  */
984 inline Image update_image_workitem_ptr(
985  __global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x, uint stride_y, uint step_y)
986 {
987  Image img = {.ptr = ptr,
988  .offset_first_element_in_bytes = offset_first_element_in_bytes,
989  .stride_x = stride_x,
990  .stride_y = stride_y};
991  img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y;
992  return img;
993 }
994 
995 /** Wrap 3D tensor information into an image structure, and make the pointer point at this workitem's data.
996  *
997  * @param[in] ptr Pointer to the starting postion of the buffer
998  * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
999  * @param[in] stride_x Stride of the image in X dimension (in bytes)
1000  * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
1001  * @param[in] stride_y Stride of the image in Y dimension (in bytes)
1002  * @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
1003  * @param[in] stride_z Stride of the image in Z dimension (in bytes)
1004  * @param[in] step_z stride_z * number of elements along Z processed per workitem(in bytes)
1005  *
1006  * @return A 3D tensor object
1007  */
1008 inline Image update_image_from_tensor3D_workitem_ptr(__global uchar *ptr,
1009  uint offset_first_element_in_bytes,
1010  uint stride_x,
1011  uint step_x,
1012  uint stride_y,
1013  uint step_y,
1014  uint stride_z,
1015  uint step_z)
1016 {
1017  Image img = {.ptr = ptr,
1018  .offset_first_element_in_bytes = offset_first_element_in_bytes,
1019  .stride_x = stride_x,
1020  .stride_y = stride_y};
1021  img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y +
1022  get_global_id(2) * step_z;
1023  return img;
1024 }
1025 
1026 /** Wrap 3D tensor information into an tensor structure, and make the pointer point at this workitem's data.
1027  *
1028  * @param[in] ptr Pointer to the starting postion of the buffer
1029  * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
1030  * @param[in] stride_x Stride of the image in X dimension (in bytes)
1031  * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
1032  * @param[in] stride_y Stride of the image in Y dimension (in bytes)
1033  * @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
1034  * @param[in] stride_z Stride of the image in Z dimension (in bytes)
1035  * @param[in] step_z stride_z * number of elements along Z processed per workitem(in bytes)
1036  *
1037  * @return A 3D tensor object
1038  */
1039 inline Tensor3D update_tensor3D_workitem_ptr(__global uchar *ptr,
1040  uint offset_first_element_in_bytes,
1041  uint stride_x,
1042  uint step_x,
1043  uint stride_y,
1044  uint step_y,
1045  uint stride_z,
1046  uint step_z)
1047 {
1048  Tensor3D tensor = {.ptr = ptr,
1049  .offset_first_element_in_bytes = offset_first_element_in_bytes,
1050  .stride_x = stride_x,
1051  .stride_y = stride_y,
1052  .stride_z = stride_z};
1053  tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y +
1054  get_global_id(2) * step_z;
1055  return tensor;
1056 }
1057 
1058 /** Wrap 3D tensor information into an tensor structure.
1059  *
1060  * @param[in] ptr Pointer to the starting postion of the buffer
1061  * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
1062  * @param[in] stride_x Stride of the image in X dimension (in bytes)
1063  * @param[in] step_x stride_x * number of elements along X processed per workitem(in bytes)
1064  * @param[in] stride_y Stride of the image in Y dimension (in bytes)
1065  * @param[in] step_y stride_y * number of elements along Y processed per workitem(in bytes)
1066  * @param[in] stride_z Stride of the image in Z dimension (in bytes)
1067  * @param[in] step_z stride_z * number of elements along Z processed per workitem(in bytes)
1068  *
1069  * @return A 3D tensor object
1070  */
1071 inline Tensor3D tensor3D_ptr_no_update(__global uchar *ptr,
1072  uint offset_first_element_in_bytes,
1073  uint stride_x,
1074  uint step_x,
1075  uint stride_y,
1076  uint step_y,
1077  uint stride_z,
1078  uint step_z)
1079 {
1080  Tensor3D tensor = {.ptr = ptr,
1081  .offset_first_element_in_bytes = offset_first_element_in_bytes,
1082  .stride_x = stride_x,
1083  .stride_y = stride_y,
1084  .stride_z = stride_z};
1085  return tensor;
1086 }
1087 
1088 inline Tensor4D update_tensor4D_workitem_ptr(__global uchar *ptr,
1089  uint offset_first_element_in_bytes,
1090  uint stride_x,
1091  uint step_x,
1092  uint stride_y,
1093  uint step_y,
1094  uint stride_z,
1095  uint step_z,
1096  uint stride_w,
1097  uint step_w,
1098  uint mod_size)
1099 {
1100  Tensor4D tensor = {.ptr = ptr,
1101  .offset_first_element_in_bytes = offset_first_element_in_bytes,
1102  .stride_x = stride_x,
1103  .stride_y = stride_y,
1104  .stride_z = stride_z,
1105  .stride_w = stride_w};
1106 
1107  tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y +
1108  (get_global_id(2) % mod_size) * step_z + (get_global_id(2) / mod_size) * step_w;
1109  return tensor;
1110 }
1111 
1112 /** Get the pointer position of a Vector
1113  *
1114  * @param[in] vec Pointer to the starting position of the buffer
1115  * @param[in] x Relative X position
1116  */
1117 inline __global const uchar *vector_offset(const Vector *vec, int x)
1118 {
1119  return vec->ptr + x * vec->stride_x;
1120 }
1121 
1122 /** Get the pointer position of a Image
1123  *
1124  * @param[in] img Pointer to the starting position of the buffer
1125  * @param[in] x Relative X position
1126  * @param[in] y Relative Y position
1127  */
1128 inline __global uchar *offset(const Image *img, int x, int y)
1129 {
1130  return img->ptr + x * img->stride_x + y * img->stride_y;
1131 }
1132 
1133 /** Get the pointer position of a Tensor3D
1134  *
1135  * @param[in] tensor Pointer to the starting position of the buffer
1136  * @param[in] x Relative X position
1137  * @param[in] y Relative Y position
1138  * @param[in] z Relative Z position
1139  */
1140 inline __global const uchar *tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
1141 {
1142  return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z;
1143 }
1144 
1145 /** Get the pointer position of a Tensor4D
1146  *
1147  * @param[in] tensor Pointer to the starting position of the buffer
1148  * @param[in] x Relative X position
1149  * @param[in] y Relative Y position
1150  * @param[in] z Relative Z position
1151  * @param[in] w Relative W position
1152  */
1153 inline __global const uchar *tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
1154 {
1155  return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + w * tensor->stride_w;
1156 }
1157 
1158 /** Get the offset for a given linear index of a Tensor3D
1159  *
1160  * @param[in] tensor Pointer to the starting position of the buffer
1161  * @param[in] width Width of the input tensor
1162  * @param[in] height Height of the input tensor
1163  * @param[in] depth Depth of the input tensor
1164  * @param[in] index Linear index
1165  */
1166 inline __global const uchar *tensor3D_index2ptr(const Tensor3D *tensor, uint width, uint height, uint depth, uint index)
1167 {
1168  uint num_elements = width * height;
1169 
1170  const uint z = index / num_elements;
1171 
1172  index %= num_elements;
1173 
1174  const uint y = index / width;
1175 
1176  index %= width;
1177 
1178  const uint x = index;
1179 
1180  return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z +
1181  tensor->offset_first_element_in_bytes;
1182 }
1183 
1184 #endif // _HELPER_H
Tensor3D
struct Tensor3D Tensor3D
Structure to hold 3D tensor information.
Vector
Structure to hold Vector information.
Definition: helpers.h:915
Image::offset_first_element_in_bytes
int offset_first_element_in_bytes
The offset of the first element in the source image.
Definition: helpers.h:926
Vector::offset_first_element_in_bytes
int offset_first_element_in_bytes
The offset of the first element in the source image.
Definition: helpers.h:918
Image::stride_y
int stride_y
Stride of the image in Y dimension (in bytes)
Definition: helpers.h:928
tensor3D_index2ptr
const __global uchar * tensor3D_index2ptr(const Tensor3D *tensor, uint width, uint height, uint depth, uint index)
Get the offset for a given linear index of a Tensor3D.
Definition: helpers.h:1166
Image
struct Image Image
Structure to hold Image information.
Tensor4D::stride_x
int stride_x
Stride of the image in X dimension (in bytes)
Definition: helpers.h:946
Image::stride_x
int stride_x
Stride of the image in X dimension (in bytes)
Definition: helpers.h:927
Tensor3D::offset_first_element_in_bytes
int offset_first_element_in_bytes
The offset of the first element in the source image.
Definition: helpers.h:935
update_tensor4D_workitem_ptr
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, uint step_w, uint mod_size)
Definition: helpers.h:1088
Image
Structure to hold Image information.
Definition: helpers.h:923
Tensor4D::stride_z
int stride_z
Stride of the image in Z dimension (in bytes)
Definition: helpers.h:948
update_tensor3D_workitem_ptr
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)
Wrap 3D tensor information into an tensor structure, and make the pointer point at this workitem's da...
Definition: helpers.h:1039
Tensor3D
Structure to hold 3D tensor information.
Definition: helpers.h:932
update_vector_workitem_ptr
Vector update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x)
Wrap vector information into an Vector structure, and make the pointer point at this workitem's data.
Definition: helpers.h:962
Tensor4D
Structure to hold 4D tensor information.
Definition: helpers.h:942
arm_compute::test::validation::w
SimpleTensor< float > w
Definition: DFT.cpp:156
Tensor3D::stride_x
int stride_x
Stride of the image in X dimension (in bytes)
Definition: helpers.h:936
offset
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
Definition: helpers.h:1128
Tensor4D::stride_w
int stride_w
Stride of the image in W dimension (in bytes)
Definition: helpers.h:949
Tensor4D::stride_y
int stride_y
Stride of the image in Y dimension (in bytes)
Definition: helpers.h:947
Tensor3D::stride_z
int stride_z
Stride of the image in Z dimension (in bytes)
Definition: helpers.h:938
tensor
CLTensor * tensor
Pointer to the auxiliary tensor.
Definition: ClWorkloadRuntime.cpp:67
Vector::stride_x
int stride_x
Stride of the image in X dimension (in bytes)
Definition: helpers.h:919
tensor4D_offset
const __global uchar * tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
Get the pointer position of a Tensor4D.
Definition: helpers.h:1153
tensor3D_ptr_no_update
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)
Wrap 3D tensor information into an tensor structure.
Definition: helpers.h:1071
Image::ptr
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:925
Tensor3D::ptr
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:934
Vector
struct Vector Vector
Structure to hold Vector information.
update_image_from_tensor3D_workitem_ptr
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)
Wrap 3D tensor information into an image structure, and make the pointer point at this workitem's dat...
Definition: helpers.h:1008
Tensor3D::stride_y
int stride_y
Stride of the image in Y dimension (in bytes)
Definition: helpers.h:937
vector_offset
const __global uchar * vector_offset(const Vector *vec, int x)
Get the pointer position of a Vector.
Definition: helpers.h:1117
update_image_workitem_ptr
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)
Wrap image information into an Image structure, and make the pointer point at this workitem's data.
Definition: helpers.h:984
Tensor4D
struct Tensor4D Tensor4D
Structure to hold 4D tensor information.
load_store_utility.h
Tensor4D::offset_first_element_in_bytes
int offset_first_element_in_bytes
The offset of the first element in the source image.
Definition: helpers.h:945
tensor3D_offset
const __global uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
Get the pointer position of a Tensor3D.
Definition: helpers.h:1140
Vector::ptr
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:917
Tensor4D::ptr
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:944