Add ocl::SURF accurate test.

The test cases are ported from gpu module (masked version is still
missing).
Key point direction definition is synchronized with cv::SURF.
nonfree module dependency is added into ocl module's cmake file.

modules/ocl/CMakeLists.txt

@@ -4,7 +4,7 @@ if(NOT HAVE_OPENCL)
 endif()
 
 set(the_description "OpenCL-accelerated Computer Vision")
-ocv_add_module(ocl opencv_core opencv_imgproc opencv_features2d opencv_objdetect opencv_video)
+ocv_add_module(ocl opencv_core opencv_imgproc opencv_features2d opencv_objdetect opencv_video opencv_nonfree)
 
 ocv_module_include_directories()
 

modules/ocl/src/kernels/nonfree_surf.cl

@@ -78,7 +78,12 @@ uchar read_imgTex(IMAGE_INT8 img, sampler_t sam, float2 coord, int rows, int col
 
 // dynamically change the precision used for floating type
 
-#if defined DOUBLE_SUPPORT
+#if defined (DOUBLE_SUPPORT)
+#ifdef cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#elif defined (cl_amd_fp64)
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#endif
 #define F double
 #else
 #define F float
@@ -892,9 +897,9 @@ __kernel
             kp_dir += 2.0f * CV_PI_F;
         kp_dir *= 180.0f / CV_PI_F;
 
-        //kp_dir = 360.0f - kp_dir;
-        //if (fabs(kp_dir - 360.f) < FLT_EPSILON)
-        //    kp_dir = 0.f;
+        kp_dir = 360.0f - kp_dir;
+        if (fabs(kp_dir - 360.f) < FLT_EPSILON)
+            kp_dir = 0.f;
 
         featureDir[get_group_id(0)] = kp_dir;
     }
@@ -913,7 +918,7 @@ __kernel
 
     if(get_global_id(0) <= nFeatures)
     {
-        featureDir[get_global_id(0)] = 90.0f;
+        featureDir[get_global_id(0)] = 270.0f;
     }
 }
 
@@ -1011,7 +1016,12 @@ void calc_dx_dy(
     const float centerX = featureX[get_group_id(0)];
     const float centerY = featureY[get_group_id(0)];
     const float size = featureSize[get_group_id(0)];
-    float descriptor_dir = featureDir[get_group_id(0)] * (float)(CV_PI_F / 180.0f);
+    float descriptor_dir = 360.0f - featureDir[get_group_id(0)];
+    if(fabs(descriptor_dir - 360.0f) < FLT_EPSILON)
+    {
+        descriptor_dir = 0.0f;
+    }
+    descriptor_dir *= (float)(CV_PI_F / 180.0f);
 
     /* The sampling intervals and wavelet sized for selecting an orientation
     and building the keypoint descriptor are defined relative to 's' */

modules/ocl/src/surf.cpp

@@ -51,6 +51,57 @@ using namespace cv;
 using namespace cv::ocl;
 using namespace std;
 
+#if !defined (HAVE_OPENCL)
+cv::ocl::SURF_OCL::SURF_OCL()
+{
+    throw_nogpu();
+}
+cv::ocl::SURF_OCL::SURF_OCL(double, int, int, bool, float, bool)
+{
+    throw_nogpu();
+}
+int cv::ocl::SURF_OCL::descriptorSize() const
+{
+    throw_nogpu();
+    return 0;
+}
+void cv::ocl::SURF_OCL::uploadKeypoints(const vector<KeyPoint> &, oclMat &)
+{
+    throw_nogpu();
+}
+void cv::ocl::SURF_OCL::downloadKeypoints(const oclMat &, vector<KeyPoint> &)
+{
+    throw_nogpu();
+}
+void cv::ocl::SURF_OCL::downloadDescriptors(const oclMat &, vector<float> &)
+{
+    throw_nogpu();
+}
+void cv::ocl::SURF_OCL::operator()(const oclMat &, const oclMat &, oclMat &)
+{
+    throw_nogpu();
+}
+void cv::ocl::SURF_OCL::operator()(const oclMat &, const oclMat &, oclMat &, oclMat &, bool)
+{
+    throw_nogpu();
+}
+void cv::ocl::SURF_OCL::operator()(const oclMat &, const oclMat &, vector<KeyPoint> &)
+{
+    throw_nogpu();
+}
+void cv::ocl::SURF_OCL::operator()(const oclMat &, const oclMat &, vector<KeyPoint> &, oclMat &, bool)
+{
+    throw_nogpu();
+}
+void cv::ocl::SURF_OCL::operator()(const oclMat &, const oclMat &, vector<KeyPoint> &, vector<float> &, bool)
+{
+    throw_nogpu();
+}
+void cv::ocl::SURF_OCL::releaseMemory()
+{
+    throw_nogpu();
+}
+#else /* !defined (HAVE_OPENCL) */
 namespace cv
 {
     namespace ocl
@@ -623,7 +674,7 @@ void SURF_OCL_Invoker::icvSetUpright_gpu(const oclMat &keypoints, int nFeatures)
     args.push_back( make_pair( sizeof(cl_int), (void *)&nFeatures));
 
     size_t localThreads[3]  = {256, 1, 1};
-    size_t globalThreads[3] = {nFeatures, 1, 1};
+    size_t globalThreads[3] = {saturate_cast<size_t>(nFeatures), 1, 1};
 
     openCLExecuteKernelSURF(clCxt, &nonfree_surf, kernelName, globalThreads, localThreads, args, -1, -1);
 }
@@ -726,3 +777,4 @@ void SURF_OCL_Invoker::compute_descriptors_gpu(const oclMat &descriptors, const
     }
 }
 
+#endif // /* !defined (HAVE_OPENCL) */

modules/ocl/test/precomp.hpp

@@ -70,7 +70,7 @@
 #include "opencv2/ts/ts.hpp"
 #include "opencv2/ts/ts_perf.hpp"
 #include "opencv2/ocl/ocl.hpp"
-//#include "opencv2/nonfree/nonfree.hpp"
+#include "opencv2/nonfree/nonfree.hpp"
 
 #include "utility.hpp"
 #include "interpolation.hpp"

modules/ocl/test/test_surf.cpp

+/*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, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Peng Xiao, pengxiao@multicorewareinc.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 oclMaterials 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*/
+
+
+#include "precomp.hpp"
+#ifdef HAVE_OPENCL
+
+extern std::string workdir;
+
+using namespace std;
+
+static bool keyPointsEquals(const cv::KeyPoint& p1, const cv::KeyPoint& p2)
+{
+    const double maxPtDif = 1.0;
+    const double maxSizeDif = 1.0;
+    const double maxAngleDif = 2.0;
+    const double maxResponseDif = 0.1;
+
+    double dist = cv::norm(p1.pt - p2.pt);
+
+    if (dist < maxPtDif &&
+        fabs(p1.size - p2.size) < maxSizeDif &&
+        abs(p1.angle - p2.angle) < maxAngleDif &&
+        abs(p1.response - p2.response) < maxResponseDif &&
+        p1.octave == p2.octave &&
+        p1.class_id == p2.class_id)
+    {
+        return true;
+    }
+
+    return false;
+}
+
+
+struct KeyPointLess : std::binary_function<cv::KeyPoint, cv::KeyPoint, bool>
+{
+    bool operator()(const cv::KeyPoint& kp1, const cv::KeyPoint& kp2) const
+    {
+        return kp1.pt.y < kp2.pt.y || (kp1.pt.y == kp2.pt.y && kp1.pt.x < kp2.pt.x);
+    }
+};
+
+
+#define ASSERT_KEYPOINTS_EQ(gold, actual) EXPECT_PRED_FORMAT2(assertKeyPointsEquals, gold, actual);
+
+static int getMatchedPointsCount(std::vector<cv::KeyPoint>& gold, std::vector<cv::KeyPoint>& actual)
+{
+    std::sort(actual.begin(), actual.end(), KeyPointLess());
+    std::sort(gold.begin(), gold.end(), KeyPointLess());
+
+    int validCount = 0;
+
+    for (size_t i = 0; i < gold.size(); ++i)
+    {
+        const cv::KeyPoint& p1 = gold[i];
+        const cv::KeyPoint& p2 = actual[i];
+
+        if (keyPointsEquals(p1, p2))
+            ++validCount;
+    }
+
+    return validCount;
+}
+
+static int getMatchedPointsCount(const std::vector<cv::KeyPoint>& keypoints1, const std::vector<cv::KeyPoint>& keypoints2, const std::vector<cv::DMatch>& matches)
+{
+    int validCount = 0;
+
+    for (size_t i = 0; i < matches.size(); ++i)
+    {
+        const cv::DMatch& m = matches[i];
+
+        const cv::KeyPoint& p1 = keypoints1[m.queryIdx];
+        const cv::KeyPoint& p2 = keypoints2[m.trainIdx];
+
+        if (keyPointsEquals(p1, p2))
+            ++validCount;
+    }
+
+    return validCount;
+}
+
+IMPLEMENT_PARAM_CLASS(SURF_HessianThreshold, double)
+IMPLEMENT_PARAM_CLASS(SURF_Octaves, int)
+IMPLEMENT_PARAM_CLASS(SURF_OctaveLayers, int)
+IMPLEMENT_PARAM_CLASS(SURF_Extended, bool)
+IMPLEMENT_PARAM_CLASS(SURF_Upright, bool)
+
+PARAM_TEST_CASE(SURF, SURF_HessianThreshold, SURF_Octaves, SURF_OctaveLayers, SURF_Extended, SURF_Upright)
+{
+    double hessianThreshold;
+    int nOctaves;
+    int nOctaveLayers;
+    bool extended;
+    bool upright;
+
+    virtual void SetUp()
+    {
+        hessianThreshold = GET_PARAM(0);
+        nOctaves = GET_PARAM(1);
+        nOctaveLayers = GET_PARAM(2);
+        extended = GET_PARAM(3);
+        upright = GET_PARAM(4);
+    }
+};
+TEST_P(SURF, Detector)
+{
+    cv::Mat image = readImage(workdir + "fruits.jpg", cv::IMREAD_GRAYSCALE);
+    ASSERT_FALSE(image.empty());
+
+    cv::ocl::SURF_OCL surf;
+    surf.hessianThreshold = static_cast<float>(hessianThreshold);
+    surf.nOctaves = nOctaves;
+    surf.nOctaveLayers = nOctaveLayers;
+    surf.extended = extended;
+    surf.upright = upright;
+    surf.keypointsRatio = 0.05f;
+
+    std::vector<cv::KeyPoint> keypoints;
+    surf(cv::ocl::oclMat(image), cv::ocl::oclMat(), keypoints);
+
+    cv::SURF surf_gold;
+    surf_gold.hessianThreshold = hessianThreshold;
+    surf_gold.nOctaves = nOctaves;
+    surf_gold.nOctaveLayers = nOctaveLayers;
+    surf_gold.extended = extended;
+    surf_gold.upright = upright;
+
+    std::vector<cv::KeyPoint> keypoints_gold;
+    surf_gold(image, cv::noArray(), keypoints_gold);
+
+    ASSERT_EQ(keypoints_gold.size(), keypoints.size());
+    int matchedCount = getMatchedPointsCount(keypoints_gold, keypoints);
+    double matchedRatio = static_cast<double>(matchedCount) / keypoints_gold.size();
+
+    EXPECT_GT(matchedRatio, 0.95);
+}
+
+TEST_P(SURF, Descriptor)
+{
+    cv::Mat image = readImage(workdir + "fruits.jpg", cv::IMREAD_GRAYSCALE);
+    ASSERT_FALSE(image.empty());
+
+    cv::ocl::SURF_OCL surf;
+    surf.hessianThreshold = static_cast<float>(hessianThreshold);
+    surf.nOctaves = nOctaves;
+    surf.nOctaveLayers = nOctaveLayers;
+    surf.extended = extended;
+    surf.upright = upright;
+    surf.keypointsRatio = 0.05f;
+
+    cv::SURF surf_gold;
+    surf_gold.hessianThreshold = hessianThreshold;
+    surf_gold.nOctaves = nOctaves;
+    surf_gold.nOctaveLayers = nOctaveLayers;
+    surf_gold.extended = extended;
+    surf_gold.upright = upright;
+
+    std::vector<cv::KeyPoint> keypoints;
+    surf_gold(image, cv::noArray(), keypoints);
+
+    cv::ocl::oclMat descriptors;
+    surf(cv::ocl::oclMat(image), cv::ocl::oclMat(), keypoints, descriptors, true);
+
+    cv::Mat descriptors_gold;
+    surf_gold(image, cv::noArray(), keypoints, descriptors_gold, true);
+
+    cv::BFMatcher matcher(cv::NORM_L2);
+    std::vector<cv::DMatch> matches;
+    matcher.match(descriptors_gold, cv::Mat(descriptors), matches);
+
+    int matchedCount = getMatchedPointsCount(keypoints, keypoints, matches);
+    double matchedRatio = static_cast<double>(matchedCount) / keypoints.size();
+
+    EXPECT_GT(matchedRatio, 0.35);
+}
+
+INSTANTIATE_TEST_CASE_P(OCL_Features2D, SURF, testing::Combine(
+    testing::Values(/*SURF_HessianThreshold(100.0), */SURF_HessianThreshold(500.0), SURF_HessianThreshold(1000.0)),
+    testing::Values(SURF_Octaves(3), SURF_Octaves(4)),
+    testing::Values(SURF_OctaveLayers(2), SURF_OctaveLayers(3)),
+    testing::Values(SURF_Extended(false), SURF_Extended(true)),
+    testing::Values(SURF_Upright(false), SURF_Upright(true))));
+
+#endif