src/core/CL/cl_kernels/helpers.h

   1 /*
   2  * Copyright (c) 2016-2018 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 #if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED)
  28 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
  29 #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED)
  30
  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)
  36
  37 #define EXPAND(x) x
  38
  39 #define CLAMP(x, min_val, max_val) min(max(x, min_val), max_val)
  40
  41 #define VLOAD_STR(size) vload##size
  42 #define VLOAD(size) VLOAD_STR(size)
  43
  44 #define VSTORE_STR(size) vstore##size
  45 #define VSTORE(size) VSTORE_STR(size)
  46
  47 #define VEC_DATA_TYPE_STR(type, size) type##size
  48 #define VEC_DATA_TYPE(type, size) VEC_DATA_TYPE_STR(type, size)
  49
  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)
  52
  53 #define CONVERT_STR(x, type) (convert_##type((x)))
  54 #define CONVERT(x, type) CONVERT_STR(x, type)
  55
  56 #define CONVERT_SAT_STR(x, type) (convert_##type##_sat((x)))
  57 #define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type)
  58
  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)
  61
  62 #define VECTOR_DECLARATION(name)     \
  63     __global uchar *name##_ptr,      \
  64     uint        name##_stride_x, \
  65     uint        name##_step_x,   \
  66     uint        name##_offset_first_element_in_bytes
  67
  68 #define IMAGE_DECLARATION(name)      \
  69     __global uchar *name##_ptr,      \
  70     uint        name##_stride_x, \
  71     uint        name##_step_x,   \
  72     uint        name##_stride_y, \
  73     uint        name##_step_y,   \
  74     uint        name##_offset_first_element_in_bytes
  75
  76 #define TENSOR3D_DECLARATION(name)   \
  77     __global uchar *name##_ptr,      \
  78     uint        name##_stride_x, \
  79     uint        name##_step_x,   \
  80     uint        name##_stride_y, \
  81     uint        name##_step_y,   \
  82     uint        name##_stride_z, \
  83     uint        name##_step_z,   \
  84     uint        name##_offset_first_element_in_bytes
  85
  86 #define TENSOR4D_DECLARATION(name)   \
  87     __global uchar *name##_ptr,      \
  88     uint        name##_stride_x, \
  89     uint        name##_step_x,   \
  90     uint        name##_stride_y, \
  91     uint        name##_step_y,   \
  92     uint        name##_stride_z, \
  93     uint        name##_step_z,   \
  94     uint        name##_stride_w, \
  95     uint        name##_step_w,   \
  96     uint        name##_offset_first_element_in_bytes
  97
  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)
 100
 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)
 103
 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)
 106
 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)
 109
 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)
 112
 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)
 115
 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)
 118
 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)
 122
 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)
 125
 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)
 129
 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)
 132
 133 /** Structure to hold Vector information */
 134 typedef struct Vector
 135 {
 136     __global uchar *ptr;                           /**< Pointer to the starting postion of the buffer */
 137     int             offset_first_element_in_bytes; /**< The offset of the first element in the source image */
 138     int             stride_x;                      /**< Stride of the image in X dimension (in bytes) */
 139 } Vector;
 140
 141 /** Structure to hold Image information */
 142 typedef struct Image
 143 {
 144     __global uchar *ptr;                           /**< Pointer to the starting postion of the buffer */
 145     int             offset_first_element_in_bytes; /**< The offset of the first element in the source image */
 146     int             stride_x;                      /**< Stride of the image in X dimension (in bytes) */
 147     int             stride_y;                      /**< Stride of the image in Y dimension (in bytes) */
 148 } Image;
 149
 150 /** Structure to hold 3D tensor information */
 151 typedef struct Tensor3D
 152 {
 153     __global uchar *ptr;                           /**< Pointer to the starting postion of the buffer */
 154     int             offset_first_element_in_bytes; /**< The offset of the first element in the source image */
 155     int             stride_x;                      /**< Stride of the image in X dimension (in bytes) */
 156     int             stride_y;                      /**< Stride of the image in Y dimension (in bytes) */
 157     int             stride_z;                      /**< Stride of the image in Z dimension (in bytes) */
 158 } Tensor3D;
 159
 160 /** Structure to hold 4D tensor information */
 161 typedef struct Tensor4D
 162 {
 163     __global uchar *ptr;                           /**< Pointer to the starting postion of the buffer */
 164     int             offset_first_element_in_bytes; /**< The offset of the first element in the source image */
 165     int             stride_x;                      /**< Stride of the image in X dimension (in bytes) */
 166     int             stride_y;                      /**< Stride of the image in Y dimension (in bytes) */
 167     int             stride_z;                      /**< Stride of the image in Z dimension (in bytes) */
 168     int             stride_w;                      /**< Stride of the image in W dimension (in bytes) */
 169 } Tensor4D;
 170
 171 /** Wrap vector information into an Vector structure, and make the pointer point at this workitem's data.
 172  *
 173  * @param[in] ptr                           Pointer to the starting postion of the buffer
 174  * @param[in] offset_first_element_in_bytes The offset of the first element in the source vector
 175  * @param[in] stride_x                      Stride of the vector in X dimension (in bytes)
 176  * @param[in] step_x                        stride_x * number of elements along X processed per workitem(in bytes)
 177  *
 178  * @return An image object
 179  */
 180 Vector inline update_vector_workitem_ptr(__global uchar *ptr, uint offset_first_element_in_bytes, uint stride_x, uint step_x)
 181 {
 182     Vector vector =
 183     {
 184         .ptr                           = ptr,
 185         .offset_first_element_in_bytes = offset_first_element_in_bytes,
 186         .stride_x                      = stride_x,
 187     };
 188     vector.ptr += vector.offset_first_element_in_bytes + get_global_id(0) * step_x;
 189     return vector;
 190 }
 191
 192 /** Wrap image information into an Image structure, and make the pointer point at this workitem's data.
 193  *
 194  * @param[in] ptr                           Pointer to the starting postion of the buffer
 195  * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
 196  * @param[in] stride_x                      Stride of the image in X dimension (in bytes)
 197  * @param[in] step_x                        stride_x * number of elements along X processed per workitem(in bytes)
 198  * @param[in] stride_y                      Stride of the image in Y dimension (in bytes)
 199  * @param[in] step_y                        stride_y * number of elements along Y processed per workitem(in bytes)
 200  *
 201  * @return An image object
 202  */
 203 Image inline 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)
 204 {
 205     Image img =
 206     {
 207         .ptr                           = ptr,
 208         .offset_first_element_in_bytes = offset_first_element_in_bytes,
 209         .stride_x                      = stride_x,
 210         .stride_y                      = stride_y
 211     };
 212     img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y;
 213     return img;
 214 }
 215
 216 /** Wrap 3D tensor information into an image structure, and make the pointer point at this workitem's data.
 217  *
 218  * @param[in] ptr                           Pointer to the starting postion of the buffer
 219  * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
 220  * @param[in] stride_x                      Stride of the image in X dimension (in bytes)
 221  * @param[in] step_x                        stride_x * number of elements along X processed per workitem(in bytes)
 222  * @param[in] stride_y                      Stride of the image in Y dimension (in bytes)
 223  * @param[in] step_y                        stride_y * number of elements along Y processed per workitem(in bytes)
 224  * @param[in] stride_z                      Stride of the image in Z dimension (in bytes)
 225  * @param[in] step_z                        stride_z * number of elements along Z processed per workitem(in bytes)
 226  *
 227  * @return A 3D tensor object
 228  */
 229 Image inline 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)
 230 {
 231     Image img =
 232     {
 233         .ptr                           = ptr,
 234         .offset_first_element_in_bytes = offset_first_element_in_bytes,
 235         .stride_x                      = stride_x,
 236         .stride_y                      = stride_y
 237     };
 238     img.ptr += img.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z;
 239     return img;
 240 }
 241
 242 /** Wrap 3D tensor information into an tensor structure, and make the pointer point at this workitem's data.
 243  *
 244  * @param[in] ptr                           Pointer to the starting postion of the buffer
 245  * @param[in] offset_first_element_in_bytes The offset of the first element in the source image
 246  * @param[in] stride_x                      Stride of the image in X dimension (in bytes)
 247  * @param[in] step_x                        stride_x * number of elements along X processed per workitem(in bytes)
 248  * @param[in] stride_y                      Stride of the image in Y dimension (in bytes)
 249  * @param[in] step_y                        stride_y * number of elements along Y processed per workitem(in bytes)
 250  * @param[in] stride_z                      Stride of the image in Z dimension (in bytes)
 251  * @param[in] step_z                        stride_z * number of elements along Z processed per workitem(in bytes)
 252  *
 253  * @return A 3D tensor object
 254  */
 255 Tensor3D inline 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)
 256 {
 257     Tensor3D tensor =
 258     {
 259         .ptr                           = ptr,
 260         .offset_first_element_in_bytes = offset_first_element_in_bytes,
 261         .stride_x                      = stride_x,
 262         .stride_y                      = stride_y,
 263         .stride_z                      = stride_z
 264     };
 265     tensor.ptr += tensor.offset_first_element_in_bytes + get_global_id(0) * step_x + get_global_id(1) * step_y + get_global_id(2) * step_z;
 266     return tensor;
 267 }
 268
 269 Tensor4D inline 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,
 270                                              uint step_w,
 271                                              uint mod_size)
 272 {
 273     Tensor4D tensor =
 274     {
 275         .ptr                           = ptr,
 276         .offset_first_element_in_bytes = offset_first_element_in_bytes,
 277         .stride_x                      = stride_x,
 278         .stride_y                      = stride_y,
 279         .stride_z                      = stride_z,
 280         .stride_w                      = stride_w
 281     };
 282
 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;
 284     return tensor;
 285 }
 286
 287 /** Get the pointer position of a Vector
 288  *
 289  * @param[in] vec Pointer to the starting position of the buffer
 290  * @param[in] x   Relative X position
 291  */
 292 __global inline const uchar *vector_offset(const Vector *vec, int x)
 293 {
 294     return vec->ptr + x * vec->stride_x;
 295 }
 296
 297 /** Get the pointer position of a Image
 298  *
 299  * @param[in] img Pointer to the starting position of the buffer
 300  * @param[in] x   Relative X position
 301  * @param[in] y   Relative Y position
 302  */
 303 __global inline uchar *offset(const Image *img, int x, int y)
 304 {
 305     return img->ptr + x * img->stride_x + y * img->stride_y;
 306 }
 307
 308 /** Get the pointer position of a Tensor3D
 309  *
 310  * @param[in] tensor Pointer to the starting position of the buffer
 311  * @param[in] x      Relative X position
 312  * @param[in] y      Relative Y position
 313  * @param[in] z      Relative Z position
 314  */
 315 __global inline const uchar *tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
 316 {
 317     return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z;
 318 }
 319
 320 /** Get the pointer position of a Tensor4D
 321  *
 322  * @param[in] tensor Pointer to the starting position of the buffer
 323  * @param[in] x      Relative X position
 324  * @param[in] y      Relative Y position
 325  * @param[in] z      Relative Z position
 326  * @param[in] w      Relative W position
 327  */
 328 __global inline const uchar *tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
 329 {
 330     return tensor->ptr + x * tensor->stride_x + y * tensor->stride_y + z * tensor->stride_z + w * tensor->stride_w;
 331 }
 332
 333 #endif // _HELPER_H