Skip to content

O
opencv
Commit 85c904a4 authored by Vladislav Vinogradov's avatar Vladislav Vinogradov
Browse files
Options

fixed build in gpu module (SURF and ORB)

parent 46248851
Show whitespace changes
Inline Side-by-side
Showing with 91 additions and 101 deletions
modules/gpu/CMakeLists.txt
View file @ 85c904a4
......@@ -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")
......
modules/gpu/include/opencv2/gpu/gpu.hpp
View file @ 85c904a4
......@@ -1566,13 +1566,12 @@ public:
void releaseMemory();
// SURF parameters;
int extended;
int upright;
// SURF parameters
double hessianThreshold;
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<size_t> n_features_per_level_;
......
modules/gpu/src/orb.cpp
View file @ 85c904a4
......@@ -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<KeyPoint>&) { 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<KeyPoint>&, 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<KeyPoint>&) { throw_nogpu(); }
void cv::gpu::ORB_GPU::convertKeyPoints(Mat&, std::vector<KeyPoint>&) { 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)
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)
{
params_ = ORB((int)n_features, scaleFactor, nlevels);
// fill the extractors and descriptors for the corresponding scales
int n_levels = static_cast<int>(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<int> 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<float>(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<int>(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<int>(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<int>(params_.first_level_))
if (level != firstLevel_)
{
if (level < static_cast<int>(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<int>(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<short2>(0), keyPointsPyr_[level].ptr<float>(2),
keyPointsCount_[level], pattern_.ptr<int>(0), pattern_.ptr<int>(1), descRange, descriptorSize(), params_.WTA_K_, 0);
keyPointsCount_[level], pattern_.ptr<int>(0), pattern_.ptr<int>(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<int>(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<int>(params_.first_level_) ? sf : 1.0f;
float locScale = level != firstLevel_ ? sf : 1.0f;
mergeLocation_gpu(keyPointsPyr_[level].ptr<short2>(0), keyPointsRange.ptr<float>(0), keyPointsRange.ptr<float>(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];
}
......
modules/gpu/src/surf.cpp
View file @ 85c904a4
......@@ -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<float>(hessianThreshold));
loadGlobalConstants(maxCandidates, maxFeatures, img_rows, img_cols, surf_.nOctaveLayers, static_cast<float>(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<int4>(), counters.ptr<unsigned int>() + 1 + octave,
img_rows, img_cols, octave, use_mask, nOctaveLayers);
icvFindMaximaInLayer_gpu(surf_.det, surf_.trace, surf_.maxPosBuffer.ptr<int4>(), counters.ptr<unsigned int>() + 1 + octave,
img_rows, img_cols, octave, use_mask, surf_.nOctaveLayers);
unsigned int maxCounter;
cudaSafeCall( cudaMemcpy(&maxCounter, counters.ptr<unsigned int>() + 1 + octave, sizeof(unsigned int), cudaMemcpyDeviceToHost) );
......@@ -188,7 +182,7 @@ namespace
if (maxCounter > 0)
{
icvInterpolateKeypoint_gpu(det, maxPosBuffer.ptr<int4>(), maxCounter,
icvInterpolateKeypoint_gpu(surf_.det, surf_.maxPosBuffer.ptr<int4>(), maxCounter,
keypoints.ptr<float>(SURF_GPU::SF_X), keypoints.ptr<float>(SURF_GPU::SF_Y),
keypoints.ptr<int>(SURF_GPU::SF_LAPLACIAN), keypoints.ptr<float>(SURF_GPU::SF_SIZE),
keypoints.ptr<float>(SURF_GPU::SF_HESSIAN), counters.ptr<unsigned int>());
......@@ -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<KeyPoint>& 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<float>::max();
......
modules/gpu/test/precomp.hpp
View file @ 85c904a4
......@@ -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"
......
modules/gpu/test/test_arithm.cpp
View file @ 85c904a4
......@@ -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(
......
modules/gpu/test/test_features2d.cpp
View file @ 85c904a4
......@@ -135,7 +135,7 @@ TEST_P(SURF, Accuracy)
dev_descriptors.download(descriptors);
cv::BruteForceMatcher< cv::L2<float> > matcher;
cv::BFMatcher matcher(cv::NORM_L2);
std::vector<cv::DMatch> matches;
matcher.match(cv::Mat(static_cast<int>(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<cv::Hamming> matcher;
cv::BFMatcher matcher(cv::NORM_HAMMING);
std::vector<cv::DMatch> matches;
matcher.match(descriptors_gold, descriptors, matches);
......
modules/gpu/test/test_filters.cpp
View file @ 85c904a4
......@@ -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;
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment