Fix OpenCL version of HoughLinesP function
authorvbystricky <vbystricky@github.com>
Wed, 5 Nov 2014 10:31:06 +0000 (14:31 +0400)
committervbystricky <vbystricky@github.com>
Wed, 5 Nov 2014 10:31:06 +0000 (14:31 +0400)
modules/imgproc/src/opencl/hough_lines.cl

index f318133..4c2d0a9 100644 (file)
@@ -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;
             }
-        }
 
+        }
     }
 }