added blendLinear into gpu module
authorAlexey Spizhevoy <no@email>
Fri, 8 Apr 2011 05:21:47 +0000 (05:21 +0000)
committerAlexey Spizhevoy <no@email>
Fri, 8 Apr 2011 05:21:47 +0000 (05:21 +0000)
modules/gpu/include/opencv2/gpu/gpu.hpp
modules/gpu/src/blend.cpp [new file with mode: 0644]
modules/gpu/src/cuda/blend.cu [new file with mode: 0644]
modules/gpu/test/test_blend.cpp [new file with mode: 0644]

index 6b324f8..a19809a 100644 (file)
@@ -786,6 +786,11 @@ 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
+        //! 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
 \r
         ////////////////////////////// Matrix reductions //////////////////////////////\r
 \r
diff --git a/modules/gpu/src/blend.cpp b/modules/gpu/src/blend.cpp
new file mode 100644 (file)
index 0000000..7c782e6
--- /dev/null
@@ -0,0 +1,101 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+//  By downloading, copying, installing or using the software you agree to this license.\r
+//  If you do not agree to this license, do not download, install,\r
+//  copy or use the software.\r
+//\r
+//\r
+//                           License Agreement\r
+//                For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+//   * Redistribution's of source code must retain the above copyright notice,\r
+//     this list of conditions and the following disclaimer.\r
+//\r
+//   * Redistribution's in binary form must reproduce the above copyright notice,\r
+//     this list of conditions and the following disclaimer in the documentation\r
+//     and/or other GpuMaterials provided with the distribution.\r
+//\r
+//   * The name of the copyright holders may not be used to endorse or promote products\r
+//     derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or bpied warranties, including, but not limited to, the bpied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include "precomp.hpp"\r
+\r
+using namespace std;\r
+using namespace cv;\r
+using namespace cv::gpu;\r
+\r
+#if !defined (HAVE_CUDA)\r
+\r
+void cv::gpu::blendLinear(const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, \r
+                          GpuMat&) { throw_nogpu(); }\r
+\r
+#else\r
+\r
+namespace cv { namespace gpu \r
+{\r
+    template <typename T>\r
+    void blendLinearCaller(int rows, int cols, int cn, const PtrStep_<T> img1, const PtrStep_<T> img2, \r
+                           const PtrStep_<float> weights1, const PtrStep_<float> weights2, PtrStep_<T> result);\r
+\r
+    void blendLinearCaller8UC4(int rows, int cols, const PtrStep img1, const PtrStep img2, \r
+                               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
+{\r
+    CV_Assert(img1.size() == img2.size());\r
+    CV_Assert(img1.type() == img2.type());\r
+    CV_Assert(weights1.size() == img1.size());\r
+    CV_Assert(weights2.size() == img2.size());\r
+    CV_Assert(weights1.type() == CV_32F);\r
+    CV_Assert(weights2.type() == CV_32F);\r
+\r
+    const Size size = img1.size();\r
+    const int depth = img1.depth();\r
+    const int cn = img1.channels();\r
+\r
+    result.create(size, CV_MAKE_TYPE(depth, cn));\r
+\r
+    switch (depth)\r
+    {\r
+    case CV_8U:\r
+        if (cn != 4)\r
+            blendLinearCaller(size.height, size.width, cn, (const PtrStep)img1, (const PtrStep)img2, \r
+                              (const PtrStepf)weights1, (const PtrStepf)weights2, (PtrStep)result);\r
+        else\r
+            blendLinearCaller8UC4(size.height, size.width, (const PtrStep)img1, (const PtrStep)img2, \r
+                                  (const PtrStepf)weights1, (const PtrStepf)weights2, (PtrStep)result);\r
+        break;\r
+    case CV_32F:\r
+        blendLinearCaller(size.height, size.width, cn, (const PtrStepf)img1, (const PtrStepf)img2, \r
+                          (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
+    }\r
+}\r
+\r
+#endif
\ No newline at end of file
diff --git a/modules/gpu/src/cuda/blend.cu b/modules/gpu/src/cuda/blend.cu
new file mode 100644 (file)
index 0000000..a9b85c1
--- /dev/null
@@ -0,0 +1,117 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+//  By downloading, copying, installing or using the software you agree to this license.\r
+//  If you do not agree to this license, do not download, install,\r
+//  copy or use the software.\r
+//\r
+//\r
+//                           License Agreement\r
+//                For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+//   * Redistribution's of source code must retain the above copyright notice,\r
+//     this list of conditions and the following disclaimer.\r
+//\r
+//   * Redistribution's in binary form must reproduce the above copyright notice,\r
+//     this list of conditions and the following disclaimer in the documentation\r
+//     and/or other materials provided with the distribution.\r
+//\r
+//   * The name of the copyright holders may not be used to endorse or promote products\r
+//     derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or bpied warranties, including, but not limited to, the bpied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include "internal_shared.hpp"\r
+\r
+using namespace cv::gpu;\r
+\r
+namespace cv { namespace gpu \r
+{\r
+\r
+    template <typename T>\r
+    __global__ void blendLinearKernel(int rows, int cols, int cn, const PtrStep_<T> img1, const PtrStep_<T> img2,\r
+                                      const PtrStepf weights1, const PtrStepf weights2, PtrStep_<T> result)\r
+    {\r
+        int x = blockIdx.x * blockDim.x + threadIdx.x;\r
+        int y = blockIdx.y * blockDim.y + threadIdx.y;\r
+\r
+        if (y < rows && x < cols)\r
+        {\r
+            int x_ = x / cn;\r
+            float w1 = weights1.ptr(y)[x_];\r
+            float w2 = weights2.ptr(y)[x_];\r
+            T p1 = img1.ptr(y)[x];\r
+            T p2 = img2.ptr(y)[x];\r
+            result.ptr(y)[x] = (p1 * w1 + p2 * w2) / (w1 + w2 + 1e-5f);\r
+        }\r
+    }\r
+\r
+\r
+    template <typename T>\r
+    void blendLinearCaller(int rows, int cols, int cn, const PtrStep_<T> img1, const PtrStep_<T> img2, \r
+                           const PtrStepf weights1, const PtrStepf weights2, PtrStep_<T> result)\r
+    {\r
+        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
+        cudaSafeCall(cudaThreadSynchronize());\r
+    }\r
+\r
+    template void blendLinearCaller<uchar>(int, int, int, const PtrStep, const PtrStep, \r
+                                           const PtrStepf, const PtrStepf, PtrStep);\r
+    template void blendLinearCaller<float>(int, int, int, const PtrStepf, const PtrStepf, \r
+                                           const PtrStepf, const PtrStepf, PtrStepf);\r
+\r
+\r
+    __global__ void blendLinearKernel8UC4(int rows, int cols, const PtrStep img1, const PtrStep img2,\r
+                                          const PtrStepf weights1, const PtrStepf weights2, PtrStep result)\r
+    {\r
+        int x = blockIdx.x * blockDim.x + threadIdx.x;\r
+        int y = blockIdx.y * blockDim.y + threadIdx.y;\r
+\r
+        if (y < rows && x < cols)\r
+        {\r
+            float w1 = weights1.ptr(y)[x];\r
+            float w2 = weights2.ptr(y)[x];\r
+            float sum_inv = 1.f / (w1 + w2 + 1e-5f);\r
+            w1 *= sum_inv;\r
+            w2 *= sum_inv;\r
+            uchar4 p1 = ((const uchar4*)img1.ptr(y))[x];\r
+            uchar4 p2 = ((const uchar4*)img2.ptr(y))[x];\r
+            ((uchar4*)result.ptr(y))[x] = make_uchar4(p1.x * w1 + p2.x * w2, p1.y * w1 + p2.y * w2,\r
+                                                      p1.z * w1 + p2.z * w2, p1.w * w1 + p2.w * w2);\r
+        }\r
+    }\r
+\r
+\r
+    void blendLinearCaller8UC4(int rows, int cols, const PtrStep img1, const PtrStep img2, \r
+                               const PtrStepf weights1, const PtrStepf weights2, PtrStep result)\r
+    {\r
+        dim3 threads(16, 16);\r
+        dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));\r
+        \r
+        blendLinearKernel8UC4<<<grid, threads>>>(rows, cols, img1, img2, weights1, weights2, result);\r
+        cudaSafeCall(cudaThreadSynchronize());\r
+    }\r
+\r
+}}
\ No newline at end of file
diff --git a/modules/gpu/test/test_blend.cpp b/modules/gpu/test/test_blend.cpp
new file mode 100644 (file)
index 0000000..60955ad
--- /dev/null
@@ -0,0 +1,96 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+//  By downloading, copying, installing or using the software you agree to this license.\r
+//  If you do not agree to this license, do not download, install,\r
+//  copy or use the software.\r
+//\r
+//\r
+//                        Intel License Agreement\r
+//                For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000, Intel Corporation, all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+//   * Redistribution's of source code must retain the above copyright notice,\r
+//     this list of conditions and the following disclaimer.\r
+//\r
+//   * Redistribution's in binary form must reproduce the above copyright notice,\r
+//     this list of conditions and the following disclaimer in the documentation\r
+//     and/or other materials provided with the distribution.\r
+//\r
+//   * The name of Intel Corporation may not be used to endorse or promote products\r
+//     derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include "test_precomp.hpp"\r
+\r
+using namespace std;\r
+using namespace cv;\r
+using namespace cv::gpu;\r
+\r
+TEST(blendLinear, accuracy_on_8U)\r
+{\r
+    Size size(607, 1021);\r
+    RNG rng(0);\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
+        Mat img2 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_8U, cn), 0, 255, false);\r
+        Mat weights1 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);\r
+        Mat weights2 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);\r
+        Mat result_gold(size, CV_MAKE_TYPE(CV_8U, cn));\r
+        for (int y = 0; y < size.height; ++y)\r
+            for (int x = 0; x < size.width * cn; ++x)\r
+            {\r
+                float w1 = weights1.at<float>(y, x / cn);\r
+                float w2 = weights2.at<float>(y, x / cn);\r
+                result_gold.at<uchar>(y, x) = static_cast<uchar>(\r
+                    (img1.at<uchar>(y, x) * w1 + img2.at<uchar>(y, x) * w2) / (w1 + w2 + 1e-5f));\r
+            }\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
+    }\r
+}\r
+\r
+TEST(blendLinear, accuracy_on_32F)\r
+{\r
+    Size size(607, 1021);\r
+    RNG rng(0);\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
+        Mat img2 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_32F, cn), 0, 1, false);\r
+        Mat weights1 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);\r
+        Mat weights2 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);\r
+        Mat result_gold(size, CV_MAKE_TYPE(CV_32F, cn));\r
+        for (int y = 0; y < size.height; ++y)\r
+            for (int x = 0; x < size.width * cn; ++x)\r
+            {\r
+                float w1 = weights1.at<float>(y, x / cn);\r
+                float w2 = weights2.at<float>(y, x / cn);\r
+                result_gold.at<float>(y, x) = \r
+                    (img1.at<float>(y, x) * w1 + img2.at<float>(y, x) * w2) / (w1 + w2 + 1e-5f);\r
+            }\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
+    }\r
+}
\ No newline at end of file