fixed build in gpu module (SURF and ORB)

diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt
index 199b520..a403b68 100644 (file)
--- a/modules/gpu/CMakeLists.txt
+++ b/modules/gpu/CMakeLists.txt
@@ -3,7 +3,7 @@ if(ANDROID OR IOS)
 endif()
 
 set(the_description "GPU-accelerated Computer Vision")
-ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video)
+ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video opencv_nonfree)
 
 ocv_module_include_directories("${CMAKE_CURRENT_SOURCE_DIR}/src/cuda")
 

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 6ecea83..bc53975 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -1566,13 +1566,12 @@ public:
 \r
     void releaseMemory();\r
 \r
-    // SURF parameters;\r
-    int    extended;\r
-    int    upright;\r
+    // SURF parameters\r
     double hessianThreshold;\r
-    \r
-    int    nOctaves;\r
-    int    nOctaveLayers;\r
+    int nOctaves;\r
+    int nOctaveLayers;\r
+    bool extended;\r
+    bool upright;\r
     \r
     //! max keypoints = min(keypointsRatio * img.size().area(), 65535)\r
     float keypointsRatio;\r
@@ -1662,9 +1661,8 @@ public:
     };\r
 \r
     //! Constructor\r
-    //! n_features - the number of desired features\r
-    //! detector_params - parameters to use\r
-    explicit ORB_GPU(size_t n_features = 500, float scaleFactor = 1.2f, int nlevels = 3);\r
+    explicit ORB_GPU(int nFeatures = 500, float scaleFactor = 1.2f, int nLevels = 3, int edgeThreshold = 31,\r
+                     int firstLevel = 0, int WTA_K = 2, int scoreType = 0, int patchSize = 31);\r
 \r
     //! Compute the ORB features on an image\r
     //! image - the image to compute the features (supports only CV_8UC1 images)\r
@@ -1690,7 +1688,6 @@ public:
     //! returns the descriptor size in bytes\r
     inline int descriptorSize() const { return kBytes; }\r
 \r
-    void setParams(size_t n_features, float scaleFactor, int nlevels);\r
     inline void setFastParams(int threshold, bool nonmaxSupression = true)\r
     {\r
         fastDetector_.threshold = threshold;\r
@@ -1714,7 +1711,14 @@ private:
 \r
     void mergeKeyPoints(GpuMat& keypoints);\r
 \r
-    ORB params_;\r
+    int nFeatures_;\r
+    float scaleFactor_;\r
+    int nLevels_;\r
+    int edgeThreshold_;\r
+    int firstLevel_;\r
+    int WTA_K_;\r
+    int scoreType_;\r
+    int patchSize_;\r
 \r
     // The number of desired features per scale\r
     std::vector<size_t> n_features_per_level_;\r

diff --git a/modules/gpu/src/orb.cpp b/modules/gpu/src/orb.cpp
index 9972fce..fb5be1a 100644 (file)
--- a/modules/gpu/src/orb.cpp
+++ b/modules/gpu/src/orb.cpp
@@ -48,14 +48,13 @@ using namespace cv::gpu;
 \r
 #if !defined (HAVE_CUDA)\r
 \r
-cv::gpu::ORB_GPU::ORB_GPU(size_t, float, int) : fastDetector_(0) { throw_nogpu(); }\r
+cv::gpu::ORB_GPU::ORB_GPU(int, float, int, int, int, int, int, int) { throw_nogpu(); }\r
 void cv::gpu::ORB_GPU::operator()(const GpuMat&, const GpuMat&, std::vector<KeyPoint>&) { throw_nogpu(); }\r
 void cv::gpu::ORB_GPU::operator()(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
 void cv::gpu::ORB_GPU::operator()(const GpuMat&, const GpuMat&, std::vector<KeyPoint>&, GpuMat&) { throw_nogpu(); }\r
 void cv::gpu::ORB_GPU::operator()(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&) { throw_nogpu(); }\r
 void cv::gpu::ORB_GPU::downloadKeyPoints(GpuMat&, std::vector<KeyPoint>&) { throw_nogpu(); }\r
 void cv::gpu::ORB_GPU::convertKeyPoints(Mat&, std::vector<KeyPoint>&) { throw_nogpu(); }\r
-void cv::gpu::ORB_GPU::setParams(size_t, float, int) { throw_nogpu(); }\r
 void cv::gpu::ORB_GPU::release() { throw_nogpu(); }\r
 void cv::gpu::ORB_GPU::buildScalePyramids(const GpuMat&, const GpuMat&) { throw_nogpu(); }\r
 void cv::gpu::ORB_GPU::computeKeyPointsPyramid() { throw_nogpu(); }\r
@@ -83,16 +82,6 @@ namespace cv { namespace gpu { namespace device
     }\r
 }}}\r
 \r
-cv::gpu::ORB_GPU::ORB_GPU(size_t n_features, float scaleFactor, int nlevels) :\r
-    fastDetector_(DEFAULT_FAST_THRESHOLD)\r
-{\r
-    setParams(n_features, scaleFactor, nlevels);\r
-    \r
-    blurFilter = createGaussianFilter_GPU(CV_8UC1, Size(7, 7), 2, 2, BORDER_REFLECT_101);\r
-\r
-    blurForDescriptor = false;\r
-}\r
-\r
 namespace\r
 {\r
     const float HARRIS_K = 0.04f;\r
@@ -407,27 +396,27 @@ namespace
     }\r
 }\r
 \r
-void cv::gpu::ORB_GPU::setParams(size_t n_features, float scaleFactor, int nlevels)\r
-{\r
-    params_ = ORB((int)n_features, scaleFactor, nlevels);\r
-    \r
+cv::gpu::ORB_GPU::ORB_GPU(int nFeatures, float scaleFactor, int nLevels, int edgeThreshold, int firstLevel, int WTA_K, int scoreType, int patchSize) :\r
+    nFeatures_(nFeatures), scaleFactor_(scaleFactor), nLevels_(nLevels), edgeThreshold_(edgeThreshold), firstLevel_(firstLevel), WTA_K_(WTA_K),
+    scoreType_(scoreType), patchSize_(patchSize),\r
+    fastDetector_(DEFAULT_FAST_THRESHOLD)\r
+{    \r
     // fill the extractors and descriptors for the corresponding scales\r
-    int n_levels = static_cast<int>(params_.n_levels_);\r
-    float factor = 1.0f / params_.scale_factor_;\r
-    float n_desired_features_per_scale = n_features * (1.0f - factor) / (1.0f - std::pow(factor, n_levels));\r
+    float factor = 1.0f / scaleFactor_;\r
+    float n_desired_features_per_scale = nFeatures_ * (1.0f - factor) / (1.0f - std::pow(factor, nLevels_));\r
     \r
-    n_features_per_level_.resize(n_levels);\r
+    n_features_per_level_.resize(nLevels_);\r
     int sum_n_features = 0;\r
-    for (int level = 0; level < n_levels - 1; ++level)\r
+    for (int level = 0; level < nLevels_ - 1; ++level)\r
     {\r
         n_features_per_level_[level] = cvRound(n_desired_features_per_scale);\r
         sum_n_features += n_features_per_level_[level];\r
         n_desired_features_per_scale *= factor;\r
     }\r
-    n_features_per_level_[n_levels - 1] = n_features - sum_n_features;\r
+    n_features_per_level_[nLevels_ - 1] = nFeatures - sum_n_features;\r
 \r
     // pre-compute the end of a row in a circular patch\r
-    int half_patch_size = params_.patch_size_ / 2;\r
+    int half_patch_size = patchSize_ / 2;\r
     vector<int> u_max(half_patch_size + 1);\r
     for (int v = 0; v <= half_patch_size * std::sqrt(2.f) / 2 + 1; ++v)\r
         u_max[v] = cvRound(std::sqrt(static_cast<float>(half_patch_size * half_patch_size - v * v)));\r
@@ -447,17 +436,17 @@ void cv::gpu::ORB_GPU::setParams(size_t n_features, float scaleFactor, int nleve
     const int npoints = 512;\r
     Point pattern_buf[npoints];\r
     const Point* pattern0 = (const Point*)bit_pattern_31_;\r
-    if (params_.patch_size_ != 31)\r
+    if (patchSize_ != 31)\r
     {\r
         pattern0 = pattern_buf;\r
-        makeRandomPattern(params_.patch_size_, pattern_buf, npoints);\r
+        makeRandomPattern(patchSize_, pattern_buf, npoints);\r
     }\r
     \r
-    CV_Assert(params_.WTA_K_ == 2 || params_.WTA_K_ == 3 || params_.WTA_K_ == 4);    \r
+    CV_Assert(WTA_K_ == 2 || WTA_K_ == 3 || WTA_K_ == 4);    \r
 \r
     Mat h_pattern;\r
 \r
-    if (params_.WTA_K_ == 2)\r
+    if (WTA_K_ == 2)\r
     {\r
         h_pattern.create(2, npoints, CV_32SC1);\r
         \r
@@ -473,17 +462,21 @@ void cv::gpu::ORB_GPU::setParams(size_t n_features, float scaleFactor, int nleve
     else\r
     {\r
         int ntuples = descriptorSize() * 4;\r
-        initializeOrbPattern(pattern0, h_pattern, ntuples, params_.WTA_K_, npoints);\r
+        initializeOrbPattern(pattern0, h_pattern, ntuples, WTA_K_, npoints);\r
     }\r
 \r
     pattern_.upload(h_pattern);\r
+    \r
+    blurFilter = createGaussianFilter_GPU(CV_8UC1, Size(7, 7), 2, 2, BORDER_REFLECT_101);\r
+\r
+    blurForDescriptor = false;\r
 }\r
 \r
 namespace\r
 {\r
-    inline float getScale(const ORB::CommonParams& params, int level)\r
+    inline float getScale(float scaleFactor, int firstLevel, int level)\r
     {\r
-        return pow(params.scale_factor_, level - static_cast<int>(params.first_level_));\r
+        return pow(scaleFactor, level - firstLevel);\r
     }\r
 }\r
 \r
@@ -492,12 +485,12 @@ void cv::gpu::ORB_GPU::buildScalePyramids(const GpuMat& image, const GpuMat& mas
     CV_Assert(image.type() == CV_8UC1);\r
     CV_Assert(mask.empty() || (mask.type() == CV_8UC1 && mask.size() == image.size()));\r
 \r
-    imagePyr_.resize(params_.n_levels_);\r
-    maskPyr_.resize(params_.n_levels_);\r
+    imagePyr_.resize(nLevels_);\r
+    maskPyr_.resize(nLevels_);\r
 \r
-    for (int level = 0; level < static_cast<int>(params_.n_levels_); ++level)\r
+    for (int level = 0; level < nLevels_; ++level)\r
     {\r
-        float scale = 1.0f / getScale(params_, level);\r
+        float scale = 1.0f / getScale(scaleFactor_, firstLevel_, level);\r
 \r
         Size sz(cvRound(image.cols * scale), cvRound(image.rows * scale));\r
 \r
@@ -506,9 +499,9 @@ void cv::gpu::ORB_GPU::buildScalePyramids(const GpuMat& image, const GpuMat& mas
         maskPyr_[level].setTo(Scalar::all(255));\r
         \r
         // Compute the resized image\r
-        if (level != static_cast<int>(params_.first_level_))\r
+        if (level != firstLevel_)\r
         {\r
-            if (level < static_cast<int>(params_.first_level_))\r
+            if (level < firstLevel_)\r
             {\r
                 resize(image, imagePyr_[level], sz, 0, 0, INTER_LINEAR);\r
 \r
@@ -534,7 +527,7 @@ void cv::gpu::ORB_GPU::buildScalePyramids(const GpuMat& image, const GpuMat& mas
         // Filter keypoints by image border\r
         ensureSizeIsEnough(sz, CV_8UC1, buf_);\r
         buf_.setTo(Scalar::all(0));\r
-        Rect inner(params_.edge_threshold_, params_.edge_threshold_, sz.width - 2 * params_.edge_threshold_, sz.height - 2 * params_.edge_threshold_);\r
+        Rect inner(edgeThreshold_, edgeThreshold_, sz.width - 2 * edgeThreshold_, sz.height - 2 * edgeThreshold_);\r
         buf_(inner).setTo(Scalar::all(255));\r
 \r
         bitwise_and(maskPyr_[level], buf_, maskPyr_[level]);\r
@@ -566,12 +559,12 @@ void cv::gpu::ORB_GPU::computeKeyPointsPyramid()
 {\r
     using namespace cv::gpu::device::orb;\r
 \r
-    int half_patch_size = params_.patch_size_ / 2;\r
+    int half_patch_size = patchSize_ / 2;\r
 \r
-    keyPointsPyr_.resize(params_.n_levels_);\r
-    keyPointsCount_.resize(params_.n_levels_);\r
+    keyPointsPyr_.resize(nLevels_);\r
+    keyPointsCount_.resize(nLevels_);\r
     \r
-    for (int level = 0; level < static_cast<int>(params_.n_levels_); ++level)\r
+    for (int level = 0; level < nLevels_; ++level)\r
     {\r
         keyPointsCount_[level] = fastDetector_.calcKeyPointsLocation(imagePyr_[level], maskPyr_[level]);\r
 \r
@@ -582,7 +575,7 @@ void cv::gpu::ORB_GPU::computeKeyPointsPyramid()
 \r
         int n_features = n_features_per_level_[level];\r
         \r
-        if (params_.score_type_ == ORB::CommonParams::HARRIS_SCORE)\r
+        if (scoreType_ == ORB::HARRIS_SCORE)\r
         {\r
             // Keep more points than necessary as FAST does not give amazing corners\r
             cull(keyPointsPyr_[level], keyPointsCount_[level], 2 * n_features);\r
@@ -605,7 +598,7 @@ void cv::gpu::ORB_GPU::computeDescriptors(GpuMat& descriptors)
 \r
     int nAllkeypoints = 0;\r
 \r
-    for (size_t level = 0; level < params_.n_levels_; ++level)\r
+    for (int level = 0; level < nLevels_; ++level)\r
         nAllkeypoints += keyPointsCount_[level];\r
 \r
     if (nAllkeypoints == 0)\r
@@ -618,7 +611,7 @@ void cv::gpu::ORB_GPU::computeDescriptors(GpuMat& descriptors)
 \r
     int offset = 0;\r
 \r
-    for (size_t level = 0; level < params_.n_levels_; ++level)\r
+    for (int level = 0; level < nLevels_; ++level)\r
     {       \r
         GpuMat descRange = descriptors.rowRange(offset, offset + keyPointsCount_[level]);\r
 \r
@@ -630,7 +623,7 @@ void cv::gpu::ORB_GPU::computeDescriptors(GpuMat& descriptors)
         }\r
 \r
         computeOrbDescriptor_gpu(blurForDescriptor ? buf_ : imagePyr_[level], keyPointsPyr_[level].ptr<short2>(0), keyPointsPyr_[level].ptr<float>(2), \r
-            keyPointsCount_[level], pattern_.ptr<int>(0), pattern_.ptr<int>(1), descRange, descriptorSize(), params_.WTA_K_, 0);\r
+            keyPointsCount_[level], pattern_.ptr<int>(0), pattern_.ptr<int>(1), descRange, descriptorSize(), WTA_K_, 0);\r
 \r
         offset += keyPointsCount_[level];\r
     }\r
@@ -642,7 +635,7 @@ void cv::gpu::ORB_GPU::mergeKeyPoints(GpuMat& keypoints)
 \r
     int nAllkeypoints = 0;\r
 \r
-    for (size_t level = 0; level < params_.n_levels_; ++level)\r
+    for (int level = 0; level < nLevels_; ++level)\r
         nAllkeypoints += keyPointsCount_[level];\r
 \r
     if (nAllkeypoints == 0)\r
@@ -655,13 +648,13 @@ void cv::gpu::ORB_GPU::mergeKeyPoints(GpuMat& keypoints)
 \r
     int offset = 0;\r
     \r
-    for (int level = 0; level < static_cast<int>(params_.n_levels_); ++level)\r
+    for (int level = 0; level < nLevels_; ++level)\r
     {\r
-        float sf = getScale(params_, level);\r
+        float sf = getScale(scaleFactor_, firstLevel_, level);\r
 \r
         GpuMat keyPointsRange = keypoints.colRange(offset, offset + keyPointsCount_[level]);        \r
         \r
-        float locScale = level != static_cast<int>(params_.first_level_) ? sf : 1.0f;\r
+        float locScale = level != firstLevel_ ? sf : 1.0f;\r
 \r
         mergeLocation_gpu(keyPointsPyr_[level].ptr<short2>(0), keyPointsRange.ptr<float>(0), keyPointsRange.ptr<float>(1), keyPointsCount_[level], locScale, 0);\r
 \r
@@ -669,7 +662,7 @@ void cv::gpu::ORB_GPU::mergeKeyPoints(GpuMat& keypoints)
         keyPointsPyr_[level](Range(1, 3), Range(0, keyPointsCount_[level])).copyTo(range);\r
         \r
         keyPointsRange.row(4).setTo(Scalar::all(level));\r
-        keyPointsRange.row(5).setTo(Scalar::all(params_.patch_size_ * sf));\r
+        keyPointsRange.row(5).setTo(Scalar::all(patchSize_ * sf));\r
 \r
         offset += keyPointsCount_[level];\r
     }\r

diff --git a/modules/gpu/src/surf.cpp b/modules/gpu/src/surf.cpp
index a083200..de592f5 100644 (file)
--- a/modules/gpu/src/surf.cpp
+++ b/modules/gpu/src/surf.cpp
@@ -109,32 +109,26 @@ namespace
         return (HAAR_SIZE0 + HAAR_SIZE_INC * layer) << octave;\r
     }\r
     \r
-    class SURF_GPU_Invoker : private CvSURFParams\r
+    class SURF_GPU_Invoker\r
     {\r
     public:\r
         SURF_GPU_Invoker(SURF_GPU& surf, const GpuMat& img, const GpuMat& mask) :\r
-            CvSURFParams(surf),\r
-\r
-            sum(surf.sum), mask1(surf.mask1), maskSum(surf.maskSum), intBuffer(surf.intBuffer), det(surf.det), trace(surf.trace),\r
-\r
-            maxPosBuffer(surf.maxPosBuffer), \r
-\r
+            surf_(surf),\r
             img_cols(img.cols), img_rows(img.rows),\r
-\r
             use_mask(!mask.empty())\r
         {\r
             CV_Assert(!img.empty() && img.type() == CV_8UC1);\r
             CV_Assert(mask.empty() || (mask.size() == img.size() && mask.type() == CV_8UC1));\r
-            CV_Assert(nOctaves > 0 && nOctaveLayers > 0);\r
+            CV_Assert(surf_.nOctaves > 0 && surf_.nOctaveLayers > 0);\r
             CV_Assert(TargetArchs::builtWith(GLOBAL_ATOMICS) && DeviceInfo().supports(GLOBAL_ATOMICS));\r
                 \r
-            const int min_size = calcSize(nOctaves - 1, 0);\r
+            const int min_size = calcSize(surf_.nOctaves - 1, 0);\r
             CV_Assert(img_rows - min_size >= 0);\r
             CV_Assert(img_cols - min_size >= 0);\r
             \r
-            const int layer_rows = img_rows >> (nOctaves - 1);\r
-            const int layer_cols = img_cols >> (nOctaves - 1);\r
-            const int min_margin = ((calcSize((nOctaves - 1), 2) >> 1) >> (nOctaves - 1)) + 1;\r
+            const int layer_rows = img_rows >> (surf_.nOctaves - 1);\r
+            const int layer_cols = img_cols >> (surf_.nOctaves - 1);\r
+            const int min_margin = ((calcSize((surf_.nOctaves - 1), 2) >> 1) >> (surf_.nOctaves - 1)) + 1;\r
             CV_Assert(layer_rows - 2 * min_margin > 0);\r
             CV_Assert(layer_cols - 2 * min_margin > 0);\r
 \r
@@ -143,44 +137,44 @@ namespace
 \r
             CV_Assert(maxFeatures > 0);\r
 \r
-            counters.create(1, nOctaves + 1, CV_32SC1);\r
+            counters.create(1, surf_.nOctaves + 1, CV_32SC1);\r
             counters.setTo(Scalar::all(0));\r
 \r
-            loadGlobalConstants(maxCandidates, maxFeatures, img_rows, img_cols, nOctaveLayers, static_cast<float>(hessianThreshold));\r
+            loadGlobalConstants(maxCandidates, maxFeatures, img_rows, img_cols, surf_.nOctaveLayers, static_cast<float>(surf_.hessianThreshold));\r
 \r
             bindImgTex(img);\r
 \r
-            integralBuffered(img, sum, intBuffer);\r
-            bindSumTex(sum);\r
+            integralBuffered(img, surf_.sum, surf_.intBuffer);\r
+            bindSumTex(surf_.sum);\r
 \r
             if (use_mask)\r
             {\r
-                min(mask, 1.0, mask1);\r
-                integralBuffered(mask1, maskSum, intBuffer);\r
-                bindMaskSumTex(maskSum);\r
+                min(mask, 1.0, surf_.mask1);\r
+                integralBuffered(surf_.mask1, surf_.maskSum, surf_.intBuffer);\r
+                bindMaskSumTex(surf_.maskSum);\r
             }\r
         }\r
 \r
         void detectKeypoints(GpuMat& keypoints)\r
         {\r
-            ensureSizeIsEnough(img_rows * (nOctaveLayers + 2), img_cols, CV_32FC1, det);\r
-            ensureSizeIsEnough(img_rows * (nOctaveLayers + 2), img_cols, CV_32FC1, trace);\r
+            ensureSizeIsEnough(img_rows * (surf_.nOctaveLayers + 2), img_cols, CV_32FC1, surf_.det);\r
+            ensureSizeIsEnough(img_rows * (surf_.nOctaveLayers + 2), img_cols, CV_32FC1, surf_.trace);\r
             \r
-            ensureSizeIsEnough(1, maxCandidates, CV_32SC4, maxPosBuffer);\r
+            ensureSizeIsEnough(1, maxCandidates, CV_32SC4, surf_.maxPosBuffer);\r
             ensureSizeIsEnough(SURF_GPU::SF_FEATURE_STRIDE, maxFeatures, CV_32FC1, keypoints);\r
             keypoints.setTo(Scalar::all(0));\r
 \r
-            for (int octave = 0; octave < nOctaves; ++octave)\r
+            for (int octave = 0; octave < surf_.nOctaves; ++octave)\r
             {\r
                 const int layer_rows = img_rows >> octave;\r
                 const int layer_cols = img_cols >> octave;\r
 \r
                 loadOctaveConstants(octave, layer_rows, layer_cols);\r
 \r
-                icvCalcLayerDetAndTrace_gpu(det, trace, img_rows, img_cols, octave, nOctaveLayers);\r
+                icvCalcLayerDetAndTrace_gpu(surf_.det, surf_.trace, img_rows, img_cols, octave, surf_.nOctaveLayers);\r
 \r
-                icvFindMaximaInLayer_gpu(det, trace, maxPosBuffer.ptr<int4>(), counters.ptr<unsigned int>() + 1 + octave,\r
-                    img_rows, img_cols, octave, use_mask, nOctaveLayers);\r
+                icvFindMaximaInLayer_gpu(surf_.det, surf_.trace, surf_.maxPosBuffer.ptr<int4>(), counters.ptr<unsigned int>() + 1 + octave,\r
+                    img_rows, img_cols, octave, use_mask, surf_.nOctaveLayers);\r
 \r
                 unsigned int maxCounter;\r
                 cudaSafeCall( cudaMemcpy(&maxCounter, counters.ptr<unsigned int>() + 1 + octave, sizeof(unsigned int), cudaMemcpyDeviceToHost) );\r
@@ -188,7 +182,7 @@ namespace
 \r
                 if (maxCounter > 0)\r
                 {\r
-                    icvInterpolateKeypoint_gpu(det, maxPosBuffer.ptr<int4>(), maxCounter, \r
+                    icvInterpolateKeypoint_gpu(surf_.det, surf_.maxPosBuffer.ptr<int4>(), maxCounter, \r
                         keypoints.ptr<float>(SURF_GPU::SF_X), keypoints.ptr<float>(SURF_GPU::SF_Y),\r
                         keypoints.ptr<int>(SURF_GPU::SF_LAPLACIAN), keypoints.ptr<float>(SURF_GPU::SF_SIZE),\r
                         keypoints.ptr<float>(SURF_GPU::SF_HESSIAN), counters.ptr<unsigned int>());\r
@@ -200,7 +194,7 @@ namespace
 \r
             keypoints.cols = featureCounter;\r
 \r
-            if (upright)\r
+            if (surf_.upright)\r
                 keypoints.row(SURF_GPU::SF_DIR).setTo(Scalar::all(90.0));\r
             else\r
                 findOrientation(keypoints);\r
@@ -228,15 +222,7 @@ namespace
         }\r
 \r
     private:\r
-        GpuMat& sum;\r
-        GpuMat& mask1;\r
-        GpuMat& maskSum;\r
-        GpuMat& intBuffer;\r
-\r
-        GpuMat& det;\r
-        GpuMat& trace;\r
-\r
-        GpuMat& maxPosBuffer;\r
+        SURF_GPU& surf_;\r
 \r
         int img_cols, img_rows;\r
 \r
@@ -306,7 +292,7 @@ void cv::gpu::SURF_GPU::uploadKeypoints(const vector<KeyPoint>& keypoints, GpuMa
 \r
 namespace\r
 {\r
-    int getPointOctave(float size, const CvSURFParams& params)\r
+    int getPointOctave(float size, const SURF_GPU& params)\r
     {\r
         int best_octave = 0;\r
         float min_diff = numeric_limits<float>::max();\r

diff --git a/modules/gpu/test/precomp.hpp b/modules/gpu/test/precomp.hpp
index b5f0bbb..f03db7b 100644 (file)
--- a/modules/gpu/test/precomp.hpp
+++ b/modules/gpu/test/precomp.hpp
@@ -61,6 +61,7 @@
 #include "opencv2/ts/ts.hpp"\r
 #include "opencv2/ts/ts_perf.hpp"\r
 #include "opencv2/gpu/gpu.hpp"\r
+#include "opencv2/nonfree/nonfree.hpp"\r
 \r
 #include "utility.hpp"\r
 #include "interpolation.hpp"\r

diff --git a/modules/gpu/test/test_arithm.cpp b/modules/gpu/test/test_arithm.cpp
index 2fa9b22..b8fed5e 100644 (file)
--- a/modules/gpu/test/test_arithm.cpp
+++ b/modules/gpu/test/test_arithm.cpp
@@ -385,7 +385,7 @@ TEST_P(Sqrt, Array)
 \r
     gpuRes.download(dst);\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-6);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, Sqrt, Combine(\r

diff --git a/modules/gpu/test/test_features2d.cpp b/modules/gpu/test/test_features2d.cpp
index 06b0189..aedd11d 100644 (file)
--- a/modules/gpu/test/test_features2d.cpp
+++ b/modules/gpu/test/test_features2d.cpp
@@ -96,16 +96,16 @@ struct SURF : TestWithParam<cv::gpu::DeviceInfo>
         devInfo = GetParam();\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-        \r
+\r
         image = readImage("features2d/aloe.png", CV_LOAD_IMAGE_GRAYSCALE);\r
-        ASSERT_FALSE(image.empty());        \r
-        \r
+        ASSERT_FALSE(image.empty());\r
+\r
         mask = cv::Mat(image.size(), CV_8UC1, cv::Scalar::all(1));\r
         mask(cv::Range(0, image.rows / 2), cv::Range(0, image.cols / 2)).setTo(cv::Scalar::all(0));\r
-                \r
-        cv::SURF fdetector_gold; \r
+\r
+        cv::SURF fdetector_gold;\r
         fdetector_gold.extended = false;\r
-        fdetector_gold(image, mask, keypoints_gold, descriptors_gold);        \r
+        fdetector_gold(image, mask, keypoints_gold, descriptors_gold);\r
     }\r
 };\r
 \r
@@ -135,7 +135,7 @@ TEST_P(SURF, Accuracy)
 \r
     dev_descriptors.download(descriptors);\r
 \r
-    cv::BruteForceMatcher< cv::L2<float> > matcher;\r
+    cv::BFMatcher matcher(cv::NORM_L2);\r
     std::vector<cv::DMatch> matches;\r
 \r
     matcher.match(cv::Mat(static_cast<int>(keypoints_gold.size()), 64, CV_32FC1, &descriptors_gold[0]), descriptors, matches);\r
@@ -696,7 +696,7 @@ TEST_P(ORB, Accuracy)
 \r
     d_descriptors.download(descriptors);\r
 \r
-    cv::BruteForceMatcher<cv::Hamming> matcher;\r
+    cv::BFMatcher matcher(cv::NORM_HAMMING);\r
     std::vector<cv::DMatch> matches;\r
 \r
     matcher.match(descriptors_gold, descriptors, matches);\r

diff --git a/modules/gpu/test/test_filters.cpp b/modules/gpu/test/test_filters.cpp
index 9f1079f..6e2c2f1 100644 (file)
--- a/modules/gpu/test/test_filters.cpp
+++ b/modules/gpu/test/test_filters.cpp
@@ -334,6 +334,9 @@ PARAM_TEST_CASE(GaussianBlur, cv::gpu::DeviceInfo, cv::Size, UseRoi)
 \r
 TEST_P(GaussianBlur, Rgba)\r
 {\r
+    if (!devInfo.supports(cv::gpu::FEATURE_SET_COMPUTE_20) && ksize.height > 16)\r
+        return;\r
+\r
     cv::Mat dst_rgba;\r
 \r
     cv::gpu::GpuMat dev_dst_rgba;\r
@@ -347,6 +350,9 @@ TEST_P(GaussianBlur, Rgba)
 \r
 TEST_P(GaussianBlur, Gray)\r
 {\r
+    if (!devInfo.supports(cv::gpu::FEATURE_SET_COMPUTE_20) && ksize.height > 16)\r
+        return;\r
+\r
     cv::Mat dst_gray;\r
 \r
     cv::gpu::GpuMat dev_dst_gray;\r