added blendLinear into gpu module

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -786,6 +786,11 @@ namespace cv
         //! computes the proximity map for the raster template and the image where the template is searched for
         CV_EXPORTS void matchTemplate(const GpuMat& image, const GpuMat& templ, GpuMat& result, int method);
 
+        //! performs linear blending of two images
+        //! to avoid accuracy errors sum of weigths shouldn't be very close to zero
+        CV_EXPORTS void blendLinear(const GpuMat& img1, const GpuMat& img2, 
+                                    const GpuMat& weights1, const GpuMat& weights2, GpuMat& result);
+
 
         ////////////////////////////// Matrix reductions //////////////////////////////
 

modules/gpu/src/blend.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other GpuMaterials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or bpied warranties, including, but not limited to, the bpied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "precomp.hpp"
+
+using namespace std;
+using namespace cv;
+using namespace cv::gpu;
+
+#if !defined (HAVE_CUDA)
+
+void cv::gpu::blendLinear(const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, 
+                          GpuMat&) { throw_nogpu(); }
+
+#else
+
+namespace cv { namespace gpu 
+{
+    template <typename T>
+    void blendLinearCaller(int rows, int cols, int cn, const PtrStep_<T> img1, const PtrStep_<T> img2, 
+                           const PtrStep_<float> weights1, const PtrStep_<float> weights2, PtrStep_<T> result);
+
+    void blendLinearCaller8UC4(int rows, int cols, const PtrStep img1, const PtrStep img2, 
+                               const PtrStepf weights1, const PtrStepf weights2, PtrStep result);
+}}
+
+void cv::gpu::blendLinear(const GpuMat& img1, const GpuMat& img2, 
+                          const GpuMat& weights1, const GpuMat& weights2, GpuMat& result)
+{
+    CV_Assert(img1.size() == img2.size());
+    CV_Assert(img1.type() == img2.type());
+    CV_Assert(weights1.size() == img1.size());
+    CV_Assert(weights2.size() == img2.size());
+    CV_Assert(weights1.type() == CV_32F);
+    CV_Assert(weights2.type() == CV_32F);
+
+    const Size size = img1.size();
+    const int depth = img1.depth();
+    const int cn = img1.channels();
+
+    result.create(size, CV_MAKE_TYPE(depth, cn));
+
+    switch (depth)
+    {
+    case CV_8U:
+        if (cn != 4)
+            blendLinearCaller(size.height, size.width, cn, (const PtrStep)img1, (const PtrStep)img2, 
+                              (const PtrStepf)weights1, (const PtrStepf)weights2, (PtrStep)result);
+        else
+            blendLinearCaller8UC4(size.height, size.width, (const PtrStep)img1, (const PtrStep)img2, 
+                                  (const PtrStepf)weights1, (const PtrStepf)weights2, (PtrStep)result);
+        break;
+    case CV_32F:
+        blendLinearCaller(size.height, size.width, cn, (const PtrStepf)img1, (const PtrStepf)img2, 
+                          (const PtrStepf)weights1, (const PtrStepf)weights2, (PtrStepf)result);
+        break;
+    default:
+        CV_Error(CV_StsBadArg, "unsupported image depth in linear blending method");
+    }
+}
+
+#endif
\ No newline at end of file

modules/gpu/src/cuda/blend.cu

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or bpied warranties, including, but not limited to, the bpied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "internal_shared.hpp"
+
+using namespace cv::gpu;
+
+namespace cv { namespace gpu 
+{
+
+    template <typename T>
+    __global__ void blendLinearKernel(int rows, int cols, int cn, const PtrStep_<T> img1, const PtrStep_<T> img2,
+                                      const PtrStepf weights1, const PtrStepf weights2, PtrStep_<T> result)
+    {
+        int x = blockIdx.x * blockDim.x + threadIdx.x;
+        int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+        if (y < rows && x < cols)
+        {
+            int x_ = x / cn;
+            float w1 = weights1.ptr(y)[x_];
+            float w2 = weights2.ptr(y)[x_];
+            T p1 = img1.ptr(y)[x];
+            T p2 = img2.ptr(y)[x];
+            result.ptr(y)[x] = (p1 * w1 + p2 * w2) / (w1 + w2 + 1e-5f);
+        }
+    }
+
+
+    template <typename T>
+    void blendLinearCaller(int rows, int cols, int cn, const PtrStep_<T> img1, const PtrStep_<T> img2, 
+                           const PtrStepf weights1, const PtrStepf weights2, PtrStep_<T> result)
+    {
+        dim3 threads(16, 16);
+        dim3 grid(divUp(cols * cn, threads.x), divUp(rows, threads.y));
+        
+        blendLinearKernel<T><<<grid, threads>>>(rows, cols * cn, cn, img1, img2, weights1, weights2, result);
+        cudaSafeCall(cudaThreadSynchronize());
+    }
+
+    template void blendLinearCaller<uchar>(int, int, int, const PtrStep, const PtrStep, 
+                                           const PtrStepf, const PtrStepf, PtrStep);
+    template void blendLinearCaller<float>(int, int, int, const PtrStepf, const PtrStepf, 
+                                           const PtrStepf, const PtrStepf, PtrStepf);
+
+
+    __global__ void blendLinearKernel8UC4(int rows, int cols, const PtrStep img1, const PtrStep img2,
+                                          const PtrStepf weights1, const PtrStepf weights2, PtrStep result)
+    {
+        int x = blockIdx.x * blockDim.x + threadIdx.x;
+        int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+        if (y < rows && x < cols)
+        {
+            float w1 = weights1.ptr(y)[x];
+            float w2 = weights2.ptr(y)[x];
+            float sum_inv = 1.f / (w1 + w2 + 1e-5f);
+            w1 *= sum_inv;
+            w2 *= sum_inv;
+            uchar4 p1 = ((const uchar4*)img1.ptr(y))[x];
+            uchar4 p2 = ((const uchar4*)img2.ptr(y))[x];
+            ((uchar4*)result.ptr(y))[x] = make_uchar4(p1.x * w1 + p2.x * w2, p1.y * w1 + p2.y * w2,
+                                                      p1.z * w1 + p2.z * w2, p1.w * w1 + p2.w * w2);
+        }
+    }
+
+
+    void blendLinearCaller8UC4(int rows, int cols, const PtrStep img1, const PtrStep img2, 
+                               const PtrStepf weights1, const PtrStepf weights2, PtrStep result)
+    {
+        dim3 threads(16, 16);
+        dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));
+        
+        blendLinearKernel8UC4<<<grid, threads>>>(rows, cols, img1, img2, weights1, weights2, result);
+        cudaSafeCall(cudaThreadSynchronize());
+    }
+
+}}
\ No newline at end of file

modules/gpu/test/test_blend.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                        Intel License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000, Intel Corporation, all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of Intel Corporation may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "test_precomp.hpp"
+
+using namespace std;
+using namespace cv;
+using namespace cv::gpu;
+
+TEST(blendLinear, accuracy_on_8U)
+{
+    Size size(607, 1021);
+    RNG rng(0);
+    for (int cn = 1; cn <= 4; ++cn)
+    {
+        Mat img1 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_8U, cn), 0, 255, false);
+        Mat img2 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_8U, cn), 0, 255, false);
+        Mat weights1 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);
+        Mat weights2 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);
+        Mat result_gold(size, CV_MAKE_TYPE(CV_8U, cn));
+        for (int y = 0; y < size.height; ++y)
+            for (int x = 0; x < size.width * cn; ++x)
+            {
+                float w1 = weights1.at<float>(y, x / cn);
+                float w2 = weights2.at<float>(y, x / cn);
+                result_gold.at<uchar>(y, x) = static_cast<uchar>(
+                    (img1.at<uchar>(y, x) * w1 + img2.at<uchar>(y, x) * w2) / (w1 + w2 + 1e-5f));
+            }
+        GpuMat d_result;
+        blendLinear(GpuMat(img1), GpuMat(img2), GpuMat(weights1), GpuMat(weights2), d_result);
+        ASSERT_LE(cvtest::norm(result_gold, Mat(d_result), NORM_INF), 1) << ", cn=" << cn;
+    }
+}
+
+TEST(blendLinear, accuracy_on_32F)
+{
+    Size size(607, 1021);
+    RNG rng(0);
+    for (int cn = 1; cn <= 4; ++cn)
+    {
+        Mat img1 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_32F, cn), 0, 1, false);
+        Mat img2 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_32F, cn), 0, 1, false);
+        Mat weights1 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);
+        Mat weights2 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);
+        Mat result_gold(size, CV_MAKE_TYPE(CV_32F, cn));
+        for (int y = 0; y < size.height; ++y)
+            for (int x = 0; x < size.width * cn; ++x)
+            {
+                float w1 = weights1.at<float>(y, x / cn);
+                float w2 = weights2.at<float>(y, x / cn);
+                result_gold.at<float>(y, x) = 
+                    (img1.at<float>(y, x) * w1 + img2.at<float>(y, x) * w2) / (w1 + w2 + 1e-5f);
+            }
+        GpuMat d_result;
+        blendLinear(GpuMat(img1), GpuMat(img2), GpuMat(weights1), GpuMat(weights2), d_result);
+        ASSERT_LE(cvtest::norm(result_gold, Mat(d_result), NORM_INF), 1e-3) << ", cn=" << cn;
+    }
+}
\ No newline at end of file