From 957e5ef8eba63f5b7f18861242715cc2bb690f7c Mon Sep 17 00:00:00 2001 From: vbystricky Date: Wed, 5 Nov 2014 14:31:06 +0400 Subject: [PATCH] Fix OpenCL version of HoughLinesP function --- modules/imgproc/src/opencl/hough_lines.cl | 229 +++++++++++++++--------------- 1 file changed, 116 insertions(+), 113 deletions(-) diff --git a/modules/imgproc/src/opencl/hough_lines.cl b/modules/imgproc/src/opencl/hough_lines.cl index f318133..4c2d0a9 100644 --- a/modules/imgproc/src/opencl/hough_lines.cl +++ b/modules/imgproc/src/opencl/hough_lines.cl @@ -184,146 +184,149 @@ __kernel void get_lines(__global const uchar * accum_ptr, int accum_step, int ac int x = get_global_id(0); int y = get_global_id(1); - __global uchar* accum = accum_ptr + mad24(y+1, accum_step, mad24(x+1, (int) sizeof(int), accum_offset)); - __global int4* lines = (__global int4*)(lines_ptr + lines_offset); - __global int* lines_index = lines_index_ptr + 1; - - int curVote = ACCUM(accum); - - if (curVote >= threshold && - curVote > ACCUM(accum - accum_step - sizeof(int)) && - curVote > ACCUM(accum - accum_step) && - curVote > ACCUM(accum - accum_step + sizeof(int)) && - curVote > ACCUM(accum - sizeof(int)) && - curVote > ACCUM(accum + sizeof(int)) && - curVote > ACCUM(accum + accum_step - sizeof(int)) && - curVote > ACCUM(accum + accum_step) && - curVote > ACCUM(accum + accum_step + sizeof(int))) + if (y < accum_rows-2) { - const float radius = (x - (accum_cols - 2 - 1) * 0.5f) * rho; - const float angle = y * theta; - - float cosa; - float sina = sincos(angle, &cosa); - - float2 p0 = (float2)(cosa * radius, sina * radius); - float2 dir = (float2)(-sina, cosa); - - float2 pb[4] = { (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1) }; - float a; + __global uchar* accum = accum_ptr + mad24(y+1, accum_step, mad24(x+1, (int) sizeof(int), accum_offset)); + __global int4* lines = (__global int4*)(lines_ptr + lines_offset); + __global int* lines_index = lines_index_ptr + 1; - if (dir.x != 0) + int curVote = ACCUM(accum); + + if (curVote >= threshold && + curVote > ACCUM(accum - accum_step - sizeof(int)) && + curVote > ACCUM(accum - accum_step) && + curVote > ACCUM(accum - accum_step + sizeof(int)) && + curVote > ACCUM(accum - sizeof(int)) && + curVote > ACCUM(accum + sizeof(int)) && + curVote > ACCUM(accum + accum_step - sizeof(int)) && + curVote > ACCUM(accum + accum_step) && + curVote > ACCUM(accum + accum_step + sizeof(int))) { - a = -p0.x / dir.x; - pb[0].x = 0; - pb[0].y = p0.y + a * dir.y; + const float radius = (x - (accum_cols - 2 - 1) * 0.5f) * rho; + const float angle = y * theta; - a = (src_cols - 1 - p0.x) / dir.x; - pb[1].x = src_cols - 1; - pb[1].y = p0.y + a * dir.y; - } + float cosa; + float sina = sincos(angle, &cosa); - if (dir.y != 0) - { - a = -p0.y / dir.y; - pb[2].x = p0.x + a * dir.x; - pb[2].y = 0; + float2 p0 = (float2)(cosa * radius, sina * radius); + float2 dir = (float2)(-sina, cosa); - a = (src_rows - 1 - p0.y) / dir.y; - pb[3].x = p0.x + a * dir.x; - pb[3].y = src_rows - 1; - } + float2 pb[4] = { (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1) }; + float a; - if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < src_rows)) - { - p0 = pb[0]; - if (dir.x < 0) - dir = -dir; - } - else if (pb[1].x == src_cols - 1 && (pb[1].y >= 0 && pb[1].y < src_rows)) - { - p0 = pb[1]; - if (dir.x > 0) - dir = -dir; - } - else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < src_cols)) - { - p0 = pb[2]; - if (dir.y < 0) - dir = -dir; - } - else if (pb[3].y == src_rows - 1 && (pb[3].x >= 0 && pb[3].x < src_cols)) - { - p0 = pb[3]; - if (dir.y > 0) - dir = -dir; - } + if (dir.x != 0) + { + a = -p0.x / dir.x; + pb[0].x = 0; + pb[0].y = p0.y + a * dir.y; - dir /= max(fabs(dir.x), fabs(dir.y)); + a = (src_cols - 1 - p0.x) / dir.x; + pb[1].x = src_cols - 1; + pb[1].y = p0.y + a * dir.y; + } - float2 line_end[2]; - int gap; - bool inLine = false; + if (dir.y != 0) + { + a = -p0.y / dir.y; + pb[2].x = p0.x + a * dir.x; + pb[2].y = 0; - if (p0.x < 0 || p0.x >= src_cols || p0.y < 0 || p0.y >= src_rows) - return; + a = (src_rows - 1 - p0.y) / dir.y; + pb[3].x = p0.x + a * dir.x; + pb[3].y = src_rows - 1; + } - for (;;) - { - if (*(src_ptr + mad24(p0.y, src_step, p0.x + src_offset))) + if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < src_rows)) { - gap = 0; - - if (!inLine) - { - line_end[0] = p0; - line_end[1] = p0; - inLine = true; - } - else - { - line_end[1] = p0; - } + p0 = pb[0]; + if (dir.x < 0) + dir = -dir; } - else if (inLine) + else if (pb[1].x == src_cols - 1 && (pb[1].y >= 0 && pb[1].y < src_rows)) { - if (++gap > lineGap) - { - bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength || - fabs(line_end[1].y - line_end[0].y) >= lineLength; + p0 = pb[1]; + if (dir.x > 0) + dir = -dir; + } + else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < src_cols)) + { + p0 = pb[2]; + if (dir.y < 0) + dir = -dir; + } + else if (pb[3].y == src_rows - 1 && (pb[3].x >= 0 && pb[3].x < src_cols)) + { + p0 = pb[3]; + if (dir.y > 0) + dir = -dir; + } - if (good_line) - { - int index = atomic_inc(lines_index); - if (index < linesMax) - lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y); - } + dir /= max(fabs(dir.x), fabs(dir.y)); - gap = 0; - inLine = false; - } - } + float2 line_end[2]; + int gap; + bool inLine = false; - p0 = p0 + dir; if (p0.x < 0 || p0.x >= src_cols || p0.y < 0 || p0.y >= src_rows) + return; + + for (;;) { - if (inLine) + if (*(src_ptr + mad24(p0.y, src_step, p0.x + src_offset))) { - bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength || - fabs(line_end[1].y - line_end[0].y) >= lineLength; + gap = 0; - if (good_line) + if (!inLine) + { + line_end[0] = p0; + line_end[1] = p0; + inLine = true; + } + else { - int index = atomic_inc(lines_index); - if (index < linesMax) - lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y); + line_end[1] = p0; } + } + else if (inLine) + { + if (++gap > lineGap) + { + bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength || + fabs(line_end[1].y - line_end[0].y) >= lineLength; + + if (good_line) + { + int index = atomic_inc(lines_index); + if (index < linesMax) + lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y); + } + + gap = 0; + inLine = false; + } + } + p0 = p0 + dir; + if (p0.x < 0 || p0.x >= src_cols || p0.y < 0 || p0.y >= src_rows) + { + if (inLine) + { + bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength || + fabs(line_end[1].y - line_end[0].y) >= lineLength; + + if (good_line) + { + int index = atomic_inc(lines_index); + if (index < linesMax) + lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y); + } + + } + break; } - break; } - } + } } } -- 2.7.4