Commit dd0fa63c authored by yao's avatar yao

fix the bug of ocl::bruteForceMatcher

parent 3b08bf6b
...@@ -53,8 +53,8 @@ using namespace perf; ...@@ -53,8 +53,8 @@ using namespace perf;
typedef TestBaseWithParam<Size> BruteForceMatcherFixture; typedef TestBaseWithParam<Size> BruteForceMatcherFixture;
PERF_TEST_P(BruteForceMatcherFixture, DISABLED_match, PERF_TEST_P(BruteForceMatcherFixture, match,
OCL_BFMATCHER_TYPICAL_MAT_SIZES) // TODO too big difference between implementations OCL_BFMATCHER_TYPICAL_MAT_SIZES)
{ {
const Size srcSize = GetParam(); const Size srcSize = GetParam();
...@@ -82,14 +82,14 @@ PERF_TEST_P(BruteForceMatcherFixture, DISABLED_match, ...@@ -82,14 +82,14 @@ PERF_TEST_P(BruteForceMatcherFixture, DISABLED_match,
oclMatcher.matchDownload(oclTrainIdx, oclDistance, matches); oclMatcher.matchDownload(oclTrainIdx, oclDistance, matches);
SANITY_CHECK_MATCHES(matches); SANITY_CHECK_MATCHES(matches, 1e-5);
} }
else else
OCL_PERF_ELSE OCL_PERF_ELSE
} }
PERF_TEST_P(BruteForceMatcherFixture, DISABLED_knnMatch, PERF_TEST_P(BruteForceMatcherFixture, knnMatch,
OCL_BFMATCHER_TYPICAL_MAT_SIZES) // TODO too big difference between implementations OCL_BFMATCHER_TYPICAL_MAT_SIZES)
{ {
const Size srcSize = GetParam(); const Size srcSize = GetParam();
...@@ -123,8 +123,8 @@ PERF_TEST_P(BruteForceMatcherFixture, DISABLED_knnMatch, ...@@ -123,8 +123,8 @@ PERF_TEST_P(BruteForceMatcherFixture, DISABLED_knnMatch,
oclMatcher.knnMatchDownload(oclTrainIdx, oclDistance, matches); oclMatcher.knnMatchDownload(oclTrainIdx, oclDistance, matches);
std::vector<DMatch> & matches0 = matches[0], & matches1 = matches[1]; std::vector<DMatch> & matches0 = matches[0], & matches1 = matches[1];
SANITY_CHECK_MATCHES(matches0); SANITY_CHECK_MATCHES(matches0, 1e-5);
SANITY_CHECK_MATCHES(matches1); SANITY_CHECK_MATCHES(matches1, 1e-5);
} }
else else
OCL_PERF_ELSE OCL_PERF_ELSE
......
...@@ -17,6 +17,7 @@ ...@@ -17,6 +17,7 @@
// @Authors // @Authors
// Nathan, liujun@multicorewareinc.com // Nathan, liujun@multicorewareinc.com
// Peng Xiao, pengxiao@outlook.com // Peng Xiao, pengxiao@outlook.com
// Baichuan Su, baichuan@multicorewareinc.com
// //
// Redistribution and use in source and binary forms, with or without modification, // Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met: // are permitted provided that the following conditions are met:
...@@ -128,7 +129,7 @@ result_type reduce_multi_block( ...@@ -128,7 +129,7 @@ result_type reduce_multi_block(
s_query[lidy * MAX_DESC_LEN + block_index * BLOCK_SIZE + j], s_query[lidy * MAX_DESC_LEN + block_index * BLOCK_SIZE + j],
s_train[j * BLOCK_SIZE + lidx]); 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 /* 2dim launch, global size: dim0 is (query rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, dim1 is BLOCK_SIZE
...@@ -187,6 +188,8 @@ __kernel void BruteForceMatch_UnrollMatch( ...@@ -187,6 +188,8 @@ __kernel void BruteForceMatch_UnrollMatch(
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
result = DIST_RES(result);
int trainIdx = t * BLOCK_SIZE + lidx; int trainIdx = t * BLOCK_SIZE + lidx;
if (queryIdx < query_rows && trainIdx < train_rows && result < myBestDistance/* && mask(queryIdx, trainIdx)*/) if (queryIdx < query_rows && trainIdx < train_rows && result < myBestDistance/* && mask(queryIdx, trainIdx)*/)
...@@ -493,6 +496,8 @@ __kernel void BruteForceMatch_knnUnrollMatch( ...@@ -493,6 +496,8 @@ __kernel void BruteForceMatch_knnUnrollMatch(
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
result = DIST_RES(result);
const int trainIdx = t * BLOCK_SIZE + lidx; const int trainIdx = t * BLOCK_SIZE + lidx;
if (queryIdx < query_rows && trainIdx < train_rows) if (queryIdx < query_rows && trainIdx < train_rows)
......
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