src/core/CL/cl_kernels/convolution_layer.cl

   1 /*
   2  * Copyright (c) 2017 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 #include "helpers.h"
  25
  26 #if defined(FIXED_POINT_POSITION)
  27 #include "fixed_point.h"
  28 #endif // FIXED_POINT_POSITION
  29
  30 /** This kernel reshapes the tensor's low three dimensions to single column
  31  *
  32  * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
  33  *
  34  * @param[in]  src_ptr                            Pointer to the source tensor. Supported data types: F16/F32
  35  * @param[in]  src_stride_x                       Stride of the source tensor in X dimension (in bytes)
  36  * @param[in]  src_step_x                         src_stride_x * number of elements along X processed per workitem(in bytes)
  37  * @param[in]  src_stride_y                       Stride of the source tensor in Y dimension (in bytes)
  38  * @param[in]  src_step_y                         src_stride_y * number of elements along Y processed per workitem(in bytes)
  39  * @param[in]  src_stride_z                       Stride of the source tensor in Z dimension (in bytes)
  40  * @param[in]  src_step_z                         src_stride_z * number of elements along Y processed per workitem(in bytes)
  41  * @param[in]  src_offset_first_element_in_bytes  The offset of the first element in the source tensor
  42  * @param[out] dst_ptr                            Pointer to the destination tensor. Same as @p src_ptr
  43  * @param[in]  dst_stride_x                       Stride of the destination tensor in X dimension (in bytes)
  44  * @param[in]  dst_step_x                         dst_stride_x * number of elements along X processed per workitem(in bytes)
  45  * @param[in]  dst_stride_y                       Stride of the destination tensor in Y dimension (in bytes)
  46  * @param[in]  dst_step_y                         dst_stride_y * number of elements along Y processed per workitem(in bytes)
  47  * @param[in]  dst_offset_first_element_in_bytes  The offset of the first element in the destination tensor
  48  * @param[in]  bias_ptr                           Pointer to the bias tensor. Same as @p src_ptr
  49  * @param[in]  bias_stride_x                      Stride of the bias tensor in X dimension (in bytes)
  50  * @param[in]  bias_step_x                        bias_stride_x * number of elements along X processed per workitem(in bytes)
  51  * @param[in]  bias_offset_first_element_in_bytes The offset of the first element in the source tensor
  52  * @param[in]  width                              The width of the input tensor
  53  * @param[in]  height                             The height of the input tensor
  54  * @param[in]  depth                              The depth of the input tensor
  55  * @param[in]  total_filters                      Total number of filters. 4th dimension of the weights matrix
  56  */
  57 __kernel void reshape_to_columns(
  58     TENSOR3D_DECLARATION(src),
  59     IMAGE_DECLARATION(dst),
  60 #ifdef HAS_BIAS
  61     VECTOR_DECLARATION(bias),
  62 #endif /* HAS_BIAS */
  63     uint width, uint height, uint depth, uint total_filters)
  64 {
  65     Tensor3D src            = CONVERT_TO_TENSOR3D_STRUCT(src);
  66     bool     is_last_thread = (get_global_id(0) == (get_global_size(0) - 1) && get_global_id(1) == (get_global_size(1) - 1) && get_global_id(2) == (get_global_size(2) - 1));
  67
  68     __global uchar *tmp_src_ptr = src.ptr;
  69     __global uchar *tmp_dst_ptr = dst_ptr + dst_offset_first_element_in_bytes + get_global_id(0) * dst_stride_y + get_global_id(1) * width * dst_stride_y + get_global_id(
  70                                       2) * width * height * dst_stride_y;
  71 #ifdef HAS_BIAS
  72     __global uchar *tmp_bias_ptr = bias_ptr + bias_offset_first_element_in_bytes;
  73 #endif /* HAS_BIAS */
  74
  75     if(is_last_thread)
  76     {
  77         for(uint i = 0; i < total_filters; ++i)
  78         {
  79             *((__global DATA_TYPE *)tmp_dst_ptr) = *((__global DATA_TYPE *)tmp_src_ptr);
  80
  81 #ifdef HAS_BIAS
  82             *((__global DATA_TYPE *)(tmp_dst_ptr + dst_stride_y)) = *((__global DATA_TYPE *)(tmp_bias_ptr));
  83             tmp_bias_ptr += bias_stride_x;
  84 #endif /* HAS_BIAS */
  85             tmp_src_ptr += depth * src_stride_z;
  86             tmp_dst_ptr += dst_stride_x;
  87         }
  88     }
  89     else
  90     {
  91         for(uint i = 0; i < total_filters; ++i)
  92         {
  93             *((__global DATA_TYPE *)tmp_dst_ptr) = *((__global DATA_TYPE *)tmp_src_ptr);
  94             tmp_src_ptr += depth * src_stride_z;
  95             tmp_dst_ptr += dst_stride_x;
  96         }
  97     }
  98 }
  99
 100 #if defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(PAD_RIGHT) && defined(PAD_BOTTOM) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(PAD_VALUE)
 101 /** This kernel performs a reshaping of the input tensor to a tensor used to perform convolution using GEMM.
 102  *
 103  * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
 104  * @note The value to use for the paddings must be passed at compile time using -DPAD_VALUE: e.g. -DPAD_VALUE=0
 105  * @note In case biases will be added to the convolution -DHAS_BIAS has to be passed to append the final matrix with 1 in each row.
 106  *
 107  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QS8/QASYMM8/QS16/F16/F32
 108  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
 109  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 110  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
 111  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 112  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
 113  * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
 114  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
 115  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
 116  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 117  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 118  * @param[in]  dst_stride_y                      Stride of the destination tensor in Y dimension (in bytes)
 119  * @param[in]  dst_step_y                        dst_stride_y * number of elements along Y processed per workitem(in bytes)
 120  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
 121  * @param[in]  src_stride_w                      Stride of the source tensor in W dimension (in bytes).
 122  * @param[in]  dst_stride_w                      Stride of the destination tensor in W dimension (in bytes).
 123  */
 124 __kernel void im2col_generic(
 125     TENSOR3D_DECLARATION(src),
 126     IMAGE_DECLARATION(dst),
 127     uint src_stride_w,
 128     uint dst_stride_w)
 129 {
 130     const int xc    = get_global_id(0);                // x coordinate in the convolved tensor
 131     const int yc    = get_global_id(1);                // y coordinate in the convolved tensor
 132     const int ch    = get_global_id(2) % KERNEL_DEPTH; // input feature map
 133     const int batch = get_global_id(2) / KERNEL_DEPTH; // batch size
 134
 135     // Calculate input indices
 136     const int xi = xc * STRIDE_X - PAD_LEFT;
 137     const int yi = yc * STRIDE_Y - PAD_TOP;
 138
 139     // Calculate output indices
 140     const int xo = ch * KERNEL_WIDTH * KERNEL_HEIGHT;
 141     const int yo = xc + yc * CONVOLVED_WIDTH; // Index of the convolution
 142
 143     __global uchar *input_ptr      = src_ptr + src_offset_first_element_in_bytes + ch * src_stride_z + batch * src_stride_w;
 144     __global DATA_TYPE *output_ptr = ((__global DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + yo * dst_stride_y + batch * dst_stride_w)) + xo;
 145
 146     // Linearize convolution elements
 147     for(int y = yi, y_e = yi + KERNEL_HEIGHT; y < y_e; ++y)
 148     {
 149         for(int x = xi, x_e = xi + KERNEL_WIDTH; x < x_e; ++x, ++output_ptr)
 150         {
 151 #if PAD_LEFT == 0 && PAD_TOP == 0 && PAD_RIGHT == 0 && PAD_BOTTOM == 0
 152             *output_ptr = *((__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y));
 153 #else  // PAD_LEFT == 0 && PAD_TOP == 0 && PAD_RIGHT == 0 && PAD_BOTTOM == 0
 154             if(x < 0 || x >= SRC_WIDTH || y < 0 || y >= SRC_HEIGHT)
 155             {
 156                 *output_ptr = PAD_VALUE;
 157             }
 158             else
 159             {
 160                 *output_ptr = *((__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y));
 161             }
 162 #endif // PAD_LEFT == 0 && PAD_TOP == 0 && PAD_RIGHT == 0 && PAD_BOTTOM == 0
 163         }
 164     }
 165
 166 #ifdef HAS_BIAS
 167     if(ch == (KERNEL_DEPTH - 1))
 168     {
 169 #ifdef FIXED_POINT_POSITION
 170         *output_ptr = (DATA_TYPE)(1 << FIXED_POINT_POSITION);
 171 #else  // FIXED_POINT_POSITION
 172         *output_ptr       = 1.0f;
 173 #endif // FIXED_POINT_POSITION
 174     }
 175 #endif // HAS_BIAS
 176 }
 177
 178 /** This kernel performs a reshaping of the input tensor to a tensor used to perform convolution using GEMM when the kernel size is 3x3 and pad_x = pad_y = 0
 179  *
 180  * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
 181  * @note In case biases will be added to the convolution -DHAS_BIAS has to be passed to append the final matrix with 1 in each row.
 182  *
 183  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QS8/QASYMM8/QS16/F16/F32
 184  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
 185  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 186  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
 187  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 188  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
 189  * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
 190  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
 191  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
 192  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 193  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 194  * @param[in]  dst_stride_y                      Stride of the destination tensor in Y dimension (in bytes)
 195  * @param[in]  dst_step_y                        dst_stride_y * number of elements along Y processed per workitem(in bytes)
 196  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
 197  * @param[in]  src_stride_w                      Stride of the source tensor in W dimension (in bytes).
 198  * @param[in]  dst_stride_w                      Stride of the destination tensor in W dimension (in bytes).
 199  */
 200 __kernel void im2col_kernel3x3_padx0_pady0(
 201     TENSOR3D_DECLARATION(src),
 202     IMAGE_DECLARATION(dst),
 203     uint src_stride_w,
 204     uint dst_stride_w)
 205 {
 206     const int xc    = get_global_id(0);                // x coordinate in the convolved tensor
 207     const int yc    = get_global_id(1);                // y coordinate in the convolved tensor
 208     const int ch    = get_global_id(2) % KERNEL_DEPTH; // input feature map
 209     const int batch = get_global_id(2) / KERNEL_DEPTH; // batch size
 210
 211     // Calculate input indices
 212     const int xi = xc * STRIDE_X;
 213     const int yi = yc * STRIDE_Y;
 214
 215     // Calculate output indices
 216     const int xo = ch * KERNEL_WIDTH * KERNEL_HEIGHT;
 217     const int yo = xc + yc * CONVOLVED_WIDTH; // Index of the convolution
 218
 219     // Get input and output address
 220     __global uchar *input_ptr = src_ptr + src_offset_first_element_in_bytes + xi * src_stride_x + yi * src_stride_y + ch * src_stride_z + batch * src_stride_w;
 221
 222     __global DATA_TYPE *output_ptr = (__global DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + yo * dst_stride_y + batch * dst_stride_w) + xo;
 223
 224     VEC_DATA_TYPE(DATA_TYPE, 3)
 225     row0 = vload3(0, (__global DATA_TYPE *)(input_ptr + 0 * src_stride_y));
 226     VEC_DATA_TYPE(DATA_TYPE, 3)
 227     row1 = vload3(0, (__global DATA_TYPE *)(input_ptr + 1 * src_stride_y));
 228     VEC_DATA_TYPE(DATA_TYPE, 3)
 229     row2 = vload3(0, (__global DATA_TYPE *)(input_ptr + 2 * src_stride_y));
 230
 231     vstore8((VEC_DATA_TYPE(DATA_TYPE, 8))(row0.s012, row1.s012, row2.s01), 0, output_ptr);
 232     *(output_ptr + 8) = row2.s2;
 233
 234 #ifdef HAS_BIAS
 235     if(ch == (KERNEL_DEPTH - 1))
 236     {
 237 #ifdef FIXED_POINT_POSITION
 238         *(output_ptr + 9) = (DATA_TYPE)(1 << FIXED_POINT_POSITION);
 239 #else  // FIXED_POINT_POSITION
 240         *(output_ptr + 9) = 1.0f;
 241 #endif // FIXED_POINT_POSITION
 242     }
 243 #endif // HAS_BIAS
 244 }
 245 #endif //defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(PAD_RIGHT) && defined(PAD_BOTTOM) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT)
 246
 247 #if defined(WIDTH_OUTPUT)
 248 /** This kernel performs a reshaping of the output of the convolution layer.
 249  *
 250  * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
 251  *
 252  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QS8/QASYMM8/QS16/F16/F32
 253  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
 254  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 255  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
 256  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 257  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
 258  * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
 259  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
 260  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
 261  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 262  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 263  * @param[in]  dst_stride_y                      Stride of the destination tensor in Y dimension (in bytes)
 264  * @param[in]  dst_step_y                        dst_stride_y * number of elements along Y processed per workitem(in bytes)
 265  * @param[in]  dst_stride_z                      Stride of the destination tensor in Z dimension (in bytes)
 266  * @param[in]  dst_step_z                        dst_stride_z * number of elements along Z processed per workitem(in bytes)
 267  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
 268  * @param[in]  dst_stride_w                      Stride of the destination tensor in W dimension (in bytes)
 269  */
 270 __kernel void col2im(
 271     TENSOR3D_DECLARATION(src),
 272     TENSOR3D_DECLARATION(dst),
 273     uint dst_stride_w)
 274 {
 275     Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
 276     Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(dst);
 277
 278     // Compute output offset
 279     int idx = get_global_id(0) * dst.stride_z + (get_global_id(1) / WIDTH_OUTPUT) * dst_stride_y + (get_global_id(1) % WIDTH_OUTPUT) * dst_stride_x + get_global_id(2) * dst_stride_w;
 280
 281     // Store value
 282     *((__global DATA_TYPE *)(dst.ptr + idx)) = *((__global DATA_TYPE *)(src.ptr));
 283 }
 284 #endif // defined(WIDTH_OUTPUT)
 285
 286 /** This kernel reshapes the tensor's low three dimensions to single row for GEMM operation
 287  *
 288  * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float
 289  * @note In case biases will be added in late stage, -DHAS_BIAS has to be passed to append the final matrix with 1 in each row.
 290  *
 291  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QS8/QASYMM8/QS16/F16/F32
 292  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
 293  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 294  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
 295  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 296  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
 297  * @param[in]  src_step_z                        src_stride_z * number of elements along Y processed per workitem(in bytes)
 298  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
 299  * @param[out] dst_ptr                           Pointer to the destination tensor. Same as @p src_ptr
 300  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 301  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 302  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
 303  * @param[in]  width                             The width of the input tensor
 304  * @param[in]  height                            The height of the input tensor
 305  */
 306 __kernel void im2col_reduced(
 307     TENSOR3D_DECLARATION(src),
 308     VECTOR_DECLARATION(dst),
 309     uint width, uint height)
 310 {
 311     Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
 312
 313     const uint image_size = width * height;
 314
 315     __global uchar *tmp_out_ptr = dst_ptr + dst_offset_first_element_in_bytes + (get_global_id(0) + get_global_id(1) * width + get_global_id(2) * image_size) * dst_stride_x;
 316
 317     *((__global DATA_TYPE *)tmp_out_ptr) = *((__global DATA_TYPE *)src.ptr);
 318
 319 #ifdef HAS_BIAS
 320     // If it is the last thread in the 3 dimensional workgroup
 321     if(get_global_id(0) == (get_global_size(0) - 1) && get_global_id(1) == (get_global_size(1) - 1) && get_global_id(2) == (get_global_size(2) - 1))
 322     {
 323         tmp_out_ptr += dst_stride_x;
 324 #ifdef FIXED_POINT_POSITION
 325         *((__global DATA_TYPE *)tmp_out_ptr) = (DATA_TYPE)(1 << FIXED_POINT_POSITION);
 326 #else  // FIXED_POINT_POSITION
 327         *((__global DATA_TYPE *)tmp_out_ptr) = (DATA_TYPE)1;
 328 #endif // FIXED_POINT_POSITION
 329     }
 330 #endif // HAS_BIAS
 331 }
 332
 333 #if defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(PAD_RIGHT) && defined(PAD_BOTTOM) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(VECTOR_SIZE) && defined(WIDTH_MOD_VECTOR_SIZE)
 334 /** This kernel reshapes the input tensor to a tensor used to perform convolution using GEMM when
 335  * the kernel width is greater than 1 (except when the kernel size is 3x3) and pad_x == pad_y == 0.
 336  *
 337  * @note The data type must be passed at compile time using -DDATA_TYPE e.g. -DDATA_TYPE=float.
 338  * @note The vector size must be passed at compile time using -DVECTOR_SIZE e.g. -DVECTOR_SIZE=4.
 339  * @note The width modulo vector size must be passed at compile time using -DWIDTH_MOD_VECTOR_SIZE e.g. -DWIDTH_MOD_VECTOR_SIZE=3.
 340  * @note In case biases will be added to the convolution -DHAS_BIAS has to be passed to append the final matrix with 1 in each row.
 341  *
 342  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QS8/QS16/F16/F32
 343  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
 344  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 345  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
 346  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 347  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
 348  * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
 349  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
 350  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
 351  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 352  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 353  * @param[in]  dst_stride_y                      Stride of the destination tensor in Y dimension (in bytes)
 354  * @param[in]  dst_step_y                        dst_stride_y * number of elements along Y processed per workitem(in bytes)
 355  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
 356  * @param[in]  src_stride_w                      Stride of the source tensor in W dimension (in bytes).
 357  * @param[in]  dst_stride_w                      Stride of the destination tensor in W dimension (in bytes).
 358  */
 359 __kernel void im2col_generic_padx0_pady0(
 360     TENSOR3D_DECLARATION(src),
 361     IMAGE_DECLARATION(dst),
 362     uint src_stride_w,
 363     uint dst_stride_w)
 364 {
 365     const int xc    = get_global_id(0);                // x coordinate in the convolved tensor
 366     const int yc    = get_global_id(1);                // y coordinate in the convolved tensor
 367     const int ch    = get_global_id(2) % KERNEL_DEPTH; // input feature map
 368     const int batch = get_global_id(2) / KERNEL_DEPTH; // batch size
 369
 370     // Calculate input indices
 371     const int xi = xc * STRIDE_X;
 372     const int yi = yc * STRIDE_Y;
 373     // Calculate output indices
 374     const int xo                   = ch * KERNEL_WIDTH * KERNEL_HEIGHT;
 375     const int yo                   = xc + yc * CONVOLVED_WIDTH; // Index of the convolution
 376     __global uchar *input_ptr      = src_ptr + src_offset_first_element_in_bytes + ch * src_stride_z + batch * src_stride_w;
 377     __global DATA_TYPE *output_ptr = ((__global DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + yo * dst_stride_y + batch * dst_stride_w)) + xo;
 378     // Linearize convolution elements
 379     for(int y = yi, y_e = yi + KERNEL_HEIGHT; y < y_e; ++y)
 380     {
 381         int last_x = 0;
 382         for(int x = xi, x_e = xi + KERNEL_WIDTH; x + VECTOR_SIZE <= x_e; x += VECTOR_SIZE, output_ptr += VECTOR_SIZE)
 383         {
 384             VEC_DATA_TYPE(DATA_TYPE, VECTOR_SIZE)
 385             row = VLOAD(VECTOR_SIZE)(0, (__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y));
 386             VSTORE(VECTOR_SIZE)
 387             (row, 0, output_ptr);
 388             last_x = x;
 389         }
 390         // Copy the remainder of the row by doing VLOAD(WIDTH_MOD_VECTOR_SIZE) and VSTORE(WIDTH_MOD_VECTOR_SIZE).
 391         // Note that x and output_ptr have already been incremented by VECTOR_SIZE by the loop just before exit.
 392 #if WIDTH_MOD_VECTOR_SIZE == 1
 393         *output_ptr = *((__global DATA_TYPE *)(input_ptr + (last_x + VECTOR_SIZE) * src_stride_x + y * src_stride_y));
 394 #elif WIDTH_MOD_VECTOR_SIZE > 1
 395         VEC_DATA_TYPE(DATA_TYPE, WIDTH_MOD_VECTOR_SIZE)
 396         row = VLOAD(WIDTH_MOD_VECTOR_SIZE)(0, (__global DATA_TYPE *)(input_ptr + (last_x + VECTOR_SIZE) * src_stride_x + y * src_stride_y));
 397         VSTORE(WIDTH_MOD_VECTOR_SIZE)
 398         (row, 0, output_ptr);
 399 #endif /* WIDTH_MOD_VECTOR_SIZE */
 400         output_ptr += WIDTH_MOD_VECTOR_SIZE;
 401     } /* End of loop over KERNEL_HEIGHT */
 402
 403 #ifdef HAS_BIAS
 404     if(ch == (KERNEL_DEPTH - 1))
 405     {
 406 #ifdef FIXED_POINT_POSITION
 407         *output_ptr = (DATA_TYPE)(1 << FIXED_POINT_POSITION);
 408 #else  // FIXED_POINT_POSITION
 409         *output_ptr       = 1.0f;
 410 #endif // FIXED_POINT_POSITION
 411     }
 412 #endif // HAS_BIAS
 413 }
 414 #endif //defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(PAD_RIGHT) && defined(PAD_BOTTOM) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(VECTOR_SIZE) && defined(WIDTH_MOD_VECTOR_SIZE)