src/core/CL/cl_kernels/direct_convolution1x1.cl

   1 /*
   2  * Copyright (c) 2016, 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
  29 #define ADD_OP(a, b) ADD_SAT_OP_EXPAND((a), (b), DATA_TYPE_PROMOTED, 8)
  30 #define MUL_OP(a, b) MUL_SAT_OP_EXPAND(CONVERT((a), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 8)), CONVERT((b), VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 8)), DATA_TYPE_PROMOTED, 8, FIXED_POINT_POSITION)
  31
  32 // There is no need to have a larger intermediate type for qs32 because all the arguments are already promoted
  33 MULQ_SAT_IMPL(qs32x8, qs32x8)
  34
  35 #else /* FIXED_POINT_POSITION */
  36 #undef CONVERT_SAT
  37
  38 #define ADD_OP(a, b) ((a) + (b))
  39 #define MUL_OP(a, b) ((a) * (b))
  40 #define CONVERT_SAT(a, b) ((a))
  41
  42 #endif /* FIXED_POINT_POSITION */
  43
  44 #if defined(DATA_TYPE) && defined(DATA_SIZE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH)
  45
  46 #if STRIDE_X == 3
  47 #define INPUT_PIXEL_STR(data_size) extract_input_stride3_##data_size
  48 #define INPUT_PIXEL(data_size) INPUT_PIXEL_STR(data_size)
  49 #elif STRIDE_X == 2
  50 #define INPUT_PIXEL(data_size) extract_input_stride2
  51 #elif STRIDE_X == 1
  52 #define INPUT_PIXEL(data_size) extract_input_stride1
  53 #else /* STRIDE_X not equals 1, 2 or 3 */
  54 #error "Only support strides 1, 2 and 3"
  55 #endif /* STRIDE_X == 3 */
  56
  57 /** Extracts a 1D horizontal vector from the input tensor with stride as 1.
  58  *
  59  * @param[in] input_pixel Pointer to the first pixel.
  60  *
  61  * @return extracted input pixels.
  62  */
  63 inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride1(__global const DATA_TYPE *input_pixel)
  64 {
  65     return vload8(0, input_pixel);
  66 }
  67
  68 /** Extracts a 1D horizontal vector from the input tensor with stride as 2.
  69  *
  70  * @param[in] input_pixel Pointer to the first pixel.
  71  *
  72  * @return extracted input pixels.
  73  */
  74 inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride2(__global const DATA_TYPE *input_pixel)
  75 {
  76     VEC_DATA_TYPE(DATA_TYPE, 16)
  77     temp = vload16(0, input_pixel);
  78     return temp.s02468ace;
  79 }
  80
  81 /** Extracts a 1D horizontal vector from the input tensor with stride as 3 and 32-bit data size.
  82  *
  83  * @param[in] input_pixel Pointer to the first pixel.
  84  *
  85  * @return extracted input pixels.
  86  */
  87 inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3_32(__global const DATA_TYPE *input_pixel)
  88 {
  89     VEC_DATA_TYPE(DATA_TYPE, 4)
  90     temp1 = vload4(0, input_pixel);
  91     VEC_DATA_TYPE(DATA_TYPE, 4)
  92     temp2 = vload4(0, input_pixel + 6);
  93     VEC_DATA_TYPE(DATA_TYPE, 4)
  94     temp3 = vload4(0, input_pixel + 12);
  95     VEC_DATA_TYPE(DATA_TYPE, 4)
  96     temp4 = vload4(0, input_pixel + 18);
  97     return (VEC_DATA_TYPE(DATA_TYPE, 8))(temp1.s03, temp2.s03, temp3.s03, temp4.s03);
  98 }
  99
 100 /** Extracts a 1D horizontal vector from the input tensor with stride as 3 and 16-bit data size.
 101  *
 102  * @param[in] input_pixel Pointer to the first pixel.
 103  *
 104  * @return extracted input pixels.
 105  */
 106 inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3_16(__global const DATA_TYPE *input_pixel)
 107 {
 108     VEC_DATA_TYPE(DATA_TYPE, 8)
 109     temp1 = vload8(0, input_pixel);
 110     VEC_DATA_TYPE(DATA_TYPE, 8)
 111     temp2 = vload8(0, input_pixel + 8);
 112     VEC_DATA_TYPE(DATA_TYPE, 8)
 113     temp3 = vload8(0, input_pixel + 16);
 114     return (VEC_DATA_TYPE(DATA_TYPE, 8))(temp1.s036, temp2.s147, temp3.s25);
 115 }
 116
 117 /** Extracts a 1D horizontal vector from the input tensor with stride as 3 and 8-bit data size.
 118  *
 119  * @param[in] input_pixel Pointer to the first pixel.
 120  *
 121  * @return extracted input pixels.
 122  */
 123 inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3_8(__global const DATA_TYPE *input_pixel)
 124 {
 125     VEC_DATA_TYPE(DATA_TYPE, 16)
 126     temp1 = vload16(0, input_pixel);
 127     VEC_DATA_TYPE(DATA_TYPE, 16)
 128     temp2 = vload16(0, input_pixel + 12);
 129     return (VEC_DATA_TYPE(DATA_TYPE, 8))(temp1.s0369, temp2.s0369);
 130 }
 131
 132 /** This kernel performs a direct convolution to convolve the low three dimensions.
 133  *
 134  * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
 135  * @note The data size must be passed at compile time using -DDATA_SIZE e.g. -DDATA_SIZE=32
 136  * @note The convolution stride x must be passed at compile time using -DSTRIDE_X e.g. -DSTRIDE_X=1
 137  * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH
 138  * @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.
 139  *
 140  * @param[in]  src_ptr                               Pointer to the source tensor. Supported data types: F16/F32
 141  * @param[in]  src_stride_x                          Stride of the source tensor in X dimension (in bytes)
 142  * @param[in]  src_step_x                            src_stride_x * number of elements along X processed per workitem(in bytes)
 143  * @param[in]  src_stride_y                          Stride of the source tensor in Y dimension (in bytes)
 144  * @param[in]  src_step_y                            src_stride_y * number of elements along Y processed per workitem(in bytes)
 145  * @param[in]  src_stride_z                          Stride of the source tensor in Z dimension (in bytes)
 146  * @param[in]  src_step_z                            src_stride_z * number of elements along Z processed per workitem(in bytes)
 147  * @param[in]  src_offset_first_element_in_bytes     The offset of the first element in the source tensor
 148  * @param[out] dst_ptr                               Pointer to the destination tensor. Supported data types: same as @p src_ptr
 149  * @param[in]  dst_stride_x                          Stride of the destination tensor in X dimension (in bytes)
 150  * @param[in]  dst_step_x                            dst_stride_x * number of elements along X processed per workitem(in bytes)
 151  * @param[in]  dst_stride_y                          Stride of the destination tensor in Y dimension (in bytes)
 152  * @param[in]  dst_step_y                            dst_stride_y * number of elements along Z processed per workitem(in bytes)
 153  * @param[in]  dst_stride_z                          Stride of the destination tensor in Z dimension (in bytes)
 154  * @param[in]  dst_step_z                            dst_stride_z * number of elements along Z processed per workitem(in bytes)
 155  * @param[in]  dst_offset_first_element_in_bytes     The offset of the first element in the destination tensor
 156  * @param[in]  weights_ptr                           Pointer to the weights tensor. Supported data types: same as @p src_ptr
 157  * @param[in]  weights_stride_x                      Stride of the weights tensor in X dimension (in bytes)
 158  * @param[in]  weights_step_x                        weights_stride_x * number of elements along X processed per workitem(in bytes)
 159  * @param[in]  weights_stride_y                      Stride of the weights tensor in Y dimension (in bytes)
 160  * @param[in]  weights_step_y                        weights_stride_y * number of elements along y processed per workitem(in bytes)
 161  * @param[in]  weights_stride_z                      Stride of the weights tensor in Z dimension (in bytes)
 162  * @param[in]  weights_step_z                        weights_stride_z * number of elements along Z processed per workitem(in bytes)
 163  * @param[in]  weights_offset_first_element_in_bytes The offset of the first element in the weights tensor
 164  * @param[in]  biases_ptr                            Pointer to the biases tensor. Same as @p src_ptr
 165  * @param[in]  biases_stride_x                       Stride of the biases tensor in X dimension (in bytes)
 166  * @param[in]  biases_step_x                         biases_stride_x * number of elements along X processed per workitem(in bytes)
 167  * @param[in]  biases_offset_first_element_in_bytes  The offset of the first element in the biases tensor
 168  * @param[in]  weights_stride_w                      Stride of the weights tensor in the 4th dimension
 169  */
 170 __kernel void direct_convolution1x1(
 171     TENSOR3D_DECLARATION(src),
 172     TENSOR3D_DECLARATION(dst),
 173     TENSOR3D_DECLARATION(weights),
 174 #ifdef HAS_BIAS
 175     VECTOR_DECLARATION(biases),
 176 #endif /* defined(HAS_BIAS) */
 177     unsigned int weights_stride_w)
 178 {
 179     Image    src     = CONVERT_TO_IMAGE_STRUCT(src);
 180     Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights);
 181     Tensor3D dst     = CONVERT_TO_TENSOR3D_STRUCT(dst);
 182
 183 #ifdef HAS_BIAS
 184     Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases);
 185 #endif /* defined(HAS_BIAS) */
 186
 187     VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 8)
 188     pixels = 0;
 189
 190     const uint z_index = get_global_id(2);
 191
 192     weights.ptr += z_index * weights_stride_w;
 193
 194     for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d)
 195     {
 196         DATA_TYPE weight = *(__global DATA_TYPE *)weights.ptr;
 197         VEC_DATA_TYPE(DATA_TYPE, 8)
 198         input_pixel = INPUT_PIXEL(DATA_SIZE)((__global DATA_TYPE *)src.ptr);
 199         pixels      = ADD_OP(pixels, MUL_OP((VEC_DATA_TYPE(DATA_TYPE, 8))weight, input_pixel));
 200         src.ptr += src_stride_z;
 201         weights.ptr += weights_stride_z;
 202     }
 203
 204 #ifdef HAS_BIAS
 205     pixels = ADD_OP(pixels, (VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 8)) * ((__global DATA_TYPE *)(vector_offset(&biases, z_index))));
 206 #endif /* defined(HAS_BIAS) */
 207
 208     vstore8(CONVERT_SAT(pixels, VEC_DATA_TYPE(DATA_TYPE, 8)), 0, (__global DATA_TYPE *)dst.ptr);
 209 }
 210 #endif // defined(DATA_TYPE) && defined(DATA_SIZE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH)
 211
 212 #if defined(WEIGHTS_DEPTH)
 213
 214 #define CONVOLUTION1x1_BIFROST(acc, src, weight_value) \
 215     ({                                                 \
 216         acc.s0 = mad(src.s0, weight_value, acc.s0);    \
 217         acc.s1 = mad(src.s1, weight_value, acc.s1);    \
 218         acc.s2 = mad(src.s2, weight_value, acc.s2);    \
 219         acc.s3 = mad(src.s3, weight_value, acc.s3);    \
 220     })
 221
 222 /** An optimized direct convolution 1x1 OpenCL kernel for Bifrost architectures when the data type is F32
 223  *
 224  * @note This OpenCL kernel works only with stride_x and stride_y equal to 1
 225  * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH
 226  * @note In case biases, -DHAS_BIAS must to be passed at compile
 227  *
 228  * @param[in]  src_ptr                               Pointer to the source tensor. Supported data types: F32
 229  * @param[in]  src_stride_x                          Stride of the source tensor in X dimension (in bytes)
 230  * @param[in]  src_step_x                            src_stride_x * number of elements along X processed per workitem(in bytes)
 231  * @param[in]  src_stride_y                          Stride of the source tensor in Y dimension (in bytes)
 232  * @param[in]  src_step_y                            src_stride_y * number of elements along Y processed per workitem(in bytes)
 233  * @param[in]  src_stride_z                          Stride of the source tensor in Z dimension (in bytes)
 234  * @param[in]  src_step_z                            src_stride_z * number of elements along Z processed per workitem(in bytes)
 235  * @param[in]  src_offset_first_element_in_bytes     The offset of the first element in the source tensor
 236  * @param[out] dst_ptr                               Pointer to the destination tensor. Supported data types: same as @p src_ptr
 237  * @param[in]  dst_stride_x                          Stride of the destination tensor in X dimension (in bytes)
 238  * @param[in]  dst_step_x                            dst_stride_x * number of elements along X processed per workitem(in bytes)
 239  * @param[in]  dst_stride_y                          Stride of the destination tensor in Y dimension (in bytes)
 240  * @param[in]  dst_step_y                            dst_stride_y * number of elements along Z processed per workitem(in bytes)
 241  * @param[in]  dst_stride_z                          Stride of the destination tensor in Z dimension (in bytes)
 242  * @param[in]  dst_step_z                            dst_stride_z * number of elements along Z processed per workitem(in bytes)
 243  * @param[in]  dst_offset_first_element_in_bytes     The offset of the first element in the destination tensor
 244  * @param[in]  weights_ptr                           Pointer to the weights tensor. Supported data types: same as @p src_ptr
 245  * @param[in]  weights_stride_x                      Stride of the weights tensor in X dimension (in bytes)
 246  * @param[in]  weights_step_x                        weights_stride_x * number of elements along X processed per workitem(in bytes)
 247  * @param[in]  weights_stride_y                      Stride of the weights tensor in Y dimension (in bytes)
 248  * @param[in]  weights_step_y                        weights_stride_y * number of elements along y processed per workitem(in bytes)
 249  * @param[in]  weights_stride_z                      Stride of the weights tensor in Z dimension (in bytes)
 250  * @param[in]  weights_step_z                        weights_stride_z * number of elements along Z processed per workitem(in bytes)
 251  * @param[in]  weights_offset_first_element_in_bytes The offset of the first element in the weights tensor
 252  * @param[in]  biases_ptr                            Pointer to the biases tensor. Same as @p src_ptr
 253  * @param[in]  biases_stride_x                       Stride of the biases tensor in X dimension (in bytes)
 254  * @param[in]  biases_step_x                         biases_stride_x * number of elements along X processed per workitem(in bytes)
 255  * @param[in]  biases_offset_first_element_in_bytes  The offset of the first element in the biases tensor
 256  * @param[in]  weights_stride_w                      Stride of the weights tensor in the 4th dimension
 257  */
 258 __kernel void direct_convolution1x1_f32_bifrost(
 259     TENSOR3D_DECLARATION(src),
 260     TENSOR3D_DECLARATION(dst),
 261     TENSOR3D_DECLARATION(weights),
 262 #ifdef HAS_BIAS
 263     VECTOR_DECLARATION(biases),
 264 #endif /* defined(HAS_BIAS) */
 265     unsigned int weights_stride_w)
 266 {
 267     // Get the kernel index
 268     const int kernel_index = get_global_id(2);
 269
 270     Image    src = CONVERT_TO_IMAGE_STRUCT(src);
 271     Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);
 272
 273     float4 acc0 = 0.0f;
 274     float4 acc1 = 0.0f;
 275     float4 acc2 = 0.0f;
 276     float4 acc3 = 0.0f;
 277
 278     __global uchar *weights_addr = (__global uchar *)(weights_ptr + weights_offset_first_element_in_bytes + kernel_index * weights_stride_w);
 279     __global uchar *src_addr     = (__global uchar *)offset(&src, 0, 0);
 280
 281     for(ushort d = 0; d < (ushort)WEIGHTS_DEPTH; ++d)
 282     {
 283         // Load the weights
 284         float weight = *((__global float *)weights_addr);
 285
 286         // Load values from row0 of input tensor
 287         float4 src0 = vload4(0, (__global float *)(src_addr + 0 * src_stride_y));
 288         float4 src1 = vload4(0, (__global float *)(src_addr + 1 * src_stride_y));
 289         float4 src2 = vload4(0, (__global float *)(src_addr + 2 * src_stride_y));
 290         float4 src3 = vload4(0, (__global float *)(src_addr + 3 * src_stride_y));
 291
 292         CONVOLUTION1x1_BIFROST(acc0, src0, weight);
 293         CONVOLUTION1x1_BIFROST(acc1, src1, weight);
 294         CONVOLUTION1x1_BIFROST(acc2, src2, weight);
 295         CONVOLUTION1x1_BIFROST(acc3, src3, weight);
 296
 297         src_addr += src_stride_z;
 298         weights_addr += weights_stride_z;
 299     }
 300
 301 #ifdef HAS_BIAS
 302     Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases);
 303
 304     float bias = (float) * ((__global float *)(vector_offset(&biases, kernel_index)));
 305
 306     acc0.s0 += bias;
 307     acc0.s1 += bias;
 308     acc0.s2 += bias;
 309     acc0.s3 += bias;
 310     acc1.s0 += bias;
 311     acc1.s1 += bias;
 312     acc1.s2 += bias;
 313     acc1.s3 += bias;
 314     acc2.s0 += bias;
 315     acc2.s1 += bias;
 316     acc2.s2 += bias;
 317     acc2.s3 += bias;
 318     acc3.s0 += bias;
 319     acc3.s1 += bias;
 320     acc3.s2 += bias;
 321     acc3.s3 += bias;
 322 #endif /* defined(HAS_BIAS) */
 323
 324     vstore4(acc0, 0, (__global float *)(dst.ptr + 0 * dst_stride_y));
 325     vstore4(acc1, 0, (__global float *)(dst.ptr + 1 * dst_stride_y));
 326     vstore4(acc2, 0, (__global float *)(dst.ptr + 2 * dst_stride_y));
 327     vstore4(acc3, 0, (__global float *)(dst.ptr + 3 * dst_stride_y));
 328 }
 329 #endif // defined(WEIGHTS_DEPTH)