fixed gpu::sum* on CC1.0, updated some tests

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index df30a8c..0862ee6 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -66,8 +66,14 @@ namespace cv
 \r
         enum GpuFeature\r
         {\r
-            NATIVE_DOUBLE,\r
-            ATOMICS\r
+            COMPUTE_10 = 10,\r
+            COMPUTE_11 = 11,\r
+            COMPUTE_12 = 12,\r
+            COMPUTE_13 = 13,\r
+            COMPUTE_20 = 20,\r
+            COMPUTE_21 = 21,\r
+            ATOMICS = COMPUTE_11,\r
+            NATIVE_DOUBLE = COMPUTE_13\r
         };\r
 \r
         class CV_EXPORTS TargetArchs\r

diff --git a/modules/gpu/src/cuda/matrix_reductions.cu b/modules/gpu/src/cuda/matrix_reductions.cu
index 398e376..396a9d7 100644 (file)
--- a/modules/gpu/src/cuda/matrix_reductions.cu
+++ b/modules/gpu/src/cuda/matrix_reductions.cu
@@ -1394,7 +1394,7 @@ namespace cv { namespace gpu { namespace mathfunc
         smem[tid] = res.x;\r
         smem[tid + nthreads] = res.y;\r
         smem[tid + 2 * nthreads] = res.z;\r
-        smem[tid + 3 * nthreads] = res.z;\r
+        smem[tid + 3 * nthreads] = res.w;\r
         __syncthreads();\r
 \r
         sumInSmem<nthreads, R>(smem, tid);\r
@@ -1432,21 +1432,25 @@ namespace cv { namespace gpu { namespace mathfunc
                     src, (typename TypeVec<R, 1>::vec_t*)buf.ptr(0));\r
             sumPass2Kernel<T, R, threads_x * threads_y><<<1, threads_x * threads_y>>>(\r
                     (typename TypeVec<R, 1>::vec_t*)buf.ptr(0), grid.x * grid.y);\r
+            break;\r
         case 2:\r
             sumKernel_C2<T, R, IdentityOp<R>, threads_x * threads_y><<<grid, threads>>>(\r
                     src, (typename TypeVec<R, 2>::vec_t*)buf.ptr(0));\r
             sumPass2Kernel_C2<T, R, threads_x * threads_y><<<1, threads_x * threads_y>>>(\r
                     (typename TypeVec<R, 2>::vec_t*)buf.ptr(0), grid.x * grid.y);\r
+            break;\r
         case 3:\r
             sumKernel_C3<T, R, IdentityOp<R>, threads_x * threads_y><<<grid, threads>>>(\r
                     src, (typename TypeVec<R, 3>::vec_t*)buf.ptr(0));\r
             sumPass2Kernel_C3<T, R, threads_x * threads_y><<<1, threads_x * threads_y>>>(\r
                     (typename TypeVec<R, 3>::vec_t*)buf.ptr(0), grid.x * grid.y);\r
+            break;\r
         case 4:\r
             sumKernel_C4<T, R, IdentityOp<R>, threads_x * threads_y><<<grid, threads>>>(\r
                     src, (typename TypeVec<R, 4>::vec_t*)buf.ptr(0));\r
             sumPass2Kernel_C4<T, R, threads_x * threads_y><<<1, threads_x * threads_y>>>(\r
                     (typename TypeVec<R, 4>::vec_t*)buf.ptr(0), grid.x * grid.y);\r
+            break;\r
         }\r
         cudaSafeCall(cudaThreadSynchronize());\r
 \r

diff --git a/modules/gpu/src/initialization.cpp b/modules/gpu/src/initialization.cpp
index d11b95c..039302c 100644 (file)
--- a/modules/gpu/src/initialization.cpp
+++ b/modules/gpu/src/initialization.cpp
@@ -71,11 +71,7 @@ namespace
 \r
 CV_EXPORTS bool cv::gpu::TargetArchs::builtWith(cv::gpu::GpuFeature feature)\r
 {\r
-    if (feature == NATIVE_DOUBLE)\r
-        return ::compareToSet(CUDA_ARCH_FEATURES, 13, std::greater_equal<int>());\r
-    if (feature == ATOMICS)\r
-        return ::compareToSet(CUDA_ARCH_FEATURES, 11, std::greater_equal<int>());\r
-    return true;\r
+    return ::compareToSet(CUDA_ARCH_FEATURES, feature, std::greater_equal<int>());\r
 }\r
 \r
 \r

diff --git a/tests/gpu/src/arithm.cpp b/tests/gpu/src/arithm.cpp
index 28cbc42..b7501a4 100644 (file)
--- a/tests/gpu/src/arithm.cpp
+++ b/tests/gpu/src/arithm.cpp
@@ -947,6 +947,16 @@ struct CV_GpuSumTest: CvTest
                 // sum\r
                 //\r
 \r
+                gen(1 + rand() % 500, 1 + rand() % 500, CV_MAKETYPE(type, 1), src);\r
+                a = sum(src);\r
+                b = sum(GpuMat(src));\r
+                if (abs(a[0] - b[0]) > src.size().area() * max_err)\r
+                {\r
+                    ts->printf(CvTS::CONSOLE, "1 cols: %d, rows: %d, expected: %f, actual: %f\n", src.cols, src.rows, a[0], b[0]);\r
+                    ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);\r
+                    return;\r
+                }\r
+\r
                 gen(1 + rand() % 500, 1 + rand() % 500, CV_MAKETYPE(type, 2), src);\r
                 a = sum(src);\r
                 b = sum(GpuMat(src));\r

diff --git a/tests/gpu/src/meanshift.cpp b/tests/gpu/src/meanshift.cpp
index dc39423..d1bd345 100644 (file)
--- a/tests/gpu/src/meanshift.cpp
+++ b/tests/gpu/src/meanshift.cpp
@@ -56,7 +56,7 @@ struct CV_GpuMeanShiftTest : public CvTest
         cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png");\r
         cv::Mat img_template;       \r
         \r
-        if (cv::gpu::TargetArchs::hasEqualOrGreater(2, 0) && cv::gpu::DeviceInfo().major() >= 2)\r
+        if (cv::gpu::TargetArchs::builtWith(cv::gpu::COMPUTE_20) && cv::gpu::DeviceInfo().major() >= 2)\r
             img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result.png");\r
         else\r
             img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result_CC1X.png");\r
@@ -199,7 +199,7 @@ struct CV_GpuMeanShiftProcTest : public CvTest
             cv::Mat spmap_template;\r
             cv::FileStorage fs;\r
 \r
-            if (cv::gpu::TargetArchs::hasEqualOrGreater(2, 0) && cv::gpu::DeviceInfo().major() >= 2)\r
+            if (cv::gpu::TargetArchs::builtWith(cv::gpu::COMPUTE_20) && cv::gpu::DeviceInfo().major() >= 2)\r
                 fs.open(std::string(ts->get_data_path()) + "meanshift/spmap.yaml", cv::FileStorage::READ);\r
             else\r
                 fs.open(std::string(ts->get_data_path()) + "meanshift/spmap_CC1X.yaml", cv::FileStorage::READ);\r

diff --git a/tests/gpu/src/mssegmentation.cpp b/tests/gpu/src/mssegmentation.cpp
index ea09b01..15a72f3 100644 (file)
--- a/tests/gpu/src/mssegmentation.cpp
+++ b/tests/gpu/src/mssegmentation.cpp
@@ -69,7 +69,7 @@ struct CV_GpuMeanShiftSegmentationTest : public CvTest {
             {\r
                 stringstream path;\r
                 path << ts->get_data_path() << "meanshift/cones_segmented_sp10_sr10_minsize" << minsize;\r
-                if (TargetArchs::hasEqualOrGreater(2, 0) && DeviceInfo().major() >= 2)\r
+                if (TargetArchs::builtWith(COMPUTE_20) && DeviceInfo().major() >= 2)\r
                     path << ".png";\r
                 else\r
                     path << "_CC1X.png";\r