Allow more input query/train types for ocl::bfmatcher

RadiusMatch for HammingDist cannot pass yet.

diff --git a/modules/ocl/src/brute_force_matcher.cpp b/modules/ocl/src/brute_force_matcher.cpp
index b883a1b..9c4a217 100644 (file)
--- a/modules/ocl/src/brute_force_matcher.cpp
+++ b/modules/ocl/src/brute_force_matcher.cpp
@@ -77,7 +77,6 @@ template < int BLOCK_SIZE, int MAX_DESC_LEN/*, typename Mask*/ >
 void matchUnrolledCached(const oclMat &query, const oclMat &train, const oclMat &/*mask*/,
                          const oclMat &trainIdx, const oclMat &distance, int distType)
 {
-    assert(query.type() == CV_32F);
     cv::ocl::Context *ctx = query.clCxt;
     size_t globalSize[] = {(query.rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, BLOCK_SIZE, 1};
     size_t localSize[] = {BLOCK_SIZE, BLOCK_SIZE, 1};
@@ -121,7 +120,6 @@ template < int BLOCK_SIZE/*, typename Mask*/ >
 void match(const oclMat &query, const oclMat &train, const oclMat &/*mask*/,
            const oclMat &trainIdx, const oclMat &distance, int distType)
 {
-    assert(query.type() == CV_32F);
     cv::ocl::Context *ctx = query.clCxt;
     size_t globalSize[] = {(query.rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, BLOCK_SIZE, 1};
     size_t localSize[] = {BLOCK_SIZE, BLOCK_SIZE, 1};
@@ -164,7 +162,6 @@ template < int BLOCK_SIZE, int MAX_DESC_LEN/*, typename Mask*/ >
 void matchUnrolledCached(const oclMat &query, const oclMat &train, float maxDistance, const oclMat &/*mask*/,
                          const oclMat &trainIdx, const oclMat &distance, const oclMat &nMatches, int distType)
 {
-    assert(query.type() == CV_32F);
     cv::ocl::Context *ctx = query.clCxt;
     size_t globalSize[] = {(train.rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, (query.rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, 1};
     size_t localSize[] = {BLOCK_SIZE, BLOCK_SIZE, 1};
@@ -207,7 +204,6 @@ template < int BLOCK_SIZE/*, typename Mask*/ >
 void radius_match(const oclMat &query, const oclMat &train, float maxDistance, const oclMat &/*mask*/,
                   const oclMat &trainIdx, const oclMat &distance, const oclMat &nMatches, int distType)
 {
-    assert(query.type() == CV_32F);
     cv::ocl::Context *ctx = query.clCxt;
     size_t globalSize[] = {(train.rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, (query.rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, 1};
     size_t localSize[] = {BLOCK_SIZE, BLOCK_SIZE, 1};
@@ -566,17 +562,6 @@ void cv::ocl::BruteForceMatcher_OCL_base::matchSingle(const oclMat &query, const
     if (query.empty() || train.empty())
         return;
     
-    // match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int
-    int callType = query.depth();
-    if (callType != 5)
-        CV_Error(CV_UNSUPPORTED_FORMAT_ERR, "BruteForceMatch OpenCL only support float type query!\n");
-
-    if ((distType == 0 && callType == 1 ) || (distType == 1 && callType != 5) || (distType == 2 && (callType != 0
-        || callType != 2 || callType != 4)))
-    {
-        CV_Error(CV_UNSUPPORTED_DEPTH_ERR, "BruteForceMatch OpenCL only support float type query!\n");
-    }
-
     CV_Assert(query.channels() == 1 && query.depth() < CV_64F);
     CV_Assert(train.cols == query.cols && train.type() == query.type());
 
@@ -687,17 +672,6 @@ void cv::ocl::BruteForceMatcher_OCL_base::matchCollection(const oclMat &query, c
     if (query.empty() || trainCollection.empty())
         return;
 
-    // match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int
-    int callType = query.depth();
-    if (callType != 5)
-        CV_Error(CV_UNSUPPORTED_FORMAT_ERR, "BruteForceMatch OpenCL only support float type query!\n");
-
-    if ((distType == 0 && callType == 1 ) || (distType == 1 && callType != 5) || (distType == 2 && (callType != 0
-        || callType != 2 || callType != 4)))
-    {
-        CV_Error(CV_UNSUPPORTED_DEPTH_ERR, "BruteForceMatch OpenCL only support float type query!\n");
-    }
-
     CV_Assert(query.channels() == 1 && query.depth() < CV_64F);
     
     const int nQuery = query.rows;
@@ -706,7 +680,6 @@ void cv::ocl::BruteForceMatcher_OCL_base::matchCollection(const oclMat &query, c
     ensureSizeIsEnough(1, nQuery, CV_32S, imgIdx);
     ensureSizeIsEnough(1, nQuery, CV_32F, distance);
 
-
     matchDispatcher(query, (const oclMat *)trainCollection.ptr(), trainCollection.cols, masks, trainIdx, imgIdx, distance, distType);
 
     return;
@@ -778,18 +751,6 @@ void cv::ocl::BruteForceMatcher_OCL_base::knnMatchSingle(const oclMat &query, co
     if (query.empty() || train.empty())
         return;
 
-    // match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int
-    int callType = query.depth();
-
-    if (callType != 5)
-        CV_Error(CV_UNSUPPORTED_FORMAT_ERR, "BruteForceMatch OpenCL only support float type query!\n");
-
-    if ((distType == 0 && callType == 1 ) || (distType == 1 && callType != 5) || (distType == 2 && (callType != 0
-        || callType != 2 || callType != 4)))
-    {
-        CV_Error(CV_UNSUPPORTED_DEPTH_ERR, "BruteForceMatch OpenCL only support float type query!\n");
-    }
-
     CV_Assert(query.channels() == 1 && query.depth() < CV_64F);
     CV_Assert(train.type() == query.type() && train.cols == query.cols);
 
@@ -886,26 +847,7 @@ void cv::ocl::BruteForceMatcher_OCL_base::knnMatch2Collection(const oclMat &quer
 
     typedef void (*caller_t)(const oclMat & query, const oclMat & trains, const oclMat & masks,
                              const oclMat & trainIdx, const oclMat & imgIdx, const oclMat & distance);
-#if 0
-    static const caller_t callers[3][6] =
-    {
-        {
-            ocl_match2L1_gpu<unsigned char>, 0/*match2L1_gpu<signed char>*/,
-            ocl_match2L1_gpu<unsigned short>, ocl_match2L1_gpu<short>,
-            ocl_match2L1_gpu<int>, ocl_match2L1_gpu<float>
-        },
-        {
-            0/*match2L2_gpu<unsigned char>*/, 0/*match2L2_gpu<signed char>*/,
-            0/*match2L2_gpu<unsigned short>*/, 0/*match2L2_gpu<short>*/,
-            0/*match2L2_gpu<int>*/, ocl_match2L2_gpu<float>
-        },
-        {
-            ocl_match2Hamming_gpu<unsigned char>, 0/*match2Hamming_gpu<signed char>*/,
-            ocl_match2Hamming_gpu<unsigned short>, 0/*match2Hamming_gpu<short>*/,
-            ocl_match2Hamming_gpu<int>, 0/*match2Hamming_gpu<float>*/
-        }
-    };
-#endif
+
     CV_Assert(query.channels() == 1 && query.depth() < CV_64F);
 
     const int nQuery = query.rows;
@@ -1051,23 +993,11 @@ void cv::ocl::BruteForceMatcher_OCL_base::knnMatch(const oclMat &query, vector<
 
 // radiusMatchSingle
 void cv::ocl::BruteForceMatcher_OCL_base::radiusMatchSingle(const oclMat &query, const oclMat &train,
-        oclMat &trainIdx,      oclMat &distance, oclMat &nMatches, float maxDistance, const oclMat &mask)
+        oclMat &trainIdx, oclMat &distance, oclMat &nMatches, float maxDistance, const oclMat &mask)
 {
     if (query.empty() || train.empty())
         return;
 
-    // match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int
-    int callType = query.depth();
-
-    if (callType != 5)
-        CV_Error(CV_UNSUPPORTED_FORMAT_ERR, "BruteForceMatch OpenCL only support float type query!\n");
-
-    if ((distType == 0 && callType == 1 ) || (distType == 1 && callType != 5) || (distType == 2 && (callType != 0
-        || callType != 2 || callType != 4)))
-    {
-        CV_Error(CV_UNSUPPORTED_DEPTH_ERR, "BruteForceMatch OpenCL only support float type query!\n");
-    }
-
     const int nQuery = query.rows;
     const int nTrain = train.rows;
 

diff --git a/modules/ocl/src/opencl/brute_force_match.cl b/modules/ocl/src/opencl/brute_force_match.cl
index 113f7d5..0812832 100644 (file)
--- a/modules/ocl/src/opencl/brute_force_match.cl
+++ b/modules/ocl/src/opencl/brute_force_match.cl
@@ -85,14 +85,17 @@ int bit1Count(int x)
         typedef int value_type;
         typedef int result_type;
 #   endif
+#define DIST_RES(x) (x)
 #elif (DIST_TYPE == 1) // L2Dist
 #define DIST(x, y) (((x) - (y)) * ((x) - (y)))
 typedef float value_type;
 typedef float result_type;
+#define DIST_RES(x) sqrt(x)
 #elif (DIST_TYPE == 2) // Hamming
-#define DIST(x, y) bit1Count(((x) ^ (y))
+#define DIST(x, y) bit1Count( (x) ^ (y) )
 typedef int value_type;
 typedef int result_type;
+#define DIST_RES(x) (x)
 #endif
 
 result_type reduce_block(
@@ -107,10 +110,10 @@ result_type reduce_block(
     for (int j = 0 ; j < BLOCK_SIZE ; j++)
     {
         result += DIST(
-            s_query[lidy * BLOCK_SIZE + j], 
+            s_query[lidy * BLOCK_SIZE + j],
             s_train[j * BLOCK_SIZE + lidx]);
     }
-    return result;
+    return DIST_RES(result);
 }
 
 result_type reduce_multi_block(
@@ -126,10 +129,10 @@ result_type reduce_multi_block(
     for (int j = 0 ; j < BLOCK_SIZE ; j++)
     {
         result += DIST(
-            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]);
     }
-    return result;
+    return DIST_RES(result);
 }
 
 /* 2dim launch, global size: dim0 is (query rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, dim1 is BLOCK_SIZE
@@ -153,8 +156,8 @@ __kernel void BruteForceMatch_UnrollMatch(
     const int lidy = get_local_id(1);
     const int groupidx = get_group_id(0);
 
-    __local value_type *s_query = sharebuffer;
-    __local value_type *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
+    __local value_type *s_query = (__local value_type *)sharebuffer;
+    __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
 
     int queryIdx = groupidx * BLOCK_SIZE + lidy;
     // load the query into local memory.
@@ -251,8 +254,8 @@ __kernel void BruteForceMatch_Match(
     float myBestDistance = MAX_FLOAT;
     int myBestTrainIdx = -1;
 
-    __local value_type *s_query = sharebuffer;
-    __local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
+    __local value_type *s_query = (__local value_type *)sharebuffer;
+    __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
 
     // loop
     for (int t = 0 ;  t < (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE ; t++)
@@ -345,8 +348,8 @@ __kernel void BruteForceMatch_RadiusUnrollMatch(
     const int queryIdx = groupidy * BLOCK_SIZE + lidy;
     const int trainIdx = groupidx * BLOCK_SIZE + lidx;
 
-    __local value_type *s_query = sharebuffer;
-    __local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
+    __local value_type *s_query = (__local value_type *)sharebuffer;
+    __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
 
     result_type result = 0;
     for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; ++i)
@@ -365,7 +368,8 @@ __kernel void BruteForceMatch_RadiusUnrollMatch(
         barrier(CLK_LOCAL_MEM_FENCE);
     }
 
-    if (queryIdx < query_rows && trainIdx < train_rows && result < maxDistance/* && mask(queryIdx, trainIdx)*/)
+    if (queryIdx < query_rows && trainIdx < train_rows &&
+        convert_float(result) < maxDistance/* && mask(queryIdx, trainIdx)*/)
     {
         unsigned int ind = atom_inc(nMatches + queryIdx/*, (unsigned int) -1*/);
 
@@ -405,8 +409,8 @@ __kernel void BruteForceMatch_RadiusMatch(
     const int queryIdx = groupidy * BLOCK_SIZE + lidy;
     const int trainIdx = groupidx * BLOCK_SIZE + lidx;
 
-    __local value_type *s_query = sharebuffer;
-    __local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
+    __local value_type *s_query = (__local value_type *)sharebuffer;
+    __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
 
     result_type result = 0;
     for (int i = 0 ; i < (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE ; ++i)
@@ -425,7 +429,8 @@ __kernel void BruteForceMatch_RadiusMatch(
         barrier(CLK_LOCAL_MEM_FENCE);
     }
 
-    if (queryIdx < query_rows && trainIdx < train_rows && result < maxDistance/* && mask(queryIdx, trainIdx)*/)
+    if (queryIdx < query_rows && trainIdx < train_rows && 
+        convert_float(result) < maxDistance/* && mask(queryIdx, trainIdx)*/)
     {
         unsigned int ind = atom_inc(nMatches + queryIdx);
 
@@ -458,8 +463,8 @@ __kernel void BruteForceMatch_knnUnrollMatch(
     const int groupidx = get_group_id(0);
 
     const int queryIdx = groupidx * BLOCK_SIZE + lidy;
-    local value_type *s_query = sharebuffer;
-    local value_type *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
+    __local value_type *s_query = (__local value_type *)sharebuffer;
+    __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
 
     // load the query into local memory.
     for (int i = 0 ;  i <  MAX_DESC_LEN / BLOCK_SIZE; i ++)
@@ -480,7 +485,6 @@ __kernel void BruteForceMatch_knnUnrollMatch(
         result_type result = 0;
         for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; i++)
         {
-            const int loadX = lidx + i * BLOCK_SIZE;
             //load a BLOCK_SIZE * BLOCK_SIZE block into local train.
             const int loadx = lidx + i * BLOCK_SIZE;
             s_train[lidx * BLOCK_SIZE + lidy] = loadx < train_cols ? train[min(t * BLOCK_SIZE + lidy, train_rows - 1) * (step / sizeof(float)) + loadx] : 0;
@@ -514,8 +518,8 @@ __kernel void BruteForceMatch_knnUnrollMatch(
 
     barrier(CLK_LOCAL_MEM_FENCE);
 
-    local float *s_distance = (local float *)sharebuffer;
-    local int *s_trainIdx = (local int *)(sharebuffer + BLOCK_SIZE * BLOCK_SIZE);
+    __local float *s_distance = (local float *)sharebuffer;
+    __local int *s_trainIdx = (local int *)(sharebuffer + BLOCK_SIZE * BLOCK_SIZE);
 
     // find BestMatch
     s_distance += lidy * BLOCK_SIZE;
@@ -604,8 +608,8 @@ __kernel void BruteForceMatch_knnMatch(
     const int groupidx = get_group_id(0);
 
     const int queryIdx = groupidx * BLOCK_SIZE + lidy;
-    local value_type *s_query = sharebuffer;
-    local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
+    __local value_type *s_query = (__local value_type *)sharebuffer;
+    __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
 
     float myBestDistance1 = MAX_FLOAT;
     float myBestDistance2 = MAX_FLOAT;
@@ -766,4 +770,4 @@ kernel void BruteForceMatch_findBestMatch(
 )
 {
     /* Todo */
-}
\ No newline at end of file
+}

diff --git a/modules/ocl/test/test_brute_force_matcher.cpp b/modules/ocl/test/test_brute_force_matcher.cpp
index d13f03e..2031762 100644 (file)
--- a/modules/ocl/test/test_brute_force_matcher.cpp
+++ b/modules/ocl/test/test_brute_force_matcher.cpp
@@ -158,11 +158,7 @@ namespace
 
     TEST_P(BruteForceMatcher, RadiusMatch_Single)
     {
-        float radius;
-        if(distType == cv::ocl::BruteForceMatcher_OCL_base::L2Dist)
-            radius = 1.f / countFactor / countFactor;
-        else
-            radius = 1.f / countFactor;
+        float radius = 1.f / countFactor;
 
         cv::ocl::BruteForceMatcher_OCL_base matcher(distType);