Compute Library
 21.02
minmaxloc.cl
Go to the documentation of this file.
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 
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(
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(
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 {
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 }
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
Definition: helpers.h:846
#define CONVERT(x, type)
Definition: helpers.h:517
#define CONVERT_TO_IMAGE_STRUCT(name)
Definition: helpers.h:617
__constant int16 idx16
Definition: minmaxloc.cl:48
#define DATA_TYPE
#define DATA_TYPE_MAX
Definition: minmaxloc.cl:32
2D Coordinates structure
Definition: types.h:28
#define IMAGE_DECLARATION(name)
Definition: helpers.h:581
SimpleTensor< float > src
Definition: DFT.cpp:155
__kernel void minmaxloc(__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global int *min_max, __global uint *min_max_count, __global Coordinates2D *min_loc, uint max_min_loc_count, __global Coordinates2D *max_loc, uint max_max_loc_count)
This function counts the min and max occurrences in an image and tags their position.
Definition: minmaxloc.cl:143
Structure to hold Image information.
Definition: helpers.h:659
__constant DATA_TYPE16 type_min
Definition: minmaxloc.cl:46
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:661
SimpleTensor< T > select(const SimpleTensor< uint8_t > &c, const SimpleTensor< T > &x, const SimpleTensor< T > &y)
Definition: Select.cpp:38
int FloatFlip(float val)
Definition: minmaxloc.cl:35
__constant DATA_TYPE16 type_max
Definition: minmaxloc.cl:47
int int_val(TokenStream &in, bool &valid)
Definition: MLGOParser.cpp:313
__kernel void minmax(__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global int *min_max, int width)
This function identifies the min and maximum value of an input image.
Definition: minmaxloc.cl:65
#define DATA_TYPE_MIN
Definition: minmaxloc.cl:28
#define VEC_DATA_TYPE(type, size)
Definition: helpers.h:514