src/core/CL/cl_kernels/roi_pooling_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 DATA_SIZE == 32
  27 #define VEC_SIZE 4
  28 #define VEC_MAX vec4_max
  29 #elif DATA_SIZE == 16
  30 #define VEC_SIZE 8
  31 #define VEC_MAX vec8_max
  32 #else /* DATA_SIZE not equals 32 or 16 */
  33 #error "Unsupported data size"
  34 #endif /* DATA_SIZE == 32 */
  35
  36 inline DATA_TYPE vec4_max(VEC_DATA_TYPE(DATA_TYPE, 4) vec)
  37 {
  38     VEC_DATA_TYPE(DATA_TYPE, 2)
  39     temp = fmax(vec.lo, vec.hi);
  40     return fmax(temp.x, temp.y);
  41 }
  42
  43 inline DATA_TYPE vec8_max(VEC_DATA_TYPE(DATA_TYPE, 8) vec)
  44 {
  45     VEC_DATA_TYPE(DATA_TYPE, 4)
  46     temp = fmax(vec.lo, vec.hi);
  47     return vec4_max(temp);
  48 }
  49
  50 /** Performs a roi pooling on a single output pixel.
  51  *
  52  * @param[in] input          Pointer to input Tensor3D struct.
  53  * @param[in] region_start_x Start x index projected onto the input tensor.
  54  * @param[in] region_end_x   End x index projected onto the input tensor.
  55  * @param[in] region_start_y Start y index projected onto the input tensor.
  56  * @param[in] region_end_y   End y index projected onto the input tensor.
  57  * @param[in] pz             z index of the input tensor.
  58  *
  59  * @return A max pooled value from the region specified in the input tensor.
  60  */
  61 inline DATA_TYPE roi_pool_1x1(const Tensor3D *input, int region_start_x, int region_end_x, int region_start_y, int region_end_y, int pz)
  62 {
  63     // Iterate through the pooling region
  64     if((region_end_x <= region_start_x) || (region_end_y <= region_start_y))
  65     {
  66         return (DATA_TYPE)0;
  67     }
  68     else
  69     {
  70         int num_iter = (int)((region_end_x - region_start_x) / VEC_SIZE);
  71         VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
  72         curr_max = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))(-FLT_MAX);
  73         for(int j = region_start_y; j < region_end_y; ++j)
  74         {
  75             int i = region_start_x;
  76             for(; i < region_start_x + num_iter * VEC_SIZE; i += VEC_SIZE)
  77             {
  78                 VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
  79                 val      = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)tensor3D_offset(input, i, j, pz));
  80                 curr_max = fmax(val, curr_max);
  81             }
  82             for(; i < region_end_x; ++i)
  83             {
  84                 DATA_TYPE val = *(__global DATA_TYPE *)tensor3D_offset(input, i, j, pz);
  85                 curr_max      = fmax(curr_max, val);
  86             }
  87         }
  88         return (DATA_TYPE)VEC_MAX(curr_max);
  89     }
  90 }
  91
  92 /** Performs a roi pooling function.
  93  *
  94  * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16, F32;
  95  * @note Datasize must be passed using -DDATA_SIZE e.g. -DDATA_SIZE=32;
  96  * @note Input dimensions must be passed using -DMAX_DIM_X, -DMAX_DIM_Y and -DMAX_DIM_Z;
  97  * @note Pooled region dimensions must be passed using -DPOOLED_DIM_X and -DPOOLED_DIM_Y;
  98  * @note Spatial scale must be passed using -DSPATIAL_SCALE;
  99  *
 100  * @param[in]  input_ptr                            Pointer to the source image. Supported data types: F16, F32
 101  * @param[in]  input_stride_x                       Stride of the source image in X dimension (in bytes)
 102  * @param[in]  input_step_x                         input_stride_x * number of elements along X processed per workitem(in bytes)
 103  * @param[in]  input_stride_y                       Stride of the source image in Y dimension (in bytes)
 104  * @param[in]  input_step_y                         input_stride_y * number of elements along Y processed per workitem(in bytes)
 105  * @param[in]  input_stride_z                       Stride of the source tensor in Z dimension (in bytes)
 106  * @param[in]  input_step_z                         input_stride_z * number of elements along Z processed per workitem(in bytes)
 107  * @param[in]  input_offset_first_element_in_bytes  The offset of the first element in the pooled region of the source image as specifed by ROI
 108  * @param[in]  rois_ptr                             Pointer to the rois array. Layout: {x, y, width, height, batch_indx}
 109  * @param[in]  rois_stride_x                        Stride of the rois array in X dimension (in bytes)
 110  * @param[in]  rois_step_x                          rois_stride_x * number of elements along X processed per workitem(in bytes)
 111  * @param[in]  rois_offset_first_element_in_bytes   The offset of the first element in the rois array
 112  * @param[out] output_ptr                           Pointer to the destination image. Supported data types: F16, F32
 113  * @param[in]  output_stride_x                      Stride of the destination image in X dimension (in bytes)
 114  * @param[in]  output_step_x                        output_stride_x * number of elements along X processed per workitem(in bytes)
 115  * @param[in]  output_stride_y                      Stride of the destination image in Y dimension (in bytes)
 116  * @param[in]  output_step_y                        output_stride_y * number of elements along Y processed per workitem(in bytes)
 117  * @param[in]  output_stride_z                      Stride of the destination tensor in Z dimension (in bytes)
 118  * @param[in]  output_step_z                        output_stride_z * number of elements along Z processed per workitem(in bytes)
 119  * @param[in]  output_offset_first_element_in_bytes The offset of the first element in the destination image
 120  * @param[in]  input_stride_w                       Stride of the source image in W dimension (in bytes)
 121  * @param[in]  output_stride_w                      Stride of the destination image in W dimension (in bytes)
 122  */
 123 __kernel void roi_pooling_layer(
 124     TENSOR3D_DECLARATION(input),
 125     VECTOR_DECLARATION(rois),
 126     TENSOR3D_DECLARATION(output),
 127     unsigned int input_stride_w, unsigned int output_stride_w)
 128 {
 129     // Get pixels pointer
 130     Tensor3D input  = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
 131     Vector   rois   = CONVERT_TO_VECTOR_STRUCT_NO_STEP(rois);
 132     Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
 133
 134     const int px = get_global_id(0);
 135     const int py = get_global_id(1);
 136     const int pw = get_global_id(2);
 137
 138     // Load roi parameters
 139     // roi is laid out as follows:
 140     // { x, y, width, height, batch_index }
 141     const ushort4 roi      = vload4(0, (__global ushort *)vector_offset(&rois, pw));
 142     const ushort roi_batch = *((__global ushort *)vector_offset(&rois, pw) + 4);
 143     const int2 roi_anchor  = convert_int2_sat(round(convert_float2(roi.s01) * (float)SPATIAL_SCALE));
 144     const int2 roi_dims    = convert_int2_sat(fmax(round(convert_float2(roi.s23) * (float)SPATIAL_SCALE), 1.f));
 145
 146     // Calculate pooled region start and end
 147     const float2 spatial_indx     = (float2)(px, py);
 148     const float2 pooled_dims      = (float2)(POOLED_DIM_X, POOLED_DIM_Y);
 149     const int2   max_spatial_dims = (int2)(MAX_DIM_X, MAX_DIM_Y);
 150     int2         region_start     = convert_int2_sat(floor(spatial_indx / pooled_dims * convert_float2(roi_dims))) + roi_anchor;
 151     int2         region_end       = convert_int2_sat(floor((spatial_indx + 1) / pooled_dims * convert_float2(roi_dims))) + roi_anchor;
 152
 153     region_start = clamp(region_start, 0, max_spatial_dims);
 154     region_end   = clamp(region_end, 0, max_spatial_dims);
 155
 156     // Move input and output pointer across the fourth dimension
 157     input.ptr += roi_batch * input_stride_w;
 158     output.ptr += pw * output_stride_w;
 159
 160     for(int pz = 0; pz < MAX_DIM_Z; ++pz)
 161     {
 162         *(__global DATA_TYPE *)tensor3D_offset(&output, px, py, pz) = (__global DATA_TYPE)roi_pool_1x1(&input,
 163                                                                                                        region_start.x,
 164                                                                                                        region_end.x,
 165                                                                                                        region_start.y,
 166                                                                                                        region_end.y, pz);
 167     }
 168 }