From e5188c7e94043846765a0c0b1c78407401211a81 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Tue, 1 Oct 2013 18:17:46 +0400 Subject: [PATCH] added GpuMat::Allocator interface --- modules/core/include/opencv2/core/cuda.hpp | 28 +++++-- .../core/include/opencv2/core/cuda.inl.hpp | 26 +++---- modules/core/src/cuda/gpu_mat.cu | 73 +++++++++++++++---- modules/core/src/cuda_gpu_mat.cpp | 22 +++++- 4 files changed, 113 insertions(+), 36 deletions(-) diff --git a/modules/core/include/opencv2/core/cuda.hpp b/modules/core/include/opencv2/core/cuda.hpp index 6bca5413f9..90c21cf64c 100644 --- a/modules/core/include/opencv2/core/cuda.hpp +++ b/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 diff --git a/modules/core/include/opencv2/core/cuda.inl.hpp b/modules/core/include/opencv2/core/cuda.inl.hpp index 7410074b5e..170d0affb3 100644 --- a/modules/core/include/opencv2/core/cuda.inl.hpp +++ b/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); } diff --git a/modules/core/src/cuda/gpu_mat.cu b/modules/core/src/cuda/gpu_mat.cu index d9332b7332..265f3c96c5 100644 --- a/modules/core/src/cuda/gpu_mat.cu +++ b/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(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; diff --git a/modules/core/src/cuda_gpu_mat.cpp b/modules/core/src/cuda_gpu_mat.cpp index 2303f11ee8..80a7462cb4 100644 --- a/modules/core/src/cuda_gpu_mat.cpp +++ b/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;