core cuda changes

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<int>;
+    using SizeArray = MatShape;
     using StepArray = std::vector<size_t>;
     using IndexArray = std::vector<int>;
 
@@ -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.

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

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<int>(tensor->dl_tensor.shape[i]);
     }
     sizes.back() = static_cast<int>(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;

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<GpuData>(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

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)

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<uchar*>(_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<Range>& 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

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<cuda::GpuMatND>& vv = *(const std::vector<cuda::GpuMatND>*)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