diff --git a/modules/stitching/doc/warpers.rst b/modules/stitching/doc/warpers.rst index 1025ffa0cb..278020f07c 100644 --- a/modules/stitching/doc/warpers.rst +++ b/modules/stitching/doc/warpers.rst @@ -14,17 +14,17 @@ Rotation-only model image warper interface. :: public: virtual ~RotationWarper() {} - virtual Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R) = 0; + virtual Point2f warpPoint(const Point2f &pt, InputArray K, InputArray R) = 0; - virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) = 0; + virtual Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) = 0; - virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Mat &dst) = 0; + virtual Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + OutputArray dst) = 0; - virtual void warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Size dst_size, Mat &dst) = 0; + virtual void warpBackward(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + Size dst_size, OutputArray dst) = 0; - virtual Rect warpRoi(Size src_size, const Mat &K, const Mat &R) = 0; + virtual Rect warpRoi(Size src_size, InputArray K, InputArray R) = 0; }; detail::RotationWarper::warpPoint @@ -32,7 +32,7 @@ detail::RotationWarper::warpPoint Projects the image point. -.. ocv:function:: Point2f detail::RotationWarper::warpPoint(const Point2f &pt, const Mat &K, const Mat &R) +.. ocv:function:: Point2f detail::RotationWarper::warpPoint(const Point2f &pt, InputArray K, InputArray R) :param pt: Source point @@ -47,7 +47,7 @@ detail::RotationWarper::buildMaps Builds the projection maps according to the given camera data. -.. ocv:function:: Rect detail::RotationWarper::buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) +.. ocv:function:: Rect detail::RotationWarper::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) :param src_size: Source image size @@ -66,7 +66,7 @@ detail::RotationWarper::warp Projects the image. -.. ocv:function:: Point detail::RotationWarper::warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Mat &dst) +.. ocv:function:: Point detail::RotationWarper::warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst) :param src: Source image @@ -87,7 +87,7 @@ detail::RotationWarper::warpBackward Projects the image backward. -.. ocv:function:: void detail::RotationWarper::warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Size dst_size, Mat &dst) +.. ocv:function:: void detail::RotationWarper::warpBackward(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, Size dst_size, OutputArray dst) :param src: Projected image @@ -106,7 +106,7 @@ Projects the image backward. detail::RotationWarper::warpRoi ------------------------------- -.. ocv:function:: Rect detail::RotationWarper::warpRoi(Size src_size, const Mat &K, const Mat &R) +.. ocv:function:: Rect detail::RotationWarper::warpRoi(Size src_size, InputArray K, InputArray R) :param src_size: Source image bounding box @@ -124,9 +124,9 @@ Base class for warping logic implementation. :: struct CV_EXPORTS ProjectorBase { - void setCameraParams(const Mat &K = Mat::eye(3, 3, CV_32F), - const Mat &R = Mat::eye(3, 3, CV_32F), - const Mat &T = Mat::zeros(3, 1, CV_32F)); + void setCameraParams(InputArray K = Mat::eye(3, 3, CV_32F), + InputArray R = Mat::eye(3, 3, CV_32F), + InputArray T = Mat::zeros(3, 1, CV_32F)); float scale; float k[9]; @@ -146,17 +146,17 @@ Base class for rotation-based warper using a `detail::ProjectorBase`_ derived cl class CV_EXPORTS RotationWarperBase : public RotationWarper { public: - Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R); + Point2f warpPoint(const Point2f &pt, InputArray K, InputArray R); - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap); + Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap); - Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Mat &dst); + Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + OutputArray dst); - void warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Size dst_size, Mat &dst); + void warpBackward(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + Size dst_size, OutputArray dst); - Rect warpRoi(Size src_size, const Mat &K, const Mat &R); + Rect warpRoi(Size src_size, InputArray K, InputArray R); protected: @@ -183,14 +183,14 @@ Warper that maps an image onto the z = 1 plane. :: void setScale(float scale) { projector_.scale = scale; } - Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R, const Mat &T); + Point2f warpPoint(const Point2f &pt, InputArray K, InputArray R, InputArray T); - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap); + Rect buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray xmap, OutputArray ymap); - Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, - Mat &dst); + Point warp(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, + OutputArray dst); - Rect warpRoi(Size src_size, const Mat &K, const Mat &R, const Mat &T); + Rect warpRoi(Size src_size, InputArray K, InputArray R, InputArray T); protected: void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br); diff --git a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp index 8906d88a3c..093f07cc1a 100644 --- a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp @@ -56,17 +56,17 @@ class CV_EXPORTS RotationWarper public: virtual ~RotationWarper() {} - virtual Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R) = 0; + virtual Point2f warpPoint(const Point2f &pt, InputArray K, InputArray R) = 0; - virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) = 0; + virtual Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) = 0; - virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Mat &dst) = 0; + virtual Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + OutputArray dst) = 0; - virtual void warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Size dst_size, Mat &dst) = 0; + virtual void warpBackward(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + Size dst_size, OutputArray dst) = 0; - virtual Rect warpRoi(Size src_size, const Mat &K, const Mat &R) = 0; + virtual Rect warpRoi(Size src_size, InputArray K, InputArray R) = 0; virtual float getScale() const { return 1.f; } virtual void setScale(float) {} @@ -75,9 +75,9 @@ public: struct CV_EXPORTS ProjectorBase { - void setCameraParams(const Mat &K = Mat::eye(3, 3, CV_32F), - const Mat &R = Mat::eye(3, 3, CV_32F), - const Mat &T = Mat::zeros(3, 1, CV_32F)); + void setCameraParams(InputArray K = Mat::eye(3, 3, CV_32F), + InputArray R = Mat::eye(3, 3, CV_32F), + InputArray T = Mat::zeros(3, 1, CV_32F)); float scale; float k[9]; @@ -92,17 +92,17 @@ template class CV_EXPORTS RotationWarperBase : public RotationWarper { public: - Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R); + Point2f warpPoint(const Point2f &pt, InputArray K, InputArray R); - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap); + Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap); - Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Mat &dst); + Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + OutputArray dst); - void warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Size dst_size, Mat &dst); + void warpBackward(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + Size dst_size, OutputArray dst); - Rect warpRoi(Size src_size, const Mat &K, const Mat &R); + Rect warpRoi(Size src_size, InputArray K, InputArray R); float getScale() const { return projector_.scale; } void setScale(float val) { projector_.scale = val; } @@ -132,14 +132,14 @@ class CV_EXPORTS PlaneWarper : public RotationWarperBase public: PlaneWarper(float scale = 1.f) { projector_.scale = scale; } - Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R, const Mat &T); + Point2f warpPoint(const Point2f &pt, InputArray K, InputArray R, InputArray T); - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap); + virtual Rect buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray xmap, OutputArray ymap); - Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, - Mat &dst); + virtual Point warp(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, + OutputArray dst); - Rect warpRoi(Size src_size, const Mat &K, const Mat &R, const Mat &T); + Rect warpRoi(Size src_size, InputArray K, InputArray R, InputArray T); protected: void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br); @@ -333,7 +333,7 @@ class CV_EXPORTS PlaneWarperGpu : public PlaneWarper public: PlaneWarperGpu(float scale = 1.f) : PlaneWarper(scale) {} - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) + Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) { Rect result = buildMaps(src_size, K, R, d_xmap_, d_ymap_); d_xmap_.download(xmap); @@ -341,7 +341,7 @@ public: return result; } - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap) + Rect buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray xmap, OutputArray ymap) { Rect result = buildMaps(src_size, K, R, T, d_xmap_, d_ymap_); d_xmap_.download(xmap); @@ -349,8 +349,8 @@ public: return result; } - Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Mat &dst) + Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + OutputArray dst) { d_src_.upload(src); Point result = warp(d_src_, K, R, interp_mode, border_mode, d_dst_); @@ -358,8 +358,8 @@ public: return result; } - Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, - Mat &dst) + Point warp(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, + OutputArray dst) { d_src_.upload(src); Point result = warp(d_src_, K, R, T, interp_mode, border_mode, d_dst_); @@ -367,15 +367,15 @@ public: return result; } - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, cuda::GpuMat &xmap, cuda::GpuMat &ymap); + Rect buildMaps(Size src_size, InputArray K, InputArray R, cuda::GpuMat & xmap, cuda::GpuMat & ymap); - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, cuda::GpuMat &xmap, cuda::GpuMat &ymap); + Rect buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, cuda::GpuMat & xmap, cuda::GpuMat & ymap); - Point warp(const cuda::GpuMat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - cuda::GpuMat &dst); + Point warp(const cuda::GpuMat & src, InputArray K, InputArray R, int interp_mode, int border_mode, + cuda::GpuMat & dst); - Point warp(const cuda::GpuMat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, - cuda::GpuMat &dst); + Point warp(const cuda::GpuMat & src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, + cuda::GpuMat & dst); private: cuda::GpuMat d_xmap_, d_ymap_, d_src_, d_dst_; @@ -387,7 +387,7 @@ class CV_EXPORTS SphericalWarperGpu : public SphericalWarper public: SphericalWarperGpu(float scale) : SphericalWarper(scale) {} - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) + Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) { Rect result = buildMaps(src_size, K, R, d_xmap_, d_ymap_); d_xmap_.download(xmap); @@ -395,8 +395,8 @@ public: return result; } - Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Mat &dst) + Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + OutputArray dst) { d_src_.upload(src); Point result = warp(d_src_, K, R, interp_mode, border_mode, d_dst_); @@ -404,10 +404,10 @@ public: return result; } - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, cuda::GpuMat &xmap, cuda::GpuMat &ymap); + Rect buildMaps(Size src_size, InputArray K, InputArray R, cuda::GpuMat & xmap, cuda::GpuMat & ymap); - Point warp(const cuda::GpuMat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - cuda::GpuMat &dst); + Point warp(const cuda::GpuMat & src, InputArray K, InputArray R, int interp_mode, int border_mode, + cuda::GpuMat & dst); private: cuda::GpuMat d_xmap_, d_ymap_, d_src_, d_dst_; @@ -419,7 +419,7 @@ class CV_EXPORTS CylindricalWarperGpu : public CylindricalWarper public: CylindricalWarperGpu(float scale) : CylindricalWarper(scale) {} - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) + Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) { Rect result = buildMaps(src_size, K, R, d_xmap_, d_ymap_); d_xmap_.download(xmap); @@ -427,8 +427,8 @@ public: return result; } - Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Mat &dst) + Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + OutputArray dst) { d_src_.upload(src); Point result = warp(d_src_, K, R, interp_mode, border_mode, d_dst_); @@ -436,10 +436,10 @@ public: return result; } - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, cuda::GpuMat &xmap, cuda::GpuMat &ymap); + Rect buildMaps(Size src_size, InputArray K, InputArray R, cuda::GpuMat & xmap, cuda::GpuMat & ymap); - Point warp(const cuda::GpuMat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - cuda::GpuMat &dst); + Point warp(const cuda::GpuMat & src, InputArray K, InputArray R, int interp_mode, int border_mode, + cuda::GpuMat & dst); private: cuda::GpuMat d_xmap_, d_ymap_, d_src_, d_dst_; @@ -503,6 +503,45 @@ protected: } }; +/////////////////////////////////////// OpenCL Accelerated Warpers ///////////////////////////////////// + +class CV_EXPORTS PlaneWarperOcl : public PlaneWarper +{ +public: + PlaneWarperOcl(float scale = 1.f) : PlaneWarper(scale) { } + + virtual Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) + { + return buildMaps(src_size, K, R, Mat::zeros(3, 1, CV_32FC1), xmap, ymap); + } + + virtual Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst) + { + return warp(src, K, R, Mat::zeros(3, 1, CV_32FC1), interp_mode, border_mode, dst); + } + + virtual Rect buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray xmap, OutputArray ymap); + virtual Point warp(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, OutputArray dst); +}; + +class CV_EXPORTS SphericalWarperOcl : public SphericalWarper +{ +public: + SphericalWarperOcl(float scale) : SphericalWarper(scale) { } + + virtual Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap); + virtual Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst); +}; + +class CV_EXPORTS CylindricalWarperOcl : public CylindricalWarper +{ +public: + CylindricalWarperOcl(float scale) : CylindricalWarper(scale) { } + + virtual Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap); + virtual Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst); +}; + } // namespace detail } // namespace cv diff --git a/modules/stitching/include/opencv2/stitching/detail/warpers_inl.hpp b/modules/stitching/include/opencv2/stitching/detail/warpers_inl.hpp index 144e9e32d4..f6eae4fa78 100644 --- a/modules/stitching/include/opencv2/stitching/detail/warpers_inl.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/warpers_inl.hpp @@ -51,7 +51,7 @@ namespace cv { namespace detail { template -Point2f RotationWarperBase

::warpPoint(const Point2f &pt, const Mat &K, const Mat &R) +Point2f RotationWarperBase

::warpPoint(const Point2f &pt, InputArray K, InputArray R) { projector_.setCameraParams(K, R); Point2f uv; @@ -61,15 +61,17 @@ Point2f RotationWarperBase

::warpPoint(const Point2f &pt, const Mat &K, const template -Rect RotationWarperBase

::buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) +Rect RotationWarperBase

::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray _xmap, OutputArray _ymap) { projector_.setCameraParams(K, R); Point dst_tl, dst_br; detectResultRoi(src_size, dst_tl, dst_br); - xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); - ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); + _xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); + _ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); + + Mat xmap = _xmap.getMat(), ymap = _ymap.getMat(); float x, y; for (int v = dst_tl.y; v <= dst_br.y; ++v) @@ -87,8 +89,8 @@ Rect RotationWarperBase

::buildMaps(Size src_size, const Mat &K, const Mat &R, template -Point RotationWarperBase

::warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Mat &dst) +Point RotationWarperBase

::warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + OutputArray dst) { Mat xmap, ymap; Rect dst_roi = buildMaps(src.size(), K, R, xmap, ymap); @@ -101,14 +103,16 @@ Point RotationWarperBase

::warp(const Mat &src, const Mat &K, const Mat &R, in template -void RotationWarperBase

::warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - Size dst_size, Mat &dst) +void RotationWarperBase

::warpBackward(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, + Size dst_size, OutputArray dst) { projector_.setCameraParams(K, R); Point src_tl, src_br; detectResultRoi(dst_size, src_tl, src_br); - CV_Assert(src_br.x - src_tl.x + 1 == src.cols && src_br.y - src_tl.y + 1 == src.rows); + + Size size = src.size(); + CV_Assert(src_br.x - src_tl.x + 1 == size.width && src_br.y - src_tl.y + 1 == size.height); Mat xmap(dst_size, CV_32F); Mat ymap(dst_size, CV_32F); @@ -130,7 +134,7 @@ void RotationWarperBase

::warpBackward(const Mat &src, const Mat &K, const Mat template -Rect RotationWarperBase

::warpRoi(Size src_size, const Mat &K, const Mat &R) +Rect RotationWarperBase

::warpRoi(Size src_size, InputArray K, InputArray R) { projector_.setCameraParams(K, R); diff --git a/modules/stitching/include/opencv2/stitching/warpers.hpp b/modules/stitching/include/opencv2/stitching/warpers.hpp index da5fe26183..cdcb35c201 100644 --- a/modules/stitching/include/opencv2/stitching/warpers.hpp +++ b/modules/stitching/include/opencv2/stitching/warpers.hpp @@ -167,6 +167,24 @@ public: }; #endif +class PlaneWarperOcl: public WarperCreator +{ +public: + Ptr create(float scale) const { return makePtr(scale); } +}; + +class SphericalWarperOcl: public WarperCreator +{ +public: + Ptr create(float scale) const { return makePtr(scale); } +}; + +class CylindricalWarperOcl: public WarperCreator +{ +public: + Ptr create(float scale) const { return makePtr(scale); } +}; + } // namespace cv #endif // __OPENCV_STITCHING_WARPER_CREATORS_HPP__ diff --git a/modules/stitching/src/opencl/warpers.cl b/modules/stitching/src/opencl/warpers.cl new file mode 100644 index 0000000000..032ddf3cee --- /dev/null +++ b/modules/stitching/src/opencl/warpers.cl @@ -0,0 +1,148 @@ +/*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) 2010-2012, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Peng Xiao, pengxiao@multicorewareinc.com +// +// 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 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*/ + +__kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset, + __global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols, + __constant float * ck_rinv, __constant float * ct, + int tl_u, int tl_v, float scale) +{ + int du = get_global_id(0); + int dv = get_global_id(1); + + if (du < cols && dv < rows) + { + __global float * xmap = (__global float *)(xmapptr + mad24(dv, xmap_step, xmap_offset + du * (int)sizeof(float))); + __global float * ymap = (__global float *)(ymapptr + mad24(dv, ymap_step, ymap_offset + du * (int)sizeof(float))); + + float u = tl_u + du; + float v = tl_v + dv; + float x, y; + + float x_ = u / scale - ct[0]; + float y_ = v / scale - ct[1]; + + float z; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * (1 - ct[2]); + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * (1 - ct[2]); + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * (1 - ct[2]); + + x /= z; + y /= z; + + xmap[0] = x; + ymap[0] = y; + } +} + +__kernel void buildWarpCylindricalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset, + __global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols, + __constant float * ck_rinv, int tl_u, int tl_v, float scale) +{ + int du = get_global_id(0); + int dv = get_global_id(1); + + if (du < cols && dv < rows) + { + __global float * xmap = (__global float *)(xmapptr + mad24(dv, xmap_step, xmap_offset + du * (int)sizeof(float))); + __global float * ymap = (__global float *)(ymapptr + mad24(dv, ymap_step, ymap_offset + du * (int)sizeof(float))); + + float u = tl_u + du; + float v = tl_v + dv; + float x, y; + + u /= scale; + float x_ = sin(u); + float y_ = v / scale; + float z_ = cos(u); + + float z; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; + + if (z > 0) x /= z, y /= z; + else x = y = -1; + + xmap[0] = x; + ymap[0] = y; + } +} + +__kernel void buildWarpSphericalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset, + __global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols, + __constant float * ck_rinv, int tl_u, int tl_v, float scale) +{ + int du = get_global_id(0); + int dv = get_global_id(1); + + if (du < cols && dv < rows) + { + __global float * xmap = (__global float *)(xmapptr + mad24(dv, xmap_step, xmap_offset + du * (int)sizeof(float))); + __global float * ymap = (__global float *)(ymapptr + mad24(dv, ymap_step, ymap_offset + du * (int)sizeof(float))); + + float u = tl_u + du; + float v = tl_v + dv; + float x, y; + + v /= scale; + u /= scale; + + float sinv = sin(v); + float x_ = sinv * sin(u); + float y_ = -cos(v); + float z_ = sinv * cos(u); + + float z; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; + + if (z > 0) x /= z, y /= z; + else x = y = -1; + + xmap[0] = x; + ymap[0] = y; + } +} diff --git a/modules/stitching/src/precomp.hpp b/modules/stitching/src/precomp.hpp index 04445176e0..499202fa0f 100644 --- a/modules/stitching/src/precomp.hpp +++ b/modules/stitching/src/precomp.hpp @@ -53,6 +53,7 @@ #include #include #include "opencv2/core.hpp" +#include "opencv2/core/ocl.hpp" #include "opencv2/core/utility.hpp" #include "opencv2/stitching.hpp" #include "opencv2/stitching/detail/autocalib.hpp" diff --git a/modules/stitching/src/warpers.cpp b/modules/stitching/src/warpers.cpp index 3f71f20405..eb15d44c04 100644 --- a/modules/stitching/src/warpers.cpp +++ b/modules/stitching/src/warpers.cpp @@ -45,8 +45,10 @@ namespace cv { namespace detail { -void ProjectorBase::setCameraParams(const Mat &K, const Mat &R, const Mat &T) +void ProjectorBase::setCameraParams(InputArray _K, InputArray _R, InputArray _T) { + Mat K = _K.getMat(), R = _R.getMat(), T = _T.getMat(); + CV_Assert(K.size() == Size(3, 3) && K.type() == CV_32F); CV_Assert(R.size() == Size(3, 3) && R.type() == CV_32F); CV_Assert((T.size() == Size(1, 3) || T.size() == Size(3, 1)) && T.type() == CV_32F); @@ -76,7 +78,7 @@ void ProjectorBase::setCameraParams(const Mat &K, const Mat &R, const Mat &T) } -Point2f PlaneWarper::warpPoint(const Point2f &pt, const Mat &K, const Mat &R, const Mat &T) +Point2f PlaneWarper::warpPoint(const Point2f &pt, InputArray K, InputArray R, InputArray T) { projector_.setCameraParams(K, R, T); Point2f uv; @@ -85,15 +87,17 @@ Point2f PlaneWarper::warpPoint(const Point2f &pt, const Mat &K, const Mat &R, co } -Rect PlaneWarper::buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap) +Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray _xmap, OutputArray _ymap) { projector_.setCameraParams(K, R, T); Point dst_tl, dst_br; detectResultRoi(src_size, dst_tl, dst_br); - xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); - ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); + _xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); + _ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); + + Mat xmap = _xmap.getMat(), ymap = _ymap.getMat(); float x, y; for (int v = dst_tl.y; v <= dst_br.y; ++v) @@ -110,8 +114,8 @@ Rect PlaneWarper::buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat } -Point PlaneWarper::warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, - Mat &dst) +Point PlaneWarper::warp(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, + OutputArray dst) { Mat xmap, ymap; Rect dst_roi = buildMaps(src.size(), K, R, T, xmap, ymap); @@ -123,7 +127,7 @@ Point PlaneWarper::warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T } -Rect PlaneWarper::warpRoi(Size src_size, const Mat &K, const Mat &R, const Mat &T) +Rect PlaneWarper::warpRoi(Size src_size, InputArray K, InputArray R, InputArray T) { projector_.setCameraParams(K, R, T); @@ -211,12 +215,12 @@ void SphericalWarper::detectResultRoi(Size src_size, Point &dst_tl, Point &dst_b #ifdef HAVE_OPENCV_CUDAWARPING -Rect PlaneWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, cuda::GpuMat &xmap, cuda::GpuMat &ymap) +Rect PlaneWarperGpu::buildMaps(Size src_size, InputArray K, InputArray R, cuda::GpuMat & xmap, cuda::GpuMat & ymap) { return buildMaps(src_size, K, R, Mat::zeros(3, 1, CV_32F), xmap, ymap); } -Rect PlaneWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, cuda::GpuMat &xmap, cuda::GpuMat &ymap) +Rect PlaneWarperGpu::buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, cuda::GpuMat & xmap, cuda::GpuMat & ymap) { projector_.setCameraParams(K, R, T); @@ -229,15 +233,15 @@ Rect PlaneWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, const return Rect(dst_tl, dst_br); } -Point PlaneWarperGpu::warp(const cuda::GpuMat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - cuda::GpuMat &dst) +Point PlaneWarperGpu::warp(const cuda::GpuMat & src, InputArray K, InputArray R, int interp_mode, int border_mode, + cuda::GpuMat & dst) { return warp(src, K, R, Mat::zeros(3, 1, CV_32F), interp_mode, border_mode, dst); } -Point PlaneWarperGpu::warp(const cuda::GpuMat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, - cuda::GpuMat &dst) +Point PlaneWarperGpu::warp(const cuda::GpuMat & src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, + cuda::GpuMat & dst) { Rect dst_roi = buildMaps(src.size(), K, R, T, d_xmap_, d_ymap_); dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); @@ -246,7 +250,7 @@ Point PlaneWarperGpu::warp(const cuda::GpuMat &src, const Mat &K, const Mat &R, } -Rect SphericalWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, cuda::GpuMat &xmap, cuda::GpuMat &ymap) +Rect SphericalWarperGpu::buildMaps(Size src_size, InputArray K, InputArray R, cuda::GpuMat & xmap, cuda::GpuMat & ymap) { projector_.setCameraParams(K, R); @@ -260,8 +264,8 @@ Rect SphericalWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, cu } -Point SphericalWarperGpu::warp(const cuda::GpuMat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - cuda::GpuMat &dst) +Point SphericalWarperGpu::warp(const cuda::GpuMat & src, InputArray K, InputArray R, int interp_mode, int border_mode, + cuda::GpuMat & dst) { Rect dst_roi = buildMaps(src.size(), K, R, d_xmap_, d_ymap_); dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); @@ -270,7 +274,7 @@ Point SphericalWarperGpu::warp(const cuda::GpuMat &src, const Mat &K, const Mat } -Rect CylindricalWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, cuda::GpuMat &xmap, cuda::GpuMat &ymap) +Rect CylindricalWarperGpu::buildMaps(Size src_size, InputArray K, InputArray R, cuda::GpuMat & xmap, cuda::GpuMat & ymap) { projector_.setCameraParams(K, R); @@ -284,8 +288,8 @@ Rect CylindricalWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, } -Point CylindricalWarperGpu::warp(const cuda::GpuMat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, - cuda::GpuMat &dst) +Point CylindricalWarperGpu::warp(const cuda::GpuMat & src, InputArray K, InputArray R, int interp_mode, int border_mode, + cuda::GpuMat & dst) { Rect dst_roi = buildMaps(src.size(), K, R, d_xmap_, d_ymap_); dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); diff --git a/modules/stitching/src/warpers_ocl.cpp b/modules/stitching/src/warpers_ocl.cpp new file mode 100644 index 0000000000..ce4b8946d8 --- /dev/null +++ b/modules/stitching/src/warpers_ocl.cpp @@ -0,0 +1,187 @@ +/*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 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 "precomp.hpp" +#include "opencl_kernels.hpp" + +namespace cv { +namespace detail { + +/////////////////////////////////////////// PlaneWarperOcl //////////////////////////////////////////// + +Rect PlaneWarperOcl::buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray xmap, OutputArray ymap) +{ + projector_.setCameraParams(K, R); + + Point dst_tl, dst_br; + detectResultRoi(src_size, dst_tl, dst_br); + + if (ocl::useOpenCL()) + { + ocl::Kernel k("buildWarpPlaneMaps", ocl::stitching::warpers_oclsrc); + if (!k.empty()) + { + Size dsize(dst_br.x - dst_tl.x + 1, dst_br.y - dst_tl.y + 1); + xmap.create(dsize, CV_32FC1); + ymap.create(dsize, CV_32FC1); + + Mat r_kinv(1, 9, CV_32FC1, projector_.r_kinv), t(1, 3, CV_32FC1, projector_.t); + UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), + ur_kinv = r_kinv.getUMat(ACCESS_READ), ut = t.getUMat(ACCESS_READ); + + k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), + ocl::KernelArg::PtrReadOnly(ur_kinv), ocl::KernelArg::PtrReadOnly(ut), + dst_tl.x, dst_tl.y, projector_.scale); + + size_t globalsize[2] = { dsize.width, dsize.height }; + if (k.run(2, globalsize, NULL, true)) + return Rect(dst_tl, dst_br); + } + } + + return PlaneWarper::buildMaps(src_size, K, R, T, xmap, ymap); +} + +Point PlaneWarperOcl::warp(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, OutputArray dst) +{ + UMat uxmap, uymap; + Rect dst_roi = buildMaps(src.size(), K, R, T, uxmap, uymap); + + dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); + UMat udst = dst.getUMat(); + remap(src, udst, uxmap, uymap, interp_mode, border_mode); + + return dst_roi.tl(); +} + +/////////////////////////////////////////// SphericalWarperOcl //////////////////////////////////////// + +Rect SphericalWarperOcl::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) +{ + projector_.setCameraParams(K, R); + + Point dst_tl, dst_br; + detectResultRoi(src_size, dst_tl, dst_br); + + if (ocl::useOpenCL()) + { + ocl::Kernel k("buildWarpSphericalMaps", ocl::stitching::warpers_oclsrc); + if (!k.empty()) + { + Size dsize(dst_br.x - dst_tl.x + 1, dst_br.y - dst_tl.y + 1); + xmap.create(dsize, CV_32FC1); + ymap.create(dsize, CV_32FC1); + + Mat r_kinv(1, 9, CV_32FC1, projector_.r_kinv); + UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), ur_kinv = r_kinv.getUMat(ACCESS_READ); + + k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), + ocl::KernelArg::PtrReadOnly(ur_kinv), dst_tl.x, dst_tl.y, projector_.scale); + + size_t globalsize[2] = { dsize.width, dsize.height }; + if (k.run(2, globalsize, NULL, true)) + return Rect(dst_tl, dst_br); + } + } + + return SphericalWarper::buildMaps(src_size, K, R, xmap, ymap); +} + +Point SphericalWarperOcl::warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst) +{ + UMat uxmap, uymap; + Rect dst_roi = buildMaps(src.size(), K, R, uxmap, uymap); + + dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); + UMat udst = dst.getUMat(); + remap(src, udst, uxmap, uymap, interp_mode, border_mode); + + return dst_roi.tl(); +} + +/////////////////////////////////////////// CylindricalWarperOcl //////////////////////////////////////// + +Rect CylindricalWarperOcl::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) +{ + projector_.setCameraParams(K, R); + + Point dst_tl, dst_br; + detectResultRoi(src_size, dst_tl, dst_br); + + if (ocl::useOpenCL()) + { + ocl::Kernel k("buildWarpCylindricalMaps", ocl::stitching::warpers_oclsrc); + if (!k.empty()) + { + Size dsize(dst_br.x - dst_tl.x + 1, dst_br.y - dst_tl.y + 1); + xmap.create(dsize, CV_32FC1); + ymap.create(dsize, CV_32FC1); + + Mat r_kinv(1, 9, CV_32FC1, projector_.r_kinv); + UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), ur_kinv = r_kinv.getUMat(ACCESS_READ); + + k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), + ocl::KernelArg::PtrReadOnly(ur_kinv), dst_tl.x, dst_tl.y, projector_.scale); + + size_t globalsize[2] = { dsize.width, dsize.height }; + if (k.run(2, globalsize, NULL, true)) + return Rect(dst_tl, dst_br); + } + } + + return CylindricalWarper::buildMaps(src_size, K, R, xmap, ymap); +} + +Point CylindricalWarperOcl::warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst) +{ + UMat uxmap, uymap; + Rect dst_roi = buildMaps(src.size(), K, R, uxmap, uymap); + + dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); + UMat udst = dst.getUMat(); + remap(src, udst, uxmap, uymap, interp_mode, border_mode); + + return dst_roi.tl(); +} + +} // namespace detail +} // namespace cv diff --git a/modules/stitching/test/ocl/test_warpers.cpp b/modules/stitching/test/ocl/test_warpers.cpp new file mode 100644 index 0000000000..d2f5dc0091 --- /dev/null +++ b/modules/stitching/test/ocl/test_warpers.cpp @@ -0,0 +1,149 @@ +/*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) 2010-2013, Advanced Micro Devices, 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 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 OpenCV Foundation 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" +#include "opencv2/ts/ocl_test.hpp" +#include "opencv2/stitching/warpers.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +///////////////////////// WarperTestBase /////////////////////////// + +struct WarperTestBase : + public Test, public TestUtils +{ + Mat src, dst, xmap, ymap; + Mat udst, uxmap, uymap; + Mat K, R; + + virtual void generateTestData() + { + Size size = randomSize(1, MAX_VALUE); + + src = randomMat(size, CV_32FC1, -500, 500); + + K = Mat::eye(3, 3, CV_32FC1); + R = Mat::eye(3, 3, CV_32FC1); + } + + void Near(double threshold = 0.) + { + EXPECT_MAT_NEAR(xmap, uxmap, threshold); + EXPECT_MAT_NEAR(ymap, uymap, threshold); + EXPECT_MAT_NEAR(dst, udst, threshold); + } +}; + +//////////////////////////////// SphericalWarperOcl ///////////////////////////////////////////////// + +typedef WarperTestBase SphericalWarperOclTest; + +OCL_TEST_F(SphericalWarperOclTest, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + Ptr creator = makePtr(); + Ptr warper = creator->create(2.0); + + OCL_OFF(warper->buildMaps(src.size(), K, R, xmap, ymap)); + OCL_ON(warper->buildMaps(src.size(), K, R, uxmap, uymap)); + + OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); + OCL_ON(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); + + Near(1e-4); + } +} + +//////////////////////////////// CylindricalWarperOcl ///////////////////////////////////////////////// + +typedef WarperTestBase CylindricalWarperOclTest; + +OCL_TEST_F(CylindricalWarperOclTest, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + Ptr creator = makePtr(); + Ptr warper = creator->create(2.0); + + OCL_OFF(warper->buildMaps(src.size(), K, R, xmap, ymap)); + OCL_ON(warper->buildMaps(src.size(), K, R, uxmap, uymap)); + + OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); + OCL_ON(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); + + Near(1e-4); + } +} + +//////////////////////////////// PlaneWarperOcl ///////////////////////////////////////////////// + +typedef WarperTestBase PlaneWarperOclTest; + +OCL_TEST_F(PlaneWarperOclTest, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + Ptr creator = makePtr(); + Ptr warper = creator->create(2.0); + + OCL_OFF(warper->buildMaps(src.size(), K, R, xmap, ymap)); + OCL_ON(warper->buildMaps(src.size(), K, R, uxmap, uymap)); + + OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); + OCL_ON(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); + + Near(1e-5); + } +} + +} } // namespace cvtest::ocl + +#endif // HAVE_OPENCL diff --git a/samples/cpp/stitching_detailed.cpp b/samples/cpp/stitching_detailed.cpp index b6eefc6f99..5eb3df46cc 100644 --- a/samples/cpp/stitching_detailed.cpp +++ b/samples/cpp/stitching_detailed.cpp @@ -71,8 +71,11 @@ static void printUsage() " --preview\n" " Run stitching in the preview mode. Works faster than usual mode,\n" " but output image will have lower resolution.\n" - " --try_gpu (yes|no)\n" - " Try to use GPU. The default value is 'no'. All default values\n" + " --try_cuda (yes|no)\n" + " Try to use CUDA. The default value is 'no'. All default values\n" + " are for CPU mode.\n" + " --try_ocl (yes|no)\n" + " Try to use OpenCL. The default value is 'no'. All default values\n" " are for CPU mode.\n" "\nMotion Estimation Flags:\n" " --work_megapix \n" @@ -123,7 +126,8 @@ static void printUsage() // Default command line args vector img_names; bool preview = false; -bool try_gpu = false; +bool try_cuda = false; +bool try_ocl = false; double work_megapix = 0.6; double seam_megapix = 0.1; double compose_megapix = -1; @@ -161,15 +165,28 @@ static int parseCmdArgs(int argc, char** argv) { preview = true; } - else if (string(argv[i]) == "--try_gpu") + else if (string(argv[i]) == "--try_cuda") { if (string(argv[i + 1]) == "no") - try_gpu = false; + try_cuda = false; else if (string(argv[i + 1]) == "yes") - try_gpu = true; + try_cuda = true; else { - cout << "Bad --try_gpu flag value\n"; + cout << "Bad --try_cuda flag value\n"; + return -1; + } + i++; + } + else if (string(argv[i]) == "--try_ocl") + { + if (string(argv[i + 1]) == "no") + try_ocl = false; + else if (string(argv[i + 1]) == "yes") + try_ocl = true; + else + { + cout << "Bad --try_ocl flag value\n"; return -1; } i++; @@ -357,7 +374,7 @@ int main(int argc, char* argv[]) if (features_type == "surf") { #ifdef HAVE_OPENCV_NONFREE - if (try_gpu && cuda::getCudaEnabledDeviceCount() > 0) + if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) finder = makePtr(); else #endif @@ -430,7 +447,7 @@ int main(int argc, char* argv[]) t = getTickCount(); #endif vector pairwise_matches; - BestOf2NearestMatcher matcher(try_gpu, match_conf); + BestOf2NearestMatcher matcher(try_cuda, match_conf); matcher(features, pairwise_matches); matcher.collectGarbage(); LOGLN("Pairwise matching, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); @@ -552,8 +569,17 @@ int main(int argc, char* argv[]) // Warp images and their masks Ptr warper_creator; + if (try_ocl) + { + if (warp_type == "plane") + warper_creator = makePtr(); + else if (warp_type == "cylindrical") + warper_creator = makePtr(); + else if (warp_type == "spherical") + warper_creator = makePtr(); + } #ifdef HAVE_OPENCV_CUDAWARPING - if (try_gpu && cuda::getCudaEnabledDeviceCount() > 0) + else if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) { if (warp_type == "plane") warper_creator = makePtr(); @@ -636,7 +662,7 @@ int main(int argc, char* argv[]) else if (seam_find_type == "gc_color") { #ifdef HAVE_OPENCV_CUDA - if (try_gpu && cuda::getCudaEnabledDeviceCount() > 0) + if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) seam_finder = makePtr(GraphCutSeamFinderBase::COST_COLOR); else #endif @@ -645,7 +671,7 @@ int main(int argc, char* argv[]) else if (seam_find_type == "gc_colorgrad") { #ifdef HAVE_OPENCV_CUDA - if (try_gpu && cuda::getCudaEnabledDeviceCount() > 0) + if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) seam_finder = makePtr(GraphCutSeamFinderBase::COST_COLOR_GRAD); else #endif @@ -755,11 +781,11 @@ int main(int argc, char* argv[]) if (!blender) { - blender = Blender::createDefault(blend_type, try_gpu); + blender = Blender::createDefault(blend_type, try_cuda); Size dst_sz = resultRoi(corners, sizes).size(); float blend_width = sqrt(static_cast(dst_sz.area())) * blend_strength / 100.f; if (blend_width < 1.f) - blender = Blender::createDefault(Blender::NO, try_gpu); + blender = Blender::createDefault(Blender::NO, try_cuda); else if (blend_type == Blender::MULTI_BAND) { MultiBandBlender* mb = dynamic_cast(blender.get());