added Harris corner detector into gpu module

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

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