Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

fixed overflow bugs, updated perf tests

Browse Source
pull/11/head
Vladislav Vinogradov 12 years ago
parent 1eefc69927
commit 9201db32eb
4 changed files with 53 additions and 43 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 37
      modules/gpu/perf/perf_imgproc.cpp
  2. 34
      modules/gpu/perf_cpu/perf_imgproc.cpp
  3. 10
      modules/gpu/src/cuda/hough.cu
  4. 15
      modules/gpu/src/hough.cpp

37
modules/gpu/perf/perf_imgproc.cpp
Unescape View File

@ -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

34
modules/gpu/perf_cpu/perf_imgproc.cpp
Unescape View File

@ -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<cv::Vec2f> 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

10
modules/gpu/src/cuda/hough.cu
Unescape View File

@ -75,12 +75,12 @@ namespace cv { namespace gpu { namespace device
}
}
void linesAccum_gpu(DevMem2Db src, PtrStep_<uint> accum, float theta, int numangle, int numrho, float irho)
void linesAccum_gpu(DevMem2Db src, DevMem2D_<uint> accum, float rho, float theta)
{
const dim3 block(32, 8);
const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
linesAccum<<<grid, block>>>(src, accum, theta, numangle, numrho, irho);
linesAccum<<<grid, block>>>(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_<uint> accum, float2* out, int* voices, int maxSize, float threshold, float theta, float rho, bool doSort)
unsigned int linesGetResult_gpu(DevMem2D_<uint> 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<float2> out_ptr(out);
thrust::device_ptr<int> voices_ptr(voices);

15
modules/gpu/src/hough.cpp
Unescape View File

@ -46,8 +46,8 @@ namespace cv { namespace gpu { namespace device
{
namespace hough
{
void linesAccum_gpu(DevMem2Db src, PtrStep_<uint> accum, float theta, int numangle, int numrho, float irho);
int linesGetResult_gpu(DevMem2D_<uint> accum, float2* out, int* voices, int maxSize, float threshold, float theta, float rho, bool doSort);
void linesAccum_gpu(DevMem2Db src, DevMem2D_<uint> accum, float rho, float theta);
unsigned int linesGetResult_gpu(DevMem2D_<uint> 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<float2>(0), lines.ptr<int>(1), maxLines, threshold, theta, rho, doSort);
ensureSizeIsEnough(2, maxLines, CV_32FC2, lines);
unsigned int count = hough::linesGetResult_gpu(accum, lines.ptr<float2>(0), lines.ptr<int>(1), maxLines, rho, theta, threshold, doSort);
if (count > 0)
lines.cols = std::min(count, maxLines);
lines.cols = count;
else
lines.release();
}

Write Preview
Loading…
Cancel
Save