1) more convenient naming for samples gpu

2) added mask support to device 'transform' function
3) sample hog gpu: waitKey(1) -> waitKey(3), in other case image is not displayed.

diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu
index 3f3fd71..d69a32a 100644 (file)
--- a/modules/gpu/src/cuda/mathfunc.cu
+++ b/modules/gpu/src/cuda/mathfunc.cu
@@ -217,7 +217,7 @@ namespace cv { namespace gpu { namespace mathfunc
     template <typename T1, typename T2>\r
     struct NotEqual\r
     {\r
-        __device__ uchar operator()(const T1& src1, const T2& src2, int, int)\r
+        __device__ uchar operator()(const T1& src1, const T2& src2)\r
         {\r
             return static_cast<uchar>(static_cast<int>(src1 != src2) * 255);\r
         }\r

diff --git a/modules/gpu/src/cuda/transform.hpp b/modules/gpu/src/cuda/transform.hpp
index 5445674..af516b3 100644 (file)
--- a/modules/gpu/src/cuda/transform.hpp
+++ b/modules/gpu/src/cuda/transform.hpp
@@ -47,31 +47,49 @@
 \r
 namespace cv { namespace gpu { namespace device\r
 {\r
-    template <typename T, typename D, typename UnOp>\r
-    static __global__ void transform(const DevMem2D_<T> src, PtrStep_<D> dst, UnOp op)\r
+    //! Mask accessor\r
+    template<class T> struct MaskReader_\r
+    {\r
+        PtrStep_<T> mask;\r
+        explicit MaskReader_(PtrStep_<T> mask): mask(mask) {}                \r
+\r
+        __device__ bool operator()(int y, int x) const { return mask.ptr(y)[x]; }\r
+    };\r
+\r
+    //! Stub mask accessor\r
+    struct NoMask \r
+    {\r
+        __device__ bool operator()(int y, int x) const { return true; } \r
+    };\r
+\r
+    //! Transform kernels\r
+\r
+    template <typename T, typename D, typename Mask, typename UnOp>\r
+    static __global__ void transform(const DevMem2D_<T> src, PtrStep_<D> dst, const Mask mask, UnOp op)\r
     {\r
                const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
                const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
 \r
-        if (x < src.cols && y < src.rows)\r
+        if (x < src.cols && y < src.rows && mask(y, x))\r
         {\r
             T src_data = src.ptr(y)[x];\r
-            dst.ptr(y)[x] = op(src_data, x, y);\r
+            dst.ptr(y)[x] = op(src_data);\r
         }\r
     }\r
-    template <typename T1, typename T2, typename D, typename BinOp>\r
-    static __global__ void transform(const DevMem2D_<T1> src1, const PtrStep_<T2> src2, PtrStep_<D> dst, BinOp op)\r
+\r
+    template <typename T1, typename T2, typename D, typename Mask, typename BinOp>\r
+    static __global__ void transform(const DevMem2D_<T1> src1, const PtrStep_<T2> src2, PtrStep_<D> dst, const Mask mask, BinOp op)\r
     {\r
                const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
                const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
 \r
-        if (x < src1.cols && y < src1.rows)\r
+        if (x < src1.cols && y < src1.rows && mask(y, x))\r
         {\r
             T1 src1_data = src1.ptr(y)[x];\r
             T2 src2_data = src2.ptr(y)[x];\r
-            dst.ptr(y)[x] = op(src1_data, src2_data, x, y);\r
+            dst.ptr(y)[x] = op(src1_data, src2_data);\r
         }\r
-    }\r
+    }  \r
 }}}\r
 \r
 namespace cv \r
@@ -87,7 +105,7 @@ namespace cv
             grid.x = divUp(src.cols, threads.x);\r
             grid.y = divUp(src.rows, threads.y);        \r
 \r
-            device::transform<T, D, UnOp><<<grid, threads, 0, stream>>>(src, dst, op);\r
+            device::transform<T, D, UnOp><<<grid, threads, 0, stream>>>(src, dst, device::NoMask(), op);\r
 \r
             if (stream == 0)\r
                 cudaSafeCall( cudaThreadSynchronize() );\r
@@ -101,7 +119,7 @@ namespace cv
             grid.x = divUp(src1.cols, threads.x);\r
             grid.y = divUp(src1.rows, threads.y);        \r
 \r
-            device::transform<T1, T2, D, BinOp><<<grid, threads, 0, stream>>>(src1, src2, dst, op);\r
+            device::transform<T1, T2, D><<<grid, threads, 0, stream>>>(src1, src2, dst, device::NoMask(), op);\r
 \r
             if (stream == 0)\r
                 cudaSafeCall( cudaThreadSynchronize() );            \r

diff --git a/samples/gpu/CMakeLists.txt b/samples/gpu/CMakeLists.txt
index c658e8c..041c0cf 100644 (file)
--- a/samples/gpu/CMakeLists.txt
+++ b/samples/gpu/CMakeLists.txt
@@ -24,11 +24,11 @@ if (BUILD_EXAMPLES)
     #      Define executable targets\r
     # ---------------------------------------------\r
     MACRO(MY_DEFINE_EXAMPLE name srcs)\r
-        set(the_target "example_${name}")\r
+        set(the_target "example_gpu_${name}")\r
        add_executable(${the_target} ${srcs})\r
                set_target_properties(${the_target} PROPERTIES\r
                    OUTPUT_NAME "${name}"\r
-                   PROJECT_LABEL "(EXAMPLE) ${name}")\r
+                   PROJECT_LABEL "(EXAMPLE_GPU) ${name}")\r
        add_dependencies(${the_target} opencv_core opencv_flann opencv_imgproc opencv_highgui\r
            opencv_ml opencv_video opencv_objdetect opencv_features2d\r
            opencv_calib3d opencv_legacy opencv_contrib opencv_gpu)\r

diff --git a/samples/gpu/hog.cpp b/samples/gpu/hog.cpp
index ed8823e..34e7c58 100644 (file)
--- a/samples/gpu/hog.cpp
+++ b/samples/gpu/hog.cpp
@@ -283,7 +283,7 @@ void App::RunOpencvGui()
             // Show results\r
             putText(img_to_show, GetPerformanceSummary(), Point(5, 25), FONT_HERSHEY_SIMPLEX, 1.0, Scalar(0, 0, 255), 2);\r
             imshow("opencv_gpu_hog", img_to_show);\r
-            HandleKey((char)waitKey(1));\r
+            HandleKey((char)waitKey(3));\r
 \r
             if (settings.src_is_video)\r
             {\r