added OpenCL accelerated warpers

modules/stitching/doc/warpers.rst~

+Images Warping
+==============
+
+.. highlight:: cpp
+
+detail::RotationWarper
+----------------------
+.. ocv:class:: detail::RotationWarper
+
+Rotation-only model image warper interface. ::
+
+    class CV_EXPORTS RotationWarper
+    {
+    public:
+        virtual ~RotationWarper() {}
+
+        virtual Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R) = 0;
+
+        virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) = 0;
+
+        virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode,
+                           Mat &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 Rect warpRoi(Size src_size, const Mat &K, const Mat &R) = 0;
+    };
+
+detail::RotationWarper::warpPoint
+---------------------------------
+
+Projects the image point.
+
+.. ocv:function:: Point2f detail::RotationWarper::warpPoint(const Point2f &pt, const Mat &K, const Mat &R)
+
+    :param pt: Source point
+
+    :param K: Camera intrinsic parameters
+
+    :param R: Camera rotation matrix
+
+    :return: Projected point
+
+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)
+
+    :param src_size: Source image size
+
+    :param K: Camera intrinsic parameters
+
+    :param R: Camera rotation matrix
+
+    :param xmap: Projection map for the x axis
+
+    :param ymap: Projection map for the y axis
+
+    :return: Projected image minimum bounding box
+
+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)
+
+    :param src: Source image
+
+    :param K: Camera intrinsic parameters
+
+    :param R: Camera rotation matrix
+
+    :param interp_mode: Interpolation mode
+
+    :param border_mode: Border extrapolation mode
+
+    :param dst: Projected image
+
+    :return: Project image top-left corner
+
+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)
+
+    :param src: Projected image
+
+    :param K: Camera intrinsic parameters
+
+    :param R: Camera rotation matrix
+
+    :param interp_mode: Interpolation mode
+
+    :param border_mode: Border extrapolation mode
+
+    :param dst_size: Backward-projected image size
+
+    :param dst: Backward-projected image
+
+detail::RotationWarper::warpRoi
+-------------------------------
+
+.. ocv:function:: Rect detail::RotationWarper::warpRoi(Size src_size, const Mat &K, const Mat &R)
+
+    :param src_size: Source image bounding box
+
+    :param K: Camera intrinsic parameters
+
+    :param R: Camera rotation matrix
+
+    :return: Projected image minimum bounding box
+
+detail::ProjectorBase
+---------------------
+.. ocv:struct:: detail::ProjectorBase
+
+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));
+
+        float scale;
+        float k[9];
+        float rinv[9];
+        float r_kinv[9];
+        float k_rinv[9];
+        float t[3];
+    };
+
+detail::RotationWarperBase
+--------------------------
+.. ocv:class:: detail::RotationWarperBase
+
+Base class for rotation-based warper using a `detail::ProjectorBase`_ derived class. ::
+
+    template <class P>
+    class CV_EXPORTS RotationWarperBase : public RotationWarper
+    {
+    public:
+        Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R);
+
+        Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap);
+
+        Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode,
+                Mat &dst);
+
+        void warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode,
+                        Size dst_size, Mat &dst);
+
+        Rect warpRoi(Size src_size, const Mat &K, const Mat &R);
+
+    protected:
+
+        // Detects ROI of the destination image. It's correct for any projection.
+        virtual void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br);
+
+        // Detects ROI of the destination image by walking over image border.
+        // Correctness for any projection isn't guaranteed.
+        void detectResultRoiByBorder(Size src_size, Point &dst_tl, Point &dst_br);
+
+        P projector_;
+    };
+
+detail::PlaneWarper
+-------------------
+.. ocv:class:: detail::PlaneWarper : public detail::RotationWarperBase<PlaneProjector>
+
+Warper that maps an image onto the z = 1 plane. ::
+
+    class CV_EXPORTS PlaneWarper : public RotationWarperBase<PlaneProjector>
+    {
+    public:
+        PlaneWarper(float scale = 1.f) { projector_.scale = scale; }
+
+        void setScale(float scale) { projector_.scale = scale; }
+
+        Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R, const Mat &T);
+
+        Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap);
+
+        Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode,
+                   Mat &dst);
+
+        Rect warpRoi(Size src_size, const Mat &K, const Mat &R, const Mat &T);
+
+    protected:
+        void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br);
+    };
+
+.. seealso:: :ocv:class:`detail::RotationWarper`
+
+detail::PlaneWarper::PlaneWarper
+--------------------------------
+
+Construct an instance of the plane warper class.
+
+.. ocv:function:: void detail::PlaneWarper::PlaneWarper(float scale = 1.f)
+
+    :param scale: Projected image scale multiplier
+
+detail::SphericalWarper
+-----------------------
+.. ocv:class:: detail::SphericalWarper : public detail::RotationWarperBase<SphericalProjector>
+
+Warper that maps an image onto the unit sphere located at the origin. ::
+
+    class CV_EXPORTS SphericalWarper : public RotationWarperBase<SphericalProjector>
+    {
+    public:
+        SphericalWarper(float scale) { projector_.scale = scale; }
+
+    protected:
+        void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br);
+    };
+
+.. seealso:: :ocv:class:`detail::RotationWarper`
+
+detail::SphericalWarper::SphericalWarper
+----------------------------------------
+
+Construct an instance of the spherical warper class.
+
+.. ocv:function:: void detail::SphericalWarper::SphericalWarper(float scale)
+
+    :param scale: Projected image scale multiplier
+
+detail::CylindricalWarper
+-------------------------
+.. ocv:class:: detail::CylindricalWarper : public detail::RotationWarperBase<CylindricalProjector>
+
+Warper that maps an image onto the x*x + z*z = 1 cylinder. ::
+
+    class CV_EXPORTS CylindricalWarper : public RotationWarperBase<CylindricalProjector>
+    {
+    public:
+        CylindricalWarper(float scale) { projector_.scale = scale; }
+
+    protected:
+        void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br)
+        {
+            RotationWarperBase<CylindricalProjector>::detectResultRoiByBorder(src_size, dst_tl, dst_br);
+        }
+    };
+
+.. seealso:: :ocv:class:`detail::RotationWarper`
+
+detail::CylindricalWarper::CylindricalWarper
+--------------------------------------------
+
+Construct an instance of the cylindrical warper class.
+
+.. ocv:function:: void detail::CylindricalWarper::CylindricalWarper(float scale)
+
+    :param scale: Projected image scale multiplier

modules/stitching/include/opencv2/stitching/detail/warpers.hpp

@@ -134,9 +134,9 @@ public:
 
     Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R, const Mat &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, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap);
 
-    Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode,
+    virtual Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode,
                Mat &dst);
 
     Rect warpRoi(Size src_size, const Mat &K, const Mat &R, const Mat &T);
@@ -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, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap)
+    {
+        return buildMaps(src_size, K, R, Mat::zeros(3, 1, CV_32FC1), xmap, ymap);
+    }
+
+    virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Mat &dst)
+    {
+        return warp(src, K, R, Mat::zeros(3, 1, CV_32FC1), interp_mode, border_mode, dst);
+    }
+
+    virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap);
+    virtual Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, Mat &dst);
+};
+
+class CV_EXPORTS SphericalWarperOcl :  public SphericalWarper
+{
+public:
+    SphericalWarperOcl(float scale) : SphericalWarper(scale) { }
+
+    virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap);
+    virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Mat &dst);
+};
+
+class CV_EXPORTS CylindricalWarperOcl :  public CylindricalWarper
+{
+public:
+    CylindricalWarperOcl(float scale) : CylindricalWarper(scale) { }
+
+    virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap);
+    virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Mat &dst);
+};
+
 } // namespace detail
 } // namespace cv
 

modules/stitching/include/opencv2/stitching/warpers.hpp

@@ -167,6 +167,24 @@ public:
 };
 #endif
 
+class PlaneWarperOcl: public WarperCreator
+{
+public:
+    Ptr<detail::RotationWarper> create(float scale) const { return makePtr<detail::PlaneWarperOcl>(scale); }
+};
+
+class SphericalWarperOcl: public WarperCreator
+{
+public:
+    Ptr<detail::RotationWarper> create(float scale) const { return makePtr<detail::SphericalWarperOcl>(scale); }
+};
+
+class CylindricalWarperOcl: public WarperCreator
+{
+public:
+    Ptr<detail::RotationWarper> create(float scale) const { return makePtr<detail::CylindricalWarperOcl>(scale); }
+};
+
 } // namespace cv
 
 #endif // __OPENCV_STITCHING_WARPER_CREATORS_HPP__

modules/stitching/src/opencl/warpers.cl

+/*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;
+    }
+}

modules/stitching/src/precomp.hpp

@@ -53,6 +53,7 @@
 #include <sstream>
 #include <cmath>
 #include "opencv2/core.hpp"
+#include "opencv2/core/ocl.hpp"
 #include "opencv2/core/utility.hpp"
 #include "opencv2/stitching.hpp"
 #include "opencv2/stitching/detail/autocalib.hpp"

modules/stitching/src/warpers_ocl.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 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, const Mat & K, const Mat & R, const Mat & T, Mat & xmap, Mat & 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())
+        {
+            xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1);
+            ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1);
+
+            Mat r_kinv(1, 9, CV_32FC1, projector_.r_kinv), t(1, 3, CV_32FC1, projector_.t);
+            UMat uxmap = xmap.getUMat(ACCESS_WRITE), uymap = ymap.getUMat(ACCESS_WRITE),
+                    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] = { xmap.cols, xmap.rows };
+            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(const Mat & src, const Mat & K, const Mat & R, const Mat & T, int interp_mode, int border_mode, Mat & dst)
+{
+    Mat 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(ACCESS_WRITE);
+    remap(src, udst, uxmap, uymap, interp_mode, border_mode);
+
+    return dst_roi.tl();
+}
+
+/////////////////////////////////////////// SphericalWarperOcl ////////////////////////////////////////
+
+Rect SphericalWarperOcl::buildMaps(Size src_size, const Mat & K, const Mat &R, Mat & xmap, Mat & 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())
+        {
+            xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1);
+            ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1);
+
+            Mat r_kinv(1, 9, CV_32FC1, projector_.r_kinv);
+            UMat uxmap = xmap.getUMat(ACCESS_WRITE), uymap = ymap.getUMat(ACCESS_WRITE), 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] = { xmap.cols, xmap.rows };
+            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(const Mat & src, const Mat & K, const Mat & R, int interp_mode, int border_mode, Mat & dst)
+{
+    Mat 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(ACCESS_WRITE);
+    remap(src, udst, uxmap, uymap, interp_mode, border_mode);
+
+    return dst_roi.tl();
+}
+
+/////////////////////////////////////////// CylindricalWarperOcl ////////////////////////////////////////
+
+Rect CylindricalWarperOcl::buildMaps(Size src_size, const Mat & K, const Mat & R, Mat & xmap, Mat & 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())
+        {
+            xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1);
+            ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1);
+
+            Mat r_kinv(1, 9, CV_32FC1, projector_.r_kinv);
+            UMat uxmap = xmap.getUMat(ACCESS_WRITE), uymap = ymap.getUMat(ACCESS_WRITE), 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] = { xmap.cols, xmap.rows };
+            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(const Mat & src, const Mat & K, const Mat & R, int interp_mode, int border_mode, Mat & dst)
+{
+    Mat 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(ACCESS_WRITE);
+    remap(src, udst, uxmap, uymap, interp_mode, border_mode);
+
+    return dst_roi.tl();
+}
+
+} // namespace detail
+} // namespace cv

modules/stitching/test/ocl/test_warpers.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) 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<WarperCreator> creator = makePtr<SphericalWarperOcl>();
+        Ptr<detail::RotationWarper> 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_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst));
+
+        Near(1e-5);
+    }
+}
+
+//////////////////////////////// CylindricalWarperOcl /////////////////////////////////////////////////
+
+typedef WarperTestBase CylindricalWarperOclTest;
+
+OCL_TEST_F(CylindricalWarperOclTest, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        Ptr<WarperCreator> creator = makePtr<CylindricalWarperOcl>();
+        Ptr<detail::RotationWarper> 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_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst));
+
+        Near(1e-5);
+    }
+}
+
+//////////////////////////////// PlaneWarperOcl /////////////////////////////////////////////////
+
+typedef WarperTestBase PlaneWarperOclTest;
+
+OCL_TEST_F(PlaneWarperOclTest, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        Ptr<WarperCreator> creator = makePtr<PlaneWarperOcl>();
+        Ptr<detail::RotationWarper> 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_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst));
+
+        Near(1e-5);
+    }
+}
+
+} } // namespace cvtest::ocl
+
+#endif // HAVE_OPENCL

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 <float>\n"
@@ -123,7 +126,8 @@ static void printUsage()
 // Default command line args
 vector<String> 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<SurfFeaturesFinderGpu>();
         else
 #endif
@@ -430,7 +447,7 @@ int main(int argc, char* argv[])
     t = getTickCount();
 #endif
     vector<MatchesInfo> 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<WarperCreator> warper_creator;
+    if (try_ocl)
+    {
+        if (warp_type == "plane")
+            warper_creator = makePtr<cv::PlaneWarperOcl>();
+        else if (warp_type == "cylindrical")
+            warper_creator = makePtr<cv::CylindricalWarperOcl>();
+        else if (warp_type == "spherical")
+            warper_creator = makePtr<cv::SphericalWarperOcl>();
+    }
 #ifdef HAVE_OPENCV_CUDAWARPING
-    if (try_gpu && cuda::getCudaEnabledDeviceCount() > 0)
+    else if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0)
     {
         if (warp_type == "plane")
             warper_creator = makePtr<cv::PlaneWarperGpu>();
@@ -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<detail::GraphCutSeamFinderGpu>(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<detail::GraphCutSeamFinderGpu>(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<float>(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<MultiBandBlender*>(blender.get());