added GpuMat::Allocator interface

modules/core/include/opencv2/core/cuda.hpp

@@ -61,16 +61,29 @@ namespace cv { namespace cuda {
 class CV_EXPORTS GpuMat
 {
 public:
+    class CV_EXPORTS Allocator
+    {
+    public:
+        virtual ~Allocator() {}
+
+        virtual bool allocate(uchar** devPtr, size_t* step, int** refcount, int rows, int cols, size_t elemSize) = 0;
+        virtual void free(uchar* devPtr, int* refcount) = 0;
+    };
+
+    //! default allocator
+    static Allocator* defaultAllocator();
+    static void setDefaultAllocator(Allocator* allocator);
+
     //! default constructor
-    GpuMat();
+    explicit GpuMat(Allocator* allocator = defaultAllocator());
 
     //! constructs GpuMat of the specified size and type
-    GpuMat(int rows, int cols, int type);
-    GpuMat(Size size, int type);
+    GpuMat(int rows, int cols, int type, Allocator* allocator = defaultAllocator());
+    GpuMat(Size size, int type, Allocator* allocator = defaultAllocator());
 
     //! constucts GpuMat and fills it with the specified value _s
-    GpuMat(int rows, int cols, int type, Scalar s);
-    GpuMat(Size size, int type, Scalar s);
+    GpuMat(int rows, int cols, int type, Scalar s, Allocator* allocator = defaultAllocator());
+    GpuMat(Size size, int type, Scalar s, Allocator* allocator = defaultAllocator());
 
     //! copy constructor
     GpuMat(const GpuMat& m);
@@ -84,7 +97,7 @@ public:
     GpuMat(const GpuMat& m, Rect roi);
 
     //! builds GpuMat from host memory (Blocking call)
-    explicit GpuMat(InputArray arr);
+    explicit GpuMat(InputArray arr, Allocator* allocator = defaultAllocator());
 
     //! destructor - calls release()
     ~GpuMat();
@@ -249,6 +262,9 @@ public:
     //! helper fields used in locateROI and adjustROI
     uchar* datastart;
     uchar* dataend;
+
+    //! allocator
+    Allocator* allocator;
 };
 
 //! creates continuous matrix

modules/core/include/opencv2/core/cuda.inl.hpp

@@ -51,29 +51,29 @@ namespace cv { namespace cuda {
 //////////////////////////////// GpuMat ///////////////////////////////
 
 inline
-GpuMat::GpuMat()
-    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
+GpuMat::GpuMat(Allocator* allocator_)
+    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), allocator(allocator_)
 {}
 
 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)
+GpuMat::GpuMat(int rows_, int cols_, int type_, Allocator* allocator_)
+    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), allocator(allocator_)
 {
     if (rows_ > 0 && cols_ > 0)
         create(rows_, cols_, type_);
 }
 
 inline
-GpuMat::GpuMat(Size size_, int type_)
-    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
+GpuMat::GpuMat(Size size_, int type_, Allocator* allocator_)
+    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), allocator(allocator_)
 {
     if (size_.height > 0 && size_.width > 0)
         create(size_.height, size_.width, type_);
 }
 
 inline
-GpuMat::GpuMat(int rows_, int cols_, int type_, Scalar s_)
-    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
+GpuMat::GpuMat(int rows_, int cols_, int type_, Scalar s_, Allocator* allocator_)
+    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), allocator(allocator_)
 {
     if (rows_ > 0 && cols_ > 0)
     {
@@ -83,8 +83,8 @@ GpuMat::GpuMat(int rows_, int cols_, int type_, Scalar s_)
 }
 
 inline
-GpuMat::GpuMat(Size size_, int type_, Scalar s_)
-    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
+GpuMat::GpuMat(Size size_, int type_, Scalar s_, Allocator* allocator_)
+    : flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), allocator(allocator_)
 {
     if (size_.height > 0 && size_.width > 0)
     {
@@ -95,15 +95,15 @@ GpuMat::GpuMat(Size size_, int type_, Scalar s_)
 
 inline
 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)
+    : flags(m.flags), rows(m.rows), cols(m.cols), step(m.step), data(m.data), refcount(m.refcount), datastart(m.datastart), dataend(m.dataend), allocator(m.allocator)
 {
     if (refcount)
         CV_XADD(refcount, 1);
 }
 
 inline
-GpuMat::GpuMat(InputArray arr) :
-    flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0)
+GpuMat::GpuMat(InputArray arr, Allocator* allocator_) :
+    flags(0), rows(0), cols(0), step(0), data(0), refcount(0), datastart(0), dataend(0), allocator(allocator_)
 {
     upload(arr);
 }

modules/core/src/cuda/gpu_mat.cu

@@ -55,6 +55,54 @@ using namespace cv;
 using namespace cv::cuda;
 using namespace cv::cudev;
 
+namespace
+{
+    class DefaultAllocator : public GpuMat::Allocator
+    {
+    public:
+        bool allocate(uchar** devPtr, size_t* step, int** refcount, int rows, int cols, size_t elemSize);
+        void free(uchar* devPtr, int* refcount);
+    };
+
+    bool DefaultAllocator::allocate(uchar** devPtr, size_t* step, int** refcount, int rows, int cols, size_t elemSize)
+    {
+        if (rows > 1 && cols > 1)
+        {
+            CV_CUDEV_SAFE_CALL( cudaMallocPitch(devPtr, step, elemSize * cols, rows) );
+        }
+        else
+        {
+            // Single row or single column must be continuous
+            CV_CUDEV_SAFE_CALL( cudaMalloc(devPtr, elemSize * cols * rows) );
+            *step = elemSize * cols;
+        }
+
+        *refcount = static_cast<int*>(fastMalloc(sizeof(int)));
+
+        return true;
+    }
+
+    void DefaultAllocator::free(uchar* devPtr, int* refcount)
+    {
+        cudaFree(devPtr);
+        fastFree(refcount);
+    }
+
+    DefaultAllocator cudaDefaultAllocator;
+    GpuMat::Allocator* g_defaultAllocator = &cudaDefaultAllocator;
+}
+
+GpuMat::Allocator* cv::cuda::GpuMat::defaultAllocator()
+{
+    return g_defaultAllocator;
+}
+
+void cv::cuda::GpuMat::setDefaultAllocator(Allocator* allocator)
+{
+    CV_Assert( allocator != 0 );
+    g_defaultAllocator = allocator;
+}
+
 /////////////////////////////////////////////////////
 /// create
 
@@ -76,19 +124,17 @@ void cv::cuda::GpuMat::create(int _rows, int _cols, int _type)
         rows = _rows;
         cols = _cols;
 
-        size_t esz = elemSize();
+        uchar* devPtr;
+        const size_t esz = elemSize();
 
-        void* devPtr;
+        bool allocSuccess = allocator->allocate(&devPtr, &step, &refcount, rows, cols, esz);
 
-        if (rows > 1 && cols > 1)
+        if (!allocSuccess)
         {
-            CV_CUDEV_SAFE_CALL( cudaMallocPitch(&devPtr, &step, esz * cols, rows) );
-        }
-        else
-        {
-            // Single row or single column must be continuous
-            CV_CUDEV_SAFE_CALL( cudaMalloc(&devPtr, esz * cols * rows) );
-            step = esz * cols;
+            // custom allocator fails, try default allocator
+            allocator = defaultAllocator();
+            allocSuccess = allocator->allocate(&devPtr, &step, &refcount, rows, cols, esz);
+            CV_Assert( allocSuccess );
         }
 
         if (esz * cols == step)
@@ -110,11 +156,10 @@ void cv::cuda::GpuMat::create(int _rows, int _cols, int _type)
 
 void cv::cuda::GpuMat::release()
 {
+    CV_DbgAssert( allocator != 0 );
+
     if (refcount && CV_XADD(refcount, -1) == 1)
-    {
-        cudaFree(datastart);
-        fastFree(refcount);
-    }
+        allocator->free(datastart, refcount);
 
     data = datastart = dataend = 0;
     step = rows = cols = 0;

modules/core/src/cuda_gpu_mat.cpp

@@ -49,7 +49,8 @@ using namespace cv::cuda;
 cv::cuda::GpuMat::GpuMat(int rows_, int cols_, int type_, void* data_, size_t step_) :
     flags(Mat::MAGIC_VAL + (type_ & Mat::TYPE_MASK)), rows(rows_), cols(cols_),
     step(step_), data((uchar*)data_), refcount(0),
-    datastart((uchar*)data_), dataend((uchar*)data_)
+    datastart((uchar*)data_), dataend((uchar*)data_),
+    allocator(defaultAllocator())
 {
     size_t minstep = cols * elemSize();
 
@@ -74,7 +75,8 @@ cv::cuda::GpuMat::GpuMat(int rows_, int cols_, int type_, void* data_, size_t st
 cv::cuda::GpuMat::GpuMat(Size size_, int type_, void* data_, size_t step_) :
     flags(Mat::MAGIC_VAL + (type_ & Mat::TYPE_MASK)), rows(size_.height), cols(size_.width),
     step(step_), data((uchar*)data_), refcount(0),
-    datastart((uchar*)data_), dataend((uchar*)data_)
+    datastart((uchar*)data_), dataend((uchar*)data_),
+    allocator(defaultAllocator())
 {
     size_t minstep = cols * elemSize();
 
@@ -92,6 +94,7 @@ cv::cuda::GpuMat::GpuMat(Size size_, int type_, void* data_, size_t step_) :
 
         flags |= step == minstep ? Mat::CONTINUOUS_FLAG : 0;
     }
+
     dataend += step * (rows - 1) + minstep;
 }
 
@@ -100,6 +103,7 @@ cv::cuda::GpuMat::GpuMat(const GpuMat& m, Range rowRange_, Range colRange_)
     flags = m.flags;
     step = m.step; refcount = m.refcount;
     data = m.data; datastart = m.datastart; dataend = m.dataend;
+    allocator = m.allocator;
 
     if (rowRange_ == Range::all())
     {
@@ -139,7 +143,8 @@ cv::cuda::GpuMat::GpuMat(const GpuMat& m, Range rowRange_, Range colRange_)
 cv::cuda::GpuMat::GpuMat(const GpuMat& m, Rect roi) :
     flags(m.flags), rows(roi.height), cols(roi.width),
     step(m.step), data(m.data + roi.y*step), refcount(m.refcount),
-    datastart(m.datastart), dataend(m.dataend)
+    datastart(m.datastart), dataend(m.dataend),
+    allocator(m.allocator)
 {
     flags &= roi.width < m.cols ? ~Mat::CONTINUOUS_FLAG : -1;
     data += roi.x * elemSize();
@@ -347,6 +352,17 @@ GpuMat cv::cuda::allocMatFromBuf(int rows, int cols, int type, GpuMat& mat)
 
 #ifndef HAVE_CUDA
 
+GpuMat::Allocator* cv::cuda::GpuMat::defaultAllocator()
+{
+    return 0;
+}
+
+void cv::cuda::GpuMat::setDefaultAllocator(Allocator* allocator)
+{
+    (void) allocator;
+    throw_no_cuda();
+}
+
 void cv::cuda::GpuMat::create(int _rows, int _cols, int _type)
 {
     (void) _rows;