diff --git a/modules/ocl/CMakeLists.txt b/modules/ocl/CMakeLists.txt index d0e254d3a2..a9ec2f4a0a 100644 --- a/modules/ocl/CMakeLists.txt +++ b/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() diff --git a/modules/ocl/src/kernels/nonfree_surf.cl b/modules/ocl/src/kernels/nonfree_surf.cl index 69f64795e9..8cffe3d93a 100644 --- a/modules/ocl/src/kernels/nonfree_surf.cl +++ b/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' */ diff --git a/modules/ocl/src/surf.cpp b/modules/ocl/src/surf.cpp index e2ac21b972..9d1372bbe0 100644 --- a/modules/ocl/src/surf.cpp +++ b/modules/ocl/src/surf.cpp @@ -160,7 +160,7 @@ public: if (use_mask) { - throw std::exception(); + CV_Error(CV_StsBadFunc, "Masked SURF detector is not implemented yet"); //!FIXME // temp fix for missing min overload //oclMat temp(mask.size(), mask.type()); @@ -623,7 +623,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(nFeatures), 1, 1}; openCLExecuteKernelSURF(clCxt, &nonfree_surf, kernelName, globalThreads, localThreads, args, -1, -1); } @@ -725,4 +725,3 @@ void SURF_OCL_Invoker::compute_descriptors_gpu(const oclMat &descriptors, const openCLExecuteKernelSURF(clCxt, &nonfree_surf, kernelName, globalThreads, localThreads, args, -1, -1); } } - diff --git a/modules/ocl/test/precomp.hpp b/modules/ocl/test/precomp.hpp index c630871014..e8c1aaa1b9 100644 --- a/modules/ocl/test/precomp.hpp +++ b/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" diff --git a/modules/ocl/test/test_surf.cpp b/modules/ocl/test/test_surf.cpp new file mode 100644 index 0000000000..c4cf60fcbc --- /dev/null +++ b/modules/ocl/test/test_surf.cpp @@ -0,0 +1,227 @@ +/*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 +{ + 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& gold, std::vector& 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& keypoints1, const std::vector& keypoints2, const std::vector& 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(hessianThreshold); + surf.nOctaves = nOctaves; + surf.nOctaveLayers = nOctaveLayers; + surf.extended = extended; + surf.upright = upright; + surf.keypointsRatio = 0.05f; + + std::vector 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 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(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(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 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 matches; + matcher.match(descriptors_gold, cv::Mat(descriptors), matches); + + int matchedCount = getMatchedPointsCount(keypoints, keypoints, matches); + double matchedRatio = static_cast(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