From 09ba133ddf42523a8cb6fd30a1290f589c0b8ad3 Mon Sep 17 00:00:00 2001
From: Vladislav Vinogradov <no@email>
Date: Mon, 31 Oct 2011 08:51:00 +0000
Subject: [PATCH] prepared GpuMat for moving to core module

---
 modules/gpu/include/opencv2/gpu/gpumat.hpp | 290 +++++--
 modules/gpu/src/gpumat.cpp                 | 950 ++++++++++-----------
 modules/stitching/src/blenders.cpp         |  10 +-
 samples/gpu/performance/tests.cpp          |  86 +-
 4 files changed, 739 insertions(+), 597 deletions(-)

diff --git a/modules/gpu/include/opencv2/gpu/gpumat.hpp b/modules/gpu/include/opencv2/gpu/gpumat.hpp
index c3dc39b0e6..780311cd4e 100644
--- a/modules/gpu/include/opencv2/gpu/gpumat.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpumat.hpp
@@ -48,21 +48,51 @@
 
 namespace cv { namespace gpu
 {
-    class Stream;
-    class CudaMem;
-
     //! Smart pointer for GPU memory with reference counting. Its interface is mostly similar with cv::Mat.
     class CV_EXPORTS GpuMat
     {
     public:
+        //! returns lightweight DevMem2D_ structure for passing to nvcc-compiled code.
+        // Contains just image size, data ptr and step.
+        template <class T> operator DevMem2D_<T>() const;
+        template <class T> operator PtrStep_<T>() const;
+		template <class T> operator PtrStep<T>() const;
+
+
+
+
+
+        //! builds GpuMat from Mat. Perfom blocking upload to device.
+        explicit GpuMat(const Mat& m);
+
+        //! pefroms blocking upload data to GpuMat.
+        void upload(const cv::Mat& m);
+
+        //! downloads data from device to host memory. Blocking calls.
+        void download(cv::Mat& m) const;
+        operator Mat() const
+        {
+            Mat m;
+            download(m);
+            return m;
+        }
+
+
+
+
+
+
         //! default constructor
         GpuMat();
+
         //! constructs GpuMatrix of the specified size and type (_type is CV_8UC1, CV_64FC3, CV_32SC(12) etc.)
         GpuMat(int rows, int cols, int type);
         GpuMat(Size size, int type);
+
         //! constucts GpuMatrix and fills it with the specified value _s.
         GpuMat(int rows, int cols, int type, const Scalar& s);
         GpuMat(Size size, int type, const Scalar& s);
+
         //! copy constructor
         GpuMat(const GpuMat& m);
 
@@ -74,35 +104,11 @@ namespace cv { namespace gpu
         GpuMat(const GpuMat& m, const Range& rowRange, const Range& colRange);
         GpuMat(const GpuMat& m, const Rect& roi);
 
-        //! builds GpuMat from Mat. Perfom blocking upload to device.
-        explicit GpuMat (const Mat& m);
-
         //! destructor - calls release()
         ~GpuMat();
 
         //! assignment operators
         GpuMat& operator = (const GpuMat& m);
-        //! assignment operator. Perfom blocking upload to device.
-        GpuMat& operator = (const Mat& m);
-
-        //! returns lightweight DevMem2D_ structure for passing to nvcc-compiled code.
-        // Contains just image size, data ptr and step.
-        template <class T> operator DevMem2D_<T>() const;
-        template <class T> operator PtrStep_<T>() const;
-		template <class T> operator PtrStep<T>() const;
-
-        //! pefroms blocking upload data to GpuMat.
-        void upload(const cv::Mat& m);
-
-        //! upload async
-        void upload(const CudaMem& m, Stream& stream);
-
-        //! downloads data from device to host memory. Blocking calls.
-        operator Mat() const;
-        void download(cv::Mat& m) const;
-
-        //! download async
-        void download(CudaMem& m, Stream& stream) const;
 
         //! returns a new GpuMatrix header for the specified row
         GpuMat row(int y) const;
@@ -119,13 +125,13 @@ namespace cv { namespace gpu
         GpuMat clone() const;
         //! copies the GpuMatrix content to "m".
         // It calls m.create(this->size(), this->type()).
-        void copyTo( GpuMat& m ) const;
+        void copyTo(GpuMat& m) const;
         //! copies those GpuMatrix elements to "m" that are marked with non-zero mask elements.
-        void copyTo( GpuMat& m, const GpuMat& mask ) const;
+        void copyTo(GpuMat& m, const GpuMat& mask) const;
         //! converts GpuMatrix to another datatype with optional scalng. See cvConvertScale.
-        void convertTo( GpuMat& m, int rtype, double alpha=1, double beta=0 ) const;
+        void convertTo(GpuMat& m, int rtype, double alpha = 1, double beta = 0) const;
 
-        void assignTo( GpuMat& m, int type=-1 ) const;
+        void assignTo(GpuMat& m, int type=-1) const;
 
         //! sets every GpuMatrix element to s
         GpuMat& operator = (const Scalar& s);
@@ -147,13 +153,13 @@ namespace cv { namespace gpu
         void swap(GpuMat& mat);
 
         //! locates GpuMatrix header within a parent GpuMatrix. See below
-        void locateROI( Size& wholeSize, Point& ofs ) const;
+        void locateROI(Size& wholeSize, Point& ofs) const;
         //! moves/resizes the current GpuMatrix ROI inside the parent GpuMatrix.
-        GpuMat& adjustROI( int dtop, int dbottom, int dleft, int dright );
+        GpuMat& adjustROI(int dtop, int dbottom, int dleft, int dright);
         //! extracts a rectangular sub-GpuMatrix
         // (this is a generalized form of row, rowRange etc.)
-        GpuMat operator()( Range rowRange, Range colRange ) const;
-        GpuMat operator()( const Rect& roi ) const;
+        GpuMat operator()(Range rowRange, Range colRange) const;
+        GpuMat operator()(const Rect& roi) const;
 
         //! returns true iff the GpuMatrix data is continuous
         // (i.e. when there are no gaps between successive rows).
@@ -186,9 +192,6 @@ namespace cv { namespace gpu
         template<typename _Tp> _Tp* ptr(int y = 0);
         template<typename _Tp> const _Tp* ptr(int y = 0) const;
 
-        //! matrix transposition
-        GpuMat t() const;
-
         /*! includes several bit-fields:
         - the magic signature
         - continuity flag
@@ -196,10 +199,13 @@ namespace cv { namespace gpu
         - number of channels
         */
         int flags;
+
         //! the number of rows and columns
         int rows, cols;
+
         //! a distance between successive rows in bytes; includes the gap if any
         size_t step;
+
         //! pointer to the data
         uchar* data;
 
@@ -223,23 +229,60 @@ namespace cv { namespace gpu
     CV_EXPORTS void ensureSizeIsEnough(int rows, int cols, int type, GpuMat& m);
     CV_EXPORTS void ensureSizeIsEnough(Size size, int type, GpuMat& m);
 
-    ////////////////////////////////////////////////////////////////////////
-    //////////////////////////////// GpuMat ////////////////////////////////
     ////////////////////////////////////////////////////////////////////////
 
-    inline GpuMat::GpuMat() : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0) {}
+    template <class T> inline GpuMat::operator DevMem2D_<T>() const { return DevMem2D_<T>(rows, cols, (T*)data, step); }
+    template <class T> inline GpuMat::operator PtrStep_<T>() const { return PtrStep_<T>(static_cast< DevMem2D_<T> >(*this)); }
+	template <class T> inline GpuMat::operator PtrStep<T>() const { return PtrStep<T>((T*)data, step); }	
 
-    inline GpuMat::GpuMat(int rows_, int cols_, int type_) : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
+
+
+
+
+
+    inline GpuMat::GpuMat() 
+        : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0) 
+    {
+    }
+
+    inline GpuMat::GpuMat(int rows_, int cols_, int type_) 
+        : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
     {
         if (rows_ > 0 && cols_ > 0)
             create(rows_, cols_, type_);
     }
 
-    inline GpuMat::~GpuMat() { release(); }
+    inline GpuMat::GpuMat(Size size_, int type_) 
+        : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
+    {
+        if (size_.height > 0 && size_.width > 0)
+            create(size_.height, size_.width, type_);
+    }
 
-    template <class T> inline GpuMat::operator DevMem2D_<T>() const { return DevMem2D_<T>(rows, cols, (T*)data, step); }
-    template <class T> inline GpuMat::operator PtrStep_<T>() const { return PtrStep_<T>(static_cast< DevMem2D_<T> >(*this)); }
-	template <class T> inline GpuMat::operator PtrStep<T>() const { return PtrStep<T>((T*)data, step); }	
+    inline GpuMat::GpuMat(int rows_, int cols_, int type_, const Scalar& s_) 
+        : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
+    {
+        if (rows_ > 0 && cols_ > 0)
+        {
+            create(rows_, cols_, type_);
+            setTo(s_);
+        }
+    }
+
+    inline GpuMat::GpuMat(Size size_, int type_, const Scalar& s_) 
+        : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
+    {
+        if (size_.height > 0 && size_.width > 0)
+        {
+            create(size_.height, size_.width, type_);
+            setTo(s_);
+        }
+    }
+
+    inline GpuMat::~GpuMat() 
+    { 
+        release(); 
+    }
 
     inline GpuMat GpuMat::clone() const
     {
@@ -250,15 +293,21 @@ namespace cv { namespace gpu
 
     inline void GpuMat::assignTo(GpuMat& m, int type) const
     {
-        if( type < 0 )
+        if (type < 0)
             m = *this;
         else
             convertTo(m, type);
     }
 
-    inline size_t GpuMat::step1() const { return step/elemSize1(); }
+    inline size_t GpuMat::step1() const 
+    { 
+        return step / elemSize1(); 
+    }
 
-    inline bool GpuMat::empty() const { return data == 0; }
+    inline bool GpuMat::empty() const 
+    { 
+        return data == 0; 
+    }
 
     template<typename _Tp> inline _Tp* GpuMat::ptr(int y)
     {
@@ -270,7 +319,148 @@ namespace cv { namespace gpu
         return (const _Tp*)ptr(y);
     }
 
-    inline void swap(GpuMat& a, GpuMat& b) { a.swap(b); }
+    inline void swap(GpuMat& a, GpuMat& b) 
+    { 
+        a.swap(b); 
+    }
+
+    inline GpuMat GpuMat::row(int y) const 
+    { 
+        return GpuMat(*this, Range(y, y+1), Range::all()); 
+    }
+
+    inline GpuMat GpuMat::col(int x) const 
+    { 
+        return GpuMat(*this, Range::all(), Range(x, x+1)); 
+    }
+
+    inline GpuMat GpuMat::rowRange(int startrow, int endrow) const 
+    { 
+        return GpuMat(*this, Range(startrow, endrow), Range::all()); 
+    }
+
+    inline GpuMat GpuMat::rowRange(const Range& r) const 
+    { 
+        return GpuMat(*this, r, Range::all()); 
+    }
+
+    inline GpuMat GpuMat::colRange(int startcol, int endcol) const 
+    { 
+        return GpuMat(*this, Range::all(), Range(startcol, endcol)); 
+    }
+
+    inline GpuMat GpuMat::colRange(const Range& r) const 
+    { 
+        return GpuMat(*this, Range::all(), r); 
+    }
+
+    inline void GpuMat::create(Size size_, int type_) 
+    { 
+        create(size_.height, size_.width, type_); 
+    }
+
+    inline GpuMat GpuMat::operator()(Range rowRange, Range colRange) const 
+    { 
+        return GpuMat(*this, rowRange, colRange); 
+    }
+
+    inline GpuMat GpuMat::operator()(const Rect& roi) const 
+    { 
+        return GpuMat(*this, roi); 
+    }
+
+    inline bool GpuMat::isContinuous() const 
+    { 
+        return (flags & Mat::CONTINUOUS_FLAG) != 0; 
+    }
+
+    inline size_t GpuMat::elemSize() const 
+    { 
+        return CV_ELEM_SIZE(flags); 
+    }
+
+    inline size_t GpuMat::elemSize1() const 
+    { 
+        return CV_ELEM_SIZE1(flags); 
+    }
+
+    inline int GpuMat::type() const 
+    { 
+        return CV_MAT_TYPE(flags); 
+    }
+
+    inline int GpuMat::depth() const 
+    { 
+        return CV_MAT_DEPTH(flags); 
+    }
+
+    inline int GpuMat::channels() const 
+    { 
+        return CV_MAT_CN(flags); 
+    }
+
+    inline Size GpuMat::size() const 
+    { 
+        return Size(cols, rows); 
+    }
+
+    inline unsigned char* GpuMat::ptr(int y)
+    {
+        CV_DbgAssert((unsigned)y < (unsigned)rows);
+        return data + step * y;
+    }
+
+    inline const unsigned char* GpuMat::ptr(int y) const
+    {
+        CV_DbgAssert((unsigned)y < (unsigned)rows);
+        return data + step * y;
+    }
+
+    inline GpuMat& GpuMat::operator = (const Scalar& s)
+    {
+        setTo(s);
+        return *this;
+    }
+
+    inline GpuMat createContinuous(int rows, int cols, int type)
+    {
+        GpuMat m;
+        createContinuous(rows, cols, type, m);
+        return m;
+    }
+
+    inline void createContinuous(Size size, int type, GpuMat& m)
+    {
+        createContinuous(size.height, size.width, type, m);
+    }
+
+    inline GpuMat createContinuous(Size size, int type)
+    {
+        GpuMat m;
+        createContinuous(size, type, m);
+        return m;
+    }
+
+    inline void ensureSizeIsEnough(Size size, int type, GpuMat& m)
+    {
+        ensureSizeIsEnough(size.height, size.width, type, m);
+    }
+
+    inline void createContinuous(int rows, int cols, int type, GpuMat& m)
+    {
+        int area = rows * cols;
+        if (!m.isContinuous() || m.type() != type || m.size().area() != area)
+            m.create(1, area, type);
+        m = m.reshape(0, rows);
+    }
+
+    inline void ensureSizeIsEnough(int rows, int cols, int type, GpuMat& m)
+    {
+        if (m.type() == type && m.rows >= rows && m.cols >= cols)
+            m = m(Rect(0, 0, cols, rows));
+        else
+            m.create(rows, cols, type);
+    }
 }}
 
 #endif // __OPENCV_GPUMAT_HPP__
diff --git a/modules/gpu/src/gpumat.cpp b/modules/gpu/src/gpumat.cpp
index 8d36057f3a..b1818a56c1 100644
--- a/modules/gpu/src/gpumat.cpp
+++ b/modules/gpu/src/gpumat.cpp
@@ -44,37 +44,7 @@
 
 using namespace cv;
 using namespace cv::gpu;
-
-////////////////////////////////////////////////////////////////////////
-//////////////////////////////// GpuMat ////////////////////////////////
-////////////////////////////////////////////////////////////////////////
-
-cv::gpu::GpuMat::GpuMat(Size size_, int type_) : 
-    flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
-{
-    if (size_.height > 0 && size_.width > 0)
-        create(size_.height, size_.width, type_);
-}
-
-cv::gpu::GpuMat::GpuMat(int rows_, int cols_, int type_, const Scalar& s_) : 
-    flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
-{
-    if (rows_ > 0 && cols_ > 0)
-    {
-        create(rows_, cols_, type_);
-        *this = s_;
-    }
-}
-
-cv::gpu::GpuMat::GpuMat(Size size_, int type_, const Scalar& s_) : 
-    flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
-{
-    if (size_.height > 0 && size_.width > 0)
-    {
-        create(size_.height, size_.width, type_);
-        *this = s_;
-    }
-}
+using namespace std;
 
 cv::gpu::GpuMat::GpuMat(const GpuMat& m) : 
     flags(m.flags), rows(m.rows), cols(m.cols), step(m.step), data(m.data), refcount(m.refcount), datastart(m.datastart), dataend(m.dataend)
@@ -84,10 +54,12 @@ cv::gpu::GpuMat::GpuMat(const GpuMat& m) :
 }
 
 cv::gpu::GpuMat::GpuMat(int rows_, int cols_, int type_, void* data_, size_t step_) : 
-    flags(Mat::MAGIC_VAL + (type_ & TYPE_MASK)), rows(rows_), cols(cols_), step(step_), data((uchar*)data_), refcount(0),
+    flags(Mat::MAGIC_VAL + (type_ & TYPE_MASK)), rows(rows_), cols(cols_), 
+    step(step_), data((uchar*)data_), refcount(0),
     datastart((uchar*)data_), dataend((uchar*)data_)
 {
     size_t minstep = cols * elemSize();
+
     if (step == Mat::AUTO_STEP)
     {
         step = minstep;
@@ -95,8 +67,11 @@ cv::gpu::GpuMat::GpuMat(int rows_, int cols_, int type_, void* data_, size_t ste
     }
     else
     {
-        if (rows == 1) step = minstep;
-        CV_DbgAssert( step >= minstep );
+        if (rows == 1) 
+            step = minstep;
+
+        CV_DbgAssert(step >= minstep);
+
         flags |= step == minstep ? Mat::CONTINUOUS_FLAG : 0;
     }
     dataend += step * (rows - 1) + minstep;
@@ -108,6 +83,7 @@ cv::gpu::GpuMat::GpuMat(Size size_, int type_, void* data_, size_t step_) :
     datastart((uchar*)data_), dataend((uchar*)data_)
 {
     size_t minstep = cols * elemSize();
+
     if (step == Mat::AUTO_STEP)
     {
         step = minstep;
@@ -115,8 +91,11 @@ cv::gpu::GpuMat::GpuMat(Size size_, int type_, void* data_, size_t step_) :
     }
     else
     {
-        if (rows == 1) step = minstep;
-        CV_DbgAssert( step >= minstep );
+        if (rows == 1) 
+            step = minstep;
+
+        CV_DbgAssert(step >= minstep);
+
         flags |= step == minstep ? Mat::CONTINUOUS_FLAG : 0;
     }
     dataend += step * (rows - 1) + minstep;
@@ -132,7 +111,8 @@ cv::gpu::GpuMat::GpuMat(const GpuMat& m, const Range& rowRange, const Range& col
         rows = m.rows;
     else
     {
-        CV_Assert( 0 <= rowRange.start && rowRange.start <= rowRange.end && rowRange.end <= m.rows );
+        CV_Assert(0 <= rowRange.start && rowRange.start <= rowRange.end && rowRange.end <= m.rows);
+
         rows = rowRange.size();
         data += step*rowRange.start;
     }
@@ -141,18 +121,20 @@ cv::gpu::GpuMat::GpuMat(const GpuMat& m, const Range& rowRange, const Range& col
         cols = m.cols;
     else
     {
-        CV_Assert( 0 <= colRange.start && colRange.start <= colRange.end && colRange.end <= m.cols );
+        CV_Assert(0 <= colRange.start && colRange.start <= colRange.end && colRange.end <= m.cols);
+
         cols = colRange.size();
         data += colRange.start*elemSize();
         flags &= cols < m.cols ? ~Mat::CONTINUOUS_FLAG : -1;
     }
 
-    if( rows == 1 )
+    if (rows == 1)
         flags |= Mat::CONTINUOUS_FLAG;
 
-    if( refcount )
+    if (refcount)
         CV_XADD(refcount, 1);
-    if( rows <= 0 || cols <= 0 )
+
+    if (rows <= 0 || cols <= 0)
         rows = cols = 0;
 }
 
@@ -162,12 +144,14 @@ cv::gpu::GpuMat::GpuMat(const GpuMat& m, const Rect& roi) :
     datastart(m.datastart), dataend(m.dataend)
 {
     flags &= roi.width < m.cols ? ~Mat::CONTINUOUS_FLAG : -1;
-    data += roi.x*elemSize();
-    CV_Assert( 0 <= roi.x && 0 <= roi.width && roi.x + roi.width <= m.cols &&
-        0 <= roi.y && 0 <= roi.height && roi.y + roi.height <= m.rows );
-    if( refcount )
+    data += roi.x * elemSize();
+
+    CV_Assert(0 <= roi.x && 0 <= roi.width && roi.x + roi.width <= m.cols && 0 <= roi.y && 0 <= roi.height && roi.y + roi.height <= m.rows);
+
+    if (refcount)
         CV_XADD(refcount, 1);
-    if( rows <= 0 || cols <= 0 )
+
+    if (rows <= 0 || cols <= 0)
         rows = cols = 0;
 }
 
@@ -179,223 +163,204 @@ cv::gpu::GpuMat::GpuMat(const Mat& m) :
 
 GpuMat& cv::gpu::GpuMat::operator = (const GpuMat& m)
 {
-    if( this != &m )
+    if (this != &m)
     {
-        if( m.refcount )
-            CV_XADD(m.refcount, 1);
-        release();
-        flags = m.flags;
-        rows = m.rows; cols = m.cols;
-        step = m.step; data = m.data;
-        datastart = m.datastart; dataend = m.dataend;
-        refcount = m.refcount;
+        GpuMat temp(m);
+        swap(temp);
     }
+
     return *this;
 }
 
-GpuMat& cv::gpu::GpuMat::operator = (const Mat& m) 
-{ 
-    upload(m); return *this; 
-}
-
-cv::gpu::GpuMat::operator Mat() const
-{
-    Mat m;
-    download(m);
-    return m;
-}
-
-GpuMat cv::gpu::GpuMat::row(int y) const 
-{ 
-    return GpuMat(*this, Range(y, y+1), Range::all()); 
-}
-
-GpuMat cv::gpu::GpuMat::col(int x) const 
-{ 
-    return GpuMat(*this, Range::all(), Range(x, x+1)); 
-}
-
-GpuMat cv::gpu::GpuMat::rowRange(int startrow, int endrow) const 
-{ 
-    return GpuMat(*this, Range(startrow, endrow), Range::all()); 
-}
-
-GpuMat cv::gpu::GpuMat::rowRange(const Range& r) const 
-{ 
-    return GpuMat(*this, r, Range::all()); 
-}
-
-GpuMat cv::gpu::GpuMat::colRange(int startcol, int endcol) const 
-{ 
-    return GpuMat(*this, Range::all(), Range(startcol, endcol)); 
-}
-
-GpuMat cv::gpu::GpuMat::colRange(const Range& r) const 
-{ 
-    return GpuMat(*this, Range::all(), r); 
-}
-
-void cv::gpu::GpuMat::create(Size size_, int type_) 
-{ 
-    create(size_.height, size_.width, type_); 
-}
-
 void cv::gpu::GpuMat::swap(GpuMat& b)
 {
-    std::swap( flags, b.flags );
-    std::swap( rows, b.rows ); 
-    std::swap( cols, b.cols );
-    std::swap( step, b.step ); 
-    std::swap( data, b.data );
-    std::swap( datastart, b.datastart );
-    std::swap( dataend, b.dataend );
-    std::swap( refcount, b.refcount );
+    std::swap(flags, b.flags);
+    std::swap(rows, b.rows); 
+    std::swap(cols, b.cols);
+    std::swap(step, b.step); 
+    std::swap(data, b.data);
+    std::swap(datastart, b.datastart);
+    std::swap(dataend, b.dataend);
+    std::swap(refcount, b.refcount);
 }
 
 void cv::gpu::GpuMat::locateROI(Size& wholeSize, Point& ofs) const
 {
-    size_t esz = elemSize(), minstep;
-    ptrdiff_t delta1 = data - datastart, delta2 = dataend - datastart;
-    CV_DbgAssert( step > 0 );
-    if( delta1 == 0 )
+    size_t esz = elemSize();
+    ptrdiff_t delta1 = data - datastart;
+    ptrdiff_t delta2 = dataend - datastart;
+
+    CV_DbgAssert(step > 0);
+
+    if (delta1 == 0)
         ofs.x = ofs.y = 0;
     else
     {
-        ofs.y = (int)(delta1/step);
-        ofs.x = (int)((delta1 - step*ofs.y)/esz);
-        CV_DbgAssert( data == datastart + ofs.y*step + ofs.x*esz );
+        ofs.y = static_cast<int>(delta1 / step);
+        ofs.x = static_cast<int>((delta1 - step * ofs.y) / esz);
+
+        CV_DbgAssert(data == datastart + ofs.y * step + ofs.x * esz);
     }
-    minstep = (ofs.x + cols)*esz;
-    wholeSize.height = (int)((delta2 - minstep)/step + 1);
-    wholeSize.height = std::max(wholeSize.height, ofs.y + rows);
-    wholeSize.width = (int)((delta2 - step*(wholeSize.height-1))/esz);
-    wholeSize.width = std::max(wholeSize.width, ofs.x + cols);
+
+    size_t minstep = (ofs.x + cols) * esz;
+
+    wholeSize.height = std::max(static_cast<int>((delta2 - minstep) / step + 1), ofs.y + rows);
+    wholeSize.width = std::max(static_cast<int>((delta2 - step * (wholeSize.height - 1)) / esz), ofs.x + cols);
 }
 
 GpuMat& cv::gpu::GpuMat::adjustROI(int dtop, int dbottom, int dleft, int dright)
 {
-    Size wholeSize; Point ofs;
+    Size wholeSize; 
+    Point ofs;
+    locateROI(wholeSize, ofs);
+
     size_t esz = elemSize();
-    locateROI( wholeSize, ofs );
-    int row1 = std::max(ofs.y - dtop, 0), row2 = std::min(ofs.y + rows + dbottom, wholeSize.height);
-    int col1 = std::max(ofs.x - dleft, 0), col2 = std::min(ofs.x + cols + dright, wholeSize.width);
-    data += (row1 - ofs.y)*step + (col1 - ofs.x)*esz;
-    rows = row2 - row1; cols = col2 - col1;
-    if( esz*cols == step || rows == 1 )
+
+    int row1 = std::max(ofs.y - dtop, 0); 
+    int row2 = std::min(ofs.y + rows + dbottom, wholeSize.height);
+
+    int col1 = std::max(ofs.x - dleft, 0);
+    int col2 = std::min(ofs.x + cols + dright, wholeSize.width);
+
+    data += (row1 - ofs.y) * step + (col1 - ofs.x) * esz;
+    rows = row2 - row1; 
+    cols = col2 - col1;
+
+    if (esz * cols == step || rows == 1)
         flags |= Mat::CONTINUOUS_FLAG;
     else
         flags &= ~Mat::CONTINUOUS_FLAG;
+
     return *this;
 }
 
-cv::gpu::GpuMat GpuMat::operator()(Range rowRange, Range colRange) const 
-{ 
-    return GpuMat(*this, rowRange, colRange); 
-}
-
-cv::gpu::GpuMat GpuMat::operator()(const Rect& roi) const 
-{ 
-    return GpuMat(*this, roi); 
-}
-
-bool cv::gpu::GpuMat::isContinuous() const 
-{ 
-    return (flags & Mat::CONTINUOUS_FLAG) != 0; 
-}
-
-size_t cv::gpu::GpuMat::elemSize() const 
-{ 
-    return CV_ELEM_SIZE(flags); 
-}
-
-size_t cv::gpu::GpuMat::elemSize1() const 
-{ 
-    return CV_ELEM_SIZE1(flags); 
-}
-
-int cv::gpu::GpuMat::type() const 
-{ 
-    return CV_MAT_TYPE(flags); 
-}
-
-int cv::gpu::GpuMat::depth() const 
-{ 
-    return CV_MAT_DEPTH(flags); 
-}
-
-int cv::gpu::GpuMat::channels() const 
-{ 
-    return CV_MAT_CN(flags); 
-}
-
-Size cv::gpu::GpuMat::size() const 
-{ 
-    return Size(cols, rows); 
-}
-
-unsigned char* cv::gpu::GpuMat::ptr(int y)
+GpuMat cv::gpu::GpuMat::reshape(int new_cn, int new_rows) const
 {
-    CV_DbgAssert( (unsigned)y < (unsigned)rows );
-    return data + step*y;
+    GpuMat hdr = *this;
+
+    int cn = channels();
+    if (new_cn == 0)
+        new_cn = cn;
+
+    int total_width = cols * cn;
+
+    if ((new_cn > total_width || total_width % new_cn != 0) && new_rows == 0)
+        new_rows = rows * total_width / new_cn;
+
+    if (new_rows != 0 && new_rows != rows)
+    {
+        int total_size = total_width * rows;
+
+        if (!isContinuous())
+            CV_Error(CV_BadStep, "The matrix is not continuous, thus its number of rows can not be changed");
+
+        if ((unsigned)new_rows > (unsigned)total_size)
+            CV_Error(CV_StsOutOfRange, "Bad new number of rows");
+
+        total_width = total_size / new_rows;
+
+        if (total_width * new_rows != total_size)
+            CV_Error(CV_StsBadArg, "The total number of matrix elements is not divisible by the new number of rows");
+
+        hdr.rows = new_rows;
+        hdr.step = total_width * elemSize1();
+    }
+
+    int new_width = total_width / new_cn;
+
+    if (new_width * new_cn != total_width)
+        CV_Error(CV_BadNumChannels, "The total width is not divisible by the new number of channels");
+
+    hdr.cols = new_width;
+    hdr.flags = (hdr.flags & ~CV_MAT_CN_MASK) | ((new_cn - 1) << CV_CN_SHIFT);
+
+    return hdr;
 }
 
-const unsigned char* cv::gpu::GpuMat::ptr(int y) const
-{
-    CV_DbgAssert( (unsigned)y < (unsigned)rows );
-    return data + step*y;
-}
 
-GpuMat cv::gpu::GpuMat::t() const
-{
-    GpuMat tmp;
-    transpose(*this, tmp);
-    return tmp;
-}
 
-GpuMat cv::gpu::createContinuous(int rows, int cols, int type)
-{
-    GpuMat m;
-    createContinuous(rows, cols, type, m);
-    return m;
-}
 
-void cv::gpu::createContinuous(Size size, int type, GpuMat& m)
-{
-    createContinuous(size.height, size.width, type, m);
-}
 
-GpuMat cv::gpu::createContinuous(Size size, int type)
-{
-    GpuMat m;
-    createContinuous(size, type, m);
-    return m;
-}
 
-void cv::gpu::ensureSizeIsEnough(Size size, int type, GpuMat& m)
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+class GpuFuncTable
 {
-    ensureSizeIsEnough(size.height, size.width, type, m);
-}
+public:
+    virtual ~GpuFuncTable() {}
+
+    virtual void copy(const Mat& src, GpuMat& dst) const = 0;
+    virtual void copy(const GpuMat& src, Mat& dst) const = 0;
+    virtual void copy(const GpuMat& src, GpuMat& dst) const = 0;
+
+    virtual void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask) const = 0;
+
+    virtual void convert(const GpuMat& src, GpuMat& dst) const = 0;
+    virtual void convert(const GpuMat& src, GpuMat& dst, double alpha, double beta) const = 0;
+
+    virtual void setTo(GpuMat& m, const Scalar& s, const GpuMat& mask) const = 0;
+
+    virtual void mallocPitch(void** devPtr, size_t* step, size_t width, size_t height) const = 0;
+    virtual void free(void* devPtr) const = 0;
+};
+
 
 #if !defined (HAVE_CUDA)
 
-void cv::gpu::GpuMat::upload(const Mat&) { throw_nogpu(); }
-void cv::gpu::GpuMat::download(cv::Mat&) const { throw_nogpu(); }
-void cv::gpu::GpuMat::copyTo(GpuMat&) const { throw_nogpu(); }
-void cv::gpu::GpuMat::copyTo(GpuMat&, const GpuMat&) const { throw_nogpu(); }
-void cv::gpu::GpuMat::convertTo(GpuMat&, int, double, double) const { throw_nogpu(); }
-GpuMat& cv::gpu::GpuMat::operator = (const Scalar&) { throw_nogpu(); return *this; }
-GpuMat& cv::gpu::GpuMat::setTo(const Scalar&, const GpuMat&) { throw_nogpu(); return *this; }
-GpuMat cv::gpu::GpuMat::reshape(int, int) const { throw_nogpu(); return GpuMat(); }
-void cv::gpu::GpuMat::create(int, int, int) { throw_nogpu(); }
-void cv::gpu::GpuMat::release() {}
-void cv::gpu::createContinuous(int, int, int, GpuMat&) { throw_nogpu(); }
-void cv::gpu::ensureSizeIsEnough(int, int, int, GpuMat&) { throw_nogpu(); }
+class EmptyFuncTable : public GpuFuncTable
+{
+public:
+    void copy(const Mat&, GpuMat&) const { throw_nogpu(); }
+    void copy(const GpuMat&, Mat&) const { throw_nogpu(); }
+    void copy(const GpuMat&, GpuMat&) const { throw_nogpu(); }
 
-#else /* !defined (HAVE_CUDA) */
+    void copyWithMask(const GpuMat&, GpuMat&, const GpuMat&) const { throw_nogpu(); }
+
+    void convert(const GpuMat&, GpuMat&) const { throw_nogpu(); }
+    void convert(const GpuMat&, GpuMat&, double, double) const { throw_nogpu(); }
+
+    void setTo(GpuMat&, const Scalar&, const GpuMat&) const { throw_nogpu(); }
+
+    void mallocPitch(void**, size_t*, size_t, size_t) const { throw_nogpu(); }
+    void free(void*) const {}
+};
+
+const GpuFuncTable* gpuFuncTable()
+{
+    static EmptyFuncTable empty;
+    return &empty;
+}
+
+#else
 
 namespace cv { namespace gpu { namespace device
 {
-    void copy_to_with_mask(const DevMem2Db& src, DevMem2Db dst, int depth, const DevMem2Db& mask, int channels, const cudaStream_t & stream = 0);
+    void copy_to_with_mask(const DevMem2Db& src, DevMem2Db dst, int depth, const DevMem2Db& mask, int channels, const cudaStream_t& stream = 0);
 
     template <typename T>
     void set_to_gpu(const DevMem2Db& mat, const T* scalar, int channels, cudaStream_t stream);
@@ -405,56 +370,11 @@ namespace cv { namespace gpu { namespace device
     void convert_gpu(const DevMem2Db& src, int sdepth, const DevMem2Db& dst, int ddepth, double alpha, double beta, cudaStream_t stream = 0);
 }}}
 
-
-void cv::gpu::GpuMat::upload(const Mat& m)
-{
-    CV_DbgAssert(!m.empty());
-    create(m.size(), m.type());
-    cudaSafeCall( cudaMemcpy2D(data, step, m.data, m.step, cols * elemSize(), rows, cudaMemcpyHostToDevice) );
-}
-
-void cv::gpu::GpuMat::upload(const CudaMem& m, Stream& stream)
-{
-    CV_DbgAssert(!m.empty());
-    stream.enqueueUpload(m, *this);
-}
-
-void cv::gpu::GpuMat::download(cv::Mat& m) const
-{
-    CV_DbgAssert(!this->empty());
-    m.create(size(), type());
-    cudaSafeCall( cudaMemcpy2D(m.data, m.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToHost) );
-}
-
-void cv::gpu::GpuMat::download(CudaMem& m, Stream& stream) const
-{
-    CV_DbgAssert(!m.empty());
-    stream.enqueueDownload(*this, m);
-}
-
-void cv::gpu::GpuMat::copyTo(GpuMat& m) const
-{
-    CV_DbgAssert(!this->empty());
-    m.create(size(), type());
-    cudaSafeCall( cudaMemcpy2D(m.data, m.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToDevice) );
-    cudaSafeCall( cudaDeviceSynchronize() );
-}
-
-void cv::gpu::GpuMat::copyTo(GpuMat& mat, const GpuMat& mask) const
-{
-    if (mask.empty())
-    {
-        copyTo(mat);
-    }
-    else
-    {
-        mat.create(size(), type());
-        device::copy_to_with_mask(*this, mat, depth(), mask, channels());
-    }
-}
-
 namespace
 {
+    //////////////////////////////////////////////////////////////////////////
+    // Convert
+
     template<int n> struct NPPTypeTraits;
     template<> struct NPPTypeTraits<CV_8U>  { typedef Npp8u npp_type; };
     template<> struct NPPTypeTraits<CV_16U> { typedef Npp16u npp_type; };
@@ -510,129 +430,10 @@ namespace
     {
         device::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), 1.0, 0.0);
     }
-}
 
-void cv::gpu::GpuMat::convertTo( GpuMat& dst, int rtype, double alpha, double beta ) const
-{
-    CV_Assert((depth() != CV_64F && CV_MAT_DEPTH(rtype) != CV_64F) || 
-        (TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE)));
-
-    bool noScale = fabs(alpha-1) < std::numeric_limits<double>::epsilon() && fabs(beta) < std::numeric_limits<double>::epsilon();
-
-    if( rtype < 0 )
-        rtype = type();
-    else
-        rtype = CV_MAKETYPE(CV_MAT_DEPTH(rtype), channels());
-
-    int scn = channels();
-    int sdepth = depth(), ddepth = CV_MAT_DEPTH(rtype);
-    if( sdepth == ddepth && noScale )
-    {
-        copyTo(dst);
-        return;
-    }
-
-    GpuMat temp;
-    const GpuMat* psrc = this;
-    if( sdepth != ddepth && psrc == &dst )
-        psrc = &(temp = *this);
-
-    dst.create( size(), rtype );
-
-    if (!noScale)
-        device::convert_gpu(psrc->reshape(1), sdepth, dst.reshape(1), ddepth, alpha, beta);
-    else
-    {
-        typedef void (*convert_caller_t)(const GpuMat& src, GpuMat& dst);
-        static const convert_caller_t convert_callers[8][8][4] =
-        {
-            {
-                {0,0,0,0},
-                {convertToKernelCaller, convertToKernelCaller, convertToKernelCaller, convertToKernelCaller},
-                {NppCvt<CV_8U, CV_16U, nppiConvert_8u16u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,NppCvt<CV_8U, CV_16U, nppiConvert_8u16u_C4R>::cvt},
-                {NppCvt<CV_8U, CV_16S, nppiConvert_8u16s_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,NppCvt<CV_8U, CV_16S, nppiConvert_8u16s_C4R>::cvt},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {NppCvt<CV_8U, CV_32F, nppiConvert_8u32f_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0}
-            },
-            {
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0}
-            },
-            {
-                {NppCvt<CV_16U, CV_8U, nppiConvert_16u8u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,NppCvt<CV_16U, CV_8U, nppiConvert_16u8u_C4R>::cvt},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {NppCvt<CV_16U, CV_32S, nppiConvert_16u32s_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {NppCvt<CV_16U, CV_32F, nppiConvert_16u32f_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0}
-            },
-            {
-                {NppCvt<CV_16S, CV_8U, nppiConvert_16s8u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,NppCvt<CV_16S, CV_8U, nppiConvert_16s8u_C4R>::cvt},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0},
-                {NppCvt<CV_16S, CV_32S, nppiConvert_16s32s_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {NppCvt<CV_16S, CV_32F, nppiConvert_16s32f_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0}
-            },
-            {
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0}
-            },
-            {
-                {NppCvt<CV_32F, CV_8U, nppiConvert_32f8u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {NppCvt<CV_32F, CV_16U, nppiConvert_32f16u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {NppCvt<CV_32F, CV_16S, nppiConvert_32f16s_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0}
-            },
-            {
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
-                {0,0,0,0},
-                {0,0,0,0}
-            },
-            {
-                {0,0,0,0},{0,0,0,0},{0,0,0,0},{0,0,0,0},{0,0,0,0},{0,0,0,0},{0,0,0,0},{0,0,0,0}
-            }
-        };
-
-        convert_callers[sdepth][ddepth][scn-1](*psrc, dst);
-    }
-}
-
-GpuMat& GpuMat::operator = (const Scalar& s)
-{
-    setTo(s);
-    return *this;
-}
-
-namespace
-{
+    //////////////////////////////////////////////////////////////////////////
+    // Set
+    
     template<int SDEPTH, int SCN> struct NppSetFunc
     {
         typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
@@ -655,7 +456,9 @@ namespace
             NppiSize sz;
             sz.width = src.cols;
             sz.height = src.rows;
+
             Scalar_<src_t> nppS = s;
+
             nppSafeCall( func(nppS.val, src.ptr<src_t>(), static_cast<int>(src.step), sz) );
 
             cudaSafeCall( cudaDeviceSynchronize() );
@@ -670,7 +473,9 @@ namespace
             NppiSize sz;
             sz.width = src.cols;
             sz.height = src.rows;
+
             Scalar_<src_t> nppS = s;
+
             nppSafeCall( func(nppS[0], src.ptr<src_t>(), static_cast<int>(src.step), sz) );
 
             cudaSafeCall( cudaDeviceSynchronize() );
@@ -706,7 +511,9 @@ namespace
             NppiSize sz;
             sz.width = src.cols;
             sz.height = src.rows;
+
             Scalar_<src_t> nppS = s;
+
             nppSafeCall( func(nppS.val, src.ptr<src_t>(), static_cast<int>(src.step), sz, mask.ptr<Npp8u>(), static_cast<int>(mask.step)) );
 
             cudaSafeCall( cudaDeviceSynchronize() );
@@ -721,7 +528,9 @@ namespace
             NppiSize sz;
             sz.width = src.cols;
             sz.height = src.rows;
+
             Scalar_<src_t> nppS = s;
+
             nppSafeCall( func(nppS[0], src.ptr<src_t>(), static_cast<int>(src.step), sz, mask.ptr<Npp8u>(), static_cast<int>(mask.step)) );
 
             cudaSafeCall( cudaDeviceSynchronize() );
@@ -736,122 +545,281 @@ namespace
     }
 }
 
-GpuMat& GpuMat::setTo(const Scalar& s, const GpuMat& mask)
+class CudaFuncTable : public GpuFuncTable
 {
-    CV_Assert(mask.type() == CV_8UC1);
+public:
+    virtual void copy(const Mat& src, GpuMat& dst) const 
+    { 
+        cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyHostToDevice) );
+    }
+    virtual void copy(const GpuMat& src, Mat& dst) const
+    { 
+        cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToHost) );
+    }
+    virtual void copy(const GpuMat& src, GpuMat& dst) const
+    { 
+        cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, src.data, src.step, src.cols * src.elemSize(), src.rows, cudaMemcpyDeviceToDevice) );
+    }
 
-    CV_Assert((depth() != CV_64F) || 
-        (TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE)));
+    virtual void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask) const 
+    { 
+        device::copy_to_with_mask(src, dst, src.depth(), mask, src.channels());
+    }
 
-    CV_DbgAssert(!this->empty());
-
-    NppiSize sz;
-    sz.width  = cols;
-    sz.height = rows;
-
-    if (mask.empty())
-    {
-        if (s[0] == 0.0 && s[1] == 0.0 && s[2] == 0.0 && s[3] == 0.0)
+    void convert(const GpuMat& src, GpuMat& dst) const 
+    { 
+        typedef void (*caller_t)(const GpuMat& src, GpuMat& dst);
+        static const caller_t callers[7][7][7] =
         {
-            cudaSafeCall( cudaMemset2D(data, step, 0, cols * elemSize(), rows) );
-            return *this;
-        }
-        if (depth() == CV_8U)
-        {
-            int cn = channels();
-
-            if (cn == 1 || (cn == 2 && s[0] == s[1]) || (cn == 3 && s[0] == s[1] && s[0] == s[2]) || (cn == 4 && s[0] == s[1] && s[0] == s[2] && s[0] == s[3]))
+            {                
+                /*  8U ->  8U */ {0, 0, 0, 0},
+                /*  8U ->  8S */ {convertToKernelCaller, convertToKernelCaller, convertToKernelCaller, convertToKernelCaller},
+                /*  8U -> 16U */ {NppCvt<CV_8U, CV_16U, nppiConvert_8u16u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,NppCvt<CV_8U, CV_16U, nppiConvert_8u16u_C4R>::cvt},
+                /*  8U -> 16S */ {NppCvt<CV_8U, CV_16S, nppiConvert_8u16s_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,NppCvt<CV_8U, CV_16S, nppiConvert_8u16s_C4R>::cvt},
+                /*  8U -> 32S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /*  8U -> 32F */ {NppCvt<CV_8U, CV_32F, nppiConvert_8u32f_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /*  8U -> 64F */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller}
+            },
             {
-                int val = saturate_cast<uchar>(s[0]);
-                cudaSafeCall( cudaMemset2D(data, step, val, cols * elemSize(), rows) );
-                return *this;
+                /*  8S ->  8U */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /*  8S ->  8S */ {0,0,0,0},
+                /*  8S -> 16U */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /*  8S -> 16S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /*  8S -> 32S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /*  8S -> 32F */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /*  8S -> 64F */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller}
+            },
+            {
+                /* 16U ->  8U */ {NppCvt<CV_16U, CV_8U, nppiConvert_16u8u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,NppCvt<CV_16U, CV_8U, nppiConvert_16u8u_C4R>::cvt},
+                /* 16U ->  8S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 16U -> 16U */ {0,0,0,0},
+                /* 16U -> 16S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 16U -> 32S */ {NppCvt<CV_16U, CV_32S, nppiConvert_16u32s_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 16U -> 32F */ {NppCvt<CV_16U, CV_32F, nppiConvert_16u32f_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 16U -> 64F */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller}
+            },
+            {
+                /* 16S ->  8U */ {NppCvt<CV_16S, CV_8U, nppiConvert_16s8u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,NppCvt<CV_16S, CV_8U, nppiConvert_16s8u_C4R>::cvt},
+                /* 16S ->  8S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 16S -> 16U */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 16S -> 16S */ {0,0,0,0},
+                /* 16S -> 32S */ {NppCvt<CV_16S, CV_32S, nppiConvert_16s32s_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 16S -> 32F */ {NppCvt<CV_16S, CV_32F, nppiConvert_16s32f_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 16S -> 64F */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller}
+            },
+            {
+                /* 32S ->  8U */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32S ->  8S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32S -> 16U */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32S -> 16S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32S -> 32S */ {0,0,0,0},
+                /* 32S -> 32F */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32S -> 64F */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller}
+            },
+            {
+                /* 32F ->  8U */ {NppCvt<CV_32F, CV_8U, nppiConvert_32f8u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32F ->  8S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32F -> 16U */ {NppCvt<CV_32F, CV_16U, nppiConvert_32f16u_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32F -> 16S */ {NppCvt<CV_32F, CV_16S, nppiConvert_32f16s_C1R>::cvt,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32F -> 32S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 32F -> 32F */ {0,0,0,0},
+                /* 32F -> 64F */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller}
+            },
+            {
+                /* 64F ->  8U */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 64F ->  8S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 64F -> 16U */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 64F -> 16S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 64F -> 32S */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 64F -> 32F */ {convertToKernelCaller,convertToKernelCaller,convertToKernelCaller,convertToKernelCaller},
+                /* 64F -> 64F */ {0,0,0,0}
             }
-        }
-        typedef void (*set_caller_t)(GpuMat& src, const Scalar& s);
-        static const set_caller_t set_callers[8][4] =
-        {
-            {NppSet<CV_8U, 1, nppiSet_8u_C1R>::set,kernelSet<uchar>,kernelSet<uchar>,NppSet<CV_8U, 4, nppiSet_8u_C4R>::set},
-            {kernelSet<schar>,kernelSet<schar>,kernelSet<schar>,kernelSet<schar>},
-            {NppSet<CV_16U, 1, nppiSet_16u_C1R>::set,NppSet<CV_16U, 2, nppiSet_16u_C2R>::set,kernelSet<ushort>,NppSet<CV_16U, 4, nppiSet_16u_C4R>::set},
-            {NppSet<CV_16S, 1, nppiSet_16s_C1R>::set,NppSet<CV_16S, 2, nppiSet_16s_C2R>::set,kernelSet<short>,NppSet<CV_16S, 4, nppiSet_16s_C4R>::set},
-            {NppSet<CV_32S, 1, nppiSet_32s_C1R>::set,kernelSet<int>,kernelSet<int>,NppSet<CV_32S, 4, nppiSet_32s_C4R>::set},
-            {NppSet<CV_32F, 1, nppiSet_32f_C1R>::set,kernelSet<float>,kernelSet<float>,NppSet<CV_32F, 4, nppiSet_32f_C4R>::set},
-            {kernelSet<double>,kernelSet<double>,kernelSet<double>,kernelSet<double>},
-            {0,0,0,0}
         };
-        set_callers[depth()][channels()-1](*this, s);
+
+        caller_t func = callers[src.depth()][dst.depth()][src.channels() - 1];
+        CV_DbgAssert(func != 0);
+
+        func(src, dst);
     }
-    else
-    {
-        typedef void (*set_caller_t)(GpuMat& src, const Scalar& s, const GpuMat& mask);
-        static const set_caller_t set_callers[8][4] =
-        {
-            {NppSetMask<CV_8U, 1, nppiSet_8u_C1MR>::set,kernelSetMask<uchar>,kernelSetMask<uchar>,NppSetMask<CV_8U, 4, nppiSet_8u_C4MR>::set},
-            {kernelSetMask<schar>,kernelSetMask<schar>,kernelSetMask<schar>,kernelSetMask<schar>},
-            {NppSetMask<CV_16U, 1, nppiSet_16u_C1MR>::set,kernelSetMask<ushort>,kernelSetMask<ushort>,NppSetMask<CV_16U, 4, nppiSet_16u_C4MR>::set},
-            {NppSetMask<CV_16S, 1, nppiSet_16s_C1MR>::set,kernelSetMask<short>,kernelSetMask<short>,NppSetMask<CV_16S, 4, nppiSet_16s_C4MR>::set},
-            {NppSetMask<CV_32S, 1, nppiSet_32s_C1MR>::set,kernelSetMask<int>,kernelSetMask<int>,NppSetMask<CV_32S, 4, nppiSet_32s_C4MR>::set},
-            {NppSetMask<CV_32F, 1, nppiSet_32f_C1MR>::set,kernelSetMask<float>,kernelSetMask<float>,NppSetMask<CV_32F, 4, nppiSet_32f_C4MR>::set},
-            {kernelSetMask<double>,kernelSetMask<double>,kernelSetMask<double>,kernelSetMask<double>},
-            {0,0,0,0}
-        };
-        set_callers[depth()][channels()-1](*this, s, mask);
+    void convert(const GpuMat& src, GpuMat& dst, double alpha, double beta) const 
+    { 
+        device::convert_gpu(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), alpha, beta);
     }
 
-    return *this;
+    void setTo(GpuMat& m, const Scalar& s, const GpuMat& mask) const
+    {
+        NppiSize sz;
+        sz.width  = m.cols;
+        sz.height = m.rows;
+
+        if (mask.empty())
+        {
+            if (s[0] == 0.0 && s[1] == 0.0 && s[2] == 0.0 && s[3] == 0.0)
+            {
+                cudaSafeCall( cudaMemset2D(m.data, m.step, 0, m.cols * m.elemSize(), m.rows) );
+                return;
+            }
+
+            if (m.depth() == CV_8U)
+            {
+                int cn = m.channels();
+
+                if (cn == 1 || (cn == 2 && s[0] == s[1]) || (cn == 3 && s[0] == s[1] && s[0] == s[2]) || (cn == 4 && s[0] == s[1] && s[0] == s[2] && s[0] == s[3]))
+                {
+                    int val = saturate_cast<gpu::uchar>(s[0]);
+                    cudaSafeCall( cudaMemset2D(m.data, m.step, val, m.cols * m.elemSize(), m.rows) );
+                    return;
+                }
+            }
+
+            typedef void (*caller_t)(GpuMat& src, const Scalar& s);
+            static const caller_t callers[7][4] =
+            {
+                {NppSet<CV_8U, 1, nppiSet_8u_C1R>::set,kernelSet<gpu::uchar>,kernelSet<gpu::uchar>,NppSet<CV_8U, 4, nppiSet_8u_C4R>::set},
+                {kernelSet<gpu::schar>,kernelSet<gpu::schar>,kernelSet<gpu::schar>,kernelSet<gpu::schar>},
+                {NppSet<CV_16U, 1, nppiSet_16u_C1R>::set,NppSet<CV_16U, 2, nppiSet_16u_C2R>::set,kernelSet<gpu::ushort>,NppSet<CV_16U, 4, nppiSet_16u_C4R>::set},
+                {NppSet<CV_16S, 1, nppiSet_16s_C1R>::set,NppSet<CV_16S, 2, nppiSet_16s_C2R>::set,kernelSet<short>,NppSet<CV_16S, 4, nppiSet_16s_C4R>::set},
+                {NppSet<CV_32S, 1, nppiSet_32s_C1R>::set,kernelSet<int>,kernelSet<int>,NppSet<CV_32S, 4, nppiSet_32s_C4R>::set},
+                {NppSet<CV_32F, 1, nppiSet_32f_C1R>::set,kernelSet<float>,kernelSet<float>,NppSet<CV_32F, 4, nppiSet_32f_C4R>::set},
+                {kernelSet<double>,kernelSet<double>,kernelSet<double>,kernelSet<double>}
+            };
+
+            callers[m.depth()][m.channels() - 1](m, s);
+        }
+        else
+        {
+            typedef void (*caller_t)(GpuMat& src, const Scalar& s, const GpuMat& mask);
+
+            static const caller_t callers[7][4] =
+            {
+                {NppSetMask<CV_8U, 1, nppiSet_8u_C1MR>::set,kernelSetMask<gpu::uchar>,kernelSetMask<gpu::uchar>,NppSetMask<CV_8U, 4, nppiSet_8u_C4MR>::set},
+                {kernelSetMask<gpu::schar>,kernelSetMask<gpu::schar>,kernelSetMask<gpu::schar>,kernelSetMask<gpu::schar>},
+                {NppSetMask<CV_16U, 1, nppiSet_16u_C1MR>::set,kernelSetMask<gpu::ushort>,kernelSetMask<gpu::ushort>,NppSetMask<CV_16U, 4, nppiSet_16u_C4MR>::set},
+                {NppSetMask<CV_16S, 1, nppiSet_16s_C1MR>::set,kernelSetMask<short>,kernelSetMask<short>,NppSetMask<CV_16S, 4, nppiSet_16s_C4MR>::set},
+                {NppSetMask<CV_32S, 1, nppiSet_32s_C1MR>::set,kernelSetMask<int>,kernelSetMask<int>,NppSetMask<CV_32S, 4, nppiSet_32s_C4MR>::set},
+                {NppSetMask<CV_32F, 1, nppiSet_32f_C1MR>::set,kernelSetMask<float>,kernelSetMask<float>,NppSetMask<CV_32F, 4, nppiSet_32f_C4MR>::set},
+                {kernelSetMask<double>,kernelSetMask<double>,kernelSetMask<double>,kernelSetMask<double>}
+            };
+
+            callers[m.depth()][m.channels() - 1](m, s, mask);
+        }
+    }
+
+    void mallocPitch(void** devPtr, size_t* step, size_t width, size_t height) const
+    {
+        cudaSafeCall( cudaMallocPitch(devPtr, step, width, height) );
+    }
+
+    void free(void* devPtr) const
+    {
+        cudaFree(devPtr);
+    }
+};
+
+const GpuFuncTable* gpuFuncTable()
+{
+    static CudaFuncTable cuda;
+    return &cuda;
 }
 
+#endif
 
-GpuMat cv::gpu::GpuMat::reshape(int new_cn, int new_rows) const
+void cv::gpu::GpuMat::upload(const Mat& m)
 {
-    GpuMat hdr = *this;
+    CV_DbgAssert(!m.empty());
 
-    int cn = channels();
-    if( new_cn == 0 )
-        new_cn = cn;
+    create(m.size(), m.type());
 
-    int total_width = cols * cn;
+    gpuFuncTable()->copy(m, *this);
+}
 
-    if( (new_cn > total_width || total_width % new_cn != 0) && new_rows == 0 )
-        new_rows = rows * total_width / new_cn;
+void cv::gpu::GpuMat::download(cv::Mat& m) const
+{
+    CV_DbgAssert(!empty());
 
-    if( new_rows != 0 && new_rows != rows )
+    m.create(size(), type());
+
+    gpuFuncTable()->copy(*this, m);
+}
+
+void cv::gpu::GpuMat::copyTo(GpuMat& m) const
+{
+    CV_DbgAssert(!empty());
+
+    m.create(size(), type());
+
+    gpuFuncTable()->copy(*this, m);
+}
+
+void cv::gpu::GpuMat::copyTo(GpuMat& mat, const GpuMat& mask) const
+{
+    if (mask.empty())
+        copyTo(mat);
+    else
     {
-        int total_size = total_width * rows;
-        if( !isContinuous() )
-            CV_Error( CV_BadStep, "The matrix is not continuous, thus its number of rows can not be changed" );
+        mat.create(size(), type());
 
-        if( (unsigned)new_rows > (unsigned)total_size )
-            CV_Error( CV_StsOutOfRange, "Bad new number of rows" );
+        gpuFuncTable()->copyWithMask(*this, mat, mask);
+    }
+}
 
-        total_width = total_size / new_rows;
+void cv::gpu::GpuMat::convertTo(GpuMat& dst, int rtype, double alpha, double beta) const
+{
+    bool noScale = fabs(alpha - 1) < numeric_limits<double>::epsilon() && fabs(beta) < numeric_limits<double>::epsilon();
 
-        if( total_width * new_rows != total_size )
-            CV_Error( CV_StsBadArg, "The total number of matrix elements is not divisible by the new number of rows" );
+    if (rtype < 0)
+        rtype = type();
+    else
+        rtype = CV_MAKETYPE(CV_MAT_DEPTH(rtype), channels());
 
-        hdr.rows = new_rows;
-        hdr.step = total_width * elemSize1();
+    int scn = channels();
+    int sdepth = depth();
+    int ddepth = CV_MAT_DEPTH(rtype);
+    if (sdepth == ddepth && noScale)
+    {
+        copyTo(dst);
+        return;
     }
 
-    int new_width = total_width / new_cn;
+    GpuMat temp;
+    const GpuMat* psrc = this;
+    if (sdepth != ddepth && psrc == &dst)
+    {
+        temp = *this;
+        psrc = &temp;
+    }
 
-    if( new_width * new_cn != total_width )
-        CV_Error( CV_BadNumChannels, "The total width is not divisible by the new number of channels" );
+    dst.create(size(), rtype);
 
-    hdr.cols = new_width;
-    hdr.flags = (hdr.flags & ~CV_MAT_CN_MASK) | ((new_cn-1) << CV_CN_SHIFT);
-    return hdr;
+    if (noScale)
+        gpuFuncTable()->convert(*psrc, dst);
+    else
+        gpuFuncTable()->convert(*psrc, dst, alpha, beta);
+}
+
+GpuMat& cv::gpu::GpuMat::setTo(const Scalar& s, const GpuMat& mask)
+{
+    CV_Assert(mask.empty() || mask.type() == CV_8UC1);
+    CV_DbgAssert(!empty());
+
+    gpuFuncTable()->setTo(*this, s, mask);    
+
+    return *this;
 }
 
 void cv::gpu::GpuMat::create(int _rows, int _cols, int _type)
 {
     _type &= TYPE_MASK;
-    if( rows == _rows && cols == _cols && type() == _type && data )
+
+    if (rows == _rows && cols == _cols && type() == _type && data)
         return;
-    if( data )
+
+    if (data)
         release();
-    CV_DbgAssert( _rows >= 0 && _cols >= 0 );
-    if( _rows > 0 && _cols > 0 )
+
+    CV_DbgAssert(_rows >= 0 && _cols >= 0);
+
+    if (_rows > 0 && _cols > 0)
     {
         flags = Mat::MAGIC_VAL + _type;
         rows = _rows;
@@ -859,8 +827,8 @@ void cv::gpu::GpuMat::create(int _rows, int _cols, int _type)
 
         size_t esz = elemSize();
 
-        void *dev_ptr;
-        cudaSafeCall( cudaMallocPitch(&dev_ptr, &step, esz * cols, rows) );
+        void* devPtr;
+        gpuFuncTable()->mallocPitch(&devPtr, &step, esz * cols, rows);
 
         // Single row must be continuous
         if (rows == 1)
@@ -869,43 +837,27 @@ void cv::gpu::GpuMat::create(int _rows, int _cols, int _type)
         if (esz * cols == step)
             flags |= Mat::CONTINUOUS_FLAG;
 
-        int64 _nettosize = (int64)step*rows;
-        size_t nettosize = (size_t)_nettosize;
+        int64 _nettosize = static_cast<int64>(step) * rows;
+        size_t nettosize = static_cast<size_t>(_nettosize);
 
-        datastart = data = (uchar*)dev_ptr;
+        datastart = data = static_cast<uchar*>(devPtr);
         dataend = data + nettosize;
 
-        refcount = (int*)fastMalloc(sizeof(*refcount));
+        refcount = static_cast<int*>(fastMalloc(sizeof(*refcount)));
         *refcount = 1;
     }
 }
 
 void cv::gpu::GpuMat::release()
 {
-    if( refcount && CV_XADD(refcount, -1) == 1 )
+    if (refcount && CV_XADD(refcount, -1) == 1)
     {
         fastFree(refcount);
-        cudaFree(datastart);
+
+        gpuFuncTable()->free(datastart);
     }
+
     data = datastart = dataend = 0;
     step = rows = cols = 0;
     refcount = 0;
 }
-
-void cv::gpu::createContinuous(int rows, int cols, int type, GpuMat& m)
-{
-    int area = rows * cols;
-    if (!m.isContinuous() || m.type() != type || m.size().area() != area)
-        m.create(1, area, type);
-    m = m.reshape(0, rows);
-}
-
-void cv::gpu::ensureSizeIsEnough(int rows, int cols, int type, GpuMat& m)
-{
-    if (m.type() == type && m.rows >= rows && m.cols >= cols)
-        m = m(Rect(0, 0, cols, rows));
-    else
-        m.create(rows, cols, type);
-}
-
-#endif /* !defined (HAVE_CUDA) */
diff --git a/modules/stitching/src/blenders.cpp b/modules/stitching/src/blenders.cpp
index 47c3ce2b25..ec6959dd94 100644
--- a/modules/stitching/src/blenders.cpp
+++ b/modules/stitching/src/blenders.cpp
@@ -387,7 +387,7 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, vector<Mat> &pyr)
     pyr.resize(num_levels + 1);
 
     vector<gpu::GpuMat> gpu_pyr(num_levels + 1);
-    gpu_pyr[0] = img;
+    gpu_pyr[0].upload(img);
     for (int i = 0; i < num_levels; ++i)
         gpu::pyrDown(gpu_pyr[i], gpu_pyr[i + 1]);
 
@@ -396,10 +396,10 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, vector<Mat> &pyr)
     {
         gpu::pyrUp(gpu_pyr[i + 1], tmp);
         gpu::subtract(gpu_pyr[i], tmp, gpu_pyr[i]);
-        pyr[i] = gpu_pyr[i];
+        gpu_pyr[i].download(pyr[i]);
     }
 
-    pyr[num_levels] = gpu_pyr[num_levels];
+    gpu_pyr[num_levels].download(pyr[num_levels]);
 #endif
 }
 
@@ -425,7 +425,7 @@ void restoreImageFromLaplacePyrGpu(vector<Mat> &pyr)
 
     vector<gpu::GpuMat> gpu_pyr(pyr.size());
     for (size_t i = 0; i < pyr.size(); ++i)
-        gpu_pyr[i] = pyr[i];
+        gpu_pyr[i].upload(pyr[i]);
 
     gpu::GpuMat tmp;
     for (size_t i = pyr.size() - 1; i > 0; --i)
@@ -434,7 +434,7 @@ void restoreImageFromLaplacePyrGpu(vector<Mat> &pyr)
         gpu::add(tmp, gpu_pyr[i - 1], gpu_pyr[i - 1]);
     }
 
-    pyr[0] = gpu_pyr[0];
+    gpu_pyr[0].download(pyr[0]);
 #endif
 }
 
diff --git a/samples/gpu/performance/tests.cpp b/samples/gpu/performance/tests.cpp
index f4a965bb1a..bb0d26fb53 100644
--- a/samples/gpu/performance/tests.cpp
+++ b/samples/gpu/performance/tests.cpp
@@ -37,7 +37,7 @@ TEST(matchTemplate)
         matchTemplate(src, templ, dst, CV_TM_CCORR);
         CPU_OFF;
 
-        d_templ = templ;
+        d_templ.upload(templ);
         d_dst.create(d_src.rows - d_templ.rows + 1, d_src.cols - d_templ.cols + 1, CV_32F);
 
         GPU_ON;
@@ -65,7 +65,7 @@ TEST(minMaxLoc)
         minMaxLoc(src, &min_val, &max_val, &min_loc, &max_loc);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
 
         GPU_ON;
         gpu::minMaxLoc(d_src, &min_val, &max_val, &min_loc, &max_loc);
@@ -107,9 +107,9 @@ TEST(remap)
         remap(src, dst, xmap, ymap, interpolation, borderMode);
         CPU_OFF;
 
-        d_src = src;
-        d_xmap = xmap;
-        d_ymap = ymap;
+        d_src.upload(src);
+        d_xmap.upload(xmap);
+        d_ymap.upload(ymap);
         d_dst.create(d_xmap.size(), d_src.type());
 
         GPU_ON;
@@ -142,9 +142,9 @@ TEST(remap)
         remap(src, dst, xmap, ymap, interpolation, borderMode);
         CPU_OFF;
 
-        d_src = src;
-        d_xmap = xmap;
-        d_ymap = ymap;
+        d_src.upload(src);
+        d_xmap.upload(xmap);
+        d_ymap.upload(ymap);
         d_dst.create(d_xmap.size(), d_src.type());
 
         GPU_ON;
@@ -177,9 +177,9 @@ TEST(remap)
         remap(src, dst, xmap, ymap, interpolation, borderMode);
         CPU_OFF;
 
-        d_src = src;
-        d_xmap = xmap;
-        d_ymap = ymap;
+        d_src.upload(src);
+        d_xmap.upload(xmap);
+        d_ymap.upload(ymap);
         d_dst.create(d_xmap.size(), d_src.type());
 
         GPU_ON;
@@ -212,9 +212,9 @@ TEST(remap)
         remap(src, dst, xmap, ymap, interpolation, borderMode);
         CPU_OFF;
 
-        d_src = src;
-        d_xmap = xmap;
-        d_ymap = ymap;
+        d_src.upload(src);
+        d_xmap.upload(xmap);
+        d_ymap.upload(ymap);
         d_dst.create(d_xmap.size(), d_src.type());
 
         GPU_ON;
@@ -240,7 +240,7 @@ TEST(dft)
         dft(src, dst);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(d_src.size(), d_src.type());
 
         GPU_ON;
@@ -266,7 +266,7 @@ TEST(cornerHarris)
         cornerHarris(src, dst, 5, 7, 0.1, BORDER_REFLECT101);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(src.size(), src.type());
 
         GPU_ON;
@@ -286,7 +286,7 @@ TEST(integral)
     gen(src, size, size, CV_8U, 0, 256);
     sum.create(size + 1, size + 1, CV_32S);
 
-    d_src = src;
+    d_src.upload(src);
     d_sum.create(size + 1, size + 1, CV_32S);
 
     for (int i = 0; i < 5; ++i)
@@ -320,7 +320,7 @@ TEST(norm)
             norm(src, NORM_INF);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
 
         GPU_ON;
         for (int i = 0; i < 5; ++i)
@@ -350,7 +350,7 @@ TEST(meanShift)
 
         gen(src, size, size, CV_8UC4, Scalar::all(0), Scalar::all(256));
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(d_src.size(), d_src.type());
 
         GPU_ON;
@@ -483,8 +483,8 @@ TEST(magnitude)
         magnitude(x, y, mag);
         CPU_OFF;
 
-        d_x = x;
-        d_y = y;
+        d_x.upload(x);
+        d_y.upload(y);
         d_mag.create(size, size, CV_32F);
 
         GPU_ON;
@@ -511,8 +511,8 @@ TEST(add)
         add(src1, src2, dst);
         CPU_OFF;
 
-        d_src1 = src1;
-        d_src2 = src2;
+        d_src1.upload(src1);
+        d_src2.upload(src2);
         d_dst.create(size, size, CV_32F);
 
         GPU_ON;
@@ -538,7 +538,7 @@ TEST(log)
         log(src, dst);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size, size, CV_32F);
 
         GPU_ON;
@@ -564,7 +564,7 @@ TEST(exp)
         exp(src, dst);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size, size, CV_32F);
 
         GPU_ON;
@@ -591,8 +591,8 @@ TEST(mulSpectrums)
         mulSpectrums(src1, src2, dst, 0, true);
         CPU_OFF;
 
-        d_src1 = src1;
-        d_src2 = src2;
+        d_src1.upload(src1);
+        d_src2.upload(src2);
         d_dst.create(size, size, CV_32FC2);
 
         GPU_ON;
@@ -618,7 +618,7 @@ TEST(resize)
         resize(src, dst, dst.size());
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size * 2, size * 2, CV_8U);
 
         GPU_ON;
@@ -636,7 +636,7 @@ TEST(resize)
         resize(src, dst, dst.size());
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size / 2, size / 2, CV_8U);
 
         GPU_ON;
@@ -654,7 +654,7 @@ TEST(resize)
         resize(src, dst, dst.size());
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size * 2, size * 2, CV_8U);
 
         GPU_ON;
@@ -672,7 +672,7 @@ TEST(resize)
         resize(src, dst, dst.size());
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size / 2, size / 2, CV_8U);
 
         GPU_ON;
@@ -690,7 +690,7 @@ TEST(resize)
         resize(src, dst, dst.size());
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size * 2, size * 2, CV_8U);
 
         GPU_ON;
@@ -708,7 +708,7 @@ TEST(resize)
         resize(src, dst, dst.size());
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size / 2, size / 2, CV_8U);
 
         GPU_ON;
@@ -726,7 +726,7 @@ TEST(resize)
         resize(src, dst, dst.size());
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size * 2, size * 2, CV_8U);
 
         GPU_ON;
@@ -744,7 +744,7 @@ TEST(resize)
         resize(src, dst, dst.size());
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size / 2, size / 2, CV_8U);
 
         GPU_ON;
@@ -900,7 +900,7 @@ TEST(erode)
         erode(src, dst, ker);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(d_src.size(), d_src.type());
 
         GPU_ON;
@@ -925,7 +925,7 @@ TEST(threshold)
         threshold(src, dst, 50.0, 0.0, THRESH_BINARY);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size, size, CV_8U);
 
         GPU_ON;
@@ -944,7 +944,7 @@ TEST(threshold)
         threshold(src, dst, 50.0, 0.0, THRESH_BINARY);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size, size, CV_32F);
 
         GPU_ON;
@@ -969,7 +969,7 @@ TEST(pow)
         pow(src, -2.0, dst);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(size, size, CV_32F);
 
         GPU_ON;
@@ -1004,7 +1004,7 @@ TEST(projectPoints)
         projectPoints(src, rvec, tvec, camera_mat, Mat::zeros(1, 8, CV_32F), dst);
         CPU_OFF;
 
-        d_src = src;
+        d_src.upload(src);
         d_dst.create(1, size, CV_32FC2);
 
         GPU_ON;
@@ -1491,9 +1491,9 @@ TEST(gemm)
         gemm(src1, src2, 1.0, src3, 1.0, dst);
         CPU_OFF;
 
-        d_src1 = src1;
-        d_src2 = src2;
-        d_src3 = src3;
+        d_src1.upload(src1);
+        d_src2.upload(src2);
+        d_src3.upload(src3);
         d_dst.create(d_src1.size(), d_src1.type());
 
         GPU_ON;