From 493d7a6e45b373b62c4f6fe4d49c187135591a8d Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Tue, 30 Nov 2010 08:44:04 +0000 Subject: [PATCH] added min eigen val based corner detector into gpu module --- modules/gpu/include/opencv2/gpu/gpu.hpp | 8 ++++- modules/gpu/src/cuda/imgproc.cu | 54 ++++++++++++++++++++++++++++++- modules/gpu/src/imgproc_gpu.cpp | 47 ++++++++++++++++++--------- tests/gpu/src/imgproc_gpu.cpp | 57 +++++++++++++++++++++++++++++++++ 4 files changed, 149 insertions(+), 17 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 043f163..50dfc2a 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -629,7 +629,13 @@ namespace cv //! 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); + CV_EXPORTS void cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, double k); + + + //! computes minimum eigen value of 2x2 derivative covariation matrix at each pixel - the cornerness criteria + // (does BORDER_CONSTANT interpolation with 0 as the fill value) + CV_EXPORTS void cornerMinEigenVal(const GpuMat& src, GpuMat& dst, int blockSize, int ksize=3); + //////////////////////////////// Filter Engine //////////////////////////////// diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index 1e632a2..3ae37b1 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -466,7 +466,8 @@ namespace cv { namespace gpu { namespace imgproc /////////////////////////////////////////// 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) + __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; @@ -511,4 +512,55 @@ namespace cv { namespace gpu { namespace imgproc cornerHarris_kernel<<>>(cols, rows, block_size / 2, k, Dx, Dy, dst); cudaSafeCall(cudaThreadSynchronize()); } + +/////////////////////////////////////////// Corner Min Eigen Val ///////////////////////////////////////////////// + + __global__ void cornerMinEigenVal_kernel(const int cols, const int rows, const int block_size, + 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; + } + } + + a *= 0.5f; + c *= 0.5f; + ((float*)dst.ptr(y))[x] = (a + c) - sqrtf((a - c) * (a - c) + b * b); + } + } + + void cornerMinEigenVal_caller(const int block_size, 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)); + + cornerMinEigenVal_kernel<<>>(cols, rows, block_size / 2, Dx, Dy, dst); + cudaSafeCall(cudaThreadSynchronize()); + } }}} diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index 00696d8..26a464e 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/modules/gpu/src/imgproc_gpu.cpp @@ -69,6 +69,7 @@ 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(); } +void cv::gpu::cornerMinEigenVal(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); } #else /* !defined (HAVE_CUDA) */ @@ -861,32 +862,48 @@ void cv::gpu::histRange(const GpuMat& src, GpuMat hist[4], const GpuMat levels[4 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 cornerMinEigenVal_caller(const int block_size, const DevMem2D Dx, const DevMem2D Dy, DevMem2D dst); }}} -void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int apertureSize, double k) +namespace { - CV_Assert(src.type() == CV_32F); + void computeGradients(const GpuMat& src, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize) + { + CV_Assert(src.type() == CV_32F); - double scale = (double)(1 << ((apertureSize > 0 ? apertureSize : 3) - 1)) * blockSize; - if (apertureSize < 0) scale *= 2.; - scale = 1./scale; + double scale = (double)(1 << ((ksize > 0 ? ksize : 3) - 1)) * blockSize; + if (ksize < 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); + if (ksize > 0) + { + Sobel(src, Dx, CV_32F, 1, 0, ksize, scale); + Sobel(src, Dy, CV_32F, 0, 1, ksize, scale); + } + else + { + Scharr(src, Dx, CV_32F, 1, 0, scale); + Scharr(src, Dy, CV_32F, 0, 1, scale); + } } +} +void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, double k) +{ + GpuMat Dx, Dy; + computeGradients(src, Dx, Dy, blockSize, ksize); dst.create(src.size(), CV_32F); imgproc::cornerHarris_caller(blockSize, (float)k, Dx, Dy, dst); } +void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, int blockSize, int ksize) +{ + GpuMat Dx, Dy; + computeGradients(src, Dx, Dy, blockSize, ksize); + dst.create(src.size(), CV_32F); + imgproc::cornerMinEigenVal_caller(blockSize, Dx, Dy, dst); +} + #endif /* !defined (HAVE_CUDA) */ diff --git a/tests/gpu/src/imgproc_gpu.cpp b/tests/gpu/src/imgproc_gpu.cpp index 12e1f32..051730a 100644 --- a/tests/gpu/src/imgproc_gpu.cpp +++ b/tests/gpu/src/imgproc_gpu.cpp @@ -659,6 +659,62 @@ struct CV_GpuCornerHarrisTest: CvTest } }; +//////////////////////////////////////////////////////////////////////// +// Corner Min Eigen Val + +struct CV_GpuCornerMinEigenValTest: CvTest +{ + CV_GpuCornerMinEigenValTest(): CvTest("GPU-CornerMinEigenValTest", "cornerMinEigenVal") {} + + 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::cornerMinEigenVal(src, dst_gold, blockSize, apertureSize, borderType); + + cv::gpu::GpuMat dst; + cv::gpu::cornerMinEigenVal(cv::gpu::GpuMat(src), dst, blockSize, apertureSize); + + 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 ///////////////////////////////////// ///////////////////////////////////////////////////////////////////////////// @@ -677,4 +733,5 @@ CV_GpuNppImageCannyTest CV_GpuNppImageCanny_test; CV_GpuCvtColorTest CV_GpuCvtColor_test; CV_GpuHistogramsTest CV_GpuHistograms_test; CV_GpuCornerHarrisTest CV_GpuCornerHarris_test; +CV_GpuCornerMinEigenValTest CV_GpuCornerMinEigenVal_test; -- 2.7.4