From 2f89729f2dd03e463661d0341976b3c5f2593734 Mon Sep 17 00:00:00 2001 From: Abhishek Gola Date: Wed, 20 May 2026 17:29:34 +0530 Subject: [PATCH] core cuda changes --- modules/core/include/opencv2/core/cuda.hpp | 28 ++++++++-- .../core/include/opencv2/core/cuda.inl.hpp | 4 +- modules/core/misc/python/pyopencv_cuda.hpp | 2 +- modules/core/src/cuda/gpu_mat_nd.cu | 56 ++++++++++++++----- modules/core/src/cuda_gpu_mat.cpp | 35 ++++++++++++ modules/core/src/cuda_gpu_mat_nd.cpp | 43 +++++++++++--- modules/core/src/matrix_wrap.cpp | 20 ++++++- 7 files changed, 156 insertions(+), 32 deletions(-) diff --git a/modules/core/include/opencv2/core/cuda.hpp b/modules/core/include/opencv2/core/cuda.hpp index 03f8dfdc67..c9b4b59ab9 100644 --- a/modules/core/include/opencv2/core/cuda.hpp +++ b/modules/core/include/opencv2/core/cuda.hpp @@ -158,6 +158,10 @@ public: //! decreases reference counter, deallocate the data when reference counter reaches 0 CV_WRAP void release(); + //! allocates or reuses underlying storage to fit the requested 2D size and type (no-op if already compatible) + void fit(int rows, int cols, int type); + void fit(Size size, int type); + //! swaps with other smart pointer CV_WRAP void swap(GpuMat& mat); @@ -394,7 +398,7 @@ struct CV_EXPORTS_W GpuData class CV_EXPORTS_W GpuMatND { public: - using SizeArray = std::vector; + using SizeArray = MatShape; using StepArray = std::vector; using IndexArray = std::vector; @@ -409,7 +413,7 @@ public: @param type Array type. Use CV_8UC1, ..., CV_16FC4 to create 1-4 channel matrices, or CV_8UC(n), ..., CV_64FC(n) to create multi-channel (up to CV_CN_MAX channels) matrices. */ - GpuMatND(SizeArray size, int type); + GpuMatND(const MatShape& shape, int type); /** @overload @param size Array of integers specifying an n-dimensional array shape. @@ -424,7 +428,7 @@ public: (if specified, the last step must be equal to the element size, otherwise it will be added as such). If not specified, the matrix is assumed to be continuous. */ - GpuMatND(SizeArray size, int type, void* data, StepArray step = StepArray()); + GpuMatND(const MatShape& shape, int type, void* data, StepArray step = StepArray()); /** @brief Allocates GPU memory. Suppose there is some GPU memory already allocated. In that case, this method may choose to reuse that @@ -433,12 +437,19 @@ public: (i.e., isSubmatrix() is false). In other words, this method guarantees that the GPU memory allocated by this method is always continuous and is not a sub-region of another GpuMatND. */ - void create(SizeArray size, int type); + void create(const MatShape& shape, int type); void release(); void swap(GpuMatND& m) noexcept; + /** @brief Allocates or reuses underlying storage to fit the requested n-D size and type. + - No-op if already compatible (sufficient capacity, continuous, not a submatrix, not external, no ROI). + - Reallocates otherwise. + Mirrors 2D GpuMat::fit semantics for multi-dimensional tensors. + */ + void fit(const MatShape& shape, int type); + /** @brief Creates a full copy of the array and the underlying data. The method creates a full copy of the array. It mimics the behavior of Mat::clone(), i.e. the original step is not taken into account. So, the array copy is a continuous array @@ -538,7 +549,7 @@ public: private: //! internal use - void setFields(SizeArray size, int type, StepArray step = StepArray()); + void setFields(MatShape size, int type, StepArray step = StepArray()); public: /*! includes several bit-fields: @@ -553,7 +564,7 @@ public: int dims; //! shape of this array - SizeArray size; + MatShape size; /*! step values Their semantics is identical to the semantics of step for Mat. @@ -580,6 +591,11 @@ private: size_t offset; }; +// Provide namespace-level aliases for nested array types used in bindings +using SizeArray = GpuMatND::SizeArray; +using StepArray = GpuMatND::StepArray; +using IndexArray = GpuMatND::IndexArray; + /** @brief Creates a continuous matrix. @param rows Row count. diff --git a/modules/core/include/opencv2/core/cuda.inl.hpp b/modules/core/include/opencv2/core/cuda.inl.hpp index 237ef21052..be5d7622a7 100644 --- a/modules/core/include/opencv2/core/cuda.inl.hpp +++ b/modules/core/include/opencv2/core/cuda.inl.hpp @@ -418,10 +418,10 @@ GpuMatND::GpuMatND() : } inline -GpuMatND::GpuMatND(SizeArray _size, int _type) : +GpuMatND::GpuMatND(const MatShape& _shape, int _type) : flags(0), dims(0), data(nullptr), offset(0) { - create(std::move(_size), _type); + create(_shape, _type); } inline diff --git a/modules/core/misc/python/pyopencv_cuda.hpp b/modules/core/misc/python/pyopencv_cuda.hpp index 2d6b284d95..ab451b0daa 100644 --- a/modules/core/misc/python/pyopencv_cuda.hpp +++ b/modules/core/misc/python/pyopencv_cuda.hpp @@ -144,7 +144,7 @@ bool parseDLPackTensor(DLManagedTensor* tensor, cv::cuda::GpuMatND& obj, bool co sizes[i] = static_cast(tensor->dl_tensor.shape[i]); } sizes.back() = static_cast(tensor->dl_tensor.shape[tensor->dl_tensor.ndim - 1]); - obj = cv::cuda::GpuMatND(sizes, type, tensor->dl_tensor.data, steps); + obj = cv::cuda::GpuMatND(cv::MatShape(sizes.begin(), sizes.end()), type, tensor->dl_tensor.data, steps); if (copy) obj = obj.clone(); return true; diff --git a/modules/core/src/cuda/gpu_mat_nd.cu b/modules/core/src/cuda/gpu_mat_nd.cu index 3f51fd8afa..58103b0b53 100644 --- a/modules/core/src/cuda/gpu_mat_nd.cu +++ b/modules/core/src/cuda/gpu_mat_nd.cu @@ -30,26 +30,25 @@ GpuData::~GpuData() ///////////////////////////////////////////////////// /// create -void GpuMatND::create(SizeArray _size, int _type) +void GpuMatND::create(const MatShape& _shape, int _type) { { - auto elements_nonzero = [](SizeArray& v) + auto elements_nonzero = [](const MatShape& v) { - return std::all_of(v.begin(), v.end(), - [](unsigned u){ return u > 0; }); + return std::all_of(v.begin(), v.end(), [](int u){ return u > 0; }); }; - CV_Assert(!_size.empty()); - CV_Assert(elements_nonzero(_size)); + CV_Assert(!_shape.empty()); + CV_Assert(elements_nonzero(_shape)); } _type &= Mat::TYPE_MASK; - if (size == _size && type() == _type && !empty() && !external() && isContinuous() && !isSubmatrix()) + if (size == _shape && type() == _type && !empty() && !external() && isContinuous() && !isSubmatrix()) return; release(); - setFields(std::move(_size), _type); + setFields(_shape, _type); data_ = std::make_shared(totalMemSize()); data = data_->data; @@ -206,10 +205,9 @@ void GpuMatND::upload(InputArray src) if (!mat.isContinuous()) mat = mat.clone(); - SizeArray _size(mat.dims); - std::copy_n(mat.size.p, mat.dims, _size.data()); + MatShape _shape(mat.size.p, mat.size.p + mat.dims); - create(std::move(_size), mat.type()); + create(_shape, mat.type()); CV_CUDEV_SAFE_CALL(cudaMemcpy(getDevicePtr(), mat.data, totalMemSize(), cudaMemcpyHostToDevice)); } @@ -223,10 +221,9 @@ void GpuMatND::upload(InputArray src, Stream& stream) if (!mat.isContinuous()) mat = mat.clone(); - SizeArray _size(mat.dims); - std::copy_n(mat.size.p, mat.dims, _size.data()); + MatShape _shape(mat.size.p, mat.size.p + mat.dims); - create(std::move(_size), mat.type()); + create(_shape, mat.type()); cudaStream_t _stream = StreamAccessor::getStream(stream); CV_CUDEV_SAFE_CALL(cudaMemcpyAsync(getDevicePtr(), mat.data, totalMemSize(), cudaMemcpyHostToDevice, _stream)); @@ -266,4 +263,35 @@ void GpuMatND::download(OutputArray dst, Stream& stream) const CV_CUDEV_SAFE_CALL(cudaMemcpyAsync(mat.data, gmat.getDevicePtr(), mat.total() * mat.elemSize(), cudaMemcpyDeviceToHost, _stream)); } +void GpuMatND::fit(const MatShape& _shape, int _type) +{ + { + auto elements_nonzero = [](const MatShape& v) + { + return std::all_of(v.begin(), v.end(), [](int u){ return u > 0; }); + }; + CV_Assert(!_shape.empty()); + CV_Assert(elements_nonzero(_shape)); + } + + _type &= Mat::TYPE_MASK; + + const size_t esz = (size_t)CV_ELEM_SIZE(_type); + size_t step0 = esz; + for (int i = (int)_shape.size() - 2; i >= 0; --i) + step0 *= (size_t)_shape[(size_t)i + 1]; + const size_t requiredBytes = (size_t)_shape[0] * step0; + + const bool canReuse = !empty() && !external() && isContinuous() && !isSubmatrix() && offset == 0; + const size_t oldCapacity = canReuse ? totalMemSize() : 0; + + if (!canReuse || requiredBytes > oldCapacity) + { + create(_shape, _type); + return; + } + + setFields(_shape, _type); +} + #endif diff --git a/modules/core/src/cuda_gpu_mat.cpp b/modules/core/src/cuda_gpu_mat.cpp index 84b5b210d1..1a4ffb63b6 100644 --- a/modules/core/src/cuda_gpu_mat.cpp +++ b/modules/core/src/cuda_gpu_mat.cpp @@ -435,6 +435,41 @@ void cv::cuda::GpuMat::create(int _rows, int _cols, int _type) void cv::cuda::GpuMat::release() { + if( refcount && CV_XADD(refcount, -1) == 1 ) + defaultAllocator()->free(this); + dataend = datastart = data = 0; + step = 0; + refcount = 0; + flags &= ~Mat::MAGIC_MASK; +} + +void cv::cuda::GpuMat::fit(int _rows, int _cols, int _type) +{ + // Emulate Mat::fit semantics for GPU: avoid reallocation when possible + size_t oldTotalBytes = (datastart && dataend) ? (size_t)(dataend - datastart) : 0; + size_t esz = CV_ELEM_SIZE(_type); + size_t newTotalBytes = (size_t)_rows * (size_t)_cols * esz; + + // Reallocate if: + // - need non-contiguous layout (GPU pitch) or ROI, or + // - requested bytes exceed current capacity, or + // - data pointer is not at the start (ROI) + if (newTotalBytes > 0 && (!isContinuous() || newTotalBytes > oldTotalBytes || data != datastart)) + { + create(_rows, _cols, _type); + return; + } + + // Reuse existing allocation, just adjust header/type + flags = (flags & ~Mat::TYPE_MASK) | CV_MAKETYPE(CV_MAT_DEPTH(_type), CV_MAT_CN(_type)); + rows = _rows; + cols = _cols; + step = (size_t)cols * esz; +} + +void cv::cuda::GpuMat::fit(Size _sz, int _type) +{ + fit(_sz.height, _sz.width, _type); } void cv::cuda::GpuMat::upload(InputArray arr) diff --git a/modules/core/src/cuda_gpu_mat_nd.cpp b/modules/core/src/cuda_gpu_mat_nd.cpp index d810ba2854..99e506311b 100644 --- a/modules/core/src/cuda_gpu_mat_nd.cpp +++ b/modules/core/src/cuda_gpu_mat_nd.cpp @@ -9,13 +9,13 @@ using namespace cv::cuda; GpuMatND::~GpuMatND() = default; -GpuMatND::GpuMatND(SizeArray _size, int _type, void* _data, StepArray _step) : +GpuMatND::GpuMatND(const MatShape& _shape, int _type, void* _data, StepArray _step) : flags(0), dims(0), data(static_cast(_data)), offset(0) { - CV_Assert(_step.empty() || _size.size() == _step.size() + 1 || - (_size.size() == _step.size() && _step.back() == (size_t)CV_ELEM_SIZE(_type))); + CV_Assert(_step.empty() || _shape.size() == _step.size() + 1 || + (_shape.size() == _step.size() && _step.back() == (size_t)CV_ELEM_SIZE(_type))); - setFields(std::move(_size), _type, std::move(_step)); + setFields(_shape, _type, std::move(_step)); } GpuMatND GpuMatND::operator()(const std::vector& ranges) const @@ -95,7 +95,7 @@ GpuMatND::operator GpuMat() const return createGpuMatHeader().clone(); } -void GpuMatND::setFields(SizeArray _size, int _type, StepArray _step) +void GpuMatND::setFields(MatShape _size, int _type, StepArray _step) { _type &= Mat::TYPE_MASK; @@ -142,9 +142,9 @@ GpuData::~GpuData() { } -void GpuMatND::create(SizeArray _size, int _type) +void GpuMatND::create(const MatShape& _shape, int _type) { - CV_UNUSED(_size); + CV_UNUSED(_shape); CV_UNUSED(_type); throw_no_cuda(); } @@ -191,4 +191,33 @@ void GpuMatND::download(OutputArray dst, Stream& stream) const throw_no_cuda(); } +void GpuMatND::fit(const MatShape& _shape, int _type) +{ + auto elements_nonzero = [](const MatShape& v) + { + return std::all_of(v.begin(), v.end(), [](int u){ return u > 0; }); + }; + CV_Assert(!_shape.empty()); + CV_Assert(elements_nonzero(_shape)); + + _type &= Mat::TYPE_MASK; + + const size_t esz = (size_t)CV_ELEM_SIZE(_type); + size_t step0 = esz; + for (int i = (int)_shape.size() - 2; i >= 0; --i) + step0 *= (size_t)_shape[(size_t)i + 1]; + const size_t requiredBytes = (size_t)_shape[0] * step0; + + const bool canReuse = !empty() && !external() && isContinuous() && !isSubmatrix() && offset == 0; + const size_t oldCapacity = canReuse ? totalMemSize() : 0; + + if (!canReuse || requiredBytes > oldCapacity) + { + create(_shape, _type); + return; + } + + setFields(_shape, _type); +} + #endif diff --git a/modules/core/src/matrix_wrap.cpp b/modules/core/src/matrix_wrap.cpp index 32f43a55ec..be914b0d6d 100644 --- a/modules/core/src/matrix_wrap.cpp +++ b/modules/core/src/matrix_wrap.cpp @@ -1025,6 +1025,22 @@ int _InputArray::type(int i) const #endif } + if (k == STD_VECTOR_CUDA_GPU_MAT_ND) + { +#ifdef HAVE_CUDA + const std::vector& vv = *(const std::vector*)obj; + if (vv.empty()) + { + CV_Assert((flags & FIXED_TYPE) != 0); + return CV_MAT_TYPE(flags); + } + CV_Assert(i < (int)vv.size()); + return vv[i >= 0 ? i : 0].type(); +#else + CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)"); +#endif + } + if( k == OPENGL_BUFFER ) return ((const ogl::Buffer*)obj)->type(); @@ -1438,7 +1454,7 @@ void _OutputArray::create(Size _sz, int mtype, int i, bool allowTransposed, _Out CV_Assert(!fixedSize() || ((((cuda::GpuMatND*)obj)->dims == 2) && (((cuda::GpuMatND*)obj)->size[0] == _sz.height) && (((cuda::GpuMatND*)obj)->size[1] == _sz.width))); CV_Assert(!fixedType() || ((cuda::GpuMatND*)obj)->type() == mtype); #ifdef HAVE_CUDA - cuda::GpuMatND::SizeArray sizes = {_sz.height, _sz.width}; + cuda::GpuMatND::SizeArray sizes({_sz.height, _sz.width}); ((cuda::GpuMatND*)obj)->create(sizes, mtype); return; #else @@ -1504,7 +1520,7 @@ void _OutputArray::create(int _rows, int _cols, int mtype, int i, bool allowTran CV_Assert(!fixedSize() || ((((cuda::GpuMatND*)obj)->dims == 2) && (((cuda::GpuMatND*)obj)->size[0] == _rows) && (((cuda::GpuMatND*)obj)->size[1] == _cols))); CV_Assert(!fixedType() || ((cuda::GpuMatND*)obj)->type() == mtype); #ifdef HAVE_CUDA - cuda::GpuMatND::SizeArray sizes = {_rows, _cols}; + cuda::GpuMatND::SizeArray sizes({_rows, _cols}); ((cuda::GpuMatND*)obj)->create(sizes, mtype); return; #else