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
}
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
}
//////////////////////////////////////////////////////////////////////
+// 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<float>(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>();
+ cv::Vec4i* end = h_lines.ptr<cv::Vec4i>() + h_lines.cols;
+ std::sort(begin, end, Vec4iComparator());
+ SANITY_CHECK(h_lines);
+ }
+ else
+ {
+ std::vector<cv::Vec4i> 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);
}
////////////////////////////////////////////////////////////////////////
+ // houghLinesProbabilistic
+
+ texture<uchar, cudaTextureType2D, cudaReadModeElementType> 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<<<grid, block>>>(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,
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(); }
}
//////////////////////////////////////////////////////////
+// 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<unsigned short>::max() );
+ CV_Assert( src.rows < std::numeric_limits<unsigned short>::max() );
+
+ ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list);
+ unsigned int* srcPoints = buf.list.ptr<unsigned int>();
+
+ 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<int4>(), maxLines, rho, theta, maxLineGap, minLineLength);
+
+ if (linesCount > 0)
+ lines.cols = linesCount;
+ else
+ lines.release();
+}
+
+//////////////////////////////////////////////////////////
// HoughCircles
namespace cv { namespace gpu { namespace device
--- /dev/null
+#include <cmath>
+#include <iostream>
+
+#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 <image_name>, 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<Vec4i> 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<Vec4i> 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;
+}
+
projects / platform / upstream / opencv.git / commitdiff