Merge branch 4.x

modules/features/perf/opencl/perf_brute_force_matcher.cpp

@@ -123,6 +123,27 @@ OCL_PERF_TEST_P(BruteForceMatcherFixture, RadiusMatch, ::testing::Combine(OCL_PE
     SANITY_CHECK_MATCHES(matches1, 1e-3);
 }
 
+OCL_PERF_TEST_P(BruteForceMatcherFixture, MatchCrossCheck, ::testing::Combine(OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), OCL_PERF_ENUM((MatType)CV_32FC1) ) )
+{
+    const Size_MatType_t params = GetParam();
+    const Size srcSize = get<0>(params);
+    const int type = get<1>(params);
+
+    checkDeviceMaxMemoryAllocSize(srcSize, type);
+
+    vector<DMatch> matches;
+    UMat uquery(srcSize, type), utrain(srcSize, type);
+
+    declare.in(uquery, utrain, WARMUP_RNG);
+
+    BFMatcher matcher(NORM_L2, true /*crossCheck*/);
+
+    OCL_TEST_CYCLE()
+        matcher.match(uquery, utrain, matches);
+
+    SANITY_CHECK_MATCHES(matches, 1e-3);
+}
+
 } // ocl
 } // cvtest
 

modules/features/src/matchers.cpp

@@ -74,7 +74,7 @@ static void ensureSizeIsEnough(int rows, int cols, int type, UMat &m)
 }
 
 static bool ocl_matchSingle(InputArray query, InputArray train,
-        UMat &trainIdx, UMat &distance, int distType)
+        UMat &trainIdx, UMat *distance, int distType)
 {
     if (query.empty() || train.empty())
         return false;
@@ -83,7 +83,8 @@ static bool ocl_matchSingle(InputArray query, InputArray train,
     const int query_cols = query.cols();
 
     ensureSizeIsEnough(1, query_rows, CV_32S, trainIdx);
-    ensureSizeIsEnough(1, query_rows, CV_32F, distance);
+    if (distance)
+        ensureSizeIsEnough(1, query_rows, CV_32F, *distance);
 
     ocl::Device devDef = ocl::Device::getDefault();
 
@@ -117,7 +118,10 @@ static bool ocl_matchSingle(InputArray query, InputArray train,
     idx = k.set(idx, ocl::KernelArg::PtrReadOnly(uquery));
     idx = k.set(idx, ocl::KernelArg::PtrReadOnly(utrain));
     idx = k.set(idx, ocl::KernelArg::PtrWriteOnly(trainIdx));
-    idx = k.set(idx, ocl::KernelArg::PtrWriteOnly(distance));
+    UMat dummyDist;
+    if (!distance)
+        ensureSizeIsEnough(1, query_rows, CV_32F, dummyDist);
+    idx = k.set(idx, ocl::KernelArg::PtrWriteOnly(distance ? *distance : dummyDist));
     idx = k.set(idx, uquery.rows);
     idx = k.set(idx, uquery.cols);
     idx = k.set(idx, utrain.rows);
@@ -734,7 +738,7 @@ Ptr<DescriptorMatcher> BFMatcher::clone( bool emptyTrainData ) const
 static bool ocl_match(InputArray query, InputArray _train, std::vector< std::vector<DMatch> > &matches, int dstType)
 {
     UMat trainIdx, distance;
-    if (!ocl_matchSingle(query, _train, trainIdx, distance, dstType))
+    if (!ocl_matchSingle(query, _train, trainIdx, &distance, dstType))
         return false;
     if (!ocl_matchDownload(trainIdx, distance, matches))
         return false;
@@ -752,6 +756,146 @@ static bool ocl_knnMatch(InputArray query, InputArray _train, std::vector< std::
         return false;
     return true;
 }
+
+static bool ocl_matchWithCrossCheckSinglePass(InputArray query, InputArray train,
+        std::vector< std::vector<DMatch> >& matches, int dstType)
+{
+    if (query.empty() || train.empty())
+        return false;
+
+    const int query_rows = query.rows();
+    const int train_rows = train.rows();
+
+    UMat fwdIdx, fwdDist;
+    ensureSizeIsEnough(1, query_rows, CV_32S, fwdIdx);
+    ensureSizeIsEnough(1, query_rows, CV_32F, fwdDist);
+
+    UMat revBest;
+    ensureSizeIsEnough(1, train_rows, CV_32SC2, revBest);
+    revBest.setTo(Scalar::all(-1));
+
+    ocl::Device devDef = ocl::Device::getDefault();
+
+    UMat uquery = query.getUMat(), utrain = train.getUMat();
+    int kercn = 1;
+    if (devDef.isIntel() &&
+        (0 == (uquery.step % 4)) && (0 == (uquery.cols % 4)) && (0 == (uquery.offset % 4)) &&
+        (0 == (utrain.step % 4)) && (0 == (utrain.cols % 4)) && (0 == (utrain.offset % 4)))
+        kercn = 4;
+
+    int block_size = 16;
+    int max_desc_len = 0;
+    bool is_cpu = devDef.type() == ocl::Device::TYPE_CPU;
+    if (uquery.cols <= 64)
+        max_desc_len = 64 / kercn;
+    else if (uquery.cols <= 128 && !is_cpu)
+        max_desc_len = 128 / kercn;
+
+    int depth = uquery.depth();
+    cv::String opts;
+    opts = cv::format("-D T=%s -D TN=%s -D kercn=%d %s -D DIST_TYPE=%d -D BLOCK_SIZE=%d -D MAX_DESC_LEN=%d -D HAVE_INT64_ATOMICS",
+        ocl::typeToStr(depth), ocl::typeToStr(CV_MAKETYPE(depth, kercn)), kercn, depth == CV_32F ? "-D T_FLOAT" : "", dstType, block_size, max_desc_len);
+    ocl::Kernel k("BruteForceMatch_CrossCheckMatch", ocl::features::brute_force_match_oclsrc, opts);
+    if(k.empty())
+        return false;
+
+    size_t globalSize[] = {((size_t)uquery.size().height + block_size - 1) / block_size * block_size, (size_t)block_size};
+    size_t localSize[] = {(size_t)block_size, (size_t)block_size};
+
+    int idx = 0;
+    idx = k.set(idx, ocl::KernelArg::PtrReadOnly(uquery));
+    idx = k.set(idx, ocl::KernelArg::PtrReadOnly(utrain));
+    idx = k.set(idx, ocl::KernelArg::PtrWriteOnly(fwdIdx));
+    idx = k.set(idx, ocl::KernelArg::PtrWriteOnly(fwdDist));
+    idx = k.set(idx, ocl::KernelArg::PtrWriteOnly(revBest));
+    idx = k.set(idx, uquery.rows);
+    idx = k.set(idx, uquery.cols);
+    idx = k.set(idx, utrain.rows);
+    idx = k.set(idx, utrain.cols);
+    idx = k.set(idx, (int)(uquery.step / sizeof(float)));
+
+    if (!k.run(2, globalSize, localSize, false))
+        return false;
+
+    Mat fwdIdxCPU = fwdIdx.getMat(ACCESS_READ);
+    Mat fwdDistCPU = fwdDist.getMat(ACCESS_READ);
+    Mat revBestCPU = revBest.getMat(ACCESS_READ);
+
+    if (fwdIdxCPU.empty() || fwdDistCPU.empty() || revBestCPU.empty())
+        return false;
+
+    const int* fwd = fwdIdxCPU.ptr<int>();
+    const float* dist = fwdDistCPU.ptr<float>();
+    const int* revBestData = revBestCPU.ptr<int>();
+
+    matches.clear();
+    matches.reserve(query_rows);
+
+    for (int q = 0; q < query_rows; ++q)
+    {
+        int t = fwd[q];
+        if (t >= 0 && t < train_rows)
+        {
+            uint64_t packed_val;
+            memcpy(&packed_val, &revBestData[2 * t], sizeof(uint64_t));
+            int revQueryIdx = (int)(packed_val & 0xFFFFFFFF);
+            if (revQueryIdx == q)
+            {
+                matches.push_back(std::vector<DMatch>(1, DMatch(q, t, 0, dist[q])));
+            }
+        }
+    }
+    return true;
+}
+
+static bool ocl_matchWithCrossCheck(InputArray query, InputArray train,
+        std::vector< std::vector<DMatch> >& matches, int dstType)
+{
+    if (query.empty() || train.empty())
+        return false;
+
+    ocl::Device devDef = ocl::Device::getDefault();
+
+    if (devDef.isExtensionSupported("cl_khr_int64_base_atomics"))
+    {
+        if (ocl_matchWithCrossCheckSinglePass(query, train, matches, dstType))
+            return true;
+    }
+
+    UMat fwdIdx, fwdDist;
+    if (!ocl_matchSingle(query, train, fwdIdx, &fwdDist, dstType))
+        return false;
+
+    UMat revIdx;
+    if (!ocl_matchSingle(train, query, revIdx, nullptr, dstType))
+        return false;
+
+    Mat fwdIdxCPU = fwdIdx.getMat(ACCESS_READ);
+    Mat revIdxCPU = revIdx.getMat(ACCESS_READ);
+    Mat fwdDistCPU = fwdDist.getMat(ACCESS_READ);
+
+    if (fwdIdxCPU.empty() || revIdxCPU.empty() || fwdDistCPU.empty())
+        return false;
+
+    const int nQuery = fwdIdxCPU.cols;
+    const int nTrain = revIdxCPU.cols;
+    const int* fwd  = fwdIdxCPU.ptr<int>();
+    const int* rev  = revIdxCPU.ptr<int>();
+    const float* dist = fwdDistCPU.ptr<float>();
+
+    matches.clear();
+    matches.reserve(nQuery);
+
+    for (int q = 0; q < nQuery; ++q)
+    {
+        int t = fwd[q];
+        if (t >= 0 && t < nTrain && rev[t] == q)
+        {
+            matches.push_back(std::vector<DMatch>(1, DMatch(q, t, 0, dist[q])));
+        }
+    }
+    return true;
+}
 #endif
 
 void BFMatcher::knnMatchImpl( InputArray _queryDescriptors, std::vector<std::vector<DMatch> >& matches, int knn,
@@ -794,21 +938,15 @@ void BFMatcher::knnMatchImpl( InputArray _queryDescriptors, std::vector<std::vec
     {
         if(knn == 1)
         {
-            if(trainDescCollection.empty())
+            InputArray train = trainDescCollection.empty() ?
+                (InputArray)utrainDescCollection[0] : (InputArray)trainDescCollection[0];
+            bool oclOk = crossCheck ?
+                ocl_matchWithCrossCheck(_queryDescriptors, train, matches, normType) :
+                ocl_match(_queryDescriptors, train, matches, normType);
+            if (oclOk)
             {
-                if(ocl_match(_queryDescriptors, utrainDescCollection[0], matches, normType))
-                {
-                    CV_IMPL_ADD(CV_IMPL_OCL);
-                    return;
-                }
-            }
-            else
-            {
-                if(ocl_match(_queryDescriptors, trainDescCollection[0], matches, normType))
-                {
-                    CV_IMPL_ADD(CV_IMPL_OCL);
-                    return;
-                }
+                CV_IMPL_ADD(CV_IMPL_OCL);
+                return;
             }
         }
         else

modules/features/src/opencl/brute_force_match.cl

@@ -557,4 +557,138 @@ __kernel void BruteForceMatch_knnMatch(
         bestTrainIdx[queryIdx] = (int2)(myBestTrainIdx1, myBestTrainIdx2);
         bestDistance[queryIdx] = (float2)(myBestDistance1, myBestDistance2);
     }
-}
+}
+
+#ifdef HAVE_INT64_ATOMICS
+#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
+
+__kernel void BruteForceMatch_CrossCheckMatch(
+    __global T *query,
+    __global T *train,
+    __global int *bestTrainIdx,
+    __global float *bestDistance,
+    __global ulong *revBest,
+    int query_rows,
+    int query_cols,
+    int train_rows,
+    int train_cols,
+    int step
+)
+{
+    const int lidx = get_local_id(0);
+    const int lidy = get_local_id(1);
+    const int groupidx = get_group_id(0);
+
+    const int queryIdx = mad24(BLOCK_SIZE, groupidx, lidy);
+    const int queryOffset = min(queryIdx, query_rows - 1) * step;
+    __global TN *query_vec = (__global TN *)(query + queryOffset);
+    query_cols /= kercn;
+
+    __local float sharebuffer[SHARED_MEM_SZ];
+    __local value_type *s_query = (__local value_type *)sharebuffer;
+
+#if 0 < MAX_DESC_LEN
+    __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
+    #pragma unroll
+    for (int i = 0; i < MAX_DESC_LEN / BLOCK_SIZE; i++)
+    {
+        const int loadx = mad24(BLOCK_SIZE, i, lidx);
+        s_query[mad24(MAX_DESC_LEN, lidy, loadx)] = loadx < query_cols ? query_vec[loadx] : 0;
+    }
+#else
+    __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE_ODD * BLOCK_SIZE;
+    const int s_query_i = mad24(BLOCK_SIZE_ODD, lidy, lidx);
+    const int s_train_i = mad24(BLOCK_SIZE_ODD, lidx, lidy);
+#endif
+
+    float myBestDistance = MAX_FLOAT;
+    int myBestTrainIdx = -1;
+
+    for (int t = 0, endt = (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE; t < endt; t++)
+    {
+        result_type result = 0;
+
+        const int trainOffset = min(mad24(BLOCK_SIZE, t, lidy), train_rows - 1) * step;
+        __global TN *train_vec = (__global TN *)(train + trainOffset);
+#if 0 < MAX_DESC_LEN
+        #pragma unroll
+        for (int i = 0; i < MAX_DESC_LEN / BLOCK_SIZE; i++)
+        {
+            const int loadx = mad24(BLOCK_SIZE, i, lidx);
+            s_train[mad24(BLOCK_SIZE, lidx, lidy)] = loadx < train_cols ? train_vec[loadx] : 0;
+
+            barrier(CLK_LOCAL_MEM_FENCE);
+
+            result += reduce_multi_block(s_query, s_train, i, lidx, lidy);
+
+            barrier(CLK_LOCAL_MEM_FENCE);
+        }
+#else
+        for (int i = 0, endq = (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE; i < endq; i++)
+        {
+            const int loadx = mad24(BLOCK_SIZE, i, lidx);
+            if (loadx < query_cols)
+            {
+                s_query[s_query_i] = query_vec[loadx];
+                s_train[s_train_i] = train_vec[loadx];
+            }
+            else
+            {
+                s_query[s_query_i] = 0;
+                s_train[s_train_i] = 0;
+            }
+
+            barrier(CLK_LOCAL_MEM_FENCE);
+
+            result += reduce_block_match(s_query, s_train, lidx, lidy);
+
+            barrier(CLK_LOCAL_MEM_FENCE);
+        }
+#endif
+        result = DIST_RES(result);
+
+        const int trainIdx = mad24(BLOCK_SIZE, t, lidx);
+
+        if (queryIdx < query_rows && trainIdx < train_rows)
+        {
+            if (result < myBestDistance)
+            {
+                myBestDistance = result;
+                myBestTrainIdx = trainIdx;
+            }
+
+            uint dist_bits = as_uint(result);
+            ulong packed = ((ulong)dist_bits << 32) | (ulong)queryIdx;
+            atom_min(&revBest[trainIdx], packed);
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    __local float *s_distance = (__local float *)sharebuffer;
+    __local int *s_trainIdx = (__local int *)(sharebuffer + BLOCK_SIZE_ODD * BLOCK_SIZE);
+
+    s_distance += lidy * BLOCK_SIZE_ODD;
+    s_trainIdx += lidy * BLOCK_SIZE_ODD;
+    s_distance[lidx] = myBestDistance;
+    s_trainIdx[lidx] = myBestTrainIdx;
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    #pragma unroll
+    for (int k = 0; k < BLOCK_SIZE; k++)
+    {
+        if (myBestDistance > s_distance[k])
+        {
+            myBestDistance = s_distance[k];
+            myBestTrainIdx = s_trainIdx[k];
+        }
+    }
+
+    if (queryIdx < query_rows && lidx == 0)
+    {
+        bestTrainIdx[queryIdx] = myBestTrainIdx;
+        bestDistance[queryIdx] = myBestDistance;
+    }
+}
+#endif

modules/features/test/test_matchers_algorithmic.cpp

@@ -629,4 +629,47 @@ TEST(Features2d_DMatch, issue_17771)
     EXPECT_NO_THROW(ubf->knnMatch(usources, utargets, match, 1, mask, true));
 }
 
+// Verify that cross-check BFMatcher gives identical results via Mat (CPU) and UMat (OCL or CPU).
+// When OpenCL is active the UMat path exercises ocl_matchWithCrossCheck; when it is not,
+// both paths fall through to the same CPU code — either way the results must match.
+TEST(Features2d_BFMatcher_CrossCheck, ocl_matches_cpu)
+{
+    RNG rng(42);
+    const int nQuery = 200;
+    const int nTrain = 400;
+    const int dim    = 128;
+
+    // Float descriptors: the OCL dispatch in knnMatchImpl requires CV_32FC1
+    Mat queryMat(nQuery, dim, CV_32FC1);
+    Mat trainMat(nTrain, dim, CV_32FC1);
+    rng.fill(queryMat, RNG::UNIFORM, 0.f, 1.f);
+    rng.fill(trainMat, RNG::UNIFORM, 0.f, 1.f);
+
+    // CPU reference: Mat inputs always take the CPU path
+    Ptr<BFMatcher> cpuMatcher = BFMatcher::create(NORM_L2, true /*crossCheck*/);
+    vector<DMatch> cpuMatches;
+    cpuMatcher->match(queryMat, trainMat, cpuMatches);
+
+    // UMat path: activates OCL dispatch when OpenCL is available
+    UMat queryUMat = queryMat.getUMat(ACCESS_READ);
+    UMat trainUMat = trainMat.getUMat(ACCESS_READ);
+    Ptr<BFMatcher> oclMatcher = BFMatcher::create(NORM_L2, true /*crossCheck*/);
+    vector<DMatch> oclMatches;
+    oclMatcher->match(queryUMat, trainUMat, oclMatches);
+
+    // Both paths must return the same set of matches (order may differ)
+    ASSERT_EQ(cpuMatches.size(), oclMatches.size());
+
+    auto byQuery = [](const DMatch& a, const DMatch& b) { return a.queryIdx < b.queryIdx; };
+    sort(cpuMatches.begin(), cpuMatches.end(), byQuery);
+    sort(oclMatches.begin(), oclMatches.end(), byQuery);
+
+    for (size_t i = 0; i < cpuMatches.size(); ++i)
+    {
+        EXPECT_EQ(cpuMatches[i].queryIdx, oclMatches[i].queryIdx) << "at index " << i;
+        EXPECT_EQ(cpuMatches[i].trainIdx, oclMatches[i].trainIdx) << "at index " << i;
+        EXPECT_NEAR(cpuMatches[i].distance, oclMatches[i].distance, 1e-3f) << "at index " << i;
+    }
+}
+
 }} // namespace