implement optimized version of gpu::BruteForceMatcher::knnMatch when k == 2

diff --git a/modules/gpu/src/brute_force_matcher.cpp b/modules/gpu/src/brute_force_matcher.cpp
index 19521a2..d6bbb7e 100644 (file)
--- a/modules/gpu/src/brute_force_matcher.cpp
+++ b/modules/gpu/src/brute_force_matcher.cpp
@@ -105,13 +105,13 @@ namespace cv { namespace gpu { namespace bfmatcher
 \r
     template <typename T>\r
     void knnMatchL1_gpu(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn,\r
-        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
+        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
     template <typename T>\r
     void knnMatchL2_gpu(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn,\r
-        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
+        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
     template <typename T>\r
     void knnMatchHamming_gpu(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn,\r
-        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
+        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
 \r
     template <typename T>\r
     void radiusMatchL1_gpu(const DevMem2D& queryDescs, const DevMem2D& trainDescs, float maxDistance,\r
@@ -428,7 +428,7 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat& queryDescs, con
     using namespace cv::gpu::bfmatcher;\r
 \r
     typedef void (*match_caller_t)(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn,\r
-        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
+        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
 \r
     static const match_caller_t match_callers[3][8] =\r
     {\r
@@ -454,23 +454,28 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat& queryDescs, con
 \r
     ensureSizeIsEnough(nQuery, k, CV_32S, trainIdx);\r
     ensureSizeIsEnough(nQuery, k, CV_32F, distance);\r
-    ensureSizeIsEnough(nQuery, nTrain, CV_32FC1, allDist);\r
+    if (k != 2)\r
+        ensureSizeIsEnough(nQuery, nTrain, CV_32FC1, allDist);\r
 \r
     if (stream)\r
     {\r
         stream.enqueueMemSet(trainIdx, Scalar::all(-1));\r
-        stream.enqueueMemSet(allDist, Scalar::all(numeric_limits<float>::max()));\r
+        if (k != 2)\r
+            stream.enqueueMemSet(allDist, Scalar::all(numeric_limits<float>::max()));\r
     }\r
     else\r
     {\r
         trainIdx.setTo(Scalar::all(-1));\r
-        allDist.setTo(Scalar::all(numeric_limits<float>::max()));\r
+        if (k != 2)\r
+            allDist.setTo(Scalar::all(numeric_limits<float>::max()));\r
     }\r
 \r
     match_caller_t func = match_callers[distType][queryDescs.depth()];\r
     CV_Assert(func != 0);\r
+    \r
+    bool cc_12 = TargetArchs::builtWith(FEATURE_SET_COMPUTE_12) && DeviceInfo().supports(FEATURE_SET_COMPUTE_12);\r
 \r
-    func(queryDescs, trainDescs, k, mask, trainIdx, distance, allDist, StreamAccessor::getStream(stream));\r
+    func(queryDescs, trainDescs, k, mask, trainIdx, distance, allDist, cc_12, StreamAccessor::getStream(stream));\r
 }\r
 \r
 void cv::gpu::BruteForceMatcher_GPU_base::knnMatchDownload(const GpuMat& trainIdx, const GpuMat& distance,\r

diff --git a/modules/gpu/src/cuda/brute_force_matcher.cu b/modules/gpu/src/cuda/brute_force_matcher.cu
index 6b13619..4cd1142 100644 (file)
--- a/modules/gpu/src/cuda/brute_force_matcher.cu
+++ b/modules/gpu/src/cuda/brute_force_matcher.cu
@@ -87,9 +87,8 @@ namespace cv { namespace gpu { namespace bfmatcher
         PtrStep curMask;\r
     };\r
 \r
-    class WithOutMask\r
+    struct WithOutMask\r
     {\r
-    public:\r
         __device__ __forceinline__ void nextMask() const\r
         {\r
         }\r
@@ -102,21 +101,19 @@ namespace cv { namespace gpu { namespace bfmatcher
     ///////////////////////////////////////////////////////////////////////////////\r
     // Reduce Sum\r
 \r
-    template <int BLOCK_DIM_X> struct SumReductor;    \r
+    template <int BLOCK_DIM_X> struct SumReductor;\r
     template <> struct SumReductor<16>\r
     {\r
-        template <typename T> static __device__ void reduce(T* sdiff_row, T& mySum)\r
+        template <typename T> static __device__ void reduce(volatile T* sdiff_row, T& mySum)\r
         {\r
-            volatile T* smem = sdiff_row;\r
-\r
-            smem[threadIdx.x] = mySum;\r
+            sdiff_row[threadIdx.x] = mySum;\r
             \r
             if (threadIdx.x < 8) \r
             {\r
-                smem[threadIdx.x] = mySum += smem[threadIdx.x + 8]; \r
-                smem[threadIdx.x] = mySum += smem[threadIdx.x + 4]; \r
-                smem[threadIdx.x] = mySum += smem[threadIdx.x + 2];\r
-                smem[threadIdx.x] = mySum += smem[threadIdx.x + 1];  \r
+                sdiff_row[threadIdx.x] = mySum += sdiff_row[threadIdx.x + 8]; \r
+                sdiff_row[threadIdx.x] = mySum += sdiff_row[threadIdx.x + 4]; \r
+                sdiff_row[threadIdx.x] = mySum += sdiff_row[threadIdx.x + 2];\r
+                sdiff_row[threadIdx.x] = mySum += sdiff_row[threadIdx.x + 1];  \r
             }\r
         }\r
     };\r
@@ -344,7 +341,7 @@ namespace cv { namespace gpu { namespace bfmatcher
     ///////////////////////////////////////////////////////////////////////////////\r
     // warpReduceMinIdxIdx\r
 \r
-    template <int BLOCK_DIM_Y> struct MinIdxIdxWarpReductor;\r
+    template <int BLOCK_DIM_Y> struct MinIdxIdxWarpReductor;    \r
     template <> struct MinIdxIdxWarpReductor<16>\r
     {\r
         template <typename T> \r
@@ -435,6 +432,7 @@ namespace cv { namespace gpu { namespace bfmatcher
         __device__ __forceinline__ void prepare(const T* queryDescs, int desc_len, U* smem)\r
         {\r
             loadDescsVals<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN>(queryDescs, desc_len, queryVals, smem);\r
+            __syncthreads();\r
         }\r
 \r
         template <typename Dist>\r
@@ -778,6 +776,173 @@ namespace cv { namespace gpu { namespace bfmatcher
 ///////////////////////////////////////////////////////////////////////////////////\r
 //////////////////////////////////// Knn Match ////////////////////////////////////\r
 ///////////////////////////////////////////////////////////////////////////////////\r
+\r
+    template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename Dist, typename ReduceDescCalculator, typename T, typename Mask>\r
+    __device__ void distanceCalcLoop(const PtrStep_<T>& query, const DevMem2D_<T>& train, const Mask& m, int queryIdx,\r
+        typename Dist::ResultType& distMin1, typename Dist::ResultType& distMin2, int& bestTrainIdx1, int& bestTrainIdx2, \r
+        typename Dist::ResultType* smem)\r
+    {\r
+        ReduceDescCalculator reduceDescCalc;\r
+\r
+        reduceDescCalc.prepare(query.ptr(queryIdx), train.cols, (typename Dist::ValueType*)smem);\r
+        \r
+        typename Dist::ResultType* sdiffRow = smem + BLOCK_DIM_X * threadIdx.y;\r
+\r
+        for (int trainIdx = threadIdx.y; trainIdx < train.rows; trainIdx += BLOCK_DIM_Y)\r
+        {\r
+            if (m(queryIdx, trainIdx))\r
+            {\r
+                Dist dist;\r
+\r
+                const T* trainRow = train.ptr(trainIdx);\r
+                \r
+                reduceDescCalc.calc(trainRow, train.cols, dist, sdiffRow);\r
+\r
+                if (threadIdx.x == 0)\r
+                {\r
+                    typename Dist::ResultType val = dist;\r
+\r
+                    if (val < distMin1)\r
+                    {\r
+                        distMin1 = val;\r
+                        bestTrainIdx1 = trainIdx;\r
+                    }\r
+                    else if (val < distMin2)\r
+                    {\r
+                        distMin2 = val;\r
+                        bestTrainIdx2 = trainIdx;\r
+                    }\r
+                }\r
+            }\r
+        }\r
+    }\r
+\r
+    template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename Dist, typename ReduceDescCalculator, typename T, typename Mask>\r
+    __global__ void knnMatch2(const PtrStep_<T> query, const DevMem2D_<T> train, const Mask m, PtrStep_<int2> trainIdx, PtrStep_<float2> distance)\r
+    {\r
+        typedef typename Dist::ResultType ResultType;\r
+        typedef typename Dist::ValueType ValueType;\r
+\r
+        __shared__ ResultType smem[BLOCK_DIM_X * BLOCK_DIM_Y];\r
+\r
+        const int queryIdx = blockIdx.x;\r
+\r
+        ResultType distMin1 = numeric_limits<ResultType>::max();\r
+        ResultType distMin2 = numeric_limits<ResultType>::max();\r
+\r
+        int bestTrainIdx1 = -1;\r
+        int bestTrainIdx2 = -1;\r
+\r
+        distanceCalcLoop<BLOCK_DIM_X, BLOCK_DIM_Y, Dist, ReduceDescCalculator>(query, train, m, queryIdx, \r
+            distMin1, distMin2, bestTrainIdx1, bestTrainIdx2, smem);\r
+        __syncthreads();\r
+\r
+        volatile ResultType* sdistMinRow = smem;\r
+        volatile int* sbestTrainIdxRow = (int*)(sdistMinRow + 2 * BLOCK_DIM_Y);\r
+\r
+        if (threadIdx.x == 0)\r
+        {\r
+            sdistMinRow[threadIdx.y] = distMin1;\r
+            sdistMinRow[threadIdx.y + BLOCK_DIM_Y] = distMin2;\r
+\r
+            sbestTrainIdxRow[threadIdx.y] = bestTrainIdx1;            \r
+            sbestTrainIdxRow[threadIdx.y + BLOCK_DIM_Y] = bestTrainIdx2;\r
+        }\r
+        __syncthreads();\r
+\r
+        if (threadIdx.x == 0 && threadIdx.y == 0)\r
+        {\r
+            distMin1 = numeric_limits<ResultType>::max();\r
+            distMin2 = numeric_limits<ResultType>::max();\r
+\r
+            bestTrainIdx1 = -1;\r
+            bestTrainIdx2 = -1;\r
+\r
+            #pragma unroll\r
+            for (int i = 0; i < BLOCK_DIM_Y; ++i)\r
+            {\r
+                ResultType val = sdistMinRow[i];\r
+\r
+                if (val < distMin1)\r
+                {\r
+                    distMin1 = val;\r
+                    bestTrainIdx1 = sbestTrainIdxRow[i];\r
+                }\r
+                else if (val < distMin2)\r
+                {\r
+                    distMin2 = val;\r
+                    bestTrainIdx2 = sbestTrainIdxRow[i];\r
+                }\r
+            }\r
+\r
+            #pragma unroll\r
+            for (int i = BLOCK_DIM_Y; i < 2 * BLOCK_DIM_Y; ++i)\r
+            {\r
+                ResultType val = sdistMinRow[i];\r
+\r
+                if (val < distMin2)\r
+                {\r
+                    distMin2 = val;\r
+                    bestTrainIdx2 = sbestTrainIdxRow[i];\r
+                }\r
+            }\r
+\r
+            trainIdx.ptr(queryIdx)[0] = make_int2(bestTrainIdx1, bestTrainIdx2);\r
+            distance.ptr(queryIdx)[0] = make_float2(distMin1, distMin2);\r
+        }\r
+    }\r
+\r
+    template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename Dist, typename T, typename Mask>\r
+    void knnMatch2Simple_caller(const DevMem2D_<T>& queryDescs, const DevMem2D_<T>& trainDescs, const Mask& mask, \r
+        const DevMem2D_<int2>& trainIdx, const DevMem2D_<float2>& distance, cudaStream_t stream)\r
+    {\r
+        dim3 grid(queryDescs.rows, 1, 1);\r
+        dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);\r
+\r
+        knnMatch2<BLOCK_DIM_X, BLOCK_DIM_Y, Dist, ReduceDescCalculatorSimple<BLOCK_DIM_X, T>, T>\r
+            <<<grid, threads, 0, stream>>>(queryDescs, trainDescs, mask, trainIdx, distance);\r
+        cudaSafeCall( cudaGetLastError() );\r
+\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
+    }\r
+    template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int MAX_DESCRIPTORS_LEN, bool DESC_LEN_EQ_MAX_LEN, typename Dist, typename T, typename Mask>\r
+    void knnMatch2Cached_caller(const DevMem2D_<T>& queryDescs, const DevMem2D_<T>& trainDescs, const Mask& mask, \r
+        const DevMem2D_<int2>& trainIdx, const DevMem2D_<float2>& distance, cudaStream_t stream)\r
+    {\r
+        StaticAssert<BLOCK_DIM_X * BLOCK_DIM_Y >= MAX_DESCRIPTORS_LEN>::check(); // block size must be greter than descriptors length\r
+        StaticAssert<MAX_DESCRIPTORS_LEN % BLOCK_DIM_X == 0>::check();           // max descriptors length must divide to blockDimX\r
+\r
+        dim3 grid(queryDescs.rows, 1, 1);\r
+        dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);\r
+\r
+        knnMatch2<BLOCK_DIM_X, BLOCK_DIM_Y, Dist, ReduceDescCalculatorCached<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN, DESC_LEN_EQ_MAX_LEN, T, typename Dist::ValueType>, T>\r
+              <<<grid, threads, 0, stream>>>(queryDescs, trainDescs, mask, trainIdx, distance);\r
+        cudaSafeCall( cudaGetLastError() );\r
+\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
+    }\r
+    \r
+    template <typename Dist, typename T, typename Mask>\r
+    void knnMatch2Dispatcher(const DevMem2D_<T>& query, const DevMem2D_<T>& train, const Mask& mask, \r
+        const DevMem2D_<int2>& trainIdx, const DevMem2D_<float2>& distance, bool cc_12, cudaStream_t stream)\r
+    {\r
+        if (query.cols < 64)\r
+            knnMatch2Cached_caller<16, 16, 64, false, Dist>(query, train, mask, trainIdx, distance, stream);\r
+        else if (query.cols == 64)\r
+            knnMatch2Cached_caller<16, 16, 64, true, Dist>(query, train, mask, trainIdx, distance, stream);\r
+        else if (query.cols < 128)\r
+            knnMatch2Cached_caller<16, 16, 128, false, Dist>(query, train, mask, trainIdx, distance, stream);\r
+        else if (query.cols == 128 && cc_12)\r
+            knnMatch2Cached_caller<16, 16, 128, true, Dist>(query, train, mask, trainIdx, distance, stream);\r
+        else if (query.cols < 256 && cc_12)\r
+            knnMatch2Cached_caller<16, 16, 256, false, Dist>(query, train, mask, trainIdx, distance, stream);\r
+        else if (query.cols == 256 && cc_12)\r
+            knnMatch2Cached_caller<16, 16, 256, true, Dist>(query, train, mask, trainIdx, distance, stream);\r
+        else\r
+            knnMatch2Simple_caller<16, 16, Dist>(query, train, mask, trainIdx, distance, stream);\r
+    }\r
     \r
     ///////////////////////////////////////////////////////////////////////////////\r
     // Calc distance kernel\r
@@ -1026,77 +1191,74 @@ namespace cv { namespace gpu { namespace bfmatcher
         findKnnMatch_caller<256>(knn, trainIdx, distance, allDist, stream);\r
     }\r
 \r
-    template <typename T>\r
-    void knnMatchL1_gpu(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn,\r
-        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream)\r
+    template < typename Dist, typename T >\r
+    void knnMatchDispatcher(const DevMem2D_<T>& queryDescs, const DevMem2D_<T>& trainDescs, int knn,\r
+        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream)\r
     {\r
         if (mask.data)\r
         {\r
-            calcDistanceDispatcher< L1Dist<T> >((DevMem2D_<T>)queryDescs, (DevMem2D_<T>)trainDescs, SingleMask(mask), allDist, stream);\r
+            if (knn == 2)\r
+            {\r
+                knnMatch2Dispatcher<Dist>(queryDescs, trainDescs, SingleMask(mask), (DevMem2D_<int2>)trainIdx, (DevMem2D_<float2>)distance, cc_12, stream);\r
+                return;\r
+            }\r
+\r
+            calcDistanceDispatcher<Dist>(queryDescs, trainDescs, SingleMask(mask), allDist, stream);\r
         }\r
         else\r
         {\r
-            calcDistanceDispatcher< L1Dist<T> >((DevMem2D_<T>)queryDescs, (DevMem2D_<T>)trainDescs, WithOutMask(), allDist, stream);\r
+            if (knn == 2)\r
+            {\r
+                knnMatch2Dispatcher<Dist>(queryDescs, trainDescs, WithOutMask(), (DevMem2D_<int2>)trainIdx, (DevMem2D_<float2>)distance, cc_12, stream);\r
+                return;\r
+            }\r
+\r
+            calcDistanceDispatcher<Dist>(queryDescs, trainDescs, WithOutMask(), allDist, stream);\r
         }\r
 \r
         findKnnMatchDispatcher(knn, trainIdx, distance, allDist, stream);\r
     }\r
 \r
-    template void knnMatchL1_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL1_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL1_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL1_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL1_gpu<int   >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL1_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
+    template <typename T>\r
+    void knnMatchL1_gpu(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn,\r
+        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream)\r
+    {\r
+        knnMatchDispatcher< L1Dist<T> >((DevMem2D_<T>)queryDescs, (DevMem2D_<T>)trainDescs, knn, mask, trainIdx, distance, allDist, cc_12, stream);\r
+    }\r
+\r
+    template void knnMatchL1_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL1_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL1_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL1_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL1_gpu<int   >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL1_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
 \r
     template <typename T>\r
     void knnMatchL2_gpu(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn,\r
-        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream)\r
+        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream)\r
     {\r
-        if (mask.data)\r
-        {\r
-            calcDistanceDispatcher<L2Dist>((DevMem2D_<T>)queryDescs, (DevMem2D_<T>)trainDescs, \r
-                SingleMask(mask), allDist, stream);\r
-        }\r
-        else\r
-        {\r
-            calcDistanceDispatcher<L2Dist>((DevMem2D_<T>)queryDescs, (DevMem2D_<T>)trainDescs, \r
-                WithOutMask(), allDist, stream);\r
-        }\r
-\r
-        findKnnMatchDispatcher(knn, trainIdx, distance, allDist, stream);\r
+        knnMatchDispatcher<L2Dist>((DevMem2D_<T>)queryDescs, (DevMem2D_<T>)trainDescs, knn, mask, trainIdx, distance, allDist, cc_12, stream);\r
     }\r
 \r
-    template void knnMatchL2_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL2_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL2_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL2_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL2_gpu<int   >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchL2_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
+    template void knnMatchL2_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL2_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL2_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL2_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL2_gpu<int   >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchL2_gpu<float >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
 \r
     template <typename T>\r
     void knnMatchHamming_gpu(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn,\r
-        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream)\r
+        const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream)\r
     {\r
-        if (mask.data)\r
-        {\r
-            calcDistanceDispatcher<HammingDist>((DevMem2D_<T>)queryDescs, (DevMem2D_<T>)trainDescs, \r
-                SingleMask(mask), allDist, stream);\r
-        }\r
-        else\r
-        {\r
-            calcDistanceDispatcher<HammingDist>((DevMem2D_<T>)queryDescs, (DevMem2D_<T>)trainDescs, \r
-                WithOutMask(), allDist, stream);\r
-        }\r
-\r
-        findKnnMatchDispatcher(knn, trainIdx, distance, allDist, stream);\r
+        knnMatchDispatcher<HammingDist>((DevMem2D_<T>)queryDescs, (DevMem2D_<T>)trainDescs, knn, mask, trainIdx, distance, allDist, cc_12, stream);\r
     }\r
 \r
-    template void knnMatchHamming_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchHamming_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchHamming_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchHamming_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
-    template void knnMatchHamming_gpu<int   >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, cudaStream_t stream);\r
+    template void knnMatchHamming_gpu<uchar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchHamming_gpu<schar >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchHamming_gpu<ushort>(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchHamming_gpu<short >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
+    template void knnMatchHamming_gpu<int   >(const DevMem2D& queryDescs, const DevMem2D& trainDescs, int knn, const DevMem2D& mask, const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2Df& allDist, bool cc_12, cudaStream_t stream);\r
 \r
 ///////////////////////////////////////////////////////////////////////////////////\r
 /////////////////////////////////// Radius Match //////////////////////////////////\r

diff --git a/modules/gpu/test/test_features2d.cpp b/modules/gpu/test/test_features2d.cpp
index 516c0f5..248ca76 100644 (file)
--- a/modules/gpu/test/test_features2d.cpp
+++ b/modules/gpu/test/test_features2d.cpp
@@ -320,7 +320,7 @@ TEST_P(BruteForceMatcher, KnnMatch)
     PRINT_PARAM(distStr);\r
     PRINT_PARAM(dim);\r
 \r
-    const int knn = 3;\r
+    const int knn = 2;\r
 \r
     std::vector< std::vector<cv::DMatch> > matches;\r
 \r

diff --git a/samples/gpu/performance/tests.cpp b/samples/gpu/performance/tests.cpp
index 1b2cbbc..ebc2dad 100644 (file)
--- a/samples/gpu/performance/tests.cpp
+++ b/samples/gpu/performance/tests.cpp
@@ -286,7 +286,7 @@ TEST(BruteForceMatcher)
 {\r
     // Init CPU matcher\r
 \r
-    int desc_len = 128;\r
+    int desc_len = 64;\r
 \r
     BruteForceMatcher< L2<float> > matcher;\r
 \r
@@ -328,7 +328,7 @@ TEST(BruteForceMatcher)
     d_matcher.knnMatch(d_query, d_train, d_matches, knn);\r
     GPU_OFF;\r
 \r
-    SUBTEST << "radiusMatch";\r
+    /*SUBTEST << "radiusMatch";\r
     float max_distance = 3.8f;\r
 \r
     CPU_ON;\r
@@ -337,7 +337,7 @@ TEST(BruteForceMatcher)
 \r
     GPU_ON;\r
     d_matcher.radiusMatch(d_query, d_train, d_matches, max_distance);\r
-    GPU_OFF;\r
+    GPU_OFF;*/\r
 }\r
 \r
 \r