Merge branch 4.x

modules/core/include/opencv2/core.hpp

@@ -1965,8 +1965,8 @@ The function solveCubic finds the real roots of a cubic equation:
 
 The roots are stored in the roots array.
 @param coeffs equation coefficients, an array of 3 or 4 elements.
-@param roots output array of real roots that has 1 or 3 elements.
-@return number of real roots. It can be 0, 1 or 2.
+@param roots output array of real roots that has 0, 1, 2 or 3 elements.
+@return number of real roots. It can be -1 (all real numbers), 0, 1, 2 or 3.
 */
 CV_EXPORTS_W int solveCubic(InputArray coeffs, OutputArray roots);
 

modules/core/include/opencv2/core/hal/intrin_cpp.hpp

@@ -225,32 +225,30 @@ These operations allow to reorder or recombine elements in one or multiple vecto
 Element-wise binary and unary operations.
 
 - Arithmetics:
-@ref v_add(const v_reg &a, const v_reg &b) "+",
-@ref v_sub(const v_reg &a, const v_reg &b) "-",
-@ref v_mul(const v_reg &a, const v_reg &b) "*",
-@ref v_div(const v_reg &a, const v_reg &b) "/",
+@ref v_add,
+@ref v_sub,
+@ref v_mul,
+@ref v_div,
 @ref v_mul_expand
 
 - Non-saturating arithmetics: @ref v_add_wrap, @ref v_sub_wrap
 
 - Bitwise shifts:
-@ref v_shl(const v_reg &a, int s) "<<",
-@ref v_shr(const v_reg &a, int s) ">>",
 @ref v_shl, @ref v_shr
 
 - Bitwise logic:
-@ref v_and(const v_reg &a, const v_reg &b) "&",
-@ref v_or(const v_reg &a, const v_reg &b) "|",
-@ref v_xor(const v_reg &a, const v_reg &b) "^",
-@ref v_not(const v_reg &a) "~"
+@ref v_and,
+@ref v_or,
+@ref v_xor,
+@ref v_not
 
 - Comparison:
-@ref v_gt(const v_reg &a, const v_reg &b) ">",
-@ref v_ge(const v_reg &a, const v_reg &b) ">=",
-@ref v_lt(const v_reg &a, const v_reg &b) "<",
-@ref v_le(const v_reg &a, const v_reg &b) "<=",
-@ref v_eq(const v_reg &a, const v_reg &b) "==",
-@ref v_ne(const v_reg &a, const v_reg &b) "!="
+@ref v_gt,
+@ref v_ge,
+@ref v_lt,
+@ref v_le,
+@ref v_eq,
+@ref v_ne
 
 - min/max: @ref v_min, @ref v_max
 

modules/core/include/opencv2/core/hal/intrin_legacy_ops.h

@@ -0,0 +1,111 @@
+// 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
+
+// This file has been created for compatibility with older versions of Universal Intrinscs
+// Binary operators for vector types has been removed since version 4.11
+// Include this file manually after OpenCV headers if you need these operators
+
+#ifndef OPENCV_HAL_INTRIN_LEGACY_OPS_HPP
+#define OPENCV_HAL_INTRIN_LEGACY_OPS_HPP
+
+#ifdef __OPENCV_BUILD
+#error "Universal Intrinsics operators are deprecated and should not be used in OpenCV library"
+#endif
+
+#ifdef __riscv
+#warning "Operators might conflict with built-in functions on RISC-V platform"
+#endif
+
+#if defined(CV_VERSION) && CV_VERSION_MAJOR == 4 && CV_VERSION_MINOR < 9
+#warning "Older versions of OpenCV (<4.9) already have Universal Intrinscs operators"
+#endif
+
+
+namespace cv { namespace hal {
+
+#define BIN_OP(OP, FUN) \
+template <typename R> R operator OP (const R & lhs, const R & rhs) { return FUN(lhs, rhs); }
+
+#define BIN_A_OP(OP, FUN) \
+template <typename R> R & operator OP (R & res, const R & val) { res = FUN(res, val); return res; }
+
+#define UN_OP(OP, FUN) \
+template <typename R> R operator OP (const R & val) { return FUN(val); }
+
+BIN_OP(+, v_add)
+BIN_OP(-, v_sub)
+BIN_OP(*, v_mul)
+BIN_OP(/, v_div)
+BIN_OP(&, v_and)
+BIN_OP(|, v_or)
+BIN_OP(^, v_xor)
+
+BIN_OP(==, v_eq)
+BIN_OP(!=, v_ne)
+BIN_OP(<, v_lt)
+BIN_OP(>, v_gt)
+BIN_OP(<=, v_le)
+BIN_OP(>=, v_ge)
+
+BIN_A_OP(+=, v_add)
+BIN_A_OP(-=, v_sub)
+BIN_A_OP(*=, v_mul)
+BIN_A_OP(/=, v_div)
+BIN_A_OP(&=, v_and)
+BIN_A_OP(|=, v_or)
+BIN_A_OP(^=, v_xor)
+
+UN_OP(~, v_not)
+
+// TODO: shift operators?
+
+}} // cv::hal::
+
+//==============================================================================
+
+#ifdef OPENCV_ENABLE_INLINE_INTRIN_OPERATOR_TEST
+
+namespace cv { namespace hal {
+
+inline static void opencv_operator_compile_test()
+{
+    using namespace cv;
+    v_float32 a, b, c;
+    uint8_t shift = 1;
+    a = b + c;
+    a = b - c;
+    a = b * c;
+    a = b / c;
+    a = b & c;
+    a = b | c;
+    a = b ^ c;
+    // a = b >> shift;
+    // a = b << shift;
+
+    a = (b == c);
+    a = (b != c);
+    a = (b < c);}}
+    a = (b > c);
+    a = (b <= c);
+    a = (b >= c);
+
+    a += b;
+    a -= b;
+    a *= b;
+    a /= b;
+    a &= b;
+    a |= b;
+    a ^= b;
+    // a <<= shift;
+    // a >>= shift;
+
+    a = ~b;
+}
+
+}} // cv::hal::
+
+#endif
+
+
+#endif // OPENCV_HAL_INTRIN_LEGACY_OPS_HPP

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

@@ -3184,6 +3184,12 @@ Mat_<_Tp>& Mat_<_Tp>::operator = (const MatExpr& e)
     return *this;
 }
 
+template<typename _Tp> inline
+MatExpr Mat_<_Tp>::zeros(int _ndims, const int* _sizes)
+{
+    return Mat::zeros(_ndims, _sizes, traits::Type<_Tp>::value);
+}
+
 template<typename _Tp> inline
 MatExpr Mat_<_Tp>::zeros(int rows, int cols)
 {

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

@@ -147,7 +147,23 @@ namespace cv { namespace cuda
         inline explicit NppStreamHandler(cudaStream_t newStream)
         {
             nppStreamContext = {};
-            nppSafeCall(nppGetStreamContext(&nppStreamContext));
+            #if CUDA_VERSION < 12090
+                nppSafeCall(nppGetStreamContext(&nppStreamContext));
+            #else
+                int device = 0;
+                cudaSafeCall(cudaGetDevice(&device));
+
+                cudaDeviceProp prop{};
+                cudaSafeCall(cudaGetDeviceProperties(&prop, device));
+
+                nppStreamContext.nCudaDeviceId = device;
+                nppStreamContext.nMultiProcessorCount = prop.multiProcessorCount;
+                nppStreamContext.nMaxThreadsPerMultiProcessor = prop.maxThreadsPerMultiProcessor;
+                nppStreamContext.nMaxThreadsPerBlock = prop.maxThreadsPerBlock;
+                nppStreamContext.nSharedMemPerBlock = prop.sharedMemPerBlock;
+                nppStreamContext.nCudaDevAttrComputeCapabilityMajor = prop.major;
+                nppStreamContext.nCudaDevAttrComputeCapabilityMinor = prop.minor;
+            #endif
             nppStreamContext.hStream = newStream;
             cudaSafeCall(cudaStreamGetFlags(nppStreamContext.hStream, &nppStreamContext.nStreamFlags));
         }