ONE - On-device Neural Engine
All Data Structures Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Modules Pages
helpers.h File Reference

Go to the source code of this file.

Data Structures

struct  Vector
 
struct  Image
 
struct  Tensor3D
 
struct  Tensor4D
 

Macros

#define GPU_ARCH_MIDGARD   0x100
 
#define GPU_ARCH_BIFROST   0x200
 
#define CONCAT(a, b)   a##b
 
#define EXPAND(x)   x
 
#define CLAMP(x, min_val, max_val)   min(max(x, min_val), max_val)
 
#define VLOAD_STR(size)   vload##size
 
#define VLOAD(size)   VLOAD_STR(size)
 
#define VSTORE_STR(size)   vstore##size
 
#define VSTORE(size)   VSTORE_STR(size)
 
#define float1   float
 
#define half1   half
 
#define char1   char
 
#define uchar1   uchar
 
#define short1   short
 
#define ushort1   ushort
 
#define int1   int
 
#define uint1   uint
 
#define long1   long
 
#define ulong1   ulong
 
#define double1   double
 
#define vload1(OFFSET, PTR)   *(OFFSET + PTR)
 
#define vstore1(DATA, OFFSET, PTR)   *(OFFSET + PTR) = DATA
 
#define convert_float_sat   convert_float
 
#define convert_float1_sat   convert_float
 
#define convert_float2_sat   convert_float2
 
#define convert_float3_sat   convert_float3
 
#define convert_float4_sat   convert_float4
 
#define convert_float8_sat   convert_float8
 
#define convert_float16_sat   convert_float16
 
#define convert_half_sat   convert_float
 
#define convert_half1_sat   convert_half
 
#define convert_half2_sat   convert_half2
 
#define convert_half3_sat   convert_half3
 
#define convert_half4_sat   convert_half4
 
#define convert_half8_sat   convert_half8
 
#define convert_half16_sat   convert_half16
 
#define convert_float1   convert_float
 
#define convert_half1   convert_half
 
#define convert_char1   convert_char
 
#define convert_uchar1   convert_uchar
 
#define convert_short1   convert_short
 
#define convert_ushort1   convert_ushort
 
#define convert_int1   convert_int
 
#define convert_uint1   convert_uint
 
#define convert_long1   convert_long
 
#define convert_ulong1   convert_ulong
 
#define convert_double1   convert_double
 
#define convert_char1_sat   convert_char_sat
 
#define convert_uchar1_sat   convert_uchar_sat
 
#define convert_short1_sat   convert_short_sat
 
#define convert_ushort1_sat   convert_ushort_sat
 
#define convert_int1_sat   convert_int_sat
 
#define convert_uint1_sat   convert_uint_sat
 
#define convert_long1_sat   convert_long_sat
 
#define convert_ulong1_sat   convert_ulong_sat
 
#define convert_double1_sat   convert_double_sat
 
#define VEC_DATA_TYPE_STR(type, size)   type##size
 
#define VEC_DATA_TYPE(type, size)   VEC_DATA_TYPE_STR(type, size)
 
#define CL_VEC_DATA_TYPE_STR(type, size)   type##size
 
#define CL_VEC_DATA_TYPE(type, size)   CL_VEC_DATA_TYPE_STR(type, size)
 
#define CONVERT_STR(x, type)   (convert_##type((x)))
 
#define CONVERT(x, type)   CONVERT_STR(x, type)
 
#define CONVERT_SAT_STR(x, type)   (convert_##type##_sat((x)))
 
#define CONVERT_SAT(x, type)   CONVERT_SAT_STR(x, type)
 
#define CONVERT_SAT_ROUND_STR(x, type, round)   (convert_##type##_sat_##round((x)))
 
#define CONVERT_SAT_ROUND(x, type, round)   CONVERT_SAT_ROUND_STR(x, type, round)
 
#define VECTOR_DECLARATION(name)
 
#define IMAGE_DECLARATION(name)
 
#define TENSOR3D_DECLARATION(name)
 
#define TENSOR4D_DECLARATION(name)
 
#define CONVERT_TO_VECTOR_STRUCT(name)
 
#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name)    update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0)
 
#define CONVERT_TO_IMAGE_STRUCT(name)
 
#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name)
 
#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)
 
#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name)
 
#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)
 
#define CONVERT_TO_TENSOR3D_STRUCT(name)
 
#define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name)
 
#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size)
 
#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size)
 
REVn

REVn reverses the given vector whose size is n.

Parameters
[in]xThe vector to be reversed
Returns
The reversed vector
#define REV1(x)   ((x))
 
#define REV2(x)   ((x).s10)
 
#define REV3(x)   ((x).s210)
 
#define REV4(x)   ((x).s3210)
 
#define REV8(x)   ((x).s76543210)
 
#define REV16(x)   ((x).sFEDCBA9876543210)
 
REVERSE

Reverse the given vector.

Parameters
[in]xThe vector to be reversed
[in]sThe size of the vector
Returns
The reversed vector
#define REVERSE_STR(x, s)   REV##s((x))
 
#define REVERSE(x, s)   REVERSE_STR(x, s)
 
ROTs_n

Circular-right-shift (rotate-right) the vector of size s by the amount of n.

Parameters
[in]xThe vector to be shifted
Returns
The shifted vector
#define ROT1_0(x)   ((x))
 
#define ROT2_0(x)   ((x))
 
#define ROT2_1(x)   ((x).s10)
 
#define ROT3_0(x)   ((x))
 
#define ROT3_1(x)   ((x).s201)
 
#define ROT3_2(x)   ((x).s120)
 
#define ROT4_0(x)   ((x))
 
#define ROT4_1(x)   ((x).s3012)
 
#define ROT4_2(x)   ((x).s2301)
 
#define ROT4_3(x)   ((x).s1230)
 
#define ROT8_0(x)   ((x))
 
#define ROT8_1(x)   ((x).s70123456)
 
#define ROT8_2(x)   ((x).s67012345)
 
#define ROT8_3(x)   ((x).s56701234)
 
#define ROT8_4(x)   ((x).s45670123)
 
#define ROT8_5(x)   ((x).s34567012)
 
#define ROT8_6(x)   ((x).s23456701)
 
#define ROT8_7(x)   ((x).s12345670)
 
#define ROT16_0(x)   ((x))
 
#define ROT16_1(x)   ((x).sF0123456789ABCDE)
 
#define ROT16_2(x)   ((x).sEF0123456789ABCD)
 
#define ROT16_3(x)   ((x).sDEF0123456789ABC)
 
#define ROT16_4(x)   ((x).sCDEF0123456789AB)
 
#define ROT16_5(x)   ((x).sBCDEF0123456789A)
 
#define ROT16_6(x)   ((x).sABCDEF0123456789)
 
#define ROT16_7(x)   ((x).s9ABCDEF012345678)
 
#define ROT16_8(x)   ((x).s89ABCDEF01234567)
 
#define ROT16_9(x)   ((x).s789ABCDEF0123456)
 
#define ROT16_10(x)   ((x).s6789ABCDEF012345)
 
#define ROT16_11(x)   ((x).s56789ABCDEF01234)
 
#define ROT16_12(x)   ((x).s456789ABCDEF0123)
 
#define ROT16_13(x)   ((x).s3456789ABCDEF012)
 
#define ROT16_14(x)   ((x).s23456789ABCDEF01)
 
#define ROT16_15(x)   ((x).s123456789ABCDEF0)
 
ROTATE

Circular-right-shift (rotate-right) the given vector by the given amount.

Parameters
[in]xThe vector to be shifted
[in]sThe size of the vector
[in]nThe amount to be shifted
Returns
The shifted vector
#define ROTATE_STR(x, s, n)   ROT##s##_##n(x)
 
#define ROTATE(x, s, n)   ROTATE_STR(x, s, n)
 
V_OFFSn

Creates a vector of size n filled with offset values corresponding to the location of each element.

Parameters
[in]dtThe data type of the output vector
Returns
The vector filled with offset values
#define V_OFFS1(dt)   (dt)(0)
 
#define V_OFFS2(dt)   (dt)(0, 1)
 
#define V_OFFS3(dt)   (dt)(0, 1, 3)
 
#define V_OFFS4(dt)   (dt)(0, 1, 2, 3)
 
#define V_OFFS8(dt)   (dt)(0, 1, 2, 3, 4, 5, 6, 7)
 
#define V_OFFS16(dt)   (dt)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)
 
VEC_OFFS

Create a vector filled with offset values corresponding to the location of each element.

Parameters
[in]dtThe data type of the output vector
[in]sThe size of the output vector
Returns
The vector filled with offset values
#define VEC_OFFS_STR(dt, s)   V_OFFS##s(dt)
 
#define VEC_OFFS(dt, s)   VEC_OFFS_STR(dt, s)
 

Typedefs

typedef struct Vector Vector
 
typedef struct Image Image
 
typedef struct Tensor3D Tensor3D
 
typedef struct Tensor4D Tensor4D
 

Functions

Vector update_vector_workitem_ptr (__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x)
 
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)
 
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)
 
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)
 
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)
 
__global const uchar * vector_offset (const Vector *vec, int x)
 
__global uchar * offset (const Image *img, int x, int y)
 
__global const uchar * tensor3D_offset (const Tensor3D *tensor, int x, int y, int z)
 
__global const uchar * tensor4D_offset (const Tensor4D *tensor, int x, int y, int z, int w)
 

Macro Definition Documentation

◆ char1

#define char1   char

Definition at line 221 of file helpers.h.

◆ CL_VEC_DATA_TYPE

#define CL_VEC_DATA_TYPE (   type,
  size 
)    CL_VEC_DATA_TYPE_STR(type, size)

Definition at line 278 of file helpers.h.

◆ CL_VEC_DATA_TYPE_STR

#define CL_VEC_DATA_TYPE_STR (   type,
  size 
)    type##size

Definition at line 277 of file helpers.h.

◆ CLAMP

#define CLAMP (   x,
  min_val,
  max_val 
)    min(max(x, min_val), max_val)

Clamp the given value between an upper and lower bound.

Parameters
[in]xThe value to be clamped
[in]min_valThe lower bound
[in]max_valThe upper bound
Returns
The clamped value.

Definition at line 89 of file helpers.h.

◆ CONCAT

#define CONCAT (   a,
 
)    a##b

Concatenate two inputs.

Parameters
[in]aThe first input to be concatenated
[in]bThe second input to be concatenated
Returns
The concatenated output

Definition at line 71 of file helpers.h.

◆ CONVERT

#define CONVERT (   x,
  type 
)    CONVERT_STR(x, type)

Definition at line 281 of file helpers.h.

◆ convert_char1

#define convert_char1   convert_char

Definition at line 254 of file helpers.h.

◆ convert_char1_sat

#define convert_char1_sat   convert_char_sat

Definition at line 264 of file helpers.h.

◆ convert_double1

#define convert_double1   convert_double

Definition at line 262 of file helpers.h.

◆ convert_double1_sat

#define convert_double1_sat   convert_double_sat

Definition at line 272 of file helpers.h.

◆ convert_float1

#define convert_float1   convert_float

Definition at line 252 of file helpers.h.

◆ convert_float16_sat

#define convert_float16_sat   convert_float16

Definition at line 243 of file helpers.h.

◆ convert_float1_sat

#define convert_float1_sat   convert_float

Definition at line 238 of file helpers.h.

◆ convert_float2_sat

#define convert_float2_sat   convert_float2

Definition at line 239 of file helpers.h.

◆ convert_float3_sat

#define convert_float3_sat   convert_float3

Definition at line 240 of file helpers.h.

◆ convert_float4_sat

#define convert_float4_sat   convert_float4

Definition at line 241 of file helpers.h.

◆ convert_float8_sat

#define convert_float8_sat   convert_float8

Definition at line 242 of file helpers.h.

◆ convert_float_sat

#define convert_float_sat   convert_float

Definition at line 237 of file helpers.h.

◆ convert_half1

#define convert_half1   convert_half

Definition at line 253 of file helpers.h.

◆ convert_half16_sat

#define convert_half16_sat   convert_half16

Definition at line 250 of file helpers.h.

◆ convert_half1_sat

#define convert_half1_sat   convert_half

Definition at line 245 of file helpers.h.

◆ convert_half2_sat

#define convert_half2_sat   convert_half2

Definition at line 246 of file helpers.h.

◆ convert_half3_sat

#define convert_half3_sat   convert_half3

Definition at line 247 of file helpers.h.

◆ convert_half4_sat

#define convert_half4_sat   convert_half4

Definition at line 248 of file helpers.h.

◆ convert_half8_sat

#define convert_half8_sat   convert_half8

Definition at line 249 of file helpers.h.

◆ convert_half_sat

#define convert_half_sat   convert_float

Definition at line 244 of file helpers.h.

◆ convert_int1

#define convert_int1   convert_int

Definition at line 258 of file helpers.h.

◆ convert_int1_sat

#define convert_int1_sat   convert_int_sat

Definition at line 268 of file helpers.h.

◆ convert_long1

#define convert_long1   convert_long

Definition at line 260 of file helpers.h.

◆ convert_long1_sat

#define convert_long1_sat   convert_long_sat

Definition at line 270 of file helpers.h.

◆ CONVERT_SAT

#define CONVERT_SAT (   x,
  type 
)    CONVERT_SAT_STR(x, type)

Definition at line 284 of file helpers.h.

◆ CONVERT_SAT_ROUND

#define CONVERT_SAT_ROUND (   x,
  type,
  round 
)    CONVERT_SAT_ROUND_STR(x, type, round)

Definition at line 287 of file helpers.h.

◆ CONVERT_SAT_ROUND_STR

#define CONVERT_SAT_ROUND_STR (   x,
  type,
  round 
)    (convert_##type##_sat_##round((x)))

Definition at line 286 of file helpers.h.

◆ CONVERT_SAT_STR

#define CONVERT_SAT_STR (   x,
  type 
)    (convert_##type##_sat((x)))

Definition at line 283 of file helpers.h.

◆ convert_short1

#define convert_short1   convert_short

Definition at line 256 of file helpers.h.

◆ convert_short1_sat

#define convert_short1_sat   convert_short_sat

Definition at line 266 of file helpers.h.

◆ CONVERT_STR

#define CONVERT_STR (   x,
  type 
)    (convert_##type((x)))

Definition at line 280 of file helpers.h.

◆ CONVERT_TENSOR3D_TO_IMAGE_STRUCT [1/2]

#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT (   name)
Value:
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)
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)
Definition helpers.h:460

Definition at line 322 of file helpers.h.

357{
358 __global uchar *ptr;
359 int offset_first_element_in_bytes;
360 int stride_x;
361} Vector;
362
364typedef struct Image
365{
366 __global uchar *ptr;
368 int stride_x;
369 int stride_y;
370} Image;
371
373typedef struct Tensor3D
374{
375 __global uchar *ptr;
377 int stride_x;
378 int stride_y;
379 int stride_z;
380} Tensor3D;
381
383typedef struct Tensor4D
384{
385 __global uchar *ptr;
387 int stride_x;
388 int stride_y;
389 int stride_z;
390 int stride_w;
391} Tensor4D;
392
404inline Vector update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes,
405 uint stride_x, uint step_x)
406{
407 Vector vector = {
408 .ptr = ptr,
409 .offset_first_element_in_bytes = offset_first_element_in_bytes,
410 .stride_x = stride_x,
411 };
412 vector.ptr += vector.offset_first_element_in_bytes + get_global_id(0) * step_x;
413 return vector;
414}
415
430inline Image update_image_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes,
431 uint stride_x, uint step_x, uint stride_y, uint step_y)
432{
433 Image img = {.ptr = ptr,
434 .offset_first_element_in_bytes = offset_first_element_in_bytes,
435 .stride_x = stride_x,
436 .stride_y = stride_y};
437 img.ptr +=
438 img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y;
439 return img;
440}
441
459inline Image update_image_from_tensor3D_workitem_ptr(__global uchar *ptr,
460 uint offset_first_element_in_bytes,
461 uint stride_x, uint step_x, uint stride_y,
462 uint step_y, uint stride_z, uint step_z)
463{
464 Image img = {.ptr = ptr,
465 .offset_first_element_in_bytes = offset_first_element_in_bytes,
466 .stride_x = stride_x,
467 .stride_y = stride_y};
468 img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x +
469 get_global_id(1) * step_y + get_global_id(2) * step_z;
470 return img;
471}
472
490inline Tensor3D update_tensor3D_workitem_ptr(__global uchar *ptr,
491 uint offset_first_element_in_bytes, uint stride_x,
492 uint step_x, uint stride_y, uint step_y, uint stride_z,
493 uint step_z)
494{
495 Tensor3D tensor = {.ptr = ptr,
496 .offset_first_element_in_bytes = offset_first_element_in_bytes,
497 .stride_x = stride_x,
498 .stride_y = stride_y,
499 .stride_z = stride_z};
500 tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x +
501 get_global_id(1) * step_y + get_global_id(2) * step_z;
502 return tensor;
503}
504
505inline Tensor4D update_tensor4D_workitem_ptr(__global uchar *ptr,
506 uint offset_first_element_in_bytes, uint stride_x,
507 uint step_x, uint stride_y, uint step_y, uint stride_z,
508 uint step_z, uint stride_w, uint step_w, uint mod_size)
509{
510 Tensor4D tensor = {.ptr = ptr,
511 .offset_first_element_in_bytes = offset_first_element_in_bytes,
512 .stride_x = stride_x,
513 .stride_y = stride_y,
514 .stride_z = stride_z,
515 .stride_w = stride_w};
516
517 tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x +
518 get_global_id(1) * step_y + (get_global_id(2) % mod_size) * step_z +
519 (get_global_id(2) / mod_size) * step_w;
520 return tensor;
521}
522
528inline __global const uchar *vector_offset(const Vector *vec, int x)
529{
530 return vec->ptr + x * vec->stride_x;
531}
532
539inline __global uchar *offset(const Image *img, int x, int y)
540{
541 return img->ptr + x * img->stride_x + y * img->stride_y;
542}
543
551inline __global const uchar *tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
552{
553 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z;
554}
555
564inline __global const uchar *tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
565{
566 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z +
567 w * tensor->stride_w;
568}
569
570#endif // _HELPER_H
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)
Definition helpers.h:431
__global const uchar * tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
Definition helpers.h:565
__global uchar * offset(const Image *img, int x, int y)
Definition helpers.h:540
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:506
Vector update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x)
Definition helpers.h:405
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)
Definition helpers.h:491
__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
Definition helpers.h:552
__global const uchar * vector_offset(const Vector *vec, int x)
Definition helpers.h:529
int stride_x
Definition helpers.h:369
int offset_first_element_in_bytes
Definition helpers.h:368
__global uchar * ptr
Definition helpers.h:367
int stride_y
Definition helpers.h:370
int stride_y
Definition helpers.h:379
__global uchar * ptr
Definition helpers.h:376
int offset_first_element_in_bytes
Definition helpers.h:377
int stride_z
Definition helpers.h:380
int stride_x
Definition helpers.h:378
int offset_first_element_in_bytes
Definition helpers.h:387
int stride_z
Definition helpers.h:390
int stride_y
Definition helpers.h:389
int stride_w
Definition helpers.h:391
__global uchar * ptr
Definition helpers.h:386
int stride_x
Definition helpers.h:388
int offset_first_element_in_bytes
Definition helpers.h:360
__global uchar * ptr
Definition helpers.h:359
int stride_x
Definition helpers.h:361

◆ CONVERT_TENSOR3D_TO_IMAGE_STRUCT [2/2]

#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT (   name)
Value:
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)

Definition at line 322 of file helpers.h.

◆ CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP

#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP (   name)
Value:
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)

Definition at line 327 of file helpers.h.

◆ CONVERT_TO_IMAGE_STRUCT

#define CONVERT_TO_IMAGE_STRUCT (   name)
Value:
update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, \
name##_step_x, name##_stride_y, name##_step_y)

Definition at line 314 of file helpers.h.

◆ CONVERT_TO_IMAGE_STRUCT_NO_STEP

#define CONVERT_TO_IMAGE_STRUCT_NO_STEP (   name)
Value:
update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, \
name##_stride_y, 0)

Definition at line 318 of file helpers.h.

◆ CONVERT_TO_TENSOR3D_STRUCT

#define CONVERT_TO_TENSOR3D_STRUCT (   name)
Value:
update_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)

Definition at line 337 of file helpers.h.

◆ CONVERT_TO_TENSOR3D_STRUCT_NO_STEP

#define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP (   name)
Value:
update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, \
0, name##_stride_y, 0, name##_stride_z, 0)

Definition at line 342 of file helpers.h.

◆ CONVERT_TO_TENSOR4D_STRUCT

#define CONVERT_TO_TENSOR4D_STRUCT (   name,
  mod_size 
)
Value:
update_tensor4D_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, name##_stride_w, name##_step_w, mod_size)

Definition at line 346 of file helpers.h.

◆ CONVERT_TO_TENSOR4D_STRUCT_NO_STEP

#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP (   name,
  mod_size 
)
Value:
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)

Definition at line 351 of file helpers.h.

◆ CONVERT_TO_VECTOR_STRUCT

#define CONVERT_TO_VECTOR_STRUCT (   name)
Value:
update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, \
name##_step_x)

Definition at line 307 of file helpers.h.

◆ CONVERT_TO_VECTOR_STRUCT_NO_STEP

#define CONVERT_TO_VECTOR_STRUCT_NO_STEP (   name)     update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0)

Definition at line 311 of file helpers.h.

◆ convert_uchar1

#define convert_uchar1   convert_uchar

Definition at line 255 of file helpers.h.

◆ convert_uchar1_sat

#define convert_uchar1_sat   convert_uchar_sat

Definition at line 265 of file helpers.h.

◆ convert_uint1

#define convert_uint1   convert_uint

Definition at line 259 of file helpers.h.

◆ convert_uint1_sat

#define convert_uint1_sat   convert_uint_sat

Definition at line 269 of file helpers.h.

◆ convert_ulong1

#define convert_ulong1   convert_ulong

Definition at line 261 of file helpers.h.

◆ convert_ulong1_sat

#define convert_ulong1_sat   convert_ulong_sat

Definition at line 271 of file helpers.h.

◆ convert_ushort1

#define convert_ushort1   convert_ushort

Definition at line 257 of file helpers.h.

◆ convert_ushort1_sat

#define convert_ushort1_sat   convert_ushort_sat

Definition at line 267 of file helpers.h.

◆ double1

#define double1   double

Definition at line 229 of file helpers.h.

◆ EXPAND

#define EXPAND (   x)    x

Expand the given vector

Parameters
[in]xThe vector to be expanded
Returns
The expanded output

Definition at line 79 of file helpers.h.

◆ float1

#define float1   float

Definition at line 219 of file helpers.h.

◆ GPU_ARCH_BIFROST

#define GPU_ARCH_BIFROST   0x200

Definition at line 62 of file helpers.h.

◆ GPU_ARCH_MIDGARD

#define GPU_ARCH_MIDGARD   0x100

Definition at line 61 of file helpers.h.

◆ half1

#define half1   half

Definition at line 220 of file helpers.h.

◆ IMAGE_DECLARATION

#define IMAGE_DECLARATION (   name)
Value:
__global uchar *name##_ptr, uint name##_stride_x, uint name##_step_x, uint name##_stride_y, \
uint name##_step_y, uint name##_offset_first_element_in_bytes

Definition at line 293 of file helpers.h.

◆ int1

#define int1   int

Definition at line 225 of file helpers.h.

◆ long1

#define long1   long

Definition at line 227 of file helpers.h.

◆ REV1

#define REV1 (   x)    ((x))

Definition at line 99 of file helpers.h.

◆ REV16

#define REV16 (   x)    ((x).sFEDCBA9876543210)

Definition at line 104 of file helpers.h.

◆ REV2

#define REV2 (   x)    ((x).s10)

Definition at line 100 of file helpers.h.

◆ REV3

#define REV3 (   x)    ((x).s210)

Definition at line 101 of file helpers.h.

◆ REV4

#define REV4 (   x)    ((x).s3210)

Definition at line 102 of file helpers.h.

◆ REV8

#define REV8 (   x)    ((x).s76543210)

Definition at line 103 of file helpers.h.

◆ REVERSE

#define REVERSE (   x,
 
)    REVERSE_STR(x, s)

Definition at line 117 of file helpers.h.

◆ REVERSE_STR

#define REVERSE_STR (   x,
 
)    REV##s((x))

Definition at line 116 of file helpers.h.

◆ ROT16_0

#define ROT16_0 (   x)    ((x))

Definition at line 151 of file helpers.h.

◆ ROT16_1

#define ROT16_1 (   x)    ((x).sF0123456789ABCDE)

Definition at line 152 of file helpers.h.

◆ ROT16_10

#define ROT16_10 (   x)    ((x).s6789ABCDEF012345)

Definition at line 161 of file helpers.h.

◆ ROT16_11

#define ROT16_11 (   x)    ((x).s56789ABCDEF01234)

Definition at line 162 of file helpers.h.

◆ ROT16_12

#define ROT16_12 (   x)    ((x).s456789ABCDEF0123)

Definition at line 163 of file helpers.h.

◆ ROT16_13

#define ROT16_13 (   x)    ((x).s3456789ABCDEF012)

Definition at line 164 of file helpers.h.

◆ ROT16_14

#define ROT16_14 (   x)    ((x).s23456789ABCDEF01)

Definition at line 165 of file helpers.h.

◆ ROT16_15

#define ROT16_15 (   x)    ((x).s123456789ABCDEF0)

Definition at line 166 of file helpers.h.

◆ ROT16_2

#define ROT16_2 (   x)    ((x).sEF0123456789ABCD)

Definition at line 153 of file helpers.h.

◆ ROT16_3

#define ROT16_3 (   x)    ((x).sDEF0123456789ABC)

Definition at line 154 of file helpers.h.

◆ ROT16_4

#define ROT16_4 (   x)    ((x).sCDEF0123456789AB)

Definition at line 155 of file helpers.h.

◆ ROT16_5

#define ROT16_5 (   x)    ((x).sBCDEF0123456789A)

Definition at line 156 of file helpers.h.

◆ ROT16_6

#define ROT16_6 (   x)    ((x).sABCDEF0123456789)

Definition at line 157 of file helpers.h.

◆ ROT16_7

#define ROT16_7 (   x)    ((x).s9ABCDEF012345678)

Definition at line 158 of file helpers.h.

◆ ROT16_8

#define ROT16_8 (   x)    ((x).s89ABCDEF01234567)

Definition at line 159 of file helpers.h.

◆ ROT16_9

#define ROT16_9 (   x)    ((x).s789ABCDEF0123456)

Definition at line 160 of file helpers.h.

◆ ROT1_0

#define ROT1_0 (   x)    ((x))

Definition at line 128 of file helpers.h.

◆ ROT2_0

#define ROT2_0 (   x)    ((x))

Definition at line 130 of file helpers.h.

◆ ROT2_1

#define ROT2_1 (   x)    ((x).s10)

Definition at line 131 of file helpers.h.

◆ ROT3_0

#define ROT3_0 (   x)    ((x))

Definition at line 133 of file helpers.h.

◆ ROT3_1

#define ROT3_1 (   x)    ((x).s201)

Definition at line 134 of file helpers.h.

◆ ROT3_2

#define ROT3_2 (   x)    ((x).s120)

Definition at line 135 of file helpers.h.

◆ ROT4_0

#define ROT4_0 (   x)    ((x))

Definition at line 137 of file helpers.h.

◆ ROT4_1

#define ROT4_1 (   x)    ((x).s3012)

Definition at line 138 of file helpers.h.

◆ ROT4_2

#define ROT4_2 (   x)    ((x).s2301)

Definition at line 139 of file helpers.h.

◆ ROT4_3

#define ROT4_3 (   x)    ((x).s1230)

Definition at line 140 of file helpers.h.

◆ ROT8_0

#define ROT8_0 (   x)    ((x))

Definition at line 142 of file helpers.h.

◆ ROT8_1

#define ROT8_1 (   x)    ((x).s70123456)

Definition at line 143 of file helpers.h.

◆ ROT8_2

#define ROT8_2 (   x)    ((x).s67012345)

Definition at line 144 of file helpers.h.

◆ ROT8_3

#define ROT8_3 (   x)    ((x).s56701234)

Definition at line 145 of file helpers.h.

◆ ROT8_4

#define ROT8_4 (   x)    ((x).s45670123)

Definition at line 146 of file helpers.h.

◆ ROT8_5

#define ROT8_5 (   x)    ((x).s34567012)

Definition at line 147 of file helpers.h.

◆ ROT8_6

#define ROT8_6 (   x)    ((x).s23456701)

Definition at line 148 of file helpers.h.

◆ ROT8_7

#define ROT8_7 (   x)    ((x).s12345670)

Definition at line 149 of file helpers.h.

◆ ROTATE

#define ROTATE (   x,
  s,
 
)    ROTATE_STR(x, s, n)

Definition at line 180 of file helpers.h.

◆ ROTATE_STR

#define ROTATE_STR (   x,
  s,
 
)    ROT##s##_##n(x)

Definition at line 179 of file helpers.h.

◆ short1

#define short1   short

Definition at line 223 of file helpers.h.

◆ TENSOR3D_DECLARATION

#define TENSOR3D_DECLARATION (   name)
Value:
__global uchar *name##_ptr, uint name##_stride_x, uint name##_step_x, uint name##_stride_y, \
uint name##_step_y, uint name##_stride_z, uint name##_step_z, \
uint name##_offset_first_element_in_bytes

Definition at line 297 of file helpers.h.

◆ TENSOR4D_DECLARATION

#define TENSOR4D_DECLARATION (   name)
Value:
__global uchar *name##_ptr, uint name##_stride_x, uint name##_step_x, uint name##_stride_y, \
uint name##_step_y, uint name##_stride_z, uint name##_step_z, uint name##_stride_w, \
uint name##_step_w, uint name##_offset_first_element_in_bytes

Definition at line 302 of file helpers.h.

◆ uchar1

#define uchar1   uchar

Definition at line 222 of file helpers.h.

◆ uint1

#define uint1   uint

Definition at line 226 of file helpers.h.

◆ ulong1

#define ulong1   ulong

Definition at line 228 of file helpers.h.

◆ ushort1

#define ushort1   ushort

Definition at line 224 of file helpers.h.

◆ V_OFFS1

#define V_OFFS1 (   dt)    (dt)(0)

Definition at line 192 of file helpers.h.

◆ V_OFFS16

#define V_OFFS16 (   dt)    (dt)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)

Definition at line 197 of file helpers.h.

◆ V_OFFS2

#define V_OFFS2 (   dt)    (dt)(0, 1)

Definition at line 193 of file helpers.h.

◆ V_OFFS3

#define V_OFFS3 (   dt)    (dt)(0, 1, 3)

Definition at line 194 of file helpers.h.

◆ V_OFFS4

#define V_OFFS4 (   dt)    (dt)(0, 1, 2, 3)

Definition at line 195 of file helpers.h.

◆ V_OFFS8

#define V_OFFS8 (   dt)    (dt)(0, 1, 2, 3, 4, 5, 6, 7)

Definition at line 196 of file helpers.h.

◆ VEC_DATA_TYPE

#define VEC_DATA_TYPE (   type,
  size 
)    VEC_DATA_TYPE_STR(type, size)

Definition at line 275 of file helpers.h.

◆ VEC_DATA_TYPE_STR

#define VEC_DATA_TYPE_STR (   type,
  size 
)    type##size

Definition at line 274 of file helpers.h.

◆ VEC_OFFS

#define VEC_OFFS (   dt,
 
)    VEC_OFFS_STR(dt, s)

Definition at line 210 of file helpers.h.

◆ VEC_OFFS_STR

#define VEC_OFFS_STR (   dt,
 
)    V_OFFS##s(dt)

Definition at line 209 of file helpers.h.

◆ VECTOR_DECLARATION

#define VECTOR_DECLARATION (   name)
Value:
__global uchar *name##_ptr, uint name##_stride_x, uint name##_step_x, \
uint name##_offset_first_element_in_bytes

Definition at line 289 of file helpers.h.

◆ VLOAD

#define VLOAD (   size)    VLOAD_STR(size)

Definition at line 214 of file helpers.h.

◆ vload1

#define vload1 (   OFFSET,
  PTR 
)    *(OFFSET + PTR)

Definition at line 231 of file helpers.h.

◆ VLOAD_STR

#define VLOAD_STR (   size)    vload##size

Definition at line 213 of file helpers.h.

◆ VSTORE

#define VSTORE (   size)    VSTORE_STR(size)

Definition at line 217 of file helpers.h.

◆ vstore1

#define vstore1 (   DATA,
  OFFSET,
  PTR 
)    *(OFFSET + PTR) = DATA

Definition at line 232 of file helpers.h.

◆ VSTORE_STR

#define VSTORE_STR (   size)    vstore##size

Definition at line 216 of file helpers.h.

Typedef Documentation

◆ Image

typedef struct Image Image

Structure to hold Image information

◆ Tensor3D

typedef struct Tensor3D Tensor3D

Structure to hold 3D tensor information

◆ Tensor4D

typedef struct Tensor4D Tensor4D

Structure to hold 4D tensor information

◆ Vector

typedef struct Vector Vector

Structure to hold Vector information

Function Documentation

◆ offset()

__global uchar * offset ( const Image img,
int  x,
int  y 
)
inline

Get the pointer position of a Image

Parameters
[in]imgPointer to the starting position of the buffer
[in]xRelative X position
[in]yRelative Y position

Definition at line 540 of file helpers.h.

541{
542 return img->ptr + x * img->stride_x + y * img->stride_y;
543}

References Image::ptr, Image::stride_x, and Image::stride_y.

Referenced by luci_interpreter::StaticMemoryManager::allocate_memory(), ANeuralNetworksExecution_setInputFromMemory(), ANeuralNetworksExecution_setOutputFromMemory(), ANeuralNetworksMemory_createFromFd(), ANeuralNetworksModel_setOperandValueFromMemory(), onert::backend::acl_common::asQuantizationInfo(), Tensor::at(), coco::FeatureLayouts::BCHW::at(), coco::FeatureLayouts::BHWC::at(), coco::FeatureLayouts::BC::at(), mir::Tensor< T >::atOffset(), mir::Tensor< T >::atOffset(), mir::TensorVariant::atOffset(), onert::backend::IPortableTensor::calcOffset(), luci_interpreter::kernels::computePaddingWithOffset(), concatenationPrepare(), onert_micro::train::pal::Conv2DBiasGrad(), onert::backend::cpu::ops::DepthwiseConvolutionLayer::convQ8iHybridPerChannel(), flatbuffers::vector_downward::data_at(), dequantize_qasymm8(), dequantize_qasymm8_signed(), luci_interpreter::kernels::testing::dequantizeTensorData(), enco::CppCode::dump(), flexbuffers::FLATBUFFERS_FINAL_CLASS::Value::ElemWidth(), onert_micro::core::OMTrainingRuntimeModule::evaluateMetric(), onert_micro::core::train::OMTrainingHandler::evaluateMetric(), luci_interpreter::kernels::MirrorPad::execute(), luci_interpreter::execute_kernel_CircleMirrorPad(), luci_interpreter::execute_kernel_CircleSlice(), luci::exportNodes(), luci::CircleReader::file_data(), tflite2circle::CircleModel::finalize(), nnfw::cker::FullyConnectedHybrid(), souschef::ExplicitDataChef< std::string >::generate(), record_minmax::getMovingAverage(), mir::TensorVariant::getOffset(), onert::exec::feature::nchw::Reader< T >::getRef(), onert::exec::feature::nhwc::Reader< T >::getRef(), onert_micro::train::pal::GRUWeightGrads(), onert_micro::core::train::OMTrainingHandler::handleError(), flexbuffers::Indirect(), flexbuffers::Indirect(), onert::loader::BaseLoader< LoaderDomain >::loadOperand(), luci_interpreter::kernels::testing::makeInputTensor(), nnfw::cker::train::MaxPool2D(), nnfw::cker::train::MSEGrad(), record_minmax::DirectoryIterator::next(), nnfw::misc::vector::Object< T >::Object(), nnfw::cker::FusedBatchNorm::operator()(), Eigen::internal::gemm_pack_rhs< Scalar, Index, TensorContractionSubMapper< Scalar, Index, Rhs, TensorEvaluator< const TensorReshapingOp< NewDimension, const TensorImagePatchOp< Rows, Cols, ArgType > >, Device >, nocontract_t, contract_t, packet_size, inner_dim_contiguous, inner_dim_reordered, Alignment >, nr, ColMajor, false, false >::operator()(), Eigen::internal::gemm_pack_rhs< Scalar, Index, TensorContractionSubMapper< Scalar, Index, Rhs, TensorEvaluator< const TensorReshapingOp< NewDimension, const TensorImagePatchOp< Rows, Cols, ArgType > >, Device >, nocontract_t, contract_t, 2, inner_dim_contiguous, inner_dim_reordered, Alignment >, nr, ColMajor, false, false >::operator()(), Eigen::internal::gemm_pack_rhs< Scalar, Index, TensorContractionSubMapper< Scalar, Index, Rhs, TensorEvaluator< const TensorReshapingOp< NewDimension, const TensorImagePatchOp< Rows, Cols, ArgType > >, Device >, nocontract_t, contract_t, 1, inner_dim_contiguous, inner_dim_reordered, Alignment >, nr, ColMajor, false, false >::operator()(), flatbuffers::VectorIterator< T, IT >::operator+(), flatbuffers::VectorIterator< T, IT >::operator+=(), flatbuffers::VectorIterator< T, IT >::operator-(), flatbuffers::VectorIterator< T, IT >::operator-=(), nnfw::cker::PortableAsymmetricQuantizeFloats(), quantize_qasymm8(), tflite::reference_ops::RankOneSelect(), nnfw::cker::RankOneSelect(), nnfw::cker::ReducedOutputOffset(), onert::backend::train::FirstFitPlanner< Index >::release(), onert::backend::basic::FirstFitPlanner::release(), flatbuffers::vector_downward::release_raw(), flatbuffers::FlatBufferBuilder::ReleaseRaw(), flatbuffers::grpc::MessageBuilder::ReleaseRaw(), coco::Arg::reorder(), nnfw::cker::RoPE(), nnfw::cker::functor::FillPhiloxRandomTask< Distribution, true >::Run(), nnfw::cker::functor::FillPhiloxRandomTask< Distribution, false >::Run(), MappedMemory::set(), luci_interpreter::Tensor::set_offset(), ModelArgumentInfo::setFromMemory(), ExecutionBuilder::setInputFromMemory(), ModelBuilder::setOperandValueFromMemory(), ExecutionBuilder::setOutputFromMemory(), and onert_micro::core::OMTrainingRuntimeModule::trainSingleStep().

◆ tensor3D_offset()

__global const uchar * tensor3D_offset ( const Tensor3D tensor,
int  x,
int  y,
int  z 
)
inline

Get the pointer position of a Tensor3D

Parameters
[in]tensorPointer to the starting position of the buffer
[in]xRelative X position
[in]yRelative Y position
[in]zRelative Z position

Definition at line 552 of file helpers.h.

553{
554 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z;
555}

◆ tensor4D_offset()

__global const uchar * tensor4D_offset ( const Tensor4D tensor,
int  x,
int  y,
int  z,
int  w 
)
inline

Get the pointer position of a Tensor4D

Parameters
[in]tensorPointer to the starting position of the buffer
[in]xRelative X position
[in]yRelative Y position
[in]zRelative Z position
[in]wRelative W position

Definition at line 565 of file helpers.h.

566{
567 return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z +
568 w * tensor->stride_w;
569}

◆ 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 
)
inline

Wrap 3D tensor information into an image structure, and make the pointer point at this workitem's data.

Parameters
[in]ptrPointer to the starting postion of the buffer
[in]offset_first_element_in_bytesThe offset of the first element in the source image
[in]stride_xStride of the image in X dimension (in bytes)
[in]step_xstride_x * number of elements along X processed per workitem(in bytes)
[in]stride_yStride of the image in Y dimension (in bytes)
[in]step_ystride_y * number of elements along Y processed per workitem(in bytes)
[in]stride_zStride of the image in Z dimension (in bytes)
[in]step_zstride_z * number of elements along Z processed per workitem(in bytes)
Returns
A 3D tensor object

Definition at line 460 of file helpers.h.

464{
465 Image img = {.ptr = ptr,
466 .offset_first_element_in_bytes = offset_first_element_in_bytes,
467 .stride_x = stride_x,
468 .stride_y = stride_y};
469 img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x +
470 get_global_id(1) * step_y + get_global_id(2) * step_z;
471 return img;
472}

References Image::offset_first_element_in_bytes, and Image::ptr.

◆ 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 
)
inline

Wrap image information into an Image structure, and make the pointer point at this workitem's data.

Parameters
[in]ptrPointer to the starting postion of the buffer
[in]offset_first_element_in_bytesThe offset of the first element in the source image
[in]stride_xStride of the image in X dimension (in bytes)
[in]step_xstride_x * number of elements along X processed per workitem(in bytes)
[in]stride_yStride of the image in Y dimension (in bytes)
[in]step_ystride_y * number of elements along Y processed per workitem(in bytes)
Returns
An image object

Definition at line 431 of file helpers.h.

433{
434 Image img = {.ptr = ptr,
435 .offset_first_element_in_bytes = offset_first_element_in_bytes,
436 .stride_x = stride_x,
437 .stride_y = stride_y};
438 img.ptr +=
439 img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y;
440 return img;
441}

References Image::offset_first_element_in_bytes, and Image::ptr.

◆ 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 
)
inline

Wrap 3D tensor information into an tensor structure, and make the pointer point at this workitem's data.

Parameters
[in]ptrPointer to the starting postion of the buffer
[in]offset_first_element_in_bytesThe offset of the first element in the source image
[in]stride_xStride of the image in X dimension (in bytes)
[in]step_xstride_x * number of elements along X processed per workitem(in bytes)
[in]stride_yStride of the image in Y dimension (in bytes)
[in]step_ystride_y * number of elements along Y processed per workitem(in bytes)
[in]stride_zStride of the image in Z dimension (in bytes)
[in]step_zstride_z * number of elements along Z processed per workitem(in bytes)
Returns
A 3D tensor object

Definition at line 491 of file helpers.h.

495{
496 Tensor3D tensor = {.ptr = ptr,
497 .offset_first_element_in_bytes = offset_first_element_in_bytes,
498 .stride_x = stride_x,
499 .stride_y = stride_y,
500 .stride_z = stride_z};
501 tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x +
502 get_global_id(1) * step_y + get_global_id(2) * step_z;
503 return tensor;
504}

◆ 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 
)
inline

Definition at line 506 of file helpers.h.

510{
511 Tensor4D tensor = {.ptr = ptr,
512 .offset_first_element_in_bytes = offset_first_element_in_bytes,
513 .stride_x = stride_x,
514 .stride_y = stride_y,
515 .stride_z = stride_z,
516 .stride_w = stride_w};
517
518 tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x +
519 get_global_id(1) * step_y + (get_global_id(2) % mod_size) * step_z +
520 (get_global_id(2) / mod_size) * step_w;
521 return tensor;
522}

◆ update_vector_workitem_ptr()

Vector update_vector_workitem_ptr ( __global uchar *  ptr,
uint  offset_first_element_in_bytes,
uint  stride_x,
uint  step_x 
)
inline

Wrap vector information into an Vector structure, and make the pointer point at this workitem's data.

Parameters
[in]ptrPointer to the starting postion of the buffer
[in]offset_first_element_in_bytesThe offset of the first element in the source vector
[in]stride_xStride of the vector in X dimension (in bytes)
[in]step_xstride_x * number of elements along X processed per workitem(in bytes)
Returns
An image object

Definition at line 405 of file helpers.h.

407{
408 Vector vector = {
409 .ptr = ptr,
410 .offset_first_element_in_bytes = offset_first_element_in_bytes,
411 .stride_x = stride_x,
412 };
413 vector.ptr += vector.offset_first_element_in_bytes + get_global_id(0) * step_x;
414 return vector;
415}

References Vector::offset_first_element_in_bytes, and Vector::ptr.

◆ vector_offset()

__global const uchar * vector_offset ( const Vector vec,
int  x 
)
inline

Get the pointer position of a Vector

Parameters
[in]vecPointer to the starting position of the buffer
[in]xRelative X position

Definition at line 529 of file helpers.h.

530{
531 return vec->ptr + x * vec->stride_x;
532}

References Vector::ptr, and Vector::stride_x.