remove softcascade host dependencies on gpu module

modules/core/include/opencv2/core/gpumat.hpp

@@ -50,6 +50,141 @@
 
 namespace cv { namespace gpu
 {
+    //////////////////////////////// CudaMem ////////////////////////////////
+    // CudaMem is limited cv::Mat with page locked memory allocation.
+    // Page locked memory is only needed for async and faster coping to GPU.
+    // It is convertable to cv::Mat header without reference counting
+    // so you can use it with other opencv functions.
+
+    // Page-locks the matrix m memory and maps it for the device(s)
+    CV_EXPORTS void registerPageLocked(Mat& m);
+    // Unmaps the memory of matrix m, and makes it pageable again.
+    CV_EXPORTS void unregisterPageLocked(Mat& m);
+
+    class CV_EXPORTS CudaMem
+    {
+    public:
+        enum  { ALLOC_PAGE_LOCKED = 1, ALLOC_ZEROCOPY = 2, ALLOC_WRITE_COMBINED = 4 };
+
+        CudaMem();
+        CudaMem(const CudaMem& m);
+
+        CudaMem(int rows, int cols, int type, int _alloc_type = ALLOC_PAGE_LOCKED);
+        CudaMem(Size size, int type, int alloc_type = ALLOC_PAGE_LOCKED);
+
+
+        //! creates from cv::Mat with coping data
+        explicit CudaMem(const Mat& m, int alloc_type = ALLOC_PAGE_LOCKED);
+
+        ~CudaMem();
+
+        CudaMem& operator = (const CudaMem& m);
+
+        //! returns deep copy of the matrix, i.e. the data is copied
+        CudaMem clone() const;
+
+        //! allocates new matrix data unless the matrix already has specified size and type.
+        void create(int rows, int cols, int type, int alloc_type = ALLOC_PAGE_LOCKED);
+        void create(Size size, int type, int alloc_type = ALLOC_PAGE_LOCKED);
+
+        //! decrements reference counter and released memory if needed.
+        void release();
+
+        //! returns matrix header with disabled reference counting for CudaMem data.
+        Mat createMatHeader() const;
+        operator Mat() const;
+
+        //! maps host memory into device address space and returns GpuMat header for it. Throws exception if not supported by hardware.
+        GpuMat createGpuMatHeader() const;
+        operator GpuMat() const;
+
+        //returns if host memory can be mapperd to gpu address space;
+        static bool canMapHostMemory();
+
+        // Please see cv::Mat for descriptions
+        bool isContinuous() const;
+        size_t elemSize() const;
+        size_t elemSize1() const;
+        int type() const;
+        int depth() const;
+        int channels() const;
+        size_t step1() const;
+        Size size() const;
+        bool empty() const;
+
+
+        // Please see cv::Mat for descriptions
+        int flags;
+        int rows, cols;
+        size_t step;
+
+        uchar* data;
+        int* refcount;
+
+        uchar* datastart;
+        uchar* dataend;
+
+        int alloc_type;
+    };
+
+
+    //////////////////////////////// CudaStream ////////////////////////////////
+    // Encapculates Cuda Stream. Provides interface for async coping.
+    // Passed to each function that supports async kernel execution.
+    // Reference counting is enabled
+
+    class CV_EXPORTS Stream
+    {
+    public:
+        Stream();
+        ~Stream();
+
+        Stream(const Stream&);
+        Stream& operator =(const Stream&);
+
+        bool queryIfComplete();
+        void waitForCompletion();
+
+        //! downloads asynchronously
+        // Warning! cv::Mat must point to page locked memory (i.e. to CudaMem data or to its subMat)
+        void enqueueDownload(const GpuMat& src, CudaMem& dst);
+        void enqueueDownload(const GpuMat& src, Mat& dst);
+
+        //! uploads asynchronously
+        // Warning! cv::Mat must point to page locked memory (i.e. to CudaMem data or to its ROI)
+        void enqueueUpload(const CudaMem& src, GpuMat& dst);
+        void enqueueUpload(const Mat& src, GpuMat& dst);
+
+        //! copy asynchronously
+        void enqueueCopy(const GpuMat& src, GpuMat& dst);
+
+        //! memory set asynchronously
+        void enqueueMemSet(GpuMat& src, Scalar val);
+        void enqueueMemSet(GpuMat& src, Scalar val, const GpuMat& mask);
+
+        //! converts matrix type, ex from float to uchar depending on type
+        void enqueueConvert(const GpuMat& src, GpuMat& dst, int dtype, double a = 1, double b = 0);
+
+        //! adds a callback to be called on the host after all currently enqueued items in the stream have completed
+        typedef void (*StreamCallback)(Stream& stream, int status, void* userData);
+        void enqueueHostCallback(StreamCallback callback, void* userData);
+
+        static Stream& Null();
+
+        operator bool() const;
+
+    private:
+        struct Impl;
+
+        explicit Stream(Impl* impl);
+        void create();
+        void release();
+
+        Impl *impl;
+
+        friend struct StreamAccessor;
+    };
+
     //////////////////////////////// Initialization & Info ////////////////////////
 
     //! This is the only function that do not throw exceptions if the library is compiled without Cuda.

modules/core/include/opencv2/core/stream_accessor.hpp

@@ -0,0 +1,64 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other GpuMaterials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or 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*/
+
+#ifndef __OPENCV_GPU_STREAM_ACCESSOR_HPP__
+#define __OPENCV_GPU_STREAM_ACCESSOR_HPP__
+
+#include "opencv2/gpu/gpu.hpp"
+#include "cuda_runtime_api.h"
+
+namespace cv
+{
+    namespace gpu
+    {
+        // This is only header file that depends on Cuda. All other headers are independent.
+        // So if you use OpenCV binaries you do noot need to install Cuda Toolkit.
+        // But of you wanna use GPU by yourself, may get cuda stream instance using the class below.
+        // In this case you have to install Cuda Toolkit.
+        struct StreamAccessor
+        {
+            CV_EXPORTS static cudaStream_t getStream(const Stream& stream);
+        };
+    }
+}
+
+#endif /* __OPENCV_GPU_STREAM_ACCESSOR_HPP__ */

modules/core/src/cudastream.cpp

@@ -0,0 +1,377 @@
+/*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"
+
+using namespace cv;
+using namespace cv::gpu;
+
+#if !defined (HAVE_CUDA)
+
+cv::gpu::Stream::Stream() { throw_nogpu(); }
+cv::gpu::Stream::~Stream() {}
+cv::gpu::Stream::Stream(const Stream&) { throw_nogpu(); }
+Stream& cv::gpu::Stream::operator=(const Stream&) { throw_nogpu(); return *this; }
+bool cv::gpu::Stream::queryIfComplete() { throw_nogpu(); return false; }
+void cv::gpu::Stream::waitForCompletion() { throw_nogpu(); }
+void cv::gpu::Stream::enqueueDownload(const GpuMat&, Mat&) { throw_nogpu(); }
+void cv::gpu::Stream::enqueueDownload(const GpuMat&, CudaMem&) { throw_nogpu(); }
+void cv::gpu::Stream::enqueueUpload(const CudaMem&, GpuMat&) { throw_nogpu(); }
+void cv::gpu::Stream::enqueueUpload(const Mat&, GpuMat&) { throw_nogpu(); }
+void cv::gpu::Stream::enqueueCopy(const GpuMat&, GpuMat&) { throw_nogpu(); }
+void cv::gpu::Stream::enqueueMemSet(GpuMat&, Scalar) { throw_nogpu(); }
+void cv::gpu::Stream::enqueueMemSet(GpuMat&, Scalar, const GpuMat&) { throw_nogpu(); }
+void cv::gpu::Stream::enqueueConvert(const GpuMat&, GpuMat&, int, double, double) { throw_nogpu(); }
+void cv::gpu::Stream::enqueueHostCallback(StreamCallback, void*) { throw_nogpu(); }
+Stream& cv::gpu::Stream::Null() { throw_nogpu(); static Stream s; return s; }
+cv::gpu::Stream::operator bool() const { throw_nogpu(); return false; }
+cv::gpu::Stream::Stream(Impl*) { throw_nogpu(); }
+void cv::gpu::Stream::create() { throw_nogpu(); }
+void cv::gpu::Stream::release() { throw_nogpu(); }
+
+#else /* !defined (HAVE_CUDA) */
+
+#include "opencv2/gpu/stream_accessor.hpp"
+
+namespace
+{
+#if defined(__GNUC__)
+    #define cudaSafeCall(expr)  ___cudaSafeCall(expr, __FILE__, __LINE__, __func__)
+    #define nppSafeCall(expr)  ___nppSafeCall(expr, __FILE__, __LINE__, __func__)
+#else /* defined(__CUDACC__) || defined(__MSVC__) */
+    #define cudaSafeCall(expr)  ___cudaSafeCall(expr, __FILE__, __LINE__)
+    #define nppSafeCall(expr)  ___nppSafeCall(expr, __FILE__, __LINE__)
+#endif
+
+    inline void ___cudaSafeCall(cudaError_t err, const char *file, const int line, const char *func = "")
+    {
+        if (cudaSuccess != err)
+            cv::gpu::error(cudaGetErrorString(err), file, line, func);
+    }
+
+    inline void ___nppSafeCall(int err, const char *file, const int line, const char *func = "")
+    {
+        if (err < 0)
+        {
+            std::ostringstream msg;
+            msg << "NPP API Call Error: " << err;
+            cv::gpu::error(msg.str().c_str(), file, line, func);
+        }
+    }
+}
+
+namespace cv { namespace gpu
+{
+    void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream);
+    void convertTo(const GpuMat& src, GpuMat& dst, double alpha, double beta, cudaStream_t stream);
+    void setTo(GpuMat& src, Scalar s, cudaStream_t stream);
+    void setTo(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream);
+}}
+
+struct Stream::Impl
+{
+    static cudaStream_t getStream(const Impl* impl)
+    {
+        return impl ? impl->stream : 0;
+    }
+
+    cudaStream_t stream;
+    int ref_counter;
+};
+
+cudaStream_t cv::gpu::StreamAccessor::getStream(const Stream& stream)
+{
+    return Stream::Impl::getStream(stream.impl);
+}
+
+cv::gpu::Stream::Stream() : impl(0)
+{
+    create();
+}
+
+cv::gpu::Stream::~Stream()
+{
+    release();
+}
+
+cv::gpu::Stream::Stream(const Stream& stream) : impl(stream.impl)
+{
+    if (impl)
+        CV_XADD(&impl->ref_counter, 1);
+}
+
+Stream& cv::gpu::Stream::operator =(const Stream& stream)
+{
+    if (this != &stream)
+    {
+        release();
+        impl = stream.impl;
+        if (impl)
+            CV_XADD(&impl->ref_counter, 1);
+    }
+
+    return *this;
+}
+
+bool cv::gpu::Stream::queryIfComplete()
+{
+    cudaStream_t stream = Impl::getStream(impl);
+    cudaError_t err = cudaStreamQuery(stream);
+
+    if (err == cudaErrorNotReady || err == cudaSuccess)
+        return err == cudaSuccess;
+
+    cudaSafeCall(err);
+    return false;
+}
+
+void cv::gpu::Stream::waitForCompletion()
+{
+    cudaStream_t stream = Impl::getStream(impl);
+    cudaSafeCall( cudaStreamSynchronize(stream) );
+}
+
+void cv::gpu::Stream::enqueueDownload(const GpuMat& src, Mat& dst)
+{
+    // if not -> allocation will be done, but after that dst will not point to page locked memory
+    CV_Assert( src.size() == dst.size() && src.type() == dst.type() );
+
+    cudaStream_t stream = Impl::getStream(impl);
+    size_t bwidth = src.cols * src.elemSize();
+    cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) );
+}
+
+void cv::gpu::Stream::enqueueDownload(const GpuMat& src, CudaMem& dst)
+{
+    dst.create(src.size(), src.type(), CudaMem::ALLOC_PAGE_LOCKED);
+
+    cudaStream_t stream = Impl::getStream(impl);
+    size_t bwidth = src.cols * src.elemSize();
+    cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToHost, stream) );
+}
+
+void cv::gpu::Stream::enqueueUpload(const CudaMem& src, GpuMat& dst)
+{
+    dst.create(src.size(), src.type());
+
+    cudaStream_t stream = Impl::getStream(impl);
+    size_t bwidth = src.cols * src.elemSize();
+    cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) );
+}
+
+void cv::gpu::Stream::enqueueUpload(const Mat& src, GpuMat& dst)
+{
+    dst.create(src.size(), src.type());
+
+    cudaStream_t stream = Impl::getStream(impl);
+    size_t bwidth = src.cols * src.elemSize();
+    cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyHostToDevice, stream) );
+}
+
+void cv::gpu::Stream::enqueueCopy(const GpuMat& src, GpuMat& dst)
+{
+    dst.create(src.size(), src.type());
+
+    cudaStream_t stream = Impl::getStream(impl);
+    size_t bwidth = src.cols * src.elemSize();
+    cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, src.data, src.step, bwidth, src.rows, cudaMemcpyDeviceToDevice, stream) );
+}
+
+void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar val)
+{
+    const int sdepth = src.depth();
+
+    if (sdepth == CV_64F)
+    {
+        if (!deviceSupports(NATIVE_DOUBLE))
+            CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");
+    }
+
+    cudaStream_t stream = Impl::getStream(impl);
+
+    if (val[0] == 0.0 && val[1] == 0.0 && val[2] == 0.0 && val[3] == 0.0)
+    {
+        cudaSafeCall( cudaMemset2DAsync(src.data, src.step, 0, src.cols * src.elemSize(), src.rows, stream) );
+        return;
+    }
+
+    if (sdepth == CV_8U)
+    {
+        int cn = src.channels();
+
+        if (cn == 1 || (cn == 2 && val[0] == val[1]) || (cn == 3 && val[0] == val[1] && val[0] == val[2]) || (cn == 4 && val[0] == val[1] && val[0] == val[2] && val[0] == val[3]))
+        {
+            int ival = saturate_cast<uchar>(val[0]);
+            cudaSafeCall( cudaMemset2DAsync(src.data, src.step, ival, src.cols * src.elemSize(), src.rows, stream) );
+            return;
+        }
+    }
+
+    setTo(src, val, stream);
+}
+
+void cv::gpu::Stream::enqueueMemSet(GpuMat& src, Scalar val, const GpuMat& mask)
+{
+    const int sdepth = src.depth();
+
+    if (sdepth == CV_64F)
+    {
+        if (!deviceSupports(NATIVE_DOUBLE))
+            CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");
+    }
+
+    CV_Assert(mask.type() == CV_8UC1);
+
+    cudaStream_t stream = Impl::getStream(impl);
+
+    setTo(src, val, mask, stream);
+}
+
+void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int dtype, double alpha, double beta)
+{
+    if (dtype < 0)
+        dtype = src.type();
+    else
+        dtype = CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src.channels());
+
+    const int sdepth = src.depth();
+    const int ddepth = CV_MAT_DEPTH(dtype);
+
+    if (sdepth == CV_64F || ddepth == CV_64F)
+    {
+        if (!deviceSupports(NATIVE_DOUBLE))
+            CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");
+    }
+
+    bool noScale = fabs(alpha - 1) < std::numeric_limits<double>::epsilon()
+                && fabs(beta) < std::numeric_limits<double>::epsilon();
+
+    if (sdepth == ddepth && noScale)
+    {
+        enqueueCopy(src, dst);
+        return;
+    }
+
+    dst.create(src.size(), dtype);
+
+    cudaStream_t stream = Impl::getStream(impl);
+    convertTo(src, dst, alpha, beta, stream);
+}
+
+#if CUDA_VERSION >= 5000
+
+namespace
+{
+    struct CallbackData
+    {
+        cv::gpu::Stream::StreamCallback callback;
+        void* userData;
+        Stream stream;
+    };
+
+    void CUDART_CB cudaStreamCallback(cudaStream_t, cudaError_t status, void* userData)
+    {
+        CallbackData* data = reinterpret_cast<CallbackData*>(userData);
+        data->callback(data->stream, static_cast<int>(status), data->userData);
+        delete data;
+    }
+}
+
+#endif
+
+void cv::gpu::Stream::enqueueHostCallback(StreamCallback callback, void* userData)
+{
+#if CUDA_VERSION >= 5000
+    CallbackData* data = new CallbackData;
+    data->callback = callback;
+    data->userData = userData;
+    data->stream = *this;
+
+    cudaStream_t stream = Impl::getStream(impl);
+
+    cudaSafeCall( cudaStreamAddCallback(stream, cudaStreamCallback, data, 0) );
+#else
+    (void) callback;
+    (void) userData;
+    CV_Error(CV_StsNotImplemented, "This function requires CUDA 5.0");
+#endif
+}
+
+cv::gpu::Stream& cv::gpu::Stream::Null()
+{
+    static Stream s((Impl*) 0);
+    return s;
+}
+
+cv::gpu::Stream::operator bool() const
+{
+    return impl && impl->stream;
+}
+
+cv::gpu::Stream::Stream(Impl* impl_) : impl(impl_)
+{
+}
+
+void cv::gpu::Stream::create()
+{
+    if (impl)
+        release();
+
+    cudaStream_t stream;
+    cudaSafeCall( cudaStreamCreate( &stream ) );
+
+    impl = (Stream::Impl*) fastMalloc(sizeof(Stream::Impl));
+
+    impl->stream = stream;
+    impl->ref_counter = 1;
+}
+
+void cv::gpu::Stream::release()
+{
+    if (impl && CV_XADD(&impl->ref_counter, -1) == 1)
+    {
+        cudaSafeCall( cudaStreamDestroy(impl->stream) );
+        cv::fastFree(impl);
+    }
+}
+
+#endif /* !defined (HAVE_CUDA) */

modules/core/src/matrix_operations.cpp

@@ -0,0 +1,324 @@
+/*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 "opencv2/core/gpumat.hpp"
+
+using namespace cv;
+using namespace cv::gpu;
+
+cv::gpu::CudaMem::CudaMem()
+    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), alloc_type(0)
+{
+}
+
+cv::gpu::CudaMem::CudaMem(int _rows, int _cols, int _type, int _alloc_type)
+    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), alloc_type(0)
+{
+    if( _rows > 0 && _cols > 0 )
+        create( _rows, _cols, _type, _alloc_type);
+}
+
+cv::gpu::CudaMem::CudaMem(Size _size, int _type, int _alloc_type)
+    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), alloc_type(0)
+{
+    if( _size.height > 0 && _size.width > 0 )
+        create( _size.height, _size.width, _type, _alloc_type);
+}
+
+cv::gpu::CudaMem::CudaMem(const CudaMem& 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), alloc_type(m.alloc_type)
+{
+    if( refcount )
+        CV_XADD(refcount, 1);
+}
+
+cv::gpu::CudaMem::CudaMem(const Mat& m, int _alloc_type)
+    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), alloc_type(0)
+{
+    if( m.rows > 0 && m.cols > 0 )
+        create( m.size(), m.type(), _alloc_type);
+
+    Mat tmp = createMatHeader();
+    m.copyTo(tmp);
+}
+
+cv::gpu::CudaMem::~CudaMem()
+{
+    release();
+}
+
+CudaMem& cv::gpu::CudaMem::operator = (const CudaMem& 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;
+        alloc_type = m.alloc_type;
+    }
+    return *this;
+}
+
+CudaMem cv::gpu::CudaMem::clone() const
+{
+    CudaMem m(size(), type(), alloc_type);
+    Mat to = m;
+    Mat from = *this;
+    from.copyTo(to);
+    return m;
+}
+
+void cv::gpu::CudaMem::create(Size _size, int _type, int _alloc_type)
+{
+    create(_size.height, _size.width, _type, _alloc_type);
+}
+
+Mat cv::gpu::CudaMem::createMatHeader() const
+{
+    return Mat(size(), type(), data, step);
+}
+
+cv::gpu::CudaMem::operator Mat() const
+{
+    return createMatHeader();
+}
+
+cv::gpu::CudaMem::operator GpuMat() const
+{
+    return createGpuMatHeader();
+}
+
+bool cv::gpu::CudaMem::isContinuous() const
+{
+    return (flags & Mat::CONTINUOUS_FLAG) != 0;
+}
+
+size_t cv::gpu::CudaMem::elemSize() const
+{
+    return CV_ELEM_SIZE(flags);
+}
+
+size_t cv::gpu::CudaMem::elemSize1() const
+{
+    return CV_ELEM_SIZE1(flags);
+}
+
+int cv::gpu::CudaMem::type() const
+{
+    return CV_MAT_TYPE(flags);
+}
+
+int cv::gpu::CudaMem::depth() const
+{
+    return CV_MAT_DEPTH(flags);
+}
+
+int cv::gpu::CudaMem::channels() const
+{
+    return CV_MAT_CN(flags);
+}
+
+size_t cv::gpu::CudaMem::step1() const
+{
+    return step/elemSize1();
+}
+
+Size cv::gpu::CudaMem::size() const
+{
+    return Size(cols, rows);
+}
+
+bool cv::gpu::CudaMem::empty() const
+{
+    return data == 0;
+}
+
+#if !defined (HAVE_CUDA)
+
+void cv::gpu::registerPageLocked(Mat&) { throw_nogpu(); }
+void cv::gpu::unregisterPageLocked(Mat&) { throw_nogpu(); }
+void cv::gpu::CudaMem::create(int /*_rows*/, int /*_cols*/, int /*_type*/, int /*type_alloc*/) { throw_nogpu(); }
+bool cv::gpu::CudaMem::canMapHostMemory() { throw_nogpu(); return false; }
+void cv::gpu::CudaMem::release() { throw_nogpu(); }
+GpuMat cv::gpu::CudaMem::createGpuMatHeader () const { throw_nogpu(); return GpuMat(); }
+
+#else /* !defined (HAVE_CUDA) */
+#include <cuda_runtime_api.h>
+namespace
+{
+#if defined(__GNUC__)
+    #define cudaSafeCall(expr)  ___cudaSafeCall(expr, __FILE__, __LINE__, __func__)
+    #define nppSafeCall(expr)  ___nppSafeCall(expr, __FILE__, __LINE__, __func__)
+#else /* defined(__CUDACC__) || defined(__MSVC__) */
+    #define cudaSafeCall(expr)  ___cudaSafeCall(expr, __FILE__, __LINE__)
+    #define nppSafeCall(expr)  ___nppSafeCall(expr, __FILE__, __LINE__)
+#endif
+
+    inline void ___cudaSafeCall(cudaError_t err, const char *file, const int line, const char *func = "")
+    {
+        if (cudaSuccess != err)
+            cv::gpu::error(cudaGetErrorString(err), file, line, func);
+    }
+
+    inline void ___nppSafeCall(int err, const char *file, const int line, const char *func = "")
+    {
+        if (err < 0)
+        {
+            std::ostringstream msg;
+            msg << "NPP API Call Error: " << err;
+            cv::gpu::error(msg.str().c_str(), file, line, func);
+        }
+    }
+}
+
+void cv::gpu::registerPageLocked(Mat& m)
+{
+    cudaSafeCall( cudaHostRegister(m.ptr(), m.step * m.rows, cudaHostRegisterPortable) );
+}
+
+void cv::gpu::unregisterPageLocked(Mat& m)
+{
+    cudaSafeCall( cudaHostUnregister(m.ptr()) );
+}
+
+bool cv::gpu::CudaMem::canMapHostMemory()
+{
+    cudaDeviceProp prop;
+    cudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) );
+    return (prop.canMapHostMemory != 0) ? true : false;
+}
+
+namespace
+{
+    size_t alignUpStep(size_t what, size_t alignment)
+    {
+        size_t alignMask = alignment-1;
+        size_t inverseAlignMask = ~alignMask;
+        size_t res = (what + alignMask) & inverseAlignMask;
+        return res;
+    }
+}
+
+void cv::gpu::CudaMem::create(int _rows, int _cols, int _type, int _alloc_type)
+{
+    if (_alloc_type == ALLOC_ZEROCOPY && !canMapHostMemory())
+            cv::gpu::error("ZeroCopy is not supported by current device", __FILE__, __LINE__);
+
+    _type &= TYPE_MASK;
+    if( rows == _rows && cols == _cols && type() == _type && data )
+        return;
+    if( data )
+        release();
+    CV_DbgAssert( _rows >= 0 && _cols >= 0 );
+    if( _rows > 0 && _cols > 0 )
+    {
+        flags = Mat::MAGIC_VAL + Mat::CONTINUOUS_FLAG + _type;
+        rows = _rows;
+        cols = _cols;
+        step = elemSize()*cols;
+        if (_alloc_type == ALLOC_ZEROCOPY)
+        {
+            cudaDeviceProp prop;
+            cudaSafeCall( cudaGetDeviceProperties(&prop, getDevice()) );
+            step = alignUpStep(step, prop.textureAlignment);
+        }
+        int64 _nettosize = (int64)step*rows;
+        size_t nettosize = (size_t)_nettosize;
+        if( _nettosize != (int64)nettosize )
+            CV_Error(CV_StsNoMem, "Too big buffer is allocated");
+        size_t datasize = alignSize(nettosize, (int)sizeof(*refcount));
+
+        //datastart = data = (uchar*)fastMalloc(datasize + sizeof(*refcount));
+        alloc_type = _alloc_type;
+        void *ptr;
+
+        switch (alloc_type)
+        {
+            case ALLOC_PAGE_LOCKED:    cudaSafeCall( cudaHostAlloc( &ptr, datasize, cudaHostAllocDefault) ); break;
+            case ALLOC_ZEROCOPY:       cudaSafeCall( cudaHostAlloc( &ptr, datasize, cudaHostAllocMapped) );  break;
+            case ALLOC_WRITE_COMBINED: cudaSafeCall( cudaHostAlloc( &ptr, datasize, cudaHostAllocWriteCombined) ); break;
+            default: cv::gpu::error("Invalid alloc type", __FILE__, __LINE__);
+        }
+
+        datastart = data =  (uchar*)ptr;
+        dataend = data + nettosize;
+
+        refcount = (int*)cv::fastMalloc(sizeof(*refcount));
+        *refcount = 1;
+    }
+}
+
+GpuMat cv::gpu::CudaMem::createGpuMatHeader () const
+{
+    GpuMat res;
+    if (alloc_type == ALLOC_ZEROCOPY)
+    {
+        void *pdev;
+        cudaSafeCall( cudaHostGetDevicePointer( &pdev, data, 0 ) );
+        res = GpuMat(rows, cols, type(), pdev, step);
+    }
+    else
+        cv::gpu::error("Zero-copy is not supported or memory was allocated without zero-copy flag", __FILE__, __LINE__);
+
+    return res;
+}
+
+void cv::gpu::CudaMem::release()
+{
+    if( refcount && CV_XADD(refcount, -1) == 1 )
+    {
+        cudaSafeCall( cudaFreeHost(datastart ) );
+        fastFree(refcount);
+    }
+    data = datastart = dataend = 0;
+    step = rows = cols = 0;
+    refcount = 0;
+}
+
+#endif /* !defined (HAVE_CUDA) */