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
}
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 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>();
+ cv::Vec4i* end = h_lines.ptr<cv::Vec4i>() + h_lines.cols;
+ std::sort(begin, end, Vec4iComparator());
+ SANITY_CHECK(h_lines);
+ }
+ else
+ {
+ cv::Mat mask;
+ cv::Canny(image, mask, 50, 100);
+
+ 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);
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);
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;
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;
}
////////////////////////////////////////////////////////////////////////
+ // 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
__global__ void circlesAccumCenters(const unsigned int* list, const int count, const PtrStepi dx, const PtrStepi dy,
}}}
-#endif /* CUDA_DISABLER */
\ No newline at end of file
+#endif /* CUDA_DISABLER */
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(); }
}
//////////////////////////////////////////////////////////
+// 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
namespace cv { namespace gpu { namespace device
{
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>());
--- /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, 50, 200, 3);
+
+ Mat dst_cpu;
+ cvtColor(mask, dst_cpu, CV_GRAY2BGR);
+ Mat dst_gpu = dst_cpu.clone();
+
+ vector<Vec4i> 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<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);
+ }
+ 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;
+}
+
projects / profile / ivi / opencv.git / commitdiff