28 #define DATA_TYPE_MIN 0x0 32 #define DATA_TYPE_MAX 0xFF 43 return (u_val.int_val >= 0) ? u_val.int_val : u_val.int_val ^ 0x7FFFFFFF;
48 __constant int16
idx16 = (int16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
67 __global
int *min_max,
80 for(; i + 16 <= width; i += 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);
88 #ifdef NON_MULTIPLE_OF_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;
103 local_min.s01234567 = min(local_min.s01234567, local_min.s89ABCDEF);
104 local_max.s01234567 = max(local_max.s01234567, local_max.s89ABCDEF);
106 local_min.s0123 = min(local_min.s0123, local_min.s4567);
107 local_max.s0123 = max(local_max.s0123, local_max.s4567);
109 local_min.s01 = min(local_min.s01, local_min.s23);
110 local_max.s01 = max(local_max.s01, local_max.s23);
112 local_min.s0 = min(local_min.s0, local_min.s1);
113 local_max.s0 = max(local_max.s0, local_max.s1);
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));
120 atomic_min(&min_max[0], local_min.s0);
121 atomic_max(&min_max[1], local_max.s0);
145 __global
int *min_max,
146 __global uint *min_max_count
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];
164 int min_value = min_max[0];
165 int max_value = min_max[1];
170 if(value == min_value)
172 uint idx = atomic_inc(&min_max_count[0]);
174 if(idx < max_min_loc_count)
176 min_loc[idx].x = get_global_id(0);
177 min_loc[idx].y = get_global_id(1);
181 if(value == max_value)
183 uint idx = atomic_inc(&min_max_count[1]);
185 if(idx < max_max_loc_count)
187 max_loc[idx].x = get_global_id(0);
188 max_loc[idx].y = get_global_id(1);
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
#define CONVERT_TO_IMAGE_STRUCT(name)
#define IMAGE_DECLARATION(name)
SimpleTensor< float > src
__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.
Structure to hold Image information.
__constant DATA_TYPE16 type_min
__global uchar * ptr
Pointer to the starting postion of the buffer.
SimpleTensor< T > select(const SimpleTensor< uint8_t > &c, const SimpleTensor< T > &x, const SimpleTensor< T > &y)
__constant DATA_TYPE16 type_max
int int_val(TokenStream &in, bool &valid)
__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.
#define VEC_DATA_TYPE(type, size)