diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 6b324f852b..a19809a050 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/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 ////////////////////////////// diff --git a/modules/gpu/src/blend.cpp b/modules/gpu/src/blend.cpp new file mode 100644 index 0000000000..7c782e6cb4 --- /dev/null +++ b/modules/gpu/src/blend.cpp @@ -0,0 +1,101 @@ +/*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 + void blendLinearCaller(int rows, int cols, int cn, const PtrStep_ img1, const PtrStep_ img2, + const PtrStep_ weights1, const PtrStep_ weights2, PtrStep_ 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 diff --git a/modules/gpu/src/cuda/blend.cu b/modules/gpu/src/cuda/blend.cu new file mode 100644 index 0000000000..a9b85c1549 --- /dev/null +++ b/modules/gpu/src/cuda/blend.cu @@ -0,0 +1,117 @@ +/*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 + __global__ void blendLinearKernel(int rows, int cols, int cn, 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) + { + 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 + void blendLinearCaller(int rows, int cols, int cn, const PtrStep_ img1, const PtrStep_ img2, + const PtrStepf weights1, const PtrStepf weights2, PtrStep_ result) + { + dim3 threads(16, 16); + dim3 grid(divUp(cols * cn, threads.x), divUp(rows, threads.y)); + + blendLinearKernel<<>>(rows, cols * cn, cn, img1, img2, weights1, weights2, result); + cudaSafeCall(cudaThreadSynchronize()); + } + + template void blendLinearCaller(int, int, int, const PtrStep, const PtrStep, + const PtrStepf, const PtrStepf, PtrStep); + template void blendLinearCaller(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<<>>(rows, cols, img1, img2, weights1, weights2, result); + cudaSafeCall(cudaThreadSynchronize()); + } + +}} \ 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 index 0000000000..60955ad81d --- /dev/null +++ b/modules/gpu/test/test_blend.cpp @@ -0,0 +1,96 @@ +/*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(y, x / cn); + float w2 = weights2.at(y, x / cn); + result_gold.at(y, x) = static_cast( + (img1.at(y, x) * w1 + img2.at(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(y, x / cn); + float w2 = weights2.at(y, x / cn); + result_gold.at(y, x) = + (img1.at(y, x) * w1 + img2.at(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