From 9201db32ebe87468d5a0548507c376db39ba35a4 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Tue, 14 Aug 2012 12:33:47 +0400 Subject: [PATCH] fixed overflow bugs, updated perf tests --- modules/gpu/perf/perf_imgproc.cpp | 37 +++++++++++++++------------ modules/gpu/perf_cpu/perf_imgproc.cpp | 34 +++++++++++++----------- modules/gpu/src/cuda/hough.cu | 10 +++++--- modules/gpu/src/hough.cpp | 15 +++++------ 4 files changed, 53 insertions(+), 43 deletions(-) diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp index 77f607452c..3baca6c9d3 100644 --- a/modules/gpu/perf/perf_imgproc.cpp +++ b/modules/gpu/perf/perf_imgproc.cpp @@ -1334,42 +1334,47 @@ INSTANTIATE_TEST_CASE_P(ImgProc, ImagePyramid_getLayer, testing::Combine( ////////////////////////////////////////////////////////////////////// // HoughLines -GPU_PERF_TEST(HoughLines, cv::gpu::DeviceInfo, std::string) +IMPLEMENT_PARAM_CLASS(DoSort, bool) + +GPU_PERF_TEST(HoughLines, cv::gpu::DeviceInfo, cv::Size, DoSort) { + declare.time(30.0); + const cv::gpu::DeviceInfo devInfo = GET_PARAM(0); cv::gpu::setDevice(devInfo.deviceID()); - const std::string fileName = GET_PARAM(1); + const cv::Size size = GET_PARAM(1); + const bool doSort = GET_PARAM(2); const float rho = 1.0f; const float theta = CV_PI / 180.0f; const int threshold = 300; - cv::Mat img_base = readImage(fileName, cv::IMREAD_GRAYSCALE); - ASSERT_FALSE(img_base.empty()); + cv::RNG rng(123456789); - cv::Mat img; - cv::resize(img_base, img, cv::Size(1920, 1080)); + cv::Mat src(size, CV_8UC1, cv::Scalar::all(0)); - cv::Mat edges; - cv::Canny(img, edges, 50, 200); + const int numLines = rng.uniform(500, 2000); + for (int i = 0; i < numLines; ++i) + { + cv::Point p1(rng.uniform(0, src.cols), rng.uniform(0, src.rows)); + cv::Point p2(rng.uniform(0, src.cols), rng.uniform(0, src.rows)); + cv::line(src, p1, p2, cv::Scalar::all(255), 2); + } - cv::gpu::GpuMat d_edges(edges); + cv::gpu::GpuMat d_src(src); cv::gpu::GpuMat d_lines; cv::gpu::GpuMat d_accum; - cv::gpu::HoughLines(d_edges, d_lines, d_accum, rho, theta, threshold); + cv::gpu::HoughLines(d_src, d_lines, d_accum, rho, theta, threshold, doSort); TEST_CYCLE() { - cv::gpu::HoughLines(d_edges, d_lines, d_accum, rho, theta, threshold); + cv::gpu::HoughLines(d_src, d_lines, d_accum, rho, theta, threshold, doSort); } } INSTANTIATE_TEST_CASE_P(ImgProc, HoughLines, testing::Combine( ALL_DEVICES, - testing::Values(std::string("cv/shared/pic1.png"), - std::string("cv/shared/pic3.png"), - std::string("cv/shared/pic4.png"), - std::string("cv/shared/pic5.png"), - std::string("cv/shared/pic6.png")))); + GPU_TYPICAL_MAT_SIZES, + testing::Values(DoSort(false), DoSort(true)))); #endif diff --git a/modules/gpu/perf_cpu/perf_imgproc.cpp b/modules/gpu/perf_cpu/perf_imgproc.cpp index bc77646832..1b3c0951c6 100644 --- a/modules/gpu/perf_cpu/perf_imgproc.cpp +++ b/modules/gpu/perf_cpu/perf_imgproc.cpp @@ -730,38 +730,42 @@ INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor, testing::Combine( ////////////////////////////////////////////////////////////////////// // HoughLines -GPU_PERF_TEST(HoughLines, cv::gpu::DeviceInfo, std::string) +IMPLEMENT_PARAM_CLASS(DoSort, bool) + +GPU_PERF_TEST(HoughLines, cv::gpu::DeviceInfo, cv::Size, DoSort) { - const std::string fileName = GET_PARAM(1); + declare.time(30.0); + + const cv::Size size = GET_PARAM(1); const float rho = 1.0f; const float theta = CV_PI / 180.0f; const int threshold = 300; - cv::Mat img_base = readImage(fileName, cv::IMREAD_GRAYSCALE); - ASSERT_FALSE(img_base.empty()); + cv::RNG rng(123456789); - cv::Mat img; - cv::resize(img_base, img, cv::Size(1920, 1080)); + cv::Mat src(size, CV_8UC1, cv::Scalar::all(0)); - cv::Mat edges; - cv::Canny(img, edges, 50, 200); + const int numLines = rng.uniform(500, 2000); + for (int i = 0; i < numLines; ++i) + { + cv::Point p1(rng.uniform(0, src.cols), rng.uniform(0, src.rows)); + cv::Point p2(rng.uniform(0, src.cols), rng.uniform(0, src.rows)); + cv::line(src, p1, p2, cv::Scalar::all(255), 2); + } std::vector lines; - cv::HoughLines(edges, lines, rho, theta, threshold); + cv::HoughLines(src, lines, rho, theta, threshold); TEST_CYCLE() { - cv::HoughLines(edges, lines, rho, theta, threshold); + cv::HoughLines(src, lines, rho, theta, threshold); } } INSTANTIATE_TEST_CASE_P(ImgProc, HoughLines, testing::Combine( ALL_DEVICES, - testing::Values(std::string("cv/shared/pic1.png"), - std::string("cv/shared/pic3.png"), - std::string("cv/shared/pic4.png"), - std::string("cv/shared/pic5.png"), - std::string("cv/shared/pic6.png")))); + GPU_TYPICAL_MAT_SIZES, + testing::Values(DoSort(false), DoSort(true)))); #endif diff --git a/modules/gpu/src/cuda/hough.cu b/modules/gpu/src/cuda/hough.cu index 0a439f45ec..9f5cd65348 100644 --- a/modules/gpu/src/cuda/hough.cu +++ b/modules/gpu/src/cuda/hough.cu @@ -75,12 +75,12 @@ namespace cv { namespace gpu { namespace device } } - void linesAccum_gpu(DevMem2Db src, PtrStep_ accum, float theta, int numangle, int numrho, float irho) + void linesAccum_gpu(DevMem2Db src, DevMem2D_ accum, float rho, float theta) { const dim3 block(32, 8); const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y)); - linesAccum<<>>(src, accum, theta, numangle, numrho, irho); + linesAccum<<>>(src, accum, theta, accum.rows - 2, accum.cols - 2, 1.0f / rho); cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaDeviceSynchronize() ); @@ -125,7 +125,7 @@ namespace cv { namespace gpu { namespace device } } - int linesGetResult_gpu(DevMem2D_ accum, float2* out, int* voices, int maxSize, float threshold, float theta, float rho, bool doSort) + unsigned int linesGetResult_gpu(DevMem2D_ accum, float2* out, int* voices, unsigned int maxSize, float rho, float theta, float threshold, bool doSort) { void* counter_ptr; cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) ); @@ -143,7 +143,9 @@ namespace cv { namespace gpu { namespace device uint total_count; cudaSafeCall( cudaMemcpy(&total_count, counter_ptr, sizeof(uint), cudaMemcpyDeviceToHost) ); - if (doSort) + total_count = ::min(total_count, maxSize); + + if (doSort && total_count > 0) { thrust::device_ptr out_ptr(out); thrust::device_ptr voices_ptr(voices); diff --git a/modules/gpu/src/hough.cpp b/modules/gpu/src/hough.cpp index 721e6802a8..b577ca5071 100644 --- a/modules/gpu/src/hough.cpp +++ b/modules/gpu/src/hough.cpp @@ -46,8 +46,8 @@ namespace cv { namespace gpu { namespace device { namespace hough { - void linesAccum_gpu(DevMem2Db src, PtrStep_ accum, float theta, int numangle, int numrho, float irho); - int linesGetResult_gpu(DevMem2D_ accum, float2* out, int* voices, int maxSize, float threshold, float theta, float rho, bool doSort); + void linesAccum_gpu(DevMem2Db src, DevMem2D_ accum, float rho, float theta); + unsigned int linesGetResult_gpu(DevMem2D_ accum, float2* out, int* voices, unsigned int maxSize, float rho, float theta, float threshold, bool doSort); } }}} @@ -59,12 +59,11 @@ void cv::gpu::HoughLinesTransform(const GpuMat& src, GpuMat& accum, float rho, f const int numangle = cvRound(CV_PI / theta); const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho); - const float irho = 1.0f / rho; - accum.create(numangle + 2, numrho + 2, CV_32SC1); + ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, accum); accum.setTo(cv::Scalar::all(0)); - hough::linesAccum_gpu(src, accum, theta, numangle, numrho, irho); + hough::linesAccum_gpu(src, accum, rho, theta); } void cv::gpu::HoughLinesGet(const GpuMat& accum, GpuMat& lines, float rho, float theta, int threshold, bool doSort, int maxLines) @@ -73,11 +72,11 @@ void cv::gpu::HoughLinesGet(const GpuMat& accum, GpuMat& lines, float rho, float CV_Assert(accum.type() == CV_32SC1); - lines.create(2, maxLines, CV_32FC2); - int count = hough::linesGetResult_gpu(accum, lines.ptr(0), lines.ptr(1), maxLines, threshold, theta, rho, doSort); + ensureSizeIsEnough(2, maxLines, CV_32FC2, lines); + unsigned int count = hough::linesGetResult_gpu(accum, lines.ptr(0), lines.ptr(1), maxLines, rho, theta, threshold, doSort); if (count > 0) - lines.cols = std::min(count, maxLines); + lines.cols = count; else lines.release(); }