implemented OpenCL version of cv::preCornerDetect

modules/imgproc/perf/opencl/perf_imgproc.cpp

@@ -132,6 +132,25 @@ OCL_PERF_TEST_P(CornerHarrisFixture, CornerHarris,
     SANITY_CHECK(dst, 5e-6, ERROR_RELATIVE);
 }
 
+///////////// PreCornerDetect ////////////////////////
+
+typedef Size_MatType PreCornerDetectFixture;
+
+OCL_PERF_TEST_P(PreCornerDetectFixture, PreCornerDetect,
+            ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1)))
+{
+    const Size_MatType_t params = GetParam();
+    const Size srcSize = get<0>(params);
+    const int type = get<1>(params), borderType = BORDER_REFLECT;
+
+    UMat src(srcSize, type), dst(srcSize, CV_32FC1);
+    declare.in(src, WARMUP_RNG).out(dst);
+
+    OCL_TEST_CYCLE() cv::preCornerDetect(src, dst, 3, borderType);
+
+    SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE);
+}
+
 ///////////// Integral ////////////////////////
 
 typedef tuple<Size, MatDepth> IntegralParams;

modules/imgproc/src/corner.cpp

@@ -363,6 +363,36 @@ static bool ocl_cornerMinEigenValVecs(InputArray _src, OutputArray _dst, int blo
     return cornelKernel.run(2, globalsize, localsize, false);
 }
 
+static bool ocl_preCornerDetect( InputArray _src, OutputArray _dst, int ksize, int borderType, int depth )
+{
+    UMat Dx, Dy, D2x, D2y, Dxy;
+
+    Sobel( _src, Dx, CV_32F, 1, 0, ksize, 1, 0, borderType );
+    Sobel( _src, Dy, CV_32F, 0, 1, ksize, 1, 0, borderType );
+    Sobel( _src, D2x, CV_32F, 2, 0, ksize, 1, 0, borderType );
+    Sobel( _src, D2y, CV_32F, 0, 2, ksize, 1, 0, borderType );
+    Sobel( _src, Dxy, CV_32F, 1, 1, ksize, 1, 0, borderType );
+
+    _dst.create( _src.size(), CV_32FC1 );
+    UMat dst = _dst.getUMat();
+
+    double factor = 1 << (ksize - 1);
+    if( depth == CV_8U )
+        factor *= 255;
+    factor = 1./(factor * factor * factor);
+
+    ocl::Kernel k("preCornerDetect", ocl::imgproc::precornerdetect_oclsrc);
+    if (k.empty())
+        return false;
+
+    k.args(ocl::KernelArg::ReadOnlyNoSize(Dx), ocl::KernelArg::ReadOnlyNoSize(Dy),
+           ocl::KernelArg::ReadOnlyNoSize(D2x), ocl::KernelArg::ReadOnlyNoSize(D2y),
+           ocl::KernelArg::ReadOnlyNoSize(Dxy), ocl::KernelArg::WriteOnly(dst), (float)factor);
+
+    size_t globalsize[2] = { dst.cols, dst.rows };
+    return k.run(2, globalsize, NULL, false);
+}
+
 }
 
 void cv::cornerMinEigenVal( InputArray _src, OutputArray _dst, int blockSize, int ksize, int borderType )
@@ -381,7 +411,7 @@ void cv::cornerMinEigenVal( InputArray _src, OutputArray _dst, int blockSize, in
 void cv::cornerHarris( InputArray _src, OutputArray _dst, int blockSize, int ksize, double k, int borderType )
 {
     if (ocl::useOpenCL() && _src.dims() <= 2 && _dst.isUMat() &&
-        ocl_cornerMinEigenValVecs(_src, _dst, blockSize, ksize, k, borderType, HARRIS))
+            ocl_cornerMinEigenValVecs(_src, _dst, blockSize, ksize, k, borderType, HARRIS))
         return;
 
     Mat src = _src.getMat();
@@ -406,10 +436,15 @@ void cv::cornerEigenValsAndVecs( InputArray _src, OutputArray _dst, int blockSiz
 
 void cv::preCornerDetect( InputArray _src, OutputArray _dst, int ksize, int borderType )
 {
-    Mat Dx, Dy, D2x, D2y, Dxy, src = _src.getMat();
+    int type = _src.type();
+    CV_Assert( type == CV_8UC1 || type == CV_32FC1 );
 
-    CV_Assert( src.type() == CV_8UC1 || src.type() == CV_32FC1 );
-    _dst.create( src.size(), CV_32F );
+    if (ocl::useOpenCL() && _src.dims() <= 2 && _dst.isUMat() &&
+            ocl_preCornerDetect(_src, _dst, ksize, borderType, CV_MAT_DEPTH(type)))
+        return;
+
+    Mat Dx, Dy, D2x, D2y, Dxy, src = _src.getMat();
+    _dst.create( src.size(), CV_32FC1 );
     Mat dst = _dst.getMat();
 
     Sobel( src, Dx, CV_32F, 1, 0, ksize, 1, 0, borderType );

modules/imgproc/src/opencl/precornerdetect.cl

@@ -0,0 +1,75 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Shengen Yan,yanshengen@gmail.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+__kernel void preCornerDetect(__global const uchar * Dxptr, int dx_step, int dx_offset,
+                              __global const uchar * Dyptr, int dy_step, int dy_offset,
+                              __global const uchar * D2xptr, int d2x_step, int d2x_offset,
+                              __global const uchar * D2yptr, int d2y_step, int d2y_offset,
+                              __global const uchar * Dxyptr, int dxy_step, int dxy_offset,
+                              __global uchar * dstptr, int dst_step, int dst_offset,
+                              int dst_rows, int dst_cols, float factor)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < dst_cols && y < dst_rows)
+    {
+        int dx_index = mad24(dx_step, y, (int)sizeof(float) * x + dx_offset);
+        int dy_index = mad24(dy_step, y, (int)sizeof(float) * x + dy_offset);
+        int d2x_index = mad24(d2x_step, y, (int)sizeof(float) * x + d2x_offset);
+        int d2y_index = mad24(d2y_step, y, (int)sizeof(float) * x + d2y_offset);
+        int dxy_index = mad24(dxy_step, y, (int)sizeof(float) * x + dxy_offset);
+        int dst_index = mad24(dst_step, y, (int)sizeof(float) * x + dst_offset);
+
+        float dx = *(__global const float *)(Dxptr + dx_index);
+        float dy = *(__global const float *)(Dyptr + dy_index);
+        float d2x = *(__global const float *)(D2xptr + d2x_index);
+        float d2y = *(__global const float *)(D2yptr + d2y_index);
+        float dxy = *(__global const float *)(Dxyptr + dxy_index);
+        __global float * dst = (__global float *)(dstptr + dst_index);
+
+        dst[0] = factor * (dx*dx*d2y + dy*dy*d2x - 2*dx*dy*dxy);
+    }
+}

modules/imgproc/test/ocl/test_imgproc.cpp

@@ -259,6 +259,26 @@ OCL_TEST_P(CornerHarris, Mat)
     }
 }
 
+//////////////////////////////// preCornerDetect //////////////////////////////////////////
+
+typedef ImgprocTestBase PreCornerDetect;
+
+OCL_TEST_P(PreCornerDetect, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        random_roi();
+
+        const int apertureSize = blockSize;
+
+        OCL_OFF(cv::preCornerDetect(src_roi, dst_roi, apertureSize, borderType));
+        OCL_ON(cv::preCornerDetect(usrc_roi, udst_roi, apertureSize, borderType));
+
+        Near(1e-6, true);
+    }
+}
+
+
 ////////////////////////////////// integral /////////////////////////////////////////////////
 
 struct Integral :
@@ -341,7 +361,6 @@ struct Threshold :
     virtual void SetUp()
     {
         type = GET_PARAM(0);
-        blockSize = GET_PARAM(1);
         thresholdType = GET_PARAM(2);
         useRoi = GET_PARAM(3);
     }
@@ -437,6 +456,13 @@ OCL_INSTANTIATE_TEST_CASE_P(Imgproc, CornerHarris, Combine(
                                     (BorderType)BORDER_REFLECT, (BorderType)BORDER_REFLECT_101),
                             Bool()));
 
+OCL_INSTANTIATE_TEST_CASE_P(Imgproc, PreCornerDetect, Combine(
+                            Values((MatType)CV_8UC1, CV_32FC1),
+                            Values(3, 5),
+                            Values( (BorderType)BORDER_CONSTANT, (BorderType)BORDER_REPLICATE,
+                                    (BorderType)BORDER_REFLECT, (BorderType)BORDER_REFLECT_101),
+                            Bool()));
+
 OCL_INSTANTIATE_TEST_CASE_P(Imgproc, Integral, Combine(
                             Values((MatType)CV_8UC1), // TODO does not work with CV_32F, CV_64F
                             Values(CV_32SC1, CV_32FC1), // desired sdepth