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_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT)
 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 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.
 105  *
 106  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QS8/QS16/F16/F32
 107  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
 108  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 109  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
 110  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 111  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
 112  * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
 113  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
 114  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
 115  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 116  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 117  * @param[in]  dst_stride_y                      Stride of the destination tensor in Y dimension (in bytes)
 118  * @param[in]  dst_step_y                        dst_stride_y * number of elements along Y processed per workitem(in bytes)
 119  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
 120  * @param[in]  filter_depth                      The depth of the used filter
 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 filter_depth,
 128     uint src_stride_w,
 129     uint dst_stride_w)
 130 {
 131     const int xc    = get_global_id(0);                // x coordinate in the convolved tensor
 132     const int yc    = get_global_id(1);                // y coordinate in the convolved tensor
 133     const int ch    = get_global_id(2) % filter_depth; // input feature map
 134     const int batch = get_global_id(2) / filter_depth; // the batch
 135
 136     // Calculate input indeces
 137     const int xi = xc * STRIDE_X - PAD_X;
 138     const int yi = yc * STRIDE_Y - PAD_Y;
 139
 140     // Calculate output indeces
 141     const int xo = ch * KERNEL_WIDTH * KERNEL_HEIGHT;
 142     const int yo = xc + yc * CONVOLVED_WIDTH; // Index of the convolution
 143
 144     __global uchar *input_ptr      = src_ptr + src_offset_first_element_in_bytes + ch * src_stride_z + batch * src_stride_w;
 145     __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;
 146
 147     // Linearize convolution elements
 148     for(int y = yi, y_e = yi + KERNEL_HEIGHT; y < y_e; ++y)
 149     {
 150         for(int x = xi, x_e = xi + KERNEL_WIDTH; x < x_e; ++x, ++output_ptr)
 151         {
 152 #if PAD_X == 0 && PAD_Y == 0
 153             *output_ptr = *((__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y));
 154 #else  // PAD_X == 0 && PAD_Y == 0
 155             if(x < 0 || x >= SRC_WIDTH || y < 0 || y >= SRC_HEIGHT)
 156             {
 157                 *output_ptr = 0;
 158             }
 159             else
 160             {
 161                 *output_ptr = *((__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y));
 162             }
 163 #endif // PAD_X == 0 && PAD_Y == 0
 164         }
 165     }
 166
 167 #ifdef HAS_BIAS
 168     if(ch == (KERNEL_DEPTH - 1))
 169     {
 170 #ifdef FIXED_POINT_POSITION
 171         *output_ptr = (DATA_TYPE)(1 << FIXED_POINT_POSITION);
 172 #else  // FIXED_POINT_POSITION
 173         *output_ptr       = 1.0f;
 174 #endif // FIXED_POINT_POSITION
 175     }
 176 #endif // HAS_BIAS
 177 }
 178
 179 /** 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
 180  *
 181  * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
 182  * @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.
 183  *
 184  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QS8/QS16/F16/F32
 185  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
 186  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 187  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
 188  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 189  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
 190  * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
 191  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
 192  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
 193  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 194  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 195  * @param[in]  dst_stride_y                      Stride of the destination tensor in Y dimension (in bytes)
 196  * @param[in]  dst_step_y                        dst_stride_y * number of elements along Y processed per workitem(in bytes)
 197  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
 198  * @param[in]  filter_depth                      The depth of the used filter
 199  * @param[in]  src_stride_w                      Stride of the source tensor in W dimension (in bytes).
 200  * @param[in]  dst_stride_w                      Stride of the destination tensor in W dimension (in bytes).
 201  */
 202 __kernel void im2col_kernel3x3_padx0_pady0(
 203     TENSOR3D_DECLARATION(src),
 204     IMAGE_DECLARATION(dst),
 205     uint filter_depth,
 206     uint src_stride_w,
 207     uint dst_stride_w)
 208 {
 209     const int xc    = get_global_id(0);                // x coordinate in the convolved tensor
 210     const int yc    = get_global_id(1);                // y coordinate in the convolved tensor
 211     const int ch    = get_global_id(2) % filter_depth; // input feature map
 212     const int batch = get_global_id(2) / filter_depth; // the batch
 213
 214     // Calculate input indeces
 215     const int xi = xc * STRIDE_X;
 216     const int yi = yc * STRIDE_Y;
 217
 218     // Calculate output indeces
 219     const int xo = ch * KERNEL_WIDTH * KERNEL_HEIGHT;
 220     const int yo = xc + yc * CONVOLVED_WIDTH; // Index of the convolution
 221
 222     // Get input and output address
 223     __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;
 224
 225     __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;
 226
 227     VEC_DATA_TYPE(DATA_TYPE, 3)
 228     row0 = vload3(0, (__global DATA_TYPE *)(input_ptr + 0 * src_stride_y));
 229     VEC_DATA_TYPE(DATA_TYPE, 3)
 230     row1 = vload3(0, (__global DATA_TYPE *)(input_ptr + 1 * src_stride_y));
 231     VEC_DATA_TYPE(DATA_TYPE, 3)
 232     row2 = vload3(0, (__global DATA_TYPE *)(input_ptr + 2 * src_stride_y));
 233
 234     vstore8((VEC_DATA_TYPE(DATA_TYPE, 8))(row0.s012, row1.s012, row2.s01), 0, output_ptr);
 235     *(output_ptr + 8) = row2.s2;
 236
 237 #ifdef HAS_BIAS
 238     if(ch == (KERNEL_DEPTH - 1))
 239     {
 240 #ifdef FIXED_POINT_POSITION
 241         *(output_ptr + 9) = (DATA_TYPE)(1 << FIXED_POINT_POSITION);
 242 #else  // FIXED_POINT_POSITION
 243         *(output_ptr + 9) = 1.0f;
 244 #endif // FIXED_POINT_POSITION
 245     }
 246 #endif // HAS_BIAS
 247 }
 248 #endif //defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT)
 249
 250 #if defined(WIDTH_OUTPUT)
 251 /** This kernel performs a reshaping of the output of the convolution layer.
 252  *
 253  * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
 254  *
 255  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QS8/QS16/F16/F32
 256  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
 257  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 258  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
 259  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 260  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
 261  * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
 262  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
 263  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
 264  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 265  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 266  * @param[in]  dst_stride_y                      Stride of the destination tensor in Y dimension (in bytes)
 267  * @param[in]  dst_step_y                        dst_stride_y * number of elements along Y processed per workitem(in bytes)
 268  * @param[in]  dst_stride_z                      Stride of the destination tensor in Z dimension (in bytes)
 269  * @param[in]  dst_step_z                        dst_stride_z * number of elements along Z processed per workitem(in bytes)
 270  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
 271  * @param[in]  dst_stride_w                      Stride of the destination tensor in W dimension (in bytes)
 272  */
 273 __kernel void col2im(
 274     TENSOR3D_DECLARATION(src),
 275     TENSOR3D_DECLARATION(dst),
 276     uint dst_stride_w)
 277 {
 278     Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
 279     Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(dst);
 280
 281     // Compute output offset
 282     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;
 283
 284     // Store value
 285     *((__global DATA_TYPE *)(dst.ptr + idx)) = *((__global DATA_TYPE *)(src.ptr));
 286 }
 287 #endif // defined(WIDTH_OUTPUT)
 288
 289 /** This kernel reshapes the tensor's low three dimensions to single row for GEMM operation
 290  *
 291  * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float
 292  * @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.
 293  *
 294  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QS8/F16/F32
 295  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
 296  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 297  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
 298  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 299  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
 300  * @param[in]  src_step_z                        src_stride_z * number of elements along Y processed per workitem(in bytes)
 301  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
 302  * @param[out] dst_ptr                           Pointer to the destination tensor. Same as @p src_ptr
 303  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 304  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 305  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
 306  * @param[in]  width                             The width of the input tensor
 307  * @param[in]  height                            The height of the input tensor
 308  */
 309 __kernel void im2col_reduced(
 310     TENSOR3D_DECLARATION(src),
 311     VECTOR_DECLARATION(dst),
 312     uint width, uint height)
 313 {
 314     Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
 315
 316     const uint image_size = width * height;
 317
 318     __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;
 319
 320     *((__global DATA_TYPE *)tmp_out_ptr) = *((__global DATA_TYPE *)src.ptr);
 321
 322 #ifdef HAS_BIAS
 323     // If it is the last thread in the 3 dimensional workgroup
 324     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))
 325     {
 326         tmp_out_ptr += dst_stride_x;
 327 #ifdef FIXED_POINT_POSITION
 328         *((__global DATA_TYPE *)tmp_out_ptr) = (DATA_TYPE)(1 << FIXED_POINT_POSITION);
 329 #else  // FIXED_POINT_POSITION
 330         *((__global DATA_TYPE *)tmp_out_ptr) = (DATA_TYPE)1;
 331 #endif // FIXED_POINT_POSITION
 332     }
 333 #endif // HAS_BIAS
 334 }