gpu::HoughLinesP

2012-12-06 15:14:20 +04:00
parent a6dc6f72b3
commit aad76090ce
7 changed files with 291 additions and 4 deletions
@@ -177,7 +177,7 @@ namespace
        const int x = blockIdx.x * blockDim.x + threadIdx.x;
        const int y = blockIdx.y * blockDim.y + threadIdx.y;

-        if (x >= dx.cols || y >= dx.rows)
+        if (x == 0 || x >= dx.cols - 1 || y == 0 || y >= dx.rows - 1)
            return;

        int dxVal = dx(y, x);
@@ -378,7 +378,7 @@ namespace
            pos.x += c_dx[threadIdx.x];
            pos.y += c_dy[threadIdx.x];

-            if (pos.x > 0 && pos.x <= map.cols && pos.y > 0 && pos.y <= map.rows && map(pos.y, pos.x) == 1)
+            if (pos.x > 0 && pos.x < map.cols && pos.y > 0 && pos.y < map.rows && map(pos.y, pos.x) == 1)
            {
                map(pos.y, pos.x) = 2;

@@ -410,7 +410,7 @@ namespace
                pos.x += c_dx[threadIdx.x & 7];
                pos.y += c_dy[threadIdx.x & 7];

-                if (pos.x > 0 && pos.x <= map.cols && pos.y > 0 && pos.y <= map.rows && map(pos.y, pos.x) == 1)
+                if (pos.x > 0 && pos.x < map.cols && pos.y > 0 && pos.y < map.rows && map(pos.y, pos.x) == 1)
                {
                    map(pos.y, pos.x) = 2;

@@ -291,6 +291,108 @@ namespace cv { namespace gpu { namespace device
            return totalCount;
        }

+        ////////////////////////////////////////////////////////////////////////
+        // houghLinesProbabilistic
+
+        texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex_mask(false, cudaFilterModePoint, cudaAddressModeClamp);
+
+        __global__ void houghLinesProbabilistic(const PtrStepSzi Dx, const PtrStepi Dy,
+                                                int4* out, const int maxSize,
+                                                const int lineGap, const int lineLength)
+        {
+            const int SHIFT = 10;
+
+            const int x = blockIdx.x * blockDim.x + threadIdx.x;
+            const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+            if (x >= Dx.cols || y >= Dx.rows || tex2D(tex_mask, x, y) == 0)
+                return;
+
+            const int dx = Dx(y, x);
+            const int dy = Dy(y, x);
+
+            if (dx == 0 && dy == 0)
+                return;
+
+            const int vx = dy;
+            const int vy = -dx;
+
+            const float mag = ::sqrtf(vx * vx + vy * vy);
+
+            const int x0 = x << SHIFT;
+            const int y0 = y << SHIFT;
+
+            int sx = __float2int_rn((vx << SHIFT) / mag);
+            int sy = __float2int_rn((vy << SHIFT) / mag);
+
+            int2 line_end[2] = {make_int2(x,y), make_int2(x,y)};
+
+            for (int k = 0; k < 2; ++k)
+            {
+                int gap = 0;
+                int x1 = x0 + sx;
+                int y1 = y0 + sy;
+
+                for (;; x1 += sx, y1 += sy)
+                {
+                    const int x2 = x1 >> SHIFT;
+                    const int y2 = y1 >> SHIFT;
+
+                    if (x2 < 0 || x2 >= Dx.cols || y2 < 0 || y2 >= Dx.rows)
+                        break;
+
+                    if (tex2D(tex_mask, x2, y2))
+                    {
+                        gap = 0;
+                        line_end[k].x = x2;
+                        line_end[k].y = y2;
+                    }
+                    else if(++gap > lineGap)
+                        break;
+                }
+
+                sx = -sx;
+                sy = -sy;
+            }
+
+            const bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength ||
+                                   ::abs(line_end[1].y - line_end[0].y) >= lineLength;
+
+            if (good_line)
+            {
+                const int ind = ::atomicAdd(&g_counter, 1);
+                if (ind < maxSize)
+                    out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
+            }
+        }
+
+        int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi Dx, PtrStepSzi Dy,
+                                        int4* out, int maxSize,
+                                        int lineGap, int lineLength)
+        {
+            void* counterPtr;
+            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
+
+            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
+
+            const dim3 block(32, 8);
+            const dim3 grid(divUp(mask.cols, block.x), divUp(mask.rows, block.y));
+
+            bindTexture(&tex_mask, mask);
+
+            houghLinesProbabilistic<<<grid, block>>>(Dx, Dy, out, maxSize, lineGap, lineLength);
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+
+            int totalCount;
+            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
+
+            totalCount = ::min(totalCount, maxSize);
+
+            return totalCount;
+        }
+
        ////////////////////////////////////////////////////////////////////////
        // circlesAccumCenters

@@ -1509,4 +1611,4 @@ namespace cv { namespace gpu { namespace device
 }}}


-#endif /* CUDA_DISABLER */
+#endif /* CUDA_DISABLER */
@@ -52,6 +52,8 @@ void cv::gpu::HoughLines(const GpuMat&, GpuMat&, float, float, int, bool, int) {
 void cv::gpu::HoughLines(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, bool, int) { throw_nogpu(); }
 void cv::gpu::HoughLinesDownload(const GpuMat&, OutputArray, OutputArray) { throw_nogpu(); }

+void cv::gpu::HoughLinesP(const GpuMat&, GpuMat&, CannyBuf&, int, int, int) { throw_nogpu(); }
+
 void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, int, float, float, int, int, int, int, int) { throw_nogpu(); }
 void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, HoughCirclesBuf&, int, float, float, int, int, int, int, int) { throw_nogpu(); }
 void cv::gpu::HoughCirclesDownload(const GpuMat&, OutputArray) { throw_nogpu(); }
@@ -155,6 +157,42 @@ void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, Ou
    }
 }

+//////////////////////////////////////////////////////////
+// HoughLinesP
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi Dx, PtrStepSzi Dy,
+                                        int4* out, int maxSize,
+                                        int lineGap, int lineLength);
+    }
+}}}
+
+void cv::gpu::HoughLinesP(const GpuMat& image, GpuMat& lines, CannyBuf& cannyBuf, int minLineLength, int maxLineGap, int maxLines)
+{
+    using namespace cv::gpu::device::hough;
+
+    CV_Assert( image.type() == CV_8UC1 );
+    CV_Assert( image.cols < std::numeric_limits<unsigned short>::max() );
+    CV_Assert( image.rows < std::numeric_limits<unsigned short>::max() );
+
+    GpuMat mask;
+    Canny(image, cannyBuf, mask, 50, 100);
+
+    ensureSizeIsEnough(1, maxLines, CV_32SC4, lines);
+
+    int linesCount = houghLinesProbabilistic_gpu(mask, cannyBuf.dx, cannyBuf.dy,
+                                                 lines.ptr<int4>(), maxLines,
+                                                 maxLineGap, minLineLength);
+
+    if (linesCount > 0)
+        lines.cols = linesCount;
+    else
+        lines.release();
+}
+
 //////////////////////////////////////////////////////////
 // HoughCircles

@@ -1484,6 +1484,7 @@ namespace
    {
        using namespace canny;

+        buf.map.setTo(Scalar::all(0));
        calcMap(dx, dy, buf.mag, buf.map, low_thresh, high_thresh);

        edgesHysteresisLocal(buf.map, buf.st1.ptr<ushort2>());