From 9adfc2cadc0627cfc5c1045f2041c25033adc347 Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Tue, 30 Nov 2010 08:04:37 +0000 Subject: [PATCH] added Harris corner detector into gpu module --- modules/gpu/include/opencv2/gpu/gpu.hpp | 4 ++ modules/gpu/src/cuda/imgproc.cu | 48 ++++++++++++++++++++ modules/gpu/src/imgproc_gpu.cpp | 33 ++++++++++++++ tests/gpu/src/imgproc_gpu.cpp | 58 +++++++++++++++++++++++++ 4 files changed, 143 insertions(+) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index cdc9e4a635..043f163ed2 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -627,6 +627,10 @@ namespace cv //! disabled until fix crash CV_EXPORTS void Canny(const GpuMat& image, GpuMat& edges, double threshold1, double threshold2, int apertureSize = 3); + //! computes Harris cornerness criteria at each image pixel + // (does BORDER_CONSTANT interpolation with 0 as the fill value) + CV_EXPORTS void cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int apertureSize, double k); + //////////////////////////////// Filter Engine //////////////////////////////// /*! diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index 232611962f..1e632a2211 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -463,4 +463,52 @@ namespace cv { namespace gpu { namespace imgproc { reprojectImageTo3D_caller(disp, xyzw, q, stream); } + +/////////////////////////////////////////// Corner Harris ///////////////////////////////////////////////// + + __global__ void cornerHarris_kernel(const int cols, const int rows, const int block_size, const float k, const PtrStep Dx, const PtrStep Dy, PtrStep dst) + { + const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; + const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x < cols && y < rows) + { + float a = 0.f; + float b = 0.f; + float c = 0.f; + + const unsigned int j_begin = max(x - block_size, 0); + const unsigned int i_begin = max(y - block_size, 0); + const unsigned int j_end = min(x + block_size + 1, cols); + const unsigned int i_end = min(y + block_size + 1, rows); + + for (unsigned int i = i_begin; i < i_end; ++i) + { + const float* dx_row = (const float*)Dx.ptr(i); + const float* dy_row = (const float*)Dy.ptr(i); + for (unsigned int j = j_begin; j < j_end; ++j) + { + float dx = dx_row[j]; + float dy = dy_row[j]; + a += dx * dx; + b += dx * dy; + c += dy * dy; + } + } + + ((float*)dst.ptr(y))[x] = a * c - b * b - k * (a + c) * (a + c); + } + } + + void cornerHarris_caller(const int block_size, const float k, const DevMem2D Dx, const DevMem2D Dy, DevMem2D dst) + { + const int rows = Dx.rows; + const int cols = Dx.cols; + + dim3 threads(32, 8); + dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); + + cornerHarris_kernel<<>>(cols, rows, block_size / 2, k, Dx, Dy, dst); + cudaSafeCall(cudaThreadSynchronize()); + } }}} diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index 74f30144ab..00696d8eaa 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/modules/gpu/src/imgproc_gpu.cpp @@ -68,6 +68,8 @@ void cv::gpu::histEven(const GpuMat&, GpuMat&, int, int, int) { throw_nogpu(); } void cv::gpu::histEven(const GpuMat&, GpuMat*, int*, int*, int*) { throw_nogpu(); } void cv::gpu::histRange(const GpuMat&, GpuMat&, const GpuMat&) { throw_nogpu(); } void cv::gpu::histRange(const GpuMat&, GpuMat*, const GpuMat*) { throw_nogpu(); } +void cv::gpu::cornerHarris(const GpuMat&, GpuMat&, int, int, double) { throw_nogpu(); } + #else /* !defined (HAVE_CUDA) */ @@ -856,4 +858,35 @@ void cv::gpu::histRange(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4 hist_callers[src.depth()](src, hist, levels); } +namespace cv { namespace gpu { namespace imgproc { + + void cornerHarris_caller(const int block_size, const float k, const DevMem2D Dx, const DevMem2D Dy, DevMem2D dst); + +}}} + +void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int apertureSize, double k) +{ + CV_Assert(src.type() == CV_32F); + + double scale = (double)(1 << ((apertureSize > 0 ? apertureSize : 3) - 1)) * blockSize; + if (apertureSize < 0) scale *= 2.; + scale = 1./scale; + + GpuMat Dx, Dy; + if (apertureSize > 0) + { + Sobel(src, Dx, CV_32F, 1, 0, apertureSize, scale); + Sobel(src, Dy, CV_32F, 0, 1, apertureSize, scale); + } + else + { + Scharr(src, Dx, CV_32F, 1, 0, scale); + Scharr(src, Dy, CV_32F, 0, 1, scale); + } + + dst.create(src.size(), CV_32F); + imgproc::cornerHarris_caller(blockSize, (float)k, Dx, Dy, dst); +} + + #endif /* !defined (HAVE_CUDA) */ diff --git a/tests/gpu/src/imgproc_gpu.cpp b/tests/gpu/src/imgproc_gpu.cpp index 2c1333422b..12e1f32eab 100644 --- a/tests/gpu/src/imgproc_gpu.cpp +++ b/tests/gpu/src/imgproc_gpu.cpp @@ -603,6 +603,62 @@ void CV_GpuHistogramsTest::run( int ) } } +//////////////////////////////////////////////////////////////////////// +// Corner Harris feature detector + +struct CV_GpuCornerHarrisTest: CvTest +{ + CV_GpuCornerHarrisTest(): CvTest("GPU-CornerHarrisTest", "cornerHarris") {} + + void run(int) + { + try + { + int rows = 1 + rand() % 300, cols = 1 + rand() % 300; + if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, -1)) return; + for (int i = 0; i < 3; ++i) + { + rows = 1 + rand() % 300; cols = 1 + rand() % 300; + if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return; + } + } + catch (const Exception& e) + { + if (!check_and_treat_gpu_exception(e, ts)) throw; + return; + } + } + + bool compareToCpuTest(int rows, int cols, int depth, int blockSize, int apertureSize) + { + RNG rng; + cv::Mat src(rows, cols, depth); + if (depth == CV_32F) + rng.fill(src, RNG::UNIFORM, cv::Scalar(0), cv::Scalar(1)); + + double k = 0.1; + int borderType = BORDER_DEFAULT; + + cv::Mat dst_gold; + cv::cornerHarris(src, dst_gold, blockSize, apertureSize, k, borderType); + + cv::gpu::GpuMat dst; + cv::gpu::cornerHarris(cv::gpu::GpuMat(src), dst, blockSize, apertureSize, k); + + cv::Mat dsth = dst; + for (int i = apertureSize + 2; i < dst.rows - apertureSize - 2; ++i) + { + for (int j = apertureSize + 2; j < dst.cols - apertureSize - 2; ++j) + { + float a = dst_gold.at(i, j); + float b = dsth.at(i, j); + if (fabs(a - b) > 1e-3f) return false; + } + } + return true; + } +}; + ///////////////////////////////////////////////////////////////////////////// /////////////////// tests registration ///////////////////////////////////// ///////////////////////////////////////////////////////////////////////////// @@ -620,3 +676,5 @@ CV_GpuNppImageIntegralTest CV_GpuNppImageIntegral_test; CV_GpuNppImageCannyTest CV_GpuNppImageCanny_test; CV_GpuCvtColorTest CV_GpuCvtColor_test; CV_GpuHistogramsTest CV_GpuHistograms_test; +CV_GpuCornerHarrisTest CV_GpuCornerHarris_test; +