added min eigen val based corner detector into gpu module

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

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<<<grid, threads>>>(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<<<grid, threads>>>(cols, rows, block_size / 2, Dx, Dy, dst);
+        cudaSafeCall(cudaThreadSynchronize());
+    }
 }}}

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;
+        double scale = (double)(1 << ((ksize > 0 ? ksize : 3) - 1)) * blockSize;
-    if (apertureSize < 0) scale *= 2.;
+        if (ksize < 0) scale *= 2.;
-    scale = 1./scale;
+        scale = 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;
-    if (apertureSize > 0)
+    computeGradients(src, Dx, Dy, blockSize, ksize);
-    {
-        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);
 }
 
+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) */

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<float>(i, j);
+                float b = dsth.at<float>(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;