generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

implemented optimized version of bf_radius_match for train collection

Browse Source
This commit is contained in:
Vladislav Vinogradov 2011-09-27 06:45:17 +00:00
parent e99b090871
commit ea4f65b570
4 changed files with 527 additions and 207 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

33
modules/gpu/include/opencv2/gpu/gpu.hpp
View File

@ -1288,16 +1288,16 @@ namespace cv
const std::vector<GpuMat>& masks = std::vector<GpuMat>(), bool compactResult = false ); const std::vector<GpuMat>& masks = std::vector<GpuMat>(), bool compactResult = false );
// Find best matches for each query descriptor which have distance less than maxDistance. // Find best matches for each query descriptor which have distance less than maxDistance.
// nMatches.at<unsigned int>(0, queruIdx) will contain matches count for queryIdx. // nMatches.at<int>(0, queryIdx) will contain matches count for queryIdx.
// carefully nMatches can be greater than trainIdx.cols - it means that matcher didn't find all matches, // carefully nMatches can be greater than trainIdx.cols - it means that matcher didn't find all matches,
// because it didn't have enough memory. // because it didn't have enough memory.
// trainIdx.at<int>(queruIdx, i) will contain ith train index (i < min(nMatches.at<unsigned int>(0, queruIdx), trainIdx.cols)) // trainIdx.at<int>(queruIdx, i) will contain ith train index (i < min(nMatches.at<int>(0, queruIdx), trainIdx.cols))
// distance.at<int>(queruIdx, i) will contain ith distance (i < min(nMatches.at<unsigned int>(0, queruIdx), trainIdx.cols)) // distance.at<int>(queruIdx, i) will contain ith distance (i < min(nMatches.at<int>(0, queruIdx), trainIdx.cols))
// If trainIdx is empty, then trainIdx and distance will be created with size nQuery x nTrain, // If trainIdx is empty, then trainIdx and distance will be created with size nQuery x (nTrain / 2),
// otherwize user can pass own allocated trainIdx and distance with size nQuery x nMaxMatches // otherwize user can pass own allocated trainIdx and distance with size nQuery x nMaxMatches
// Matches doesn't sorted. // Matches doesn't sorted.
void radiusMatch(const GpuMat& queryDescs, const GpuMat& trainDescs, void radiusMatchSingle(const GpuMat& queryDescs, const GpuMat& trainDescs,
GpuMat& trainIdx, GpuMat& nMatches, GpuMat& distance, float maxDistance, GpuMat& trainIdx, GpuMat& distance, GpuMat& nMatches, float maxDistance,
const GpuMat& mask = GpuMat(), Stream& stream = Stream::Null()); const GpuMat& mask = GpuMat(), Stream& stream = Stream::Null());
// Download trainIdx, nMatches and distance and convert it to vector with DMatch. // Download trainIdx, nMatches and distance and convert it to vector with DMatch.
@ -1305,10 +1305,10 @@ namespace cv
// compactResult is used when mask is not empty. If compactResult is false matches // compactResult is used when mask is not empty. If compactResult is false matches
// vector will have the same size as queryDescriptors rows. If compactResult is true // vector will have the same size as queryDescriptors rows. If compactResult is true
// matches vector will not contain matches for fully masked out query descriptors. // matches vector will not contain matches for fully masked out query descriptors.
static void radiusMatchDownload(const GpuMat& trainIdx, const GpuMat& nMatches, const GpuMat& distance, static void radiusMatchDownload(const GpuMat& trainIdx, const GpuMat& distance, const GpuMat& nMatches,
std::vector< std::vector<DMatch> >& matches, bool compactResult = false); std::vector< std::vector<DMatch> >& matches, bool compactResult = false);
// Convert trainIdx, nMatches and distance to vector with DMatch. // Convert trainIdx, nMatches and distance to vector with DMatch.
static void radiusMatchConvert(const Mat& trainIdx, const Mat& nMatches, const Mat& distance, static void radiusMatchConvert(const Mat& trainIdx, const Mat& distance, const Mat& nMatches,
std::vector< std::vector<DMatch> >& matches, bool compactResult = false); std::vector< std::vector<DMatch> >& matches, bool compactResult = false);
// Find best matches for each query descriptor which have distance less than maxDistance // Find best matches for each query descriptor which have distance less than maxDistance
@ -1317,6 +1317,23 @@ namespace cv
std::vector< std::vector<DMatch> >& matches, float maxDistance, std::vector< std::vector<DMatch> >& matches, float maxDistance,
const GpuMat& mask = GpuMat(), bool compactResult = false); const GpuMat& mask = GpuMat(), bool compactResult = false);
// Find best matches for each query descriptor which have distance less than maxDistance.
// Matches doesn't sorted.
void radiusMatchCollection(const GpuMat& queryDescs, const GpuMat& trainCollection,
GpuMat& trainIdx, GpuMat& imgIdx, GpuMat& distance, GpuMat& nMatches, float maxDistance,
const GpuMat& maskCollection, Stream& stream = Stream::Null());
// Download trainIdx, imgIdx, nMatches and distance and convert it to vector with DMatch.
// matches will be sorted in increasing order of distances.
// compactResult is used when mask is not empty. If compactResult is false matches
// vector will have the same size as queryDescriptors rows. If compactResult is true
// matches vector will not contain matches for fully masked out query descriptors.
static void radiusMatchDownload(const GpuMat& trainIdx, const GpuMat& imgIdx, const GpuMat& distance, const GpuMat& nMatches,
std::vector< std::vector<DMatch> >& matches, bool compactResult = false);
// Convert trainIdx, nMatches and distance to vector with DMatch.
static void radiusMatchConvert(const Mat& trainIdx, const Mat& imgIdx, const Mat& distance, const Mat& nMatches,
std::vector< std::vector<DMatch> >& matches, bool compactResult = false);
// Find best matches from train collection for each query descriptor which have distance less than // Find best matches from train collection for each query descriptor which have distance less than
// maxDistance (in increasing order of distances). // maxDistance (in increasing order of distances).
void radiusMatch(const GpuMat& queryDescs, std::vector< std::vector<DMatch> >& matches, float maxDistance, void radiusMatch(const GpuMat& queryDescs, std::vector< std::vector<DMatch> >& matches, float maxDistance,

2
modules/gpu/perf/perf_features2d.cpp
View File

@ -89,7 +89,7 @@ PERF_TEST_P(DevInfo_DescSize, BruteForceMatcher_radiusMatch, testing::Combine(te
SIMPLE_TEST_CYCLE() SIMPLE_TEST_CYCLE()
{ {
matcher.radiusMatch(query, train, trainIdx, nMatches, distance, 2.0); matcher.radiusMatchSingle(query, train, trainIdx, distance, nMatches, 2.0);
} }
Mat trainIdx_host(trainIdx); Mat trainIdx_host(trainIdx);

269
modules/gpu/src/brute_force_matcher.cpp
View File

@ -68,10 +68,13 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatchDownload(const GpuMat&, const
void cv::gpu::BruteForceMatcher_GPU_base::knnMatchConvert(const Mat&, const Mat&, std::vector< std::vector<DMatch> >&, bool) { throw_nogpu(); } void cv::gpu::BruteForceMatcher_GPU_base::knnMatchConvert(const Mat&, const Mat&, std::vector< std::vector<DMatch> >&, bool) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat&, const GpuMat&, std::vector< std::vector<DMatch> >&, int, const GpuMat&, bool) { throw_nogpu(); } void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat&, const GpuMat&, std::vector< std::vector<DMatch> >&, int, const GpuMat&, bool) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat&, std::vector< std::vector<DMatch> >&, int, const std::vector<GpuMat>&, bool) { throw_nogpu(); } void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat&, std::vector< std::vector<DMatch> >&, int, const std::vector<GpuMat>&, bool) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, float, const GpuMat&, Stream&) { throw_nogpu(); } void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchSingle(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, float, const GpuMat&, Stream&) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchDownload(const GpuMat&, const GpuMat&, const GpuMat&, std::vector< std::vector<DMatch> >&, bool) { throw_nogpu(); } void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchDownload(const GpuMat&, const GpuMat&, const GpuMat&, std::vector< std::vector<DMatch> >&, bool) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchConvert(const Mat&, const Mat&, const Mat&, std::vector< std::vector<DMatch> >&, bool) { throw_nogpu(); } void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchConvert(const Mat&, const Mat&, const Mat&, std::vector< std::vector<DMatch> >&, bool) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat&, const GpuMat&, std::vector< std::vector<DMatch> >&, float, const GpuMat&, bool) { throw_nogpu(); } void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat&, const GpuMat&, std::vector< std::vector<DMatch> >&, float, const GpuMat&, bool) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchCollection(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, float, const GpuMat&, Stream&) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchDownload(const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, vector< vector<DMatch> >&, bool) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchConvert(const Mat&, const Mat&, const Mat&, const Mat&, vector< vector<DMatch> >&, bool) { throw_nogpu(); }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat&, std::vector< std::vector<DMatch> >&, float, const std::vector<GpuMat>&, bool) { throw_nogpu(); } void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat&, std::vector< std::vector<DMatch> >&, float, const std::vector<GpuMat>&, bool) { throw_nogpu(); }
#else /* !defined (HAVE_CUDA) */ #else /* !defined (HAVE_CUDA) */
@ -114,27 +117,27 @@ namespace cv { namespace gpu { namespace bf_knnmatch
namespace cv { namespace gpu { namespace bf_radius_match namespace cv { namespace gpu { namespace bf_radius_match
{ {
template <typename T> void radiusMatchL1_gpu(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask, template <typename T> void radiusMatchSingleL1_gpu(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask,
const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream); cudaStream_t stream);
template <typename T> void radiusMatchL2_gpu(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask, template <typename T> void radiusMatchSingleL2_gpu(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask,
const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream); cudaStream_t stream);
template <typename T> void radiusMatchHamming_gpu(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask, template <typename T> void radiusMatchSingleHamming_gpu(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask,
const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream);
template <typename T> void radiusMatchCollectionL1_gpu(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection,
const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream);
template <typename T> void radiusMatchCollectionL2_gpu(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection,
const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream);
template <typename T> void radiusMatchCollectionHamming_gpu(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection,
const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream); cudaStream_t stream);
}}} }}}
namespace
{
struct ImgIdxSetter
{
explicit inline ImgIdxSetter(int imgIdx_) : imgIdx(imgIdx_) {}
inline void operator()(DMatch& m) const {m.imgIdx = imgIdx;}
int imgIdx;
};
}
cv::gpu::BruteForceMatcher_GPU_base::BruteForceMatcher_GPU_base(DistType distType_) : distType(distType_) cv::gpu::BruteForceMatcher_GPU_base::BruteForceMatcher_GPU_base(DistType distType_) : distType(distType_)
{ {
} }
@ -551,6 +554,16 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat& queryDescs, con
knnMatchDownload(trainIdx, distance, matches, compactResult); knnMatchDownload(trainIdx, distance, matches, compactResult);
} }
namespace
{
struct ImgIdxSetter
{
explicit inline ImgIdxSetter(int imgIdx_) : imgIdx(imgIdx_) {}
inline void operator()(DMatch& m) const {m.imgIdx = imgIdx;}
int imgIdx;
};
}
void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat& queryDescs, void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat& queryDescs,
vector< vector<DMatch> >& matches, int knn, const vector<GpuMat>& masks, bool compactResult) vector< vector<DMatch> >& matches, int knn, const vector<GpuMat>& masks, bool compactResult)
{ {
@ -596,8 +609,8 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat& queryDescs,
//////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////
// RadiusMatch // RadiusMatch
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat& queryDescs, const GpuMat& trainDescs, void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchSingle(const GpuMat& queryDescs, const GpuMat& trainDescs,
GpuMat& trainIdx, GpuMat& nMatches, GpuMat& distance, float maxDistance, const GpuMat& mask, Stream& stream) GpuMat& trainIdx, GpuMat& distance, GpuMat& nMatches, float maxDistance, const GpuMat& mask, Stream& stream)
{ {
if (queryDescs.empty() || trainDescs.empty()) if (queryDescs.empty() || trainDescs.empty())
return; return;
@ -605,26 +618,26 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat& queryDescs,
using namespace cv::gpu::bf_radius_match; using namespace cv::gpu::bf_radius_match;
typedef void (*radiusMatch_caller_t)(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask, typedef void (*radiusMatch_caller_t)(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask,
const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream); cudaStream_t stream);
static const radiusMatch_caller_t radiusMatch_callers[3][8] = static const radiusMatch_caller_t radiusMatch_callers[3][8] =
{ {
{ {
radiusMatchL1_gpu<unsigned char>, 0/*radiusMatchL1_gpu<signed char>*/, radiusMatchL1_gpu<unsigned short>, radiusMatchSingleL1_gpu<unsigned char>, 0/*radiusMatchSingleL1_gpu<signed char>*/, radiusMatchSingleL1_gpu<unsigned short>,
radiusMatchL1_gpu<short>, radiusMatchL1_gpu<int>, radiusMatchL1_gpu<float>, 0, 0 radiusMatchSingleL1_gpu<short>, radiusMatchSingleL1_gpu<int>, radiusMatchSingleL1_gpu<float>, 0, 0
}, },
{ {
0/*radiusMatchL2_gpu<unsigned char>*/, 0/*radiusMatchL2_gpu<signed char>*/, 0/*radiusMatchL2_gpu<unsigned short>*/, 0/*radiusMatchSingleL2_gpu<unsigned char>*/, 0/*radiusMatchSingleL2_gpu<signed char>*/, 0/*radiusMatchSingleL2_gpu<unsigned short>*/,
0/*radiusMatchL2_gpu<short>*/, 0/*radiusMatchL2_gpu<int>*/, radiusMatchL2_gpu<float>, 0, 0 0/*radiusMatchSingleL2_gpu<short>*/, 0/*radiusMatchSingleL2_gpu<int>*/, radiusMatchSingleL2_gpu<float>, 0, 0
}, },
{ {
radiusMatchHamming_gpu<unsigned char>, 0/*radiusMatchHamming_gpu<signed char>*/, radiusMatchHamming_gpu<unsigned short>, radiusMatchSingleHamming_gpu<unsigned char>, 0/*radiusMatchSingleHamming_gpu<signed char>*/, radiusMatchSingleHamming_gpu<unsigned short>,
0/*radiusMatchHamming_gpu<short>*/, radiusMatchHamming_gpu<int>, 0, 0, 0 0/*radiusMatchSingleHamming_gpu<short>*/, radiusMatchSingleHamming_gpu<int>, 0, 0, 0
} }
}; };
CV_Assert(TargetArchs::builtWith(SHARED_ATOMICS) && DeviceInfo().supports(GLOBAL_ATOMICS)); CV_Assert(TargetArchs::builtWith(SHARED_ATOMICS) && DeviceInfo().supports(SHARED_ATOMICS));
const int nQuery = queryDescs.rows; const int nQuery = queryDescs.rows;
const int nTrain = trainDescs.rows; const int nTrain = trainDescs.rows;
@ -636,38 +649,33 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat& queryDescs,
ensureSizeIsEnough(1, nQuery, CV_32SC1, nMatches); ensureSizeIsEnough(1, nQuery, CV_32SC1, nMatches);
if (trainIdx.empty()) if (trainIdx.empty())
{ {
ensureSizeIsEnough(nQuery, nTrain, CV_32SC1, trainIdx); ensureSizeIsEnough(nQuery, nTrain / 2, CV_32SC1, trainIdx);
ensureSizeIsEnough(nQuery, nTrain, CV_32FC1, distance); ensureSizeIsEnough(nQuery, nTrain / 2, CV_32FC1, distance);
} }
if (stream)
stream.enqueueMemSet(nMatches, Scalar::all(0));
else
nMatches.setTo(Scalar::all(0));
radiusMatch_caller_t func = radiusMatch_callers[distType][queryDescs.depth()]; radiusMatch_caller_t func = radiusMatch_callers[distType][queryDescs.depth()];
CV_Assert(func != 0); CV_Assert(func != 0);
func(queryDescs, trainDescs, maxDistance, mask, trainIdx, nMatches, distance, StreamAccessor::getStream(stream)); func(queryDescs, trainDescs, maxDistance, mask, trainIdx, distance, nMatches, StreamAccessor::getStream(stream));
} }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchDownload(const GpuMat& trainIdx, const GpuMat& nMatches, void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchDownload(const GpuMat& trainIdx, const GpuMat& distance, const GpuMat& nMatches,
const GpuMat& distance, std::vector< std::vector<DMatch> >& matches, bool compactResult) vector< vector<DMatch> >& matches, bool compactResult)
{ {
if (trainIdx.empty() || nMatches.empty() || distance.empty()) if (trainIdx.empty() || distance.empty() || nMatches.empty())
return; return;
Mat trainIdxCPU = trainIdx; Mat trainIdxCPU = trainIdx;
Mat nMatchesCPU = nMatches;
Mat distanceCPU = distance; Mat distanceCPU = distance;
Mat nMatchesCPU = nMatches;
radiusMatchConvert(trainIdxCPU, nMatchesCPU, distanceCPU, matches, compactResult); radiusMatchConvert(trainIdxCPU, distanceCPU, nMatchesCPU, matches, compactResult);
} }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchConvert(const Mat& trainIdx, const Mat& nMatches, const Mat& distance, void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchConvert(const Mat& trainIdx, const Mat& distance, const Mat& nMatches,
std::vector< std::vector<DMatch> >& matches, bool compactResult) vector< vector<DMatch> >& matches, bool compactResult)
{ {
if (trainIdx.empty() || nMatches.empty() || distance.empty()) if (trainIdx.empty() || distance.empty() || nMatches.empty())
return; return;
CV_Assert(trainIdx.type() == CV_32SC1); CV_Assert(trainIdx.type() == CV_32SC1);
@ -679,13 +687,135 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchConvert(const Mat& trainIdx
matches.clear(); matches.clear();
matches.reserve(nQuery); matches.reserve(nQuery);
const unsigned int* nMatches_ptr = nMatches.ptr<unsigned int>(); const int* nMatches_ptr = nMatches.ptr<int>();
for (int queryIdx = 0; queryIdx < nQuery; ++queryIdx) for (int queryIdx = 0; queryIdx < nQuery; ++queryIdx)
{ {
const int* trainIdx_ptr = trainIdx.ptr<int>(queryIdx); const int* trainIdx_ptr = trainIdx.ptr<int>(queryIdx);
const float* distance_ptr = distance.ptr<float>(queryIdx); const float* distance_ptr = distance.ptr<float>(queryIdx);
const int nMatches = std::min(static_cast<int>(nMatches_ptr[queryIdx]), trainIdx.cols); const int nMatches = std::min(nMatches_ptr[queryIdx], trainIdx.cols);
if (nMatches == 0)
{
if (!compactResult)
matches.push_back(vector<DMatch>());
continue;
}
matches.push_back(vector<DMatch>(nMatches));
vector<DMatch>& curMatches = matches.back();
for (int i = 0; i < nMatches; ++i, ++trainIdx_ptr, ++distance_ptr)
{
int trainIdx = *trainIdx_ptr;
float distance = *distance_ptr;
DMatch m(queryIdx, trainIdx, 0, distance);
curMatches[i] = m;
}
sort(curMatches.begin(), curMatches.end());
}
}
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat& queryDescs, const GpuMat& trainDescs,
vector< vector<DMatch> >& matches, float maxDistance, const GpuMat& mask, bool compactResult)
{
GpuMat trainIdx, distance, nMatches;
radiusMatchSingle(queryDescs, trainDescs, trainIdx, distance, nMatches, maxDistance, mask);
radiusMatchDownload(trainIdx, distance, nMatches, matches, compactResult);
}
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchCollection(const GpuMat& queryDescs, const GpuMat& trainCollection,
GpuMat& trainIdx, GpuMat& imgIdx, GpuMat& distance, GpuMat& nMatches, float maxDistance,
const GpuMat& maskCollection, Stream& stream)
{
if (queryDescs.empty() || trainCollection.empty())
return;
using namespace cv::gpu::bf_radius_match;
typedef void (*radiusMatch_caller_t)(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection,
const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream);
static const radiusMatch_caller_t radiusMatch_callers[3][8] =
{
{
radiusMatchCollectionL1_gpu<unsigned char>, 0/*radiusMatchCollectionL1_gpu<signed char>*/, radiusMatchCollectionL1_gpu<unsigned short>,
radiusMatchCollectionL1_gpu<short>, radiusMatchCollectionL1_gpu<int>, radiusMatchCollectionL1_gpu<float>, 0, 0
},
{
0/*radiusMatchCollectionL2_gpu<unsigned char>*/, 0/*radiusMatchCollectionL2_gpu<signed char>*/, 0/*radiusMatchCollectionL2_gpu<unsigned short>*/,
0/*radiusMatchCollectionL2_gpu<short>*/, 0/*radiusMatchCollectionL2_gpu<int>*/, radiusMatchCollectionL2_gpu<float>, 0, 0
},
{
radiusMatchCollectionHamming_gpu<unsigned char>, 0/*radiusMatchCollectionHamming_gpu<signed char>*/, radiusMatchCollectionHamming_gpu<unsigned short>,
0/*radiusMatchCollectionHamming_gpu<short>*/, radiusMatchCollectionHamming_gpu<int>, 0, 0, 0
}
};
CV_Assert(TargetArchs::builtWith(SHARED_ATOMICS) && DeviceInfo().supports(SHARED_ATOMICS));
const int nQuery = queryDescs.rows;
CV_Assert(queryDescs.channels() == 1 && queryDescs.depth() < CV_64F);
CV_Assert(trainIdx.empty() || (trainIdx.rows == nQuery && trainIdx.size() == distance.size() && trainIdx.size() == imgIdx.size()));
ensureSizeIsEnough(1, nQuery, CV_32SC1, nMatches);
if (trainIdx.empty())
{
ensureSizeIsEnough(nQuery, nQuery / 2, CV_32SC1, trainIdx);
ensureSizeIsEnough(nQuery, nQuery / 2, CV_32SC1, imgIdx);
ensureSizeIsEnough(nQuery, nQuery / 2, CV_32FC1, distance);
}
radiusMatch_caller_t func = radiusMatch_callers[distType][queryDescs.depth()];
CV_Assert(func != 0);
func(queryDescs, trainCollection, maxDistance, maskCollection, trainIdx, imgIdx, distance, nMatches, StreamAccessor::getStream(stream));
}
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchDownload(const GpuMat& trainIdx, const GpuMat& imgIdx, const GpuMat& distance, const GpuMat& nMatches,
vector< vector<DMatch> >& matches, bool compactResult)
{
if (trainIdx.empty() || imgIdx.empty() || distance.empty() || nMatches.empty())
return;
Mat trainIdxCPU = trainIdx;
Mat imgIdxCPU = imgIdx;
Mat distanceCPU = distance;
Mat nMatchesCPU = nMatches;
radiusMatchConvert(trainIdxCPU, imgIdxCPU, distanceCPU, nMatchesCPU, matches, compactResult);
}
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchConvert(const Mat& trainIdx, const Mat& imgIdx, const Mat& distance, const Mat& nMatches,
vector< vector<DMatch> >& matches, bool compactResult)
{
if (trainIdx.empty() || imgIdx.empty() || distance.empty() || nMatches.empty())
return;
CV_Assert(trainIdx.type() == CV_32SC1);
CV_Assert(imgIdx.type() == CV_32SC1 && imgIdx.size() == trainIdx.size());
CV_Assert(distance.type() == CV_32FC1 && distance.size() == trainIdx.size());
CV_Assert(nMatches.type() == CV_32SC1 && nMatches.isContinuous() && nMatches.cols >= trainIdx.rows);
const int nQuery = trainIdx.rows;
matches.clear();
matches.reserve(nQuery);
const int* nMatches_ptr = nMatches.ptr<int>();
for (int queryIdx = 0; queryIdx < nQuery; ++queryIdx)
{
const int* trainIdx_ptr = trainIdx.ptr<int>(queryIdx);
const int* imgIdx_ptr = imgIdx.ptr<int>(queryIdx);
const float* distance_ptr = distance.ptr<float>(queryIdx);
const int nMatches = std::min(nMatches_ptr[queryIdx], trainIdx.cols);
if (nMatches == 0) if (nMatches == 0)
{ {
@ -698,63 +828,34 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchConvert(const Mat& trainIdx
vector<DMatch>& curMatches = matches.back(); vector<DMatch>& curMatches = matches.back();
curMatches.reserve(nMatches); curMatches.reserve(nMatches);
for (int i = 0; i < nMatches; ++i, ++trainIdx_ptr, ++distance_ptr) for (int i = 0; i < nMatches; ++i, ++trainIdx_ptr, ++imgIdx_ptr, ++distance_ptr)
{ {
int trainIdx = *trainIdx_ptr; int trainIdx = *trainIdx_ptr;
int imgIdx = *imgIdx_ptr;
float distance = *distance_ptr; float distance = *distance_ptr;
DMatch m(queryIdx, trainIdx, 0, distance); DMatch m(queryIdx, trainIdx, imgIdx, distance);
curMatches.push_back(m); curMatches.push_back(m);
} }
sort(curMatches.begin(), curMatches.end()); sort(curMatches.begin(), curMatches.end());
} }
} }
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat& queryDescs, const GpuMat& trainDescs,
vector< vector<DMatch> >& matches, float maxDistance, const GpuMat& mask, bool compactResult)
{
GpuMat trainIdx, nMatches, distance;
radiusMatch(queryDescs, trainDescs, trainIdx, nMatches, distance, maxDistance, mask);
radiusMatchDownload(trainIdx, nMatches, distance, matches, compactResult);
}
void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat& queryDescs, vector< vector<DMatch> >& matches, void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat& queryDescs, vector< vector<DMatch> >& matches,
float maxDistance, const vector<GpuMat>& masks, bool compactResult) float maxDistance, const vector<GpuMat>& masks, bool compactResult)
{ {
if (queryDescs.empty() || empty()) GpuMat trainCollection;
return; GpuMat maskCollection;
matches.resize(queryDescs.rows); makeGpuCollection(trainCollection, maskCollection, masks);
vector< vector<DMatch> > curMatches; GpuMat trainIdx, imgIdx, distance, nMatches;
for (size_t imgIdx = 0; imgIdx < trainDescCollection.size(); ++imgIdx) radiusMatchCollection(queryDescs, trainCollection, trainIdx, imgIdx, distance, nMatches, maxDistance, maskCollection);
{
radiusMatch(queryDescs, trainDescCollection[imgIdx], curMatches, maxDistance,
masks.empty() ? GpuMat() : masks[imgIdx]);
for (int queryIdx = 0; queryIdx < queryDescs.rows; ++queryIdx) radiusMatchDownload(trainIdx, imgIdx, distance, nMatches, matches, compactResult);
{
vector<DMatch>& localMatch = curMatches[queryIdx];
vector<DMatch>& globalMatch = matches[queryIdx];
for_each(localMatch.begin(), localMatch.end(), ImgIdxSetter(static_cast<int>(imgIdx)));
const size_t oldSize = globalMatch.size();
copy(localMatch.begin(), localMatch.end(), back_inserter(globalMatch));
inplace_merge(globalMatch.begin(), globalMatch.begin() + oldSize, globalMatch.end());
}
}
if (compactResult)
{
vector< vector<DMatch> >::iterator new_end = remove_if(matches.begin(), matches.end(),
mem_fun_ref(&vector<DMatch>::empty));
matches.erase(new_end, matches.end());
}
} }
#endif /* !defined (HAVE_CUDA) */ #endif /* !defined (HAVE_CUDA) */

410
modules/gpu/src/cuda/bf_radius_match.cu
View File

@ -49,94 +49,210 @@ using namespace cv::gpu::device;
namespace cv { namespace gpu { namespace bf_radius_match namespace cv { namespace gpu { namespace bf_radius_match
{ {
__device__ __forceinline__ void store(const int* sidx, const float* sdist, const unsigned int scount, int* trainIdx, float* distance, int& sglob_ind, const int tid) template <typename T> struct SingleTrain
{ {
if (tid < scount) enum {USE_IMG_IDX = 0};
explicit SingleTrain(const DevMem2D_<T>& train_) : train(train_)
{ {
trainIdx[sglob_ind + tid] = sidx[tid]; }
distance[sglob_ind + tid] = sdist[tid];
static __device__ __forceinline__ void store(const int* s_trainIdx, const int* s_imgIdx, const float* s_dist, unsigned int& s_count, int& s_globInd,
int* trainIdx, int* imgIdx, float* distance, int maxCount)
{
const int tid = threadIdx.y * blockDim.x + threadIdx.x;
if (tid < s_count && s_globInd + tid < maxCount)
{
trainIdx[s_globInd + tid] = s_trainIdx[tid];
distance[s_globInd + tid] = s_dist[tid];
} }
if (tid == 0) if (tid == 0)
sglob_ind += scount; {
s_globInd += s_count;
s_count = 0;
}
} }
template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int BLOCK_STACK, typename VecDiff, typename Dist, typename T, typename Mask> template <int BLOCK_STACK, typename Dist, typename VecDiff, typename Mask>
__global__ void radiusMatch(const PtrStep_<T> query, const DevMem2D_<T> train, const float maxDistance, const Mask mask, __device__ __forceinline__ void loop(float maxDistance, Mask& mask, const VecDiff& vecDiff,
DevMem2Di trainIdx_, PtrStepf distance, unsigned int* nMatches) int* s_trainIdx, int* s_imgIdx, float* s_dist, unsigned int& s_count, int& s_globInd,
int* trainIdxRow, int* imgIdxRow, float* distanceRow, int maxCount,
typename Dist::result_type* s_diffRow) const
{ {
#if __CUDA_ARCH__ >= 120 #if __CUDA_ARCH__ >= 120
typedef typename Dist::result_type result_type; for (int i = 0; i < train.rows; i += blockDim.y)
typedef typename Dist::value_type value_type;
__shared__ result_type smem[BLOCK_DIM_X * BLOCK_DIM_Y];
__shared__ int sidx[BLOCK_STACK];
__shared__ float sdist[BLOCK_STACK];
__shared__ unsigned int scount;
__shared__ int sglob_ind;
const int queryIdx = blockIdx.x;
const int tid = threadIdx.y * BLOCK_DIM_X + threadIdx.x;
if (tid == 0)
{ {
scount = 0; int trainIdx = i + threadIdx.y;
sglob_ind = 0;
}
__syncthreads();
int* trainIdx_row = trainIdx_.ptr(queryIdx); if (trainIdx < train.rows && mask(blockIdx.x, trainIdx))
float* distance_row = distance.ptr(queryIdx);
const VecDiff vecDiff(query.ptr(queryIdx), train.cols, (typename Dist::value_type*)smem, tid, threadIdx.x);
typename Dist::result_type* sdiffRow = smem + BLOCK_DIM_X * threadIdx.y;
for (int trainIdx = threadIdx.y; trainIdx < train.rows; trainIdx += BLOCK_DIM_Y)
{
if (mask(queryIdx, trainIdx))
{ {
Dist dist; Dist dist;
const T* trainRow = train.ptr(trainIdx); vecDiff.calc(train.ptr(trainIdx), train.cols, dist, s_diffRow, threadIdx.x);
vecDiff.calc(trainRow, train.cols, dist, sdiffRow, threadIdx.x);
const typename Dist::result_type val = dist; const typename Dist::result_type val = dist;
if (threadIdx.x == 0 && val < maxDistance) if (threadIdx.x == 0 && val < maxDistance)
{ {
unsigned int i = atomicInc(&scount, (unsigned int) -1); unsigned int ind = atomicInc(&s_count, (unsigned int) -1);
sidx[i] = trainIdx; s_trainIdx[ind] = trainIdx;
sdist[i] = val; s_dist[ind] = val;
} }
} }
__syncthreads(); __syncthreads();
if (scount > BLOCK_STACK - BLOCK_DIM_Y) if (s_count >= BLOCK_STACK - blockDim.y)
{ store(s_trainIdx, s_imgIdx, s_dist, s_count, s_globInd, trainIdxRow, imgIdxRow, distanceRow, maxCount);
store(sidx, sdist, scount, trainIdx_row, distance_row, sglob_ind, tid);
if (tid == 0)
scount = 0;
}
__syncthreads(); __syncthreads();
} }
store(sidx, sdist, scount, trainIdx_row, distance_row, sglob_ind, tid); store(s_trainIdx, s_imgIdx, s_dist, s_count, s_globInd, trainIdxRow, imgIdxRow, distanceRow, maxCount);
if (tid == 0)
nMatches[queryIdx] = sglob_ind;
#endif #endif
} }
__device__ __forceinline__ int descLen() const
{
return train.cols;
}
const DevMem2D_<T> train;
};
template <typename T> struct TrainCollection
{
enum {USE_IMG_IDX = 1};
TrainCollection(const DevMem2D_<T>* trainCollection_, int nImg_, int desclen_) :
trainCollection(trainCollection_), nImg(nImg_), desclen(desclen_)
{
}
static __device__ __forceinline__ void store(const int* s_trainIdx, const int* s_imgIdx, const float* s_dist, unsigned int& s_count, int& s_globInd,
int* trainIdx, int* imgIdx, float* distance, int maxCount)
{
const int tid = threadIdx.y * blockDim.x + threadIdx.x;
if (tid < s_count && s_globInd + tid < maxCount)
{
trainIdx[s_globInd + tid] = s_trainIdx[tid];
imgIdx[s_globInd + tid] = s_imgIdx[tid];
distance[s_globInd + tid] = s_dist[tid];
}
if (tid == 0)
{
s_globInd += s_count;
s_count = 0;
}
}
template <int BLOCK_STACK, typename Dist, typename VecDiff, typename Mask>
__device__ void loop(float maxDistance, Mask& mask, const VecDiff& vecDiff,
int* s_trainIdx, int* s_imgIdx, float* s_dist, unsigned int& s_count, int& s_globInd,
int* trainIdxRow, int* imgIdxRow, float* distanceRow, int maxCount,
typename Dist::result_type* s_diffRow) const
{
#if __CUDA_ARCH__ >= 120
for (int imgIdx = 0; imgIdx < nImg; ++imgIdx)
{
const DevMem2D_<T> train = trainCollection[imgIdx];
mask.next();
for (int i = 0; i < train.rows; i += blockDim.y)
{
int trainIdx = i + threadIdx.y;
if (trainIdx < train.rows && mask(blockIdx.x, trainIdx))
{
Dist dist;
vecDiff.calc(train.ptr(trainIdx), desclen, dist, s_diffRow, threadIdx.x);
const typename Dist::result_type val = dist;
if (threadIdx.x == 0 && val < maxDistance)
{
unsigned int ind = atomicInc(&s_count, (unsigned int) -1);
s_trainIdx[ind] = trainIdx;
s_imgIdx[ind] = imgIdx;
s_dist[ind] = val;
}
}
__syncthreads();
if (s_count >= BLOCK_STACK - blockDim.y)
store(s_trainIdx, s_imgIdx, s_dist, s_count, s_globInd, trainIdxRow, imgIdxRow, distanceRow, maxCount);
__syncthreads();
}
}
store(s_trainIdx, s_imgIdx, s_dist, s_count, s_globInd, trainIdxRow, imgIdxRow, distanceRow, maxCount);
#endif
}
__device__ __forceinline__ int descLen() const
{
return desclen;
}
const DevMem2D_<T>* trainCollection;
const int nImg;
const int desclen;
};
template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int BLOCK_STACK, typename VecDiff, typename Dist, typename T, typename Train, typename Mask>
__global__ void radiusMatch(const PtrStep_<T> query, const Train train, float maxDistance, const Mask mask,
PtrStepi trainIdx, PtrStepi imgIdx, PtrStepf distance, int* nMatches, int maxCount)
{
typedef typename Dist::result_type result_type;
typedef typename Dist::value_type value_type;
__shared__ result_type s_mem[BLOCK_DIM_X * BLOCK_DIM_Y];
__shared__ int s_trainIdx[BLOCK_STACK];
__shared__ int s_imgIdx[Train::USE_IMG_IDX ? BLOCK_STACK : 1];
__shared__ float s_dist[BLOCK_STACK];
__shared__ unsigned int s_count;
__shared__ int s_globInd;
if (threadIdx.x == 0 && threadIdx.y == 0)
{
s_count = 0;
s_globInd = 0;
}
__syncthreads();
const VecDiff vecDiff(query.ptr(blockIdx.x), train.descLen(), (typename Dist::value_type*)s_mem, threadIdx.y * BLOCK_DIM_X + threadIdx.x, threadIdx.x);
Mask m = mask;
train.template loop<BLOCK_STACK, Dist>(maxDistance, m, vecDiff,
s_trainIdx, s_imgIdx, s_dist, s_count, s_globInd,
trainIdx.ptr(blockIdx.x), imgIdx.ptr(blockIdx.x), distance.ptr(blockIdx.x), maxCount,
s_mem + BLOCK_DIM_X * threadIdx.y);
if (threadIdx.x == 0 && threadIdx.y == 0)
nMatches[blockIdx.x] = s_globInd;
}
/////////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////////////////////
// Radius Match kernel caller // Radius Match kernel caller
template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int BLOCK_STACK, typename Dist, typename T, typename Mask> template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int BLOCK_STACK, typename Dist, typename T, typename Train, typename Mask>
void radiusMatchSimple_caller(const DevMem2D_<T>& query, const DevMem2D_<T>& train, float maxDistance, const Mask& mask, void radiusMatchSimple_caller(const DevMem2D_<T>& query, const Train& train, float maxDistance, const Mask& mask,
const DevMem2Di& trainIdx, const DevMem2Df& distance, unsigned int* nMatches, const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, int* nMatches,
cudaStream_t stream) cudaStream_t stream)
{ {
StaticAssert<BLOCK_STACK >= BLOCK_DIM_Y>::check(); StaticAssert<BLOCK_STACK >= BLOCK_DIM_Y>::check();
@ -146,16 +262,16 @@ namespace cv { namespace gpu { namespace bf_radius_match
const dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1); const dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);
radiusMatch<BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_STACK, VecDiffGlobal<BLOCK_DIM_X, T>, Dist, T> radiusMatch<BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_STACK, VecDiffGlobal<BLOCK_DIM_X, T>, Dist, T>
<<<grid, threads, 0, stream>>>(query, train, maxDistance, mask, trainIdx, distance, nMatches); <<<grid, threads, 0, stream>>>(query, train, maxDistance, mask, trainIdx, imgIdx, distance, nMatches, trainIdx.cols);
cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaGetLastError() );
if (stream == 0) if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() ); cudaSafeCall( cudaDeviceSynchronize() );
} }
template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int BLOCK_STACK, int MAX_LEN, bool LEN_EQ_MAX_LEN, typename Dist, typename T, typename Mask> template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int BLOCK_STACK, int MAX_LEN, bool LEN_EQ_MAX_LEN, typename Dist, typename T, typename Train, typename Mask>
void radiusMatchCached_caller(const DevMem2D_<T>& query, const DevMem2D_<T>& train, float maxDistance, const Mask& mask, void radiusMatchCached_caller(const DevMem2D_<T>& query, const Train& train, float maxDistance, const Mask& mask,
const DevMem2Di& trainIdx, const DevMem2Df& distance, unsigned int* nMatches, const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, int* nMatches,
cudaStream_t stream) cudaStream_t stream)
{ {
StaticAssert<BLOCK_STACK >= BLOCK_DIM_Y>::check(); StaticAssert<BLOCK_STACK >= BLOCK_DIM_Y>::check();
@ -167,7 +283,7 @@ namespace cv { namespace gpu { namespace bf_radius_match
const dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1); const dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);
radiusMatch<BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_STACK, VecDiffCachedRegister<BLOCK_DIM_X, MAX_LEN, LEN_EQ_MAX_LEN, typename Dist::value_type>, Dist, T> radiusMatch<BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_STACK, VecDiffCachedRegister<BLOCK_DIM_X, MAX_LEN, LEN_EQ_MAX_LEN, typename Dist::value_type>, Dist, T>
<<<grid, threads, 0, stream>>>(query, train, maxDistance, mask, trainIdx, distance, nMatches); <<<grid, threads, 0, stream>>>(query, train, maxDistance, mask, trainIdx, imgIdx, distance, nMatches, trainIdx.cols);
cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaGetLastError() );
if (stream == 0) if (stream == 0)
@ -177,58 +293,58 @@ namespace cv { namespace gpu { namespace bf_radius_match
/////////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////////////////////
// Radius Match Dispatcher // Radius Match Dispatcher
template <typename Dist, typename T, typename Mask> template <typename Dist, typename T, typename Train, typename Mask>
void radiusMatchDispatcher(const DevMem2D_<T>& query, const DevMem2D_<T>& train, float maxDistance, const Mask& mask, void radiusMatchDispatcher(const DevMem2D_<T>& query, const Train& train, float maxDistance, const Mask& mask,
const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream) cudaStream_t stream)
{ {
if (query.cols < 64) if (query.cols < 64)
{ {
radiusMatchCached_caller<16, 16, 64, 64, false, Dist>( radiusMatchCached_caller<16, 16, 64, 64, false, Dist>(
query, train, maxDistance, mask, query, train, maxDistance, mask,
static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Df>(distance), (unsigned int*)nMatches.data, static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Di>(imgIdx), static_cast<DevMem2Df>(distance), (int*)nMatches.data,
stream); stream);
} }
else if (query.cols == 64) else if (query.cols == 64)
{ {
radiusMatchCached_caller<16, 16, 64, 64, true, Dist>( radiusMatchCached_caller<16, 16, 64, 64, true, Dist>(
query, train, maxDistance, mask, query, train, maxDistance, mask,
static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Df>(distance), (unsigned int*)nMatches.data, static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Di>(imgIdx), static_cast<DevMem2Df>(distance), (int*)nMatches.data,
stream); stream);
} }
else if (query.cols < 128) else if (query.cols < 128)
{ {
radiusMatchCached_caller<16, 16, 64, 128, false, Dist>( radiusMatchCached_caller<16, 16, 64, 128, false, Dist>(
query, train, maxDistance, mask, query, train, maxDistance, mask,
static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Df>(distance), (unsigned int*)nMatches.data, static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Di>(imgIdx), static_cast<DevMem2Df>(distance), (int*)nMatches.data,
stream); stream);
} }
else if (query.cols == 128) else if (query.cols == 128)
{ {
radiusMatchCached_caller<16, 16, 64, 128, true, Dist>( radiusMatchCached_caller<16, 16, 64, 128, true, Dist>(
query, train, maxDistance, mask, query, train, maxDistance, mask,
static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Df>(distance), (unsigned int*)nMatches.data, static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Di>(imgIdx), static_cast<DevMem2Df>(distance), (int*)nMatches.data,
stream); stream);
} }
else if (query.cols < 256) else if (query.cols < 256)
{ {
radiusMatchCached_caller<16, 16, 64, 256, false, Dist>( radiusMatchCached_caller<16, 16, 64, 256, false, Dist>(
query, train, maxDistance, mask, query, train, maxDistance, mask,
static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Df>(distance), (unsigned int*)nMatches.data, static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Di>(imgIdx), static_cast<DevMem2Df>(distance), (int*)nMatches.data,
stream); stream);
} }
else if (query.cols == 256) else if (query.cols == 256)
{ {
radiusMatchCached_caller<16, 16, 64, 256, true, Dist>( radiusMatchCached_caller<16, 16, 64, 256, true, Dist>(
query, train, maxDistance, mask, query, train, maxDistance, mask,
static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Df>(distance), (unsigned int*)nMatches.data, static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Di>(imgIdx), static_cast<DevMem2Df>(distance), (int*)nMatches.data,
stream); stream);
} }
else else
{ {
radiusMatchSimple_caller<16, 16, 64, Dist>( radiusMatchSimple_caller<16, 16, 64, Dist>(
query, train, maxDistance, mask, query, train, maxDistance, mask,
static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Df>(distance), (unsigned int*)nMatches.data, static_cast<DevMem2Di>(trainIdx), static_cast<DevMem2Di>(imgIdx), static_cast<DevMem2Df>(distance), (int*)nMatches.data,
stream); stream);
} }
} }
@ -236,77 +352,163 @@ namespace cv { namespace gpu { namespace bf_radius_match
/////////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////////////////////
// Radius Match caller // Radius Match caller
template <typename T> void radiusMatchL1_gpu(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask, template <typename T> void radiusMatchSingleL1_gpu(const DevMem2D& query, const DevMem2D& train_, float maxDistance, const DevMem2D& mask,
const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream) cudaStream_t stream)
{ {
SingleTrain<T> train(static_cast< DevMem2D_<T> >(train_));
if (mask.data) if (mask.data)
{ {
radiusMatchDispatcher< L1Dist<T> >(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), maxDistance, SingleMask(mask), radiusMatchDispatcher< L1Dist<T> >(static_cast< DevMem2D_<T> >(query), train, maxDistance, SingleMask(mask),
trainIdx, distance, nMatches, trainIdx, DevMem2D(), distance, nMatches,
stream); stream);
} }
else else
{ {
radiusMatchDispatcher< L1Dist<T> >(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), maxDistance, WithOutMask(), radiusMatchDispatcher< L1Dist<T> >(static_cast< DevMem2D_<T> >(query), train, maxDistance, WithOutMask(),
trainIdx, distance, nMatches, trainIdx, DevMem2D(), distance, nMatches,
stream); stream);
} }
} }
template void radiusMatchL1_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); template void radiusMatchSingleL1_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchL1_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); //template void radiusMatchSingleL1_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchL1_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); template void radiusMatchSingleL1_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchL1_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); template void radiusMatchSingleL1_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchL1_gpu<int >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); template void radiusMatchSingleL1_gpu<int >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchL1_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); template void radiusMatchSingleL1_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template <typename T> void radiusMatchL2_gpu(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask, template <typename T> void radiusMatchSingleL2_gpu(const DevMem2D& query, const DevMem2D& train_, float maxDistance, const DevMem2D& mask,
const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream) cudaStream_t stream)
{ {
SingleTrain<T> train(static_cast< DevMem2D_<T> >(train_));
if (mask.data) if (mask.data)
{ {
radiusMatchDispatcher<L2Dist>(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), maxDistance, SingleMask(mask), radiusMatchDispatcher<L2Dist>(static_cast< DevMem2D_<T> >(query), train, maxDistance, SingleMask(mask),
trainIdx, distance, nMatches, trainIdx, DevMem2D(), distance, nMatches,
stream); stream);
} }
else else
{ {
radiusMatchDispatcher<L2Dist>(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), maxDistance, WithOutMask(), radiusMatchDispatcher<L2Dist>(static_cast< DevMem2D_<T> >(query), train, maxDistance, WithOutMask(),
trainIdx, distance, nMatches, trainIdx, DevMem2D(), distance, nMatches,
stream); stream);
} }
} }
//template void radiusMatchL2_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); //template void radiusMatchSingleL2_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchL2_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); //template void radiusMatchSingleL2_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchL2_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); //template void radiusMatchSingleL2_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchL2_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); //template void radiusMatchSingleL2_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchL2_gpu<int >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); //template void radiusMatchSingleL2_gpu<int >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchL2_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); template void radiusMatchSingleL2_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template <typename T> void radiusMatchHamming_gpu(const DevMem2D& query, const DevMem2D& train, float maxDistance, const DevMem2D& mask, template <typename T> void radiusMatchSingleHamming_gpu(const DevMem2D& query, const DevMem2D& train_, float maxDistance, const DevMem2D& mask,
const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream) cudaStream_t stream)
{ {
SingleTrain<T> train(static_cast< DevMem2D_<T> >(train_));
if (mask.data) if (mask.data)
{ {
radiusMatchDispatcher<HammingDist>(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), maxDistance, SingleMask(mask), radiusMatchDispatcher<HammingDist>(static_cast< DevMem2D_<T> >(query), train, maxDistance, SingleMask(mask),
trainIdx, distance, nMatches, trainIdx, DevMem2D(), distance, nMatches,
stream); stream);
} }
else else
{ {
radiusMatchDispatcher<HammingDist>(static_cast< DevMem2D_<T> >(query), static_cast< DevMem2D_<T> >(train), maxDistance, WithOutMask(), radiusMatchDispatcher<HammingDist>(static_cast< DevMem2D_<T> >(query), train, maxDistance, WithOutMask(),
trainIdx, distance, nMatches, trainIdx, DevMem2D(), distance, nMatches,
stream); stream);
} }
} }
template void radiusMatchHamming_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); template void radiusMatchSingleHamming_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchHamming_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); //template void radiusMatchSingleHamming_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchHamming_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); template void radiusMatchSingleHamming_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchHamming_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); //template void radiusMatchSingleHamming_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchHamming_gpu<int >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& nMatches, const DevMem2D& distance, cudaStream_t stream); template void radiusMatchSingleHamming_gpu<int >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance, const DevMem2D& mask, const DevMem2D& trainIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template <typename T> void radiusMatchCollectionL1_gpu(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection,
const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream)
{
TrainCollection<T> train((DevMem2D_<T>*)trainCollection.ptr(), trainCollection.cols, query.cols);
if (maskCollection.data)
{
radiusMatchDispatcher< L1Dist<T> >(static_cast< DevMem2D_<T> >(query), train, maxDistance, MaskCollection(maskCollection.data),
trainIdx, imgIdx, distance, nMatches,
stream);
}
else
{
radiusMatchDispatcher< L1Dist<T> >(static_cast< DevMem2D_<T> >(query), train, maxDistance, WithOutMask(),
trainIdx, imgIdx, distance, nMatches,
stream);
}
}
template void radiusMatchCollectionL1_gpu<uchar >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchCollectionL1_gpu<schar >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchCollectionL1_gpu<ushort>(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchCollectionL1_gpu<short >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchCollectionL1_gpu<int >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchCollectionL1_gpu<float >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template <typename T> void radiusMatchCollectionL2_gpu(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection,
const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream)
{
TrainCollection<T> train((DevMem2D_<T>*)trainCollection.ptr(), trainCollection.cols, query.cols);
if (maskCollection.data)
{
radiusMatchDispatcher<L2Dist>(static_cast< DevMem2D_<T> >(query), train, maxDistance, MaskCollection(maskCollection.data),
trainIdx, imgIdx, distance, nMatches,
stream);
}
else
{
radiusMatchDispatcher<L2Dist>(static_cast< DevMem2D_<T> >(query), train, maxDistance, WithOutMask(),
trainIdx, imgIdx, distance, nMatches,
stream);
}
}
//template void radiusMatchCollectionL2_gpu<uchar >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchCollectionL2_gpu<schar >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchCollectionL2_gpu<ushort>(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchCollectionL2_gpu<short >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchCollectionL2_gpu<int >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchCollectionL2_gpu<float >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template <typename T> void radiusMatchCollectionHamming_gpu(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection,
const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches,
cudaStream_t stream)
{
TrainCollection<T> train((DevMem2D_<T>*)trainCollection.ptr(), trainCollection.cols, query.cols);
if (maskCollection.data)
{
radiusMatchDispatcher<HammingDist>(static_cast< DevMem2D_<T> >(query), train, maxDistance, MaskCollection(maskCollection.data),
trainIdx, imgIdx, distance, nMatches,
stream);
}
else
{
radiusMatchDispatcher<HammingDist>(static_cast< DevMem2D_<T> >(query), train, maxDistance, WithOutMask(),
trainIdx, imgIdx, distance, nMatches,
stream);
}
}
template void radiusMatchCollectionHamming_gpu<uchar >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchCollectionHamming_gpu<schar >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchCollectionHamming_gpu<ushort>(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
//template void radiusMatchCollectionHamming_gpu<short >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
template void radiusMatchCollectionHamming_gpu<int >(const DevMem2D& query, const DevMem2D& trainCollection, float maxDistance, const DevMem2D_<PtrStep>& maskCollection, const DevMem2D& trainIdx, const DevMem2D& imgIdx, const DevMem2D& distance, const DevMem2D& nMatches, cudaStream_t stream);
}}} }}}