From aad76090ce96967f07cb87f8b7152b23190496d8 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Thu, 6 Dec 2012 15:14:20 +0400 Subject: [PATCH] gpu::HoughLinesP --- modules/gpu/include/opencv2/gpu/gpu.hpp | 5 ++ modules/gpu/perf/perf_imgproc.cpp | 66 ++++++++++++++++++++ modules/gpu/src/cuda/canny.cu | 6 +- modules/gpu/src/cuda/hough.cu | 104 +++++++++++++++++++++++++++++++- modules/gpu/src/hough.cpp | 38 ++++++++++++ modules/gpu/src/imgproc.cpp | 1 + samples/gpu/houghlines.cpp | 75 +++++++++++++++++++++++ 7 files changed, 291 insertions(+), 4 deletions(-) create mode 100644 samples/gpu/houghlines.cpp diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 6a477dc..2bbfe53 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -847,6 +847,11 @@ CV_EXPORTS void HoughLines(const GpuMat& src, GpuMat& lines, float rho, float th CV_EXPORTS void HoughLines(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int threshold, bool doSort = false, int maxLines = 4096); CV_EXPORTS void HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines, OutputArray h_votes = noArray()); +//! HoughLinesP + +//! finds line segments in the black-n-white image using probabalistic Hough transform +CV_EXPORTS void HoughLinesP(const GpuMat& image, GpuMat& lines, CannyBuf& cannyBuf, int minLineLength, int maxLineGap, int maxLines = 4096); + //! HoughCircles struct HoughCirclesBuf diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp index a9c8b59..b8feab8 100644 --- a/modules/gpu/perf/perf_imgproc.cpp +++ b/modules/gpu/perf/perf_imgproc.cpp @@ -1706,6 +1706,16 @@ PERF_TEST_P(Sz_Depth_Cn, ImgProc_ImagePyramidGetLayer, Combine(GPU_TYPICAL_MAT_S } namespace { + struct Vec4iComparator + { + bool operator()(const cv::Vec4i& a, const cv::Vec4i b) const + { + if (a[0] != b[0]) return a[0] < b[0]; + else if(a[1] != b[1]) return a[1] < b[1]; + else if(a[2] != b[2]) return a[2] < b[2]; + else return a[3] < b[3]; + } + }; struct Vec3fComparator { bool operator()(const cv::Vec3f& a, const cv::Vec3f b) const @@ -1785,6 +1795,62 @@ PERF_TEST_P(Sz, ImgProc_HoughLines, GPU_TYPICAL_MAT_SIZES) } ////////////////////////////////////////////////////////////////////// +// HoughLinesP + +DEF_PARAM_TEST_1(Image, std::string); + +PERF_TEST_P(Image, ImgProc_HoughLinesP, testing::Values("cv/shared/pic5.png", "stitching/a1.png")) +{ + declare.time(30.0); + + std::string fileName = getDataPath(GetParam()); + + const double rho = 1.0f; + const double theta = CV_PI / 180.0; + const int threshold = 100; + const int minLineLenght = 50; + const int maxLineGap = 5; + + cv::Mat image = cv::imread(fileName, cv::IMREAD_GRAYSCALE); + + if (PERF_RUN_GPU()) + { + cv::gpu::GpuMat d_image(image); + cv::gpu::GpuMat d_lines; + cv::gpu::CannyBuf d_buf; + + cv::gpu::HoughLinesP(d_image, d_lines, d_buf, minLineLenght, maxLineGap); + + TEST_CYCLE() + { + cv::gpu::HoughLinesP(d_image, d_lines, d_buf, minLineLenght, maxLineGap); + } + + cv::Mat h_lines(d_lines); + cv::Vec4i* begin = h_lines.ptr(); + cv::Vec4i* end = h_lines.ptr() + h_lines.cols; + std::sort(begin, end, Vec4iComparator()); + SANITY_CHECK(h_lines); + } + else + { + cv::Mat mask; + cv::Canny(image, mask, 50, 100); + + std::vector lines; + cv::HoughLinesP(mask, lines, rho, theta, threshold, minLineLenght, maxLineGap); + + TEST_CYCLE() + { + cv::HoughLinesP(mask, lines, rho, theta, threshold, minLineLenght, maxLineGap); + } + + std::sort(lines.begin(), lines.end(), Vec4iComparator()); + SANITY_CHECK(lines); + } +} + +////////////////////////////////////////////////////////////////////// // HoughCircles DEF_PARAM_TEST(Sz_Dp_MinDist, cv::Size, float, float); diff --git a/modules/gpu/src/cuda/canny.cu b/modules/gpu/src/cuda/canny.cu index b08a61c..ae9cac8 100644 --- a/modules/gpu/src/cuda/canny.cu +++ b/modules/gpu/src/cuda/canny.cu @@ -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; diff --git a/modules/gpu/src/cuda/hough.cu b/modules/gpu/src/cuda/hough.cu index ee4d025..4835e35 100644 --- a/modules/gpu/src/cuda/hough.cu +++ b/modules/gpu/src/cuda/hough.cu @@ -292,6 +292,108 @@ namespace cv { namespace gpu { namespace device } //////////////////////////////////////////////////////////////////////// + // houghLinesProbabilistic + + texture 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<<>>(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 __global__ void circlesAccumCenters(const unsigned int* list, const int count, const PtrStepi dx, const PtrStepi dy, @@ -1509,4 +1611,4 @@ namespace cv { namespace gpu { namespace device }}} -#endif /* CUDA_DISABLER */ \ No newline at end of file +#endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/hough.cpp b/modules/gpu/src/hough.cpp index 6e2170d..daeb625 100644 --- a/modules/gpu/src/hough.cpp +++ b/modules/gpu/src/hough.cpp @@ -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(); } @@ -156,6 +158,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::max() ); + CV_Assert( image.rows < std::numeric_limits::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(), maxLines, + maxLineGap, minLineLength); + + if (linesCount > 0) + lines.cols = linesCount; + else + lines.release(); +} + +////////////////////////////////////////////////////////// // HoughCircles namespace cv { namespace gpu { namespace device diff --git a/modules/gpu/src/imgproc.cpp b/modules/gpu/src/imgproc.cpp index b216de8..2c14f67 100644 --- a/modules/gpu/src/imgproc.cpp +++ b/modules/gpu/src/imgproc.cpp @@ -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()); diff --git a/samples/gpu/houghlines.cpp b/samples/gpu/houghlines.cpp new file mode 100644 index 0000000..e9bb969 --- /dev/null +++ b/samples/gpu/houghlines.cpp @@ -0,0 +1,75 @@ +#include +#include + +#include "opencv2/core/core.hpp" +#include "opencv2/highgui/highgui.hpp" +#include "opencv2/imgproc/imgproc.hpp" +#include "opencv2/gpu/gpu.hpp" + +using namespace std; +using namespace cv; +using namespace cv::gpu; + +static void help() +{ + cout << "This program demonstrates line finding with the Hough transform." << endl; + cout << "Usage:" << endl; + cout << "./gpu-example-houghlines , Default is pic1.png\n" << endl; +} + +int main(int argc, const char* argv[]) +{ + const string filename = argc >= 2 ? argv[1] : "pic1.png"; + + Mat src = imread(filename, IMREAD_GRAYSCALE); + if (src.empty()) + { + help(); + cout << "can not open " << filename << endl; + return -1; + } + + Mat mask; + Canny(src, mask, 50, 200, 3); + + Mat dst_cpu; + cvtColor(mask, dst_cpu, CV_GRAY2BGR); + Mat dst_gpu = dst_cpu.clone(); + + vector lines_cpu; + HoughLinesP(mask, lines_cpu, 1, CV_PI / 180, 50, 50, 5); + cout << lines_cpu.size() << endl; + + for (size_t i = 0; i < lines_cpu.size(); ++i) + { + Vec4i l = lines_cpu[i]; + line(dst_cpu, Point(l[0], l[1]), Point(l[2], l[3]), Scalar(0, 0, 255), 3, CV_AA); + } + + GpuMat d_src(src); + GpuMat d_lines; + CannyBuf d_buf; + gpu::HoughLinesP(d_src, d_lines, d_buf, 50, 5); + vector lines_gpu; + if (!d_lines.empty()) + { + lines_gpu.resize(d_lines.cols); + Mat h_lines(1, d_lines.cols, CV_32SC4, &lines_gpu[0]); + d_lines.download(h_lines); + } + cout << lines_gpu.size() << endl; + + for (size_t i = 0; i < lines_gpu.size(); ++i) + { + Vec4i l = lines_gpu[i]; + line(dst_gpu, Point(l[0], l[1]), Point(l[2], l[3]), Scalar(0, 0, 255), 3, CV_AA); + } + + imshow("source", src); + imshow("detected lines [CPU]", dst_cpu); + imshow("detected lines [GPU]", dst_gpu); + waitKey(); + + return 0; +} + -- 2.7.4