added downsample function into gpu module, refactored it a little bit, added guard...

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index a19809a..874666c 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -786,11 +786,13 @@ namespace cv
         //! computes the proximity map for the raster template and the image where the template is searched for\r
         CV_EXPORTS void matchTemplate(const GpuMat& image, const GpuMat& templ, GpuMat& result, int method);\r
 \r
+        //! downsamples image\r
+        CV_EXPORTS void downsample(const GpuMat& src, GpuMat& dst, int k=2);\r
+\r
         //! performs linear blending of two images\r
         //! to avoid accuracy errors sum of weigths shouldn't be very close to zero\r
-        CV_EXPORTS void blendLinear(const GpuMat& img1, const GpuMat& img2, \r
-                                    const GpuMat& weights1, const GpuMat& weights2, GpuMat& result);\r
-\r
+        CV_EXPORTS void blendLinear(const GpuMat& img1, const GpuMat& img2, const GpuMat& weights1, const GpuMat& weights2, \r
+                                    GpuMat& result);\r
 \r
         ////////////////////////////// Matrix reductions //////////////////////////////\r
 \r

diff --git a/modules/gpu/src/blend.cpp b/modules/gpu/src/blend.cpp
index 7c782e6..c5dfd58 100644 (file)
--- a/modules/gpu/src/blend.cpp
+++ b/modules/gpu/src/blend.cpp
@@ -63,8 +63,8 @@ namespace cv { namespace gpu
                                const PtrStepf weights1, const PtrStepf weights2, PtrStep result);\r
 }}\r
 \r
-void cv::gpu::blendLinear(const GpuMat& img1, const GpuMat& img2, \r
-                          const GpuMat& weights1, const GpuMat& weights2, GpuMat& result)\r
+void cv::gpu::blendLinear(const GpuMat& img1, const GpuMat& img2, const GpuMat& weights1, const GpuMat& weights2, \r
+                          GpuMat& result)\r
 {\r
     CV_Assert(img1.size() == img2.size());\r
     CV_Assert(img1.type() == img2.type());\r
@@ -94,7 +94,7 @@ void cv::gpu::blendLinear(const GpuMat& img1, const GpuMat& img2,
                           (const PtrStepf)weights1, (const PtrStepf)weights2, (PtrStepf)result);\r
         break;\r
     default:\r
-        CV_Error(CV_StsBadArg, "unsupported image depth in linear blending method");\r
+        CV_Error(CV_StsUnsupportedFormat, "bad image depth in linear blending function");\r
     }\r
 }\r
 \r

diff --git a/modules/gpu/src/cuda/blend.cu b/modules/gpu/src/cuda/blend.cu
index a9b85c1..42b5a63 100644 (file)
--- a/modules/gpu/src/cuda/blend.cu
+++ b/modules/gpu/src/cuda/blend.cu
@@ -73,7 +73,7 @@ namespace cv { namespace gpu
         dim3 threads(16, 16);\r
         dim3 grid(divUp(cols * cn, threads.x), divUp(rows, threads.y));\r
         \r
-        blendLinearKernel<T><<<grid, threads>>>(rows, cols * cn, cn, img1, img2, weights1, weights2, result);\r
+        blendLinearKernel<<<grid, threads>>>(rows, cols * cn, cn, img1, img2, weights1, weights2, result);\r
         cudaSafeCall(cudaThreadSynchronize());\r
     }\r
 \r

diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu
index 2c94f83..dad5335 100644 (file)
--- a/modules/gpu/src/cuda/imgproc.cu
+++ b/modules/gpu/src/cuda/imgproc.cu
@@ -883,5 +883,32 @@ namespace cv { namespace gpu { namespace imgproc
         cudaSafeCall(cudaThreadSynchronize());\r
     }\r
 \r
+    /////////////////////////////////////////////////////////////////////////\r
+    // downsample\r
+\r
+    template <typename T>\r
+    __global__ void downsampleKernel(const PtrStep_<T> src, int rows, int cols, int k, PtrStep_<T> dst)\r
+    {\r
+        int x = blockIdx.x * blockDim.x + threadIdx.x;\r
+        int y = blockIdx.y * blockDim.y + threadIdx.y;\r
+\r
+        if (x < cols && y < rows)\r
+            dst.ptr(y)[x] = src.ptr(y * k)[x * k];\r
+    }\r
+\r
+\r
+    template <typename T>\r
+    void downsampleCaller(const PtrStep_<T> src, int rows, int cols, int k, PtrStep_<T> dst)\r
+    {\r
+        dim3 threads(16, 16);\r
+        dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));\r
+\r
+        downsampleKernel<<<grid, threads>>>(src, rows, cols, k, dst);\r
+        cudaSafeCall(cudaThreadSynchronize());\r
+    }\r
+\r
+    template void downsampleCaller(const PtrStep src, int rows, int cols, int k, PtrStep dst);\r
+    template void downsampleCaller(const PtrStepf src, int rows, int cols, int k, PtrStepf dst);\r
+\r
 }}}\r
 \r

diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp
index 486ca45..0db8f5a 100644 (file)
--- a/modules/gpu/src/imgproc_gpu.cpp
+++ b/modules/gpu/src/imgproc_gpu.cpp
@@ -82,6 +82,7 @@ void cv::gpu::dft(const GpuMat&, GpuMat&, Size, int) { throw_nogpu(); }
 void cv::gpu::ConvolveBuf::create(Size, Size) { throw_nogpu(); }\r
 void cv::gpu::convolve(const GpuMat&, const GpuMat&, GpuMat&, bool) { throw_nogpu(); }\r
 void cv::gpu::convolve(const GpuMat&, const GpuMat&, GpuMat&, bool, ConvolveBuf&) { throw_nogpu(); }\r
+void cv::gpu::downsample(const GpuMat&, GpuMat&, int) { throw_nogpu(); }\r
 \r
 \r
 #else /* !defined (HAVE_CUDA) */\r
@@ -1355,7 +1356,33 @@ void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result,
     cufftSafeCall(cufftDestroy(planC2R));\r
 }\r
 \r
+////////////////////////////////////////////////////////////////////\r
+// downsample\r
 \r
+namespace cv { namespace gpu { namespace imgproc\r
+{\r
+    template <typename T>\r
+    void downsampleCaller(const PtrStep_<T> src, int rows, int cols, int k, PtrStep_<T> dst);\r
+}}}\r
+\r
+void cv::gpu::downsample(const GpuMat& src, GpuMat& dst, int k)\r
+{\r
+    CV_Assert(src.channels() == 1);    \r
+\r
+    dst.create((src.rows + k - 1) / k, (src.cols + k - 1) / k, src.type());\r
+\r
+    switch (src.depth())\r
+    {\r
+    case CV_8U:\r
+        imgproc::downsampleCaller((const PtrStep)src, dst.rows, dst.cols, k, (PtrStep)dst);\r
+        break;\r
+    case CV_32F:\r
+        imgproc::downsampleCaller((const PtrStepf)src, dst.rows, dst.cols, k, (PtrStepf)dst);\r
+        break;\r
+    default:\r
+        CV_Error(CV_StsUnsupportedFormat, "bad image depth in downsample function");\r
+    }\r
+}\r
 \r
 #endif /* !defined (HAVE_CUDA) */\r
 \r

diff --git a/modules/gpu/test/test_blend.cpp b/modules/gpu/test/test_blend.cpp
index 60955ad..9479695 100644 (file)
--- a/modules/gpu/test/test_blend.cpp
+++ b/modules/gpu/test/test_blend.cpp
@@ -47,8 +47,9 @@ using namespace cv::gpu;
 \r
 TEST(blendLinear, accuracy_on_8U)\r
 {\r
-    Size size(607, 1021);\r
-    RNG rng(0);\r
+    RNG& rng = cvtest::TS::ptr()->get_rng();\r
+    Size size(200 + cvtest::randInt(rng) % 1000,\r
+              200 + cvtest::randInt(rng) % 1000);\r
     for (int cn = 1; cn <= 4; ++cn)\r
     {\r
         Mat img1 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_8U, cn), 0, 255, false);\r
@@ -66,14 +67,16 @@ TEST(blendLinear, accuracy_on_8U)
             }\r
         GpuMat d_result;\r
         blendLinear(GpuMat(img1), GpuMat(img2), GpuMat(weights1), GpuMat(weights2), d_result);\r
-        ASSERT_LE(cvtest::norm(result_gold, Mat(d_result), NORM_INF), 1) << ", cn=" << cn;\r
+        ASSERT_LE(cvtest::norm(result_gold, Mat(d_result), NORM_INF), 1) \r
+            << "rows=" << size.height << ", cols=" << size.width << ", cn=" << cn;\r
     }\r
 }\r
 \r
 TEST(blendLinear, accuracy_on_32F)\r
 {\r
-    Size size(607, 1021);\r
-    RNG rng(0);\r
+    RNG& rng = cvtest::TS::ptr()->get_rng();\r
+    Size size(200 + cvtest::randInt(rng) % 1000,\r
+              200 + cvtest::randInt(rng) % 1000);\r
     for (int cn = 1; cn <= 4; ++cn)\r
     {\r
         Mat img1 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_32F, cn), 0, 1, false);\r
@@ -91,6 +94,7 @@ TEST(blendLinear, accuracy_on_32F)
             }\r
         GpuMat d_result;\r
         blendLinear(GpuMat(img1), GpuMat(img2), GpuMat(weights1), GpuMat(weights2), d_result);\r
-        ASSERT_LE(cvtest::norm(result_gold, Mat(d_result), NORM_INF), 1e-3) << ", cn=" << cn;\r
+        ASSERT_LE(cvtest::norm(result_gold, Mat(d_result), NORM_INF), 1e-3)\r
+            << "rows=" << size.height << ", cols=" << size.width << ", cn=" << cn;\r
     }\r
-}
\ No newline at end of file
+}\r

diff --git a/modules/gpu/test/test_imgproc_gpu.cpp b/modules/gpu/test/test_imgproc_gpu.cpp
index b8363a0..f61698e 100644 (file)
--- a/modules/gpu/test/test_imgproc_gpu.cpp
+++ b/modules/gpu/test/test_imgproc_gpu.cpp
@@ -914,3 +914,53 @@ TEST(minEigen, accuracy) { CV_GpuCornerMinEigenValTest test; test.safe_run(); }
 TEST(columnSum, accuracy) { CV_GpuColumnSumTest test; test.safe_run(); }\r
 TEST(norm, accuracy) { CV_GpuNormTest test; test.safe_run(); }\r
 TEST(reprojectImageTo3D, accuracy) { CV_GpuReprojectImageTo3DTest test; test.safe_run(); }\r
+\r
+TEST(downsample, accuracy_on_8U)\r
+{\r
+    RNG& rng = cvtest::TS::ptr()->get_rng();\r
+    Size size(200 + cvtest::randInt(rng) % 1000, 200 + cvtest::randInt(rng) % 1000);\r
+    Mat src = cvtest::randomMat(rng, size, CV_8U, 0, 255, false);\r
+\r
+    for (int k = 2; k <= 5; ++k)\r
+    {\r
+        GpuMat d_dst;\r
+        downsample(GpuMat(src), d_dst, k);       \r
+\r
+        Size dst_gold_size((src.cols + k - 1) / k, (src.rows + k - 1) / k);\r
+        ASSERT_EQ(dst_gold_size.width, d_dst.cols) \r
+            << "rows=" << size.height << ", cols=" << size.width << ", k=" << k;\r
+        ASSERT_EQ(dst_gold_size.height, d_dst.rows) \r
+            << "rows=" << size.height << ", cols=" << size.width << ", k=" << k;\r
+\r
+        Mat dst = d_dst;\r
+        for (int y = 0; y < dst.rows; ++y)\r
+            for (int x = 0; x < dst.cols; ++x)\r
+                ASSERT_EQ(src.at<uchar>(y * k, x * k), dst.at<uchar>(y, x))\r
+                    << "rows=" << size.height << ", cols=" << size.width << ", k=" << k;\r
+    }\r
+}\r
+\r
+TEST(downsample, accuracy_on_32F)\r
+{\r
+    RNG& rng = cvtest::TS::ptr()->get_rng();\r
+    Size size(200 + cvtest::randInt(rng) % 1000, 200 + cvtest::randInt(rng) % 1000);\r
+    Mat src = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);\r
+\r
+    for (int k = 2; k <= 5; ++k)\r
+    {\r
+        GpuMat d_dst;\r
+        downsample(GpuMat(src), d_dst, k);       \r
+\r
+        Size dst_gold_size((src.cols + k - 1) / k, (src.rows + k - 1) / k);\r
+        ASSERT_EQ(dst_gold_size.width, d_dst.cols) \r
+            << "rows=" << size.height << ", cols=" << size.width << ", k=" << k;\r
+        ASSERT_EQ(dst_gold_size.height, d_dst.rows) \r
+            << "rows=" << size.height << ", cols=" << size.width << ", k=" << k;\r
+\r
+        Mat dst = d_dst;\r
+        for (int y = 0; y < dst.rows; ++y)\r
+            for (int x = 0; x < dst.cols; ++x)\r
+                ASSERT_FLOAT_EQ(src.at<float>(y * k, x * k), dst.at<float>(y, x))\r
+                    << "rows=" << size.height << ", cols=" << size.width << ", k=" << k;\r
+    }\r
+}\r

diff --git a/samples/gpu/cascadeclassifier_nvidia_api.cpp b/samples/gpu/cascadeclassifier_nvidia_api.cpp
index 2dd3945..2faabb3 100644 (file)
--- a/samples/gpu/cascadeclassifier_nvidia_api.cpp
+++ b/samples/gpu/cascadeclassifier_nvidia_api.cpp
@@ -5,7 +5,10 @@
 #include <iomanip>\r
 #include <opencv2/opencv.hpp>\r
 #include <opencv2/gpu/gpu.hpp>\r
+\r
+#ifdef HAVE_CUDA\r
 #include "NCVHaarObjectDetection.hpp"\r
+#endif\r
 \r
 using namespace std;\r
 using namespace cv;\r