71 h = vload4(0, (__global DATA_TYPE_IN *)src1.
ptr);
73 v = vload4(0, (__global DATA_TYPE_IN *)src2.
ptr);
80 float4 p = atan2pi(convert_float4(v), convert_float4(h));
83 p = select(p, p + 2, p < 0.0f) * 128.0f;
87 vstore4(convert_uchar4_sat_rte(p), 0, angle.
ptr);
134 float4 h = convert_float4(vload4(0, (__global DATA_TYPE_IN *)src1.
ptr));
135 float4 v = convert_float4(vload4(0, (__global DATA_TYPE_IN *)src2.
ptr));
138 float4 m = sqrt(h * h + v * v);
141 float4 p = atan2pi(v, h);
144 p = select(p, p + 2, p < 0.0f) * 128.0f;
148 vstore4(convert_uchar4_sat_rte(p), 0, angle.
ptr);
203 DATA_TYPE_IN gradient = *((__global DATA_TYPE_IN *)grad.
ptr);
204 uchar an = convert_ushort(*angle.
ptr);
206 if(gradient <= lower_thr)
212 uchar ang = 127 - an;
216 short2 x_p = neighbours_coords[q_an].s02;
217 short2 y_p = neighbours_coords[q_an].s13;
218 DATA_TYPE_IN g1 = *((global DATA_TYPE_IN *)
offset(&grad, x_p.x, y_p.x));
219 DATA_TYPE_IN g2 = *((global DATA_TYPE_IN *)
offset(&grad, x_p.y, y_p.y));
221 if((gradient > g1) && (gradient > g2))
228 #define hysteresis_local_stack_L1 8 // The size of level 1 stack. This has to agree with the host side 229 #define hysteresis_local_stack_L2 16 // The size of level 2 stack, adjust this can impact the match rate with VX implementation 242 #define check_pixel(early_test, x_pos, y_pos, x_cur, y_cur) \ 247 c = *((__global char *)offset(&l1_stack_counter, x_cur, y_cur)); \ 249 if(c > (hysteresis_local_stack_L1 - 1)) \ 253 if(!atomic_or((__global uint *)offset(&recorded, x_pos, y_pos), 1)) \ 255 l1_ptr[c] = (short2)(x_pos, y_pos); \ 256 *((__global char *)offset(&l1_stack_counter, x_cur, y_cur)) += 1; \ 327 int x = get_global_id(0);
328 int y = get_global_id(1);
331 DATA_TYPE_IN val = *((__global DATA_TYPE_IN *)
offset(&src, x, y));
351 if(atomic_or((__global uint *)
offset(&visited, x, y), 1) != 0)
373 int N = max(y - 1, 0), S = min(y + 1, height - 2), W = max(x - 1, 0), E = min(x + 1, width - 2);
376 x_tmp = vload4(0, (__global DATA_TYPE_IN *)
offset(&src, W, N));
377 v_tmp = vload4(0, (__global uint *)
offset(&visited, W, N));
378 check_pixel(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, N, x, y);
379 check_pixel(((x_tmp.s1 <= low_thr) || v_tmp.s1 || (x_tmp.s1 > up_thr)), x, N, x, y);
380 check_pixel(((x_tmp.s2 <= low_thr) || v_tmp.s2 || (x_tmp.s2 > up_thr)), E, N, x, y);
382 x_tmp = vload4(0, (__global DATA_TYPE_IN *)
offset(&src, W, y));
383 v_tmp = vload4(0, (__global uint *)
offset(&visited, W, y));
384 check_pixel(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, y, x, y);
385 check_pixel(((x_tmp.s2 <= low_thr) || v_tmp.s2 || (x_tmp.s2 > up_thr)), E, y, x, y);
387 x_tmp = vload4(0, (__global DATA_TYPE_IN *)
offset(&src, W, S));
388 v_tmp = vload4(0, (__global uint *)
offset(&visited, W, S));
389 check_pixel(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, S, x, y);
390 check_pixel(((x_tmp.s1 <= low_thr) || v_tmp.s1 || (x_tmp.s1 > up_thr)), x, S, x, y);
391 check_pixel(((x_tmp.s2 <= low_thr) || v_tmp.s2 || (x_tmp.s2 > up_thr)), E, S, x, y);
396 c = *((__global
char *)
offset(&l1_stack_counter, x, y));
400 *((__global
char *)
offset(&l1_stack_counter, x, y)) -= 1;
401 int2 l_c = convert_int2(l1_ptr[c - 1]);
404 stack_L2[L2_counter].x = x;
405 stack_L2[L2_counter].y = y;
426 x = stack_L2[L2_counter].x;
427 y = stack_L2[L2_counter].y;
#define hysteresis_local_stack_L1
__kernel void suppress_non_maximum(__global uchar *grad_ptr, uint grad_stride_x, uint grad_step_x, uint grad_stride_y, uint grad_step_y, uint grad_offset_first_element_in_bytes, __global uchar *angle_ptr, uint angle_stride_x, uint angle_step_x, uint angle_stride_y, uint angle_step_y, uint angle_offset_first_element_in_bytes, __global uchar *non_max_ptr, uint non_max_stride_x, uint non_max_step_x, uint non_max_stride_y, uint non_max_step_y, uint non_max_offset_first_element_in_bytes, uint lower_thr)
Perform non maximum suppression.
#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name)
#define IMAGE_DECLARATION(name)
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
__constant short4 neighbours_coords[]
Array that holds the relative coordinates offset for the neighbouring pixels.
#define CONVERT_TO_IMAGE_STRUCT(name)
#define VEC_DATA_TYPE(type, size)
Structure to hold Image information.
#define CONVERT_SAT(x, type)
__global uchar * ptr
Pointer to the starting postion of the buffer.
#define CONVERT_SAT_ROUND(x, type, round)
#define check_pixel(early_test, x_pos, y_pos, x_cur, y_cur)
Check whether pixel is valid.
#define hysteresis_local_stack_L2
__kernel void combine_gradients_L2(__global uchar *src1_ptr, uint src1_stride_x, uint src1_step_x, uint src1_stride_y, uint src1_step_y, uint src1_offset_first_element_in_bytes, __global uchar *src2_ptr, uint src2_stride_x, uint src2_step_x, uint src2_stride_y, uint src2_step_y, uint src2_offset_first_element_in_bytes, __global uchar *grad_ptr, uint grad_stride_x, uint grad_step_x, uint grad_stride_y, uint grad_step_y, uint grad_offset_first_element_in_bytes, __global uchar *angle_ptr, uint angle_stride_x, uint angle_step_x, uint angle_stride_y, uint angle_step_y, uint angle_offset_first_element_in_bytes)
Calculate the gradient and angle from horizontal and vertical result of sobel result.
__kernel void combine_gradients_L1(__global uchar *src1_ptr, uint src1_stride_x, uint src1_step_x, uint src1_stride_y, uint src1_step_y, uint src1_offset_first_element_in_bytes, __global uchar *src2_ptr, uint src2_stride_x, uint src2_step_x, uint src2_stride_y, uint src2_step_y, uint src2_offset_first_element_in_bytes, __global uchar *grad_ptr, uint grad_stride_x, uint grad_step_x, uint grad_stride_y, uint grad_step_y, uint grad_offset_first_element_in_bytes, __global uchar *angle_ptr, uint angle_stride_x, uint angle_step_x, uint angle_stride_y, uint angle_step_y, uint angle_offset_first_element_in_bytes)
Calculate the magnitude and phase from horizontal and vertical result of sobel result.
int stride_y
Stride of the image in Y dimension (in bytes)
kernel void hysteresis(__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 uchar *out_ptr, uint out_stride_x, uint out_step_x, uint out_stride_y, uint out_step_y, uint out_offset_first_element_in_bytes, __global uchar *visited_ptr, uint visited_stride_x, uint visited_step_x, uint visited_stride_y, uint visited_step_y, uint visited_offset_first_element_in_bytes, __global uchar *recorded_ptr, uint recorded_stride_x, uint recorded_step_x, uint recorded_stride_y, uint recorded_step_y, uint recorded_offset_first_element_in_bytes, __global uchar *l1_stack_ptr, uint l1_stack_stride_x, uint l1_stack_step_x, uint l1_stack_stride_y, uint l1_stack_step_y, uint l1_stack_offset_first_element_in_bytes, __global uchar *l1_stack_counter_ptr, uint l1_stack_counter_stride_x, uint l1_stack_counter_step_x, uint l1_stack_counter_stride_y, uint l1_stack_counter_step_y, uint l1_stack_counter_offset_first_element_in_bytes, uint low_thr, uint up_thr, int width, int height)
Perform hysteresis.
int stride_x
Stride of the image in X dimension (in bytes)