From 85c904a4ba81001f1d1f2cfbf5fe6d8729c77057 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Sun, 18 Mar 2012 12:05:49 +0000 Subject: [PATCH] fixed build in gpu module (SURF and ORB) --- modules/gpu/CMakeLists.txt | 2 +- modules/gpu/include/opencv2/gpu/gpu.hpp | 26 ++++--- modules/gpu/src/orb.cpp | 93 ++++++++++++------------- modules/gpu/src/surf.cpp | 64 +++++++---------- modules/gpu/test/precomp.hpp | 1 + modules/gpu/test/test_arithm.cpp | 2 +- modules/gpu/test/test_features2d.cpp | 16 ++--- modules/gpu/test/test_filters.cpp | 6 ++ 8 files changed, 100 insertions(+), 110 deletions(-) diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt index 199b520963..a403b68b5d 100644 --- 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 6ecea83298..bc53975f77 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1566,13 +1566,12 @@ public: void releaseMemory(); - // SURF parameters; - int extended; - int upright; + // SURF parameters double hessianThreshold; - - int nOctaves; - int nOctaveLayers; + int nOctaves; + int nOctaveLayers; + bool extended; + bool upright; //! max keypoints = min(keypointsRatio * img.size().area(), 65535) float keypointsRatio; @@ -1662,9 +1661,8 @@ public: }; //! Constructor - //! n_features - the number of desired features - //! detector_params - parameters to use - explicit ORB_GPU(size_t n_features = 500, float scaleFactor = 1.2f, int nlevels = 3); + explicit ORB_GPU(int nFeatures = 500, float scaleFactor = 1.2f, int nLevels = 3, int edgeThreshold = 31, + int firstLevel = 0, int WTA_K = 2, int scoreType = 0, int patchSize = 31); //! Compute the ORB features on an image //! image - the image to compute the features (supports only CV_8UC1 images) @@ -1690,7 +1688,6 @@ public: //! returns the descriptor size in bytes inline int descriptorSize() const { return kBytes; } - void setParams(size_t n_features, float scaleFactor, int nlevels); inline void setFastParams(int threshold, bool nonmaxSupression = true) { fastDetector_.threshold = threshold; @@ -1714,7 +1711,14 @@ private: void mergeKeyPoints(GpuMat& keypoints); - ORB params_; + int nFeatures_; + float scaleFactor_; + int nLevels_; + int edgeThreshold_; + int firstLevel_; + int WTA_K_; + int scoreType_; + int patchSize_; // The number of desired features per scale std::vector n_features_per_level_; diff --git a/modules/gpu/src/orb.cpp b/modules/gpu/src/orb.cpp index 9972fced3e..fb5be1a6dc 100644 --- a/modules/gpu/src/orb.cpp +++ b/modules/gpu/src/orb.cpp @@ -48,14 +48,13 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) -cv::gpu::ORB_GPU::ORB_GPU(size_t, float, int) : fastDetector_(0) { throw_nogpu(); } +cv::gpu::ORB_GPU::ORB_GPU(int, float, int, int, int, int, int, int) { throw_nogpu(); } void cv::gpu::ORB_GPU::operator()(const GpuMat&, const GpuMat&, std::vector&) { throw_nogpu(); } void cv::gpu::ORB_GPU::operator()(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::ORB_GPU::operator()(const GpuMat&, const GpuMat&, std::vector&, GpuMat&) { throw_nogpu(); } void cv::gpu::ORB_GPU::operator()(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::ORB_GPU::downloadKeyPoints(GpuMat&, std::vector&) { throw_nogpu(); } void cv::gpu::ORB_GPU::convertKeyPoints(Mat&, std::vector&) { throw_nogpu(); } -void cv::gpu::ORB_GPU::setParams(size_t, float, int) { throw_nogpu(); } void cv::gpu::ORB_GPU::release() { throw_nogpu(); } void cv::gpu::ORB_GPU::buildScalePyramids(const GpuMat&, const GpuMat&) { throw_nogpu(); } void cv::gpu::ORB_GPU::computeKeyPointsPyramid() { throw_nogpu(); } @@ -83,16 +82,6 @@ namespace cv { namespace gpu { namespace device } }}} -cv::gpu::ORB_GPU::ORB_GPU(size_t n_features, float scaleFactor, int nlevels) : - fastDetector_(DEFAULT_FAST_THRESHOLD) -{ - setParams(n_features, scaleFactor, nlevels); - - blurFilter = createGaussianFilter_GPU(CV_8UC1, Size(7, 7), 2, 2, BORDER_REFLECT_101); - - blurForDescriptor = false; -} - namespace { const float HARRIS_K = 0.04f; @@ -407,27 +396,27 @@ namespace } } -void cv::gpu::ORB_GPU::setParams(size_t n_features, float scaleFactor, int nlevels) -{ - params_ = ORB((int)n_features, scaleFactor, nlevels); - +cv::gpu::ORB_GPU::ORB_GPU(int nFeatures, float scaleFactor, int nLevels, int edgeThreshold, int firstLevel, int WTA_K, int scoreType, int patchSize) : + nFeatures_(nFeatures), scaleFactor_(scaleFactor), nLevels_(nLevels), edgeThreshold_(edgeThreshold), firstLevel_(firstLevel), WTA_K_(WTA_K), + scoreType_(scoreType), patchSize_(patchSize), + fastDetector_(DEFAULT_FAST_THRESHOLD) +{ // fill the extractors and descriptors for the corresponding scales - int n_levels = static_cast(params_.n_levels_); - float factor = 1.0f / params_.scale_factor_; - float n_desired_features_per_scale = n_features * (1.0f - factor) / (1.0f - std::pow(factor, n_levels)); + float factor = 1.0f / scaleFactor_; + float n_desired_features_per_scale = nFeatures_ * (1.0f - factor) / (1.0f - std::pow(factor, nLevels_)); - n_features_per_level_.resize(n_levels); + n_features_per_level_.resize(nLevels_); int sum_n_features = 0; - for (int level = 0; level < n_levels - 1; ++level) + for (int level = 0; level < nLevels_ - 1; ++level) { n_features_per_level_[level] = cvRound(n_desired_features_per_scale); sum_n_features += n_features_per_level_[level]; n_desired_features_per_scale *= factor; } - n_features_per_level_[n_levels - 1] = n_features - sum_n_features; + n_features_per_level_[nLevels_ - 1] = nFeatures - sum_n_features; // pre-compute the end of a row in a circular patch - int half_patch_size = params_.patch_size_ / 2; + int half_patch_size = patchSize_ / 2; vector u_max(half_patch_size + 1); for (int v = 0; v <= half_patch_size * std::sqrt(2.f) / 2 + 1; ++v) u_max[v] = cvRound(std::sqrt(static_cast(half_patch_size * half_patch_size - v * v))); @@ -447,17 +436,17 @@ void cv::gpu::ORB_GPU::setParams(size_t n_features, float scaleFactor, int nleve const int npoints = 512; Point pattern_buf[npoints]; const Point* pattern0 = (const Point*)bit_pattern_31_; - if (params_.patch_size_ != 31) + if (patchSize_ != 31) { pattern0 = pattern_buf; - makeRandomPattern(params_.patch_size_, pattern_buf, npoints); + makeRandomPattern(patchSize_, pattern_buf, npoints); } - CV_Assert(params_.WTA_K_ == 2 || params_.WTA_K_ == 3 || params_.WTA_K_ == 4); + CV_Assert(WTA_K_ == 2 || WTA_K_ == 3 || WTA_K_ == 4); Mat h_pattern; - if (params_.WTA_K_ == 2) + if (WTA_K_ == 2) { h_pattern.create(2, npoints, CV_32SC1); @@ -473,17 +462,21 @@ void cv::gpu::ORB_GPU::setParams(size_t n_features, float scaleFactor, int nleve else { int ntuples = descriptorSize() * 4; - initializeOrbPattern(pattern0, h_pattern, ntuples, params_.WTA_K_, npoints); + initializeOrbPattern(pattern0, h_pattern, ntuples, WTA_K_, npoints); } pattern_.upload(h_pattern); + + blurFilter = createGaussianFilter_GPU(CV_8UC1, Size(7, 7), 2, 2, BORDER_REFLECT_101); + + blurForDescriptor = false; } namespace { - inline float getScale(const ORB::CommonParams& params, int level) + inline float getScale(float scaleFactor, int firstLevel, int level) { - return pow(params.scale_factor_, level - static_cast(params.first_level_)); + return pow(scaleFactor, level - firstLevel); } } @@ -492,12 +485,12 @@ void cv::gpu::ORB_GPU::buildScalePyramids(const GpuMat& image, const GpuMat& mas CV_Assert(image.type() == CV_8UC1); CV_Assert(mask.empty() || (mask.type() == CV_8UC1 && mask.size() == image.size())); - imagePyr_.resize(params_.n_levels_); - maskPyr_.resize(params_.n_levels_); + imagePyr_.resize(nLevels_); + maskPyr_.resize(nLevels_); - for (int level = 0; level < static_cast(params_.n_levels_); ++level) + for (int level = 0; level < nLevels_; ++level) { - float scale = 1.0f / getScale(params_, level); + float scale = 1.0f / getScale(scaleFactor_, firstLevel_, level); Size sz(cvRound(image.cols * scale), cvRound(image.rows * scale)); @@ -506,9 +499,9 @@ void cv::gpu::ORB_GPU::buildScalePyramids(const GpuMat& image, const GpuMat& mas maskPyr_[level].setTo(Scalar::all(255)); // Compute the resized image - if (level != static_cast(params_.first_level_)) + if (level != firstLevel_) { - if (level < static_cast(params_.first_level_)) + if (level < firstLevel_) { resize(image, imagePyr_[level], sz, 0, 0, INTER_LINEAR); @@ -534,7 +527,7 @@ void cv::gpu::ORB_GPU::buildScalePyramids(const GpuMat& image, const GpuMat& mas // Filter keypoints by image border ensureSizeIsEnough(sz, CV_8UC1, buf_); buf_.setTo(Scalar::all(0)); - Rect inner(params_.edge_threshold_, params_.edge_threshold_, sz.width - 2 * params_.edge_threshold_, sz.height - 2 * params_.edge_threshold_); + Rect inner(edgeThreshold_, edgeThreshold_, sz.width - 2 * edgeThreshold_, sz.height - 2 * edgeThreshold_); buf_(inner).setTo(Scalar::all(255)); bitwise_and(maskPyr_[level], buf_, maskPyr_[level]); @@ -566,12 +559,12 @@ void cv::gpu::ORB_GPU::computeKeyPointsPyramid() { using namespace cv::gpu::device::orb; - int half_patch_size = params_.patch_size_ / 2; + int half_patch_size = patchSize_ / 2; - keyPointsPyr_.resize(params_.n_levels_); - keyPointsCount_.resize(params_.n_levels_); + keyPointsPyr_.resize(nLevels_); + keyPointsCount_.resize(nLevels_); - for (int level = 0; level < static_cast(params_.n_levels_); ++level) + for (int level = 0; level < nLevels_; ++level) { keyPointsCount_[level] = fastDetector_.calcKeyPointsLocation(imagePyr_[level], maskPyr_[level]); @@ -582,7 +575,7 @@ void cv::gpu::ORB_GPU::computeKeyPointsPyramid() int n_features = n_features_per_level_[level]; - if (params_.score_type_ == ORB::CommonParams::HARRIS_SCORE) + if (scoreType_ == ORB::HARRIS_SCORE) { // Keep more points than necessary as FAST does not give amazing corners cull(keyPointsPyr_[level], keyPointsCount_[level], 2 * n_features); @@ -605,7 +598,7 @@ void cv::gpu::ORB_GPU::computeDescriptors(GpuMat& descriptors) int nAllkeypoints = 0; - for (size_t level = 0; level < params_.n_levels_; ++level) + for (int level = 0; level < nLevels_; ++level) nAllkeypoints += keyPointsCount_[level]; if (nAllkeypoints == 0) @@ -618,7 +611,7 @@ void cv::gpu::ORB_GPU::computeDescriptors(GpuMat& descriptors) int offset = 0; - for (size_t level = 0; level < params_.n_levels_; ++level) + for (int level = 0; level < nLevels_; ++level) { GpuMat descRange = descriptors.rowRange(offset, offset + keyPointsCount_[level]); @@ -630,7 +623,7 @@ void cv::gpu::ORB_GPU::computeDescriptors(GpuMat& descriptors) } computeOrbDescriptor_gpu(blurForDescriptor ? buf_ : imagePyr_[level], keyPointsPyr_[level].ptr(0), keyPointsPyr_[level].ptr(2), - keyPointsCount_[level], pattern_.ptr(0), pattern_.ptr(1), descRange, descriptorSize(), params_.WTA_K_, 0); + keyPointsCount_[level], pattern_.ptr(0), pattern_.ptr(1), descRange, descriptorSize(), WTA_K_, 0); offset += keyPointsCount_[level]; } @@ -642,7 +635,7 @@ void cv::gpu::ORB_GPU::mergeKeyPoints(GpuMat& keypoints) int nAllkeypoints = 0; - for (size_t level = 0; level < params_.n_levels_; ++level) + for (int level = 0; level < nLevels_; ++level) nAllkeypoints += keyPointsCount_[level]; if (nAllkeypoints == 0) @@ -655,13 +648,13 @@ void cv::gpu::ORB_GPU::mergeKeyPoints(GpuMat& keypoints) int offset = 0; - for (int level = 0; level < static_cast(params_.n_levels_); ++level) + for (int level = 0; level < nLevels_; ++level) { - float sf = getScale(params_, level); + float sf = getScale(scaleFactor_, firstLevel_, level); GpuMat keyPointsRange = keypoints.colRange(offset, offset + keyPointsCount_[level]); - float locScale = level != static_cast(params_.first_level_) ? sf : 1.0f; + float locScale = level != firstLevel_ ? sf : 1.0f; mergeLocation_gpu(keyPointsPyr_[level].ptr(0), keyPointsRange.ptr(0), keyPointsRange.ptr(1), keyPointsCount_[level], locScale, 0); @@ -669,7 +662,7 @@ void cv::gpu::ORB_GPU::mergeKeyPoints(GpuMat& keypoints) keyPointsPyr_[level](Range(1, 3), Range(0, keyPointsCount_[level])).copyTo(range); keyPointsRange.row(4).setTo(Scalar::all(level)); - keyPointsRange.row(5).setTo(Scalar::all(params_.patch_size_ * sf)); + keyPointsRange.row(5).setTo(Scalar::all(patchSize_ * sf)); offset += keyPointsCount_[level]; } diff --git a/modules/gpu/src/surf.cpp b/modules/gpu/src/surf.cpp index a083200763..de592f5e5c 100644 --- 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; } - class SURF_GPU_Invoker : private CvSURFParams + class SURF_GPU_Invoker { public: SURF_GPU_Invoker(SURF_GPU& surf, const GpuMat& img, const GpuMat& mask) : - CvSURFParams(surf), - - sum(surf.sum), mask1(surf.mask1), maskSum(surf.maskSum), intBuffer(surf.intBuffer), det(surf.det), trace(surf.trace), - - maxPosBuffer(surf.maxPosBuffer), - + surf_(surf), img_cols(img.cols), img_rows(img.rows), - use_mask(!mask.empty()) { CV_Assert(!img.empty() && img.type() == CV_8UC1); CV_Assert(mask.empty() || (mask.size() == img.size() && mask.type() == CV_8UC1)); - CV_Assert(nOctaves > 0 && nOctaveLayers > 0); + CV_Assert(surf_.nOctaves > 0 && surf_.nOctaveLayers > 0); CV_Assert(TargetArchs::builtWith(GLOBAL_ATOMICS) && DeviceInfo().supports(GLOBAL_ATOMICS)); - const int min_size = calcSize(nOctaves - 1, 0); + const int min_size = calcSize(surf_.nOctaves - 1, 0); CV_Assert(img_rows - min_size >= 0); CV_Assert(img_cols - min_size >= 0); - const int layer_rows = img_rows >> (nOctaves - 1); - const int layer_cols = img_cols >> (nOctaves - 1); - const int min_margin = ((calcSize((nOctaves - 1), 2) >> 1) >> (nOctaves - 1)) + 1; + const int layer_rows = img_rows >> (surf_.nOctaves - 1); + const int layer_cols = img_cols >> (surf_.nOctaves - 1); + const int min_margin = ((calcSize((surf_.nOctaves - 1), 2) >> 1) >> (surf_.nOctaves - 1)) + 1; CV_Assert(layer_rows - 2 * min_margin > 0); CV_Assert(layer_cols - 2 * min_margin > 0); @@ -143,44 +137,44 @@ namespace CV_Assert(maxFeatures > 0); - counters.create(1, nOctaves + 1, CV_32SC1); + counters.create(1, surf_.nOctaves + 1, CV_32SC1); counters.setTo(Scalar::all(0)); - loadGlobalConstants(maxCandidates, maxFeatures, img_rows, img_cols, nOctaveLayers, static_cast(hessianThreshold)); + loadGlobalConstants(maxCandidates, maxFeatures, img_rows, img_cols, surf_.nOctaveLayers, static_cast(surf_.hessianThreshold)); bindImgTex(img); - integralBuffered(img, sum, intBuffer); - bindSumTex(sum); + integralBuffered(img, surf_.sum, surf_.intBuffer); + bindSumTex(surf_.sum); if (use_mask) { - min(mask, 1.0, mask1); - integralBuffered(mask1, maskSum, intBuffer); - bindMaskSumTex(maskSum); + min(mask, 1.0, surf_.mask1); + integralBuffered(surf_.mask1, surf_.maskSum, surf_.intBuffer); + bindMaskSumTex(surf_.maskSum); } } void detectKeypoints(GpuMat& keypoints) { - ensureSizeIsEnough(img_rows * (nOctaveLayers + 2), img_cols, CV_32FC1, det); - ensureSizeIsEnough(img_rows * (nOctaveLayers + 2), img_cols, CV_32FC1, trace); + ensureSizeIsEnough(img_rows * (surf_.nOctaveLayers + 2), img_cols, CV_32FC1, surf_.det); + ensureSizeIsEnough(img_rows * (surf_.nOctaveLayers + 2), img_cols, CV_32FC1, surf_.trace); - ensureSizeIsEnough(1, maxCandidates, CV_32SC4, maxPosBuffer); + ensureSizeIsEnough(1, maxCandidates, CV_32SC4, surf_.maxPosBuffer); ensureSizeIsEnough(SURF_GPU::SF_FEATURE_STRIDE, maxFeatures, CV_32FC1, keypoints); keypoints.setTo(Scalar::all(0)); - for (int octave = 0; octave < nOctaves; ++octave) + for (int octave = 0; octave < surf_.nOctaves; ++octave) { const int layer_rows = img_rows >> octave; const int layer_cols = img_cols >> octave; loadOctaveConstants(octave, layer_rows, layer_cols); - icvCalcLayerDetAndTrace_gpu(det, trace, img_rows, img_cols, octave, nOctaveLayers); + icvCalcLayerDetAndTrace_gpu(surf_.det, surf_.trace, img_rows, img_cols, octave, surf_.nOctaveLayers); - icvFindMaximaInLayer_gpu(det, trace, maxPosBuffer.ptr(), counters.ptr() + 1 + octave, - img_rows, img_cols, octave, use_mask, nOctaveLayers); + icvFindMaximaInLayer_gpu(surf_.det, surf_.trace, surf_.maxPosBuffer.ptr(), counters.ptr() + 1 + octave, + img_rows, img_cols, octave, use_mask, surf_.nOctaveLayers); unsigned int maxCounter; cudaSafeCall( cudaMemcpy(&maxCounter, counters.ptr() + 1 + octave, sizeof(unsigned int), cudaMemcpyDeviceToHost) ); @@ -188,7 +182,7 @@ namespace if (maxCounter > 0) { - icvInterpolateKeypoint_gpu(det, maxPosBuffer.ptr(), maxCounter, + icvInterpolateKeypoint_gpu(surf_.det, surf_.maxPosBuffer.ptr(), maxCounter, keypoints.ptr(SURF_GPU::SF_X), keypoints.ptr(SURF_GPU::SF_Y), keypoints.ptr(SURF_GPU::SF_LAPLACIAN), keypoints.ptr(SURF_GPU::SF_SIZE), keypoints.ptr(SURF_GPU::SF_HESSIAN), counters.ptr()); @@ -200,7 +194,7 @@ namespace keypoints.cols = featureCounter; - if (upright) + if (surf_.upright) keypoints.row(SURF_GPU::SF_DIR).setTo(Scalar::all(90.0)); else findOrientation(keypoints); @@ -228,15 +222,7 @@ namespace } private: - GpuMat& sum; - GpuMat& mask1; - GpuMat& maskSum; - GpuMat& intBuffer; - - GpuMat& det; - GpuMat& trace; - - GpuMat& maxPosBuffer; + SURF_GPU& surf_; int img_cols, img_rows; @@ -306,7 +292,7 @@ void cv::gpu::SURF_GPU::uploadKeypoints(const vector& keypoints, GpuMa namespace { - int getPointOctave(float size, const CvSURFParams& params) + int getPointOctave(float size, const SURF_GPU& params) { int best_octave = 0; float min_diff = numeric_limits::max(); diff --git a/modules/gpu/test/precomp.hpp b/modules/gpu/test/precomp.hpp index b5f0bbb6f7..f03db7ba87 100644 --- a/modules/gpu/test/precomp.hpp +++ b/modules/gpu/test/precomp.hpp @@ -61,6 +61,7 @@ #include "opencv2/ts/ts.hpp" #include "opencv2/ts/ts_perf.hpp" #include "opencv2/gpu/gpu.hpp" +#include "opencv2/nonfree/nonfree.hpp" #include "utility.hpp" #include "interpolation.hpp" diff --git a/modules/gpu/test/test_arithm.cpp b/modules/gpu/test/test_arithm.cpp index 2fa9b2200f..b8fed5ef86 100644 --- a/modules/gpu/test/test_arithm.cpp +++ b/modules/gpu/test/test_arithm.cpp @@ -385,7 +385,7 @@ TEST_P(Sqrt, Array) gpuRes.download(dst); - EXPECT_MAT_NEAR(dst_gold, dst, 0.0); + EXPECT_MAT_NEAR(dst_gold, dst, 1e-6); } INSTANTIATE_TEST_CASE_P(Arithm, Sqrt, Combine( diff --git a/modules/gpu/test/test_features2d.cpp b/modules/gpu/test/test_features2d.cpp index 06b018954b..aedd11d246 100644 --- a/modules/gpu/test/test_features2d.cpp +++ b/modules/gpu/test/test_features2d.cpp @@ -96,16 +96,16 @@ struct SURF : TestWithParam devInfo = GetParam(); cv::gpu::setDevice(devInfo.deviceID()); - + image = readImage("features2d/aloe.png", CV_LOAD_IMAGE_GRAYSCALE); - ASSERT_FALSE(image.empty()); - + ASSERT_FALSE(image.empty()); + mask = cv::Mat(image.size(), CV_8UC1, cv::Scalar::all(1)); mask(cv::Range(0, image.rows / 2), cv::Range(0, image.cols / 2)).setTo(cv::Scalar::all(0)); - - cv::SURF fdetector_gold; + + cv::SURF fdetector_gold; fdetector_gold.extended = false; - fdetector_gold(image, mask, keypoints_gold, descriptors_gold); + fdetector_gold(image, mask, keypoints_gold, descriptors_gold); } }; @@ -135,7 +135,7 @@ TEST_P(SURF, Accuracy) dev_descriptors.download(descriptors); - cv::BruteForceMatcher< cv::L2 > matcher; + cv::BFMatcher matcher(cv::NORM_L2); std::vector matches; matcher.match(cv::Mat(static_cast(keypoints_gold.size()), 64, CV_32FC1, &descriptors_gold[0]), descriptors, matches); @@ -696,7 +696,7 @@ TEST_P(ORB, Accuracy) d_descriptors.download(descriptors); - cv::BruteForceMatcher matcher; + cv::BFMatcher matcher(cv::NORM_HAMMING); std::vector matches; matcher.match(descriptors_gold, descriptors, matches); diff --git a/modules/gpu/test/test_filters.cpp b/modules/gpu/test/test_filters.cpp index 9f1079f33d..6e2c2f1137 100644 --- 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) TEST_P(GaussianBlur, Rgba) { + if (!devInfo.supports(cv::gpu::FEATURE_SET_COMPUTE_20) && ksize.height > 16) + return; + cv::Mat dst_rgba; cv::gpu::GpuMat dev_dst_rgba; @@ -347,6 +350,9 @@ TEST_P(GaussianBlur, Rgba) TEST_P(GaussianBlur, Gray) { + if (!devInfo.supports(cv::gpu::FEATURE_SET_COMPUTE_20) && ksize.height > 16) + return; + cv::Mat dst_gray; cv::gpu::GpuMat dev_dst_gray;