Merge pull request #3185 from northvolt:hough_cuda_texture_obj

Use texture Object to make cv::cuda::HoughSegmentDetectorImpl::detect() thread-safe

* Use texture Object to make HoughSegmentDetectorImpl::detect() thread-safe

* Use parameter useRoi correctly in HoughLinesProbabilistic cuda test

Signed-off-by: Ron Evans <[email protected]>

Co-authored-by: airtop-bast <[email protected]>
Co-authored-by: Ron Evans <[email protected]>

Author: 3 people

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<uchar, cudaTextureType2D, cudaReadModeElementType> tex_mask(false, cudaFilterModePoint, cudaAddressModeClamp);
-
-        __global__ void houghLinesProbabilistic(const PtrStepSzi accum,
+        __global__ void houghLinesProbabilistic(cv::cudev::Texture<uchar> 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_<uchar> src_(mask);
+            cv::cudev::Texture<uchar> tex(src_, false, cudaFilterModePoint, cudaAddressModeClamp);
 
-            bindTexture(&tex_mask, mask);
-
-            houghLinesProbabilistic<<<grid, block, 0, stream>>>(accum,
+            houghLinesProbabilistic<<<grid, block, 0, stream>>>(tex, accum,
                                                      out, maxSize,
                                                      rho, theta,
                                                      lineGap, lineLength,

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);
     }
 }}}
 

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<cv::Vec4i>& 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<cv::cuda::HoughSegmentDetector> hough = cv::cuda::createHoughSegmentDetector(rho, theta, minLineLength, maxLineGap);
+
+    cv::cuda::GpuMat d_lines;
+    hough->detect(loadMat(src, useRoi), d_lines);
+
+    std::vector<cv::Vec4i> linesSegment;
+    std::vector<cv::Vec2f> 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