From ca35ed2f1c4d7956f2a84d4543f1a9cb5188fb61 Mon Sep 17 00:00:00 2001 From: cudawarped <12133430+cudawarped@users.noreply.github.com> Date: Sat, 23 Aug 2025 09:24:05 +0300 Subject: [PATCH] cuda: add compatibility layer for depreciated vector types --- .../include/opencv2/core/cuda/cuda_compat.hpp | 38 +++++++++++++++++++ .../include/opencv2/core/cuda/vec_math.hpp | 3 ++ .../include/opencv2/core/cuda/vec_traits.hpp | 4 ++ modules/core/src/cuda/gpu_mat.cu | 8 ++-- modules/dnn/src/cuda/vector_traits.hpp | 5 ++- 5 files changed, 54 insertions(+), 4 deletions(-) create mode 100644 modules/core/include/opencv2/core/cuda/cuda_compat.hpp diff --git a/modules/core/include/opencv2/core/cuda/cuda_compat.hpp b/modules/core/include/opencv2/core/cuda/cuda_compat.hpp new file mode 100644 index 0000000000..b40b2ea4f9 --- /dev/null +++ b/modules/core/include/opencv2/core/cuda/cuda_compat.hpp @@ -0,0 +1,38 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. + +#ifndef OPENCV_CUDA_CUDA_COMPAT_HPP +#define OPENCV_CUDA_CUDA_COMPAT_HPP + +#include + +namespace cv { namespace cuda { namespace device { namespace compat +{ +#if CUDA_VERSION >= 13000 + using ulonglong4 = ::ulonglong4_16a; + using double4 = ::double4_16a; + __host__ __device__ __forceinline__ + double4 make_double4(double x, double y, double z, double w) + { + return ::make_double4_16a(x, y, z, w); + } +#else + using ulonglong4 = ::ulonglong4; + using double4 = ::double4; + __host__ __device__ __forceinline__ + double4 make_double4(double x, double y, double z, double w) + { + return ::make_double4(x, y, z, w); + } +#endif + using ulonglong4Compat = ulonglong4; + using double4Compat = double4; + __host__ __device__ __forceinline__ + double4Compat make_double4_compat(double x, double y, double z, double w) + { + return make_double4(x, y, z, w); + } +}}}} + +#endif // OPENCV_CUDA_CUDA_COMPAT_HPP \ No newline at end of file diff --git a/modules/core/include/opencv2/core/cuda/vec_math.hpp b/modules/core/include/opencv2/core/cuda/vec_math.hpp index 80b1303681..afb8cc0c4c 100644 --- a/modules/core/include/opencv2/core/cuda/vec_math.hpp +++ b/modules/core/include/opencv2/core/cuda/vec_math.hpp @@ -45,6 +45,7 @@ #include "vec_traits.hpp" #include "saturate_cast.hpp" +#include "cuda_compat.hpp" /** @file * @deprecated Use @ref cudev instead. @@ -54,6 +55,8 @@ namespace cv { namespace cuda { namespace device { + using cv::cuda::device::compat::double4; + using cv::cuda::device::compat::make_double4; // saturate_cast diff --git a/modules/core/include/opencv2/core/cuda/vec_traits.hpp b/modules/core/include/opencv2/core/cuda/vec_traits.hpp index b5ff281a0b..786a9314a7 100644 --- a/modules/core/include/opencv2/core/cuda/vec_traits.hpp +++ b/modules/core/include/opencv2/core/cuda/vec_traits.hpp @@ -44,6 +44,7 @@ #define OPENCV_CUDA_VEC_TRAITS_HPP #include "common.hpp" +#include "cuda_compat.hpp" /** @file * @deprecated Use @ref cudev instead. @@ -53,6 +54,9 @@ namespace cv { namespace cuda { namespace device { + using cv::cuda::device::compat::double4; + using cv::cuda::device::compat::make_double4; + template struct TypeVec; struct __align__(8) uchar8 diff --git a/modules/core/src/cuda/gpu_mat.cu b/modules/core/src/cuda/gpu_mat.cu index b6f95445db..c70b9fc96f 100644 --- a/modules/core/src/cuda/gpu_mat.cu +++ b/modules/core/src/cuda/gpu_mat.cu @@ -51,10 +51,12 @@ #include "opencv2/core/cuda.hpp" #include "opencv2/cudev.hpp" #include "opencv2/core/cuda/utility.hpp" +#include "opencv2/core/cuda/cuda_compat.hpp" using namespace cv; using namespace cv::cuda; using namespace cv::cudev; +using cv::cuda::device::compat::double4Compat; device::ThrustAllocator::~ThrustAllocator() { @@ -341,7 +343,7 @@ void cv::cuda::GpuMat::copyTo(OutputArray _dst, InputArray _mask, Stream& stream {0,0,0,0}, {0,0,0,0}, {0,0,0,0}, - {copyWithMask, copyWithMask, copyWithMask, copyWithMask} + {copyWithMask, copyWithMask, copyWithMask, copyWithMask} }; if (mask.channels() == channels()) @@ -424,7 +426,7 @@ GpuMat& cv::cuda::GpuMat::setTo(Scalar value, Stream& stream) {setToWithOutMask,setToWithOutMask,setToWithOutMask,setToWithOutMask}, {setToWithOutMask,setToWithOutMask,setToWithOutMask,setToWithOutMask}, {setToWithOutMask,setToWithOutMask,setToWithOutMask,setToWithOutMask}, - {setToWithOutMask,setToWithOutMask,setToWithOutMask,setToWithOutMask} + {setToWithOutMask,setToWithOutMask,setToWithOutMask,setToWithOutMask} }; funcs[depth()][channels() - 1](*this, value, stream); @@ -455,7 +457,7 @@ GpuMat& cv::cuda::GpuMat::setTo(Scalar value, InputArray _mask, Stream& stream) {setToWithMask,setToWithMask,setToWithMask,setToWithMask}, {setToWithMask,setToWithMask,setToWithMask,setToWithMask}, {setToWithMask,setToWithMask,setToWithMask,setToWithMask}, - {setToWithMask,setToWithMask,setToWithMask,setToWithMask} + {setToWithMask,setToWithMask,setToWithMask,setToWithMask} }; funcs[depth()][channels() - 1](*this, mask, value, stream); diff --git a/modules/dnn/src/cuda/vector_traits.hpp b/modules/dnn/src/cuda/vector_traits.hpp index 1b9b76980c..d7d2214035 100644 --- a/modules/dnn/src/cuda/vector_traits.hpp +++ b/modules/dnn/src/cuda/vector_traits.hpp @@ -11,6 +11,7 @@ #include "memory.hpp" #include "../cuda4dnn/csl/pointer.hpp" +#include "opencv2/core/cuda/cuda_compat.hpp" #include @@ -34,9 +35,11 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace de * v_store(output_vPtr, vec); */ + using cv::cuda::device::compat::ulonglong4Compat; + namespace detail { template struct raw_type_ { }; - template <> struct raw_type_<256> { typedef ulonglong4 type; }; + template <> struct raw_type_<256> { typedef ulonglong4Compat type; }; template <> struct raw_type_<128> { typedef uint4 type; }; template <> struct raw_type_<64> { typedef uint2 type; }; template <> struct raw_type_<32> { typedef uint1 type; };