Merge branch 4.x

Author: asmorkalov

.github/workflows/PR-5.x.yaml

@@ -14,6 +14,9 @@ jobs:
   Ubuntu2004-ARM64:
     uses: opencv/ci-gha-workflow/.github/workflows/OCV-Contrib-PR-5.x-ARM64.yaml@main
 
+  Ubuntu2004-ARM64-FastCV:
+    uses: opencv/ci-gha-workflow/.github/workflows/OCV-Contrib-PR-4.x-ARM64-FastCV.yaml@main
+
   Ubuntu2004-x64-CUDA:
     uses: opencv/ci-gha-workflow/.github/workflows/OCV-Contrib-PR-5.x-U20-Cuda.yaml@main
 

modules/bgsegm/tutorials/bgsegm_bg_subtraction/bgsegm_bg_subtraction.markdown

@@ -30,8 +30,8 @@ In the following, we will have a look at two algorithms from the `bgsegm` module
 ### BackgroundSubtractorMOG
 
 It is a Gaussian Mixture-based Background/Foreground Segmentation Algorithm. It was introduced in
-the paper "An improved adaptive background mixture model for real-time tracking with shadow
-detection" by P. KadewTraKuPong and R. Bowden in 2001. It uses a method to model each background
+the paper "An Improved Adaptive Background Mixture Model for Realtime Tracking with Shadow
+Detection" by P. KaewTraKulPong and R. Bowden in 2001. It uses a method to model each background
 pixel by a mixture of K Gaussian distributions (K = 3 to 5). The weights of the mixture represent
 the time proportions that those colours stay in the scene. The probable background colours are the
 ones which stay longer and more static.

modules/cudaimgproc/src/color.cpp

@@ -70,8 +70,8 @@ namespace cv { namespace cuda {
         template <int cn>
         void Bayer2BGR_16u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
 
-        template <int cn>
-        void MHCdemosaic(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
+        template <int cn, typename Depth>
+        void MHCdemosaic(PtrStepSz<Depth> src, int2 sourceOffset, PtrStepSz<Depth> dst, int2 firstRed, cudaStream_t stream);
     }
 }}
 
@@ -2136,7 +2136,7 @@ void cv::cuda::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn,
         GpuMat src = _src.getGpuMat();
         const int depth = _src.depth();
 
-        CV_Assert( depth == CV_8U );
+        CV_Assert( depth == CV_8U || depth == CV_16U);
         CV_Assert( src.channels() == 1 );
         CV_Assert( dcn == 3 || dcn == 4 );
 
@@ -2148,16 +2148,27 @@ void cv::cuda::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn,
         Size wholeSize;
         Point ofs;
         src.locateROI(wholeSize, ofs);
-        PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);
 
         const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1,
                                         code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1);
 
-        if (dcn == 3)
-            cv::cuda::device::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
-        else
-            cv::cuda::device::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
-
+        if (dcn == 3) {
+            if (depth == CV_8U) {
+                PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);
+                cv::cuda::device::MHCdemosaic<3, uchar>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+            } else {
+                PtrStepSz<ushort> srcWhole(wholeSize.height, wholeSize.width, src.ptr<ushort>(), src.step);
+                cv::cuda::device::MHCdemosaic<3, ushort>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+            }
+        } else {
+            if (depth == CV_8U) {
+                PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);
+                cv::cuda::device::MHCdemosaic<4, uchar>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+            } else {
+                PtrStepSz<ushort> srcWhole(wholeSize.height, wholeSize.width, src.ptr<ushort>(), src.step);
+                cv::cuda::device::MHCdemosaic<4, ushort>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+            }
+        }
         break;
     }
 
@@ -2166,7 +2177,7 @@ void cv::cuda::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn,
         GpuMat src = _src.getGpuMat();
         const int depth = _src.depth();
 
-        CV_Assert( depth == CV_8U );
+        CV_Assert( depth == CV_8U || depth == CV_16U);
 
         _dst.create(_src.size(), CV_MAKE_TYPE(depth, 1));
         GpuMat dst = _dst.getGpuMat();
@@ -2176,12 +2187,17 @@ void cv::cuda::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn,
         Size wholeSize;
         Point ofs;
         src.locateROI(wholeSize, ofs);
-        PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);
 
         const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1,
                                         code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1);
 
-        cv::cuda::device::MHCdemosaic<1>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+        if (depth == CV_8U) {
+            PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);
+            cv::cuda::device::MHCdemosaic<1, uchar>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+        } else {
+            PtrStepSz<ushort> srcWhole(wholeSize.height, wholeSize.width, src.ptr<ushort>(), src.step);
+            cv::cuda::device::MHCdemosaic<1, ushort>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+        }
 
         break;
     }

modules/cudaimgproc/src/cuda/debayer.cu

@@ -390,6 +390,17 @@ namespace cv { namespace cuda { namespace device
     //
     // ported to CUDA
 
+    template<typename Depth> __device__
+    typename TypeVec<Depth, 3>::vec_type make_3(Depth x, Depth y, Depth z);
+
+    template<> __device__ TypeVec<uchar, 3>::vec_type make_3<uchar>(uchar x, uchar y, uchar z) {
+        return make_uchar3(x, y, z);
+    }
+
+    template<> __device__ TypeVec<ushort, 3>::vec_type make_3<ushort>(ushort x, ushort y, ushort z) {
+        return make_ushort3(x, y, z);
+    }
+
     template <typename DstType, class Ptr2D>
     __global__ void MHCdemosaic(PtrStepSz<DstType> dst, Ptr2D src, const int2 firstRed)
     {
@@ -506,34 +517,36 @@ namespace cv { namespace cuda { namespace device
         alternate.x = (x + firstRed.x) % 2;
         alternate.y = (y + firstRed.y) % 2;
 
-        // in BGR sequence;
-        uchar3 pixelColor =
+        typedef typename VecTraits<DstType>::elem_type SrcElemType;
+        typedef typename TypeVec<SrcElemType, 3>::vec_type SrcType;
+
+        SrcType pixelColor =
             (alternate.y == 0) ?
                 ((alternate.x == 0) ?
-                    make_uchar3(saturate_cast<uchar>(PATTERN.y), saturate_cast<uchar>(PATTERN.x), saturate_cast<uchar>(C)) :
-                    make_uchar3(saturate_cast<uchar>(PATTERN.w), saturate_cast<uchar>(C), saturate_cast<uchar>(PATTERN.z))) :
+                    make_3<SrcElemType>(saturate_cast<SrcElemType>(PATTERN.y), saturate_cast<SrcElemType>(PATTERN.x), saturate_cast<SrcElemType>(C)) :
+                    make_3<SrcElemType>(saturate_cast<SrcElemType>(PATTERN.w), saturate_cast<SrcElemType>(C), saturate_cast<SrcElemType>(PATTERN.z))) :
                 ((alternate.x == 0) ?
-                    make_uchar3(saturate_cast<uchar>(PATTERN.z), saturate_cast<uchar>(C), saturate_cast<uchar>(PATTERN.w)) :
-                    make_uchar3(saturate_cast<uchar>(C), saturate_cast<uchar>(PATTERN.x), saturate_cast<uchar>(PATTERN.y)));
+                    make_3<SrcElemType>(saturate_cast<SrcElemType>(PATTERN.z), saturate_cast<SrcElemType>(C), saturate_cast<SrcElemType>(PATTERN.w)) :
+                    make_3<SrcElemType>(saturate_cast<SrcElemType>(C), saturate_cast<SrcElemType>(PATTERN.x), saturate_cast<SrcElemType>(PATTERN.y)));
 
         dst(y, x) = toDst<DstType>(pixelColor);
     }
 
-    template <int cn>
-    void MHCdemosaic(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream)
+    template <int cn, typename Depth>
+    void MHCdemosaic(PtrStepSz<Depth> src, int2 sourceOffset, PtrStepSz<Depth> dst, int2 firstRed, cudaStream_t stream)
     {
-        typedef typename TypeVec<uchar, cn>::vec_type dst_t;
+        typedef typename TypeVec<Depth, cn>::vec_type dst_t;
 
         const dim3 block(32, 8);
         const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
 
         if (sourceOffset.x || sourceOffset.y) {
-            cv::cudev::TextureOff<uchar> texSrc(src, sourceOffset.y, sourceOffset.x);
-            MHCdemosaic<dst_t, cv::cudev::TextureOffPtr<uchar>><<<grid, block, 0, stream>>>((PtrStepSz<dst_t>)dst, texSrc, firstRed);
+            cv::cudev::TextureOff<Depth> texSrc(src, sourceOffset.y, sourceOffset.x);
+            MHCdemosaic<dst_t, cv::cudev::TextureOffPtr<Depth>><<<grid, block, 0, stream>>>((PtrStepSz<dst_t>)dst, texSrc, firstRed);
         }
         else {
-            cv::cudev::Texture<uchar> texSrc(src);
-            MHCdemosaic<dst_t, cv::cudev::TexturePtr<uchar>><<<grid, block, 0, stream>>>((PtrStepSz<dst_t>)dst, texSrc, firstRed);
+            cv::cudev::Texture<Depth> texSrc(src);
+            MHCdemosaic<dst_t, cv::cudev::TexturePtr<Depth>><<<grid, block, 0, stream>>>((PtrStepSz<dst_t>)dst, texSrc, firstRed);
         }
 
         cudaSafeCall( cudaGetLastError() );
@@ -542,9 +555,12 @@ namespace cv { namespace cuda { namespace device
             cudaSafeCall( cudaDeviceSynchronize() );
     }
 
-    template void MHCdemosaic<1>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
-    template void MHCdemosaic<3>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
-    template void MHCdemosaic<4>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
+    template void MHCdemosaic<1, uchar>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
+    template void MHCdemosaic<3, uchar>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
+    template void MHCdemosaic<4, uchar>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
+    template void MHCdemosaic<1, ushort>(PtrStepSz<ushort> src, int2 sourceOffset, PtrStepSz<ushort> dst, int2 firstRed, cudaStream_t stream);
+    template void MHCdemosaic<3, ushort>(PtrStepSz<ushort> src, int2 sourceOffset, PtrStepSz<ushort> dst, int2 firstRed, cudaStream_t stream);
+    template void MHCdemosaic<4, ushort>(PtrStepSz<ushort> src, int2 sourceOffset, PtrStepSz<ushort> dst, int2 firstRed, cudaStream_t stream);
 }}}
 
 #endif /* CUDA_DISABLER */

modules/cudev/test/CMakeLists.txt

@@ -20,9 +20,16 @@ if(OCV_DEPENDENCIES_FOUND)
     ocv_check_windows_crt_linkage()
     set(target_libs ${target_libs} ${CUDA_LIBRARIES})
     if(CUDA_VERSION VERSION_LESS "11.0")
-      ocv_update(OPENCV_CUDA_OPTIONS_opencv_test_cudev "-std=c++11")
+      # Windows version does not support --std option
+      if(UNIX OR APPLE)
+        ocv_update(OPENCV_CUDA_OPTIONS_opencv_test_cudev "-std=c++11")
+      endif()
     else()
-      ocv_update(OPENCV_CUDA_OPTIONS_opencv_test_cudev "-std=c++14")
+      if(CUDA_VERSION VERSION_LESS "12.8")
+        ocv_update(OPENCV_CUDA_OPTIONS_opencv_test_cudev "-std=c++14")
+      else()
+        ocv_update(OPENCV_CUDA_OPTIONS_opencv_test_cudev "-std=c++17")
+      endif()
       ocv_warnings_disable(CMAKE_CXX_FLAGS -Wdeprecated-declarations)
     endif()
     CUDA_ADD_EXECUTABLE(${the_target} ${OPENCV_TEST_${the_module}_SOURCES} OPTIONS ${OPENCV_CUDA_OPTIONS_opencv_test_cudev})

modules/fastcv/include/opencv2/fastcv.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024 Qualcomm Innovation Center, Inc. All rights reserved.
+ * Copyright (c) 2024-2025 Qualcomm Innovation Center, Inc. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
 */
 
@@ -11,6 +11,7 @@
 #include "opencv2/fastcv/arithm.hpp"
 #include "opencv2/fastcv/bilateralFilter.hpp"
 #include "opencv2/fastcv/blur.hpp"
+#include "opencv2/fastcv/channel.hpp"
 #include "opencv2/fastcv/cluster.hpp"
 #include "opencv2/fastcv/draw.hpp"
 #include "opencv2/fastcv/edges.hpp"

modules/fastcv/include/opencv2/fastcv/arithm.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024 Qualcomm Innovation Center, Inc. All rights reserved.
+ * Copyright (c) 2024-2025 Qualcomm Innovation Center, Inc. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
 */
 
@@ -8,6 +8,10 @@
 
 #include <opencv2/core.hpp>
 
+#define FCV_CMP_EQ(val1,val2) (fabs(val1 - val2) < FLT_EPSILON)
+
+#define FCV_OPTYPE(depth,op) ((depth<<3) + op)
+
 namespace cv {
 namespace fastcv {
 
@@ -26,6 +30,59 @@ CV_EXPORTS_W void matmuls8s32(InputArray src1, InputArray src2, OutputArray dst)
 
 //! @}
 
+//! @addtogroup fastcv
+//! @{
+
+/**
+ * @brief Arithmetic add and subtract operations for two matrices
+ *        It is optimized for Qualcomm's processors
+ * @param src1 First source matrix, can be of type CV_8U, CV_16S, CV_32F.
+ *             Note: CV_32F not supported for subtract
+ * @param src2 Second source matrix of same type and size as src1
+ * @param dst Resulting matrix of type as src mats
+ * @param op  type of operation - 0 for add and 1 for subtract
+ */
+CV_EXPORTS_W void arithmetic_op(InputArray src1, InputArray src2, OutputArray dst, int op);
+
+//! @}
+
+//! @addtogroup fastcv
+//! @{
+
+/**
+ * @brief Matrix multiplication of two float type matrices
+ *        R = a*A*B + b*C where A,B,C,R are matrices and a,b are constants
+ *        It is optimized for Qualcomm's processors
+ * @param src1 First source matrix of type CV_32F
+ * @param src2 Second source matrix of type CV_32F with same rows as src1 cols
+ * @param dst Resulting matrix of type CV_32F
+ * @param alpha multiplying factor for src1 and src2
+ * @param src3 Optional third matrix of type CV_32F to be added to matrix product
+ * @param beta multiplying factor for src3
+ */
+CV_EXPORTS_W void gemm(InputArray src1, InputArray src2, OutputArray dst, float alpha = 1.0,
+                           InputArray src3 = noArray(), float beta = 0.0);
+
+//! @}
+
+//! @addtogroup fastcv
+//! @{
+
+/**
+ * @brief Integral of a YCbCr420 image.
+ *        Note: Input height should be multiple of 2. Input width and stride should be multiple of 16.
+ *              Output stride should be multiple of 8.
+ *              It is optimized for Qualcomm's processors
+ * @param Y Input Y component of 8UC1 YCbCr420 image.
+ * @param CbCr Input CbCr component(interleaved) of 8UC1 YCbCr420 image.
+ * @param IY Output Y integral of CV_32S one channel, size (Y height + 1)*(Y width + 1)
+ * @param ICb Output Cb integral of CV_32S one channel, size (Y height/2 + 1)*(Y width/2 + 1)
+ * @param ICr Output Cr integral of CV_32S one channel, size (Y height/2 + 1)*(Y width/2 + 1)
+ */
+CV_EXPORTS_W void integrateYUV(InputArray Y, InputArray CbCr, OutputArray IY, OutputArray ICb, OutputArray ICr);
+
+//! @}
+
 } // fastcv::
 } // cv::
 

modules/fastcv/include/opencv2/fastcv/blur.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024 Qualcomm Innovation Center, Inc. All rights reserved.
+ * Copyright (c) 2024-2025 Qualcomm Innovation Center, Inc. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
 */
 
@@ -58,6 +58,22 @@ CV_EXPORTS_W void filter2D(InputArray _src, OutputArray _dst, int ddepth, InputA
 CV_EXPORTS_W void sepFilter2D(InputArray _src, OutputArray _dst, int ddepth, InputArray _kernelX, InputArray _kernelY);
 //! @}
 
+//! @addtogroup fastcv
+//! @{
+
+/**
+ * @brief Calculates the local subtractive and contrastive normalization of the image.
+ *        Each pixel of the image is normalized by the mean and standard deviation of the patch centred at the pixel.
+ *        It is optimized for Qualcomm's processors.
+ * @param _src Input image, should have one channel CV_8U or CV_32F
+ * @param _dst Output array, should be one channel, CV_8S if src of type CV_8U, or CV_32F if src of CV_32F
+ * @param pSize Patch size for mean and std dev calculation
+ * @param useStdDev If 1, bot mean and std dev will be used for normalization, if 0, only mean used
+ */
+CV_EXPORTS_W void normalizeLocalBox(InputArray _src, OutputArray _dst, Size pSize, bool useStdDev);
+
+//! @}
+
 } // fastcv::
 } // cv::
 

modules/fastcv/include/opencv2/fastcv/channel.hpp

@@ -0,0 +1,45 @@
+/*
+ * Copyright (c) 2025 Qualcomm Innovation Center, Inc. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+*/
+
+#ifndef OPENCV_FASTCV_CHANNEL_HPP
+#define OPENCV_FASTCV_CHANNEL_HPP
+
+#include <opencv2/core.hpp>
+
+namespace cv {
+namespace fastcv {
+
+//! @addtogroup fastcv
+//! @{
+
+/**
+ * @brief Creates one multi-channel mat out of several single-channel CV_8U mats.
+ *        Optimized for Qualcomm's processors
+ * @param mv input vector of matrices to be merged; all the matrices in mv must be of CV_8UC1 and have the same size
+ *           Note: numbers of mats can be 2,3 or 4.
+ * @param dst output array of depth CV_8U and same size as mv[0]; The number of channels
+ *            will be the total number of matrices in the matrix array
+ */
+CV_EXPORTS_W void merge(InputArrayOfArrays mv, OutputArray dst);
+
+//! @}
+
+//! @addtogroup fastcv
+//! @{
+
+/**
+ * @brief Splits an CV_8U multi-channel mat into several CV_8UC1 mats
+ *        Optimized for Qualcomm's processors
+ * @param src input 2,3 or 4 channel mat of depth CV_8U
+ * @param mv  output vector of size src.channels() of CV_8UC1 mats
+ */
+CV_EXPORTS_W void split(InputArray src, OutputArrayOfArrays mv);
+
+//! @}
+
+} // fastcv::
+} // cv::
+
+#endif // OPENCV_FASTCV_CHANNEL_HPP

modules/fastcv/include/opencv2/fastcv/pyramid.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024 Qualcomm Innovation Center, Inc. All rights reserved.
+ * Copyright (c) 2024-2025 Qualcomm Innovation Center, Inc. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
 */
 
@@ -16,7 +16,7 @@ namespace fastcv {
 
 /**
  * @brief Creates a gradient pyramid from an image pyramid
- *
+ *        Note: The borders are ignored during gradient calculation.
  * @param pyr Input pyramid of 1-channel 8-bit images. Only continuous images are supported.
  * @param dx Horizontal Sobel gradient pyramid of the same size as pyr
  * @param dy Verical Sobel gradient pyramid of the same size as pyr