Merge branch 4.x

modules/core/src/alloc.cpp

@@ -82,7 +82,7 @@ cv::utils::AllocatorStatisticsInterface& getAllocatorStatistics()
     return allocator_stats;
 }
 
-#if defined HAVE_POSIX_MEMALIGN || defined HAVE_MEMALIGN
+#if defined HAVE_POSIX_MEMALIGN || defined HAVE_MEMALIGN || defined HAVE_WIN32_ALIGNED_MALLOC
 static bool readMemoryAlignmentParameter()
 {
     bool value = true;
@@ -100,25 +100,27 @@ static bool readMemoryAlignmentParameter()
     // TODO add checks for valgrind, ASAN if value == false
     return value;
 }
+
+#if defined _MSC_VER
+#pragma warning(suppress:4714)  // preventive: const marked as __forceinline not inlined
+static __forceinline
+#else
 static inline
+#endif
 bool isAlignedAllocationEnabled()
 {
-    static bool initialized = false;
-    static bool useMemalign = true;
-    if (!initialized)
-    {
-        initialized = true;  // trick to avoid stuck in acquire (works only if allocations are scope based)
-        useMemalign = readMemoryAlignmentParameter();
-    }
+    // use construct on first use idiom https://isocpp.org/wiki/faq/ctors#static-init-order-on-first-use
+    // details: https://github.com/opencv/opencv/issues/15691
+    static bool useMemalign = readMemoryAlignmentParameter();
     return useMemalign;
 }
-// do not use variable directly, details: https://github.com/opencv/opencv/issues/15691
+
+// need for this static const is disputed; retaining as it doesn't cause harm
 static const bool g_force_initialization_memalign_flag
 #if defined __GNUC__
     __attribute__((unused))
 #endif
     = isAlignedAllocationEnabled();
-
 #endif
 
 #ifdef OPENCV_ALLOC_ENABLE_STATISTICS
@@ -146,6 +148,14 @@ void* fastMalloc(size_t size)
             return OutOfMemoryError(size);
         return ptr;
     }
+#elif defined HAVE_WIN32_ALIGNED_MALLOC
+    if (isAlignedAllocationEnabled())
+    {
+        void* ptr = _aligned_malloc(size, CV_MALLOC_ALIGN);
+        if(!ptr)
+            return OutOfMemoryError(size);
+        return ptr;
+    }
 #endif
     uchar* udata = (uchar*)malloc(size + sizeof(void*) + CV_MALLOC_ALIGN);
     if(!udata)
@@ -168,6 +178,12 @@ void fastFree(void* ptr)
         free(ptr);
         return;
     }
+#elif defined HAVE_WIN32_ALIGNED_MALLOC
+    if (isAlignedAllocationEnabled())
+    {
+        _aligned_free(ptr);
+        return;
+    }
 #endif
     if(ptr)
     {

modules/core/src/arithm.cpp

@@ -57,26 +57,6 @@ namespace cv
 *                                   logical operations                                   *
 \****************************************************************************************/
 
-void convertAndUnrollScalar( const Mat& sc, int buftype, uchar* scbuf, size_t blocksize )
-{
-    int scn = (int)sc.total(), cn = CV_MAT_CN(buftype);
-    size_t esz = CV_ELEM_SIZE(buftype);
-    BinaryFunc cvtFn = getConvertFunc(sc.depth(), buftype);
-    CV_Assert(cvtFn);
-    cvtFn(sc.ptr(), 1, 0, 1, scbuf, 1, Size(std::min(cn, scn), 1), 0);
-    // unroll the scalar
-    if( scn < cn )
-    {
-        CV_Assert( scn == 1 );
-        size_t esz1 = CV_ELEM_SIZE1(buftype);
-        for( size_t i = esz1; i < esz; i++ )
-            scbuf[i] = scbuf[i - esz1];
-    }
-    for( size_t i = esz; i < blocksize*esz; i++ )
-        scbuf[i] = scbuf[i - esz];
-}
-
-
 enum { OCL_OP_ADD=0, OCL_OP_SUB=1, OCL_OP_RSUB=2, OCL_OP_ABSDIFF=3, OCL_OP_MUL=4,
        OCL_OP_MUL_SCALE=5, OCL_OP_DIV_SCALE=6, OCL_OP_RECIP_SCALE=7, OCL_OP_ADDW=8,
        OCL_OP_AND=9, OCL_OP_OR=10, OCL_OP_XOR=11, OCL_OP_NOT=12, OCL_OP_MIN=13, OCL_OP_MAX=14,
@@ -647,7 +627,8 @@ static void arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
         (kind1 == _InputArray::MATX && (sz1 == Size(1,4) || sz1 == Size(1,1))) ||
         (kind2 == _InputArray::MATX && (sz2 == Size(1,4) || sz2 == Size(1,1))) )
     {
-        if( checkScalar(*psrc1, type2, kind1, kind2) )
+        if ((type1 == CV_64F && (sz1.height == 1 || sz1.height == 4)) &&
+            checkScalar(*psrc1, type2, kind1, kind2))
         {
             // src1 is a scalar; swap it with src2
             swap(psrc1, psrc2);
@@ -1002,9 +983,7 @@ static BinaryFuncC* getRecipTab()
     return recipTab;
 }
 
-}
-
-void cv::multiply(InputArray src1, InputArray src2,
+void multiply(InputArray src1, InputArray src2,
                   OutputArray dst, double scale, int dtype)
 {
     CV_INSTRUMENT_REGION();
@@ -1013,7 +992,7 @@ void cv::multiply(InputArray src1, InputArray src2,
               true, &scale, std::abs(scale - 1.0) < DBL_EPSILON ? OCL_OP_MUL : OCL_OP_MUL_SCALE);
 }
 
-void cv::divide(InputArray src1, InputArray src2,
+void divide(InputArray src1, InputArray src2,
                 OutputArray dst, double scale, int dtype)
 {
     CV_INSTRUMENT_REGION();
@@ -1021,7 +1000,7 @@ void cv::divide(InputArray src1, InputArray src2,
     arithm_op(src1, src2, dst, noArray(), dtype, getDivTab(), true, &scale, OCL_OP_DIV_SCALE);
 }
 
-void cv::divide(double scale, InputArray src2,
+void divide(double scale, InputArray src2,
                 OutputArray dst, int dtype)
 {
     CV_INSTRUMENT_REGION();
@@ -1029,13 +1008,17 @@ void cv::divide(double scale, InputArray src2,
     arithm_op(src2, src2, dst, noArray(), dtype, getRecipTab(), true, &scale, OCL_OP_RECIP_SCALE);
 }
 
+UMat UMat::mul(InputArray m, double scale) const
+{
+    UMat dst;
+    multiply(*this, m, dst, scale);
+    return dst;
+}
+
 /****************************************************************************************\
 *                                      addWeighted                                       *
 \****************************************************************************************/
 
-namespace cv
-{
-
 static BinaryFuncC* getAddWeightedTab()
 {
     static BinaryFuncC addWeightedTab[] =
@@ -1849,6 +1832,9 @@ void cv::inRange(InputArray _src, InputArray _lowerb,
     }
 }
 
+
+#ifndef OPENCV_EXCLUDE_C_API
+
 /****************************************************************************************\
 *                                Earlier API: cvAdd etc.                                 *
 \****************************************************************************************/
@@ -2008,4 +1994,5 @@ cvCmpS( const void* srcarr1, double value, void* dstarr, int cmp_op )
     cv::compare( src1, value, dst, cmp_op );
 }
 
+#endif  // OPENCV_EXCLUDE_C_API
 /* End of file. */

modules/core/src/array.cpp

@@ -48,6 +48,8 @@
 
 #include "precomp.hpp"
 
+#ifndef OPENCV_EXCLUDE_C_API
+
 #define  CV_ORIGIN_TL  0
 #define  CV_ORIGIN_BL  1
 
@@ -2748,53 +2750,6 @@ void DefaultDeleter<CvMatND>::operator ()(CvMatND* obj) const { cvReleaseMatND(&
 void DefaultDeleter<CvSparseMat>::operator ()(CvSparseMat* obj) const { cvReleaseSparseMat(&obj); }
 void DefaultDeleter<CvMemStorage>::operator ()(CvMemStorage* obj) const { cvReleaseMemStorage(&obj); }
 
-template <typename T> static inline
-void scalarToRawData_(const Scalar& s, T * const buf, const int cn, const int unroll_to)
-{
-    int i = 0;
-    for(; i < cn; i++)
-        buf[i] = saturate_cast<T>(s.val[i]);
-    for(; i < unroll_to; i++)
-        buf[i] = buf[i-cn];
-}
-
-void scalarToRawData(const Scalar& s, void* _buf, int type, int unroll_to)
-{
-    CV_INSTRUMENT_REGION();
-
-    const int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
-    CV_Assert(cn <= 4);
-    switch(depth)
-    {
-    case CV_8U:
-        scalarToRawData_<uchar>(s, (uchar*)_buf, cn, unroll_to);
-        break;
-    case CV_8S:
-        scalarToRawData_<schar>(s, (schar*)_buf, cn, unroll_to);
-        break;
-    case CV_16U:
-        scalarToRawData_<ushort>(s, (ushort*)_buf, cn, unroll_to);
-        break;
-    case CV_16S:
-        scalarToRawData_<short>(s, (short*)_buf, cn, unroll_to);
-        break;
-    case CV_32S:
-        scalarToRawData_<int>(s, (int*)_buf, cn, unroll_to);
-        break;
-    case CV_32F:
-        scalarToRawData_<float>(s, (float*)_buf, cn, unroll_to);
-        break;
-    case CV_64F:
-        scalarToRawData_<double>(s, (double*)_buf, cn, unroll_to);
-        break;
-    case CV_16F:
-        scalarToRawData_<float16_t>(s, (float16_t*)_buf, cn, unroll_to);
-        break;
-    default:
-        CV_Error(CV_StsUnsupportedFormat,"");
-    }
-}
-
 } // cv::
 
 
@@ -2817,4 +2772,5 @@ cvRelease( void** struct_ptr )
 }
 
 
+#endif  // OPENCV_EXCLUDE_C_API
 /* End of file. */

modules/core/src/bindings_utils.cpp

@@ -5,6 +5,8 @@
 #include "precomp.hpp"
 #include "opencv2/core/bindings_utils.hpp"
 #include <sstream>
+#include <opencv2/core/utils/filesystem.hpp>
+#include <opencv2/core/utils/filesystem.private.hpp>
 
 namespace cv { namespace utils {
 
@@ -208,4 +210,15 @@ CV_EXPORTS_W String dumpInputOutputArrayOfArrays(InputOutputArrayOfArrays argume
     return ss.str();
 }
 
+namespace fs {
+cv::String getCacheDirectoryForDownloads()
+{
+#if OPENCV_HAVE_FILESYSTEM_SUPPORT
+    return cv::utils::fs::getCacheDirectory("downloads", "OPENCV_DOWNLOADS_CACHE_DIR");
+#else
+    CV_Error(Error::StsNotImplemented, "File system support is disabled in this OpenCV build!");
+#endif
+}
+} // namespace fs
+
 }} // namespace

modules/core/src/convert.dispatch.cpp

@@ -154,7 +154,7 @@ static bool ocl_convertFp16( InputArray _src, OutputArray _dst, int sdepth, int
                               sdepth == CV_32F ? "half" : "float",
                               rowsPerWI,
                               sdepth == CV_32F ? " -D FLOAT_TO_HALF " : "");
-    ocl::Kernel k("convertFp16", ocl::core::halfconvert_oclsrc, build_opt);
+    ocl::Kernel k(sdepth == CV_32F ? "convertFp16_FP32_to_FP16" : "convertFp16_FP16_to_FP32", ocl::core::halfconvert_oclsrc, build_opt);
     if (k.empty())
         return false;
 

modules/core/src/convert_c.cpp

@@ -5,6 +5,7 @@
 
 #include "precomp.hpp"
 
+#ifndef OPENCV_EXCLUDE_C_API
 
 CV_IMPL void
 cvSplit( const void* srcarr, void* dstarr0, void* dstarr1, void* dstarr2, void* dstarr3 )
@@ -132,3 +133,5 @@ CV_IMPL void cvNormalize( const CvArr* srcarr, CvArr* dstarr,
     CV_Assert( dst.size() == src.size() && src.channels() == dst.channels() );
     cv::normalize( src, dst, a, b, norm_type, dst.type(), mask );
 }
+
+#endif  // OPENCV_EXCLUDE_C_API

modules/core/src/convert_scale.dispatch.cpp

@@ -9,7 +9,6 @@
 #include "convert_scale.simd.hpp"
 #include "convert_scale.simd_declarations.hpp" // defines CV_CPU_DISPATCH_MODES_ALL=AVX2,...,BASELINE based on CMakeLists.txt content
 
-
 namespace cv
 {
 
@@ -117,143 +116,4 @@ void convertScaleAbs(InputArray _src, OutputArray _dst, double alpha, double bet
     }
 }
 
-//==================================================================================================
-
-#ifdef HAVE_OPENCL
-
-static bool ocl_normalize( InputArray _src, InputOutputArray _dst, InputArray _mask, int dtype,
-                           double scale, double delta )
-{
-    UMat src = _src.getUMat();
-
-    if( _mask.empty() )
-        src.convertTo( _dst, dtype, scale, delta );
-    else if (src.channels() <= 4)
-    {
-        const ocl::Device & dev = ocl::Device::getDefault();
-
-        int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype),
-                ddepth = CV_MAT_DEPTH(dtype), wdepth = std::max(CV_32F, std::max(sdepth, ddepth)),
-                rowsPerWI = dev.isIntel() ? 4 : 1;
-
-        float fscale = static_cast<float>(scale), fdelta = static_cast<float>(delta);
-        bool haveScale = std::fabs(scale - 1) > DBL_EPSILON,
-                haveZeroScale = !(std::fabs(scale) > DBL_EPSILON),
-                haveDelta = std::fabs(delta) > DBL_EPSILON,
-                doubleSupport = dev.doubleFPConfig() > 0;
-
-        if (!haveScale && !haveDelta && stype == dtype)
-        {
-            _src.copyTo(_dst, _mask);
-            return true;
-        }
-        if (haveZeroScale)
-        {
-            _dst.setTo(Scalar(delta), _mask);
-            return true;
-        }
-
-        if ((sdepth == CV_64F || ddepth == CV_64F) && !doubleSupport)
-            return false;
-
-        char cvt[2][40];
-        String opts = format("-D srcT=%s -D dstT=%s -D convertToWT=%s -D cn=%d -D rowsPerWI=%d"
-                             " -D convertToDT=%s -D workT=%s%s%s%s -D srcT1=%s -D dstT1=%s",
-                             ocl::typeToStr(stype), ocl::typeToStr(dtype),
-                             ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]), cn,
-                             rowsPerWI, ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]),
-                             ocl::typeToStr(CV_MAKE_TYPE(wdepth, cn)),
-                             doubleSupport ? " -D DOUBLE_SUPPORT" : "",
-                             haveScale ? " -D HAVE_SCALE" : "",
-                             haveDelta ? " -D HAVE_DELTA" : "",
-                             ocl::typeToStr(sdepth), ocl::typeToStr(ddepth));
-
-        ocl::Kernel k("normalizek", ocl::core::normalize_oclsrc, opts);
-        if (k.empty())
-            return false;
-
-        UMat mask = _mask.getUMat(), dst = _dst.getUMat();
-
-        ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
-                maskarg = ocl::KernelArg::ReadOnlyNoSize(mask),
-                dstarg = ocl::KernelArg::ReadWrite(dst);
-
-        if (haveScale)
-        {
-            if (haveDelta)
-                k.args(srcarg, maskarg, dstarg, fscale, fdelta);
-            else
-                k.args(srcarg, maskarg, dstarg, fscale);
-        }
-        else
-        {
-            if (haveDelta)
-                k.args(srcarg, maskarg, dstarg, fdelta);
-            else
-                k.args(srcarg, maskarg, dstarg);
-        }
-
-        size_t globalsize[2] = { (size_t)src.cols, ((size_t)src.rows + rowsPerWI - 1) / rowsPerWI };
-        return k.run(2, globalsize, NULL, false);
-    }
-    else
-    {
-        UMat temp;
-        src.convertTo( temp, dtype, scale, delta );
-        temp.copyTo( _dst, _mask );
-    }
-
-    return true;
-}
-
-#endif
-
-void normalize(InputArray _src, InputOutputArray _dst, double a, double b,
-               int norm_type, int rtype, InputArray _mask)
-{
-    CV_INSTRUMENT_REGION();
-
-    double scale = 1, shift = 0;
-    int type = _src.type(), depth = CV_MAT_DEPTH(type);
-
-    if( rtype < 0 )
-        rtype = _dst.fixedType() ? _dst.depth() : depth;
-
-    if( norm_type == CV_MINMAX )
-    {
-        double smin = 0, smax = 0;
-        double dmin = MIN( a, b ), dmax = MAX( a, b );
-        minMaxIdx( _src, &smin, &smax, 0, 0, _mask );
-        scale = (dmax - dmin)*(smax - smin > DBL_EPSILON ? 1./(smax - smin) : 0);
-        if( rtype == CV_32F )
-        {
-            scale = (float)scale;
-            shift = (float)dmin - (float)(smin*scale);
-        }
-        else
-            shift = dmin - smin*scale;
-    }
-    else if( norm_type == CV_L2 || norm_type == CV_L1 || norm_type == CV_C )
-    {
-        scale = norm( _src, norm_type, _mask );
-        scale = scale > DBL_EPSILON ? a/scale : 0.;
-        shift = 0;
-    }
-    else
-        CV_Error( CV_StsBadArg, "Unknown/unsupported norm type" );
-
-    CV_OCL_RUN(_dst.isUMat(),
-               ocl_normalize(_src, _dst, _mask, rtype, scale, shift))
-
-    Mat src = _src.getMat();
-    if( _mask.empty() )
-        src.convertTo( _dst, rtype, scale, shift );
-    else
-    {
-        Mat temp;
-        src.convertTo( temp, rtype, scale, shift );
-        temp.copyTo( _dst, _mask );
-    }
-}
-
 } // namespace

modules/core/src/copy.cpp

@@ -53,6 +53,72 @@
 namespace cv
 {
 
+template <typename T> static inline
+void scalarToRawData_(const Scalar& s, T * const buf, const int cn, const int unroll_to)
+{
+    int i = 0;
+    for(; i < cn; i++)
+        buf[i] = saturate_cast<T>(s.val[i]);
+    for(; i < unroll_to; i++)
+        buf[i] = buf[i-cn];
+}
+
+void scalarToRawData(const Scalar& s, void* _buf, int type, int unroll_to)
+{
+    CV_INSTRUMENT_REGION();
+
+    const int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+    CV_Assert(cn <= 4);
+    switch(depth)
+    {
+    case CV_8U:
+        scalarToRawData_<uchar>(s, (uchar*)_buf, cn, unroll_to);
+        break;
+    case CV_8S:
+        scalarToRawData_<schar>(s, (schar*)_buf, cn, unroll_to);
+        break;
+    case CV_16U:
+        scalarToRawData_<ushort>(s, (ushort*)_buf, cn, unroll_to);
+        break;
+    case CV_16S:
+        scalarToRawData_<short>(s, (short*)_buf, cn, unroll_to);
+        break;
+    case CV_32S:
+        scalarToRawData_<int>(s, (int*)_buf, cn, unroll_to);
+        break;
+    case CV_32F:
+        scalarToRawData_<float>(s, (float*)_buf, cn, unroll_to);
+        break;
+    case CV_64F:
+        scalarToRawData_<double>(s, (double*)_buf, cn, unroll_to);
+        break;
+    case CV_16F:
+        scalarToRawData_<float16_t>(s, (float16_t*)_buf, cn, unroll_to);
+        break;
+    default:
+        CV_Error(CV_StsUnsupportedFormat,"");
+    }
+}
+
+void convertAndUnrollScalar( const Mat& sc, int buftype, uchar* scbuf, size_t blocksize )
+{
+    int scn = (int)sc.total(), cn = CV_MAT_CN(buftype);
+    size_t esz = CV_ELEM_SIZE(buftype);
+    BinaryFunc cvtFn = getConvertFunc(sc.depth(), buftype);
+    CV_Assert(cvtFn);
+    cvtFn(sc.ptr(), 1, 0, 1, scbuf, 1, Size(std::min(cn, scn), 1), 0);
+    // unroll the scalar
+    if( scn < cn )
+    {
+        CV_Assert( scn == 1 );
+        size_t esz1 = CV_ELEM_SIZE1(buftype);
+        for( size_t i = esz1; i < esz; i++ )
+            scbuf[i] = scbuf[i - esz1];
+    }
+    for( size_t i = esz; i < blocksize*esz; i++ )
+        scbuf[i] = scbuf[i - esz];
+}
+
 template<typename T> static void
 copyMask_(const uchar* _src, size_t sstep, const uchar* mask, size_t mstep, uchar* _dst, size_t dstep, Size size)
 {
@@ -594,490 +660,6 @@ Mat& Mat::setTo(InputArray _value, InputArray _mask)
     return *this;
 }
 
-#if CV_SIMD128
-template<typename V> CV_ALWAYS_INLINE void flipHoriz_single( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size size, size_t esz )
-{
-    typedef typename V::lane_type T;
-    int end = (int)(size.width*esz);
-    int width = (end + 1)/2;
-    int width_1 = width & -v_uint8x16::nlanes;
-    int i, j;
-
-#if CV_STRONG_ALIGNMENT
-    CV_Assert(isAligned<sizeof(T)>(src, dst));
-#endif
-
-    for( ; size.height--; src += sstep, dst += dstep )
-    {
-        for( i = 0, j = end; i < width_1; i += v_uint8x16::nlanes, j -= v_uint8x16::nlanes )
-        {
-            V t0, t1;
-
-            t0 = v_load((T*)((uchar*)src + i));
-            t1 = v_load((T*)((uchar*)src + j - v_uint8x16::nlanes));
-            t0 = v_reverse(t0);
-            t1 = v_reverse(t1);
-            v_store((T*)(dst + j - v_uint8x16::nlanes), t0);
-            v_store((T*)(dst + i), t1);
-        }
-        if (isAligned<sizeof(T)>(src, dst))
-        {
-            for ( ; i < width; i += sizeof(T), j -= sizeof(T) )
-            {
-                T t0, t1;
-
-                t0 = *((T*)((uchar*)src + i));
-                t1 = *((T*)((uchar*)src + j - sizeof(T)));
-                *((T*)(dst + j - sizeof(T))) = t0;
-                *((T*)(dst + i)) = t1;
-            }
-        }
-        else
-        {
-            for ( ; i < width; i += sizeof(T), j -= sizeof(T) )
-            {
-                for (int k = 0; k < (int)sizeof(T); k++)
-                {
-                    uchar t0, t1;
-
-                    t0 = *((uchar*)src + i + k);
-                    t1 = *((uchar*)src + j + k - sizeof(T));
-                    *(dst + j + k - sizeof(T)) = t0;
-                    *(dst + i + k) = t1;
-                }
-            }
-        }
-    }
-}
-
-template<typename T1, typename T2> CV_ALWAYS_INLINE void flipHoriz_double( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size size, size_t esz )
-{
-    int end = (int)(size.width*esz);
-    int width = (end + 1)/2;
-
-#if CV_STRONG_ALIGNMENT
-    CV_Assert(isAligned<sizeof(T1)>(src, dst));
-    CV_Assert(isAligned<sizeof(T2)>(src, dst));
-#endif
-
-    for( ; size.height--; src += sstep, dst += dstep )
-    {
-        for ( int i = 0, j = end; i < width; i += sizeof(T1) + sizeof(T2), j -= sizeof(T1) + sizeof(T2) )
-        {
-            T1 t0, t1;
-            T2 t2, t3;
-
-            t0 = *((T1*)((uchar*)src + i));
-            t2 = *((T2*)((uchar*)src + i + sizeof(T1)));
-            t1 = *((T1*)((uchar*)src + j - sizeof(T1) - sizeof(T2)));
-            t3 = *((T2*)((uchar*)src + j - sizeof(T2)));
-            *((T1*)(dst + j - sizeof(T1) - sizeof(T2))) = t0;
-            *((T2*)(dst + j - sizeof(T2))) = t2;
-            *((T1*)(dst + i)) = t1;
-            *((T2*)(dst + i + sizeof(T1))) = t3;
-        }
-    }
-}
-#endif
-
-static void
-flipHoriz( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size size, size_t esz )
-{
-#if CV_SIMD
-#if CV_STRONG_ALIGNMENT
-    size_t alignmentMark = ((size_t)src)|((size_t)dst)|sstep|dstep;
-#endif
-    if (esz == 2 * v_uint8x16::nlanes)
-    {
-        int end = (int)(size.width*esz);
-        int width = end/2;
-
-        for( ; size.height--; src += sstep, dst += dstep )
-        {
-            for( int i = 0, j = end - 2 * v_uint8x16::nlanes; i < width; i += 2 * v_uint8x16::nlanes, j -= 2 * v_uint8x16::nlanes )
-            {
-#if CV_SIMD256
-                v_uint8x32 t0, t1;
-
-                t0 = v256_load((uchar*)src + i);
-                t1 = v256_load((uchar*)src + j);
-                v_store(dst + j, t0);
-                v_store(dst + i, t1);
-#else
-                v_uint8x16 t0, t1, t2, t3;
-
-                t0 = v_load((uchar*)src + i);
-                t1 = v_load((uchar*)src + i + v_uint8x16::nlanes);
-                t2 = v_load((uchar*)src + j);
-                t3 = v_load((uchar*)src + j + v_uint8x16::nlanes);
-                v_store(dst + j, t0);
-                v_store(dst + j + v_uint8x16::nlanes, t1);
-                v_store(dst + i, t2);
-                v_store(dst + i + v_uint8x16::nlanes, t3);
-#endif
-            }
-        }
-    }
-    else if (esz == v_uint8x16::nlanes)
-    {
-        int end = (int)(size.width*esz);
-        int width = end/2;
-
-        for( ; size.height--; src += sstep, dst += dstep )
-        {
-            for( int i = 0, j = end - v_uint8x16::nlanes; i < width; i += v_uint8x16::nlanes, j -= v_uint8x16::nlanes )
-            {
-                v_uint8x16 t0, t1;
-
-                t0 = v_load((uchar*)src + i);
-                t1 = v_load((uchar*)src + j);
-                v_store(dst + j, t0);
-                v_store(dst + i, t1);
-            }
-        }
-    }
-    else if (esz == 8
-#if CV_STRONG_ALIGNMENT
-            && isAligned<sizeof(uint64)>(alignmentMark)
-#endif
-    )
-    {
-        flipHoriz_single<v_uint64x2>(src, sstep, dst, dstep, size, esz);
-    }
-    else if (esz == 4
-#if CV_STRONG_ALIGNMENT
-            && isAligned<sizeof(unsigned)>(alignmentMark)
-#endif
-    )
-    {
-        flipHoriz_single<v_uint32x4>(src, sstep, dst, dstep, size, esz);
-    }
-    else if (esz == 2
-#if CV_STRONG_ALIGNMENT
-            && isAligned<sizeof(ushort)>(alignmentMark)
-#endif
-    )
-    {
-        flipHoriz_single<v_uint16x8>(src, sstep, dst, dstep, size, esz);
-    }
-    else if (esz == 1)
-    {
-        flipHoriz_single<v_uint8x16>(src, sstep, dst, dstep, size, esz);
-    }
-    else if (esz == 24
-#if CV_STRONG_ALIGNMENT
-            && isAligned<sizeof(uint64_t)>(alignmentMark)
-#endif
-    )
-    {
-        int end = (int)(size.width*esz);
-        int width = (end + 1)/2;
-
-        for( ; size.height--; src += sstep, dst += dstep )
-        {
-            for ( int i = 0, j = end; i < width; i += v_uint8x16::nlanes + sizeof(uint64_t), j -= v_uint8x16::nlanes + sizeof(uint64_t) )
-            {
-                v_uint8x16 t0, t1;
-                uint64_t t2, t3;
-
-                t0 = v_load((uchar*)src + i);
-                t2 = *((uint64_t*)((uchar*)src + i + v_uint8x16::nlanes));
-                t1 = v_load((uchar*)src + j - v_uint8x16::nlanes - sizeof(uint64_t));
-                t3 = *((uint64_t*)((uchar*)src + j - sizeof(uint64_t)));
-                v_store(dst + j - v_uint8x16::nlanes - sizeof(uint64_t), t0);
-                *((uint64_t*)(dst + j - sizeof(uint64_t))) = t2;
-                v_store(dst + i, t1);
-                *((uint64_t*)(dst + i + v_uint8x16::nlanes)) = t3;
-            }
-        }
-    }
-#if !CV_STRONG_ALIGNMENT
-    else if (esz == 12)
-    {
-        flipHoriz_double<uint64_t,uint>(src, sstep, dst, dstep, size, esz);
-    }
-    else if (esz == 6)
-    {
-        flipHoriz_double<uint,ushort>(src, sstep, dst, dstep, size, esz);
-    }
-    else if (esz == 3)
-    {
-        flipHoriz_double<ushort,uchar>(src, sstep, dst, dstep, size, esz);
-    }
-#endif
-    else
-#endif // CV_SIMD
-    {
-        int i, j, limit = (int)(((size.width + 1)/2)*esz);
-        AutoBuffer<int> _tab(size.width*esz);
-        int* tab = _tab.data();
-
-        for( i = 0; i < size.width; i++ )
-            for( size_t k = 0; k < esz; k++ )
-                tab[i*esz + k] = (int)((size.width - i - 1)*esz + k);
-
-        for( ; size.height--; src += sstep, dst += dstep )
-        {
-            for( i = 0; i < limit; i++ )
-            {
-                j = tab[i];
-                uchar t0 = src[i], t1 = src[j];
-                dst[i] = t1; dst[j] = t0;
-            }
-        }
-    }
-}
-
-static void
-flipVert( const uchar* src0, size_t sstep, uchar* dst0, size_t dstep, Size size, size_t esz )
-{
-    const uchar* src1 = src0 + (size.height - 1)*sstep;
-    uchar* dst1 = dst0 + (size.height - 1)*dstep;
-    size.width *= (int)esz;
-
-    for( int y = 0; y < (size.height + 1)/2; y++, src0 += sstep, src1 -= sstep,
-                                                  dst0 += dstep, dst1 -= dstep )
-    {
-        int i = 0;
-#if CV_SIMD
-#if CV_STRONG_ALIGNMENT
-        if (isAligned<sizeof(int)>(src0, src1, dst0, dst1))
-#endif
-        {
-            for (; i <= size.width - CV_SIMD_WIDTH; i += CV_SIMD_WIDTH)
-            {
-                v_int32 t0 = vx_load((int*)(src0 + i));
-                v_int32 t1 = vx_load((int*)(src1 + i));
-                vx_store((int*)(dst0 + i), t1);
-                vx_store((int*)(dst1 + i), t0);
-            }
-        }
-#if CV_STRONG_ALIGNMENT
-        else
-        {
-            for (; i <= size.width - CV_SIMD_WIDTH; i += CV_SIMD_WIDTH)
-            {
-                v_uint8 t0 = vx_load(src0 + i);
-                v_uint8 t1 = vx_load(src1 + i);
-                vx_store(dst0 + i, t1);
-                vx_store(dst1 + i, t0);
-            }
-        }
-#endif
-#endif
-
-        if (isAligned<sizeof(int)>(src0, src1, dst0, dst1))
-        {
-            for( ; i <= size.width - 16; i += 16 )
-            {
-                int t0 = ((int*)(src0 + i))[0];
-                int t1 = ((int*)(src1 + i))[0];
-
-                ((int*)(dst0 + i))[0] = t1;
-                ((int*)(dst1 + i))[0] = t0;
-
-                t0 = ((int*)(src0 + i))[1];
-                t1 = ((int*)(src1 + i))[1];
-
-                ((int*)(dst0 + i))[1] = t1;
-                ((int*)(dst1 + i))[1] = t0;
-
-                t0 = ((int*)(src0 + i))[2];
-                t1 = ((int*)(src1 + i))[2];
-
-                ((int*)(dst0 + i))[2] = t1;
-                ((int*)(dst1 + i))[2] = t0;
-
-                t0 = ((int*)(src0 + i))[3];
-                t1 = ((int*)(src1 + i))[3];
-
-                ((int*)(dst0 + i))[3] = t1;
-                ((int*)(dst1 + i))[3] = t0;
-            }
-
-            for( ; i <= size.width - 4; i += 4 )
-            {
-                int t0 = ((int*)(src0 + i))[0];
-                int t1 = ((int*)(src1 + i))[0];
-
-                ((int*)(dst0 + i))[0] = t1;
-                ((int*)(dst1 + i))[0] = t0;
-            }
-        }
-
-        for( ; i < size.width; i++ )
-        {
-            uchar t0 = src0[i];
-            uchar t1 = src1[i];
-
-            dst0[i] = t1;
-            dst1[i] = t0;
-        }
-    }
-}
-
-#ifdef HAVE_OPENCL
-
-enum { FLIP_COLS = 1 << 0, FLIP_ROWS = 1 << 1, FLIP_BOTH = FLIP_ROWS | FLIP_COLS };
-
-static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode )
-{
-    CV_Assert(flipCode >= -1 && flipCode <= 1);
-
-    const ocl::Device & dev = ocl::Device::getDefault();
-    int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type),
-            flipType, kercn = std::min(ocl::predictOptimalVectorWidth(_src, _dst), 4);
-
-    bool doubleSupport = dev.doubleFPConfig() > 0;
-    if (!doubleSupport && depth == CV_64F)
-        kercn = cn;
-
-    if (cn > 4)
-        return false;
-
-    const char * kernelName;
-    if (flipCode == 0)
-        kernelName = "arithm_flip_rows", flipType = FLIP_ROWS;
-    else if (flipCode > 0)
-        kernelName = "arithm_flip_cols", flipType = FLIP_COLS;
-    else
-        kernelName = "arithm_flip_rows_cols", flipType = FLIP_BOTH;
-
-    int pxPerWIy = (dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU)) ? 4 : 1;
-    kercn = (cn!=3 || flipType == FLIP_ROWS) ? std::max(kercn, cn) : cn;
-
-    ocl::Kernel k(kernelName, ocl::core::flip_oclsrc,
-        format( "-D T=%s -D T1=%s -D DEPTH=%d -D cn=%d -D PIX_PER_WI_Y=%d -D kercn=%d",
-                kercn != cn ? ocl::typeToStr(CV_MAKE_TYPE(depth, kercn)) : ocl::vecopTypeToStr(CV_MAKE_TYPE(depth, kercn)),
-                kercn != cn ? ocl::typeToStr(depth) : ocl::vecopTypeToStr(depth), depth, cn, pxPerWIy, kercn));
-    if (k.empty())
-        return false;
-
-    Size size = _src.size();
-    _dst.create(size, type);
-    UMat src = _src.getUMat(), dst = _dst.getUMat();
-
-    int cols = size.width * cn / kercn, rows = size.height;
-    cols = flipType == FLIP_COLS ? (cols + 1) >> 1 : cols;
-    rows = flipType & FLIP_ROWS ? (rows + 1) >> 1 : rows;
-
-    k.args(ocl::KernelArg::ReadOnlyNoSize(src),
-           ocl::KernelArg::WriteOnly(dst, cn, kercn), rows, cols);
-
-    size_t maxWorkGroupSize = dev.maxWorkGroupSize();
-    CV_Assert(maxWorkGroupSize % 4 == 0);
-
-    size_t globalsize[2] = { (size_t)cols, ((size_t)rows + pxPerWIy - 1) / pxPerWIy },
-            localsize[2] = { maxWorkGroupSize / 4, 4 };
-    return k.run(2, globalsize, (flipType == FLIP_COLS) && !dev.isIntel() ? localsize : NULL, false);
-}
-
-#endif
-
-#if defined HAVE_IPP
-static bool ipp_flip(Mat &src, Mat &dst, int flip_mode)
-{
-#ifdef HAVE_IPP_IW
-    CV_INSTRUMENT_REGION_IPP();
-
-    // Details: https://github.com/opencv/opencv/issues/12943
-    if (flip_mode <= 0 /* swap rows */
-        && cv::ipp::getIppTopFeatures() != ippCPUID_SSE42
-        && (int64_t)(src.total()) * src.elemSize() >= CV_BIG_INT(0x80000000)/*2Gb*/
-    )
-        return false;
-
-    IppiAxis ippMode;
-    if(flip_mode < 0)
-        ippMode = ippAxsBoth;
-    else if(flip_mode == 0)
-        ippMode = ippAxsHorizontal;
-    else
-        ippMode = ippAxsVertical;
-
-    try
-    {
-        ::ipp::IwiImage iwSrc = ippiGetImage(src);
-        ::ipp::IwiImage iwDst = ippiGetImage(dst);
-
-        CV_INSTRUMENT_FUN_IPP(::ipp::iwiMirror, iwSrc, iwDst, ippMode);
-    }
-    catch(const ::ipp::IwException &)
-    {
-        return false;
-    }
-
-    return true;
-#else
-    CV_UNUSED(src); CV_UNUSED(dst); CV_UNUSED(flip_mode);
-    return false;
-#endif
-}
-#endif
-
-
-void flip( InputArray _src, OutputArray _dst, int flip_mode )
-{
-    CV_INSTRUMENT_REGION();
-
-    CV_Assert( _src.dims() <= 2 );
-    Size size = _src.size();
-
-    if (flip_mode < 0)
-    {
-        if (size.width == 1)
-            flip_mode = 0;
-        if (size.height == 1)
-            flip_mode = 1;
-    }
-
-    if ((size.width == 1 && flip_mode > 0) ||
-        (size.height == 1 && flip_mode == 0))
-    {
-        return _src.copyTo(_dst);
-    }
-
-    CV_OCL_RUN( _dst.isUMat(), ocl_flip(_src, _dst, flip_mode))
-
-    Mat src = _src.getMat();
-    int type = src.type();
-    _dst.create( size, type );
-    Mat dst = _dst.getMat();
-
-    CV_IPP_RUN_FAST(ipp_flip(src, dst, flip_mode));
-
-    size_t esz = CV_ELEM_SIZE(type);
-
-    if( flip_mode <= 0 )
-        flipVert( src.ptr(), src.step, dst.ptr(), dst.step, src.size(), esz );
-    else
-        flipHoriz( src.ptr(), src.step, dst.ptr(), dst.step, src.size(), esz );
-
-    if( flip_mode < 0 )
-        flipHoriz( dst.ptr(), dst.step, dst.ptr(), dst.step, dst.size(), esz );
-}
-
-void rotate(InputArray _src, OutputArray _dst, int rotateMode)
-{
-    CV_Assert(_src.dims() <= 2);
-
-    switch (rotateMode)
-    {
-    case ROTATE_90_CLOCKWISE:
-        transpose(_src, _dst);
-        flip(_dst, _dst, 1);
-        break;
-    case ROTATE_180:
-        flip(_src, _dst, -1);
-        break;
-    case ROTATE_90_COUNTERCLOCKWISE:
-        transpose(_src, _dst);
-        flip(_dst, _dst, 0);
-        break;
-    default:
-        break;
-    }
-}
 
 #if defined HAVE_OPENCL && !defined __APPLE__
 
@@ -1325,13 +907,12 @@ void copyMakeConstBorder_8u( const uchar* src, size_t srcstep, cv::Size srcroi,
         memcpy( dstInner + srcroi.width, constBuf, right );
     }
 
-    dst += dststep*top;
-
     for( i = 0; i < top; i++ )
-        memcpy(dst + (i - top)*dststep, constBuf, dstroi.width);
+        memcpy(dst + i * dststep, constBuf, dstroi.width);
 
+    dst += (top + srcroi.height) * dststep;
     for( i = 0; i < bottom; i++ )
-        memcpy(dst + (i + srcroi.height)*dststep, constBuf, dstroi.width);
+        memcpy(dst + i * dststep, constBuf, dstroi.width);
 }
 
 }
@@ -1500,6 +1081,9 @@ void cv::copyMakeBorder( InputArray _src, OutputArray _dst, int top, int bottom,
     }
 }
 
+
+#ifndef OPENCV_EXCLUDE_C_API
+
 /* dst = src */
 CV_IMPL void
 cvCopy( const void* srcarr, void* dstarr, const void* maskarr )
@@ -1606,4 +1190,5 @@ cvFlip( const CvArr* srcarr, CvArr* dstarr, int flip_mode )
     cv::flip( src, dst, flip_mode );
 }
 
+#endif  // OPENCV_EXCLUDE_C_API
 /* End of file. */

modules/core/src/count_non_zero.dispatch.cpp

@@ -62,11 +62,9 @@ static bool ipp_countNonZero( Mat &src, int &res )
 {
     CV_INSTRUMENT_REGION_IPP();
 
-#if defined __APPLE__ || (defined _MSC_VER && defined _M_IX86)
     // see https://github.com/opencv/opencv/issues/17453
-    if (src.dims <= 2 && src.step > 520000)
+    if (src.dims <= 2 && src.step > 520000 && cv::ipp::getIppTopFeatures() == ippCPUID_SSE42)
         return false;
-#endif
 
 #if IPP_VERSION_X100 < 201801
     // Poor performance of SSE42

modules/core/src/cuda/gpu_mat_nd.cu

@@ -0,0 +1,269 @@
+// 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.
+
+#include "opencv2/opencv_modules.hpp"
+
+#ifndef HAVE_OPENCV_CUDEV
+
+#error "opencv_cudev is required"
+
+#else
+
+#include "opencv2/core/cuda.hpp"
+#include "opencv2/cudev.hpp"
+
+using namespace cv;
+using namespace cv::cuda;
+
+GpuData::GpuData(const size_t _size)
+    : data(nullptr), size(_size)
+{
+    CV_CUDEV_SAFE_CALL(cudaMalloc(&data, _size));
+}
+
+GpuData::~GpuData()
+{
+    CV_CUDEV_SAFE_CALL(cudaFree(data));
+}
+
+/////////////////////////////////////////////////////
+/// create
+
+void GpuMatND::create(SizeArray _size, int _type)
+{
+    {
+        auto elements_nonzero = [](SizeArray& v)
+        {
+            return std::all_of(v.begin(), v.end(),
+                [](unsigned u){ return u > 0; });
+        };
+        CV_Assert(!_size.empty());
+        CV_Assert(elements_nonzero(_size));
+    }
+
+    _type &= Mat::TYPE_MASK;
+
+    if (size == _size && type() == _type && !empty() && !external() && isContinuous() && !isSubmatrix())
+        return;
+
+    release();
+
+    setFields(std::move(_size), _type);
+
+    data_ = std::make_shared<GpuData>(totalMemSize());
+    data = data_->data;
+    offset = 0;
+}
+
+/////////////////////////////////////////////////////
+/// release
+
+void GpuMatND::release()
+{
+    data = nullptr;
+    data_.reset();
+
+    flags = dims = offset = 0;
+    size.clear();
+    step.clear();
+}
+
+/////////////////////////////////////////////////////
+/// clone
+
+static bool next(uchar*& d, const uchar*& s, std::vector<int>& idx, const int dims, const GpuMatND& dst, const GpuMatND& src)
+{
+    int inc = dims-3;
+
+    while (true)
+    {
+        if (idx[inc] == src.size[inc] - 1)
+        {
+            if (inc == 0)
+            {
+                return false;
+            }
+
+            idx[inc] = 0;
+            d -= (dst.size[inc] - 1) * dst.step[inc];
+            s -= (src.size[inc] - 1) * src.step[inc];
+            inc--;
+        }
+        else
+        {
+            idx[inc]++;
+            d += dst.step[inc];
+            s += src.step[inc];
+            break;
+        }
+    }
+
+    return true;
+}
+
+GpuMatND GpuMatND::clone() const
+{
+    CV_DbgAssert(!empty());
+
+    GpuMatND ret(size, type());
+
+    if (isContinuous())
+    {
+        CV_CUDEV_SAFE_CALL(cudaMemcpy(ret.getDevicePtr(), getDevicePtr(), ret.totalMemSize(), cudaMemcpyDeviceToDevice));
+    }
+    else
+    {
+        // 1D arrays are always continuous
+
+        if (dims == 2)
+        {
+            CV_CUDEV_SAFE_CALL(
+                cudaMemcpy2D(ret.getDevicePtr(), ret.step[0], getDevicePtr(), step[0],
+                    size[1]*step[1], size[0], cudaMemcpyDeviceToDevice)
+            );
+        }
+        else
+        {
+            std::vector<int> idx(dims-2, 0);
+
+            uchar* d = ret.getDevicePtr();
+            const uchar* s = getDevicePtr();
+
+            // iterate each 2D plane
+            do
+            {
+                CV_CUDEV_SAFE_CALL(
+                    cudaMemcpy2DAsync(
+                        d, ret.step[dims-2], s, step[dims-2],
+                        size[dims-1]*step[dims-1], size[dims-2], cudaMemcpyDeviceToDevice)
+                );
+            }
+            while (next(d, s, idx, dims, ret, *this));
+
+            CV_CUDEV_SAFE_CALL(cudaStreamSynchronize(0));
+        }
+    }
+
+    return ret;
+}
+
+GpuMatND GpuMatND::clone(Stream& stream) const
+{
+    CV_DbgAssert(!empty());
+
+    GpuMatND ret(size, type());
+
+    cudaStream_t _stream = StreamAccessor::getStream(stream);
+
+    if (isContinuous())
+    {
+        CV_CUDEV_SAFE_CALL(cudaMemcpyAsync(ret.getDevicePtr(), getDevicePtr(), ret.totalMemSize(), cudaMemcpyDeviceToDevice, _stream));
+    }
+    else
+    {
+        // 1D arrays are always continuous
+
+        if (dims == 2)
+        {
+            CV_CUDEV_SAFE_CALL(
+                cudaMemcpy2DAsync(ret.getDevicePtr(), ret.step[0], getDevicePtr(), step[0],
+                    size[1]*step[1], size[0], cudaMemcpyDeviceToDevice, _stream)
+            );
+        }
+        else
+        {
+            std::vector<int> idx(dims-2, 0);
+
+            uchar* d = ret.getDevicePtr();
+            const uchar* s = getDevicePtr();
+
+            // iterate each 2D plane
+            do
+            {
+                CV_CUDEV_SAFE_CALL(
+                    cudaMemcpy2DAsync(
+                        d, ret.step[dims-2], s, step[dims-2],
+                        size[dims-1]*step[dims-1], size[dims-2], cudaMemcpyDeviceToDevice, _stream)
+                );
+            }
+            while (next(d, s, idx, dims, ret, *this));
+        }
+    }
+
+    return ret;
+}
+
+/////////////////////////////////////////////////////
+/// upload
+
+void GpuMatND::upload(InputArray src)
+{
+    Mat mat = src.getMat();
+
+    CV_DbgAssert(!mat.empty());
+
+    if (!mat.isContinuous())
+        mat = mat.clone();
+
+    SizeArray _size(mat.dims);
+    std::copy_n(mat.size.p, mat.dims, _size.data());
+
+    create(std::move(_size), mat.type());
+
+    CV_CUDEV_SAFE_CALL(cudaMemcpy(getDevicePtr(), mat.data, totalMemSize(), cudaMemcpyHostToDevice));
+}
+
+void GpuMatND::upload(InputArray src, Stream& stream)
+{
+    Mat mat = src.getMat();
+
+    CV_DbgAssert(!mat.empty());
+
+    if (!mat.isContinuous())
+        mat = mat.clone();
+
+    SizeArray _size(mat.dims);
+    std::copy_n(mat.size.p, mat.dims, _size.data());
+
+    create(std::move(_size), mat.type());
+
+    cudaStream_t _stream = StreamAccessor::getStream(stream);
+    CV_CUDEV_SAFE_CALL(cudaMemcpyAsync(getDevicePtr(), mat.data, totalMemSize(), cudaMemcpyHostToDevice, _stream));
+}
+
+/////////////////////////////////////////////////////
+/// download
+
+void GpuMatND::download(OutputArray dst) const
+{
+    CV_DbgAssert(!empty());
+
+    dst.create(dims, size.data(), type());
+    Mat mat = dst.getMat();
+
+    GpuMatND gmat = *this;
+
+    if (!gmat.isContinuous())
+        gmat = gmat.clone();
+
+    CV_CUDEV_SAFE_CALL(cudaMemcpy(mat.data, gmat.getDevicePtr(), mat.total() * mat.elemSize(), cudaMemcpyDeviceToHost));
+}
+
+void GpuMatND::download(OutputArray dst, Stream& stream) const
+{
+    CV_DbgAssert(!empty());
+
+    dst.create(dims, size.data(), type());
+    Mat mat = dst.getMat();
+
+    GpuMatND gmat = *this;
+
+    if (!gmat.isContinuous())
+        gmat = gmat.clone(stream);
+
+    cudaStream_t _stream = StreamAccessor::getStream(stream);
+    CV_CUDEV_SAFE_CALL(cudaMemcpyAsync(mat.data, gmat.getDevicePtr(), mat.total() * mat.elemSize(), cudaMemcpyDeviceToHost, _stream));
+}
+
+#endif

modules/core/src/cuda_gpu_mat_nd.cpp

@@ -0,0 +1,180 @@
+// 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.
+
+#include "precomp.hpp"
+
+using namespace cv;
+using namespace cv::cuda;
+
+GpuMatND::~GpuMatND() = default;
+
+GpuMatND::GpuMatND(SizeArray _size, 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);
+
+    setFields(std::move(_size), _type, std::move(_step));
+}
+
+GpuMatND GpuMatND::operator()(const std::vector<Range>& ranges) const
+{
+    CV_Assert(dims == (int)ranges.size());
+
+    for (int i = 0; i < dims; ++i)
+    {
+        Range r = ranges[i];
+        CV_Assert(r == Range::all() || (0 <= r.start && r.start < r.end && r.end <= size[i]));
+    }
+
+    GpuMatND ret = *this;
+
+    for (int i = 0; i < dims; ++i)
+    {
+        Range r = ranges[i];
+        if (r != Range::all() && r != Range(0, ret.size[i]))
+        {
+            ret.offset += r.start * ret.step[i];
+            ret.size[i] = r.size();
+            ret.flags |= Mat::SUBMATRIX_FLAG;
+        }
+    }
+
+    ret.flags = cv::updateContinuityFlag(ret.flags, dims, ret.size.data(), ret.step.data());
+
+    return ret;
+}
+
+GpuMat GpuMatND::createGpuMatHeader(IndexArray idx, Range rowRange, Range colRange) const
+{
+    CV_Assert((int)idx.size() == dims - 2);
+
+    std::vector<Range> ranges;
+    for (int i : idx)
+        ranges.emplace_back(i, i+1);
+    ranges.push_back(rowRange);
+    ranges.push_back(colRange);
+
+    return (*this)(ranges).createGpuMatHeader();
+}
+
+GpuMat GpuMatND::createGpuMatHeader() const
+{
+    auto Effectively2D = [](GpuMatND m)
+    {
+        for (int i = 0; i < m.dims - 2; ++i)
+            if (m.size[i] > 1)
+                return false;
+        return true;
+    };
+    CV_Assert(Effectively2D(*this));
+
+    return GpuMat(size[dims-2], size[dims-1], type(), getDevicePtr(), step[dims-2]);
+}
+
+GpuMat GpuMatND::operator()(IndexArray idx, Range rowRange, Range colRange) const
+{
+    return createGpuMatHeader(idx, rowRange, colRange).clone();
+}
+
+GpuMatND::operator GpuMat() const
+{
+    return createGpuMatHeader().clone();
+}
+
+void GpuMatND::setFields(SizeArray _size, int _type, StepArray _step)
+{
+    _type &= Mat::TYPE_MASK;
+
+    flags = Mat::MAGIC_VAL + _type;
+    dims = static_cast<int>(_size.size());
+    size = std::move(_size);
+
+    if (_step.empty())
+    {
+        step = StepArray(dims);
+
+        step.back() = elemSize();
+        for (int _i = dims - 2; _i >= 0; --_i)
+        {
+            const size_t i = _i;
+            step[i] = step[i+1] * size[i+1];
+        }
+
+        flags |= Mat::CONTINUOUS_FLAG;
+    }
+    else
+    {
+        step = std::move(_step);
+        step.push_back(elemSize());
+
+        flags = cv::updateContinuityFlag(flags, dims, size.data(), step.data());
+    }
+
+    CV_Assert(size.size() == step.size());
+    CV_Assert(step.back() == elemSize());
+}
+
+#ifndef HAVE_CUDA
+
+GpuData::GpuData(const size_t _size)
+    : data(nullptr), size(0)
+{
+    CV_UNUSED(_size);
+    throw_no_cuda();
+}
+
+GpuData::~GpuData()
+{
+}
+
+void GpuMatND::create(SizeArray _size, int _type)
+{
+    CV_UNUSED(_size);
+    CV_UNUSED(_type);
+    throw_no_cuda();
+}
+
+void GpuMatND::release()
+{
+    throw_no_cuda();
+}
+
+GpuMatND GpuMatND::clone() const
+{
+    throw_no_cuda();
+}
+
+GpuMatND GpuMatND::clone(Stream& stream) const
+{
+    CV_UNUSED(stream);
+    throw_no_cuda();
+}
+
+void GpuMatND::upload(InputArray src)
+{
+    CV_UNUSED(src);
+    throw_no_cuda();
+}
+
+void GpuMatND::upload(InputArray src, Stream& stream)
+{
+    CV_UNUSED(src);
+    CV_UNUSED(stream);
+    throw_no_cuda();
+}
+
+void GpuMatND::download(OutputArray dst) const
+{
+    CV_UNUSED(dst);
+    throw_no_cuda();
+}
+
+void GpuMatND::download(OutputArray dst, Stream& stream) const
+{
+    CV_UNUSED(dst);
+    CV_UNUSED(stream);
+    throw_no_cuda();
+}
+
+#endif

modules/core/src/cuda_stream.cpp

@@ -41,6 +41,7 @@
 //M*/
 
 #include "precomp.hpp"
+#include <cstdint>
 
 using namespace cv;
 using namespace cv::cuda;
@@ -293,6 +294,7 @@ public:
 
     Impl();
     Impl(const Ptr<GpuMat::Allocator>& allocator);
+    Impl(const unsigned int cudaFlags);
     explicit Impl(cudaStream_t stream);
 
     ~Impl();
@@ -312,6 +314,13 @@ cv::cuda::Stream::Impl::Impl(const Ptr<GpuMat::Allocator>& allocator) : stream(0
     ownStream = true;
 }
 
+cv::cuda::Stream::Impl::Impl(const unsigned int cudaFlags) : stream(0), ownStream(false)
+{
+    cudaSafeCall(cudaStreamCreateWithFlags(&stream, cudaFlags));
+    ownStream = true;
+    allocator = makePtr<StackAllocator>(stream);
+}
+
 cv::cuda::Stream::Impl::Impl(cudaStream_t stream_) : stream(stream_), ownStream(false)
 {
     allocator = makePtr<StackAllocator>(stream);
@@ -450,6 +459,16 @@ cv::cuda::Stream::Stream(const Ptr<GpuMat::Allocator>& allocator)
 #endif
 }
 
+cv::cuda::Stream::Stream(const size_t cudaFlags)
+{
+#ifndef HAVE_CUDA
+    CV_UNUSED(cudaFlags);
+    throw_no_cuda();
+#else
+    impl_ = makePtr<Impl>(cudaFlags & UINT_MAX);
+#endif
+}
+
 bool cv::cuda::Stream::queryIfComplete() const
 {
 #ifndef HAVE_CUDA

modules/core/src/datastructs.cpp

@@ -40,6 +40,8 @@
 //M*/
 #include "precomp.hpp"
 
+#ifndef OPENCV_EXCLUDE_C_API
+
 /* default alignment for dynamic data strucutures, resided in storages. */
 #define  CV_STRUCT_ALIGN    ((int)sizeof(double))
 
@@ -3585,4 +3587,5 @@ void  seqInsertSlice( CvSeq* seq, int before_index, const CvArr* from_arr )
 
 }
 
+#endif  // OPENCV_EXCLUDE_C_API
 /* End of file. */

modules/core/src/directx.cpp

@@ -1050,7 +1050,7 @@ bool ocl_convert_nv12_to_bgr(
 
     k.args(clImageY, clImageUV, clBuffer, step, cols, rows);
 
-    size_t globalsize[] = { (size_t)cols, (size_t)rows };
+    size_t globalsize[] = { (size_t)cols/2, (size_t)rows/2 };
     return k.run(2, globalsize, 0, false);
 }
 
@@ -1071,7 +1071,7 @@ bool ocl_convert_bgr_to_nv12(
 
     k.args(clBuffer, step, cols, rows, clImageY, clImageUV);
 
-    size_t globalsize[] = { (size_t)cols, (size_t)rows };
+    size_t globalsize[] = { (size_t)cols/2, (size_t)rows/2 };
     return k.run(2, globalsize, 0, false);
 }
 

modules/core/src/dxt.cpp

@@ -4640,6 +4640,9 @@ int cv::getOptimalDFTSize( int size0 )
     return optimalDFTSizeTab[b];
 }
 
+
+#ifndef OPENCV_EXCLUDE_C_API
+
 CV_IMPL void
 cvDFT( const CvArr* srcarr, CvArr* dstarr, int flags, int nonzero_rows )
 {
@@ -4695,4 +4698,5 @@ cvGetOptimalDFTSize( int size0 )
     return cv::getOptimalDFTSize(size0);
 }
 
+#endif  // OPENCV_EXCLUDE_C_API
 /* End of file. */

modules/core/src/hal_internal.cpp

@@ -42,6 +42,7 @@
 //
 //M*/
 
+#include "precomp.hpp"
 #include "hal_internal.hpp"
 
 #ifdef HAVE_LAPACK

modules/core/src/hal_internal.hpp

@@ -45,8 +45,6 @@
 #ifndef OPENCV_CORE_HAL_INTERNAL_HPP
 #define OPENCV_CORE_HAL_INTERNAL_HPP
 
-#include "precomp.hpp"
-
 #ifdef HAVE_LAPACK
 
 int lapack_LU32f(float* a, size_t a_step, int m, float* b, size_t b_step, int n, int* info);

modules/core/src/intel_gpu_gemm.inl.hpp

@@ -25,7 +25,6 @@
 #ifdef HAVE_OPENCL
 
 #include <sstream>
-#include "precomp.hpp"
 #include "opencl_kernels_core.hpp"
 #include "opencv2/core/opencl/runtime/opencl_clamdblas.hpp"
 #include "opencv2/core/opencl/runtime/opencl_core.hpp"

modules/core/src/lapack.cpp

@@ -753,8 +753,6 @@ SVBkSb( int m, int n, const double* w, size_t wstep,
                 (double*)alignPtr(buffer, sizeof(double)), DBL_EPSILON*2 );
 }
 
-}
-
 /****************************************************************************************\
 *                                 Determinant of the matrix                              *
 \****************************************************************************************/
@@ -764,7 +762,7 @@ SVBkSb( int m, int n, const double* w, size_t wstep,
                    m(0,1)*((double)m(1,0)*m(2,2) - (double)m(1,2)*m(2,0)) +  \
                    m(0,2)*((double)m(1,0)*m(2,1) - (double)m(1,1)*m(2,0)))
 
-double cv::determinant( InputArray _mat )
+double determinant( InputArray _mat )
 {
     CV_INSTRUMENT_REGION();
 
@@ -842,7 +840,7 @@ double cv::determinant( InputArray _mat )
 #define Df( y, x ) ((float*)(dstdata + y*dststep))[x]
 #define Dd( y, x ) ((double*)(dstdata + y*dststep))[x]
 
-double cv::invert( InputArray _src, OutputArray _dst, int method )
+double invert( InputArray _src, OutputArray _dst, int method )
 {
     CV_INSTRUMENT_REGION();
 
@@ -1069,13 +1067,19 @@ double cv::invert( InputArray _src, OutputArray _dst, int method )
     return result;
 }
 
+UMat UMat::inv(int method) const
+{
+    UMat m;
+    invert(*this, m, method);
+    return m;
+}
 
 
 /****************************************************************************************\
 *                              Solving a linear system                                   *
 \****************************************************************************************/
 
-bool cv::solve( InputArray _src, InputArray _src2arg, OutputArray _dst, int method )
+bool solve( InputArray _src, InputArray _src2arg, OutputArray _dst, int method )
 {
     CV_INSTRUMENT_REGION();
 
@@ -1374,7 +1378,7 @@ bool cv::solve( InputArray _src, InputArray _src2arg, OutputArray _dst, int meth
 
 /////////////////// finding eigenvalues and eigenvectors of a symmetric matrix ///////////////
 
-bool cv::eigen( InputArray _src, OutputArray _evals, OutputArray _evects )
+bool eigen( InputArray _src, OutputArray _evals, OutputArray _evects )
 {
     CV_INSTRUMENT_REGION();
 
@@ -1396,7 +1400,7 @@ bool cv::eigen( InputArray _src, OutputArray _evals, OutputArray _evects )
     const bool evecNeeded = _evects.needed();
     const int esOptions = evecNeeded ? Eigen::ComputeEigenvectors : Eigen::EigenvaluesOnly;
     _evals.create(n, 1, type);
-    cv::Mat evals = _evals.getMat();
+    Mat evals = _evals.getMat();
     if ( type == CV_64F )
     {
         Eigen::MatrixXd src_eig, zeros_eig;
@@ -1448,9 +1452,6 @@ bool cv::eigen( InputArray _src, OutputArray _evals, OutputArray _evects )
 #endif
 }
 
-namespace cv
-{
-
 static void _SVDcompute( InputArray _aarr, OutputArray _w,
                          OutputArray _u, OutputArray _vt, int flags )
 {
@@ -1598,6 +1599,9 @@ void cv::SVBackSubst(InputArray w, InputArray u, InputArray vt, InputArray rhs,
 }
 
 
+
+#ifndef OPENCV_EXCLUDE_C_API
+
 CV_IMPL double
 cvDet( const CvArr* arr )
 {
@@ -1789,3 +1793,4 @@ cvSVBkSb( const CvArr* warr, const CvArr* uarr,
     cv::SVD::backSubst(w, u, v, rhs, dst);
     CV_Assert( dst.data == dst0.data );
 }
+#endif  // OPENCV_EXCLUDE_C_API

modules/core/src/mathfuncs.cpp

@@ -1638,6 +1638,9 @@ void patchNaNs( InputOutputArray _a, double _val )
 
 }
 
+
+#ifndef OPENCV_EXCLUDE_C_API
+
 CV_IMPL void cvExp( const CvArr* srcarr, CvArr* dstarr )
 {
     cv::Mat src = cv::cvarrToMat(srcarr), dst = cv::cvarrToMat(dstarr);
@@ -1660,6 +1663,7 @@ CV_IMPL int cvCheckArr( const CvArr* arr, int flags,
     return cv::checkRange(cv::cvarrToMat(arr), (flags & CV_CHECK_QUIET) != 0, 0, minVal, maxVal );
 }
 
+#endif  // OPENCV_EXCLUDE_C_API
 
 /*
   Finds real roots of cubic, quadratic or linear equation.
@@ -1955,6 +1959,8 @@ double cv::solvePoly( InputArray _coeffs0, OutputArray _roots0, int maxIters )
 }
 
 
+#ifndef OPENCV_EXCLUDE_C_API
+
 void cvSolvePoly(const CvMat* a, CvMat *r, int maxiter, int)
 {
     cv::Mat _a = cv::cvarrToMat(a);
@@ -1964,6 +1970,7 @@ void cvSolvePoly(const CvMat* a, CvMat *r, int maxiter, int)
     CV_Assert( _r.data == _r0.data ); // check that the array of roots was not reallocated
 }
 
+#endif  // OPENCV_EXCLUDE_C_API
 
 
 // Common constants for dispatched code

modules/core/src/mathfuncs_core.dispatch.cpp

@@ -7,6 +7,10 @@
 #include "mathfuncs_core.simd.hpp"
 #include "mathfuncs_core.simd_declarations.hpp" // defines CV_CPU_DISPATCH_MODES_ALL=AVX2,...,BASELINE based on CMakeLists.txt content
 
+
+#define IPP_DISABLE_MAGNITUDE_32F 1  // accuracy: https://github.com/opencv/opencv/issues/19506
+
+
 namespace cv { namespace hal {
 
 ///////////////////////////////////// ATAN2 ////////////////////////////////////
@@ -44,8 +48,25 @@ void magnitude32f(const float* x, const float* y, float* mag, int len)
     CV_INSTRUMENT_REGION();
 
     CALL_HAL(magnitude32f, cv_hal_magnitude32f, x, y, mag, len);
+
+#ifdef HAVE_IPP
+    bool allowIPP = true;
+#ifdef IPP_DISABLE_MAGNITUDE_32F
+    if (cv::ipp::getIppTopFeatures() & (
+#if IPP_VERSION_X100 >= 201700
+            ippCPUID_AVX512F |
+#endif
+            ippCPUID_AVX2)
+    )
+    {
+        allowIPP = (len & 7) == 0;
+    }
+#endif
+
     // SSE42 performance issues
-    CV_IPP_RUN(IPP_VERSION_X100 > 201800 || cv::ipp::getIppTopFeatures() != ippCPUID_SSE42, CV_INSTRUMENT_FUN_IPP(ippsMagnitude_32f, x, y, mag, len) >= 0);
+    CV_IPP_RUN((IPP_VERSION_X100 > 201800 || cv::ipp::getIppTopFeatures() != ippCPUID_SSE42) && allowIPP,
+        CV_INSTRUMENT_FUN_IPP(ippsMagnitude_32f, x, y, mag, len) >= 0);
+#endif
 
     CV_CPU_DISPATCH(magnitude32f, (x, y, mag, len),
         CV_CPU_DISPATCH_MODES_ALL);

modules/core/src/matmul.dispatch.cpp

@@ -999,8 +999,79 @@ double Mat::dot(InputArray _mat) const
     return r;
 }
 
+
+#ifdef HAVE_OPENCL
+
+static bool ocl_dot( InputArray _src1, InputArray _src2, double & res )
+{
+    UMat src1 = _src1.getUMat().reshape(1), src2 = _src2.getUMat().reshape(1);
+
+    int type = src1.type(), depth = CV_MAT_DEPTH(type),
+            kercn = ocl::predictOptimalVectorWidth(src1, src2);
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if ( !doubleSupport && depth == CV_64F )
+        return false;
+
+    int dbsize = ocl::Device::getDefault().maxComputeUnits();
+    size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
+    int ddepth = std::max(CV_32F, depth);
+
+    int wgs2_aligned = 1;
+    while (wgs2_aligned < (int)wgs)
+        wgs2_aligned <<= 1;
+    wgs2_aligned >>= 1;
+
+    char cvt[40];
+    ocl::Kernel k("reduce", ocl::core::reduce_oclsrc,
+                  format("-D srcT=%s -D srcT1=%s -D dstT=%s -D dstTK=%s -D ddepth=%d -D convertToDT=%s -D OP_DOT "
+                         "-D WGS=%d -D WGS2_ALIGNED=%d%s%s%s -D kercn=%d",
+                         ocl::typeToStr(CV_MAKE_TYPE(depth, kercn)), ocl::typeToStr(depth),
+                         ocl::typeToStr(ddepth), ocl::typeToStr(CV_MAKE_TYPE(ddepth, kercn)),
+                         ddepth, ocl::convertTypeStr(depth, ddepth, kercn, cvt),
+                         (int)wgs, wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : "",
+                         _src1.isContinuous() ? " -D HAVE_SRC_CONT" : "",
+                         _src2.isContinuous() ? " -D HAVE_SRC2_CONT" : "", kercn));
+    if (k.empty())
+        return false;
+
+    UMat db(1, dbsize, ddepth);
+
+    ocl::KernelArg src1arg = ocl::KernelArg::ReadOnlyNoSize(src1),
+            src2arg = ocl::KernelArg::ReadOnlyNoSize(src2),
+            dbarg = ocl::KernelArg::PtrWriteOnly(db);
+
+    k.args(src1arg, src1.cols, (int)src1.total(), dbsize, dbarg, src2arg);
+
+    size_t globalsize = dbsize * wgs;
+    if (k.run(1, &globalsize, &wgs, true))
+    {
+        res = sum(db.getMat(ACCESS_READ))[0];
+        return true;
+    }
+    return false;
+}
+
+#endif
+
+double UMat::dot(InputArray m) const
+{
+    CV_INSTRUMENT_REGION();
+
+    CV_Assert(m.sameSize(*this) && m.type() == type());
+
+#ifdef HAVE_OPENCL
+    double r = 0;
+    CV_OCL_RUN_(dims <= 2, ocl_dot(*this, m, r), r)
+#endif
+
+    return getMat(ACCESS_READ).dot(m);
+}
+
 }  // namespace cv::
 
+
+#ifndef OPENCV_EXCLUDE_C_API
 /****************************************************************************************\
 *                                    Earlier API                                         *
 \****************************************************************************************/
@@ -1225,4 +1296,6 @@ cvBackProjectPCA( const CvArr* proj_arr, const CvArr* avg_arr,
     CV_Assert(dst0.data == dst.data);
 }
 
+#endif  // OPENCV_EXCLUDE_C_API
+
 /* End of file. */

modules/core/src/matmul.simd.hpp

@@ -1537,7 +1537,7 @@ transform_8u( const uchar* src, uchar* dst, const float* m, int len, int scn, in
 static void
 transform_16u( const ushort* src, ushort* dst, const float* m, int len, int scn, int dcn )
 {
-#if CV_SIMD && !defined(__aarch64__) && !defined(_M_ARM64)
+#if CV_SIMD
     if( scn == 3 && dcn == 3 )
     {
         int x = 0;

modules/core/src/matrix.cpp

@@ -204,7 +204,7 @@ MatAllocator* Mat::getStdAllocator()
 
 //==================================================================================================
 
-bool MatSize::operator==(const MatSize& sz) const
+bool MatSize::operator==(const MatSize& sz) const CV_NOEXCEPT
 {
     int d = dims();
     int dsz = sz.dims();
@@ -337,7 +337,7 @@ void finalizeHdr(Mat& m)
 
 //======================================= Mat ======================================================
 
-Mat::Mat()
+Mat::Mat() CV_NOEXCEPT
     : flags(MAGIC_VAL), dims(0), rows(0), cols(0), data(0), datastart(0), dataend(0),
       datalimit(0), allocator(0), u(0), size(&rows), step(0)
 {}

modules/core/src/matrix_c.cpp

@@ -1,7 +1,12 @@
+// 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
+
+#include "precomp.hpp"
 #include "opencv2/core/mat.hpp"
 #include "opencv2/core/types_c.h"
-#include "precomp.hpp"
 
+#ifndef OPENCV_EXCLUDE_C_API
 // glue
 
 CvMatND cvMatND(const cv::Mat& m)
@@ -342,3 +347,5 @@ cvSort( const CvArr* _src, CvArr* _dst, CvArr* _idx, int flags )
         CV_Assert( dst0.data == dst.data );
     }
 }
+
+#endif  // OPENCV_EXCLUDE_C_API

modules/core/src/matrix_iterator.cpp

@@ -2,9 +2,8 @@
 // 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
 
-
-#include "opencv2/core/mat.hpp"
 #include "precomp.hpp"
+#include "opencv2/core/mat.hpp"
 
 namespace cv {
 

modules/core/src/matrix_operations.cpp

@@ -2,11 +2,10 @@
 // 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
 
-
+#include "precomp.hpp"
 #include "opencv2/core/mat.hpp"
 #include "opencv2/core/types_c.h"
 #include "opencl_kernels_core.hpp"
-#include "precomp.hpp"
 
 #undef HAVE_IPP
 #undef CV_IPP_RUN_FAST
@@ -227,6 +226,23 @@ void cv::setIdentity( InputOutputArray _m, const Scalar& s )
     }
 }
 
+
+namespace cv {
+
+UMat UMat::eye(int rows, int cols, int type)
+{
+    return UMat::eye(Size(cols, rows), type);
+}
+
+UMat UMat::eye(Size size, int type)
+{
+    UMat m(size, type);
+    setIdentity(m);
+    return m;
+}
+
+}  // namespace
+
 //////////////////////////////////////////// trace ///////////////////////////////////////////
 
 cv::Scalar cv::trace( InputArray _m )
@@ -261,285 +277,6 @@ cv::Scalar cv::trace( InputArray _m )
     return cv::sum(m.diag());
 }
 
-////////////////////////////////////// transpose /////////////////////////////////////////
-
-namespace cv
-{
-
-template<typename T> static void
-transpose_( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size sz )
-{
-    int i=0, j, m = sz.width, n = sz.height;
-
-    #if CV_ENABLE_UNROLLED
-    for(; i <= m - 4; i += 4 )
-    {
-        T* d0 = (T*)(dst + dstep*i);
-        T* d1 = (T*)(dst + dstep*(i+1));
-        T* d2 = (T*)(dst + dstep*(i+2));
-        T* d3 = (T*)(dst + dstep*(i+3));
-
-        for( j = 0; j <= n - 4; j += 4 )
-        {
-            const T* s0 = (const T*)(src + i*sizeof(T) + sstep*j);
-            const T* s1 = (const T*)(src + i*sizeof(T) + sstep*(j+1));
-            const T* s2 = (const T*)(src + i*sizeof(T) + sstep*(j+2));
-            const T* s3 = (const T*)(src + i*sizeof(T) + sstep*(j+3));
-
-            d0[j] = s0[0]; d0[j+1] = s1[0]; d0[j+2] = s2[0]; d0[j+3] = s3[0];
-            d1[j] = s0[1]; d1[j+1] = s1[1]; d1[j+2] = s2[1]; d1[j+3] = s3[1];
-            d2[j] = s0[2]; d2[j+1] = s1[2]; d2[j+2] = s2[2]; d2[j+3] = s3[2];
-            d3[j] = s0[3]; d3[j+1] = s1[3]; d3[j+2] = s2[3]; d3[j+3] = s3[3];
-        }
-
-        for( ; j < n; j++ )
-        {
-            const T* s0 = (const T*)(src + i*sizeof(T) + j*sstep);
-            d0[j] = s0[0]; d1[j] = s0[1]; d2[j] = s0[2]; d3[j] = s0[3];
-        }
-    }
-    #endif
-    for( ; i < m; i++ )
-    {
-        T* d0 = (T*)(dst + dstep*i);
-        j = 0;
-        #if CV_ENABLE_UNROLLED
-        for(; j <= n - 4; j += 4 )
-        {
-            const T* s0 = (const T*)(src + i*sizeof(T) + sstep*j);
-            const T* s1 = (const T*)(src + i*sizeof(T) + sstep*(j+1));
-            const T* s2 = (const T*)(src + i*sizeof(T) + sstep*(j+2));
-            const T* s3 = (const T*)(src + i*sizeof(T) + sstep*(j+3));
-
-            d0[j] = s0[0]; d0[j+1] = s1[0]; d0[j+2] = s2[0]; d0[j+3] = s3[0];
-        }
-        #endif
-        for( ; j < n; j++ )
-        {
-            const T* s0 = (const T*)(src + i*sizeof(T) + j*sstep);
-            d0[j] = s0[0];
-        }
-    }
-}
-
-template<typename T> static void
-transposeI_( uchar* data, size_t step, int n )
-{
-    for( int i = 0; i < n; i++ )
-    {
-        T* row = (T*)(data + step*i);
-        uchar* data1 = data + i*sizeof(T);
-        for( int j = i+1; j < n; j++ )
-            std::swap( row[j], *(T*)(data1 + step*j) );
-    }
-}
-
-typedef void (*TransposeFunc)( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size sz );
-typedef void (*TransposeInplaceFunc)( uchar* data, size_t step, int n );
-
-#define DEF_TRANSPOSE_FUNC(suffix, type) \
-static void transpose_##suffix( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size sz ) \
-{ transpose_<type>(src, sstep, dst, dstep, sz); } \
-\
-static void transposeI_##suffix( uchar* data, size_t step, int n ) \
-{ transposeI_<type>(data, step, n); }
-
-DEF_TRANSPOSE_FUNC(8u, uchar)
-DEF_TRANSPOSE_FUNC(16u, ushort)
-DEF_TRANSPOSE_FUNC(8uC3, Vec3b)
-DEF_TRANSPOSE_FUNC(32s, int)
-DEF_TRANSPOSE_FUNC(16uC3, Vec3s)
-DEF_TRANSPOSE_FUNC(32sC2, Vec2i)
-DEF_TRANSPOSE_FUNC(32sC3, Vec3i)
-DEF_TRANSPOSE_FUNC(32sC4, Vec4i)
-DEF_TRANSPOSE_FUNC(32sC6, Vec6i)
-DEF_TRANSPOSE_FUNC(32sC8, Vec8i)
-
-static TransposeFunc transposeTab[] =
-{
-    0, transpose_8u, transpose_16u, transpose_8uC3, transpose_32s, 0, transpose_16uC3, 0,
-    transpose_32sC2, 0, 0, 0, transpose_32sC3, 0, 0, 0, transpose_32sC4,
-    0, 0, 0, 0, 0, 0, 0, transpose_32sC6, 0, 0, 0, 0, 0, 0, 0, transpose_32sC8
-};
-
-static TransposeInplaceFunc transposeInplaceTab[] =
-{
-    0, transposeI_8u, transposeI_16u, transposeI_8uC3, transposeI_32s, 0, transposeI_16uC3, 0,
-    transposeI_32sC2, 0, 0, 0, transposeI_32sC3, 0, 0, 0, transposeI_32sC4,
-    0, 0, 0, 0, 0, 0, 0, transposeI_32sC6, 0, 0, 0, 0, 0, 0, 0, transposeI_32sC8
-};
-
-#ifdef HAVE_OPENCL
-
-static bool ocl_transpose( InputArray _src, OutputArray _dst )
-{
-    const ocl::Device & dev = ocl::Device::getDefault();
-    const int TILE_DIM = 32, BLOCK_ROWS = 8;
-    int type = _src.type(), cn = CV_MAT_CN(type), depth = CV_MAT_DEPTH(type),
-        rowsPerWI = dev.isIntel() ? 4 : 1;
-
-    UMat src = _src.getUMat();
-    _dst.create(src.cols, src.rows, type);
-    UMat dst = _dst.getUMat();
-
-    String kernelName("transpose");
-    bool inplace = dst.u == src.u;
-
-    if (inplace)
-    {
-        CV_Assert(dst.cols == dst.rows);
-        kernelName += "_inplace";
-    }
-    else
-    {
-        // check required local memory size
-        size_t required_local_memory = (size_t) TILE_DIM*(TILE_DIM+1)*CV_ELEM_SIZE(type);
-        if (required_local_memory > ocl::Device::getDefault().localMemSize())
-            return false;
-    }
-
-    ocl::Kernel k(kernelName.c_str(), ocl::core::transpose_oclsrc,
-                  format("-D T=%s -D T1=%s -D cn=%d -D TILE_DIM=%d -D BLOCK_ROWS=%d -D rowsPerWI=%d%s",
-                         ocl::memopTypeToStr(type), ocl::memopTypeToStr(depth),
-                         cn, TILE_DIM, BLOCK_ROWS, rowsPerWI, inplace ? " -D INPLACE" : ""));
-    if (k.empty())
-        return false;
-
-    if (inplace)
-        k.args(ocl::KernelArg::ReadWriteNoSize(dst), dst.rows);
-    else
-        k.args(ocl::KernelArg::ReadOnly(src),
-               ocl::KernelArg::WriteOnlyNoSize(dst));
-
-    size_t localsize[2]  = { TILE_DIM, BLOCK_ROWS };
-    size_t globalsize[2] = { (size_t)src.cols, inplace ? ((size_t)src.rows + rowsPerWI - 1) / rowsPerWI : (divUp((size_t)src.rows, TILE_DIM) * BLOCK_ROWS) };
-
-    if (inplace && dev.isIntel())
-    {
-        localsize[0] = 16;
-        localsize[1] = dev.maxWorkGroupSize() / localsize[0];
-    }
-
-    return k.run(2, globalsize, localsize, false);
-}
-
-#endif
-
-#ifdef HAVE_IPP
-static bool ipp_transpose( Mat &src, Mat &dst )
-{
-    CV_INSTRUMENT_REGION_IPP();
-
-    int type = src.type();
-    typedef IppStatus (CV_STDCALL * IppiTranspose)(const void * pSrc, int srcStep, void * pDst, int dstStep, IppiSize roiSize);
-    typedef IppStatus (CV_STDCALL * IppiTransposeI)(const void * pSrcDst, int srcDstStep, IppiSize roiSize);
-    IppiTranspose ippiTranspose = 0;
-    IppiTransposeI ippiTranspose_I = 0;
-
-    if (dst.data == src.data && dst.cols == dst.rows)
-    {
-        CV_SUPPRESS_DEPRECATED_START
-        ippiTranspose_I =
-            type == CV_8UC1 ? (IppiTransposeI)ippiTranspose_8u_C1IR :
-            type == CV_8UC3 ? (IppiTransposeI)ippiTranspose_8u_C3IR :
-            type == CV_8UC4 ? (IppiTransposeI)ippiTranspose_8u_C4IR :
-            type == CV_16UC1 ? (IppiTransposeI)ippiTranspose_16u_C1IR :
-            type == CV_16UC3 ? (IppiTransposeI)ippiTranspose_16u_C3IR :
-            type == CV_16UC4 ? (IppiTransposeI)ippiTranspose_16u_C4IR :
-            type == CV_16SC1 ? (IppiTransposeI)ippiTranspose_16s_C1IR :
-            type == CV_16SC3 ? (IppiTransposeI)ippiTranspose_16s_C3IR :
-            type == CV_16SC4 ? (IppiTransposeI)ippiTranspose_16s_C4IR :
-            type == CV_32SC1 ? (IppiTransposeI)ippiTranspose_32s_C1IR :
-            type == CV_32SC3 ? (IppiTransposeI)ippiTranspose_32s_C3IR :
-            type == CV_32SC4 ? (IppiTransposeI)ippiTranspose_32s_C4IR :
-            type == CV_32FC1 ? (IppiTransposeI)ippiTranspose_32f_C1IR :
-            type == CV_32FC3 ? (IppiTransposeI)ippiTranspose_32f_C3IR :
-            type == CV_32FC4 ? (IppiTransposeI)ippiTranspose_32f_C4IR : 0;
-        CV_SUPPRESS_DEPRECATED_END
-    }
-    else
-    {
-        ippiTranspose =
-            type == CV_8UC1 ? (IppiTranspose)ippiTranspose_8u_C1R :
-            type == CV_8UC3 ? (IppiTranspose)ippiTranspose_8u_C3R :
-            type == CV_8UC4 ? (IppiTranspose)ippiTranspose_8u_C4R :
-            type == CV_16UC1 ? (IppiTranspose)ippiTranspose_16u_C1R :
-            type == CV_16UC3 ? (IppiTranspose)ippiTranspose_16u_C3R :
-            type == CV_16UC4 ? (IppiTranspose)ippiTranspose_16u_C4R :
-            type == CV_16SC1 ? (IppiTranspose)ippiTranspose_16s_C1R :
-            type == CV_16SC3 ? (IppiTranspose)ippiTranspose_16s_C3R :
-            type == CV_16SC4 ? (IppiTranspose)ippiTranspose_16s_C4R :
-            type == CV_32SC1 ? (IppiTranspose)ippiTranspose_32s_C1R :
-            type == CV_32SC3 ? (IppiTranspose)ippiTranspose_32s_C3R :
-            type == CV_32SC4 ? (IppiTranspose)ippiTranspose_32s_C4R :
-            type == CV_32FC1 ? (IppiTranspose)ippiTranspose_32f_C1R :
-            type == CV_32FC3 ? (IppiTranspose)ippiTranspose_32f_C3R :
-            type == CV_32FC4 ? (IppiTranspose)ippiTranspose_32f_C4R : 0;
-    }
-
-    IppiSize roiSize = { src.cols, src.rows };
-    if (ippiTranspose != 0)
-    {
-        if (CV_INSTRUMENT_FUN_IPP(ippiTranspose, src.ptr(), (int)src.step, dst.ptr(), (int)dst.step, roiSize) >= 0)
-            return true;
-    }
-    else if (ippiTranspose_I != 0)
-    {
-        if (CV_INSTRUMENT_FUN_IPP(ippiTranspose_I, dst.ptr(), (int)dst.step, roiSize) >= 0)
-            return true;
-    }
-    return false;
-}
-#endif
-
-}
-
-
-void cv::transpose( InputArray _src, OutputArray _dst )
-{
-    CV_INSTRUMENT_REGION();
-
-    int type = _src.type(), esz = CV_ELEM_SIZE(type);
-    CV_Assert( _src.dims() <= 2 && esz <= 32 );
-
-    CV_OCL_RUN(_dst.isUMat(),
-               ocl_transpose(_src, _dst))
-
-    Mat src = _src.getMat();
-    if( src.empty() )
-    {
-        _dst.release();
-        return;
-    }
-
-    _dst.create(src.cols, src.rows, src.type());
-    Mat dst = _dst.getMat();
-
-    // handle the case of single-column/single-row matrices, stored in STL vectors.
-    if( src.rows != dst.cols || src.cols != dst.rows )
-    {
-        CV_Assert( src.size() == dst.size() && (src.cols == 1 || src.rows == 1) );
-        src.copyTo(dst);
-        return;
-    }
-
-    CV_IPP_RUN_FAST(ipp_transpose(src, dst))
-
-    if( dst.data == src.data )
-    {
-        TransposeInplaceFunc func = transposeInplaceTab[esz];
-        CV_Assert( func != 0 );
-        CV_Assert( dst.cols == dst.rows );
-        func( dst.ptr(), dst.step, dst.rows );
-    }
-    else
-    {
-        TransposeFunc func = transposeTab[esz];
-        CV_Assert( func != 0 );
-        func( src.ptr(), src.step, dst.ptr(), dst.step, src.size() );
-    }
-}
-
 
 ////////////////////////////////////// completeSymm /////////////////////////////////////////
 

modules/core/src/matrix_sparse.cpp

@@ -2,10 +2,9 @@
 // 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
 
-
+#include "precomp.hpp"
 #include "opencv2/core/mat.hpp"
 #include "opencv2/core/types_c.h"
-#include "precomp.hpp"
 
 namespace cv {
 

modules/core/src/matrix_transform.cpp

@@ -0,0 +1,770 @@
+// 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
+
+#include "precomp.hpp"
+#include "opencl_kernels_core.hpp"
+
+namespace cv {
+
+////////////////////////////////////// transpose /////////////////////////////////////////
+
+template<typename T> static void
+transpose_( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size sz )
+{
+    int i=0, j, m = sz.width, n = sz.height;
+
+    #if CV_ENABLE_UNROLLED
+    for(; i <= m - 4; i += 4 )
+    {
+        T* d0 = (T*)(dst + dstep*i);
+        T* d1 = (T*)(dst + dstep*(i+1));
+        T* d2 = (T*)(dst + dstep*(i+2));
+        T* d3 = (T*)(dst + dstep*(i+3));
+
+        for( j = 0; j <= n - 4; j += 4 )
+        {
+            const T* s0 = (const T*)(src + i*sizeof(T) + sstep*j);
+            const T* s1 = (const T*)(src + i*sizeof(T) + sstep*(j+1));
+            const T* s2 = (const T*)(src + i*sizeof(T) + sstep*(j+2));
+            const T* s3 = (const T*)(src + i*sizeof(T) + sstep*(j+3));
+
+            d0[j] = s0[0]; d0[j+1] = s1[0]; d0[j+2] = s2[0]; d0[j+3] = s3[0];
+            d1[j] = s0[1]; d1[j+1] = s1[1]; d1[j+2] = s2[1]; d1[j+3] = s3[1];
+            d2[j] = s0[2]; d2[j+1] = s1[2]; d2[j+2] = s2[2]; d2[j+3] = s3[2];
+            d3[j] = s0[3]; d3[j+1] = s1[3]; d3[j+2] = s2[3]; d3[j+3] = s3[3];
+        }
+
+        for( ; j < n; j++ )
+        {
+            const T* s0 = (const T*)(src + i*sizeof(T) + j*sstep);
+            d0[j] = s0[0]; d1[j] = s0[1]; d2[j] = s0[2]; d3[j] = s0[3];
+        }
+    }
+    #endif
+    for( ; i < m; i++ )
+    {
+        T* d0 = (T*)(dst + dstep*i);
+        j = 0;
+        #if CV_ENABLE_UNROLLED
+        for(; j <= n - 4; j += 4 )
+        {
+            const T* s0 = (const T*)(src + i*sizeof(T) + sstep*j);
+            const T* s1 = (const T*)(src + i*sizeof(T) + sstep*(j+1));
+            const T* s2 = (const T*)(src + i*sizeof(T) + sstep*(j+2));
+            const T* s3 = (const T*)(src + i*sizeof(T) + sstep*(j+3));
+
+            d0[j] = s0[0]; d0[j+1] = s1[0]; d0[j+2] = s2[0]; d0[j+3] = s3[0];
+        }
+        #endif
+        for( ; j < n; j++ )
+        {
+            const T* s0 = (const T*)(src + i*sizeof(T) + j*sstep);
+            d0[j] = s0[0];
+        }
+    }
+}
+
+template<typename T> static void
+transposeI_( uchar* data, size_t step, int n )
+{
+    for( int i = 0; i < n; i++ )
+    {
+        T* row = (T*)(data + step*i);
+        uchar* data1 = data + i*sizeof(T);
+        for( int j = i+1; j < n; j++ )
+            std::swap( row[j], *(T*)(data1 + step*j) );
+    }
+}
+
+typedef void (*TransposeFunc)( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size sz );
+typedef void (*TransposeInplaceFunc)( uchar* data, size_t step, int n );
+
+#define DEF_TRANSPOSE_FUNC(suffix, type) \
+static void transpose_##suffix( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size sz ) \
+{ transpose_<type>(src, sstep, dst, dstep, sz); } \
+\
+static void transposeI_##suffix( uchar* data, size_t step, int n ) \
+{ transposeI_<type>(data, step, n); }
+
+DEF_TRANSPOSE_FUNC(8u, uchar)
+DEF_TRANSPOSE_FUNC(16u, ushort)
+DEF_TRANSPOSE_FUNC(8uC3, Vec3b)
+DEF_TRANSPOSE_FUNC(32s, int)
+DEF_TRANSPOSE_FUNC(16uC3, Vec3s)
+DEF_TRANSPOSE_FUNC(32sC2, Vec2i)
+DEF_TRANSPOSE_FUNC(32sC3, Vec3i)
+DEF_TRANSPOSE_FUNC(32sC4, Vec4i)
+DEF_TRANSPOSE_FUNC(32sC6, Vec6i)
+DEF_TRANSPOSE_FUNC(32sC8, Vec8i)
+
+static TransposeFunc transposeTab[] =
+{
+    0, transpose_8u, transpose_16u, transpose_8uC3, transpose_32s, 0, transpose_16uC3, 0,
+    transpose_32sC2, 0, 0, 0, transpose_32sC3, 0, 0, 0, transpose_32sC4,
+    0, 0, 0, 0, 0, 0, 0, transpose_32sC6, 0, 0, 0, 0, 0, 0, 0, transpose_32sC8
+};
+
+static TransposeInplaceFunc transposeInplaceTab[] =
+{
+    0, transposeI_8u, transposeI_16u, transposeI_8uC3, transposeI_32s, 0, transposeI_16uC3, 0,
+    transposeI_32sC2, 0, 0, 0, transposeI_32sC3, 0, 0, 0, transposeI_32sC4,
+    0, 0, 0, 0, 0, 0, 0, transposeI_32sC6, 0, 0, 0, 0, 0, 0, 0, transposeI_32sC8
+};
+
+#ifdef HAVE_OPENCL
+
+static bool ocl_transpose( InputArray _src, OutputArray _dst )
+{
+    const ocl::Device & dev = ocl::Device::getDefault();
+    const int TILE_DIM = 32, BLOCK_ROWS = 8;
+    int type = _src.type(), cn = CV_MAT_CN(type), depth = CV_MAT_DEPTH(type),
+        rowsPerWI = dev.isIntel() ? 4 : 1;
+
+    UMat src = _src.getUMat();
+    _dst.create(src.cols, src.rows, type);
+    UMat dst = _dst.getUMat();
+
+    String kernelName("transpose");
+    bool inplace = dst.u == src.u;
+
+    if (inplace)
+    {
+        CV_Assert(dst.cols == dst.rows);
+        kernelName += "_inplace";
+    }
+    else
+    {
+        // check required local memory size
+        size_t required_local_memory = (size_t) TILE_DIM*(TILE_DIM+1)*CV_ELEM_SIZE(type);
+        if (required_local_memory > ocl::Device::getDefault().localMemSize())
+            return false;
+    }
+
+    ocl::Kernel k(kernelName.c_str(), ocl::core::transpose_oclsrc,
+                  format("-D T=%s -D T1=%s -D cn=%d -D TILE_DIM=%d -D BLOCK_ROWS=%d -D rowsPerWI=%d%s",
+                         ocl::memopTypeToStr(type), ocl::memopTypeToStr(depth),
+                         cn, TILE_DIM, BLOCK_ROWS, rowsPerWI, inplace ? " -D INPLACE" : ""));
+    if (k.empty())
+        return false;
+
+    if (inplace)
+        k.args(ocl::KernelArg::ReadWriteNoSize(dst), dst.rows);
+    else
+        k.args(ocl::KernelArg::ReadOnly(src),
+               ocl::KernelArg::WriteOnlyNoSize(dst));
+
+    size_t localsize[2]  = { TILE_DIM, BLOCK_ROWS };
+    size_t globalsize[2] = { (size_t)src.cols, inplace ? ((size_t)src.rows + rowsPerWI - 1) / rowsPerWI : (divUp((size_t)src.rows, TILE_DIM) * BLOCK_ROWS) };
+
+    if (inplace && dev.isIntel())
+    {
+        localsize[0] = 16;
+        localsize[1] = dev.maxWorkGroupSize() / localsize[0];
+    }
+
+    return k.run(2, globalsize, localsize, false);
+}
+
+#endif
+
+#ifdef HAVE_IPP
+static bool ipp_transpose( Mat &src, Mat &dst )
+{
+    CV_INSTRUMENT_REGION_IPP();
+
+    int type = src.type();
+    typedef IppStatus (CV_STDCALL * IppiTranspose)(const void * pSrc, int srcStep, void * pDst, int dstStep, IppiSize roiSize);
+    typedef IppStatus (CV_STDCALL * IppiTransposeI)(const void * pSrcDst, int srcDstStep, IppiSize roiSize);
+    IppiTranspose ippiTranspose = 0;
+    IppiTransposeI ippiTranspose_I = 0;
+
+    if (dst.data == src.data && dst.cols == dst.rows)
+    {
+        CV_SUPPRESS_DEPRECATED_START
+        ippiTranspose_I =
+            type == CV_8UC1 ? (IppiTransposeI)ippiTranspose_8u_C1IR :
+            type == CV_8UC3 ? (IppiTransposeI)ippiTranspose_8u_C3IR :
+            type == CV_8UC4 ? (IppiTransposeI)ippiTranspose_8u_C4IR :
+            type == CV_16UC1 ? (IppiTransposeI)ippiTranspose_16u_C1IR :
+            type == CV_16UC3 ? (IppiTransposeI)ippiTranspose_16u_C3IR :
+            type == CV_16UC4 ? (IppiTransposeI)ippiTranspose_16u_C4IR :
+            type == CV_16SC1 ? (IppiTransposeI)ippiTranspose_16s_C1IR :
+            type == CV_16SC3 ? (IppiTransposeI)ippiTranspose_16s_C3IR :
+            type == CV_16SC4 ? (IppiTransposeI)ippiTranspose_16s_C4IR :
+            type == CV_32SC1 ? (IppiTransposeI)ippiTranspose_32s_C1IR :
+            type == CV_32SC3 ? (IppiTransposeI)ippiTranspose_32s_C3IR :
+            type == CV_32SC4 ? (IppiTransposeI)ippiTranspose_32s_C4IR :
+            type == CV_32FC1 ? (IppiTransposeI)ippiTranspose_32f_C1IR :
+            type == CV_32FC3 ? (IppiTransposeI)ippiTranspose_32f_C3IR :
+            type == CV_32FC4 ? (IppiTransposeI)ippiTranspose_32f_C4IR : 0;
+        CV_SUPPRESS_DEPRECATED_END
+    }
+    else
+    {
+        ippiTranspose =
+            type == CV_8UC1 ? (IppiTranspose)ippiTranspose_8u_C1R :
+            type == CV_8UC3 ? (IppiTranspose)ippiTranspose_8u_C3R :
+            type == CV_8UC4 ? (IppiTranspose)ippiTranspose_8u_C4R :
+            type == CV_16UC1 ? (IppiTranspose)ippiTranspose_16u_C1R :
+            type == CV_16UC3 ? (IppiTranspose)ippiTranspose_16u_C3R :
+            type == CV_16UC4 ? (IppiTranspose)ippiTranspose_16u_C4R :
+            type == CV_16SC1 ? (IppiTranspose)ippiTranspose_16s_C1R :
+            type == CV_16SC3 ? (IppiTranspose)ippiTranspose_16s_C3R :
+            type == CV_16SC4 ? (IppiTranspose)ippiTranspose_16s_C4R :
+            type == CV_32SC1 ? (IppiTranspose)ippiTranspose_32s_C1R :
+            type == CV_32SC3 ? (IppiTranspose)ippiTranspose_32s_C3R :
+            type == CV_32SC4 ? (IppiTranspose)ippiTranspose_32s_C4R :
+            type == CV_32FC1 ? (IppiTranspose)ippiTranspose_32f_C1R :
+            type == CV_32FC3 ? (IppiTranspose)ippiTranspose_32f_C3R :
+            type == CV_32FC4 ? (IppiTranspose)ippiTranspose_32f_C4R : 0;
+    }
+
+    IppiSize roiSize = { src.cols, src.rows };
+    if (ippiTranspose != 0)
+    {
+        if (CV_INSTRUMENT_FUN_IPP(ippiTranspose, src.ptr(), (int)src.step, dst.ptr(), (int)dst.step, roiSize) >= 0)
+            return true;
+    }
+    else if (ippiTranspose_I != 0)
+    {
+        if (CV_INSTRUMENT_FUN_IPP(ippiTranspose_I, dst.ptr(), (int)dst.step, roiSize) >= 0)
+            return true;
+    }
+    return false;
+}
+#endif
+
+
+void transpose( InputArray _src, OutputArray _dst )
+{
+    CV_INSTRUMENT_REGION();
+
+    int type = _src.type(), esz = CV_ELEM_SIZE(type);
+    CV_Assert( _src.dims() <= 2 && esz <= 32 );
+
+    CV_OCL_RUN(_dst.isUMat(),
+               ocl_transpose(_src, _dst))
+
+    Mat src = _src.getMat();
+    if( src.empty() )
+    {
+        _dst.release();
+        return;
+    }
+
+    _dst.create(src.cols, src.rows, src.type());
+    Mat dst = _dst.getMat();
+
+    // handle the case of single-column/single-row matrices, stored in STL vectors.
+    if( src.rows != dst.cols || src.cols != dst.rows )
+    {
+        CV_Assert( src.size() == dst.size() && (src.cols == 1 || src.rows == 1) );
+        src.copyTo(dst);
+        return;
+    }
+
+    CV_IPP_RUN_FAST(ipp_transpose(src, dst))
+
+    if( dst.data == src.data )
+    {
+        TransposeInplaceFunc func = transposeInplaceTab[esz];
+        CV_Assert( func != 0 );
+        CV_Assert( dst.cols == dst.rows );
+        func( dst.ptr(), dst.step, dst.rows );
+    }
+    else
+    {
+        TransposeFunc func = transposeTab[esz];
+        CV_Assert( func != 0 );
+        func( src.ptr(), src.step, dst.ptr(), dst.step, src.size() );
+    }
+}
+
+
+#if CV_SIMD128
+template<typename V> CV_ALWAYS_INLINE void flipHoriz_single( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size size, size_t esz )
+{
+    typedef typename V::lane_type T;
+    int end = (int)(size.width*esz);
+    int width = (end + 1)/2;
+    int width_1 = width & -v_uint8x16::nlanes;
+    int i, j;
+
+#if CV_STRONG_ALIGNMENT
+    CV_Assert(isAligned<sizeof(T)>(src, dst));
+#endif
+
+    for( ; size.height--; src += sstep, dst += dstep )
+    {
+        for( i = 0, j = end; i < width_1; i += v_uint8x16::nlanes, j -= v_uint8x16::nlanes )
+        {
+            V t0, t1;
+
+            t0 = v_load((T*)((uchar*)src + i));
+            t1 = v_load((T*)((uchar*)src + j - v_uint8x16::nlanes));
+            t0 = v_reverse(t0);
+            t1 = v_reverse(t1);
+            v_store((T*)(dst + j - v_uint8x16::nlanes), t0);
+            v_store((T*)(dst + i), t1);
+        }
+        if (isAligned<sizeof(T)>(src, dst))
+        {
+            for ( ; i < width; i += sizeof(T), j -= sizeof(T) )
+            {
+                T t0, t1;
+
+                t0 = *((T*)((uchar*)src + i));
+                t1 = *((T*)((uchar*)src + j - sizeof(T)));
+                *((T*)(dst + j - sizeof(T))) = t0;
+                *((T*)(dst + i)) = t1;
+            }
+        }
+        else
+        {
+            for ( ; i < width; i += sizeof(T), j -= sizeof(T) )
+            {
+                for (int k = 0; k < (int)sizeof(T); k++)
+                {
+                    uchar t0, t1;
+
+                    t0 = *((uchar*)src + i + k);
+                    t1 = *((uchar*)src + j + k - sizeof(T));
+                    *(dst + j + k - sizeof(T)) = t0;
+                    *(dst + i + k) = t1;
+                }
+            }
+        }
+    }
+}
+
+template<typename T1, typename T2> CV_ALWAYS_INLINE void flipHoriz_double( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size size, size_t esz )
+{
+    int end = (int)(size.width*esz);
+    int width = (end + 1)/2;
+
+#if CV_STRONG_ALIGNMENT
+    CV_Assert(isAligned<sizeof(T1)>(src, dst));
+    CV_Assert(isAligned<sizeof(T2)>(src, dst));
+#endif
+
+    for( ; size.height--; src += sstep, dst += dstep )
+    {
+        for ( int i = 0, j = end; i < width; i += sizeof(T1) + sizeof(T2), j -= sizeof(T1) + sizeof(T2) )
+        {
+            T1 t0, t1;
+            T2 t2, t3;
+
+            t0 = *((T1*)((uchar*)src + i));
+            t2 = *((T2*)((uchar*)src + i + sizeof(T1)));
+            t1 = *((T1*)((uchar*)src + j - sizeof(T1) - sizeof(T2)));
+            t3 = *((T2*)((uchar*)src + j - sizeof(T2)));
+            *((T1*)(dst + j - sizeof(T1) - sizeof(T2))) = t0;
+            *((T2*)(dst + j - sizeof(T2))) = t2;
+            *((T1*)(dst + i)) = t1;
+            *((T2*)(dst + i + sizeof(T1))) = t3;
+        }
+    }
+}
+#endif
+
+static void
+flipHoriz( const uchar* src, size_t sstep, uchar* dst, size_t dstep, Size size, size_t esz )
+{
+#if CV_SIMD
+#if CV_STRONG_ALIGNMENT
+    size_t alignmentMark = ((size_t)src)|((size_t)dst)|sstep|dstep;
+#endif
+    if (esz == 2 * v_uint8x16::nlanes)
+    {
+        int end = (int)(size.width*esz);
+        int width = end/2;
+
+        for( ; size.height--; src += sstep, dst += dstep )
+        {
+            for( int i = 0, j = end - 2 * v_uint8x16::nlanes; i < width; i += 2 * v_uint8x16::nlanes, j -= 2 * v_uint8x16::nlanes )
+            {
+#if CV_SIMD256
+                v_uint8x32 t0, t1;
+
+                t0 = v256_load((uchar*)src + i);
+                t1 = v256_load((uchar*)src + j);
+                v_store(dst + j, t0);
+                v_store(dst + i, t1);
+#else
+                v_uint8x16 t0, t1, t2, t3;
+
+                t0 = v_load((uchar*)src + i);
+                t1 = v_load((uchar*)src + i + v_uint8x16::nlanes);
+                t2 = v_load((uchar*)src + j);
+                t3 = v_load((uchar*)src + j + v_uint8x16::nlanes);
+                v_store(dst + j, t0);
+                v_store(dst + j + v_uint8x16::nlanes, t1);
+                v_store(dst + i, t2);
+                v_store(dst + i + v_uint8x16::nlanes, t3);
+#endif
+            }
+        }
+    }
+    else if (esz == v_uint8x16::nlanes)
+    {
+        int end = (int)(size.width*esz);
+        int width = end/2;
+
+        for( ; size.height--; src += sstep, dst += dstep )
+        {
+            for( int i = 0, j = end - v_uint8x16::nlanes; i < width; i += v_uint8x16::nlanes, j -= v_uint8x16::nlanes )
+            {
+                v_uint8x16 t0, t1;
+
+                t0 = v_load((uchar*)src + i);
+                t1 = v_load((uchar*)src + j);
+                v_store(dst + j, t0);
+                v_store(dst + i, t1);
+            }
+        }
+    }
+    else if (esz == 8
+#if CV_STRONG_ALIGNMENT
+            && isAligned<sizeof(uint64)>(alignmentMark)
+#endif
+    )
+    {
+        flipHoriz_single<v_uint64x2>(src, sstep, dst, dstep, size, esz);
+    }
+    else if (esz == 4
+#if CV_STRONG_ALIGNMENT
+            && isAligned<sizeof(unsigned)>(alignmentMark)
+#endif
+    )
+    {
+        flipHoriz_single<v_uint32x4>(src, sstep, dst, dstep, size, esz);
+    }
+    else if (esz == 2
+#if CV_STRONG_ALIGNMENT
+            && isAligned<sizeof(ushort)>(alignmentMark)
+#endif
+    )
+    {
+        flipHoriz_single<v_uint16x8>(src, sstep, dst, dstep, size, esz);
+    }
+    else if (esz == 1)
+    {
+        flipHoriz_single<v_uint8x16>(src, sstep, dst, dstep, size, esz);
+    }
+    else if (esz == 24
+#if CV_STRONG_ALIGNMENT
+            && isAligned<sizeof(uint64_t)>(alignmentMark)
+#endif
+    )
+    {
+        int end = (int)(size.width*esz);
+        int width = (end + 1)/2;
+
+        for( ; size.height--; src += sstep, dst += dstep )
+        {
+            for ( int i = 0, j = end; i < width; i += v_uint8x16::nlanes + sizeof(uint64_t), j -= v_uint8x16::nlanes + sizeof(uint64_t) )
+            {
+                v_uint8x16 t0, t1;
+                uint64_t t2, t3;
+
+                t0 = v_load((uchar*)src + i);
+                t2 = *((uint64_t*)((uchar*)src + i + v_uint8x16::nlanes));
+                t1 = v_load((uchar*)src + j - v_uint8x16::nlanes - sizeof(uint64_t));
+                t3 = *((uint64_t*)((uchar*)src + j - sizeof(uint64_t)));
+                v_store(dst + j - v_uint8x16::nlanes - sizeof(uint64_t), t0);
+                *((uint64_t*)(dst + j - sizeof(uint64_t))) = t2;
+                v_store(dst + i, t1);
+                *((uint64_t*)(dst + i + v_uint8x16::nlanes)) = t3;
+            }
+        }
+    }
+#if !CV_STRONG_ALIGNMENT
+    else if (esz == 12)
+    {
+        flipHoriz_double<uint64_t,uint>(src, sstep, dst, dstep, size, esz);
+    }
+    else if (esz == 6)
+    {
+        flipHoriz_double<uint,ushort>(src, sstep, dst, dstep, size, esz);
+    }
+    else if (esz == 3)
+    {
+        flipHoriz_double<ushort,uchar>(src, sstep, dst, dstep, size, esz);
+    }
+#endif
+    else
+#endif // CV_SIMD
+    {
+        int i, j, limit = (int)(((size.width + 1)/2)*esz);
+        AutoBuffer<int> _tab(size.width*esz);
+        int* tab = _tab.data();
+
+        for( i = 0; i < size.width; i++ )
+            for( size_t k = 0; k < esz; k++ )
+                tab[i*esz + k] = (int)((size.width - i - 1)*esz + k);
+
+        for( ; size.height--; src += sstep, dst += dstep )
+        {
+            for( i = 0; i < limit; i++ )
+            {
+                j = tab[i];
+                uchar t0 = src[i], t1 = src[j];
+                dst[i] = t1; dst[j] = t0;
+            }
+        }
+    }
+}
+
+static void
+flipVert( const uchar* src0, size_t sstep, uchar* dst0, size_t dstep, Size size, size_t esz )
+{
+    const uchar* src1 = src0 + (size.height - 1)*sstep;
+    uchar* dst1 = dst0 + (size.height - 1)*dstep;
+    size.width *= (int)esz;
+
+    for( int y = 0; y < (size.height + 1)/2; y++, src0 += sstep, src1 -= sstep,
+                                                  dst0 += dstep, dst1 -= dstep )
+    {
+        int i = 0;
+#if CV_SIMD
+#if CV_STRONG_ALIGNMENT
+        if (isAligned<sizeof(int)>(src0, src1, dst0, dst1))
+#endif
+        {
+            for (; i <= size.width - CV_SIMD_WIDTH; i += CV_SIMD_WIDTH)
+            {
+                v_int32 t0 = vx_load((int*)(src0 + i));
+                v_int32 t1 = vx_load((int*)(src1 + i));
+                v_store((int*)(dst0 + i), t1);
+                v_store((int*)(dst1 + i), t0);
+            }
+        }
+#if CV_STRONG_ALIGNMENT
+        else
+        {
+            for (; i <= size.width - CV_SIMD_WIDTH; i += CV_SIMD_WIDTH)
+            {
+                v_uint8 t0 = vx_load(src0 + i);
+                v_uint8 t1 = vx_load(src1 + i);
+                v_store(dst0 + i, t1);
+                v_store(dst1 + i, t0);
+            }
+        }
+#endif
+#endif
+
+        if (isAligned<sizeof(int)>(src0, src1, dst0, dst1))
+        {
+            for( ; i <= size.width - 16; i += 16 )
+            {
+                int t0 = ((int*)(src0 + i))[0];
+                int t1 = ((int*)(src1 + i))[0];
+
+                ((int*)(dst0 + i))[0] = t1;
+                ((int*)(dst1 + i))[0] = t0;
+
+                t0 = ((int*)(src0 + i))[1];
+                t1 = ((int*)(src1 + i))[1];
+
+                ((int*)(dst0 + i))[1] = t1;
+                ((int*)(dst1 + i))[1] = t0;
+
+                t0 = ((int*)(src0 + i))[2];
+                t1 = ((int*)(src1 + i))[2];
+
+                ((int*)(dst0 + i))[2] = t1;
+                ((int*)(dst1 + i))[2] = t0;
+
+                t0 = ((int*)(src0 + i))[3];
+                t1 = ((int*)(src1 + i))[3];
+
+                ((int*)(dst0 + i))[3] = t1;
+                ((int*)(dst1 + i))[3] = t0;
+            }
+
+            for( ; i <= size.width - 4; i += 4 )
+            {
+                int t0 = ((int*)(src0 + i))[0];
+                int t1 = ((int*)(src1 + i))[0];
+
+                ((int*)(dst0 + i))[0] = t1;
+                ((int*)(dst1 + i))[0] = t0;
+            }
+        }
+
+        for( ; i < size.width; i++ )
+        {
+            uchar t0 = src0[i];
+            uchar t1 = src1[i];
+
+            dst0[i] = t1;
+            dst1[i] = t0;
+        }
+    }
+}
+
+#ifdef HAVE_OPENCL
+
+enum { FLIP_COLS = 1 << 0, FLIP_ROWS = 1 << 1, FLIP_BOTH = FLIP_ROWS | FLIP_COLS };
+
+static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode )
+{
+    CV_Assert(flipCode >= -1 && flipCode <= 1);
+
+    const ocl::Device & dev = ocl::Device::getDefault();
+    int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type),
+            flipType, kercn = std::min(ocl::predictOptimalVectorWidth(_src, _dst), 4);
+
+    bool doubleSupport = dev.doubleFPConfig() > 0;
+    if (!doubleSupport && depth == CV_64F)
+        kercn = cn;
+
+    if (cn > 4)
+        return false;
+
+    const char * kernelName;
+    if (flipCode == 0)
+        kernelName = "arithm_flip_rows", flipType = FLIP_ROWS;
+    else if (flipCode > 0)
+        kernelName = "arithm_flip_cols", flipType = FLIP_COLS;
+    else
+        kernelName = "arithm_flip_rows_cols", flipType = FLIP_BOTH;
+
+    int pxPerWIy = (dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU)) ? 4 : 1;
+    kercn = (cn!=3 || flipType == FLIP_ROWS) ? std::max(kercn, cn) : cn;
+
+    ocl::Kernel k(kernelName, ocl::core::flip_oclsrc,
+        format( "-D T=%s -D T1=%s -D DEPTH=%d -D cn=%d -D PIX_PER_WI_Y=%d -D kercn=%d",
+                kercn != cn ? ocl::typeToStr(CV_MAKE_TYPE(depth, kercn)) : ocl::vecopTypeToStr(CV_MAKE_TYPE(depth, kercn)),
+                kercn != cn ? ocl::typeToStr(depth) : ocl::vecopTypeToStr(depth), depth, cn, pxPerWIy, kercn));
+    if (k.empty())
+        return false;
+
+    Size size = _src.size();
+    _dst.create(size, type);
+    UMat src = _src.getUMat(), dst = _dst.getUMat();
+
+    int cols = size.width * cn / kercn, rows = size.height;
+    cols = flipType == FLIP_COLS ? (cols + 1) >> 1 : cols;
+    rows = flipType & FLIP_ROWS ? (rows + 1) >> 1 : rows;
+
+    k.args(ocl::KernelArg::ReadOnlyNoSize(src),
+           ocl::KernelArg::WriteOnly(dst, cn, kercn), rows, cols);
+
+    size_t maxWorkGroupSize = dev.maxWorkGroupSize();
+    CV_Assert(maxWorkGroupSize % 4 == 0);
+
+    size_t globalsize[2] = { (size_t)cols, ((size_t)rows + pxPerWIy - 1) / pxPerWIy },
+            localsize[2] = { maxWorkGroupSize / 4, 4 };
+    return k.run(2, globalsize, (flipType == FLIP_COLS) && !dev.isIntel() ? localsize : NULL, false);
+}
+
+#endif
+
+#if defined HAVE_IPP
+static bool ipp_flip(Mat &src, Mat &dst, int flip_mode)
+{
+#ifdef HAVE_IPP_IW
+    CV_INSTRUMENT_REGION_IPP();
+
+    // Details: https://github.com/opencv/opencv/issues/12943
+    if (flip_mode <= 0 /* swap rows */
+        && cv::ipp::getIppTopFeatures() != ippCPUID_SSE42
+        && (int64_t)(src.total()) * src.elemSize() >= CV_BIG_INT(0x80000000)/*2Gb*/
+    )
+        return false;
+
+    IppiAxis ippMode;
+    if(flip_mode < 0)
+        ippMode = ippAxsBoth;
+    else if(flip_mode == 0)
+        ippMode = ippAxsHorizontal;
+    else
+        ippMode = ippAxsVertical;
+
+    try
+    {
+        ::ipp::IwiImage iwSrc = ippiGetImage(src);
+        ::ipp::IwiImage iwDst = ippiGetImage(dst);
+
+        CV_INSTRUMENT_FUN_IPP(::ipp::iwiMirror, iwSrc, iwDst, ippMode);
+    }
+    catch(const ::ipp::IwException &)
+    {
+        return false;
+    }
+
+    return true;
+#else
+    CV_UNUSED(src); CV_UNUSED(dst); CV_UNUSED(flip_mode);
+    return false;
+#endif
+}
+#endif
+
+
+void flip( InputArray _src, OutputArray _dst, int flip_mode )
+{
+    CV_INSTRUMENT_REGION();
+
+    CV_Assert( _src.dims() <= 2 );
+    Size size = _src.size();
+
+    if (flip_mode < 0)
+    {
+        if (size.width == 1)
+            flip_mode = 0;
+        if (size.height == 1)
+            flip_mode = 1;
+    }
+
+    if ((size.width == 1 && flip_mode > 0) ||
+        (size.height == 1 && flip_mode == 0))
+    {
+        return _src.copyTo(_dst);
+    }
+
+    CV_OCL_RUN( _dst.isUMat(), ocl_flip(_src, _dst, flip_mode))
+
+    Mat src = _src.getMat();
+    int type = src.type();
+    _dst.create( size, type );
+    Mat dst = _dst.getMat();
+
+    CV_IPP_RUN_FAST(ipp_flip(src, dst, flip_mode));
+
+    size_t esz = CV_ELEM_SIZE(type);
+
+    if( flip_mode <= 0 )
+        flipVert( src.ptr(), src.step, dst.ptr(), dst.step, src.size(), esz );
+    else
+        flipHoriz( src.ptr(), src.step, dst.ptr(), dst.step, src.size(), esz );
+
+    if( flip_mode < 0 )
+        flipHoriz( dst.ptr(), dst.step, dst.ptr(), dst.step, dst.size(), esz );
+}
+
+void rotate(InputArray _src, OutputArray _dst, int rotateMode)
+{
+    CV_Assert(_src.dims() <= 2);
+
+    switch (rotateMode)
+    {
+    case ROTATE_90_CLOCKWISE:
+        transpose(_src, _dst);
+        flip(_dst, _dst, 1);
+        break;
+    case ROTATE_180:
+        flip(_src, _dst, -1);
+        break;
+    case ROTATE_90_COUNTERCLOCKWISE:
+        transpose(_src, _dst);
+        flip(_dst, _dst, 0);
+        break;
+    default:
+        break;
+    }
+}
+
+}  // namespace

modules/core/src/matrix_wrap.cpp

@@ -2,9 +2,8 @@
 // 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
 
-
-#include "opencv2/core/mat.hpp"
 #include "precomp.hpp"
+#include "opencv2/core/mat.hpp"
 
 namespace cv {
 
@@ -33,7 +32,7 @@ Mat _InputArray::getMat_(int i) const
         return m->getMat(accessFlags).row(i);
     }
 
-    if( k == MATX || k == STD_ARRAY )
+    if (k == MATX)
     {
         CV_Assert( i < 0 );
         return Mat(sz, flags, obj);
@@ -173,7 +172,7 @@ void _InputArray::getMatVector(std::vector<Mat>& mv) const
         return;
     }
 
-    if( k == MATX || k == STD_ARRAY )
+    if (k == MATX)
     {
         size_t n = sz.height, esz = CV_ELEM_SIZE(flags);
         mv.resize(n);
@@ -317,6 +316,7 @@ void _InputArray::getUMatVector(std::vector<UMat>& umv) const
 
 cuda::GpuMat _InputArray::getGpuMat() const
 {
+#ifdef HAVE_CUDA
     _InputArray::KindFlag k = kind();
 
     if (k == CUDA_GPU_MAT)
@@ -340,14 +340,22 @@ cuda::GpuMat _InputArray::getGpuMat() const
         return cuda::GpuMat();
 
     CV_Error(cv::Error::StsNotImplemented, "getGpuMat is available only for cuda::GpuMat and cuda::HostMem");
+#else
+    CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
 }
 void _InputArray::getGpuMatVector(std::vector<cuda::GpuMat>& gpumv) const
 {
+#ifdef HAVE_CUDA
     _InputArray::KindFlag k = kind();
     if (k == STD_VECTOR_CUDA_GPU_MAT)
     {
         gpumv = *(std::vector<cuda::GpuMat>*)obj;
     }
+#else
+    CV_UNUSED(gpumv);
+    CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
 }
 ogl::Buffer _InputArray::getOGlBuffer() const
 {
@@ -362,7 +370,10 @@ ogl::Buffer _InputArray::getOGlBuffer() const
 _InputArray::KindFlag _InputArray::kind() const
 {
     KindFlag k = flags & KIND_MASK;
+#if CV_VERSION_MAJOR < 5
     CV_DbgAssert(k != EXPR);
+    CV_DbgAssert(k != STD_ARRAY);
+#endif
     return k;
 }
 
@@ -392,7 +403,7 @@ Size _InputArray::size(int i) const
         return ((const UMat*)obj)->size();
     }
 
-    if( k == MATX || k == STD_ARRAY )
+    if (k == MATX)
     {
         CV_Assert( i < 0 );
         return sz;
@@ -451,11 +462,15 @@ Size _InputArray::size(int i) const
 
     if (k == STD_VECTOR_CUDA_GPU_MAT)
     {
+#ifdef HAVE_CUDA
         const std::vector<cuda::GpuMat>& vv = *(const std::vector<cuda::GpuMat>*)obj;
         if (i < 0)
             return vv.empty() ? Size() : Size((int)vv.size(), 1);
         CV_Assert(i < (int)vv.size());
         return vv[i].size();
+#else
+        CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
     }
 
     if( k == STD_VECTOR_UMAT )
@@ -612,7 +627,7 @@ int _InputArray::dims(int i) const
         return ((const UMat*)obj)->dims;
     }
 
-    if( k == MATX || k == STD_ARRAY )
+    if (k == MATX)
     {
         CV_Assert( i < 0 );
         return 2;
@@ -746,7 +761,7 @@ int _InputArray::type(int i) const
     if( k == UMAT )
         return ((const UMat*)obj)->type();
 
-    if( k == MATX || k == STD_VECTOR || k == STD_ARRAY || k == STD_VECTOR_VECTOR || k == STD_BOOL_VECTOR )
+    if( k == MATX || k == STD_VECTOR || k == STD_VECTOR_VECTOR || k == STD_BOOL_VECTOR )
         return CV_MAT_TYPE(flags);
 
     if( k == NONE )
@@ -790,6 +805,7 @@ int _InputArray::type(int i) const
 
     if (k == STD_VECTOR_CUDA_GPU_MAT)
     {
+#ifdef HAVE_CUDA
         const std::vector<cuda::GpuMat>& vv = *(const std::vector<cuda::GpuMat>*)obj;
         if (vv.empty())
         {
@@ -798,6 +814,9 @@ int _InputArray::type(int i) const
         }
         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 )
@@ -832,7 +851,7 @@ bool _InputArray::empty() const
     if( k == UMAT )
         return ((const UMat*)obj)->empty();
 
-    if( k == MATX || k == STD_ARRAY )
+    if (k == MATX)
         return false;
 
     if( k == STD_VECTOR )
@@ -901,7 +920,7 @@ bool _InputArray::isContinuous(int i) const
     if( k == UMAT )
         return i < 0 ? ((const UMat*)obj)->isContinuous() : true;
 
-    if( k == MATX || k == STD_VECTOR || k == STD_ARRAY ||
+    if( k == MATX || k == STD_VECTOR ||
         k == NONE || k == STD_VECTOR_VECTOR || k == STD_BOOL_VECTOR )
         return true;
 
@@ -942,7 +961,7 @@ bool _InputArray::isSubmatrix(int i) const
     if( k == UMAT )
         return i < 0 ? ((const UMat*)obj)->isSubmatrix() : false;
 
-    if( k == MATX || k == STD_VECTOR || k == STD_ARRAY ||
+    if( k == MATX || k == STD_VECTOR ||
         k == NONE || k == STD_VECTOR_VECTOR || k == STD_BOOL_VECTOR )
         return false;
 
@@ -987,7 +1006,7 @@ size_t _InputArray::offset(int i) const
         return ((const UMat*)obj)->offset;
     }
 
-    if( k == MATX || k == STD_VECTOR || k == STD_ARRAY ||
+    if( k == MATX || k == STD_VECTOR ||
         k == NONE || k == STD_VECTOR_VECTOR || k == STD_BOOL_VECTOR )
         return 0;
 
@@ -1046,7 +1065,7 @@ size_t _InputArray::step(int i) const
         return ((const UMat*)obj)->step;
     }
 
-    if( k == MATX || k == STD_VECTOR || k == STD_ARRAY ||
+    if( k == MATX || k == STD_VECTOR ||
         k == NONE || k == STD_VECTOR_VECTOR || k == STD_BOOL_VECTOR )
         return 0;
 
@@ -1092,7 +1111,7 @@ void _InputArray::copyTo(const _OutputArray& arr) const
 
     if( k == NONE )
         arr.release();
-    else if( k == MAT || k == MATX || k == STD_VECTOR || k == STD_ARRAY || k == STD_BOOL_VECTOR )
+    else if( k == MAT || k == MATX || k == STD_VECTOR || k == STD_BOOL_VECTOR )
     {
         Mat m = getMat();
         m.copyTo(arr);
@@ -1113,7 +1132,7 @@ void _InputArray::copyTo(const _OutputArray& arr, const _InputArray & mask) cons
 
     if( k == NONE )
         arr.release();
-    else if( k == MAT || k == MATX || k == STD_VECTOR || k == STD_ARRAY || k == STD_BOOL_VECTOR )
+    else if( k == MAT || k == MATX || k == STD_VECTOR || k == STD_BOOL_VECTOR )
     {
         Mat m = getMat();
         m.copyTo(arr, mask);
@@ -1159,22 +1178,34 @@ void _OutputArray::create(Size _sz, int mtype, int i, bool allowTransposed, _Out
     {
         CV_Assert(!fixedSize() || ((cuda::GpuMat*)obj)->size() == _sz);
         CV_Assert(!fixedType() || ((cuda::GpuMat*)obj)->type() == mtype);
+#ifdef HAVE_CUDA
         ((cuda::GpuMat*)obj)->create(_sz, mtype);
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
     }
     if( k == OPENGL_BUFFER && i < 0 && !allowTransposed && fixedDepthMask == 0 )
     {
         CV_Assert(!fixedSize() || ((ogl::Buffer*)obj)->size() == _sz);
         CV_Assert(!fixedType() || ((ogl::Buffer*)obj)->type() == mtype);
+#ifdef HAVE_OPENGL
         ((ogl::Buffer*)obj)->create(_sz, mtype);
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "OpenGL support is not enabled in this OpenCV build (missing HAVE_OPENGL)");
+#endif
     }
     if( k == CUDA_HOST_MEM && i < 0 && !allowTransposed && fixedDepthMask == 0 )
     {
         CV_Assert(!fixedSize() || ((cuda::HostMem*)obj)->size() == _sz);
         CV_Assert(!fixedType() || ((cuda::HostMem*)obj)->type() == mtype);
+#ifdef HAVE_CUDA
         ((cuda::HostMem*)obj)->create(_sz, mtype);
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
     }
     int sizes[] = {_sz.height, _sz.width};
     create(2, sizes, mtype, i, allowTransposed, fixedDepthMask);
@@ -1201,22 +1232,34 @@ void _OutputArray::create(int _rows, int _cols, int mtype, int i, bool allowTran
     {
         CV_Assert(!fixedSize() || ((cuda::GpuMat*)obj)->size() == Size(_cols, _rows));
         CV_Assert(!fixedType() || ((cuda::GpuMat*)obj)->type() == mtype);
+#ifdef HAVE_CUDA
         ((cuda::GpuMat*)obj)->create(_rows, _cols, mtype);
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
     }
     if( k == OPENGL_BUFFER && i < 0 && !allowTransposed && fixedDepthMask == 0 )
     {
         CV_Assert(!fixedSize() || ((ogl::Buffer*)obj)->size() == Size(_cols, _rows));
         CV_Assert(!fixedType() || ((ogl::Buffer*)obj)->type() == mtype);
+#ifdef HAVE_OPENGL
         ((ogl::Buffer*)obj)->create(_rows, _cols, mtype);
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "OpenGL support is not enabled in this OpenCV build (missing HAVE_OPENGL)");
+#endif
     }
     if( k == CUDA_HOST_MEM && i < 0 && !allowTransposed && fixedDepthMask == 0 )
     {
         CV_Assert(!fixedSize() || ((cuda::HostMem*)obj)->size() == Size(_cols, _rows));
         CV_Assert(!fixedType() || ((cuda::HostMem*)obj)->type() == mtype);
+#ifdef HAVE_CUDA
         ((cuda::HostMem*)obj)->create(_rows, _cols, mtype);
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
     }
     int sizes[] = {_rows, _cols};
     create(2, sizes, mtype, i, allowTransposed, fixedDepthMask);
@@ -1301,16 +1344,27 @@ void _OutputArray::create(int d, const int* sizes, int mtype, int i,
         CV_Assert( i < 0 );
         int type0 = CV_MAT_TYPE(flags);
         CV_Assert( mtype == type0 || (CV_MAT_CN(mtype) == 1 && ((1 << type0) & fixedDepthMask) != 0) );
-        CV_Assert( d == 2 && ((sizes[0] == sz.height && sizes[1] == sz.width) ||
-                                 (allowTransposed && sizes[0] == sz.width && sizes[1] == sz.height)));
-        return;
-    }
-
-    if( k == STD_ARRAY )
-    {
-        int type0 = CV_MAT_TYPE(flags);
-        CV_Assert( mtype == type0 || (CV_MAT_CN(mtype) == 1 && ((1 << type0) & fixedDepthMask) != 0) );
-        CV_Assert( d == 2 && sz.area() == sizes[0]*sizes[1]);
+        CV_CheckLE(d, 2, "");
+        Size requested_size(d == 2 ? sizes[1] : 1, d >= 1 ? sizes[0] : 1);
+        if (sz.width == 1 || sz.height == 1)
+        {
+            // NB: 1D arrays assume allowTransposed=true (see #4159)
+            int total_1d = std::max(sz.width, sz.height);
+            CV_Check(requested_size, std::max(requested_size.width, requested_size.height) == total_1d, "");
+        }
+        else
+        {
+            if (!allowTransposed)
+            {
+                CV_CheckEQ(requested_size, sz, "");
+            }
+            else
+            {
+                CV_Check(requested_size,
+                        (requested_size == sz || (requested_size.height == sz.width && requested_size.width == sz.height)),
+                        "");
+            }
+        }
         return;
     }
 
@@ -1628,20 +1682,32 @@ void _OutputArray::release() const
 
     if( k == CUDA_GPU_MAT )
     {
+#ifdef HAVE_CUDA
         ((cuda::GpuMat*)obj)->release();
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
     }
 
     if( k == CUDA_HOST_MEM )
     {
+#ifdef HAVE_CUDA
         ((cuda::HostMem*)obj)->release();
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
     }
 
     if( k == OPENGL_BUFFER )
     {
+#ifdef HAVE_OPENGL
         ((ogl::Buffer*)obj)->release();
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "OpenGL support is not enabled in this OpenCV build (missing HAVE_OPENGL)");
+#endif
     }
 
     if( k == NONE )
@@ -1672,8 +1738,12 @@ void _OutputArray::release() const
     }
     if (k == STD_VECTOR_CUDA_GPU_MAT)
     {
+#ifdef HAVE_CUDA
         ((std::vector<cuda::GpuMat>*)obj)->clear();
         return;
+#else
+        CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
     }
     CV_Error(Error::StsNotImplemented, "Unknown/unsupported array type");
 }
@@ -1772,7 +1842,7 @@ void _OutputArray::setTo(const _InputArray& arr, const _InputArray & mask) const
 
     if( k == NONE )
         ;
-    else if( k == MAT || k == MATX || k == STD_VECTOR || k == STD_ARRAY )
+    else if (k == MAT || k == MATX || k == STD_VECTOR)
     {
         Mat m = getMat();
         m.setTo(arr, mask);
@@ -1781,9 +1851,13 @@ void _OutputArray::setTo(const _InputArray& arr, const _InputArray & mask) const
         ((UMat*)obj)->setTo(arr, mask);
     else if( k == CUDA_GPU_MAT )
     {
+#ifdef HAVE_CUDA
         Mat value = arr.getMat();
         CV_Assert( checkScalar(value, type(), arr.kind(), _InputArray::CUDA_GPU_MAT) );
         ((cuda::GpuMat*)obj)->setTo(Scalar(Vec<double, 4>(value.ptr<double>())), mask);
+#else
+        CV_Error(Error::StsNotImplemented, "CUDA support is not enabled in this OpenCV build (missing HAVE_CUDA)");
+#endif
     }
     else
         CV_Error(Error::StsNotImplemented, "");

modules/core/src/norm.cpp

@@ -152,10 +152,10 @@ float normL2Sqr_(const float* a, const float* b, int n)
     {
         v_float32 t0 = vx_load(a + j) - vx_load(b + j);
         v_float32 t1 = vx_load(a + j + v_float32::nlanes) - vx_load(b + j + v_float32::nlanes);
-        v_float32 t2 = vx_load(a + j + 2 * v_float32::nlanes) - vx_load(b + j + 2 * v_float32::nlanes);
-        v_float32 t3 = vx_load(a + j + 3 * v_float32::nlanes) - vx_load(b + j + 3 * v_float32::nlanes);
         v_d0 = v_muladd(t0, t0, v_d0);
+        v_float32 t2 = vx_load(a + j + 2 * v_float32::nlanes) - vx_load(b + j + 2 * v_float32::nlanes);
         v_d1 = v_muladd(t1, t1, v_d1);
+        v_float32 t3 = vx_load(a + j + 3 * v_float32::nlanes) - vx_load(b + j + 3 * v_float32::nlanes);
         v_d2 = v_muladd(t2, t2, v_d2);
         v_d3 = v_muladd(t3, t3, v_d3);
     }
@@ -205,13 +205,10 @@ int normL1_(const uchar* a, const uchar* b, int n)
     return d;
 }
 
-}} //cv::hal
+} //cv::hal
 
 //==================================================================================================
 
-namespace cv
-{
-
 template<typename T, typename ST> int
 normInf_(const T* src, const uchar* mask, ST* _result, int len, int cn)
 {
@@ -594,12 +591,10 @@ static bool ipp_norm(Mat &src, int normType, Mat &mask, double &result)
     CV_UNUSED(src); CV_UNUSED(normType); CV_UNUSED(mask); CV_UNUSED(result);
 #endif
     return false;
-}
-#endif
+}  // ipp_norm()
+#endif  // HAVE_IPP
 
-} // cv::
-
-double cv::norm( InputArray _src, int normType, InputArray _mask )
+double norm( InputArray _src, int normType, InputArray _mask )
 {
     CV_INSTRUMENT_REGION();
 
@@ -792,9 +787,6 @@ double cv::norm( InputArray _src, int normType, InputArray _mask )
 //==================================================================================================
 
 #ifdef HAVE_OPENCL
-
-namespace cv {
-
 static bool ocl_norm( InputArray _src1, InputArray _src2, int normType, InputArray _mask, double & result )
 {
 #ifdef __ANDROID__
@@ -849,15 +841,10 @@ static bool ocl_norm( InputArray _src1, InputArray _src2, int normType, InputArr
         result /= (s2 + DBL_EPSILON);
 
     return true;
-}
-
-}
-
-#endif
+}  // ocl_norm()
+#endif  // HAVE_OPENCL
 
 #ifdef HAVE_IPP
-namespace cv
-{
 static bool ipp_norm(InputArray _src1, InputArray _src2, int normType, InputArray _mask, double &result)
 {
     CV_INSTRUMENT_REGION_IPP();
@@ -1083,12 +1070,11 @@ static bool ipp_norm(InputArray _src1, InputArray _src2, int normType, InputArra
     CV_UNUSED(_src1); CV_UNUSED(_src2); CV_UNUSED(normType); CV_UNUSED(_mask); CV_UNUSED(result);
 #endif
     return false;
-}
-}
-#endif
+}  // ipp_norm
+#endif  // HAVE_IPP
 
 
-double cv::norm( InputArray _src1, InputArray _src2, int normType, InputArray _mask )
+double norm( InputArray _src1, InputArray _src2, int normType, InputArray _mask )
 {
     CV_INSTRUMENT_REGION();
 
@@ -1280,12 +1266,12 @@ double cv::norm( InputArray _src1, InputArray _src2, int normType, InputArray _m
     return result.d;
 }
 
-cv::Hamming::ResultType cv::Hamming::operator()( const unsigned char* a, const unsigned char* b, int size ) const
+cv::Hamming::ResultType Hamming::operator()( const unsigned char* a, const unsigned char* b, int size ) const
 {
     return cv::hal::normHamming(a, b, size);
 }
 
-double cv::PSNR(InputArray _src1, InputArray _src2, double R)
+double PSNR(InputArray _src1, InputArray _src2, double R)
 {
     CV_INSTRUMENT_REGION();
 
@@ -1295,3 +1281,141 @@ double cv::PSNR(InputArray _src1, InputArray _src2, double R)
     double diff = std::sqrt(norm(_src1, _src2, NORM_L2SQR)/(_src1.total()*_src1.channels()));
     return 20*log10(R/(diff+DBL_EPSILON));
 }
+
+
+#ifdef HAVE_OPENCL
+static bool ocl_normalize( InputArray _src, InputOutputArray _dst, InputArray _mask, int dtype,
+                           double scale, double delta )
+{
+    UMat src = _src.getUMat();
+
+    if( _mask.empty() )
+        src.convertTo( _dst, dtype, scale, delta );
+    else if (src.channels() <= 4)
+    {
+        const ocl::Device & dev = ocl::Device::getDefault();
+
+        int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype),
+                ddepth = CV_MAT_DEPTH(dtype), wdepth = std::max(CV_32F, std::max(sdepth, ddepth)),
+                rowsPerWI = dev.isIntel() ? 4 : 1;
+
+        float fscale = static_cast<float>(scale), fdelta = static_cast<float>(delta);
+        bool haveScale = std::fabs(scale - 1) > DBL_EPSILON,
+                haveZeroScale = !(std::fabs(scale) > DBL_EPSILON),
+                haveDelta = std::fabs(delta) > DBL_EPSILON,
+                doubleSupport = dev.doubleFPConfig() > 0;
+
+        if (!haveScale && !haveDelta && stype == dtype)
+        {
+            _src.copyTo(_dst, _mask);
+            return true;
+        }
+        if (haveZeroScale)
+        {
+            _dst.setTo(Scalar(delta), _mask);
+            return true;
+        }
+
+        if ((sdepth == CV_64F || ddepth == CV_64F) && !doubleSupport)
+            return false;
+
+        char cvt[2][40];
+        String opts = format("-D srcT=%s -D dstT=%s -D convertToWT=%s -D cn=%d -D rowsPerWI=%d"
+                             " -D convertToDT=%s -D workT=%s%s%s%s -D srcT1=%s -D dstT1=%s",
+                             ocl::typeToStr(stype), ocl::typeToStr(dtype),
+                             ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]), cn,
+                             rowsPerWI, ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]),
+                             ocl::typeToStr(CV_MAKE_TYPE(wdepth, cn)),
+                             doubleSupport ? " -D DOUBLE_SUPPORT" : "",
+                             haveScale ? " -D HAVE_SCALE" : "",
+                             haveDelta ? " -D HAVE_DELTA" : "",
+                             ocl::typeToStr(sdepth), ocl::typeToStr(ddepth));
+
+        ocl::Kernel k("normalizek", ocl::core::normalize_oclsrc, opts);
+        if (k.empty())
+            return false;
+
+        UMat mask = _mask.getUMat(), dst = _dst.getUMat();
+
+        ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
+                maskarg = ocl::KernelArg::ReadOnlyNoSize(mask),
+                dstarg = ocl::KernelArg::ReadWrite(dst);
+
+        if (haveScale)
+        {
+            if (haveDelta)
+                k.args(srcarg, maskarg, dstarg, fscale, fdelta);
+            else
+                k.args(srcarg, maskarg, dstarg, fscale);
+        }
+        else
+        {
+            if (haveDelta)
+                k.args(srcarg, maskarg, dstarg, fdelta);
+            else
+                k.args(srcarg, maskarg, dstarg);
+        }
+
+        size_t globalsize[2] = { (size_t)src.cols, ((size_t)src.rows + rowsPerWI - 1) / rowsPerWI };
+        return k.run(2, globalsize, NULL, false);
+    }
+    else
+    {
+        UMat temp;
+        src.convertTo( temp, dtype, scale, delta );
+        temp.copyTo( _dst, _mask );
+    }
+
+    return true;
+}  // ocl_normalize
+#endif  // HAVE_OPENCL
+
+void normalize(InputArray _src, InputOutputArray _dst, double a, double b,
+               int norm_type, int rtype, InputArray _mask)
+{
+    CV_INSTRUMENT_REGION();
+
+    double scale = 1, shift = 0;
+    int type = _src.type(), depth = CV_MAT_DEPTH(type);
+
+    if( rtype < 0 )
+        rtype = _dst.fixedType() ? _dst.depth() : depth;
+
+    if( norm_type == CV_MINMAX )
+    {
+        double smin = 0, smax = 0;
+        double dmin = MIN( a, b ), dmax = MAX( a, b );
+        minMaxIdx( _src, &smin, &smax, 0, 0, _mask );
+        scale = (dmax - dmin)*(smax - smin > DBL_EPSILON ? 1./(smax - smin) : 0);
+        if( rtype == CV_32F )
+        {
+            scale = (float)scale;
+            shift = (float)dmin - (float)(smin*scale);
+        }
+        else
+            shift = dmin - smin*scale;
+    }
+    else if( norm_type == CV_L2 || norm_type == CV_L1 || norm_type == CV_C )
+    {
+        scale = norm( _src, norm_type, _mask );
+        scale = scale > DBL_EPSILON ? a/scale : 0.;
+        shift = 0;
+    }
+    else
+        CV_Error( CV_StsBadArg, "Unknown/unsupported norm type" );
+
+    CV_OCL_RUN(_dst.isUMat(),
+               ocl_normalize(_src, _dst, _mask, rtype, scale, shift))
+
+    Mat src = _src.getMat();
+    if( _mask.empty() )
+        src.convertTo( _dst, rtype, scale, shift );
+    else
+    {
+        Mat temp;
+        src.convertTo( temp, rtype, scale, shift );
+        temp.copyTo( _dst, _mask );
+    }
+}
+
+}  // namespace

modules/core/src/ocl.cpp

@@ -1149,14 +1149,14 @@ void OpenCLExecutionContext::release()
 }
 
 
+
 // true if we have initialized OpenCL subsystem with available platforms
-static bool g_isOpenCVActivated = false;
+static bool g_isOpenCLInitialized = false;
+static bool g_isOpenCLAvailable = false;
 
 bool haveOpenCL()
 {
     CV_TRACE_FUNCTION();
-    static bool g_isOpenCLInitialized = false;
-    static bool g_isOpenCLAvailable = false;
 
     if (!g_isOpenCLInitialized)
     {
@@ -1178,7 +1178,7 @@ bool haveOpenCL()
         {
             cl_uint n = 0;
             g_isOpenCLAvailable = ::clGetPlatformIDs(0, NULL, &n) == CL_SUCCESS;
-            g_isOpenCVActivated = n > 0;
+            g_isOpenCLAvailable &= n > 0;
             CV_LOG_INFO(NULL, "OpenCL: found " << n << " platforms");
         }
         catch (...)
@@ -1214,7 +1214,7 @@ bool useOpenCL()
 
 bool isOpenCLActivated()
 {
-    if (!g_isOpenCVActivated)
+    if (!g_isOpenCLAvailable)
         return false; // prevent unnecessary OpenCL activation via useOpenCL()->haveOpenCL() calls
     return useOpenCL();
 }
@@ -1451,7 +1451,7 @@ struct Platform::Impl
     bool initialized;
 };
 
-Platform::Platform()
+Platform::Platform() CV_NOEXCEPT
 {
     p = 0;
 }
@@ -1480,6 +1480,23 @@ Platform& Platform::operator = (const Platform& pl)
     return *this;
 }
 
+Platform::Platform(Platform&& pl) CV_NOEXCEPT
+{
+    p = pl.p;
+    pl.p = nullptr;
+}
+
+Platform& Platform::operator = (Platform&& pl) CV_NOEXCEPT
+{
+    if (this != &pl) {
+        if(p)
+            p->release();
+        p = pl.p;
+        pl.p = nullptr;
+    }
+    return *this;
+}
+
 void* Platform::ptr() const
 {
     return p ? p->handle : 0;
@@ -1499,25 +1516,27 @@ Platform& Platform::getDefault()
 
 /////////////////////////////////////// Device ////////////////////////////////////////////
 
-// deviceVersion has format
+// Version has format:
 //   OpenCL<space><major_version.minor_version><space><vendor-specific information>
 // by specification
 //   http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetDeviceInfo.html
 //   http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clGetDeviceInfo.html
-static void parseDeviceVersion(const String &deviceVersion, int &major, int &minor)
+//   https://www.khronos.org/registry/OpenCL/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html
+//   https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/clGetPlatformInfo.html
+static void parseOpenCLVersion(const String &version, int &major, int &minor)
 {
     major = minor = 0;
-    if (10 >= deviceVersion.length())
+    if (10 >= version.length())
         return;
-    const char *pstr = deviceVersion.c_str();
+    const char *pstr = version.c_str();
     if (0 != strncmp(pstr, "OpenCL ", 7))
         return;
-    size_t ppos = deviceVersion.find('.', 7);
+    size_t ppos = version.find('.', 7);
     if (String::npos == ppos)
         return;
-    String temp = deviceVersion.substr(7, ppos - 7);
+    String temp = version.substr(7, ppos - 7);
     major = atoi(temp.c_str());
-    temp = deviceVersion.substr(ppos + 1);
+    temp = version.substr(ppos + 1);
     minor = atoi(temp.c_str());
 }
 
@@ -1555,7 +1574,7 @@ struct Device::Impl
         addressBits_ = getProp<cl_uint, int>(CL_DEVICE_ADDRESS_BITS);
 
         String deviceVersion_ = getStrProp(CL_DEVICE_VERSION);
-        parseDeviceVersion(deviceVersion_, deviceVersionMajor_, deviceVersionMinor_);
+        parseOpenCLVersion(deviceVersion_, deviceVersionMajor_, deviceVersionMinor_);
 
         size_t pos = 0;
         while (pos < extensions_.size())
@@ -1675,7 +1694,7 @@ struct Device::Impl
 };
 
 
-Device::Device()
+Device::Device() CV_NOEXCEPT
 {
     p = 0;
 }
@@ -1704,6 +1723,23 @@ Device& Device::operator = (const Device& d)
     return *this;
 }
 
+Device::Device(Device&& d) CV_NOEXCEPT
+{
+    p = d.p;
+    d.p = nullptr;
+}
+
+Device& Device::operator = (Device&& d) CV_NOEXCEPT
+{
+    if (this != &d) {
+        if(p)
+            p->release();
+        p = d.p;
+        d.p = nullptr;
+    }
+    return *this;
+}
+
 Device::~Device()
 {
     if(p)
@@ -2832,7 +2868,7 @@ public:
 };
 
 
-Context::Context()
+Context::Context() CV_NOEXCEPT
 {
     p = 0;
 }
@@ -2917,6 +2953,23 @@ Context& Context::operator = (const Context& c)
     return *this;
 }
 
+Context::Context(Context&& c) CV_NOEXCEPT
+{
+    p = c.p;
+    c.p = nullptr;
+}
+
+Context& Context::operator = (Context&& c) CV_NOEXCEPT
+{
+    if (this != &c) {
+        if(p)
+            p->release();
+        p = c.p;
+        c.p = nullptr;
+    }
+    return *this;
+}
+
 void* Context::ptr() const
 {
     return p == NULL ? NULL : p->handle;
@@ -3229,7 +3282,7 @@ struct Queue::Impl
     cv::ocl::Queue profiling_queue_;
 };
 
-Queue::Queue()
+Queue::Queue() CV_NOEXCEPT
 {
     p = 0;
 }
@@ -3258,6 +3311,23 @@ Queue& Queue::operator = (const Queue& q)
     return *this;
 }
 
+Queue::Queue(Queue&& q) CV_NOEXCEPT
+{
+    p = q.p;
+    q.p = nullptr;
+}
+
+Queue& Queue::operator = (Queue&& q) CV_NOEXCEPT
+{
+    if (this != &q) {
+        if(p)
+            p->release();
+        p = q.p;
+        q.p = nullptr;
+    }
+    return *this;
+}
+
 Queue::~Queue()
 {
     if(p)
@@ -3313,7 +3383,7 @@ static cl_command_queue getQueue(const Queue& q)
 
 /////////////////////////////////////////// KernelArg /////////////////////////////////////////////
 
-KernelArg::KernelArg()
+KernelArg::KernelArg() CV_NOEXCEPT
     : flags(0), m(0), obj(0), sz(0), wscale(1), iwscale(1)
 {
 }
@@ -3380,16 +3450,24 @@ struct Kernel::Impl
             haveTempSrcUMats = true;  // UMat is created on RAW memory (without proper lifetime management, even from Mat)
     }
 
-    void addImage(const Image2D& image)
+    /// Preserve image lifetime (while it is specified as Kernel argument)
+    void registerImageArgument(int arg, const Image2D& image)
     {
-        images.push_back(image);
+        CV_CheckGE(arg, 0, "");
+        CV_CheckLT(arg, (int)MAX_ARRS, "");
+        if (arg < (int)shadow_images.size() && shadow_images[arg].ptr() != image.ptr())  // TODO future: replace ptr => impl (more strong check)
+        {
+            CV_Check(arg, !isInProgress, "ocl::Kernel: clearing of pending Image2D arguments is not allowed");
+        }
+        shadow_images.reserve(MAX_ARRS);
+        shadow_images.resize(std::max(shadow_images.size(), (size_t)arg + 1));
+        shadow_images[arg] = image;
     }
 
     void finit(cl_event e)
     {
         CV_UNUSED(e);
         cleanupUMats();
-        images.clear();
         isInProgress = false;
         release();
     }
@@ -3414,7 +3492,7 @@ struct Kernel::Impl
     bool isInProgress;
     bool isAsyncRun;  // true if kernel was scheduled in async mode
     int nu;
-    std::list<Image2D> images;
+    std::vector<Image2D> shadow_images;
     bool haveTempDstUMats;
     bool haveTempSrcUMats;
 };
@@ -3447,7 +3525,7 @@ static void CL_CALLBACK oclCleanupCallback(cl_event e, cl_int, void *p)
 
 namespace cv { namespace ocl {
 
-Kernel::Kernel()
+Kernel::Kernel() CV_NOEXCEPT
 {
     p = 0;
 }
@@ -3483,6 +3561,23 @@ Kernel& Kernel::operator = (const Kernel& k)
     return *this;
 }
 
+Kernel::Kernel(Kernel&& k) CV_NOEXCEPT
+{
+    p = k.p;
+    k.p = nullptr;
+}
+
+Kernel& Kernel::operator = (Kernel&& k) CV_NOEXCEPT
+{
+    if (this != &k) {
+        if(p)
+            p->release();
+        p = k.p;
+        k.p = nullptr;
+    }
+    return *this;
+}
+
 Kernel::~Kernel()
 {
     if(p)
@@ -3529,6 +3624,15 @@ bool Kernel::empty() const
     return ptr() == 0;
 }
 
+static cv::String dumpValue(size_t sz, const void* p)
+{
+    if (sz == 4)
+        return cv::format("%d / %uu / 0x%08x / %g", *(int*)p, *(int*)p, *(int*)p, *(float*)p);
+    if (sz == 8)
+        return cv::format("%lld / %lluu / 0x%16llx / %g", *(long long*)p, *(long long*)p, *(long long*)p, *(double*)p);
+    return cv::format("%p", p);
+}
+
 int Kernel::set(int i, const void* value, size_t sz)
 {
     if (!p || !p->handle)
@@ -3539,7 +3643,7 @@ int Kernel::set(int i, const void* value, size_t sz)
         p->cleanupUMats();
 
     cl_int retval = clSetKernelArg(p->handle, (cl_uint)i, sz, value);
-    CV_OCL_DBG_CHECK_RESULT(retval, cv::format("clSetKernelArg('%s', arg_index=%d, size=%d, value=%p)", p->name.c_str(), (int)i, (int)sz, (void*)value).c_str());
+    CV_OCL_DBG_CHECK_RESULT(retval, cv::format("clSetKernelArg('%s', arg_index=%d, size=%d, value=%s)", p->name.c_str(), (int)i, (int)sz, dumpValue(sz, value).c_str()).c_str());
     if (retval != CL_SUCCESS)
         return -1;
     return i+1;
@@ -3547,9 +3651,11 @@ int Kernel::set(int i, const void* value, size_t sz)
 
 int Kernel::set(int i, const Image2D& image2D)
 {
-    p->addImage(image2D);
     cl_mem h = (cl_mem)image2D.ptr();
-    return set(i, &h, sizeof(h));
+    int res = set(i, &h, sizeof(h));
+    if (res >= 0)
+        p->registerImageArgument(i, image2D);
+    return res;
 }
 
 int Kernel::set(int i, const UMat& m)
@@ -4026,7 +4132,7 @@ struct ProgramSource::Impl
 };
 
 
-ProgramSource::ProgramSource()
+ProgramSource::ProgramSource() CV_NOEXCEPT
 {
     p = 0;
 }
@@ -4070,6 +4176,23 @@ ProgramSource& ProgramSource::operator = (const ProgramSource& prog)
     return *this;
 }
 
+ProgramSource::ProgramSource(ProgramSource&& prog) CV_NOEXCEPT
+{
+    p = prog.p;
+    prog.p = nullptr;
+}
+
+ProgramSource& ProgramSource::operator = (ProgramSource&& prog) CV_NOEXCEPT
+{
+    if (this != &prog) {
+        if(p)
+            p->release();
+        p = prog.p;
+        prog.p = nullptr;
+    }
+    return *this;
+}
+
 const String& ProgramSource::source() const
 {
     CV_Assert(p);
@@ -4535,7 +4658,10 @@ struct Program::Impl
 };
 
 
-Program::Program() { p = 0; }
+Program::Program() CV_NOEXCEPT
+{
+    p = 0;
+}
 
 Program::Program(const ProgramSource& src,
         const String& buildflags, String& errmsg)
@@ -4562,6 +4688,23 @@ Program& Program::operator = (const Program& prog)
     return *this;
 }
 
+Program::Program(Program&& prog) CV_NOEXCEPT
+{
+    p = prog.p;
+    prog.p = nullptr;
+}
+
+Program& Program::operator = (Program&& prog) CV_NOEXCEPT
+{
+    if (this != &prog) {
+        if(p)
+            p->release();
+        p = prog.p;
+        prog.p = nullptr;
+    }
+    return *this;
+}
+
 Program::~Program()
 {
     if(p)
@@ -6370,7 +6513,6 @@ public:
 static OpenCLAllocator* getOpenCLAllocator_() // call once guarantee
 {
     static OpenCLAllocator* g_allocator = new OpenCLAllocator(); // avoid destructor call (using of this object is too wide)
-    g_isOpenCVActivated = true;
     return g_allocator;
 }
 MatAllocator* getOpenCLAllocator()
@@ -6566,6 +6708,9 @@ struct PlatformInfo::Impl
         refcount = 1;
         handle = *(cl_platform_id*)id;
         getDevices(devices, handle);
+
+        version_ = getStrProp(CL_PLATFORM_VERSION);
+        parseOpenCLVersion(version_, versionMajor_, versionMinor_);
     }
 
     String getStrProp(cl_platform_info prop) const
@@ -6579,9 +6724,13 @@ struct PlatformInfo::Impl
     IMPLEMENT_REFCOUNTABLE();
     std::vector<cl_device_id> devices;
     cl_platform_id handle;
+
+    String version_;
+    int versionMajor_;
+    int versionMinor_;
 };
 
-PlatformInfo::PlatformInfo()
+PlatformInfo::PlatformInfo() CV_NOEXCEPT
 {
     p = 0;
 }
@@ -6617,6 +6766,23 @@ PlatformInfo& PlatformInfo::operator =(const PlatformInfo& i)
     return *this;
 }
 
+PlatformInfo::PlatformInfo(PlatformInfo&& i) CV_NOEXCEPT
+{
+    p = i.p;
+    i.p = nullptr;
+}
+
+PlatformInfo& PlatformInfo::operator = (PlatformInfo&& i) CV_NOEXCEPT
+{
+    if (this != &i) {
+        if(p)
+            p->release();
+        p = i.p;
+        i.p = nullptr;
+    }
+    return *this;
+}
+
 int PlatformInfo::deviceNumber() const
 {
     return p ? (int)p->devices.size() : 0;
@@ -6641,7 +6807,19 @@ String PlatformInfo::vendor() const
 
 String PlatformInfo::version() const
 {
-    return p ? p->getStrProp(CL_PLATFORM_VERSION) : String();
+    return p ? p->version_ : String();
+}
+
+int PlatformInfo::versionMajor() const
+{
+    CV_Assert(p);
+    return p->versionMajor_;
+}
+
+int PlatformInfo::versionMinor() const
+{
+    CV_Assert(p);
+    return p->versionMinor_;
 }
 
 static void getPlatforms(std::vector<cl_platform_id>& platforms)
@@ -7145,7 +7323,7 @@ struct Image2D::Impl
     cl_mem handle;
 };
 
-Image2D::Image2D()
+Image2D::Image2D() CV_NOEXCEPT
 {
     p = NULL;
 }
@@ -7203,6 +7381,23 @@ Image2D & Image2D::operator = (const Image2D & i)
     return *this;
 }
 
+Image2D::Image2D(Image2D&& i) CV_NOEXCEPT
+{
+    p = i.p;
+    i.p = nullptr;
+}
+
+Image2D& Image2D::operator = (Image2D&& i) CV_NOEXCEPT
+{
+    if (this != &i) {
+        if (p)
+            p->release();
+        p = i.p;
+        i.p = nullptr;
+    }
+    return *this;
+}
+
 Image2D::~Image2D()
 {
     if (p)

modules/core/src/ocl_disabled.impl.hpp

@@ -34,10 +34,12 @@ CV_EXPORTS_W void finish() { /* nothing */ }
 
 CV_EXPORTS bool haveSVM() { return false; }
 
-Device::Device() : p(NULL) { }
+Device::Device() CV_NOEXCEPT : p(NULL) { }
 Device::Device(void* d) : p(NULL) { OCL_NOT_AVAILABLE(); }
 Device::Device(const Device& d) : p(NULL) { }
 Device& Device::operator=(const Device& d) { return *this; }
+Device::Device(Device&&) CV_NOEXCEPT : p(NULL) { }
+Device& Device::operator=(Device&&) CV_NOEXCEPT { return *this; }
 Device::~Device() { }
 
 void Device::set(void* d) { OCL_NOT_AVAILABLE(); }
@@ -147,11 +149,13 @@ const Device& Device::getDefault()
 /* static */ Device Device::fromHandle(void* d) { OCL_NOT_AVAILABLE(); }
 
 
-Context::Context() : p(NULL) { }
+Context::Context() CV_NOEXCEPT : p(NULL) { }
 Context::Context(int dtype) : p(NULL) { }
 Context::~Context() { }
 Context::Context(const Context& c) : p(NULL) { }
 Context& Context::operator=(const Context& c) { return *this; }
+Context::Context(Context&&) CV_NOEXCEPT : p(NULL) { }
+Context& Context::operator=(Context&&) CV_NOEXCEPT { return *this; }
 
 bool Context::create() { return false; }
 bool Context::create(int dtype) { return false; }
@@ -178,10 +182,12 @@ void Context::setUseSVM(bool enabled) { }
 void Context::release() { }
 
 
-Platform::Platform() : p(NULL) { }
+Platform::Platform() CV_NOEXCEPT : p(NULL) { }
 Platform::~Platform() { }
 Platform::Platform(const Platform&) : p(NULL) { }
 Platform& Platform::operator=(const Platform&) { return *this; }
+Platform::Platform(Platform&&) CV_NOEXCEPT : p(NULL) { }
+Platform& Platform::operator=(Platform&&) CV_NOEXCEPT { return *this; }
 
 void* Platform::ptr() const { return NULL; }
 
@@ -198,11 +204,13 @@ void convertFromImage(void* cl_mem_image, UMat& dst) { OCL_NOT_AVAILABLE(); }
 
 void initializeContextFromHandle(Context& ctx, void* platform, void* context, void* device) { OCL_NOT_AVAILABLE(); }
 
-Queue::Queue() : p(NULL) { }
+Queue::Queue() CV_NOEXCEPT : p(NULL) { }
 Queue::Queue(const Context& c, const Device& d) : p(NULL) { OCL_NOT_AVAILABLE(); }
 Queue::~Queue() { }
 Queue::Queue(const Queue& q) {}
 Queue& Queue::operator=(const Queue& q) { return *this; }
+Queue::Queue(Queue&&) CV_NOEXCEPT : p(NULL) { }
+Queue& Queue::operator=(Queue&&) CV_NOEXCEPT { return *this; }
 
 bool Queue::create(const Context& c, const Device& d) { OCL_NOT_AVAILABLE(); }
 void Queue::finish() {}
@@ -218,7 +226,7 @@ Queue& Queue::getDefault()
 const Queue& Queue::getProfilingQueue() const { OCL_NOT_AVAILABLE(); }
 
 
-KernelArg::KernelArg()
+KernelArg::KernelArg() CV_NOEXCEPT
     : flags(0), m(0), obj(0), sz(0), wscale(1), iwscale(1)
 {
 }
@@ -235,12 +243,14 @@ KernelArg KernelArg::Constant(const Mat& m)
 }
 
 
-Kernel::Kernel() : p(NULL) { }
+Kernel::Kernel() CV_NOEXCEPT : p(NULL) { }
 Kernel::Kernel(const char* kname, const Program& prog) : p(NULL) { OCL_NOT_AVAILABLE(); }
 Kernel::Kernel(const char* kname, const ProgramSource& prog, const String& buildopts, String* errmsg) : p(NULL) { OCL_NOT_AVAILABLE(); }
 Kernel::~Kernel() { }
 Kernel::Kernel(const Kernel& k) : p(NULL) { }
 Kernel& Kernel::operator=(const Kernel& k) { return *this; }
+Kernel::Kernel(Kernel&&) CV_NOEXCEPT : p(NULL) { }
+Kernel& Kernel::operator=(Kernel&&) CV_NOEXCEPT { return *this; }
 
 bool Kernel::empty() const { return true; }
 bool Kernel::create(const char* kname, const Program& prog) { OCL_NOT_AVAILABLE(); }
@@ -264,10 +274,12 @@ size_t Kernel::localMemSize() const { OCL_NOT_AVAILABLE(); }
 void* Kernel::ptr() const { return NULL; }
 
 
-Program::Program() : p(NULL) { }
+Program::Program() CV_NOEXCEPT : p(NULL) { }
 Program::Program(const ProgramSource& src, const String& buildflags, String& errmsg) : p(NULL) { OCL_NOT_AVAILABLE(); }
 Program::Program(const Program& prog) : p(NULL) { }
 Program& Program::operator=(const Program& prog) { return *this; }
+Program::Program(Program&&) CV_NOEXCEPT : p(NULL) { }
+Program& Program::operator=(Program&&) CV_NOEXCEPT { return *this; }
 Program::~Program() { }
 
 bool Program::create(const ProgramSource& src, const String& buildflags, String& errmsg) { OCL_NOT_AVAILABLE(); }
@@ -283,13 +295,15 @@ String Program::getPrefix() const { OCL_NOT_AVAILABLE(); }
 /* static */ String Program::getPrefix(const String& buildflags) { OCL_NOT_AVAILABLE(); }
 
 
-ProgramSource::ProgramSource() : p(NULL) { }
+ProgramSource::ProgramSource() CV_NOEXCEPT : p(NULL) { }
 ProgramSource::ProgramSource(const String& module, const String& name, const String& codeStr, const String& codeHash) : p(NULL) { }
 ProgramSource::ProgramSource(const String& prog) : p(NULL) { }
 ProgramSource::ProgramSource(const char* prog) : p(NULL) { }
 ProgramSource::~ProgramSource() { }
 ProgramSource::ProgramSource(const ProgramSource& prog) : p(NULL) { }
 ProgramSource& ProgramSource::operator=(const ProgramSource& prog) { return *this; }
+ProgramSource::ProgramSource(ProgramSource&&) CV_NOEXCEPT : p(NULL) { }
+ProgramSource& ProgramSource::operator=(ProgramSource&&) CV_NOEXCEPT { return *this; }
 
 const String& ProgramSource::source() const { OCL_NOT_AVAILABLE(); }
 ProgramSource::hash_t ProgramSource::hash() const { OCL_NOT_AVAILABLE(); }
@@ -298,12 +312,14 @@ ProgramSource::hash_t ProgramSource::hash() const { OCL_NOT_AVAILABLE(); }
 /* static */ ProgramSource ProgramSource::fromSPIR(const String& module, const String& name, const unsigned char* binary, const size_t size, const cv::String& buildOptions) { OCL_NOT_AVAILABLE(); }
 
 
-PlatformInfo::PlatformInfo() : p(NULL) { }
+PlatformInfo::PlatformInfo() CV_NOEXCEPT : p(NULL) { }
 PlatformInfo::PlatformInfo(void* id) : p(NULL) { OCL_NOT_AVAILABLE(); }
 PlatformInfo::~PlatformInfo() { }
 
 PlatformInfo::PlatformInfo(const PlatformInfo& i) : p(NULL) { }
 PlatformInfo& PlatformInfo::operator=(const PlatformInfo& i) { return *this; }
+PlatformInfo::PlatformInfo(PlatformInfo&&) CV_NOEXCEPT : p(NULL) { }
+PlatformInfo& PlatformInfo::operator=(PlatformInfo&&) CV_NOEXCEPT { return *this; }
 
 String PlatformInfo::name() const { OCL_NOT_AVAILABLE(); }
 String PlatformInfo::vendor() const { OCL_NOT_AVAILABLE(); }
@@ -341,11 +357,13 @@ int predictOptimalVectorWidthMax(InputArray src1, InputArray src2, InputArray sr
 void buildOptionsAddMatrixDescription(String& buildOptions, const String& name, InputArray _m) { OCL_NOT_AVAILABLE(); }
 
 
-Image2D::Image2D() : p(NULL) { }
+Image2D::Image2D() CV_NOEXCEPT : p(NULL) { }
 Image2D::Image2D(const UMat &src, bool norm, bool alias) { OCL_NOT_AVAILABLE(); }
 Image2D::Image2D(const Image2D & i) : p(NULL) { OCL_NOT_AVAILABLE(); }
 Image2D::~Image2D() { }
 Image2D& Image2D::operator=(const Image2D & i) { return *this; }
+Image2D::Image2D(Image2D&&) CV_NOEXCEPT : p(NULL) { }
+Image2D& Image2D::operator=(Image2D&&) CV_NOEXCEPT { return *this; }
 
 /* static */ bool Image2D::canCreateAlias(const UMat &u) { OCL_NOT_AVAILABLE(); }
 /* static */ bool Image2D::isFormatSupported(int depth, int cn, bool norm) { OCL_NOT_AVAILABLE(); }

modules/core/src/opencl/cvtclr_dx.cl

@@ -91,63 +91,50 @@ void YUV2BGR_NV12_8u(
 {
     int x = get_global_id(0);
     int y = get_global_id(1);
+    // each iteration computes 2*2=4 pixels
+    int x2 = x*2;
+    int y2 = y*2;
 
-    if (x + 1 < cols)
-    {
-        if (y + 1 < rows)
-        {
-            __global uchar* pDstRow1 = pBGR + mad24(y, bgrStep, mad24(x, NCHANNELS, 0));
-            __global uchar* pDstRow2 = pDstRow1 + bgrStep;
+    if (x2 + 1 < cols) {
+        if (y2 + 1 < rows) {
+            __global uchar *pDstRow1 = pBGR + mad24(y2, bgrStep, mad24(x2, NCHANNELS, 0));
+            __global uchar *pDstRow2 = pDstRow1 + bgrStep;
 
-            float4 Y1 = read_imagef(imgY, (int2)(x+0, y+0));
-            float4 Y2 = read_imagef(imgY, (int2)(x+1, y+0));
-            float4 Y3 = read_imagef(imgY, (int2)(x+0, y+1));
-            float4 Y4 = read_imagef(imgY, (int2)(x+1, y+1));
+            float4 Y1 = read_imagef(imgY, (int2)(x2 + 0, y2 + 0));
+            float4 Y2 = read_imagef(imgY, (int2)(x2 + 1, y2 + 0));
+            float4 Y3 = read_imagef(imgY, (int2)(x2 + 0, y2 + 1));
+            float4 Y4 = read_imagef(imgY, (int2)(x2 + 1, y2 + 1));
+            float4 Y = (float4)(Y1.x, Y2.x, Y3.x, Y4.x);
 
-            float4 UV = read_imagef(imgUV, (int2)(x/2, y/2)) - d2;
+            float4 UV = read_imagef(imgUV, (int2)(x, y)) - d2;
 
-            __constant float* coeffs = c_YUV2RGBCoeffs_420;
+            __constant float *coeffs = c_YUV2RGBCoeffs_420;
 
-            Y1 = max(0.f, Y1 - d1) * coeffs[0];
-            Y2 = max(0.f, Y2 - d1) * coeffs[0];
-            Y3 = max(0.f, Y3 - d1) * coeffs[0];
-            Y4 = max(0.f, Y4 - d1) * coeffs[0];
+            Y = max(0.f, Y - d1) * coeffs[0];
 
             float ruv = fma(coeffs[4], UV.y, 0.0f);
             float guv = fma(coeffs[3], UV.y, fma(coeffs[2], UV.x, 0.0f));
             float buv = fma(coeffs[1], UV.x, 0.0f);
 
-            float R1 = (Y1.x + ruv) * CV_8U_MAX;
-            float G1 = (Y1.x + guv) * CV_8U_MAX;
-            float B1 = (Y1.x + buv) * CV_8U_MAX;
+            float4 R = (Y + ruv) * CV_8U_MAX;
+            float4 G = (Y + guv) * CV_8U_MAX;
+            float4 B = (Y + buv) * CV_8U_MAX;
 
-            float R2 = (Y2.x + ruv) * CV_8U_MAX;
-            float G2 = (Y2.x + guv) * CV_8U_MAX;
-            float B2 = (Y2.x + buv) * CV_8U_MAX;
+            pDstRow1[0*NCHANNELS + 0] = convert_uchar_sat(B.x);
+            pDstRow1[0*NCHANNELS + 1] = convert_uchar_sat(G.x);
+            pDstRow1[0*NCHANNELS + 2] = convert_uchar_sat(R.x);
 
-            float R3 = (Y3.x + ruv) * CV_8U_MAX;
-            float G3 = (Y3.x + guv) * CV_8U_MAX;
-            float B3 = (Y3.x + buv) * CV_8U_MAX;
+            pDstRow1[1*NCHANNELS + 0] = convert_uchar_sat(B.y);
+            pDstRow1[1*NCHANNELS + 1] = convert_uchar_sat(G.y);
+            pDstRow1[1*NCHANNELS + 2] = convert_uchar_sat(R.y);
 
-            float R4 = (Y4.x + ruv) * CV_8U_MAX;
-            float G4 = (Y4.x + guv) * CV_8U_MAX;
-            float B4 = (Y4.x + buv) * CV_8U_MAX;
+            pDstRow2[0*NCHANNELS + 0] = convert_uchar_sat(B.z);
+            pDstRow2[0*NCHANNELS + 1] = convert_uchar_sat(G.z);
+            pDstRow2[0*NCHANNELS + 2] = convert_uchar_sat(R.z);
 
-            pDstRow1[0*NCHANNELS + 0] = convert_uchar_sat(B1);
-            pDstRow1[0*NCHANNELS + 1] = convert_uchar_sat(G1);
-            pDstRow1[0*NCHANNELS + 2] = convert_uchar_sat(R1);
-
-            pDstRow1[1*NCHANNELS + 0] = convert_uchar_sat(B2);
-            pDstRow1[1*NCHANNELS + 1] = convert_uchar_sat(G2);
-            pDstRow1[1*NCHANNELS + 2] = convert_uchar_sat(R2);
-
-            pDstRow2[0*NCHANNELS + 0] = convert_uchar_sat(B3);
-            pDstRow2[0*NCHANNELS + 1] = convert_uchar_sat(G3);
-            pDstRow2[0*NCHANNELS + 2] = convert_uchar_sat(R3);
-
-            pDstRow2[1*NCHANNELS + 0] = convert_uchar_sat(B4);
-            pDstRow2[1*NCHANNELS + 1] = convert_uchar_sat(G4);
-            pDstRow2[1*NCHANNELS + 2] = convert_uchar_sat(R4);
+            pDstRow2[1*NCHANNELS + 0] = convert_uchar_sat(B.w);
+            pDstRow2[1*NCHANNELS + 1] = convert_uchar_sat(G.w);
+            pDstRow2[1*NCHANNELS + 2] = convert_uchar_sat(R.w);
         }
     }
 }
@@ -172,12 +159,15 @@ void BGR2YUV_NV12_8u(
 {
     int x = get_global_id(0);
     int y = get_global_id(1);
+    // each iteration computes 2*2=4 pixels
+    int x2 = x*2;
+    int y2 = y*2;
 
-    if (x < cols)
+    if (x2 + 1 < cols)
     {
-        if (y < rows)
+        if (y2 + 1 < rows)
         {
-            __global const uchar* pSrcRow1 = pBGR + mad24(y, bgrStep, mad24(x, NCHANNELS, 0));
+            __global const uchar* pSrcRow1 = pBGR + mad24(y2, bgrStep, mad24(x2, NCHANNELS, 0));
             __global const uchar* pSrcRow2 = pSrcRow1 + bgrStep;
 
             float4 src_pix1 = convert_float4(vload4(0, pSrcRow1 + 0*NCHANNELS)) * CV_8U_SCALE;
@@ -196,12 +186,12 @@ void BGR2YUV_NV12_8u(
             UV.x = fma(coeffs[3], src_pix1.z, fma(coeffs[4], src_pix1.y, fma(coeffs[5], src_pix1.x, d2)));
             UV.y = fma(coeffs[5], src_pix1.z, fma(coeffs[6], src_pix1.y, fma(coeffs[7], src_pix1.x, d2)));
 
-            write_imagef(imgY, (int2)(x+0, y+0), Y1);
-            write_imagef(imgY, (int2)(x+1, y+0), Y2);
-            write_imagef(imgY, (int2)(x+0, y+1), Y3);
-            write_imagef(imgY, (int2)(x+1, y+1), Y4);
+            write_imagef(imgY, (int2)(x2+0, y2+0), Y1);
+            write_imagef(imgY, (int2)(x2+1, y2+0), Y2);
+            write_imagef(imgY, (int2)(x2+0, y2+1), Y3);
+            write_imagef(imgY, (int2)(x2+1, y2+1), Y4);
 
-            write_imagef(imgUV, (int2)((x/2), (y/2)), UV);
+            write_imagef(imgUV, (int2)(x, y), UV);
         }
     }
 }

modules/core/src/opencl/halfconvert.cl

@@ -47,8 +47,17 @@
 #endif
 #endif
 
-__kernel void convertFp16(__global const uchar * srcptr, int src_step, int src_offset,
-                          __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols)
+__kernel void
+#ifdef FLOAT_TO_HALF
+    convertFp16_FP32_to_FP16
+#else
+    convertFp16_FP16_to_FP32
+#endif
+(
+    __global const uchar * srcptr, int src_step, int src_offset,
+    __global uchar * dstptr, int dst_step, int dst_offset,
+    int dst_rows, int dst_cols
+)
 {
     int x = get_global_id(0);
     int y0 = get_global_id(1) * rowsPerWI;

modules/core/src/opengl.cpp

@@ -1575,6 +1575,7 @@ void cv::ogl::render(const ogl::Arrays& arr, InputArray indices, int mode, Scala
 // CL-GL Interoperability
 
 #ifdef HAVE_OPENCL
+#  include "opencv2/core/opencl/runtime/opencl_core.hpp"
 #  include "opencv2/core/opencl/runtime/opencl_gl.hpp"
 #  ifdef cl_khr_gl_sharing
 #    define HAVE_OPENCL_OPENGL_SHARING
@@ -1595,6 +1596,34 @@ void cv::ogl::render(const ogl::Arrays& arr, InputArray indices, int mode, Scala
 
 namespace cv { namespace ogl {
 
+#if defined(HAVE_OPENCL) && defined(HAVE_OPENGL) && defined(HAVE_OPENCL_OPENGL_SHARING)
+// Check to avoid crash in OpenCL runtime: https://github.com/opencv/opencv/issues/5209
+static void checkOpenCLVersion()
+{
+    using namespace cv::ocl;
+    const Device& device = Device::getDefault();
+    //CV_Assert(!device.empty());
+    cl_device_id dev = (cl_device_id)device.ptr();
+    CV_Assert(dev);
+
+    cl_platform_id platform_id = 0;
+    size_t sz = 0;
+
+    cl_int status = clGetDeviceInfo(dev, CL_DEVICE_PLATFORM, sizeof(platform_id), &platform_id, &sz);
+    CV_Assert(status == CL_SUCCESS && sz == sizeof(cl_platform_id));
+    CV_Assert(platform_id);
+
+    PlatformInfo pi(&platform_id);
+    int versionMajor = pi.versionMajor();
+    int versionMinor = pi.versionMinor();
+    if (versionMajor < 1 || (versionMajor == 1 && versionMinor <= 1))
+        CV_Error_(cv::Error::OpenCLApiCallError,
+            ("OpenCL: clCreateFromGLTexture requires OpenCL 1.2+ version: %d.%d - %s (%s)",
+                versionMajor, versionMinor, pi.name().c_str(), pi.version().c_str())
+        );
+}
+#endif
+
 namespace ocl {
 
 Context& initializeContextFromGL()
@@ -1719,6 +1748,8 @@ void convertToGLTexture2D(InputArray src, Texture2D& texture)
     Context& ctx = Context::getDefault();
     cl_context context = (cl_context)ctx.ptr();
 
+    checkOpenCLVersion();  // clCreateFromGLTexture requires OpenCL 1.2
+
     UMat u = src.getUMat();
 
     // TODO Add support for roi
@@ -1777,6 +1808,8 @@ void convertFromGLTexture2D(const Texture2D& texture, OutputArray dst)
     Context& ctx = Context::getDefault();
     cl_context context = (cl_context)ctx.ptr();
 
+    checkOpenCLVersion();  // clCreateFromGLTexture requires OpenCL 1.2
+
     // TODO Need to specify ACCESS_WRITE here somehow to prevent useless data copying!
     dst.create(texture.size(), textureType);
     UMat u = dst.getUMat();

modules/core/src/parallel.cpp

@@ -45,6 +45,9 @@
 #include <opencv2/core/utils/configuration.private.hpp>
 #include <opencv2/core/utils/trace.private.hpp>
 
+#include "opencv2/core/parallel/parallel_backend.hpp"
+#include "parallel/parallel.hpp"
+
 #if defined _WIN32 || defined WINCE
     #include <windows.h>
     #undef small
@@ -101,7 +104,6 @@
     #endif
     #include "tbb/tbb.h"
     #include "tbb/task.h"
-    #include "tbb/tbb_stddef.h"
     #if TBB_INTERFACE_VERSION >= 8000
         #include "tbb/task_arena.h"
     #endif
@@ -145,9 +147,7 @@
 #  define CV_PARALLEL_FRAMEWORK "pthreads"
 #endif
 
-#ifdef CV_PARALLEL_FRAMEWORK
 #include <atomic>
-#endif
 
 #include "parallel_impl.hpp"
 
@@ -159,9 +159,10 @@ namespace cv {
 
 ParallelLoopBody::~ParallelLoopBody() {}
 
+using namespace cv::parallel;
+
 namespace {
 
-#ifdef CV_PARALLEL_FRAMEWORK
 #ifdef ENABLE_INSTRUMENTATION
     static void SyncNodes(cv::instr::InstrNode *pNode)
     {
@@ -430,8 +431,6 @@ namespace {
     typedef ParallelLoopBodyWrapper ProxyLoopBody;
 #endif
 
-static int numThreads = -1;
-
 #if defined HAVE_TBB
     #if TBB_INTERFACE_VERSION >= 8000
         static tbb::task_arena tbbArena(tbb::task_arena::automatic);
@@ -446,7 +445,7 @@ static inline int _initMaxThreads()
     int maxThreads = omp_get_max_threads();
     if (!utils::getConfigurationParameterBool("OPENCV_FOR_OPENMP_DYNAMIC_DISABLE", false))
     {
-        omp_set_dynamic(maxThreads);
+        omp_set_dynamic(1);
     }
     return maxThreads;
 }
@@ -477,15 +476,11 @@ static SchedPtr pplScheduler;
 
 #endif
 
-#endif // CV_PARALLEL_FRAMEWORK
-
 } // namespace anon
 
 /* ================================   parallel_for_  ================================ */
 
-#ifdef CV_PARALLEL_FRAMEWORK
 static void parallel_for_impl(const cv::Range& range, const cv::ParallelLoopBody& body, double nstripes); // forward declaration
-#endif
 
 void parallel_for_(const cv::Range& range, const cv::ParallelLoopBody& body, double nstripes)
 {
@@ -500,7 +495,6 @@ void parallel_for_(const cv::Range& range, const cv::ParallelLoopBody& body, dou
     if (range.empty())
         return;
 
-#ifdef CV_PARALLEL_FRAMEWORK
     static std::atomic<bool> flagNestedParallelFor(false);
     bool isNotNestedRegion = !flagNestedParallelFor.load();
     if (isNotNestedRegion)
@@ -519,16 +513,23 @@ void parallel_for_(const cv::Range& range, const cv::ParallelLoopBody& body, dou
         }
     }
     else // nested parallel_for_() calls are not parallelized
-#endif // CV_PARALLEL_FRAMEWORK
     {
         CV_UNUSED(nstripes);
         body(range);
     }
 }
 
-#ifdef CV_PARALLEL_FRAMEWORK
+static
+void parallel_for_cb(int start, int end, void* data)
+{
+    CV_DbgAssert(data);
+    const cv::ParallelLoopBody& body = *(const cv::ParallelLoopBody*)data;
+    body(Range(start, end));
+}
+
 static void parallel_for_impl(const cv::Range& range, const cv::ParallelLoopBody& body, double nstripes)
 {
+    using namespace cv::parallel;
     if ((numThreads < 0 || numThreads > 1) && range.end - range.start > 1)
     {
         ParallelLoopBodyWrapperContext ctx(body, range, nstripes);
@@ -540,6 +541,16 @@ static void parallel_for_impl(const cv::Range& range, const cv::ParallelLoopBody
             return;
         }
 
+        std::shared_ptr<ParallelForAPI>& api = getCurrentParallelForAPI();
+        if (api)
+        {
+            CV_CheckEQ(stripeRange.start, 0, "");
+            api->parallel_for(stripeRange.end, parallel_for_cb, (void*)&pbody);
+            ctx.finalize();  // propagate exceptions if exists
+            return;
+        }
+
+#ifdef CV_PARALLEL_FRAMEWORK
 #if defined HAVE_TBB
 
 #if TBB_INTERFACE_VERSION >= 8000
@@ -590,24 +601,25 @@ static void parallel_for_impl(const cv::Range& range, const cv::ParallelLoopBody
 #endif
 
         ctx.finalize();  // propagate exceptions if exists
-    }
-    else
-    {
-        body(range);
-    }
-}
+        return;
 #endif // CV_PARALLEL_FRAMEWORK
+    }
+
+    body(range);
+}
 
 
 int getNumThreads(void)
 {
-#ifdef CV_PARALLEL_FRAMEWORK
+    std::shared_ptr<ParallelForAPI>& api = getCurrentParallelForAPI();
+    if (api)
+    {
+        return api->getNumThreads();
+    }
 
-    if(numThreads == 0)
+    if (numThreads == 0)
         return 1;
 
-#endif
-
 #if defined HAVE_TBB
 
 #if TBB_INTERFACE_VERSION >= 9100
@@ -682,10 +694,15 @@ unsigned defaultNumberOfThreads()
 void setNumThreads( int threads_ )
 {
     CV_UNUSED(threads_);
-#ifdef CV_PARALLEL_FRAMEWORK
+
     int threads = (threads_ < 0) ? defaultNumberOfThreads() : (unsigned)threads_;
     numThreads = threads;
-#endif
+
+    std::shared_ptr<ParallelForAPI>& api = getCurrentParallelForAPI();
+    if (api)
+    {
+        api->setNumThreads(numThreads);
+    }
 
 #ifdef HAVE_TBB
 
@@ -741,6 +758,12 @@ void setNumThreads( int threads_ )
 
 int getThreadNum()
 {
+    std::shared_ptr<ParallelForAPI>& api = getCurrentParallelForAPI();
+    if (api)
+    {
+        return api->getThreadNum();
+    }
+
 #if defined HAVE_TBB
     #if TBB_INTERFACE_VERSION >= 9100
         return tbb::this_task_arena::current_thread_index();
@@ -963,7 +986,13 @@ int getNumberOfCPUs()
     return nCPUs;  // cached value
 }
 
-const char* currentParallelFramework() {
+const char* currentParallelFramework()
+{
+    std::shared_ptr<ParallelForAPI>& api = getCurrentParallelForAPI();
+    if (api)
+    {
+        return api->getName();
+    }
 #ifdef CV_PARALLEL_FRAMEWORK
     return CV_PARALLEL_FRAMEWORK;
 #else

modules/core/src/parallel/factory_parallel.hpp

@@ -0,0 +1,48 @@
+// 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_CORE_PARALLEL_FACTORY_HPP
+#define OPENCV_CORE_PARALLEL_FACTORY_HPP
+
+#include "opencv2/core/parallel/parallel_backend.hpp"
+
+namespace cv { namespace parallel {
+
+class IParallelBackendFactory
+{
+public:
+    virtual ~IParallelBackendFactory() {}
+    virtual std::shared_ptr<cv::parallel::ParallelForAPI> create() const = 0;
+};
+
+
+class StaticBackendFactory CV_FINAL: public IParallelBackendFactory
+{
+protected:
+    std::function<std::shared_ptr<cv::parallel::ParallelForAPI>(void)> create_fn_;
+
+public:
+    StaticBackendFactory(std::function<std::shared_ptr<cv::parallel::ParallelForAPI>(void)>&& create_fn)
+        : create_fn_(create_fn)
+    {
+        // nothing
+    }
+
+    ~StaticBackendFactory() CV_OVERRIDE {}
+
+    std::shared_ptr<cv::parallel::ParallelForAPI> create() const CV_OVERRIDE
+    {
+        return create_fn_();
+    }
+};
+
+//
+// PluginBackendFactory is implemented in plugin_wrapper.cpp
+//
+
+std::shared_ptr<IParallelBackendFactory> createPluginParallelBackendFactory(const std::string& baseName);
+
+}}  // namespace
+
+#endif  // OPENCV_CORE_PARALLEL_FACTORY_HPP

modules/core/src/parallel/parallel.cpp

@@ -0,0 +1,177 @@
+// 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.
+#include "../precomp.hpp"
+#include "parallel.hpp"
+
+#include <opencv2/core/utils/configuration.private.hpp>
+#include <opencv2/core/utils/logger.defines.hpp>
+#ifdef NDEBUG
+#define CV_LOG_STRIP_LEVEL CV_LOG_LEVEL_DEBUG + 1
+#else
+#define CV_LOG_STRIP_LEVEL CV_LOG_LEVEL_VERBOSE + 1
+#endif
+#include <opencv2/core/utils/logger.hpp>
+
+
+#include "registry_parallel.hpp"
+#include "registry_parallel.impl.hpp"
+
+#include "plugin_parallel_api.hpp"
+#include "plugin_parallel_wrapper.impl.hpp"
+
+
+namespace cv { namespace parallel {
+
+int numThreads = -1;
+
+ParallelForAPI::~ParallelForAPI()
+{
+    // nothing
+}
+
+static
+std::string& getParallelBackendName()
+{
+    static std::string g_backendName = toUpperCase(cv::utils::getConfigurationParameterString("OPENCV_PARALLEL_BACKEND", ""));
+    return g_backendName;
+}
+
+static bool g_initializedParallelForAPI = false;
+
+static
+std::shared_ptr<ParallelForAPI> createParallelForAPI()
+{
+    const std::string& name = getParallelBackendName();
+    bool isKnown = false;
+    const auto& backends = getParallelBackendsInfo();
+    if (!name.empty())
+    {
+        CV_LOG_INFO(NULL, "core(parallel): requested backend name: " << name);
+    }
+    for (size_t i = 0; i < backends.size(); i++)
+    {
+        const auto& info = backends[i];
+        if (!name.empty())
+        {
+            if (name != info.name)
+            {
+                continue;
+            }
+            isKnown = true;
+        }
+        try
+        {
+            CV_LOG_DEBUG(NULL, "core(parallel): trying backend: " << info.name << " (priority=" << info.priority << ")");
+            if (!info.backendFactory)
+            {
+                CV_LOG_DEBUG(NULL, "core(parallel): factory is not available (plugins require filesystem support): " << info.name);
+                continue;
+            }
+            std::shared_ptr<ParallelForAPI> backend = info.backendFactory->create();
+            if (!backend)
+            {
+                CV_LOG_VERBOSE(NULL, 0, "core(parallel): not available: " << info.name);
+                continue;
+            }
+            CV_LOG_INFO(NULL, "core(parallel): using backend: " << info.name << " (priority=" << info.priority << ")");
+            g_initializedParallelForAPI = true;
+            getParallelBackendName() = info.name;
+            return backend;
+        }
+        catch (const std::exception& e)
+        {
+            CV_LOG_WARNING(NULL, "core(parallel): can't initialize " << info.name << " backend: " << e.what());
+        }
+        catch (...)
+        {
+            CV_LOG_WARNING(NULL, "core(parallel): can't initialize " << info.name << " backend: Unknown C++ exception");
+        }
+    }
+    if (name.empty())
+    {
+        CV_LOG_DEBUG(NULL, "core(parallel): fallback on builtin code");
+    }
+    else
+    {
+        if (!isKnown)
+            CV_LOG_INFO(NULL, "core(parallel): unknown backend: " << name);
+    }
+    g_initializedParallelForAPI = true;
+    return std::shared_ptr<ParallelForAPI>();
+}
+
+static inline
+std::shared_ptr<ParallelForAPI> createDefaultParallelForAPI()
+{
+    CV_LOG_DEBUG(NULL, "core(parallel): Initializing parallel backend...");
+    return createParallelForAPI();
+}
+
+std::shared_ptr<ParallelForAPI>& getCurrentParallelForAPI()
+{
+    static std::shared_ptr<ParallelForAPI> g_currentParallelForAPI = createDefaultParallelForAPI();
+    return g_currentParallelForAPI;
+}
+
+void setParallelForBackend(const std::shared_ptr<ParallelForAPI>& api, bool propagateNumThreads)
+{
+    getCurrentParallelForAPI() = api;
+    if (propagateNumThreads && api)
+    {
+        setNumThreads(numThreads);
+    }
+}
+
+bool setParallelForBackend(const std::string& backendName, bool propagateNumThreads)
+{
+    CV_TRACE_FUNCTION();
+
+    std::string backendName_u = toUpperCase(backendName);
+    if (g_initializedParallelForAPI)
+    {
+        // ... already initialized
+        if (getParallelBackendName() == backendName_u)
+        {
+            CV_LOG_INFO(NULL, "core(parallel): backend is already activated: " << (backendName.empty() ? "builtin(legacy)" : backendName));
+            return true;
+        }
+        else
+        {
+            // ... re-create new
+            CV_LOG_DEBUG(NULL, "core(parallel): replacing parallel backend...");
+            getParallelBackendName() = backendName_u;
+            getCurrentParallelForAPI() = createParallelForAPI();
+        }
+    }
+    else
+    {
+        // ... no backend exists, just specify the name (initialization is triggered by getCurrentParallelForAPI() call)
+        getParallelBackendName() = backendName_u;
+    }
+    std::shared_ptr<ParallelForAPI> api = getCurrentParallelForAPI();
+    if (!api)
+    {
+        if (!backendName.empty())
+        {
+            CV_LOG_WARNING(NULL, "core(parallel): backend is not available: " << backendName << " (using builtin legacy code)");
+            return false;
+        }
+        else
+        {
+            CV_LOG_WARNING(NULL, "core(parallel): switched to builtin code (legacy)");
+        }
+    }
+    if (!backendName_u.empty())
+    {
+        CV_Assert(backendName_u == getParallelBackendName());  // data race?
+    }
+
+    if (propagateNumThreads)
+    {
+        setNumThreads(numThreads);
+    }
+    return true;
+}
+
+}}  // namespace

modules/core/src/parallel/parallel.hpp

@@ -0,0 +1,29 @@
+// 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_CORE_SRC_PARALLEL_PARALLEL_HPP
+#define OPENCV_CORE_SRC_PARALLEL_PARALLEL_HPP
+
+#include "opencv2/core/parallel/parallel_backend.hpp"
+
+namespace cv { namespace parallel {
+
+extern int numThreads;
+
+std::shared_ptr<ParallelForAPI>& getCurrentParallelForAPI();
+
+#ifndef BUILD_PLUGIN
+
+#ifdef HAVE_TBB
+std::shared_ptr<cv::parallel::ParallelForAPI> createParallelBackendTBB();
+#endif
+
+#ifdef HAVE_OPENMP
+std::shared_ptr<cv::parallel::ParallelForAPI> createParallelBackendOpenMP();
+#endif
+
+#endif  // BUILD_PLUGIN
+
+}}  // namespace
+
+#endif // OPENCV_CORE_SRC_PARALLEL_PARALLEL_HPP

modules/core/src/parallel/parallel_openmp.cpp

@@ -0,0 +1,72 @@
+// 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.
+#include "../precomp.hpp"
+
+#ifdef HAVE_OPENMP
+
+#include "parallel.hpp"
+#include "opencv2/core/parallel/backend/parallel_for.openmp.hpp"
+
+namespace cv { namespace parallel {
+
+static
+std::shared_ptr<cv::parallel::openmp::ParallelForBackend>& getInstance()
+{
+    static std::shared_ptr<cv::parallel::openmp::ParallelForBackend> g_instance = std::make_shared<cv::parallel::openmp::ParallelForBackend>();
+    return g_instance;
+}
+
+#ifndef BUILD_PLUGIN
+std::shared_ptr<cv::parallel::ParallelForAPI> createParallelBackendOpenMP()
+{
+    return getInstance();
+}
+#endif
+
+}}  // namespace
+
+#ifdef BUILD_PLUGIN
+
+#define ABI_VERSION 0
+#define API_VERSION 0
+#include "plugin_parallel_api.hpp"
+
+static
+CvResult cv_getInstance(CV_OUT CvPluginParallelBackendAPI* handle) CV_NOEXCEPT
+{
+    try
+    {
+        if (!handle)
+            return CV_ERROR_FAIL;
+        *handle = cv::parallel::getInstance().get();
+        return CV_ERROR_OK;
+    }
+    catch (...)
+    {
+        return CV_ERROR_FAIL;
+    }
+}
+
+static const OpenCV_Core_Parallel_Plugin_API plugin_api =
+{
+    {
+        sizeof(OpenCV_Core_Parallel_Plugin_API), ABI_VERSION, API_VERSION,
+        CV_VERSION_MAJOR, CV_VERSION_MINOR, CV_VERSION_REVISION, CV_VERSION_STATUS,
+        "OpenMP (" CVAUX_STR(_OPENMP) ") OpenCV parallel plugin"
+    },
+    {
+        /*  1*/cv_getInstance
+    }
+};
+
+const OpenCV_Core_Parallel_Plugin_API* CV_API_CALL opencv_core_parallel_plugin_init_v0(int requested_abi_version, int requested_api_version, void* /*reserved=NULL*/) CV_NOEXCEPT
+{
+    if (requested_abi_version == ABI_VERSION && requested_api_version <= API_VERSION)
+        return &plugin_api;
+    return NULL;
+}
+
+#endif  // BUILD_PLUGIN
+
+#endif  // HAVE_TBB

modules/core/src/parallel/parallel_tbb.cpp

@@ -0,0 +1,74 @@
+// 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.
+#include "../precomp.hpp"
+
+#include "factory_parallel.hpp"
+
+#ifdef HAVE_TBB
+
+#include "parallel.hpp"
+#include "opencv2/core/parallel/backend/parallel_for.tbb.hpp"
+
+namespace cv { namespace parallel {
+
+static
+std::shared_ptr<cv::parallel::tbb::ParallelForBackend>& getInstance()
+{
+    static std::shared_ptr<cv::parallel::tbb::ParallelForBackend> g_instance = std::make_shared<cv::parallel::tbb::ParallelForBackend>();
+    return g_instance;
+}
+
+#ifndef BUILD_PLUGIN
+std::shared_ptr<cv::parallel::ParallelForAPI> createParallelBackendTBB()
+{
+    return getInstance();
+}
+#endif
+
+}}  // namespace
+
+#ifdef BUILD_PLUGIN
+
+#define ABI_VERSION 0
+#define API_VERSION 0
+#include "plugin_parallel_api.hpp"
+
+static
+CvResult cv_getInstance(CV_OUT CvPluginParallelBackendAPI* handle) CV_NOEXCEPT
+{
+    try
+    {
+        if (!handle)
+            return CV_ERROR_FAIL;
+        *handle = cv::parallel::getInstance().get();
+        return CV_ERROR_OK;
+    }
+    catch (...)
+    {
+        return CV_ERROR_FAIL;
+    }
+}
+
+static const OpenCV_Core_Parallel_Plugin_API plugin_api =
+{
+    {
+        sizeof(OpenCV_Core_Parallel_Plugin_API), ABI_VERSION, API_VERSION,
+        CV_VERSION_MAJOR, CV_VERSION_MINOR, CV_VERSION_REVISION, CV_VERSION_STATUS,
+        "TBB (interface " CVAUX_STR(TBB_INTERFACE_VERSION) ") OpenCV parallel plugin"
+    },
+    {
+        /*  1*/cv_getInstance
+    }
+};
+
+const OpenCV_Core_Parallel_Plugin_API* CV_API_CALL opencv_core_parallel_plugin_init_v0(int requested_abi_version, int requested_api_version, void* /*reserved=NULL*/) CV_NOEXCEPT
+{
+    if (requested_abi_version == ABI_VERSION && requested_api_version <= API_VERSION)
+        return &plugin_api;
+    return NULL;
+}
+
+#endif  // BUILD_PLUGIN
+
+#endif  // HAVE_TBB

modules/core/src/parallel/plugin_parallel_api.hpp

@@ -0,0 +1,72 @@
+// 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 PARALLEL_PLUGIN_API_HPP
+#define PARALLEL_PLUGIN_API_HPP
+
+#include <opencv2/core/cvdef.h>
+#include <opencv2/core/llapi/llapi.h>
+
+#include "opencv2/core/parallel/parallel_backend.hpp"
+
+#if !defined(BUILD_PLUGIN)
+
+/// increased for backward-compatible changes, e.g. add new function
+/// Caller API <= Plugin API -> plugin is fully compatible
+/// Caller API > Plugin API -> plugin is not fully compatible, caller should use extra checks to use plugins with older API
+#define API_VERSION 0 // preview
+
+/// increased for incompatible changes, e.g. remove function argument
+/// Caller ABI == Plugin ABI -> plugin is compatible
+/// Caller ABI > Plugin ABI -> plugin is not compatible, caller should use shim code to use old ABI plugins (caller may know how lower ABI works, so it is possible)
+/// Caller ABI < Plugin ABI -> plugin can't be used (plugin should provide interface with lower ABI to handle that)
+#define ABI_VERSION 0 // preview
+
+#else // !defined(BUILD_PLUGIN)
+
+#if !defined(ABI_VERSION) || !defined(API_VERSION)
+#error "Plugin must define ABI_VERSION and API_VERSION before including parallel_plugin_api.hpp"
+#endif
+
+#endif // !defined(BUILD_PLUGIN)
+
+typedef cv::parallel::ParallelForAPI* CvPluginParallelBackendAPI;
+
+struct OpenCV_Core_Parallel_Plugin_API_v0_0_api_entries
+{
+    /** @brief Get parallel backend API instance
+
+    @param[out] handle pointer on backend API handle
+
+    @note API-CALL 1, API-Version == 0
+     */
+    CvResult (CV_API_CALL *getInstance)(CV_OUT CvPluginParallelBackendAPI* handle) CV_NOEXCEPT;
+}; // OpenCV_Core_Parallel_Plugin_API_v0_0_api_entries
+
+typedef struct OpenCV_Core_Parallel_Plugin_API_v0
+{
+    OpenCV_API_Header api_header;
+    struct OpenCV_Core_Parallel_Plugin_API_v0_0_api_entries v0;
+} OpenCV_Core_Parallel_Plugin_API_v0;
+
+#if ABI_VERSION == 0 && API_VERSION == 0
+typedef OpenCV_Core_Parallel_Plugin_API_v0 OpenCV_Core_Parallel_Plugin_API;
+#else
+#error "Not supported configuration: check ABI_VERSION/API_VERSION"
+#endif
+
+#ifdef BUILD_PLUGIN
+extern "C" {
+
+CV_PLUGIN_EXPORTS
+const OpenCV_Core_Parallel_Plugin_API* CV_API_CALL opencv_core_parallel_plugin_init_v0
+        (int requested_abi_version, int requested_api_version, void* reserved /*NULL*/) CV_NOEXCEPT;
+
+}  // extern "C"
+#else  // BUILD_PLUGIN
+typedef const OpenCV_Core_Parallel_Plugin_API* (CV_API_CALL *FN_opencv_core_parallel_plugin_init_t)
+        (int requested_abi_version, int requested_api_version, void* reserved /*NULL*/);
+#endif  // BUILD_PLUGIN
+
+#endif // PARALLEL_PLUGIN_API_HPP

modules/core/src/parallel/plugin_parallel_wrapper.impl.hpp

@@ -0,0 +1,287 @@
+// 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.
+
+//
+// Not a standalone header, part of parallel.cpp
+//
+
+//==================================================================================================
+// Dynamic backend implementation
+
+#include "opencv2/core/utils/plugin_loader.private.hpp"
+
+namespace cv { namespace impl {
+
+using namespace cv::parallel;
+
+#if OPENCV_HAVE_FILESYSTEM_SUPPORT && defined(PARALLEL_ENABLE_PLUGINS)
+
+using namespace cv::plugin::impl;  // plugin_loader.hpp
+
+class PluginParallelBackend CV_FINAL: public std::enable_shared_from_this<PluginParallelBackend>
+{
+protected:
+    void initPluginAPI()
+    {
+        const char* init_name = "opencv_core_parallel_plugin_init_v0";
+        FN_opencv_core_parallel_plugin_init_t fn_init = reinterpret_cast<FN_opencv_core_parallel_plugin_init_t>(lib_->getSymbol(init_name));
+        if (fn_init)
+        {
+            CV_LOG_DEBUG(NULL, "Found entry: '" << init_name << "'");
+            for (int supported_api_version = API_VERSION; supported_api_version >= 0; supported_api_version--)
+            {
+                plugin_api_ = fn_init(ABI_VERSION, supported_api_version, NULL);
+                if (plugin_api_)
+                    break;
+            }
+            if (!plugin_api_)
+            {
+                CV_LOG_INFO(NULL, "core(parallel): plugin is incompatible (can't be initialized): " << lib_->getName());
+                return;
+            }
+            if (!checkCompatibility(plugin_api_->api_header, ABI_VERSION, API_VERSION, false))
+            {
+                plugin_api_ = NULL;
+                return;
+            }
+            CV_LOG_INFO(NULL, "core(parallel): plugin is ready to use '" << plugin_api_->api_header.api_description << "'");
+        }
+        else
+        {
+            CV_LOG_INFO(NULL, "core(parallel): plugin is incompatible, missing init function: '" << init_name << "', file: " << lib_->getName());
+        }
+    }
+
+
+    bool checkCompatibility(const OpenCV_API_Header& api_header, unsigned int abi_version, unsigned int api_version, bool checkMinorOpenCVVersion)
+    {
+        if (api_header.opencv_version_major != CV_VERSION_MAJOR)
+        {
+            CV_LOG_ERROR(NULL, "core(parallel): wrong OpenCV major version used by plugin '" << api_header.api_description << "': " <<
+                cv::format("%d.%d, OpenCV version is '" CV_VERSION "'", api_header.opencv_version_major, api_header.opencv_version_minor))
+            return false;
+        }
+        if (!checkMinorOpenCVVersion)
+        {
+            // no checks for OpenCV minor version
+        }
+        else if (api_header.opencv_version_minor != CV_VERSION_MINOR)
+        {
+            CV_LOG_ERROR(NULL, "core(parallel): wrong OpenCV minor version used by plugin '" << api_header.api_description << "': " <<
+                cv::format("%d.%d, OpenCV version is '" CV_VERSION "'", api_header.opencv_version_major, api_header.opencv_version_minor))
+            return false;
+        }
+        CV_LOG_DEBUG(NULL, "core(parallel): initialized '" << api_header.api_description << "': built with "
+            << cv::format("OpenCV %d.%d (ABI/API = %d/%d)",
+                 api_header.opencv_version_major, api_header.opencv_version_minor,
+                 api_header.min_api_version, api_header.api_version)
+            << ", current OpenCV version is '" CV_VERSION "' (ABI/API = " << abi_version << "/" << api_version << ")"
+        );
+        if (api_header.min_api_version != abi_version)  // future: range can be here
+        {
+            // actually this should never happen due to checks in plugin's init() function
+            CV_LOG_ERROR(NULL, "core(parallel): plugin is not supported due to incompatible ABI = " << api_header.min_api_version);
+            return false;
+        }
+        if (api_header.api_version != api_version)
+        {
+            CV_LOG_INFO(NULL, "core(parallel): NOTE: plugin is supported, but there is API version mismath: "
+                << cv::format("plugin API level (%d) != OpenCV API level (%d)", api_header.api_version, api_version));
+            if (api_header.api_version < api_version)
+            {
+                CV_LOG_INFO(NULL, "core(parallel): NOTE: some functionality may be unavailable due to lack of support by plugin implementation");
+            }
+        }
+        return true;
+    }
+
+public:
+    std::shared_ptr<cv::plugin::impl::DynamicLib> lib_;
+    const OpenCV_Core_Parallel_Plugin_API* plugin_api_;
+
+    PluginParallelBackend(const std::shared_ptr<cv::plugin::impl::DynamicLib>& lib)
+        : lib_(lib)
+        , plugin_api_(NULL)
+    {
+        initPluginAPI();
+    }
+
+    std::shared_ptr<cv::parallel::ParallelForAPI> create() const
+    {
+        CV_Assert(plugin_api_);
+
+        CvPluginParallelBackendAPI instancePtr = NULL;
+
+        if (plugin_api_->v0.getInstance)
+        {
+            if (CV_ERROR_OK == plugin_api_->v0.getInstance(&instancePtr))
+            {
+                CV_Assert(instancePtr);
+                // TODO C++20 "aliasing constructor"
+                return std::shared_ptr<cv::parallel::ParallelForAPI>(instancePtr, [](cv::parallel::ParallelForAPI*){});  // empty deleter
+            }
+        }
+        return std::shared_ptr<cv::parallel::ParallelForAPI>();
+    }
+};
+
+
+class PluginParallelBackendFactory CV_FINAL: public IParallelBackendFactory
+{
+public:
+    std::string baseName_;
+    std::shared_ptr<PluginParallelBackend> backend;
+    bool initialized;
+public:
+    PluginParallelBackendFactory(const std::string& baseName)
+        : baseName_(baseName)
+        , initialized(false)
+    {
+        // nothing, plugins are loaded on demand
+    }
+
+    std::shared_ptr<cv::parallel::ParallelForAPI> create() const CV_OVERRIDE
+    {
+        if (!initialized)
+        {
+            const_cast<PluginParallelBackendFactory*>(this)->initBackend();
+        }
+        if (backend)
+            return backend->create();
+        return std::shared_ptr<cv::parallel::ParallelForAPI>();
+    }
+protected:
+    void initBackend()
+    {
+        AutoLock lock(getInitializationMutex());
+        try
+        {
+            if (!initialized)
+                loadPlugin();
+        }
+        catch (...)
+        {
+            CV_LOG_INFO(NULL, "core(parallel): exception during plugin loading: " << baseName_ << ". SKIP");
+        }
+        initialized = true;
+    }
+    void loadPlugin();
+};
+
+static
+std::vector<FileSystemPath_t> getPluginCandidates(const std::string& baseName)
+{
+    using namespace cv::utils;
+    using namespace cv::utils::fs;
+    const std::string baseName_l = toLowerCase(baseName);
+    const std::string baseName_u = toUpperCase(baseName);
+    const FileSystemPath_t baseName_l_fs = toFileSystemPath(baseName_l);
+    std::vector<FileSystemPath_t> paths;
+    // TODO OPENCV_PLUGIN_PATH
+    const std::vector<std::string> paths_ = getConfigurationParameterPaths("OPENCV_CORE_PLUGIN_PATH", std::vector<std::string>());
+    if (paths_.size() != 0)
+    {
+        for (size_t i = 0; i < paths_.size(); i++)
+        {
+            paths.push_back(toFileSystemPath(paths_[i]));
+        }
+    }
+    else
+    {
+        FileSystemPath_t binaryLocation;
+        if (getBinLocation(binaryLocation))
+        {
+            binaryLocation = getParent(binaryLocation);
+#ifndef CV_CORE_PARALLEL_PLUGIN_SUBDIRECTORY
+            paths.push_back(binaryLocation);
+#else
+            paths.push_back(binaryLocation + toFileSystemPath("/") + toFileSystemPath(CV_CORE_PARALLEL_PLUGIN_SUBDIRECTORY_STR));
+#endif
+        }
+    }
+    const std::string default_expr = libraryPrefix() + "opencv_core_parallel_" + baseName_l + "*" + librarySuffix();
+    const std::string plugin_expr = getConfigurationParameterString((std::string("OPENCV_CORE_PARALLEL_PLUGIN_") + baseName_u).c_str(), default_expr.c_str());
+    std::vector<FileSystemPath_t> results;
+#ifdef _WIN32
+    FileSystemPath_t moduleName = toFileSystemPath(libraryPrefix() + "opencv_core_parallel_" + baseName_l + librarySuffix());
+    if (plugin_expr != default_expr)
+    {
+        moduleName = toFileSystemPath(plugin_expr);
+        results.push_back(moduleName);
+    }
+    for (const FileSystemPath_t& path : paths)
+    {
+        results.push_back(path + L"\\" + moduleName);
+    }
+    results.push_back(moduleName);
+#else
+    CV_LOG_DEBUG(NULL, "core(parallel): " << baseName << " plugin's glob is '" << plugin_expr << "', " << paths.size() << " location(s)");
+    for (const std::string& path : paths)
+    {
+        if (path.empty())
+            continue;
+        std::vector<std::string> candidates;
+        cv::glob(utils::fs::join(path, plugin_expr), candidates);
+        CV_LOG_DEBUG(NULL, "    - " << path << ": " << candidates.size());
+        copy(candidates.begin(), candidates.end(), back_inserter(results));
+    }
+#endif
+    CV_LOG_DEBUG(NULL, "Found " << results.size() << " plugin(s) for " << baseName);
+    return results;
+}
+
+void PluginParallelBackendFactory::loadPlugin()
+{
+    for (const FileSystemPath_t& plugin : getPluginCandidates(baseName_))
+    {
+        auto lib = std::make_shared<cv::plugin::impl::DynamicLib>(plugin);
+        if (!lib->isLoaded())
+        {
+            continue;
+        }
+        try
+        {
+            auto pluginBackend = std::make_shared<PluginParallelBackend>(lib);
+            if (!pluginBackend)
+            {
+                continue;
+            }
+            if (pluginBackend->plugin_api_ == NULL)
+            {
+                CV_LOG_ERROR(NULL, "core(parallel): no compatible plugin API for backend: " << baseName_ << " in " << toPrintablePath(plugin));
+                continue;
+            }
+#if !defined(_WIN32)
+            // NB: we are going to use parallel backend, so prevent automatic library unloading
+            // (avoid uncontrolled crashes in worker threads of underlying libraries: libgomp, libtbb)
+            // details: https://github.com/opencv/opencv/pull/19470#pullrequestreview-589834777
+            lib->disableAutomaticLibraryUnloading();
+#endif
+            backend = pluginBackend;
+            return;
+        }
+        catch (...)
+        {
+            CV_LOG_WARNING(NULL, "core(parallel): exception during plugin initialization: " << toPrintablePath(plugin) << ". SKIP");
+        }
+    }
+}
+
+#endif  // OPENCV_HAVE_FILESYSTEM_SUPPORT && defined(PARALLEL_ENABLE_PLUGINS)
+
+}  // namespace
+
+namespace parallel {
+
+std::shared_ptr<IParallelBackendFactory> createPluginParallelBackendFactory(const std::string& baseName)
+{
+#if OPENCV_HAVE_FILESYSTEM_SUPPORT && defined(PARALLEL_ENABLE_PLUGINS)
+    return std::make_shared<impl::PluginParallelBackendFactory>(baseName);
+#else
+    CV_UNUSED(baseName);
+    return std::shared_ptr<IParallelBackendFactory>();
+#endif
+}
+
+}}  // namespace

modules/core/src/parallel/registry_parallel.hpp

@@ -0,0 +1,25 @@
+// 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_CORE_PARALLEL_REGISTRY_HPP
+#define OPENCV_CORE_PARALLEL_REGISTRY_HPP
+
+#include "factory_parallel.hpp"
+
+namespace cv { namespace parallel {
+
+struct ParallelBackendInfo
+{
+    int priority;     // 1000-<index*10> - default builtin priority
+                      // 0 - disabled (OPENCV_PARALLEL_PRIORITY_<name> = 0)
+                      // >10000 - prioritized list (OPENCV_PARALLEL_PRIORITY_LIST)
+    std::string name;
+    std::shared_ptr<IParallelBackendFactory> backendFactory;
+};
+
+const std::vector<ParallelBackendInfo>& getParallelBackendsInfo();
+
+}} // namespace
+
+#endif // OPENCV_CORE_PARALLEL_REGISTRY_HPP

modules/core/src/parallel/registry_parallel.impl.hpp

@@ -0,0 +1,173 @@
+// 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.
+
+//
+// Not a standalone header, part of parallel.cpp
+//
+
+#include "opencv2/core/utils/filesystem.private.hpp"  // OPENCV_HAVE_FILESYSTEM_SUPPORT
+
+namespace cv { namespace parallel {
+
+#if OPENCV_HAVE_FILESYSTEM_SUPPORT && defined(PARALLEL_ENABLE_PLUGINS)
+#define DECLARE_DYNAMIC_BACKEND(name) \
+ParallelBackendInfo { \
+    1000, name, createPluginParallelBackendFactory(name) \
+},
+#else
+#define DECLARE_DYNAMIC_BACKEND(name) /* nothing */
+#endif
+
+#define DECLARE_STATIC_BACKEND(name, createBackendAPI) \
+ParallelBackendInfo { \
+    1000, name, std::make_shared<cv::parallel::StaticBackendFactory>([=] () -> std::shared_ptr<cv::parallel::ParallelForAPI> { return createBackendAPI(); }) \
+},
+
+static
+std::vector<ParallelBackendInfo>& getBuiltinParallelBackendsInfo()
+{
+    static std::vector<ParallelBackendInfo> g_backends
+    {
+#ifdef HAVE_TBB
+        DECLARE_STATIC_BACKEND("TBB", createParallelBackendTBB)
+#elif defined(PARALLEL_ENABLE_PLUGINS)
+        DECLARE_DYNAMIC_BACKEND("ONETBB")   // dedicated oneTBB plugin (interface >= 12000, binary incompatibe with TBB 2017-2020)
+        DECLARE_DYNAMIC_BACKEND("TBB")      // generic TBB plugins
+#endif
+
+#ifdef HAVE_OPENMP
+        DECLARE_STATIC_BACKEND("OPENMP", createParallelBackendOpenMP)
+#elif defined(PARALLEL_ENABLE_PLUGINS)
+        DECLARE_DYNAMIC_BACKEND("OPENMP")  // TODO Intel OpenMP?
+#endif
+    };
+    return g_backends;
+};
+
+static
+bool sortByPriority(const ParallelBackendInfo &lhs, const ParallelBackendInfo &rhs)
+{
+    return lhs.priority > rhs.priority;
+}
+
+/** @brief Manages list of enabled backends
+ */
+class ParallelBackendRegistry
+{
+protected:
+    std::vector<ParallelBackendInfo> enabledBackends;
+    ParallelBackendRegistry()
+    {
+        enabledBackends = getBuiltinParallelBackendsInfo();
+        int N = (int)enabledBackends.size();
+        for (int i = 0; i < N; i++)
+        {
+            ParallelBackendInfo& info = enabledBackends[i];
+            info.priority = 1000 - i * 10;
+        }
+        CV_LOG_DEBUG(NULL, "core(parallel): Builtin backends(" << N << "): " << dumpBackends());
+        if (readPrioritySettings())
+        {
+            CV_LOG_INFO(NULL, "core(parallel): Updated backends priorities: " << dumpBackends());
+            N = (int)enabledBackends.size();
+        }
+        int enabled = 0;
+        for (int i = 0; i < N; i++)
+        {
+            ParallelBackendInfo& info = enabledBackends[enabled];
+            if (enabled != i)
+                info = enabledBackends[i];
+            size_t param_priority = utils::getConfigurationParameterSizeT(cv::format("OPENCV_PARALLEL_PRIORITY_%s", info.name.c_str()).c_str(), (size_t)info.priority);
+            CV_Assert(param_priority == (size_t)(int)param_priority); // overflow check
+            if (param_priority > 0)
+            {
+                info.priority = (int)param_priority;
+                enabled++;
+            }
+            else
+            {
+                CV_LOG_INFO(NULL, "core(parallel): Disable backend: " << info.name);
+            }
+        }
+        enabledBackends.resize(enabled);
+        CV_LOG_DEBUG(NULL, "core(parallel): Available backends(" << enabled << "): " << dumpBackends());
+        std::sort(enabledBackends.begin(), enabledBackends.end(), sortByPriority);
+        CV_LOG_INFO(NULL, "core(parallel): Enabled backends(" << enabled << ", sorted by priority): " << (enabledBackends.empty() ? std::string("N/A") : dumpBackends()));
+    }
+
+    static std::vector<std::string> tokenize_string(const std::string& input, char token)
+    {
+        std::vector<std::string> result;
+        std::string::size_type prev_pos = 0, pos = 0;
+        while((pos = input.find(token, pos)) != std::string::npos)
+        {
+            result.push_back(input.substr(prev_pos, pos-prev_pos));
+            prev_pos = ++pos;
+        }
+        result.push_back(input.substr(prev_pos));
+        return result;
+    }
+    bool readPrioritySettings()
+    {
+        bool hasChanges = false;
+        cv::String prioritized_backends = utils::getConfigurationParameterString("OPENCV_PARALLEL_PRIORITY_LIST", NULL);
+        if (prioritized_backends.empty())
+            return hasChanges;
+        CV_LOG_INFO(NULL, "core(parallel): Configured priority list (OPENCV_PARALLEL_PRIORITY_LIST): " << prioritized_backends);
+        const std::vector<std::string> names = tokenize_string(prioritized_backends, ',');
+        for (size_t i = 0; i < names.size(); i++)
+        {
+            const std::string& name = names[i];
+            int priority = (int)(100000 + (names.size() - i) * 1000);
+            bool found = false;
+            for (size_t k = 0; k < enabledBackends.size(); k++)
+            {
+                ParallelBackendInfo& info = enabledBackends[k];
+                if (name == info.name)
+                {
+                    info.priority = priority;
+                    CV_LOG_DEBUG(NULL, "core(parallel): New backend priority: '" << name << "' => " << info.priority);
+                    found = true;
+                    hasChanges = true;
+                    break;
+                }
+            }
+            if (!found)
+            {
+                CV_LOG_INFO(NULL, "core(parallel): Adding parallel backend (plugin): '" << name << "'");
+                enabledBackends.push_back(ParallelBackendInfo{priority, name, createPluginParallelBackendFactory(name)});
+                hasChanges = true;
+            }
+        }
+        return hasChanges;
+    }
+public:
+    std::string dumpBackends() const
+    {
+        std::ostringstream os;
+        for (size_t i = 0; i < enabledBackends.size(); i++)
+        {
+            if (i > 0) os << "; ";
+            const ParallelBackendInfo& info = enabledBackends[i];
+            os << info.name << '(' << info.priority << ')';
+        }
+        return os.str();
+    }
+
+    static ParallelBackendRegistry& getInstance()
+    {
+        static ParallelBackendRegistry g_instance;
+        return g_instance;
+    }
+
+    inline const std::vector<ParallelBackendInfo>& getEnabledBackends() const { return enabledBackends; }
+};
+
+
+const std::vector<ParallelBackendInfo>& getParallelBackendsInfo()
+{
+    return cv::parallel::ParallelBackendRegistry::getInstance().getEnabledBackends();
+}
+
+}} // namespace

modules/core/src/parallel_impl.cpp

@@ -356,6 +356,16 @@ public:
 };
 
 
+// Disable thread sanitization check when CV_USE_GLOBAL_WORKERS_COND_VAR is not
+// set because it triggers as the main thread reads isActive while the children
+// thread writes it (but it all works out because a mutex is locked in the main
+// thread and isActive re-checked).
+// This is to solve issue #19463.
+#if !defined(CV_USE_GLOBAL_WORKERS_COND_VAR) && defined(__clang__) && defined(__has_feature)
+#if __has_feature(thread_sanitizer)
+__attribute__((no_sanitize("thread")))
+#endif
+#endif
 void WorkerThread::thread_body()
 {
     (void)cv::utils::getThreadID(); // notify OpenCV about new thread

modules/core/src/precomp.hpp

@@ -43,6 +43,10 @@
 #ifndef __OPENCV_PRECOMP_H__
 #define __OPENCV_PRECOMP_H__
 
+#ifdef BUILD_PLUGIN
+#include "opencv2/core/utility.hpp"
+#else  // BUILD_PLUGIN
+
 #include "opencv2/opencv_modules.hpp"
 #include "cvconfig.h"
 
@@ -375,4 +379,5 @@ int cv_snprintf(char* buf, int len, const char* fmt, ...);
 int cv_vsnprintf(char* buf, int len, const char* fmt, va_list args);
 }
 
-#endif /*_CXCORE_INTERNAL_H_*/
+#endif  // BUILD_PLUGIN
+#endif  // __OPENCV_PRECOMP_H__

modules/core/src/rand.cpp

@@ -750,6 +750,9 @@ void cv::randShuffle( InputOutputArray _dst, double iterFactor, RNG* _rng )
     func( dst, rng, iterFactor );
 }
 
+
+#ifndef OPENCV_EXCLUDE_C_API
+
 CV_IMPL void
 cvRandArr( CvRNG* _rng, CvArr* arr, int disttype, CvScalar param1, CvScalar param2 )
 {
@@ -767,6 +770,9 @@ CV_IMPL void cvRandShuffle( CvArr* arr, CvRNG* _rng, double iter_factor )
     cv::randShuffle( dst, iter_factor, &rng );
 }
 
+#endif  // OPENCV_EXCLUDE_C_API
+
+
 // Mersenne Twister random number generator.
 // Inspired by http://www.math.sci.hiroshima-u.ac.jp/~m-mat/MT/MT2002/CODES/mt19937ar.c
 

modules/core/src/stat_c.cpp

@@ -5,6 +5,8 @@
 
 #include "precomp.hpp"
 
+#ifndef OPENCV_EXCLUDE_C_API
+
 CV_IMPL CvScalar cvSum( const CvArr* srcarr )
 {
     cv::Scalar sum = cv::sum(cv::cvarrToMat(srcarr, false, true, 1));
@@ -117,3 +119,5 @@ cvNorm( const void* imgA, const void* imgB, int normType, const void* maskarr )
 
     return !maskarr ? cv::norm(a, b, normType) : cv::norm(a, b, normType, mask);
 }
+
+#endif  // OPENCV_EXCLUDE_C_API

modules/core/src/system.cpp

@@ -128,11 +128,14 @@ void* allocSingletonNewBuffer(size_t size) { return malloc(size); }
 #endif
 
 
-#if CV_VSX && defined __linux__
+#if (defined __ppc64__ || defined __PPC64__) && defined __linux__
 # include "sys/auxv.h"
 # ifndef AT_HWCAP2
 #   define AT_HWCAP2 26
 # endif
+# ifndef PPC_FEATURE2_ARCH_2_07
+#   define PPC_FEATURE2_ARCH_2_07 0x80000000
+# endif
 # ifndef PPC_FEATURE2_ARCH_3_00
 #   define PPC_FEATURE2_ARCH_3_00 0x00800000
 # endif
@@ -345,7 +348,6 @@ struct HWFeatures
 
     HWFeatures(bool run_initialize = false)
     {
-        memset( have, 0, sizeof(have[0]) * MAX_FEATURE );
         if (run_initialize)
             initialize();
     }
@@ -589,14 +591,25 @@ struct HWFeatures
     #ifdef __mips_msa
         have[CV_CPU_MSA] = true;
     #endif
-    // there's no need to check VSX availability in runtime since it's always available on ppc64le CPUs
-    have[CV_CPU_VSX] = (CV_VSX);
-    // TODO: Check VSX3 availability in runtime for other platforms
-    #if CV_VSX && defined __linux__
-        uint64 hwcap2 = getauxval(AT_HWCAP2);
-        have[CV_CPU_VSX3] = (hwcap2 & PPC_FEATURE2_ARCH_3_00);
+
+    #if (defined __ppc64__ || defined __PPC64__) && defined __linux__
+        unsigned int hwcap = getauxval(AT_HWCAP);
+        if (hwcap & PPC_FEATURE_HAS_VSX) {
+            hwcap = getauxval(AT_HWCAP2);
+            if (hwcap & PPC_FEATURE2_ARCH_3_00) {
+                have[CV_CPU_VSX] = have[CV_CPU_VSX3] = true;
+            } else {
+                have[CV_CPU_VSX] = (hwcap & PPC_FEATURE2_ARCH_2_07) != 0;
+            }
+        }
     #else
-        have[CV_CPU_VSX3] = (CV_VSX3);
+        // TODO: AIX, FreeBSD
+        #if CV_VSX || defined _ARCH_PWR8 || defined __POWER9_VECTOR__
+            have[CV_CPU_VSX] = true;
+        #endif
+        #if CV_VSX3 || defined __POWER9_VECTOR__
+            have[CV_CPU_VSX3] = true;
+        #endif
     #endif
 
     #if defined __riscv && defined __riscv_vector
@@ -730,7 +743,7 @@ struct HWFeatures
         }
     }
 
-    bool have[MAX_FEATURE+1];
+    bool have[MAX_FEATURE+1]{};
 };
 
 static HWFeatures  featuresEnabled(true), featuresDisabled = HWFeatures(false);
@@ -1810,7 +1823,7 @@ class ParseError
 {
     std::string bad_value;
 public:
-    ParseError(const std::string bad_value_) :bad_value(bad_value_) {}
+    ParseError(const std::string &bad_value_) :bad_value(bad_value_) {}
     std::string toString(const std::string &param) const
     {
         std::ostringstream out;
@@ -2313,6 +2326,13 @@ public:
             ippTopFeatures = ippCPUID_SSE42;
 
         pIppLibInfo = ippiGetLibVersion();
+
+        // workaround: https://github.com/opencv/opencv/issues/12959
+        std::string ippName(pIppLibInfo->Name ? pIppLibInfo->Name : "");
+        if (ippName.find("SSE4.2") != std::string::npos)
+        {
+            ippTopFeatures = ippCPUID_SSE42;
+        }
     }
 
 public:
@@ -2344,16 +2364,12 @@ unsigned long long getIppFeatures()
 #endif
 }
 
-unsigned long long getIppTopFeatures();
-
+#ifdef HAVE_IPP
 unsigned long long getIppTopFeatures()
 {
-#ifdef HAVE_IPP
     return getIPPSingleton().ippTopFeatures;
-#else
-    return 0;
-#endif
 }
+#endif
 
 void setIppStatus(int status, const char * const _funcname, const char * const _filename, int _line)
 {

modules/core/src/umatrix.cpp

@@ -230,7 +230,7 @@ UMatDataAutoLock::~UMatDataAutoLock()
 
 //////////////////////////////// UMat ////////////////////////////////
 
-UMat::UMat(UMatUsageFlags _usageFlags)
+UMat::UMat(UMatUsageFlags _usageFlags) CV_NOEXCEPT
 : flags(MAGIC_VAL), dims(0), rows(0), cols(0), allocator(0), usageFlags(_usageFlags), u(0), offset(0), size(&rows)
 {}
 
@@ -1318,88 +1318,6 @@ UMat UMat::t() const
     return m;
 }
 
-UMat UMat::inv(int method) const
-{
-    UMat m;
-    invert(*this, m, method);
-    return m;
-}
-
-UMat UMat::mul(InputArray m, double scale) const
-{
-    UMat dst;
-    multiply(*this, m, dst, scale);
-    return dst;
-}
-
-#ifdef HAVE_OPENCL
-
-static bool ocl_dot( InputArray _src1, InputArray _src2, double & res )
-{
-    UMat src1 = _src1.getUMat().reshape(1), src2 = _src2.getUMat().reshape(1);
-
-    int type = src1.type(), depth = CV_MAT_DEPTH(type),
-            kercn = ocl::predictOptimalVectorWidth(src1, src2);
-    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
-
-    if ( !doubleSupport && depth == CV_64F )
-        return false;
-
-    int dbsize = ocl::Device::getDefault().maxComputeUnits();
-    size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
-    int ddepth = std::max(CV_32F, depth);
-
-    int wgs2_aligned = 1;
-    while (wgs2_aligned < (int)wgs)
-        wgs2_aligned <<= 1;
-    wgs2_aligned >>= 1;
-
-    char cvt[40];
-    ocl::Kernel k("reduce", ocl::core::reduce_oclsrc,
-                  format("-D srcT=%s -D srcT1=%s -D dstT=%s -D dstTK=%s -D ddepth=%d -D convertToDT=%s -D OP_DOT "
-                         "-D WGS=%d -D WGS2_ALIGNED=%d%s%s%s -D kercn=%d",
-                         ocl::typeToStr(CV_MAKE_TYPE(depth, kercn)), ocl::typeToStr(depth),
-                         ocl::typeToStr(ddepth), ocl::typeToStr(CV_MAKE_TYPE(ddepth, kercn)),
-                         ddepth, ocl::convertTypeStr(depth, ddepth, kercn, cvt),
-                         (int)wgs, wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : "",
-                         _src1.isContinuous() ? " -D HAVE_SRC_CONT" : "",
-                         _src2.isContinuous() ? " -D HAVE_SRC2_CONT" : "", kercn));
-    if (k.empty())
-        return false;
-
-    UMat db(1, dbsize, ddepth);
-
-    ocl::KernelArg src1arg = ocl::KernelArg::ReadOnlyNoSize(src1),
-            src2arg = ocl::KernelArg::ReadOnlyNoSize(src2),
-            dbarg = ocl::KernelArg::PtrWriteOnly(db);
-
-    k.args(src1arg, src1.cols, (int)src1.total(), dbsize, dbarg, src2arg);
-
-    size_t globalsize = dbsize * wgs;
-    if (k.run(1, &globalsize, &wgs, true))
-    {
-        res = sum(db.getMat(ACCESS_READ))[0];
-        return true;
-    }
-    return false;
-}
-
-#endif
-
-double UMat::dot(InputArray m) const
-{
-    CV_INSTRUMENT_REGION();
-
-    CV_Assert(m.sameSize(*this) && m.type() == type());
-
-#ifdef HAVE_OPENCL
-    double r = 0;
-    CV_OCL_RUN_(dims <= 2, ocl_dot(*this, m, r), r)
-#endif
-
-    return getMat(ACCESS_READ).dot(m);
-}
-
 UMat UMat::zeros(int rows, int cols, int type)
 {
     return UMat(rows, cols, type, Scalar::all(0));
@@ -1430,18 +1348,6 @@ UMat UMat::ones(int ndims, const int* sz, int type)
     return UMat(ndims, sz, type, Scalar(1));
 }
 
-UMat UMat::eye(int rows, int cols, int type)
-{
-    return UMat::eye(Size(cols, rows), type);
-}
-
-UMat UMat::eye(Size size, int type)
-{
-    UMat m(size, type);
-    setIdentity(m);
-    return m;
-}
-
 }
 
 /* End of file. */

modules/core/src/utils/filesystem.cpp

@@ -587,3 +587,8 @@ cv::String getCacheDirectory(const char* /*sub_directory_name*/, const char* /*c
 #endif // OPENCV_HAVE_FILESYSTEM_SUPPORT
 
 }}} // namespace
+
+
+#if OPENCV_HAVE_FILESYSTEM_SUPPORT
+#include "plugin_loader.impl.hpp"
+#endif

modules/core/src/utils/plugin_loader.impl.hpp

@@ -0,0 +1,80 @@
+// 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.
+
+//
+// Not a standalone header, part of filesystem.cpp
+//
+
+#include "opencv2/core/utils/plugin_loader.private.hpp"
+
+#if !OPENCV_HAVE_FILESYSTEM_SUPPORT
+#error "Invalid build configuration"
+#endif
+
+#if 0  // TODO
+#ifdef NDEBUG
+#define CV_LOG_STRIP_LEVEL CV_LOG_LEVEL_DEBUG + 1
+#else
+#define CV_LOG_STRIP_LEVEL CV_LOG_LEVEL_VERBOSE + 1
+#endif
+#include <opencv2/core/utils/logger.hpp>
+#endif
+
+namespace cv { namespace plugin { namespace impl {
+
+DynamicLib::DynamicLib(const FileSystemPath_t& filename)
+    : handle(0), fname(filename), disableAutoUnloading_(false)
+{
+    libraryLoad(filename);
+}
+
+DynamicLib::~DynamicLib()
+{
+    if (!disableAutoUnloading_)
+    {
+        libraryRelease();
+    }
+    else if (handle)
+    {
+        CV_LOG_INFO(NULL, "skip auto unloading (disabled): " << toPrintablePath(fname));
+        handle = 0;
+    }
+}
+
+void* DynamicLib::getSymbol(const char* symbolName) const
+{
+    if (!handle)
+    {
+        return 0;
+    }
+    void* res = getSymbol_(handle, symbolName);
+    if (!res)
+    {
+        CV_LOG_DEBUG(NULL, "No symbol '" << symbolName << "' in " << toPrintablePath(fname));
+    }
+    return res;
+}
+
+const std::string DynamicLib::getName() const
+{
+    return toPrintablePath(fname);
+}
+
+void DynamicLib::libraryLoad(const FileSystemPath_t& filename)
+{
+    handle = libraryLoad_(filename);
+    CV_LOG_INFO(NULL, "load " << toPrintablePath(filename) << " => " << (handle ? "OK" : "FAILED"));
+}
+
+void DynamicLib::libraryRelease()
+{
+    if (handle)
+    {
+        CV_LOG_INFO(NULL, "unload "<< toPrintablePath(fname));
+        libraryRelease_(handle);
+        handle = 0;
+    }
+}
+
+}}}  // namespace

modules/core/src/va_intel.cpp

@@ -33,6 +33,17 @@ using namespace cv;
 #endif
 #endif
 
+#ifdef HAVE_VA
+#ifndef OPENCV_LIBVA_LINK
+#include "va_wrapper.impl.hpp"
+#else
+namespace cv { namespace detail {
+static void init_libva() { /* nothing */ }
+}}  // namespace
+#endif
+using namespace cv::detail;
+#endif
+
 namespace cv { namespace va_intel {
 
 #ifdef HAVE_VA_INTEL
@@ -54,6 +65,8 @@ Context& initializeContextFromVA(VADisplay display, bool tryInterop)
 #if !defined(HAVE_VA)
     NO_VA_SUPPORT_ERROR;
 #else  // !HAVE_VA
+    init_libva();
+
 #   ifdef HAVE_VA_INTEL
     contextInitialized = false;
     if (tryInterop)
@@ -176,7 +189,7 @@ static bool ocl_convert_nv12_to_bgr(cl_mem clImageY, cl_mem clImageUV, cl_mem cl
 
     k.args(clImageY, clImageUV, clBuffer, step, cols, rows);
 
-    size_t globalsize[] = { (size_t)cols, (size_t)rows };
+    size_t globalsize[] = { (size_t)cols/2, (size_t)rows/2 };
     return k.run(2, globalsize, 0, false);
 }
 
@@ -189,7 +202,7 @@ static bool ocl_convert_bgr_to_nv12(cl_mem clBuffer, int step, int cols, int row
 
     k.args(clBuffer, step, cols, rows, clImageY, clImageUV);
 
-    size_t globalsize[] = { (size_t)cols, (size_t)rows };
+    size_t globalsize[] = { (size_t)cols/2, (size_t)rows/2 };
     return k.run(2, globalsize, 0, false);
 }
 #endif // HAVE_VA_INTEL
@@ -507,6 +520,8 @@ void convertToVASurface(VADisplay display, InputArray src, VASurfaceID surface,
 #if !defined(HAVE_VA)
     NO_VA_SUPPORT_ERROR;
 #else  // !HAVE_VA
+    init_libva();
+
     const int stype = CV_8UC3;
 
     int srcType = src.type();
@@ -611,6 +626,8 @@ void convertFromVASurface(VADisplay display, VASurfaceID surface, Size size, Out
 #if !defined(HAVE_VA)
     NO_VA_SUPPORT_ERROR;
 #else  // !HAVE_VA
+    init_libva();
+
     const int dtype = CV_8UC3;
 
     // TODO Need to specify ACCESS_WRITE here somehow to prevent useless data copying!

modules/core/src/va_wrapper.impl.hpp

@@ -0,0 +1,85 @@
+// 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.
+
+//
+// Not a standalone header, part of va_intel.cpp
+//
+
+#include "opencv2/core/utils/plugin_loader.private.hpp"  // DynamicLib
+
+namespace cv { namespace detail {
+
+typedef VAStatus (*FN_vaDeriveImage)(VADisplay dpy, VASurfaceID surface, VAImage *image);
+typedef VAStatus (*FN_vaDestroyImage)(VADisplay dpy, VAImageID image);
+typedef VAStatus (*FN_vaMapBuffer)(VADisplay dpy, VABufferID buf_id, void **pbuf);
+typedef VAStatus (*FN_vaSyncSurface)(VADisplay dpy, VASurfaceID render_target);
+typedef VAStatus (*FN_vaUnmapBuffer)(VADisplay dpy, VABufferID buf_id);
+
+static FN_vaDeriveImage fn_vaDeriveImage = NULL;
+static FN_vaDestroyImage fn_vaDestroyImage = NULL;
+static FN_vaMapBuffer fn_vaMapBuffer = NULL;
+static FN_vaSyncSurface fn_vaSyncSurface = NULL;
+static FN_vaUnmapBuffer fn_vaUnmapBuffer = NULL;
+
+#define vaDeriveImage fn_vaDeriveImage
+#define vaDestroyImage fn_vaDestroyImage
+#define vaMapBuffer fn_vaMapBuffer
+#define vaSyncSurface fn_vaSyncSurface
+#define vaUnmapBuffer fn_vaUnmapBuffer
+
+
+static std::shared_ptr<cv::plugin::impl::DynamicLib> loadLibVA()
+{
+    std::shared_ptr<cv::plugin::impl::DynamicLib> lib;
+    const char* envPath = getenv("OPENCV_LIBVA_RUNTIME");
+    if (envPath)
+    {
+        lib = std::make_shared<cv::plugin::impl::DynamicLib>(envPath);
+        return lib;
+    }
+    static const char* const candidates[] = {
+        "libva.so",
+        "libva.so.2",
+        "libva.so.1",
+    };
+    for (int i = 0; i < 3; ++i)
+    {
+        lib = std::make_shared<cv::plugin::impl::DynamicLib>(candidates[i]);
+        if (lib->isLoaded())
+            break;
+    }
+    return lib;
+}
+static void init_libva()
+{
+    static bool initialized = false;
+    static auto library = loadLibVA();
+    if (!initialized)
+    {
+        if (!library || !library->isLoaded())
+        {
+            library.reset();
+            CV_Error(cv::Error::StsBadFunc, "OpenCV can't load VA library (libva)");
+        }
+        auto& lib = *library.get();
+#define VA_LOAD_SYMBOL(name) fn_ ## name = reinterpret_cast<FN_ ## name>(lib.getSymbol(#name)); \
+        if (!fn_ ## name) \
+        { \
+            library.reset(); \
+            initialized = true; \
+            CV_Error_(cv::Error::StsBadFunc, ("OpenCV can't load VA library (libva), missing symbol: %s", #name)); \
+        }
+
+        VA_LOAD_SYMBOL(vaDeriveImage);
+        VA_LOAD_SYMBOL(vaDestroyImage);
+        VA_LOAD_SYMBOL(vaMapBuffer);
+        VA_LOAD_SYMBOL(vaSyncSurface);
+        VA_LOAD_SYMBOL(vaUnmapBuffer);
+        initialized = true;
+    }
+    if (!library)
+        CV_Error(cv::Error::StsBadFunc, "OpenCV can't load/initialize VA library (libva)");
+}
+
+}}  // namespace