src/core/CL/cl_kernels/minmaxloc.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 #include "types.h"
  26
  27 #ifndef DATA_TYPE_MIN
  28 #define DATA_TYPE_MIN 0x0
  29 #endif /* DATA_TYPE_MIN */
  30
  31 #ifndef DATA_TYPE_MAX
  32 #define DATA_TYPE_MAX 0xFF
  33 #endif /* DATA_TYPE_MAX */
  34
  35 inline int FloatFlip(float val)
  36 {
  37     union
  38     {
  39         int   int_val;
  40         float flt_val;
  41     } u_val;
  42     u_val.flt_val = val;
  43     return (u_val.int_val >= 0) ? u_val.int_val : u_val.int_val ^ 0x7FFFFFFF;
  44 }
  45
  46 __constant VEC_DATA_TYPE(DATA_TYPE, 16) type_min = (VEC_DATA_TYPE(DATA_TYPE, 16))(DATA_TYPE_MIN);
  47 __constant VEC_DATA_TYPE(DATA_TYPE, 16) type_max = (VEC_DATA_TYPE(DATA_TYPE, 16))(DATA_TYPE_MAX);
  48 __constant int16 idx16 = (int16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
  49
  50 /** This function identifies the min and maximum value of an input image.
  51  *
  52  * @note Input image data type must be passed as a preprocessor argument using -DDATA_TYPE.
  53  * Moreover, the minimum and maximum value of the given data type must be provided using -DDATA_TYPE_MIN and -DDATA_TYPE_MAX respectively.
  54  * @note In case image width is not a multiple of 16 then -DNON_MULTIPLE_OF_16 must be passed.
  55  *
  56  * @param[in]  src_ptr                           Pointer to the source image. Supported data types: U8
  57  * @param[in]  src_stride_x                      Stride of the source image in X dimension (in bytes)
  58  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
  59  * @param[in]  src_stride_y                      Stride of the source image in Y dimension (in bytes)
  60  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
  61  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source image
  62  * @param[out] min_max                           Pointer to buffer with minimum value in position 0 and maximum value in position 1
  63  * @param[in]  width                             Input image width
  64  */
  65 __kernel void minmax(
  66     IMAGE_DECLARATION(src),
  67     __global int *min_max,
  68     int           width)
  69 {
  70     Image src = CONVERT_TO_IMAGE_STRUCT(src);
  71
  72     // Initialize local minimum and local maximum
  73     VEC_DATA_TYPE(DATA_TYPE, 16)
  74     local_min = type_max;
  75     VEC_DATA_TYPE(DATA_TYPE, 16)
  76     local_max = type_min;
  77
  78     // Calculate min/max of row
  79     int i = 0;
  80     for(; i + 16 <= width; i += 16)
  81     {
  82         VEC_DATA_TYPE(DATA_TYPE, 16)
  83         data      = vload16(0, (__global DATA_TYPE *)offset(&src, i, 0));
  84         local_min = min(data, local_min);
  85         local_max = max(data, local_max);
  86     }
  87
  88 #ifdef NON_MULTIPLE_OF_16
  89     // Handle non multiple of 16
  90     VEC_DATA_TYPE(DATA_TYPE, 16)
  91     data = vload16(0, (__global DATA_TYPE *)offset(&src, i, 0));
  92 #ifdef IS_DATA_TYPE_FLOAT
  93     int16 valid_indices = (i + idx16) < width;
  94 #else  /* IS_DATA_TYPE_FLOAT */
  95     VEC_DATA_TYPE(DATA_TYPE, 16)
  96     valid_indices = CONVERT((i + idx16) < width, VEC_DATA_TYPE(DATA_TYPE, 16));
  97 #endif /* IS_DATA_TYPE_FLOAT */
  98     local_max = max(local_max, select(type_min, data, valid_indices));
  99     local_min = min(local_min, select(type_max, data, valid_indices));
 100 #endif /* NON_MULTIPLE_OF_16 */
 101
 102     // Perform min/max reduction
 103     local_min.s01234567 = min(local_min.s01234567, local_min.s89ABCDEF);
 104     local_max.s01234567 = max(local_max.s01234567, local_max.s89ABCDEF);
 105
 106     local_min.s0123 = min(local_min.s0123, local_min.s4567);
 107     local_max.s0123 = max(local_max.s0123, local_max.s4567);
 108
 109     local_min.s01 = min(local_min.s01, local_min.s23);
 110     local_max.s01 = max(local_max.s01, local_max.s23);
 111
 112     local_min.s0 = min(local_min.s0, local_min.s1);
 113     local_max.s0 = max(local_max.s0, local_max.s1);
 114
 115     // Update global min/max
 116 #ifdef IS_DATA_TYPE_FLOAT
 117     atomic_min(&min_max[0], FloatFlip(local_min.s0));
 118     atomic_max(&min_max[1], FloatFlip(local_max.s0));
 119 #else  /* IS_DATA_TYPE_FLOAT */
 120     atomic_min(&min_max[0], local_min.s0);
 121     atomic_max(&min_max[1], local_max.s0);
 122 #endif /* IS_DATA_TYPE_FLOAT */
 123 }
 124
 125 /** This function counts the min and max occurrences in an image and tags their position.
 126  *
 127  * @note -DCOUNT_MIN_MAX should be specified if we want to count the occurrences of the minimum and maximum values.
 128  * @note -DLOCATE_MIN and/or -DLOCATE_MAX should be specified if we want to store the position of each occurrence on the given array.
 129  *
 130  * @param[in]  src_ptr                           Pointer to the source image. Supported data types: U8
 131  * @param[in]  src_stride_x                      Stride of the source image in X dimension (in bytes)
 132  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 133  * @param[in]  src_stride_y                      Stride of the source image in Y dimension (in bytes)
 134  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 135  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source image
 136  * @param[in]  min_max                           Pointer to buffer with minimum value in position 0 and maximum value in position 1
 137  * @param[out] min_max_count                     Pointer to buffer with minimum value occurrences in position 0 and maximum value occurrences in position 1
 138  * @param[out] min_loc                           Array that holds the location of the minimum value occurrences
 139  * @param[in]  max_min_loc_count                 The maximum number of min value occurrences coordinates the array can hold
 140  * @param[out] max_loc                           Array that holds the location of the maximum value occurrences
 141  * @param[in]  max_max_loc_count                 The maximum number of max value occurrences coordinates the array can hold
 142  */
 143 __kernel void minmaxloc(
 144     IMAGE_DECLARATION(src),
 145     __global int *min_max,
 146     __global uint *min_max_count
 147 #ifdef LOCATE_MIN
 148     ,
 149     __global Coordinates2D *min_loc, uint max_min_loc_count
 150 #endif /* LOCATE_MIN */
 151 #ifdef LOCATE_MAX
 152     ,
 153     __global Coordinates2D *max_loc, uint max_max_loc_count
 154 #endif /* LOCATE_MAX */
 155 )
 156 {
 157     Image src = CONVERT_TO_IMAGE_STRUCT(src);
 158
 159 #ifdef IS_DATA_TYPE_FLOAT
 160     __global float *min_max_ptr = (__global float *)min_max;
 161     float           min_value   = min_max_ptr[0];
 162     float           max_value   = min_max_ptr[1];
 163 #else  /* IS_DATA_TYPE_FLOAT */
 164     int min_value = min_max[0];
 165     int max_value = min_max[1];
 166 #endif /* IS_DATA_TYPE_FLOAT */
 167
 168     DATA_TYPE value = *((__global DATA_TYPE *)src.ptr);
 169 #ifdef COUNT_MIN_MAX
 170     if(value == min_value)
 171     {
 172         uint idx = atomic_inc(&min_max_count[0]);
 173 #ifdef LOCATE_MIN
 174         if(idx < max_min_loc_count)
 175         {
 176             min_loc[idx].x = get_global_id(0);
 177             min_loc[idx].y = get_global_id(1);
 178         }
 179 #endif /* LOCATE_MIN */
 180     }
 181     if(value == max_value)
 182     {
 183         uint idx = atomic_inc(&min_max_count[1]);
 184 #ifdef LOCATE_MAX
 185         if(idx < max_max_loc_count)
 186         {
 187             max_loc[idx].x = get_global_id(0);
 188             max_loc[idx].y = get_global_id(1);
 189         }
 190 #endif /* LOCATE_MAX */
 191     }
 192 #endif /* COUNT_MIN_MAX */
 193 }