From fe2e89df1bed15ce2171a3ec5fd9370dcf236a41 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Wed, 13 Feb 2013 15:56:58 +0400 Subject: [PATCH] added gpu::HoughLinesP function --- modules/gpu/include/opencv2/gpu/gpu.hpp | 5 + modules/gpu/perf/perf_imgproc.cpp | 66 +++++++++++ modules/gpu/src/cuda/hough.cu | 195 ++++++++++++++++++++++++++++++++ modules/gpu/src/hough.cpp | 53 +++++++++ samples/gpu/houghlines.cpp | 89 +++++++++++++++ 5 files changed, 408 insertions(+) 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 38eafdb..090392b 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -867,6 +867,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, HoughLinesBuf& buf, float rho, float theta, 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 3f399cd..e3d488e 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 float rho = 1.0f; + const float theta = static_cast(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); + + cv::Mat mask; + cv::Canny(image, mask, 50, 100); + + if (PERF_RUN_GPU()) + { + cv::gpu::GpuMat d_mask(mask); + cv::gpu::GpuMat d_lines; + cv::gpu::HoughLinesBuf d_buf; + + cv::gpu::HoughLinesP(d_mask, d_lines, d_buf, rho, theta, minLineLenght, maxLineGap); + + TEST_CYCLE() + { + cv::gpu::HoughLinesP(d_mask, d_lines, d_buf, rho, theta, 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 + { + 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/hough.cu b/modules/gpu/src/cuda/hough.cu index 8e35aa8..695a47d 100644 --- a/modules/gpu/src/cuda/hough.cu +++ b/modules/gpu/src/cuda/hough.cu @@ -294,6 +294,201 @@ namespace cv { namespace gpu { namespace device } //////////////////////////////////////////////////////////////////////// + // houghLinesProbabilistic + + texture tex_mask(false, cudaFilterModePoint, cudaAddressModeClamp); + + __global__ void houghLinesProbabilistic(const PtrStepSzi accum, + int4* out, const int maxSize, + const float rho, const float theta, + const int lineGap, const int lineLength, + const int rows, const int cols) + { + const int r = blockIdx.x * blockDim.x + threadIdx.x; + const int n = blockIdx.y * blockDim.y + threadIdx.y; + + if (r >= accum.cols - 2 || n >= accum.rows - 2) + return; + + const int curVotes = accum(n + 1, r + 1); + + if (curVotes >= lineLength && + curVotes > accum(n, r) && + curVotes > accum(n, r + 1) && + curVotes > accum(n, r + 2) && + curVotes > accum(n + 1, r) && + curVotes > accum(n + 1, r + 2) && + curVotes > accum(n + 2, r) && + curVotes > accum(n + 2, r + 1) && + curVotes > accum(n + 2, r + 2)) + { + const float radius = (r - (accum.cols - 2 - 1) * 0.5f) * rho; + const float angle = n * theta; + + float cosa; + float sina; + sincosf(angle, &sina, &cosa); + + float2 p0 = make_float2(cosa * radius, sina * radius); + float2 dir = make_float2(-sina, cosa); + + float2 pb[4] = {make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1)}; + float a; + + if (dir.x != 0) + { + a = -p0.x / dir.x; + pb[0].x = 0; + pb[0].y = p0.y + a * dir.y; + + a = (cols - 1 - p0.x) / dir.x; + pb[1].x = cols - 1; + pb[1].y = p0.y + a * dir.y; + } + if (dir.y != 0) + { + a = -p0.y / dir.y; + pb[2].x = p0.x + a * dir.x; + pb[2].y = 0; + + a = (rows - 1 - p0.y) / dir.y; + pb[3].x = p0.x + a * dir.x; + pb[3].y = rows - 1; + } + + if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < rows)) + { + p0 = pb[0]; + if (dir.x < 0) + dir = -dir; + } + else if (pb[1].x == cols - 1 && (pb[0].y >= 0 && pb[0].y < rows)) + { + p0 = pb[1]; + if (dir.x > 0) + dir = -dir; + } + else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < cols)) + { + p0 = pb[2]; + if (dir.y < 0) + dir = -dir; + } + else if (pb[3].y == rows - 1 && (pb[3].x >= 0 && pb[3].x < cols)) + { + p0 = pb[3]; + if (dir.y > 0) + dir = -dir; + } + + float2 d; + if (::fabsf(dir.x) > ::fabsf(dir.y)) + { + d.x = dir.x > 0 ? 1 : -1; + d.y = dir.y / ::fabsf(dir.x); + } + else + { + d.x = dir.x / ::fabsf(dir.y); + d.y = dir.y > 0 ? 1 : -1; + } + + float2 line_end[2]; + int gap; + bool inLine = false; + + float2 p1 = p0; + if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows) + return; + + for (;;) + { + if (tex2D(tex_mask, p1.x, p1.y)) + { + gap = 0; + + if (!inLine) + { + line_end[0] = p1; + line_end[1] = p1; + inLine = true; + } + else + { + line_end[1] = p1; + } + } + else if (inLine) + { + if (++gap > lineGap) + { + 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); + } + + gap = 0; + inLine = false; + } + } + + p1 = p1 + d; + if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows) + { + if (inLine) + { + 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); + } + + } + break; + } + } + } + } + + int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, 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(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); + + bindTexture(&tex_mask, mask); + + houghLinesProbabilistic<<>>(accum, + out, maxSize, + rho, theta, + lineGap, lineLength, + mask.rows, mask.cols); + 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, diff --git a/modules/gpu/src/hough.cpp b/modules/gpu/src/hough.cpp index 888c325..fecb717 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&, HoughLinesBuf&, float, float, 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(); } @@ -158,6 +160,57 @@ 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 accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength); + } +}}} + +void cv::gpu::HoughLinesP(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int minLineLength, int maxLineGap, int maxLines) +{ + using namespace cv::gpu::device::hough; + + CV_Assert( src.type() == CV_8UC1 ); + CV_Assert( src.cols < std::numeric_limits::max() ); + CV_Assert( src.rows < std::numeric_limits::max() ); + + ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list); + unsigned int* srcPoints = buf.list.ptr(); + + const int pointsCount = buildPointList_gpu(src, srcPoints); + if (pointsCount == 0) + { + lines.release(); + return; + } + + const int numangle = cvRound(CV_PI / theta); + const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho); + CV_Assert( numangle > 0 && numrho > 0 ); + + ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum); + buf.accum.setTo(Scalar::all(0)); + + DeviceInfo devInfo; + cudaDeviceProp prop; + cudaSafeCall(cudaGetDeviceProperties(&prop, devInfo.deviceID())); + linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, prop.sharedMemPerBlock, devInfo.supports(FEATURE_SET_COMPUTE_20)); + + ensureSizeIsEnough(1, maxLines, CV_32SC4, lines); + + int linesCount = houghLinesProbabilistic_gpu(src, buf.accum, lines.ptr(), maxLines, rho, theta, maxLineGap, minLineLength); + + if (linesCount > 0) + lines.cols = linesCount; + else + lines.release(); +} + +////////////////////////////////////////////////////////// // HoughCircles namespace cv { namespace gpu { namespace device diff --git a/samples/gpu/houghlines.cpp b/samples/gpu/houghlines.cpp new file mode 100644 index 0000000..e98dcc6 --- /dev/null +++ b/samples/gpu/houghlines.cpp @@ -0,0 +1,89 @@ +#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, 100, 200, 3); + + Mat dst_cpu; + cvtColor(mask, dst_cpu, CV_GRAY2BGR); + Mat dst_gpu = dst_cpu.clone(); + + vector lines_cpu; + { + const int64 start = getTickCount(); + + HoughLinesP(mask, lines_cpu, 1, CV_PI / 180, 50, 60, 5); + + const double timeSec = (getTickCount() - start) / getTickFrequency(); + cout << "CPU Time : " << timeSec * 1000 << " ms" << endl; + cout << "CPU Found : " << 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(mask); + GpuMat d_lines; + HoughLinesBuf d_buf; + { + const int64 start = getTickCount(); + + gpu::HoughLinesP(d_src, d_lines, d_buf, 1.0f, (float) (CV_PI / 180.0f), 50, 5); + + const double timeSec = (getTickCount() - start) / getTickFrequency(); + cout << "GPU Time : " << timeSec * 1000 << " ms" << endl; + cout << "GPU Found : " << d_lines.cols << endl; + } + 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); + } + + 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