mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-24 19:20:28 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Merge pull request #1687 from bitwangyaoyao:2.4_fix

Browse Source
This commit is contained in:
Andrey Pavlenko 2013-10-29 11:03:33 +04:00 committed by OpenCV Buildbot
commit ea64589164
2 changed files with 37 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
modules/ocl/perf/perf_brute_force_matcher.cpp
View File

@ -53,8 +53,8 @@ using namespace perf;
typedef TestBaseWithParam<Size> BruteForceMatcherFixture;
PERF_TEST_P(BruteForceMatcherFixture, DISABLED_match,
OCL_BFMATCHER_TYPICAL_MAT_SIZES) // TODO too big difference between implementations
PERF_TEST_P(BruteForceMatcherFixture, match,
OCL_BFMATCHER_TYPICAL_MAT_SIZES)
{
const Size srcSize = GetParam();
@ -82,14 +82,14 @@ PERF_TEST_P(BruteForceMatcherFixture, DISABLED_match,
oclMatcher.matchDownload(oclTrainIdx, oclDistance, matches);
SANITY_CHECK_MATCHES(matches);
SANITY_CHECK_MATCHES(matches, 1e-5);
}
else
OCL_PERF_ELSE
}
PERF_TEST_P(BruteForceMatcherFixture, DISABLED_knnMatch,
OCL_BFMATCHER_TYPICAL_MAT_SIZES) // TODO too big difference between implementations
PERF_TEST_P(BruteForceMatcherFixture, knnMatch,
OCL_BFMATCHER_TYPICAL_MAT_SIZES)
{
const Size srcSize = GetParam();
@ -123,8 +123,8 @@ PERF_TEST_P(BruteForceMatcherFixture, DISABLED_knnMatch,
oclMatcher.knnMatchDownload(oclTrainIdx, oclDistance, matches);
std::vector<DMatch> & matches0 = matches[0], & matches1 = matches[1];
SANITY_CHECK_MATCHES(matches0);
SANITY_CHECK_MATCHES(matches1);
SANITY_CHECK_MATCHES(matches0, 1e-5);
SANITY_CHECK_MATCHES(matches1, 1e-5);
}
else
OCL_PERF_ELSE

33
modules/ocl/src/opencl/brute_force_match.cl
View File

@ -17,6 +17,7 @@
// @Authors
// Nathan, liujun@multicorewareinc.com
// Peng Xiao, pengxiao@outlook.com
// Baichuan Su, baichuan@multicorewareinc.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
@ -112,6 +113,24 @@ result_type reduce_block(
return DIST_RES(result);
}
result_type reduce_block_match(
__local value_type *s_query,
__local value_type *s_train,
int lidx,
int lidy
)
{
result_type result = 0;
#pragma unroll
for (int j = 0 ; j < BLOCK_SIZE ; j++)
{
result += DIST(
s_query[lidy * BLOCK_SIZE + j],
s_train[j * BLOCK_SIZE + lidx]);
}
return (result);
}
result_type reduce_multi_block(
__local value_type *s_query,
__local value_type *s_train,
@ -128,7 +147,7 @@ result_type reduce_multi_block(
s_query[lidy * MAX_DESC_LEN + block_index * BLOCK_SIZE + j],
s_train[j * BLOCK_SIZE + lidx]);
}
return DIST_RES(result);
return result;
}
/* 2dim launch, global size: dim0 is (query rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, dim1 is BLOCK_SIZE
@ -187,6 +206,8 @@ __kernel void BruteForceMatch_UnrollMatch(
barrier(CLK_LOCAL_MEM_FENCE);
}
result = DIST_RES(result);
int trainIdx = t * BLOCK_SIZE + lidx;
if (queryIdx < query_rows && trainIdx < train_rows && result < myBestDistance/* && mask(queryIdx, trainIdx)*/)
@ -272,11 +293,13 @@ __kernel void BruteForceMatch_Match(
barrier(CLK_LOCAL_MEM_FENCE);
result += reduce_block(s_query, s_train, lidx, lidy);
result += reduce_block_match(s_query, s_train, lidx, lidy);
barrier(CLK_LOCAL_MEM_FENCE);
}
result = DIST_RES(result);
const int trainIdx = t * BLOCK_SIZE + lidx;
if (queryIdx < query_rows && trainIdx < train_rows && result < myBestDistance /*&& mask(queryIdx, trainIdx)*/)
@ -493,6 +516,8 @@ __kernel void BruteForceMatch_knnUnrollMatch(
barrier(CLK_LOCAL_MEM_FENCE);
}
result = DIST_RES(result);
const int trainIdx = t * BLOCK_SIZE + lidx;
if (queryIdx < query_rows && trainIdx < train_rows)
@ -631,11 +656,13 @@ __kernel void BruteForceMatch_knnMatch(
barrier(CLK_LOCAL_MEM_FENCE);
result += reduce_block(s_query, s_train, lidx, lidy);
result += reduce_block_match(s_query, s_train, lidx, lidy);
barrier(CLK_LOCAL_MEM_FENCE);
}
result = DIST_RES(result);
const int trainIdx = t * BLOCK_SIZE + lidx;
if (queryIdx < query_rows && trainIdx < train_rows /*&& mask(queryIdx, trainIdx)*/)