added assertion on device features (global atomics) into gpu tests

diff --git a/modules/gpu/src/brute_force_matcher.cpp b/modules/gpu/src/brute_force_matcher.cpp
index 7f11282..8270dfe 100644 (file)
--- a/modules/gpu/src/brute_force_matcher.cpp
+++ b/modules/gpu/src/brute_force_matcher.cpp
@@ -82,76 +82,76 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat&, vector< vec
 \r
 #else /* !defined (HAVE_CUDA) */\r
 \r
-namespace cv { namespace gpu { namespace device \r
+namespace cv { namespace gpu { namespace device\r
 {\r
     namespace bf_match\r
     {\r
-        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db& train, const DevMem2Db& mask, \r
-            const DevMem2Di& trainIdx, const DevMem2Df& distance, \r
+        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db& train, const DevMem2Db& mask,\r
+            const DevMem2Di& trainIdx, const DevMem2Df& distance,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db& train, const DevMem2Db& mask, \r
-            const DevMem2Di& trainIdx, const DevMem2Df& distance, \r
+        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db& train, const DevMem2Db& mask,\r
+            const DevMem2Di& trainIdx, const DevMem2Df& distance,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db& train, const DevMem2Db& mask, \r
-            const DevMem2Di& trainIdx, const DevMem2Df& distance, \r
+        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db& train, const DevMem2Db& mask,\r
+            const DevMem2Di& trainIdx, const DevMem2Df& distance,\r
             int cc, cudaStream_t stream);\r
 \r
-        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks, \r
-            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, \r
+        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks,\r
+            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks, \r
-            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, \r
+        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks,\r
+            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks, \r
+        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks,\r
             const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance,\r
             int cc, cudaStream_t stream);\r
     }\r
 \r
     namespace bf_knnmatch\r
     {\r
-        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db& train, int k, const DevMem2Db& mask, \r
-            const DevMem2Db& trainIdx, const DevMem2Db& distance, const DevMem2Df& allDist, \r
+        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db& train, int k, const DevMem2Db& mask,\r
+            const DevMem2Db& trainIdx, const DevMem2Db& distance, const DevMem2Df& allDist,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db& train, int k, const DevMem2Db& mask, \r
-            const DevMem2Db& trainIdx, const DevMem2Db& distance, const DevMem2Df& allDist, \r
+        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db& train, int k, const DevMem2Db& mask,\r
+            const DevMem2Db& trainIdx, const DevMem2Db& distance, const DevMem2Df& allDist,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db& train, int k, const DevMem2Db& mask, \r
-            const DevMem2Db& trainIdx, const DevMem2Db& distance, const DevMem2Df& allDist, \r
+        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db& train, int k, const DevMem2Db& mask,\r
+            const DevMem2Db& trainIdx, const DevMem2Db& distance, const DevMem2Df& allDist,\r
             int cc, cudaStream_t stream);\r
 \r
-        template <typename T> void match2L1_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks, \r
-            const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance, \r
+        template <typename T> void match2L1_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks,\r
+            const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void match2L2_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks, \r
-            const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance, \r
+        template <typename T> void match2L2_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks,\r
+            const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void match2Hamming_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks, \r
-            const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance, \r
+        template <typename T> void match2Hamming_gpu(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks,\r
+            const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance,\r
             int cc, cudaStream_t stream);\r
     }\r
 \r
-    namespace bf_radius_match \r
+    namespace bf_radius_match\r
     {\r
-        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db& train, float maxDistance, const DevMem2Db& mask, \r
-            const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, \r
+        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db& train, float maxDistance, const DevMem2Db& mask,\r
+            const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db& train, float maxDistance, const DevMem2Db& mask, \r
-            const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, \r
+        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db& train, float maxDistance, const DevMem2Db& mask,\r
+            const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,\r
             int cc, cudaStream_t stream);\r
-        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db& train, float maxDistance, const DevMem2Db& mask, \r
-            const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, \r
+        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db& train, float maxDistance, const DevMem2Db& mask,\r
+            const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,\r
             int cc, cudaStream_t stream);\r
 \r
-        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db* trains, int n, float maxDistance, const DevMem2Db* masks, \r
-            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, \r
+        template <typename T> void matchL1_gpu(const DevMem2Db& query, const DevMem2Db* trains, int n, float maxDistance, const DevMem2Db* masks,\r
+            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,\r
             int cc, cudaStream_t stream);\r
 \r
-        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db* trains, int n, float maxDistance, const DevMem2Db* masks, \r
-            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, \r
+        template <typename T> void matchL2_gpu(const DevMem2Db& query, const DevMem2Db* trains, int n, float maxDistance, const DevMem2Db* masks,\r
+            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,\r
             int cc, cudaStream_t stream);\r
 \r
-        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db* trains, int n, float maxDistance, const DevMem2Db* masks, \r
-            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, \r
+        template <typename T> void matchHamming_gpu(const DevMem2Db& query, const DevMem2Db* trains, int n, float maxDistance, const DevMem2Db* masks,\r
+            const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,\r
             int cc, cudaStream_t stream);\r
     }\r
 }}}\r
@@ -192,7 +192,7 @@ bool cv::gpu::BruteForceMatcher_GPU_base::isMaskSupported() const
 // Match\r
 \r
 void cv::gpu::BruteForceMatcher_GPU_base::matchSingle(const GpuMat& query, const GpuMat& train,\r
-    GpuMat& trainIdx, GpuMat& distance, \r
+    GpuMat& trainIdx, GpuMat& distance,\r
     const GpuMat& mask, Stream& stream)\r
 {\r
     if (query.empty() || train.empty())\r
@@ -200,25 +200,25 @@ void cv::gpu::BruteForceMatcher_GPU_base::matchSingle(const GpuMat& query, const
 \r
     using namespace ::cv::gpu::device::bf_match;\r
 \r
-    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& train, const DevMem2Db& mask, \r
+    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& train, const DevMem2Db& mask,\r
                              const DevMem2Di& trainIdx, const DevMem2Df& distance,\r
                              int cc, cudaStream_t stream);\r
 \r
     static const caller_t callers[3][6] =\r
     {\r
         {\r
-            matchL1_gpu<unsigned char>, 0/*matchL1_gpu<signed char>*/, \r
-            matchL1_gpu<unsigned short>, matchL1_gpu<short>, \r
+            matchL1_gpu<unsigned char>, 0/*matchL1_gpu<signed char>*/,\r
+            matchL1_gpu<unsigned short>, matchL1_gpu<short>,\r
             matchL1_gpu<int>, matchL1_gpu<float>\r
         },\r
         {\r
-            0/*matchL2_gpu<unsigned char>*/, 0/*matchL2_gpu<signed char>*/, \r
-            0/*matchL2_gpu<unsigned short>*/, 0/*matchL2_gpu<short>*/, \r
+            0/*matchL2_gpu<unsigned char>*/, 0/*matchL2_gpu<signed char>*/,\r
+            0/*matchL2_gpu<unsigned short>*/, 0/*matchL2_gpu<short>*/,\r
             0/*matchL2_gpu<int>*/, matchL2_gpu<float>\r
         },\r
         {\r
-            matchHamming_gpu<unsigned char>, 0/*matchHamming_gpu<signed char>*/, \r
-            matchHamming_gpu<unsigned short>, 0/*matchHamming_gpu<short>*/, \r
+            matchHamming_gpu<unsigned char>, 0/*matchHamming_gpu<signed char>*/,\r
+            matchHamming_gpu<unsigned short>, 0/*matchHamming_gpu<short>*/,\r
             matchHamming_gpu<int>, 0/*matchHamming_gpu<float>*/\r
         }\r
     };\r
@@ -334,7 +334,7 @@ void cv::gpu::BruteForceMatcher_GPU_base::makeGpuCollection(GpuMat& trainCollect
 }\r
 \r
 void cv::gpu::BruteForceMatcher_GPU_base::matchCollection(const GpuMat& query, const GpuMat& trainCollection,\r
-    GpuMat& trainIdx, GpuMat& imgIdx, GpuMat& distance, \r
+    GpuMat& trainIdx, GpuMat& imgIdx, GpuMat& distance,\r
     const GpuMat& masks, Stream& stream)\r
 {\r
     if (query.empty() || trainCollection.empty())\r
@@ -342,8 +342,8 @@ void cv::gpu::BruteForceMatcher_GPU_base::matchCollection(const GpuMat& query, c
 \r
     using namespace ::cv::gpu::device::bf_match;\r
 \r
-    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks, \r
-                             const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, \r
+    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks,\r
+                             const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance,\r
                              int cc, cudaStream_t stream);\r
 \r
     static const caller_t callers[3][6] =\r
@@ -453,25 +453,25 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatchSingle(const GpuMat& query, co
 \r
     using namespace ::cv::gpu::device::bf_knnmatch;\r
 \r
-    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& train, int k, const DevMem2Db& mask, \r
-                             const DevMem2Db& trainIdx, const DevMem2Db& distance, const DevMem2Df& allDist, \r
+    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& train, int k, const DevMem2Db& mask,\r
+                             const DevMem2Db& trainIdx, const DevMem2Db& distance, const DevMem2Df& allDist,\r
                              int cc, cudaStream_t stream);\r
 \r
     static const caller_t callers[3][6] =\r
     {\r
         {\r
-            matchL1_gpu<unsigned char>, 0/*matchL1_gpu<signed char>*/, \r
-            matchL1_gpu<unsigned short>, matchL1_gpu<short>, \r
+            matchL1_gpu<unsigned char>, 0/*matchL1_gpu<signed char>*/,\r
+            matchL1_gpu<unsigned short>, matchL1_gpu<short>,\r
             matchL1_gpu<int>, matchL1_gpu<float>\r
         },\r
         {\r
-            0/*matchL2_gpu<unsigned char>*/, 0/*matchL2_gpu<signed char>*/, \r
-            0/*matchL2_gpu<unsigned short>*/, 0/*matchL2_gpu<short>*/, \r
+            0/*matchL2_gpu<unsigned char>*/, 0/*matchL2_gpu<signed char>*/,\r
+            0/*matchL2_gpu<unsigned short>*/, 0/*matchL2_gpu<short>*/,\r
             0/*matchL2_gpu<int>*/, matchL2_gpu<float>\r
         },\r
         {\r
-            matchHamming_gpu<unsigned char>, 0/*matchHamming_gpu<signed char>*/, \r
-            matchHamming_gpu<unsigned short>, 0/*matchHamming_gpu<short>*/, \r
+            matchHamming_gpu<unsigned char>, 0/*matchHamming_gpu<signed char>*/,\r
+            matchHamming_gpu<unsigned short>, 0/*matchHamming_gpu<short>*/,\r
             matchHamming_gpu<int>, 0/*matchHamming_gpu<float>*/\r
         }\r
     };\r
@@ -501,7 +501,7 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatchSingle(const GpuMat& query, co
 \r
     caller_t func = callers[distType][query.depth()];\r
     CV_Assert(func != 0);\r
-    \r
+\r
     DeviceInfo info;\r
     int cc = info.majorVersion() * 10 + info.minorVersion();\r
 \r
@@ -520,7 +520,7 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatchDownload(const GpuMat& trainId
     knnMatchConvert(trainIdxCPU, distanceCPU, matches, compactResult);\r
 }\r
 \r
-void cv::gpu::BruteForceMatcher_GPU_base::knnMatchConvert(const Mat& trainIdx, const Mat& distance, \r
+void cv::gpu::BruteForceMatcher_GPU_base::knnMatchConvert(const Mat& trainIdx, const Mat& distance,\r
     vector< vector<DMatch> >& matches, bool compactResult)\r
 {\r
     if (trainIdx.empty() || distance.empty())\r
@@ -536,7 +536,7 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatchConvert(const Mat& trainIdx, c
 \r
     matches.clear();\r
     matches.reserve(nQuery);\r
-    \r
+\r
     const int* trainIdx_ptr = trainIdx.ptr<int>();\r
     const float* distance_ptr = distance.ptr<float>();\r
 \r
@@ -582,25 +582,25 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatch2Collection(const GpuMat& quer
 \r
     using namespace ::cv::gpu::device::bf_knnmatch;\r
 \r
-    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks, \r
-                             const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance, \r
+    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& trains, const DevMem2D_<PtrStepb>& masks,\r
+                             const DevMem2Db& trainIdx, const DevMem2Db& imgIdx, const DevMem2Db& distance,\r
                              int cc, cudaStream_t stream);\r
 \r
     static const caller_t callers[3][6] =\r
     {\r
         {\r
-            match2L1_gpu<unsigned char>, 0/*match2L1_gpu<signed char>*/, \r
-            match2L1_gpu<unsigned short>, match2L1_gpu<short>, \r
+            match2L1_gpu<unsigned char>, 0/*match2L1_gpu<signed char>*/,\r
+            match2L1_gpu<unsigned short>, match2L1_gpu<short>,\r
             match2L1_gpu<int>, match2L1_gpu<float>\r
         },\r
         {\r
-            0/*match2L2_gpu<unsigned char>*/, 0/*match2L2_gpu<signed char>*/, \r
-            0/*match2L2_gpu<unsigned short>*/, 0/*match2L2_gpu<short>*/, \r
+            0/*match2L2_gpu<unsigned char>*/, 0/*match2L2_gpu<signed char>*/,\r
+            0/*match2L2_gpu<unsigned short>*/, 0/*match2L2_gpu<short>*/,\r
             0/*match2L2_gpu<int>*/, match2L2_gpu<float>\r
         },\r
         {\r
-            match2Hamming_gpu<unsigned char>, 0/*match2Hamming_gpu<signed char>*/, \r
-            match2Hamming_gpu<unsigned short>, 0/*match2Hamming_gpu<short>*/, \r
+            match2Hamming_gpu<unsigned char>, 0/*match2Hamming_gpu<signed char>*/,\r
+            match2Hamming_gpu<unsigned short>, 0/*match2Hamming_gpu<short>*/,\r
             match2Hamming_gpu<int>, 0/*match2Hamming_gpu<float>*/\r
         }\r
     };\r
@@ -620,7 +620,7 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatch2Collection(const GpuMat& quer
 \r
     caller_t func = callers[distType][query.depth()];\r
     CV_Assert(func != 0);\r
-    \r
+\r
     DeviceInfo info;\r
     int cc = info.majorVersion() * 10 + info.minorVersion();\r
 \r
@@ -654,7 +654,7 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatch2Convert(const Mat& trainIdx,
 \r
     matches.clear();\r
     matches.reserve(nQuery);\r
-    \r
+\r
     const int* trainIdx_ptr = trainIdx.ptr<int>();\r
     const int* imgIdx_ptr = imgIdx.ptr<int>();\r
     const float* distance_ptr = distance.ptr<float>();\r
@@ -755,33 +755,33 @@ void cv::gpu::BruteForceMatcher_GPU_base::knnMatch(const GpuMat& query, vector<
 // RadiusMatch\r
 \r
 void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchSingle(const GpuMat& query, const GpuMat& train,\r
-    GpuMat& trainIdx, GpuMat& distance, GpuMat& nMatches, float maxDistance, \r
+    GpuMat& trainIdx, GpuMat& distance, GpuMat& nMatches, float maxDistance,\r
     const GpuMat& mask, Stream& stream)\r
 {\r
     if (query.empty() || train.empty())\r
         return;\r
 \r
-    using namespace ::cv::gpu::device::bf_radius_match;\r
+    using namespace cv::gpu::device::bf_radius_match;\r
 \r
-    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& train, float maxDistance, const DevMem2Db& mask, \r
-                             const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, \r
+    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db& train, float maxDistance, const DevMem2Db& mask,\r
+                             const DevMem2Di& trainIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,\r
                              int cc, cudaStream_t stream);\r
 \r
     static const caller_t callers[3][6] =\r
     {\r
         {\r
-            matchL1_gpu<unsigned char>, 0/*matchL1_gpu<signed char>*/, \r
-            matchL1_gpu<unsigned short>, matchL1_gpu<short>, \r
+            matchL1_gpu<unsigned char>, 0/*matchL1_gpu<signed char>*/,\r
+            matchL1_gpu<unsigned short>, matchL1_gpu<short>,\r
             matchL1_gpu<int>, matchL1_gpu<float>\r
         },\r
         {\r
-            0/*matchL2_gpu<unsigned char>*/, 0/*matchL2_gpu<signed char>*/, \r
-            0/*matchL2_gpu<unsigned short>*/, 0/*matchL2_gpu<short>*/, \r
+            0/*matchL2_gpu<unsigned char>*/, 0/*matchL2_gpu<signed char>*/,\r
+            0/*matchL2_gpu<unsigned short>*/, 0/*matchL2_gpu<short>*/,\r
             0/*matchL2_gpu<int>*/, matchL2_gpu<float>\r
         },\r
         {\r
-            matchHamming_gpu<unsigned char>, 0/*matchHamming_gpu<signed char>*/, \r
-            matchHamming_gpu<unsigned short>, 0/*matchHamming_gpu<short>*/, \r
+            matchHamming_gpu<unsigned char>, 0/*matchHamming_gpu<signed char>*/,\r
+            matchHamming_gpu<unsigned short>, 0/*matchHamming_gpu<short>*/,\r
             matchHamming_gpu<int>, 0/*matchHamming_gpu<float>*/\r
         }\r
     };\r
@@ -789,7 +789,8 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchSingle(const GpuMat& query,
     DeviceInfo info;\r
     int cc = info.majorVersion() * 10 + info.minorVersion();\r
 \r
-    CV_Assert(TargetArchs::builtWith(GLOBAL_ATOMICS) && info.supports(GLOBAL_ATOMICS));\r
+    if (!TargetArchs::builtWith(GLOBAL_ATOMICS) || !DeviceInfo().supports(GLOBAL_ATOMICS))\r
+        CV_Error(CV_StsNotImplemented, "The device doesn't support global atomics");\r
 \r
     const int nQuery = query.rows;\r
     const int nTrain = train.rows;\r
@@ -804,19 +805,19 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchSingle(const GpuMat& query,
         ensureSizeIsEnough(nQuery, std::max((nTrain / 100), 10), CV_32SC1, trainIdx);\r
         ensureSizeIsEnough(nQuery, std::max((nTrain / 100), 10), CV_32FC1, distance);\r
     }\r
-    \r
+\r
     if (stream)\r
         stream.enqueueMemSet(nMatches, Scalar::all(0));\r
     else\r
         nMatches.setTo(Scalar::all(0));\r
 \r
     caller_t func = callers[distType][query.depth()];\r
-    CV_Assert(func != 0);    \r
+    CV_Assert(func != 0);\r
 \r
     func(query, train, maxDistance, mask, trainIdx, distance, nMatches, cc, StreamAccessor::getStream(stream));\r
 }\r
 \r
-void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchDownload(const GpuMat& trainIdx, const GpuMat& distance, const GpuMat& nMatches, \r
+void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchDownload(const GpuMat& trainIdx, const GpuMat& distance, const GpuMat& nMatches,\r
     vector< vector<DMatch> >& matches, bool compactResult)\r
 {\r
     if (trainIdx.empty() || distance.empty() || nMatches.empty())\r
@@ -886,33 +887,33 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatch(const GpuMat& query, const
     radiusMatchDownload(trainIdx, distance, nMatches, matches, compactResult);\r
 }\r
 \r
-void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchCollection(const GpuMat& query, GpuMat& trainIdx, GpuMat& imgIdx, GpuMat& distance, GpuMat& nMatches, \r
+void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchCollection(const GpuMat& query, GpuMat& trainIdx, GpuMat& imgIdx, GpuMat& distance, GpuMat& nMatches,\r
     float maxDistance, const vector<GpuMat>& masks, Stream& stream)\r
 {\r
     if (query.empty() || empty())\r
         return;\r
 \r
-    using namespace ::cv::gpu::device::bf_radius_match;\r
+    using namespace cv::gpu::device::bf_radius_match;\r
 \r
-    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db* trains, int n, float maxDistance, const DevMem2Db* masks, \r
-                             const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches, \r
+    typedef void (*caller_t)(const DevMem2Db& query, const DevMem2Db* trains, int n, float maxDistance, const DevMem2Db* masks,\r
+                             const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance, const DevMem2D_<unsigned int>& nMatches,\r
                              int cc, cudaStream_t stream);\r
 \r
     static const caller_t callers[3][6] =\r
     {\r
         {\r
-            matchL1_gpu<unsigned char>, 0/*matchL1_gpu<signed char>*/, \r
-            matchL1_gpu<unsigned short>, matchL1_gpu<short>, \r
+            matchL1_gpu<unsigned char>, 0/*matchL1_gpu<signed char>*/,\r
+            matchL1_gpu<unsigned short>, matchL1_gpu<short>,\r
             matchL1_gpu<int>, matchL1_gpu<float>\r
         },\r
         {\r
-            0/*matchL2_gpu<unsigned char>*/, 0/*matchL2_gpu<signed char>*/, \r
-            0/*matchL2_gpu<unsigned short>*/, 0/*matchL2_gpu<short>*/, \r
+            0/*matchL2_gpu<unsigned char>*/, 0/*matchL2_gpu<signed char>*/,\r
+            0/*matchL2_gpu<unsigned short>*/, 0/*matchL2_gpu<short>*/,\r
             0/*matchL2_gpu<int>*/, matchL2_gpu<float>\r
         },\r
         {\r
-            matchHamming_gpu<unsigned char>, 0/*matchHamming_gpu<signed char>*/, \r
-            matchHamming_gpu<unsigned short>, 0/*matchHamming_gpu<short>*/, \r
+            matchHamming_gpu<unsigned char>, 0/*matchHamming_gpu<signed char>*/,\r
+            matchHamming_gpu<unsigned short>, 0/*matchHamming_gpu<short>*/,\r
             matchHamming_gpu<int>, 0/*matchHamming_gpu<float>*/\r
         }\r
     };\r
@@ -920,7 +921,8 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchCollection(const GpuMat& qu
     DeviceInfo info;\r
     int cc = info.majorVersion() * 10 + info.minorVersion();\r
 \r
-    CV_Assert(TargetArchs::builtWith(GLOBAL_ATOMICS) && info.supports(GLOBAL_ATOMICS));\r
+    if (!TargetArchs::builtWith(GLOBAL_ATOMICS) || !DeviceInfo().supports(GLOBAL_ATOMICS))\r
+        CV_Error(CV_StsNotImplemented, "The device doesn't support global atomics");\r
 \r
     const int nQuery = query.rows;\r
 \r
@@ -934,7 +936,7 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchCollection(const GpuMat& qu
         ensureSizeIsEnough(nQuery, std::max((nQuery / 100), 10), CV_32SC1, imgIdx);\r
         ensureSizeIsEnough(nQuery, std::max((nQuery / 100), 10), CV_32FC1, distance);\r
     }\r
-    \r
+\r
     if (stream)\r
         stream.enqueueMemSet(nMatches, Scalar::all(0));\r
     else\r
@@ -946,7 +948,7 @@ void cv::gpu::BruteForceMatcher_GPU_base::radiusMatchCollection(const GpuMat& qu
     vector<DevMem2Db> trains_(trainDescCollection.begin(), trainDescCollection.end());\r
     vector<DevMem2Db> masks_(masks.begin(), masks.end());\r
 \r
-    func(query, &trains_[0], static_cast<int>(trains_.size()), maxDistance, masks_.size() == 0 ? 0 : &masks_[0], \r
+    func(query, &trains_[0], static_cast<int>(trains_.size()), maxDistance, masks_.size() == 0 ? 0 : &masks_[0],\r
         trainIdx, imgIdx, distance, nMatches, cc, StreamAccessor::getStream(stream));\r
 }\r
 \r

diff --git a/modules/gpu/src/cuda/surf.cu b/modules/gpu/src/cuda/surf.cu
index 76181bc..ac7b6c2 100644 (file)
--- a/modules/gpu/src/cuda/surf.cu
+++ b/modules/gpu/src/cuda/surf.cu
@@ -40,7 +40,7 @@
 //\r
 // Copyright (c) 2010, Paul Furgale, Chi Hay Tong\r
 //\r
-// The original code was written by Paul Furgale and Chi Hay Tong \r
+// The original code was written by Paul Furgale and Chi Hay Tong\r
 // and later optimized and prepared for integration into OpenCV by Itseez.\r
 //\r
 //M*/\r
@@ -52,9 +52,9 @@
 #include "opencv2/gpu/device/functional.hpp"\r
 #include "opencv2/gpu/device/filters.hpp"\r
 \r
-namespace cv { namespace gpu { namespace device \r
+namespace cv { namespace gpu { namespace device\r
 {\r
-    namespace surf \r
+    namespace surf\r
     {\r
         ////////////////////////////////////////////////////////////////////////\r
         // Global parameters\r
@@ -123,7 +123,7 @@ namespace cv { namespace gpu { namespace device
         #endif\r
 \r
             float ratio = (float)newSize / oldSize;\r
-            \r
+\r
             real_t d = 0;\r
 \r
             #pragma unroll\r
@@ -225,7 +225,7 @@ namespace cv { namespace gpu { namespace device
             static __device__ bool check(int sum_i, int sum_j, int size)\r
             {\r
                 float ratio = (float)size / 9.0f;\r
-                \r
+\r
                 float d = 0;\r
 \r
                 int dx1 = __float2int_rn(ratio * c_DM[0]);\r
@@ -423,12 +423,12 @@ namespace cv { namespace gpu { namespace device
                     if (::fabs(x[0]) <= 1.f && ::fabs(x[1]) <= 1.f && ::fabs(x[2]) <= 1.f)\r
                     {\r
                         // if the step is within the interpolation region, perform it\r
-                        \r
+\r
                         const int size = calcSize(c_octave, maxPos.z);\r
 \r
                         const int sum_i = (maxPos.y - ((size >> 1) >> c_octave)) << c_octave;\r
                         const int sum_j = (maxPos.x - ((size >> 1) >> c_octave)) << c_octave;\r
-                        \r
+\r
                         const float center_i = sum_i + (float)(size - 1) / 2;\r
                         const float center_j = sum_j + (float)(size - 1) / 2;\r
 \r
@@ -471,8 +471,8 @@ namespace cv { namespace gpu { namespace device
             #endif\r
         }\r
 \r
-        void icvInterpolateKeypoint_gpu(const PtrStepf& det, const int4* maxPosBuffer, unsigned int maxCounter, \r
-            float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian, \r
+        void icvInterpolateKeypoint_gpu(const PtrStepf& det, const int4* maxPosBuffer, unsigned int maxCounter,\r
+            float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian,\r
             unsigned int* featureCounter)\r
         {\r
             dim3 threads;\r
@@ -509,7 +509,8 @@ namespace cv { namespace gpu { namespace device
             __shared__ float s_Y[128];\r
             __shared__ float s_angle[128];\r
 \r
-            __shared__ float s_sum[32 * 4];\r
+            __shared__ float s_sumx[32 * 4];\r
+            __shared__ float s_sumy[32 * 4];\r
 \r
             /* The sampling intervals and wavelet sized for selecting an orientation\r
              and building the keypoint descriptor are defined relative to 's' */\r
@@ -522,126 +523,109 @@ namespace cv { namespace gpu { namespace device
             const int grad_wav_size = 2 * __float2int_rn(2.0f * s);\r
 \r
             // check when grad_wav_size is too big\r
-            if ((c_img_rows + 1) >= grad_wav_size && (c_img_cols + 1) >= grad_wav_size)\r
-            {\r
-                // Calc X, Y, angle and store it to shared memory\r
-                const int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
-\r
-                float X = 0.0f, Y = 0.0f, angle = 0.0f;\r
+            if ((c_img_rows + 1) < grad_wav_size || (c_img_cols + 1) < grad_wav_size)\r
+                return;\r
 \r
-                if (tid < ORI_SAMPLES)\r
-                {\r
-                    const float margin = (float)(grad_wav_size - 1) / 2.0f;\r
-                    const int x = __float2int_rn(featureX[blockIdx.x] + c_aptX[tid] * s - margin);\r
-                    const int y = __float2int_rn(featureY[blockIdx.x] + c_aptY[tid] * s - margin);\r
+            // Calc X, Y, angle and store it to shared memory\r
+            const int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
 \r
-                    if ((unsigned)y < (unsigned)((c_img_rows + 1) - grad_wav_size) && (unsigned)x < (unsigned)((c_img_cols + 1) - grad_wav_size))\r
-                    {\r
-                        X = c_aptW[tid] * icvCalcHaarPatternSum<2>(c_NX, 4, grad_wav_size, y, x);\r
-                        Y = c_aptW[tid] * icvCalcHaarPatternSum<2>(c_NY, 4, grad_wav_size, y, x);\r
-                    \r
-                        angle = atan2f(Y, X);\r
-                        if (angle < 0)\r
-                            angle += 2.0f * CV_PI_F;\r
-                        angle *= 180.0f / CV_PI_F;\r
-                    }\r
-                }\r
-                s_X[tid] = X;\r
-                s_Y[tid] = Y;\r
-                s_angle[tid] = angle;\r
-                __syncthreads();\r
+            float X = 0.0f, Y = 0.0f, angle = 0.0f;\r
 \r
-                float bestx = 0, besty = 0, best_mod = 0;\r
+            if (tid < ORI_SAMPLES)\r
+            {\r
+                const float margin = (float)(grad_wav_size - 1) / 2.0f;\r
+                const int x = __float2int_rn(featureX[blockIdx.x] + c_aptX[tid] * s - margin);\r
+                const int y = __float2int_rn(featureY[blockIdx.x] + c_aptY[tid] * s - margin);\r
 \r
-                #pragma unroll\r
-                for (int i = 0; i < 18; ++i)\r
+                if (y >= 0 && y < (c_img_rows + 1) - grad_wav_size &&\r
+                    x >= 0 && x < (c_img_cols + 1) - grad_wav_size)\r
                 {\r
-                    const int dir = (i * 4 + threadIdx.y) * ORI_SEARCH_INC;\r
+                    X = c_aptW[tid] * icvCalcHaarPatternSum<2>(c_NX, 4, grad_wav_size, y, x);\r
+                    Y = c_aptW[tid] * icvCalcHaarPatternSum<2>(c_NY, 4, grad_wav_size, y, x);\r
 \r
-                    float sumx = 0.0f, sumy = 0.0f;\r
-                    int d = ::abs(__float2int_rn(s_angle[threadIdx.x]) - dir);\r
-                    if (d < ORI_WIN / 2 || d > 360 - ORI_WIN / 2)\r
-                    {\r
-                        sumx = s_X[threadIdx.x];\r
-                        sumy = s_Y[threadIdx.x];\r
-                    }\r
-                    d = ::abs(__float2int_rn(s_angle[threadIdx.x + 32]) - dir);\r
-                    if (d < ORI_WIN / 2 || d > 360 - ORI_WIN / 2)\r
-                    {\r
-                        sumx += s_X[threadIdx.x + 32];\r
-                        sumy += s_Y[threadIdx.x + 32];\r
-                    }\r
-                    d = ::abs(__float2int_rn(s_angle[threadIdx.x + 64]) - dir);\r
-                    if (d < ORI_WIN / 2 || d > 360 - ORI_WIN / 2)\r
-                    {\r
-                        sumx += s_X[threadIdx.x + 64];\r
-                        sumy += s_Y[threadIdx.x + 64];\r
-                    }\r
-                    d = ::abs(__float2int_rn(s_angle[threadIdx.x + 96]) - dir);\r
-                    if (d < ORI_WIN / 2 || d > 360 - ORI_WIN / 2)\r
-                    {\r
-                        sumx += s_X[threadIdx.x + 96];\r
-                        sumy += s_Y[threadIdx.x + 96];\r
-                    }\r
-\r
-                    float* s_sum_row = s_sum + threadIdx.y * 32;\r
+                    angle = atan2f(Y, X);\r
+                    if (angle < 0)\r
+                        angle += 2.0f * CV_PI_F;\r
+                    angle *= 180.0f / CV_PI_F;\r
+                }\r
+            }\r
+            s_X[tid] = X;\r
+            s_Y[tid] = Y;\r
+            s_angle[tid] = angle;\r
+            __syncthreads();\r
 \r
-                    device::reduce<32>(s_sum_row, sumx, threadIdx.x, plus<volatile float>());\r
-                    device::reduce<32>(s_sum_row, sumy, threadIdx.x, plus<volatile float>());\r
+            float bestx = 0, besty = 0, best_mod = 0;\r
 \r
-                    const float temp_mod = sumx * sumx + sumy * sumy;\r
-                    if (temp_mod > best_mod)\r
-                    {\r
-                        best_mod = temp_mod;\r
-                        bestx = sumx;\r
-                        besty = sumy;\r
-                    }\r
+            #pragma unroll\r
+            for (int i = 0; i < 18; ++i)\r
+            {\r
+                const int dir = (i * 4 + threadIdx.y) * ORI_SEARCH_INC;\r
 \r
-                    __syncthreads();\r
+                float sumx = 0.0f, sumy = 0.0f;\r
+                int d = ::abs(__float2int_rn(s_angle[threadIdx.x]) - dir);\r
+                if (d < ORI_WIN / 2 || d > 360 - ORI_WIN / 2)\r
+                {\r
+                    sumx = s_X[threadIdx.x];\r
+                    sumy = s_Y[threadIdx.x];\r
+                }\r
+                d = ::abs(__float2int_rn(s_angle[threadIdx.x + 32]) - dir);\r
+                if (d < ORI_WIN / 2 || d > 360 - ORI_WIN / 2)\r
+                {\r
+                    sumx += s_X[threadIdx.x + 32];\r
+                    sumy += s_Y[threadIdx.x + 32];\r
+                }\r
+                d = ::abs(__float2int_rn(s_angle[threadIdx.x + 64]) - dir);\r
+                if (d < ORI_WIN / 2 || d > 360 - ORI_WIN / 2)\r
+                {\r
+                    sumx += s_X[threadIdx.x + 64];\r
+                    sumy += s_Y[threadIdx.x + 64];\r
+                }\r
+                d = ::abs(__float2int_rn(s_angle[threadIdx.x + 96]) - dir);\r
+                if (d < ORI_WIN / 2 || d > 360 - ORI_WIN / 2)\r
+                {\r
+                    sumx += s_X[threadIdx.x + 96];\r
+                    sumy += s_Y[threadIdx.x + 96];\r
                 }\r
 \r
-                if (threadIdx.x == 0)\r
+                device::reduce<32>(s_sumx + threadIdx.y * 32, sumx, threadIdx.x, plus<volatile float>());\r
+                device::reduce<32>(s_sumy + threadIdx.y * 32, sumy, threadIdx.x, plus<volatile float>());\r
+\r
+                const float temp_mod = sumx * sumx + sumy * sumy;\r
+                if (temp_mod > best_mod)\r
                 {\r
-                    s_X[threadIdx.y] = bestx;\r
-                    s_Y[threadIdx.y] = besty;\r
-                    s_angle[threadIdx.y] = best_mod;\r
+                    best_mod = temp_mod;\r
+                    bestx = sumx;\r
+                    besty = sumy;\r
                 }\r
+\r
                 __syncthreads();\r
+            }\r
 \r
-                if (threadIdx.x < 2 && threadIdx.y == 0)\r
-                {\r
-                    volatile float* v_x = s_X;\r
-                    volatile float* v_y = s_Y;\r
-                    volatile float* v_mod = s_angle;\r
+            if (threadIdx.x == 0)\r
+            {\r
+                s_X[threadIdx.y] = bestx;\r
+                s_Y[threadIdx.y] = besty;\r
+                s_angle[threadIdx.y] = best_mod;\r
+            }\r
+            __syncthreads();\r
 \r
-                    bestx = v_x[threadIdx.x];\r
-                    besty = v_y[threadIdx.x];\r
-                    best_mod = v_mod[threadIdx.x];\r
+            if (threadIdx.x == 0 && threadIdx.y == 0)\r
+            {\r
+                int bestIdx = 0;\r
 \r
-                    float temp_mod = v_mod[threadIdx.x + 2];\r
-                    if (temp_mod > best_mod)\r
-                    {\r
-                        v_x[threadIdx.x] = bestx = v_x[threadIdx.x + 2];\r
-                        v_y[threadIdx.x] = besty = v_y[threadIdx.x + 2];\r
-                        v_mod[threadIdx.x] = best_mod = temp_mod;\r
-                    }\r
-                    temp_mod = v_mod[threadIdx.x + 1];\r
-                    if (temp_mod > best_mod)\r
-                    {\r
-                        v_x[threadIdx.x] = bestx = v_x[threadIdx.x + 1];\r
-                        v_y[threadIdx.x] = besty = v_y[threadIdx.x + 1];\r
-                    }\r
-                }\r
+                if (s_angle[1] > s_angle[bestIdx])\r
+                    bestIdx = 1;\r
+                if (s_angle[2] > s_angle[bestIdx])\r
+                    bestIdx = 2;\r
+                if (s_angle[3] > s_angle[bestIdx])\r
+                    bestIdx = 3;\r
 \r
-                if (threadIdx.x == 0 && threadIdx.y == 0 && best_mod != 0)\r
-                {\r
-                    float kp_dir = atan2f(besty, bestx);\r
-                    if (kp_dir < 0)\r
-                        kp_dir += 2.0f * CV_PI_F;\r
-                    kp_dir *= 180.0f / CV_PI_F;\r
+                float kp_dir = atan2f(s_Y[bestIdx], s_X[bestIdx]);\r
+                if (kp_dir < 0)\r
+                    kp_dir += 2.0f * CV_PI_F;\r
+                kp_dir *= 180.0f / CV_PI_F;\r
 \r
-                    featureDir[blockIdx.x] = kp_dir;\r
-                }\r
+                featureDir[blockIdx.x] = kp_dir;\r
             }\r
         }\r
 \r
@@ -649,7 +633,7 @@ namespace cv { namespace gpu { namespace device
         #undef ORI_WIN\r
         #undef ORI_SAMPLES\r
 \r
-        void icvCalcOrientation_gpu(const float* featureX, const float* featureY, const float* featureSize, float* featureDir, int nFeatures) \r
+        void icvCalcOrientation_gpu(const float* featureX, const float* featureY, const float* featureSize, float* featureDir, int nFeatures)\r
         {\r
             dim3 threads;\r
             threads.x = 32;\r
@@ -669,27 +653,27 @@ namespace cv { namespace gpu { namespace device
 \r
         #define PATCH_SZ 20\r
 \r
-        __constant__ float c_DW[PATCH_SZ * PATCH_SZ] = \r
+        __constant__ float c_DW[PATCH_SZ * PATCH_SZ] =\r
         {\r
-            3.695352233989979e-006f, 8.444558261544444e-006f, 1.760426494001877e-005f, 3.34794785885606e-005f, 5.808438800158911e-005f, 9.193058212986216e-005f, 0.0001327334757661447f, 0.0001748319627949968f, 0.0002100782439811155f, 0.0002302826324012131f, 0.0002302826324012131f, 0.0002100782439811155f, 0.0001748319627949968f, 0.0001327334757661447f, 9.193058212986216e-005f, 5.808438800158911e-005f, 3.34794785885606e-005f, 1.760426494001877e-005f, 8.444558261544444e-006f, 3.695352233989979e-006f, \r
-            8.444558261544444e-006f, 1.929736572492402e-005f, 4.022897701361217e-005f, 7.650675252079964e-005f, 0.0001327334903180599f, 0.0002100782585330308f, 0.0003033203829545528f, 0.0003995231236331165f, 0.0004800673632416874f, 0.0005262381164357066f, 0.0005262381164357066f, 0.0004800673632416874f, 0.0003995231236331165f, 0.0003033203829545528f, 0.0002100782585330308f, 0.0001327334903180599f, 7.650675252079964e-005f, 4.022897701361217e-005f, 1.929736572492402e-005f, 8.444558261544444e-006f, \r
-            1.760426494001877e-005f, 4.022897701361217e-005f, 8.386484114453197e-005f, 0.0001594926579855382f, 0.0002767078403849155f, 0.0004379475140012801f, 0.0006323281559161842f, 0.0008328808471560478f, 0.001000790391117334f, 0.001097041997127235f, 0.001097041997127235f, 0.001000790391117334f, 0.0008328808471560478f, 0.0006323281559161842f, 0.0004379475140012801f, 0.0002767078403849155f, 0.0001594926579855382f, 8.386484114453197e-005f, 4.022897701361217e-005f, 1.760426494001877e-005f, \r
-            3.34794785885606e-005f, 7.650675252079964e-005f, 0.0001594926579855382f, 0.0003033203247468919f, 0.0005262380582280457f, 0.0008328807889483869f, 0.001202550483867526f, 0.001583957928232849f, 0.001903285388834775f, 0.002086334861814976f, 0.002086334861814976f, 0.001903285388834775f, 0.001583957928232849f, 0.001202550483867526f, 0.0008328807889483869f, 0.0005262380582280457f, 0.0003033203247468919f, 0.0001594926579855382f, 7.650675252079964e-005f, 3.34794785885606e-005f, \r
-            5.808438800158911e-005f, 0.0001327334903180599f, 0.0002767078403849155f, 0.0005262380582280457f, 0.0009129836107604206f, 0.001444985857233405f, 0.002086335094645619f, 0.002748048631474376f, 0.00330205773934722f, 0.003619635012000799f, 0.003619635012000799f, 0.00330205773934722f, 0.002748048631474376f, 0.002086335094645619f, 0.001444985857233405f, 0.0009129836107604206f, 0.0005262380582280457f, 0.0002767078403849155f, 0.0001327334903180599f, 5.808438800158911e-005f, \r
-            9.193058212986216e-005f, 0.0002100782585330308f, 0.0004379475140012801f, 0.0008328807889483869f, 0.001444985857233405f, 0.002286989474669099f, 0.00330205773934722f, 0.004349356517195702f, 0.00522619066759944f, 0.005728822201490402f, 0.005728822201490402f, 0.00522619066759944f, 0.004349356517195702f, 0.00330205773934722f, 0.002286989474669099f, 0.001444985857233405f, 0.0008328807889483869f, 0.0004379475140012801f, 0.0002100782585330308f, 9.193058212986216e-005f, \r
-            0.0001327334757661447f, 0.0003033203829545528f, 0.0006323281559161842f, 0.001202550483867526f, 0.002086335094645619f, 0.00330205773934722f, 0.004767658654600382f, 0.006279794964939356f, 0.007545807864516974f, 0.008271530270576477f, 0.008271530270576477f, 0.007545807864516974f, 0.006279794964939356f, 0.004767658654600382f, 0.00330205773934722f, 0.002086335094645619f, 0.001202550483867526f, 0.0006323281559161842f, 0.0003033203829545528f, 0.0001327334757661447f, \r
-            0.0001748319627949968f, 0.0003995231236331165f, 0.0008328808471560478f, 0.001583957928232849f, 0.002748048631474376f, 0.004349356517195702f, 0.006279794964939356f, 0.008271529339253902f, 0.009939077310264111f, 0.01089497376233339f, 0.01089497376233339f, 0.009939077310264111f, 0.008271529339253902f, 0.006279794964939356f, 0.004349356517195702f, 0.002748048631474376f, 0.001583957928232849f, 0.0008328808471560478f, 0.0003995231236331165f, 0.0001748319627949968f, \r
-            0.0002100782439811155f, 0.0004800673632416874f, 0.001000790391117334f, 0.001903285388834775f, 0.00330205773934722f, 0.00522619066759944f, 0.007545807864516974f, 0.009939077310264111f, 0.01194280479103327f, 0.01309141051024199f, 0.01309141051024199f, 0.01194280479103327f, 0.009939077310264111f, 0.007545807864516974f, 0.00522619066759944f, 0.00330205773934722f, 0.001903285388834775f, 0.001000790391117334f, 0.0004800673632416874f, 0.0002100782439811155f, \r
-            0.0002302826324012131f, 0.0005262381164357066f, 0.001097041997127235f, 0.002086334861814976f, 0.003619635012000799f, 0.005728822201490402f, 0.008271530270576477f, 0.01089497376233339f, 0.01309141051024199f, 0.01435048412531614f, 0.01435048412531614f, 0.01309141051024199f, 0.01089497376233339f, 0.008271530270576477f, 0.005728822201490402f, 0.003619635012000799f, 0.002086334861814976f, 0.001097041997127235f, 0.0005262381164357066f, 0.0002302826324012131f, \r
-            0.0002302826324012131f, 0.0005262381164357066f, 0.001097041997127235f, 0.002086334861814976f, 0.003619635012000799f, 0.005728822201490402f, 0.008271530270576477f, 0.01089497376233339f, 0.01309141051024199f, 0.01435048412531614f, 0.01435048412531614f, 0.01309141051024199f, 0.01089497376233339f, 0.008271530270576477f, 0.005728822201490402f, 0.003619635012000799f, 0.002086334861814976f, 0.001097041997127235f, 0.0005262381164357066f, 0.0002302826324012131f, \r
-            0.0002100782439811155f, 0.0004800673632416874f, 0.001000790391117334f, 0.001903285388834775f, 0.00330205773934722f, 0.00522619066759944f, 0.007545807864516974f, 0.009939077310264111f, 0.01194280479103327f, 0.01309141051024199f, 0.01309141051024199f, 0.01194280479103327f, 0.009939077310264111f, 0.007545807864516974f, 0.00522619066759944f, 0.00330205773934722f, 0.001903285388834775f, 0.001000790391117334f, 0.0004800673632416874f, 0.0002100782439811155f, \r
-            0.0001748319627949968f, 0.0003995231236331165f, 0.0008328808471560478f, 0.001583957928232849f, 0.002748048631474376f, 0.004349356517195702f, 0.006279794964939356f, 0.008271529339253902f, 0.009939077310264111f, 0.01089497376233339f, 0.01089497376233339f, 0.009939077310264111f, 0.008271529339253902f, 0.006279794964939356f, 0.004349356517195702f, 0.002748048631474376f, 0.001583957928232849f, 0.0008328808471560478f, 0.0003995231236331165f, 0.0001748319627949968f, \r
-            0.0001327334757661447f, 0.0003033203829545528f, 0.0006323281559161842f, 0.001202550483867526f, 0.002086335094645619f, 0.00330205773934722f, 0.004767658654600382f, 0.006279794964939356f, 0.007545807864516974f, 0.008271530270576477f, 0.008271530270576477f, 0.007545807864516974f, 0.006279794964939356f, 0.004767658654600382f, 0.00330205773934722f, 0.002086335094645619f, 0.001202550483867526f, 0.0006323281559161842f, 0.0003033203829545528f, 0.0001327334757661447f, \r
-            9.193058212986216e-005f, 0.0002100782585330308f, 0.0004379475140012801f, 0.0008328807889483869f, 0.001444985857233405f, 0.002286989474669099f, 0.00330205773934722f, 0.004349356517195702f, 0.00522619066759944f, 0.005728822201490402f, 0.005728822201490402f, 0.00522619066759944f, 0.004349356517195702f, 0.00330205773934722f, 0.002286989474669099f, 0.001444985857233405f, 0.0008328807889483869f, 0.0004379475140012801f, 0.0002100782585330308f, 9.193058212986216e-005f, \r
-            5.808438800158911e-005f, 0.0001327334903180599f, 0.0002767078403849155f, 0.0005262380582280457f, 0.0009129836107604206f, 0.001444985857233405f, 0.002086335094645619f, 0.002748048631474376f, 0.00330205773934722f, 0.003619635012000799f, 0.003619635012000799f, 0.00330205773934722f, 0.002748048631474376f, 0.002086335094645619f, 0.001444985857233405f, 0.0009129836107604206f, 0.0005262380582280457f, 0.0002767078403849155f, 0.0001327334903180599f, 5.808438800158911e-005f, \r
-            3.34794785885606e-005f, 7.650675252079964e-005f, 0.0001594926579855382f, 0.0003033203247468919f, 0.0005262380582280457f, 0.0008328807889483869f, 0.001202550483867526f, 0.001583957928232849f, 0.001903285388834775f, 0.002086334861814976f, 0.002086334861814976f, 0.001903285388834775f, 0.001583957928232849f, 0.001202550483867526f, 0.0008328807889483869f, 0.0005262380582280457f, 0.0003033203247468919f, 0.0001594926579855382f, 7.650675252079964e-005f, 3.34794785885606e-005f, \r
-            1.760426494001877e-005f, 4.022897701361217e-005f, 8.386484114453197e-005f, 0.0001594926579855382f, 0.0002767078403849155f, 0.0004379475140012801f, 0.0006323281559161842f, 0.0008328808471560478f, 0.001000790391117334f, 0.001097041997127235f, 0.001097041997127235f, 0.001000790391117334f, 0.0008328808471560478f, 0.0006323281559161842f, 0.0004379475140012801f, 0.0002767078403849155f, 0.0001594926579855382f, 8.386484114453197e-005f, 4.022897701361217e-005f, 1.760426494001877e-005f, \r
-            8.444558261544444e-006f, 1.929736572492402e-005f, 4.022897701361217e-005f, 7.650675252079964e-005f, 0.0001327334903180599f, 0.0002100782585330308f, 0.0003033203829545528f, 0.0003995231236331165f, 0.0004800673632416874f, 0.0005262381164357066f, 0.0005262381164357066f, 0.0004800673632416874f, 0.0003995231236331165f, 0.0003033203829545528f, 0.0002100782585330308f, 0.0001327334903180599f, 7.650675252079964e-005f, 4.022897701361217e-005f, 1.929736572492402e-005f, 8.444558261544444e-006f, \r
+            3.695352233989979e-006f, 8.444558261544444e-006f, 1.760426494001877e-005f, 3.34794785885606e-005f, 5.808438800158911e-005f, 9.193058212986216e-005f, 0.0001327334757661447f, 0.0001748319627949968f, 0.0002100782439811155f, 0.0002302826324012131f, 0.0002302826324012131f, 0.0002100782439811155f, 0.0001748319627949968f, 0.0001327334757661447f, 9.193058212986216e-005f, 5.808438800158911e-005f, 3.34794785885606e-005f, 1.760426494001877e-005f, 8.444558261544444e-006f, 3.695352233989979e-006f,\r
+            8.444558261544444e-006f, 1.929736572492402e-005f, 4.022897701361217e-005f, 7.650675252079964e-005f, 0.0001327334903180599f, 0.0002100782585330308f, 0.0003033203829545528f, 0.0003995231236331165f, 0.0004800673632416874f, 0.0005262381164357066f, 0.0005262381164357066f, 0.0004800673632416874f, 0.0003995231236331165f, 0.0003033203829545528f, 0.0002100782585330308f, 0.0001327334903180599f, 7.650675252079964e-005f, 4.022897701361217e-005f, 1.929736572492402e-005f, 8.444558261544444e-006f,\r
+            1.760426494001877e-005f, 4.022897701361217e-005f, 8.386484114453197e-005f, 0.0001594926579855382f, 0.0002767078403849155f, 0.0004379475140012801f, 0.0006323281559161842f, 0.0008328808471560478f, 0.001000790391117334f, 0.001097041997127235f, 0.001097041997127235f, 0.001000790391117334f, 0.0008328808471560478f, 0.0006323281559161842f, 0.0004379475140012801f, 0.0002767078403849155f, 0.0001594926579855382f, 8.386484114453197e-005f, 4.022897701361217e-005f, 1.760426494001877e-005f,\r
+            3.34794785885606e-005f, 7.650675252079964e-005f, 0.0001594926579855382f, 0.0003033203247468919f, 0.0005262380582280457f, 0.0008328807889483869f, 0.001202550483867526f, 0.001583957928232849f, 0.001903285388834775f, 0.002086334861814976f, 0.002086334861814976f, 0.001903285388834775f, 0.001583957928232849f, 0.001202550483867526f, 0.0008328807889483869f, 0.0005262380582280457f, 0.0003033203247468919f, 0.0001594926579855382f, 7.650675252079964e-005f, 3.34794785885606e-005f,\r
+            5.808438800158911e-005f, 0.0001327334903180599f, 0.0002767078403849155f, 0.0005262380582280457f, 0.0009129836107604206f, 0.001444985857233405f, 0.002086335094645619f, 0.002748048631474376f, 0.00330205773934722f, 0.003619635012000799f, 0.003619635012000799f, 0.00330205773934722f, 0.002748048631474376f, 0.002086335094645619f, 0.001444985857233405f, 0.0009129836107604206f, 0.0005262380582280457f, 0.0002767078403849155f, 0.0001327334903180599f, 5.808438800158911e-005f,\r
+            9.193058212986216e-005f, 0.0002100782585330308f, 0.0004379475140012801f, 0.0008328807889483869f, 0.001444985857233405f, 0.002286989474669099f, 0.00330205773934722f, 0.004349356517195702f, 0.00522619066759944f, 0.005728822201490402f, 0.005728822201490402f, 0.00522619066759944f, 0.004349356517195702f, 0.00330205773934722f, 0.002286989474669099f, 0.001444985857233405f, 0.0008328807889483869f, 0.0004379475140012801f, 0.0002100782585330308f, 9.193058212986216e-005f,\r
+            0.0001327334757661447f, 0.0003033203829545528f, 0.0006323281559161842f, 0.001202550483867526f, 0.002086335094645619f, 0.00330205773934722f, 0.004767658654600382f, 0.006279794964939356f, 0.007545807864516974f, 0.008271530270576477f, 0.008271530270576477f, 0.007545807864516974f, 0.006279794964939356f, 0.004767658654600382f, 0.00330205773934722f, 0.002086335094645619f, 0.001202550483867526f, 0.0006323281559161842f, 0.0003033203829545528f, 0.0001327334757661447f,\r
+            0.0001748319627949968f, 0.0003995231236331165f, 0.0008328808471560478f, 0.001583957928232849f, 0.002748048631474376f, 0.004349356517195702f, 0.006279794964939356f, 0.008271529339253902f, 0.009939077310264111f, 0.01089497376233339f, 0.01089497376233339f, 0.009939077310264111f, 0.008271529339253902f, 0.006279794964939356f, 0.004349356517195702f, 0.002748048631474376f, 0.001583957928232849f, 0.0008328808471560478f, 0.0003995231236331165f, 0.0001748319627949968f,\r
+            0.0002100782439811155f, 0.0004800673632416874f, 0.001000790391117334f, 0.001903285388834775f, 0.00330205773934722f, 0.00522619066759944f, 0.007545807864516974f, 0.009939077310264111f, 0.01194280479103327f, 0.01309141051024199f, 0.01309141051024199f, 0.01194280479103327f, 0.009939077310264111f, 0.007545807864516974f, 0.00522619066759944f, 0.00330205773934722f, 0.001903285388834775f, 0.001000790391117334f, 0.0004800673632416874f, 0.0002100782439811155f,\r
+            0.0002302826324012131f, 0.0005262381164357066f, 0.001097041997127235f, 0.002086334861814976f, 0.003619635012000799f, 0.005728822201490402f, 0.008271530270576477f, 0.01089497376233339f, 0.01309141051024199f, 0.01435048412531614f, 0.01435048412531614f, 0.01309141051024199f, 0.01089497376233339f, 0.008271530270576477f, 0.005728822201490402f, 0.003619635012000799f, 0.002086334861814976f, 0.001097041997127235f, 0.0005262381164357066f, 0.0002302826324012131f,\r
+            0.0002302826324012131f, 0.0005262381164357066f, 0.001097041997127235f, 0.002086334861814976f, 0.003619635012000799f, 0.005728822201490402f, 0.008271530270576477f, 0.01089497376233339f, 0.01309141051024199f, 0.01435048412531614f, 0.01435048412531614f, 0.01309141051024199f, 0.01089497376233339f, 0.008271530270576477f, 0.005728822201490402f, 0.003619635012000799f, 0.002086334861814976f, 0.001097041997127235f, 0.0005262381164357066f, 0.0002302826324012131f,\r
+            0.0002100782439811155f, 0.0004800673632416874f, 0.001000790391117334f, 0.001903285388834775f, 0.00330205773934722f, 0.00522619066759944f, 0.007545807864516974f, 0.009939077310264111f, 0.01194280479103327f, 0.01309141051024199f, 0.01309141051024199f, 0.01194280479103327f, 0.009939077310264111f, 0.007545807864516974f, 0.00522619066759944f, 0.00330205773934722f, 0.001903285388834775f, 0.001000790391117334f, 0.0004800673632416874f, 0.0002100782439811155f,\r
+            0.0001748319627949968f, 0.0003995231236331165f, 0.0008328808471560478f, 0.001583957928232849f, 0.002748048631474376f, 0.004349356517195702f, 0.006279794964939356f, 0.008271529339253902f, 0.009939077310264111f, 0.01089497376233339f, 0.01089497376233339f, 0.009939077310264111f, 0.008271529339253902f, 0.006279794964939356f, 0.004349356517195702f, 0.002748048631474376f, 0.001583957928232849f, 0.0008328808471560478f, 0.0003995231236331165f, 0.0001748319627949968f,\r
+            0.0001327334757661447f, 0.0003033203829545528f, 0.0006323281559161842f, 0.001202550483867526f, 0.002086335094645619f, 0.00330205773934722f, 0.004767658654600382f, 0.006279794964939356f, 0.007545807864516974f, 0.008271530270576477f, 0.008271530270576477f, 0.007545807864516974f, 0.006279794964939356f, 0.004767658654600382f, 0.00330205773934722f, 0.002086335094645619f, 0.001202550483867526f, 0.0006323281559161842f, 0.0003033203829545528f, 0.0001327334757661447f,\r
+            9.193058212986216e-005f, 0.0002100782585330308f, 0.0004379475140012801f, 0.0008328807889483869f, 0.001444985857233405f, 0.002286989474669099f, 0.00330205773934722f, 0.004349356517195702f, 0.00522619066759944f, 0.005728822201490402f, 0.005728822201490402f, 0.00522619066759944f, 0.004349356517195702f, 0.00330205773934722f, 0.002286989474669099f, 0.001444985857233405f, 0.0008328807889483869f, 0.0004379475140012801f, 0.0002100782585330308f, 9.193058212986216e-005f,\r
+            5.808438800158911e-005f, 0.0001327334903180599f, 0.0002767078403849155f, 0.0005262380582280457f, 0.0009129836107604206f, 0.001444985857233405f, 0.002086335094645619f, 0.002748048631474376f, 0.00330205773934722f, 0.003619635012000799f, 0.003619635012000799f, 0.00330205773934722f, 0.002748048631474376f, 0.002086335094645619f, 0.001444985857233405f, 0.0009129836107604206f, 0.0005262380582280457f, 0.0002767078403849155f, 0.0001327334903180599f, 5.808438800158911e-005f,\r
+            3.34794785885606e-005f, 7.650675252079964e-005f, 0.0001594926579855382f, 0.0003033203247468919f, 0.0005262380582280457f, 0.0008328807889483869f, 0.001202550483867526f, 0.001583957928232849f, 0.001903285388834775f, 0.002086334861814976f, 0.002086334861814976f, 0.001903285388834775f, 0.001583957928232849f, 0.001202550483867526f, 0.0008328807889483869f, 0.0005262380582280457f, 0.0003033203247468919f, 0.0001594926579855382f, 7.650675252079964e-005f, 3.34794785885606e-005f,\r
+            1.760426494001877e-005f, 4.022897701361217e-005f, 8.386484114453197e-005f, 0.0001594926579855382f, 0.0002767078403849155f, 0.0004379475140012801f, 0.0006323281559161842f, 0.0008328808471560478f, 0.001000790391117334f, 0.001097041997127235f, 0.001097041997127235f, 0.001000790391117334f, 0.0008328808471560478f, 0.0006323281559161842f, 0.0004379475140012801f, 0.0002767078403849155f, 0.0001594926579855382f, 8.386484114453197e-005f, 4.022897701361217e-005f, 1.760426494001877e-005f,\r
+            8.444558261544444e-006f, 1.929736572492402e-005f, 4.022897701361217e-005f, 7.650675252079964e-005f, 0.0001327334903180599f, 0.0002100782585330308f, 0.0003033203829545528f, 0.0003995231236331165f, 0.0004800673632416874f, 0.0005262381164357066f, 0.0005262381164357066f, 0.0004800673632416874f, 0.0003995231236331165f, 0.0003033203829545528f, 0.0002100782585330308f, 0.0001327334903180599f, 7.650675252079964e-005f, 4.022897701361217e-005f, 1.929736572492402e-005f, 8.444558261544444e-006f,\r
             3.695352233989979e-006f, 8.444558261544444e-006f, 1.760426494001877e-005f, 3.34794785885606e-005f, 5.808438800158911e-005f, 9.193058212986216e-005f, 0.0001327334757661447f, 0.0001748319627949968f, 0.0002100782439811155f, 0.0002302826324012131f, 0.0002302826324012131f, 0.0002100782439811155f, 0.0001748319627949968f, 0.0001327334757661447f, 9.193058212986216e-005f, 5.808438800158911e-005f, 3.34794785885606e-005f, 1.760426494001877e-005f, 8.444558261544444e-006f, 3.695352233989979e-006f\r
         };\r
 \r
@@ -697,7 +681,7 @@ namespace cv { namespace gpu { namespace device
         {\r
             typedef uchar elem_type;\r
 \r
-            __device__ __forceinline__ WinReader(float centerX_, float centerY_, float win_offset_, float cos_dir_, float sin_dir_) : \r
+            __device__ __forceinline__ WinReader(float centerX_, float centerY_, float win_offset_, float cos_dir_, float sin_dir_) :\r
                 centerX(centerX_), centerY(centerY_), win_offset(win_offset_), cos_dir(cos_dir_), sin_dir(sin_dir_)\r
             {\r
             }\r
@@ -710,14 +694,14 @@ namespace cv { namespace gpu { namespace device
                 return tex2D(imgTex, pixel_x, pixel_y);\r
             }\r
 \r
-            float centerX; \r
+            float centerX;\r
             float centerY;\r
-            float win_offset; \r
-            float cos_dir; \r
+            float win_offset;\r
+            float cos_dir;\r
             float sin_dir;\r
         };\r
 \r
-        __device__ void calc_dx_dy(float s_dx_bin[25], float s_dy_bin[25], \r
+        __device__ void calc_dx_dy(float s_dx_bin[25], float s_dy_bin[25],\r
             const float* featureX, const float* featureY, const float* featureSize, const float* featureDir)\r
         {\r
             __shared__ float s_PATCH[6][6];\r
@@ -739,7 +723,7 @@ namespace cv { namespace gpu { namespace device
             sincosf(descriptor_dir, &sin_dir, &cos_dir);\r
 \r
             /* Nearest neighbour version (faster) */\r
-            const float win_offset = -(float)(win_size - 1) / 2; \r
+            const float win_offset = -(float)(win_size - 1) / 2;\r
 \r
             // Compute sampling points\r
             // since grids are 2D, need to compute xBlock and yBlock indices\r
@@ -966,11 +950,11 @@ namespace cv { namespace gpu { namespace device
             descriptor_base[threadIdx.x] = lookup / len;\r
         }\r
 \r
-        void compute_descriptors_gpu(const DevMem2Df& descriptors, \r
+        void compute_descriptors_gpu(const DevMem2Df& descriptors,\r
             const float* featureX, const float* featureY, const float* featureSize, const float* featureDir, int nFeatures)\r
         {\r
             // compute unnormalized descriptors, then normalize them - odd indexing since grid must be 2D\r
-            \r
+\r
             if (descriptors.cols == 64)\r
             {\r
                 compute_descriptors64<<<dim3(nFeatures, 16, 1), dim3(6, 6, 1)>>>(descriptors, featureX, featureY, featureSize, featureDir);\r
@@ -985,12 +969,12 @@ namespace cv { namespace gpu { namespace device
             }\r
             else\r
             {\r
-                compute_descriptors128<<<dim3(nFeatures, 16, 1), dim3(6, 6, 1)>>>(descriptors, featureX, featureY, featureSize, featureDir);            \r
+                compute_descriptors128<<<dim3(nFeatures, 16, 1), dim3(6, 6, 1)>>>(descriptors, featureX, featureY, featureSize, featureDir);\r
                 cudaSafeCall( cudaGetLastError() );\r
 \r
                 cudaSafeCall( cudaDeviceSynchronize() );\r
 \r
-                normalize_descriptors<128><<<dim3(nFeatures, 1, 1), dim3(128, 1, 1)>>>(descriptors);            \r
+                normalize_descriptors<128><<<dim3(nFeatures, 1, 1), dim3(128, 1, 1)>>>(descriptors);\r
                 cudaSafeCall( cudaGetLastError() );\r
 \r
                 cudaSafeCall( cudaDeviceSynchronize() );\r

diff --git a/modules/gpu/src/fast.cpp b/modules/gpu/src/fast.cpp
index 5a7c73b..270548d 100644 (file)
--- a/modules/gpu/src/fast.cpp
+++ b/modules/gpu/src/fast.cpp
@@ -59,7 +59,7 @@ int cv::gpu::FAST_GPU::getKeyPoints(GpuMat&) { throw_nogpu(); return 0; }
 \r
 #else /* !defined (HAVE_CUDA) */\r
 \r
-cv::gpu::FAST_GPU::FAST_GPU(int _threshold, bool _nonmaxSupression, double _keypointsRatio) : \r
+cv::gpu::FAST_GPU::FAST_GPU(int _threshold, bool _nonmaxSupression, double _keypointsRatio) :\r
     nonmaxSupression(_nonmaxSupression), threshold(_threshold), keypointsRatio(_keypointsRatio), count_(0)\r
 {\r
 }\r
@@ -109,9 +109,9 @@ void cv::gpu::FAST_GPU::operator ()(const GpuMat& img, const GpuMat& mask, GpuMa
     keypoints.cols = getKeyPoints(keypoints);\r
 }\r
 \r
-namespace cv { namespace gpu { namespace device \r
+namespace cv { namespace gpu { namespace device\r
 {\r
-    namespace fast \r
+    namespace fast\r
     {\r
         int calcKeypoints_gpu(DevMem2Db img, DevMem2Db mask, short2* kpLoc, int maxKeypoints, DevMem2Di score, int threshold);\r
         int nonmaxSupression_gpu(const short2* kpLoc, int count, DevMem2Di score, short2* loc, float* response);\r
@@ -124,7 +124,9 @@ int cv::gpu::FAST_GPU::calcKeyPointsLocation(const GpuMat& img, const GpuMat& ma
 \r
     CV_Assert(img.type() == CV_8UC1);\r
     CV_Assert(mask.empty() || (mask.type() == CV_8UC1 && mask.size() == img.size()));\r
-    CV_Assert(TargetArchs::builtWith(GLOBAL_ATOMICS) && DeviceInfo().supports(GLOBAL_ATOMICS));\r
+\r
+    if (!TargetArchs::builtWith(GLOBAL_ATOMICS) || !DeviceInfo().supports(GLOBAL_ATOMICS))\r
+        CV_Error(CV_StsNotImplemented, "The device doesn't support global atomics");\r
 \r
     int maxKeypoints = static_cast<int>(keypointsRatio * img.size().area());\r
 \r
@@ -146,7 +148,8 @@ int cv::gpu::FAST_GPU::getKeyPoints(GpuMat& keypoints)
 {\r
     using namespace cv::gpu::device::fast;\r
 \r
-    CV_Assert(TargetArchs::builtWith(GLOBAL_ATOMICS) && DeviceInfo().supports(GLOBAL_ATOMICS));\r
+    if (!TargetArchs::builtWith(GLOBAL_ATOMICS) || !DeviceInfo().supports(GLOBAL_ATOMICS))\r
+        CV_Error(CV_StsNotImplemented, "The device doesn't support global atomics");\r
 \r
     if (count_ == 0)\r
         return 0;\r
@@ -160,7 +163,7 @@ int cv::gpu::FAST_GPU::getKeyPoints(GpuMat& keypoints)
     kpLoc_.colRange(0, count_).copyTo(locRow);\r
     keypoints.row(1).setTo(Scalar::all(0));\r
 \r
-    return count_;    \r
+    return count_;\r
 }\r
 \r
 void cv::gpu::FAST_GPU::release()\r

diff --git a/modules/gpu/src/surf.cpp b/modules/gpu/src/surf.cpp
index dea51a9..9a9efc4 100644 (file)
--- a/modules/gpu/src/surf.cpp
+++ b/modules/gpu/src/surf.cpp
@@ -120,7 +120,9 @@ namespace
             CV_Assert(!img.empty() && img.type() == CV_8UC1);\r
             CV_Assert(mask.empty() || (mask.size() == img.size() && mask.type() == CV_8UC1));\r
             CV_Assert(surf_.nOctaves > 0 && surf_.nOctaveLayers > 0);\r
-            CV_Assert(TargetArchs::builtWith(GLOBAL_ATOMICS) && DeviceInfo().supports(GLOBAL_ATOMICS));\r
+\r
+            if (!TargetArchs::builtWith(GLOBAL_ATOMICS) || !DeviceInfo().supports(GLOBAL_ATOMICS))\r
+                CV_Error(CV_StsNotImplemented, "The device doesn't support global atomics");\r
 \r
             const int min_size = calcSize(surf_.nOctaves - 1, 0);\r
             CV_Assert(img_rows - min_size >= 0);\r
@@ -184,8 +186,8 @@ namespace
                 {\r
                     icvInterpolateKeypoint_gpu(surf_.det, surf_.maxPosBuffer.ptr<int4>(), maxCounter,\r
                         keypoints.ptr<float>(SURF_GPU::X_ROW), keypoints.ptr<float>(SURF_GPU::Y_ROW),\r
-                        keypoints.ptr<int>(SURF_GPU::LAPLACIAN_ROW), keypoints.ptr<int>(SURF_GPU::OCTAVE_ROW), \r
-                        keypoints.ptr<float>(SURF_GPU::SIZE_ROW), keypoints.ptr<float>(SURF_GPU::HESSIAN_ROW), \r
+                        keypoints.ptr<int>(SURF_GPU::LAPLACIAN_ROW), keypoints.ptr<int>(SURF_GPU::OCTAVE_ROW),\r
+                        keypoints.ptr<float>(SURF_GPU::SIZE_ROW), keypoints.ptr<float>(SURF_GPU::HESSIAN_ROW),\r
                         counters.ptr<unsigned int>());\r
                 }\r
             }\r
@@ -306,7 +308,7 @@ void cv::gpu::SURF_GPU::downloadKeypoints(const GpuMat& keypointsGPU, vector<Key
         Mat keypointsCPU(keypointsGPU);\r
 \r
         keypoints.resize(nFeatures);\r
-        \r
+\r
         float* kp_x = keypointsCPU.ptr<float>(SURF_GPU::X_ROW);\r
         float* kp_y = keypointsCPU.ptr<float>(SURF_GPU::Y_ROW);\r
         int* kp_laplacian = keypointsCPU.ptr<int>(SURF_GPU::LAPLACIAN_ROW);\r

diff --git a/modules/gpu/test/test_features2d.cpp b/modules/gpu/test/test_features2d.cpp
index 2121820..9af8ad9 100644 (file)
--- a/modules/gpu/test/test_features2d.cpp
+++ b/modules/gpu/test/test_features2d.cpp
@@ -108,6 +108,25 @@ testing::AssertionResult assertKeyPointsEquals(const char* gold_expr, const char
 \r
 #define ASSERT_KEYPOINTS_EQ(gold, actual) EXPECT_PRED_FORMAT2(assertKeyPointsEquals, gold, actual);\r
 \r
+int getMatchedPointsCount(std::vector<cv::KeyPoint>& gold, std::vector<cv::KeyPoint>& actual)\r
+{\r
+    std::sort(actual.begin(), actual.end(), KeyPointLess());\r
+    std::sort(gold.begin(), gold.end(), KeyPointLess());\r
+\r
+    int validCount = 0;\r
+\r
+    for (size_t i = 0; i < gold.size(); ++i)\r
+    {\r
+        const cv::KeyPoint& p1 = gold[i];\r
+        const cv::KeyPoint& p2 = actual[i];\r
+\r
+        if (keyPointsEquals(p1, p2))\r
+            ++validCount;\r
+    }\r
+\r
+    return validCount;\r
+}\r
+\r
 int getMatchedPointsCount(const std::vector<cv::KeyPoint>& keypoints1, const std::vector<cv::KeyPoint>& keypoints2, const std::vector<cv::DMatch>& matches)\r
 {\r
     int validCount = 0;\r
@@ -170,20 +189,39 @@ TEST_P(SURF, Detector)
     surf.upright = upright;\r
     surf.keypointsRatio = 0.05f;\r
 \r
-    std::vector<cv::KeyPoint> keypoints;\r
-    surf(loadMat(image), cv::gpu::GpuMat(), keypoints);\r
+    if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS))\r
+    {\r
+        try\r
+        {\r
+            std::vector<cv::KeyPoint> keypoints;\r
+            surf(loadMat(image), cv::gpu::GpuMat(), keypoints);\r
+        }\r
+        catch (const cv::Exception& e)\r
+        {\r
+            ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+        }\r
+    }\r
+    else\r
+    {\r
+        std::vector<cv::KeyPoint> keypoints;\r
+        surf(loadMat(image), cv::gpu::GpuMat(), keypoints);\r
 \r
-    cv::SURF surf_gold;\r
-    surf_gold.hessianThreshold = hessianThreshold;\r
-    surf_gold.nOctaves = nOctaves;\r
-    surf_gold.nOctaveLayers = nOctaveLayers;\r
-    surf_gold.extended = extended;\r
-    surf_gold.upright = upright;\r
+        cv::SURF surf_gold;\r
+        surf_gold.hessianThreshold = hessianThreshold;\r
+        surf_gold.nOctaves = nOctaves;\r
+        surf_gold.nOctaveLayers = nOctaveLayers;\r
+        surf_gold.extended = extended;\r
+        surf_gold.upright = upright;\r
 \r
-    std::vector<cv::KeyPoint> keypoints_gold;\r
-    surf_gold(image, cv::noArray(), keypoints_gold);\r
+        std::vector<cv::KeyPoint> keypoints_gold;\r
+        surf_gold(image, cv::noArray(), keypoints_gold);\r
 \r
-    ASSERT_KEYPOINTS_EQ(keypoints_gold, keypoints);\r
+        ASSERT_EQ(keypoints_gold.size(), keypoints.size());\r
+        int matchedCount = getMatchedPointsCount(keypoints_gold, keypoints);\r
+        double matchedRatio = static_cast<double>(matchedCount) / keypoints_gold.size();\r
+\r
+        EXPECT_GT(matchedRatio, 0.95);\r
+    }\r
 }\r
 \r
 TEST_P(SURF, Detector_Masked)\r
@@ -202,20 +240,39 @@ TEST_P(SURF, Detector_Masked)
     surf.upright = upright;\r
     surf.keypointsRatio = 0.05f;\r
 \r
-    std::vector<cv::KeyPoint> keypoints;\r
-    surf(loadMat(image), loadMat(mask), keypoints);\r
+    if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS))\r
+    {\r
+        try\r
+        {\r
+            std::vector<cv::KeyPoint> keypoints;\r
+            surf(loadMat(image), loadMat(mask), keypoints);\r
+        }\r
+        catch (const cv::Exception& e)\r
+        {\r
+            ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+        }\r
+    }\r
+    else\r
+    {\r
+        std::vector<cv::KeyPoint> keypoints;\r
+        surf(loadMat(image), loadMat(mask), keypoints);\r
+\r
+        cv::SURF surf_gold;\r
+        surf_gold.hessianThreshold = hessianThreshold;\r
+        surf_gold.nOctaves = nOctaves;\r
+        surf_gold.nOctaveLayers = nOctaveLayers;\r
+        surf_gold.extended = extended;\r
+        surf_gold.upright = upright;\r
 \r
-    cv::SURF surf_gold;\r
-    surf_gold.hessianThreshold = hessianThreshold;\r
-    surf_gold.nOctaves = nOctaves;\r
-    surf_gold.nOctaveLayers = nOctaveLayers;\r
-    surf_gold.extended = extended;\r
-    surf_gold.upright = upright;\r
+        std::vector<cv::KeyPoint> keypoints_gold;\r
+        surf_gold(image, mask, keypoints_gold);\r
 \r
-    std::vector<cv::KeyPoint> keypoints_gold;\r
-    surf_gold(image, mask, keypoints_gold);\r
+        ASSERT_EQ(keypoints_gold.size(), keypoints.size());\r
+        int matchedCount = getMatchedPointsCount(keypoints_gold, keypoints);\r
+        double matchedRatio = static_cast<double>(matchedCount) / keypoints_gold.size();\r
 \r
-    ASSERT_KEYPOINTS_EQ(keypoints_gold, keypoints);\r
+        EXPECT_GT(matchedRatio, 0.95);\r
+    }\r
 }\r
 \r
 TEST_P(SURF, Descriptor)\r
@@ -238,23 +295,39 @@ TEST_P(SURF, Descriptor)
     surf_gold.extended = extended;\r
     surf_gold.upright = upright;\r
 \r
-    std::vector<cv::KeyPoint> keypoints;\r
-    surf_gold(image, cv::noArray(), keypoints);\r
+    if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS))\r
+    {\r
+        try\r
+        {\r
+            std::vector<cv::KeyPoint> keypoints;\r
+            cv::gpu::GpuMat descriptors;\r
+            surf(loadMat(image), cv::gpu::GpuMat(), keypoints, descriptors);\r
+        }\r
+        catch (const cv::Exception& e)\r
+        {\r
+            ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+        }\r
+    }\r
+    else\r
+    {\r
+        std::vector<cv::KeyPoint> keypoints;\r
+        surf_gold(image, cv::noArray(), keypoints);\r
 \r
-    cv::gpu::GpuMat descriptors;\r
-    surf(loadMat(image), cv::gpu::GpuMat(), keypoints, descriptors, true);\r
+        cv::gpu::GpuMat descriptors;\r
+        surf(loadMat(image), cv::gpu::GpuMat(), keypoints, descriptors, true);\r
 \r
-    cv::Mat descriptors_gold;\r
-    surf_gold(image, cv::noArray(), keypoints, descriptors_gold, true);\r
+        cv::Mat descriptors_gold;\r
+        surf_gold(image, cv::noArray(), keypoints, descriptors_gold, true);\r
 \r
-    cv::BFMatcher matcher(cv::NORM_L2);\r
-    std::vector<cv::DMatch> matches;\r
-    matcher.match(descriptors_gold, cv::Mat(descriptors), matches);\r
+        cv::BFMatcher matcher(cv::NORM_L2);\r
+        std::vector<cv::DMatch> matches;\r
+        matcher.match(descriptors_gold, cv::Mat(descriptors), matches);\r
 \r
-    int matchedCount = getMatchedPointsCount(keypoints, keypoints, matches);\r
-    double matchedRatio = static_cast<double>(matchedCount) / keypoints.size();\r
+        int matchedCount = getMatchedPointsCount(keypoints, keypoints, matches);\r
+        double matchedRatio = static_cast<double>(matchedCount) / keypoints.size();\r
 \r
-    EXPECT_GT(matchedRatio, 0.35);\r
+        EXPECT_GT(matchedRatio, 0.35);\r
+    }\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Features2D, SURF, testing::Combine(\r
@@ -295,13 +368,28 @@ TEST_P(FAST, Accuracy)
     cv::gpu::FAST_GPU fast(threshold);\r
     fast.nonmaxSupression = nonmaxSupression;\r
 \r
-    std::vector<cv::KeyPoint> keypoints;\r
-    fast(loadMat(image), cv::gpu::GpuMat(), keypoints);\r
+    if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS))\r
+    {\r
+        try\r
+        {\r
+            std::vector<cv::KeyPoint> keypoints;\r
+            fast(loadMat(image), cv::gpu::GpuMat(), keypoints);\r
+        }\r
+        catch (const cv::Exception& e)\r
+        {\r
+            ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+        }\r
+    }\r
+    else\r
+    {\r
+        std::vector<cv::KeyPoint> keypoints;\r
+        fast(loadMat(image), cv::gpu::GpuMat(), keypoints);\r
 \r
-    std::vector<cv::KeyPoint> keypoints_gold;\r
-    cv::FAST(image, keypoints_gold, threshold, nonmaxSupression);\r
+        std::vector<cv::KeyPoint> keypoints_gold;\r
+        cv::FAST(image, keypoints_gold, threshold, nonmaxSupression);\r
 \r
-    ASSERT_KEYPOINTS_EQ(keypoints_gold, keypoints);\r
+        ASSERT_KEYPOINTS_EQ(keypoints_gold, keypoints);\r
+    }\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Features2D, FAST, testing::Combine(\r
@@ -364,24 +452,40 @@ TEST_P(ORB, Accuracy)
     cv::gpu::ORB_GPU orb(nFeatures, scaleFactor, nLevels, edgeThreshold, firstLevel, WTA_K, scoreType, patchSize);\r
     orb.blurForDescriptor = blurForDescriptor;\r
 \r
-    std::vector<cv::KeyPoint> keypoints;\r
-    cv::gpu::GpuMat descriptors;\r
-    orb(loadMat(image), loadMat(mask), keypoints, descriptors);\r
+    if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS))\r
+    {\r
+        try\r
+        {\r
+            std::vector<cv::KeyPoint> keypoints;\r
+            cv::gpu::GpuMat descriptors;\r
+            orb(loadMat(image), loadMat(mask), keypoints, descriptors);\r
+        }\r
+        catch (const cv::Exception& e)\r
+        {\r
+            ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+        }\r
+    }\r
+    else\r
+    {\r
+        std::vector<cv::KeyPoint> keypoints;\r
+        cv::gpu::GpuMat descriptors;\r
+        orb(loadMat(image), loadMat(mask), keypoints, descriptors);\r
 \r
-    cv::ORB orb_gold(nFeatures, scaleFactor, nLevels, edgeThreshold, firstLevel, WTA_K, scoreType, patchSize);\r
+        cv::ORB orb_gold(nFeatures, scaleFactor, nLevels, edgeThreshold, firstLevel, WTA_K, scoreType, patchSize);\r
 \r
-    std::vector<cv::KeyPoint> keypoints_gold;\r
-    cv::Mat descriptors_gold;\r
-    orb_gold(image, mask, keypoints_gold, descriptors_gold);\r
+        std::vector<cv::KeyPoint> keypoints_gold;\r
+        cv::Mat descriptors_gold;\r
+        orb_gold(image, mask, keypoints_gold, descriptors_gold);\r
 \r
-    cv::BFMatcher matcher(cv::NORM_HAMMING);\r
-    std::vector<cv::DMatch> matches;\r
-    matcher.match(descriptors_gold, cv::Mat(descriptors), matches);\r
+        cv::BFMatcher matcher(cv::NORM_HAMMING);\r
+        std::vector<cv::DMatch> matches;\r
+        matcher.match(descriptors_gold, cv::Mat(descriptors), matches);\r
 \r
-    int matchedCount = getMatchedPointsCount(keypoints_gold, keypoints, matches);\r
-    double matchedRatio = static_cast<double>(matchedCount) / keypoints.size();\r
+        int matchedCount = getMatchedPointsCount(keypoints_gold, keypoints, matches);\r
+        double matchedRatio = static_cast<double>(matchedCount) / keypoints.size();\r
 \r
-    EXPECT_GT(matchedRatio, 0.35);\r
+        EXPECT_GT(matchedRatio, 0.35);\r
+    }\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Features2D, ORB,  testing::Combine(\r
@@ -713,25 +817,40 @@ TEST_P(BruteForceMatcher, RadiusMatch)
 \r
     cv::gpu::BruteForceMatcher_GPU_base matcher(distType);\r
 \r
-    std::vector< std::vector<cv::DMatch> > matches;\r
-    matcher.radiusMatch(loadMat(query), loadMat(train), matches, radius);\r
+    if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS))\r
+    {\r
+        try\r
+        {\r
+            std::vector< std::vector<cv::DMatch> > matches;\r
+            matcher.radiusMatch(loadMat(query), loadMat(train), matches, radius);\r
+        }\r
+        catch (const cv::Exception& e)\r
+        {\r
+            ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+        }\r
+    }\r
+    else\r
+    {\r
+        std::vector< std::vector<cv::DMatch> > matches;\r
+        matcher.radiusMatch(loadMat(query), loadMat(train), matches, radius);\r
 \r
-    ASSERT_EQ(static_cast<size_t>(queryDescCount), matches.size());\r
+        ASSERT_EQ(static_cast<size_t>(queryDescCount), matches.size());\r
 \r
-    int badCount = 0;\r
-    for (size_t i = 0; i < matches.size(); i++)\r
-    {\r
-        if ((int)matches[i].size() != 1)\r
-            badCount++;\r
-        else\r
+        int badCount = 0;\r
+        for (size_t i = 0; i < matches.size(); i++)\r
         {\r
-            cv::DMatch match = matches[i][0];\r
-            if ((match.queryIdx != (int)i) || (match.trainIdx != (int)i*countFactor) || (match.imgIdx != 0))\r
+            if ((int)matches[i].size() != 1)\r
                 badCount++;\r
+            else\r
+            {\r
+                cv::DMatch match = matches[i][0];\r
+                if ((match.queryIdx != (int)i) || (match.trainIdx != (int)i*countFactor) || (match.imgIdx != 0))\r
+                    badCount++;\r
+            }\r
         }\r
-    }\r
 \r
-    ASSERT_EQ(0, badCount);\r
+        ASSERT_EQ(0, badCount);\r
+    }\r
 }\r
 \r
 TEST_P(BruteForceMatcher, RadiusMatchAdd)\r
@@ -756,42 +875,57 @@ TEST_P(BruteForceMatcher, RadiusMatchAdd)
             masks[mi].col(di * countFactor).setTo(cv::Scalar::all(0));\r
     }\r
 \r
-    std::vector< std::vector<cv::DMatch> > matches;\r
-    matcher.radiusMatch(cv::gpu::GpuMat(query), matches, radius, masks);\r
+    if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS))\r
+    {\r
+        try\r
+        {\r
+            std::vector< std::vector<cv::DMatch> > matches;\r
+            matcher.radiusMatch(cv::gpu::GpuMat(query), matches, radius, masks);\r
+        }\r
+        catch (const cv::Exception& e)\r
+        {\r
+            ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+        }\r
+    }\r
+    else\r
+    {\r
+        std::vector< std::vector<cv::DMatch> > matches;\r
+        matcher.radiusMatch(cv::gpu::GpuMat(query), matches, radius, masks);\r
 \r
-    ASSERT_EQ(static_cast<size_t>(queryDescCount), matches.size());\r
+        ASSERT_EQ(static_cast<size_t>(queryDescCount), matches.size());\r
 \r
-    int badCount = 0;\r
-    int shift = matcher.isMaskSupported() ? 1 : 0;\r
-    int needMatchCount = matcher.isMaskSupported() ? n-1 : n;\r
-    for (size_t i = 0; i < matches.size(); i++)\r
-    {\r
-        if ((int)matches[i].size() != needMatchCount)\r
-            badCount++;\r
-        else\r
+        int badCount = 0;\r
+        int shift = matcher.isMaskSupported() ? 1 : 0;\r
+        int needMatchCount = matcher.isMaskSupported() ? n-1 : n;\r
+        for (size_t i = 0; i < matches.size(); i++)\r
         {\r
-            int localBadCount = 0;\r
-            for (int k = 0; k < needMatchCount; k++)\r
+            if ((int)matches[i].size() != needMatchCount)\r
+                badCount++;\r
+            else\r
             {\r
-                cv::DMatch match = matches[i][k];\r
+                int localBadCount = 0;\r
+                for (int k = 0; k < needMatchCount; k++)\r
                 {\r
-                    if ((int)i < queryDescCount / 2)\r
+                    cv::DMatch match = matches[i][k];\r
                     {\r
-                        if ((match.queryIdx != (int)i) || (match.trainIdx != (int)i * countFactor + k + shift) || (match.imgIdx != 0) )\r
-                            localBadCount++;\r
-                    }\r
-                    else\r
-                    {\r
-                        if ((match.queryIdx != (int)i) || (match.trainIdx != ((int)i - queryDescCount / 2) * countFactor + k + shift) || (match.imgIdx != 1) )\r
-                            localBadCount++;\r
+                        if ((int)i < queryDescCount / 2)\r
+                        {\r
+                            if ((match.queryIdx != (int)i) || (match.trainIdx != (int)i * countFactor + k + shift) || (match.imgIdx != 0) )\r
+                                localBadCount++;\r
+                        }\r
+                        else\r
+                        {\r
+                            if ((match.queryIdx != (int)i) || (match.trainIdx != ((int)i - queryDescCount / 2) * countFactor + k + shift) || (match.imgIdx != 1) )\r
+                                localBadCount++;\r
+                        }\r
                     }\r
                 }\r
+                badCount += localBadCount > 0 ? 1 : 0;\r
             }\r
-            badCount += localBadCount > 0 ? 1 : 0;\r
         }\r
-    }\r
 \r
-    ASSERT_EQ(0, badCount);\r
+        ASSERT_EQ(0, badCount);\r
+    }\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Features2D, BruteForceMatcher, testing::Combine(\r

diff --git a/modules/gpu/test/test_filters.cpp b/modules/gpu/test/test_filters.cpp
index 6b9f23c..66325c3 100644 (file)
--- a/modules/gpu/test/test_filters.cpp
+++ b/modules/gpu/test/test_filters.cpp
@@ -258,13 +258,28 @@ TEST_P(GaussianBlur, Accuracy)
     double sigma1 = randomDouble(0.1, 1.0);\r
     double sigma2 = randomDouble(0.1, 1.0);\r
 \r
-    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
-    cv::gpu::GaussianBlur(loadMat(src, useRoi), dst, ksize, sigma1, sigma2, borderType);\r
+    if (ksize.height > 16 && !supportFeature(devInfo, cv::gpu::FEATURE_SET_COMPUTE_20))\r
+    {\r
+        try\r
+        {\r
+            cv::gpu::GpuMat dst;\r
+            cv::gpu::GaussianBlur(loadMat(src), dst, ksize, sigma1, sigma2, borderType);\r
+        }\r
+        catch (const cv::Exception& e)\r
+        {\r
+            ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+        }\r
+    }\r
+    else\r
+    {\r
+        cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
+        cv::gpu::GaussianBlur(loadMat(src, useRoi), dst, ksize, sigma1, sigma2, borderType);\r
 \r
-    cv::Mat dst_gold;\r
-    cv::GaussianBlur(src, dst_gold, ksize, sigma1, sigma2, borderType);\r
+        cv::Mat dst_gold;\r
+        cv::GaussianBlur(src, dst_gold, ksize, sigma1, sigma2, borderType);\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 4.0);\r
+        EXPECT_MAT_NEAR(dst_gold, dst, 4.0);\r
+    }\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, GaussianBlur, testing::Combine(\r