Compute Library
 21.02
minmaxloc.cl File Reference
#include "helpers.h"
#include "types.h"

Go to the source code of this file.

Macros

#define DATA_TYPE_MIN   0x0
 
#define DATA_TYPE_MAX   0xFF
 

Functions

int FloatFlip (float val)
 
__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. More...
 
__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. More...
 

Variables

__constant DATA_TYPE16 type_min = ( DATA_TYPE16 )( 0x0 )
 
__constant DATA_TYPE16 type_max = ( DATA_TYPE16 )( 0xFF )
 
__constant int16 idx16 = (int16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)
 

Macro Definition Documentation

◆ DATA_TYPE_MAX

#define DATA_TYPE_MAX   0xFF

Definition at line 32 of file minmaxloc.cl.

◆ DATA_TYPE_MIN

#define DATA_TYPE_MIN   0x0

Definition at line 28 of file minmaxloc.cl.

Function Documentation

◆ FloatFlip()

int FloatFlip ( float  val)
inline

Definition at line 35 of file minmaxloc.cl.

References arm_compute::mlgo::parser::int_val().

Referenced by minmax().

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 }
int int_val(TokenStream &in, bool &valid)
Definition: MLGOParser.cpp:313

◆ minmax()

__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.

Note
Input image data type must be passed as a preprocessor argument using -DDATA_TYPE. Moreover, the minimum and maximum value of the given data type must be provided using -DDATA_TYPE_MIN and -DDATA_TYPE_MAX respectively.
In case image width is not a multiple of 16 then -DNON_MULTIPLE_OF_16 must be passed.
Parameters
[in]src_ptrPointer to the source image. Supported data types: U8
[in]src_stride_xStride of the source image in X dimension (in bytes)
[in]src_step_xsrc_stride_x * number of elements along X processed per workitem(in bytes)
[in]src_stride_yStride of the source image in Y dimension (in bytes)
[in]src_step_ysrc_stride_y * number of elements along Y processed per workitem(in bytes)
[in]src_offset_first_element_in_bytesThe offset of the first element in the source image
[out]min_maxPointer to buffer with minimum value in position 0 and maximum value in position 1
[in]widthInput image width

Definition at line 65 of file minmaxloc.cl.

References CONVERT, CONVERT_TO_IMAGE_STRUCT, FloatFlip(), idx16, offset(), arm_compute::test::validation::reference::select(), type_max, type_min, and VEC_DATA_TYPE.

69 {
71 
72  // Initialize local minimum and local maximum
74  local_min = type_max;
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 }
__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
for(size_t k=0;k< _target.size();++k)
Definition: Unstack.cpp:91
SimpleTensor< float > src
Definition: DFT.cpp:155
Structure to hold Image information.
Definition: helpers.h:659
__constant DATA_TYPE16 type_min
Definition: minmaxloc.cl:46
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
#define VEC_DATA_TYPE(type, size)
Definition: helpers.h:514

◆ minmaxloc()

__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.

Note
-DCOUNT_MIN_MAX should be specified if we want to count the occurrences of the minimum and maximum values.
-DLOCATE_MIN and/or -DLOCATE_MAX should be specified if we want to store the position of each occurrence on the given array.
Parameters
[in]src_ptrPointer to the source image. Supported data types: U8
[in]src_stride_xStride of the source image in X dimension (in bytes)
[in]src_step_xsrc_stride_x * number of elements along X processed per workitem(in bytes)
[in]src_stride_yStride of the source image in Y dimension (in bytes)
[in]src_step_ysrc_stride_y * number of elements along Y processed per workitem(in bytes)
[in]src_offset_first_element_in_bytesThe offset of the first element in the source image
[in]min_maxPointer to buffer with minimum value in position 0 and maximum value in position 1
[out]min_max_countPointer to buffer with minimum value occurrences in position 0 and maximum value occurrences in position 1
[out]min_locArray that holds the location of the minimum value occurrences
[in]max_min_loc_countThe maximum number of min value occurrences coordinates the array can hold
[out]max_locArray that holds the location of the maximum value occurrences
[in]max_max_loc_countThe maximum number of max value occurrences coordinates the array can hold

Definition at line 143 of file minmaxloc.cl.

References CONVERT_TO_IMAGE_STRUCT, DATA_TYPE, and arm_compute::test::validation::src.

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 }
#define CONVERT_TO_IMAGE_STRUCT(name)
Definition: helpers.h:617
#define DATA_TYPE
SimpleTensor< float > src
Definition: DFT.cpp:155
Structure to hold Image information.
Definition: helpers.h:659
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:661

Variable Documentation

◆ idx16

__constant int16 idx16 = (int16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)

Definition at line 48 of file minmaxloc.cl.

Referenced by minmax().

◆ type_max

__constant DATA_TYPE16 type_max = ( DATA_TYPE16 )( 0xFF )

Definition at line 47 of file minmaxloc.cl.

Referenced by minmax(), and NEGEMMLowpOffsetContributionOutputStageKernel::run().

◆ type_min

__constant DATA_TYPE16 type_min = ( DATA_TYPE16 )( 0x0 )