24 #ifndef ARM_COMPUTE_HELPER_H 25 #define ARM_COMPUTE_HELPER_H 27 #if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) 28 #pragma OPENCL EXTENSION cl_khr_fp16 : enable 29 #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) 31 #if defined(ARM_COMPUTE_DEBUG_ENABLED) 32 #if defined(cl_arm_printf) 33 #pragma OPENCL EXTENSION cl_arm_printf : enable 34 #endif // defined(cl_arm_printf) 35 #endif // defined(ARM_COMPUTE_DEBUG_ENABLED) 39 #define CLAMP(x, min_val, max_val) min(max(x, min_val), max_val) 41 #define VLOAD_STR(size) vload##size 42 #define VLOAD(size) VLOAD_STR(size) 44 #define VSTORE_STR(size) vstore##size 45 #define VSTORE(size) VSTORE_STR(size) 47 #define VEC_DATA_TYPE_STR(type, size) type##size 48 #define VEC_DATA_TYPE(type, size) VEC_DATA_TYPE_STR(type, size) 50 #define CL_VEC_DATA_TYPE_STR(type, size) type##size 51 #define CL_VEC_DATA_TYPE(type, size) CL_VEC_DATA_TYPE_STR(type, size) 53 #define CONVERT_STR(x, type) (convert_##type((x))) 54 #define CONVERT(x, type) CONVERT_STR(x, type) 56 #define CONVERT_SAT_STR(x, type) (convert_##type##_sat((x))) 57 #define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type) 59 #define CONVERT_SAT_ROUND_STR(x, type, round) (convert_##type##_sat_##round((x))) 60 #define CONVERT_SAT_ROUND(x, type, round) CONVERT_SAT_ROUND_STR(x, type, round) 62 #define VECTOR_DECLARATION(name) \ 63 __global uchar *name##_ptr, \ 64 uint name##_stride_x, \ 66 uint name##_offset_first_element_in_bytes 68 #define IMAGE_DECLARATION(name) \ 69 __global uchar *name##_ptr, \ 70 uint name##_stride_x, \ 72 uint name##_stride_y, \ 74 uint name##_offset_first_element_in_bytes 76 #define TENSOR3D_DECLARATION(name) \ 77 __global uchar *name##_ptr, \ 78 uint name##_stride_x, \ 80 uint name##_stride_y, \ 82 uint name##_stride_z, \ 84 uint name##_offset_first_element_in_bytes 86 #define TENSOR4D_DECLARATION(name) \ 87 __global uchar *name##_ptr, \ 88 uint name##_stride_x, \ 90 uint name##_stride_y, \ 92 uint name##_stride_z, \ 94 uint name##_stride_w, \ 96 uint name##_offset_first_element_in_bytes 98 #define CONVERT_TO_VECTOR_STRUCT(name) \ 99 update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x) 101 #define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name) \ 102 update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0) 104 #define CONVERT_TO_IMAGE_STRUCT(name) \ 105 update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y) 107 #define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name) \ 108 update_image_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0) 110 #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \ 111 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) 113 #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(name) \ 114 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) 116 #define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name) \ 117 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) 119 #define CONVERT_TO_TENSOR3D_STRUCT(name) \ 120 update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \ 121 name##_stride_z, name##_step_z) 123 #define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name) \ 124 update_tensor3D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, 0, name##_stride_y, 0, name##_stride_z, 0) 126 #define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size) \ 127 update_tensor4D_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x, name##_stride_y, name##_step_y, \ 128 name##_stride_z, name##_step_z, name##_stride_w, name##_step_w, mod_size) 130 #define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size) \ 131 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) 262 .stride_y = stride_y,
278 .stride_y = stride_y,
279 .stride_z = stride_z,
283 tensor.
ptr += tensor.
offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + (get_global_id(2) % mod_size) * step_z + (get_global_id(2) / mod_size) * step_w;
Structure to hold Vector information.
struct Tensor4D Tensor4D
Structure to hold 4D tensor information.
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...
__global uchar * ptr
Pointer to the starting postion of the buffer.
int stride_z
Stride of the image in Z dimension (in bytes)
int stride_x
Stride of the image in X dimension (in bytes)
int offset_first_element_in_bytes
The offset of the first element in the source image.
Structure to hold 3D tensor information.
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
struct Tensor3D Tensor3D
Structure to hold 3D tensor information.
Structure to hold 4D tensor information.
__global const uchar * tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
Get the pointer position of a Tensor4D.
int stride_w
Stride of the image in W dimension (in bytes)
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...
int stride_x
Stride of the image in X dimension (in bytes)
struct Image Image
Structure to hold Image information.
__global uchar * ptr
Pointer to the starting postion of the buffer.
__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
Get the pointer position of a Tensor3D.
int stride_y
Stride of the image in Y dimension (in bytes)
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...
Structure to hold Image information.
int offset_first_element_in_bytes
The offset of the first element in the source image.
int offset_first_element_in_bytes
The offset of the first element in the source image.
__global uchar * ptr
Pointer to the starting postion of the buffer.
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...
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)
Get the pointer position of a Vector.
__global uchar * ptr
Pointer to the starting postion of the buffer.
int stride_x
Stride of the image in X dimension (in bytes)
struct Vector Vector
Structure to hold Vector information.
int stride_y
Stride of the image in Y dimension (in bytes)
int stride_z
Stride of the image in Z dimension (in bytes)
int offset_first_element_in_bytes
The offset of the first element in the source image.
int stride_y
Stride of the image in Y dimension (in bytes)
int stride_x
Stride of the image in X dimension (in bytes)