From 2edcd7eaef2c8122b461ecd65ada488f5cd3513e Mon Sep 17 00:00:00 2001 From: airtop-bast Date: Wed, 23 Feb 2022 11:10:20 +0100 Subject: [PATCH 1/2] Use texture Object to make HoughSegmentDetectorImpl::detect() thread-safe --- .../cudaimgproc/src/cuda/hough_segments.cu | 16 ++--- modules/cudaimgproc/src/hough_segments.cpp | 2 +- modules/cudaimgproc/test/test_hough.cpp | 62 +++++++++++++++++++ 3 files changed, 71 insertions(+), 9 deletions(-) diff --git a/modules/cudaimgproc/src/cuda/hough_segments.cu b/modules/cudaimgproc/src/cuda/hough_segments.cu index 988f14c1d13..59eb78f6996 100644 --- a/modules/cudaimgproc/src/cuda/hough_segments.cu +++ b/modules/cudaimgproc/src/cuda/hough_segments.cu @@ -44,14 +44,13 @@ #include "opencv2/core/cuda/common.hpp" #include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/cudev.hpp" namespace cv { namespace cuda { namespace device { namespace hough_segments { - texture tex_mask(false, cudaFilterModePoint, cudaAddressModeClamp); - - __global__ void houghLinesProbabilistic(const PtrStepSzi accum, + __global__ void houghLinesProbabilistic(cv::cudev::Texture src, const PtrStepSzi accum, int4* out, const int maxSize, const float rho, const float theta, const int lineGap, const int lineLength, @@ -157,7 +156,7 @@ namespace cv { namespace cuda { namespace device for (;;) { - if (tex2D(tex_mask, p1.x, p1.y)) + if (src(p1.y, p1.x)) { gap = 0; @@ -213,16 +212,17 @@ namespace cv { namespace cuda { namespace device } } - int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength, int* counterPtr, cudaStream_t stream) + int houghLinesProbabilistic_gpu(GpuMat &mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength, int* counterPtr, cudaStream_t stream) { cudaSafeCall( cudaMemsetAsync(counterPtr, 0, sizeof(int), stream) ); const dim3 block(32, 8); const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); + + cv::cudev::GpuMat_ src_(mask); + cv::cudev::Texture tex(src_, false, cudaFilterModePoint, cudaAddressModeClamp); - bindTexture(&tex_mask, mask); - - houghLinesProbabilistic<<>>(accum, + houghLinesProbabilistic<<>>(tex, accum, out, maxSize, rho, theta, lineGap, lineLength, diff --git a/modules/cudaimgproc/src/hough_segments.cpp b/modules/cudaimgproc/src/hough_segments.cpp index f0eb0beecd7..70bd84f00ae 100644 --- a/modules/cudaimgproc/src/hough_segments.cpp +++ b/modules/cudaimgproc/src/hough_segments.cpp @@ -65,7 +65,7 @@ namespace cv { namespace cuda { namespace device namespace hough_segments { - int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength, int* counterPtr, cudaStream_t stream); + int houghLinesProbabilistic_gpu(GpuMat &mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength, int* counterPtr, cudaStream_t stream); } }}} diff --git a/modules/cudaimgproc/test/test_hough.cpp b/modules/cudaimgproc/test/test_hough.cpp index 05c5bba2390..80080cbb2ed 100644 --- a/modules/cudaimgproc/test/test_hough.cpp +++ b/modules/cudaimgproc/test/test_hough.cpp @@ -113,6 +113,68 @@ INSTANTIATE_TEST_CASE_P(CUDA_ImgProc, HoughLines, testing::Combine( DIFFERENT_SIZES, WHOLE_SUBMAT)); +/////////////////////////////////////////////////////////////////////////////////////////////////////// +// HoughLines Probabilistic +PARAM_TEST_CASE(HoughLinesProbabilistic, cv::cuda::DeviceInfo, cv::Size, UseRoi) +{ + static void generateLines(cv::Mat& img) + { + img.setTo(cv::Scalar::all(0)); + + cv::line(img, cv::Point(10, 0), cv::Point(10, 20), cv::Scalar::all(255)); + cv::line(img, cv::Point(20, 50), cv::Point(40, 50), cv::Scalar::all(255)); + } + + static void drawLines(cv::Mat& dst, const std::vector& linesSegment) + { + dst.setTo(cv::Scalar::all(0)); + + for (size_t i = 0; i < linesSegment.size(); ++i) + { + Vec4i l = linesSegment[i]; + cv::line(dst, cv::Point(l[0], l[1]), cv::Point(l[2], l[3]), cv::Scalar::all(255)); + } + + } +}; + +CUDA_TEST_P(HoughLinesProbabilistic, Accuracy) +{ + const cv::cuda::DeviceInfo devInfo = GET_PARAM(0); + cv::cuda::setDevice(devInfo.deviceID()); + const cv::Size size = GET_PARAM(1); + const bool useRoi = GET_PARAM(2); + + const float rho = 1.0f; + const float theta = (float) (1.0 * CV_PI / 180.0); + const int minLineLength = 15; + const int maxLineGap = 8; + + cv::Mat src(size, CV_8UC1); + generateLines(src); + + Ptr hough = cv::cuda::createHoughSegmentDetector(rho, theta, minLineLength, maxLineGap); + + cv::cuda::GpuMat d_lines; + hough->detect(loadMat(src, false), d_lines); + + std::vector linesSegment; + std::vector lines; + d_lines.download(linesSegment); + + cv::Mat dst(size, CV_8UC1); + drawLines(dst, linesSegment); + + ASSERT_MAT_NEAR(src, dst, 0.0); + +} + +INSTANTIATE_TEST_CASE_P(CUDA_ImgProc, HoughLinesProbabilistic, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + WHOLE_SUBMAT)); + + /////////////////////////////////////////////////////////////////////////////////////////////////////// // HoughCircles From 4d05fb56276a9aa6ce32cd1a421a57b4d69e91ce Mon Sep 17 00:00:00 2001 From: Ron Evans Date: Thu, 3 Mar 2022 18:44:36 +0100 Subject: [PATCH 2/2] Use parameter useRoi correctly in HoughLinesProbabilistic cuda test Signed-off-by: Ron Evans --- modules/cudaimgproc/test/test_hough.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/cudaimgproc/test/test_hough.cpp b/modules/cudaimgproc/test/test_hough.cpp index 80080cbb2ed..e6a05f578f6 100644 --- a/modules/cudaimgproc/test/test_hough.cpp +++ b/modules/cudaimgproc/test/test_hough.cpp @@ -156,7 +156,7 @@ CUDA_TEST_P(HoughLinesProbabilistic, Accuracy) Ptr hough = cv::cuda::createHoughSegmentDetector(rho, theta, minLineLength, maxLineGap); cv::cuda::GpuMat d_lines; - hough->detect(loadMat(src, false), d_lines); + hough->detect(loadMat(src, useRoi), d_lines); std::vector linesSegment; std::vector lines;