Version++.

(cherry picked from commit d58d277707)

cmake/OpenCVPackaging.cmake

@@ -104,6 +104,14 @@ if(HAVE_CUDA)
   endif()
 endif()
 
+if(HAVE_TBB AND NOT BUILD_TBB)
+  if(CPACK_DEB_DEV_PACKAGE_DEPENDS)
+    set(CPACK_DEB_DEV_PACKAGE_DEPENDS "${CPACK_DEB_DEV_PACKAGE_DEPENDS}, libtbb-dev")
+  else()
+    set(CPACK_DEB_DEV_PACKAGE_DEPENDS "libtbb-dev")
+  endif()
+endif()
+
 if(NOT OPENCV_CUSTOM_PACKAGE_INFO)
   set(CPACK_COMPONENT_LIBS_DESCRIPTION "Open Computer Vision Library")
   set(CPACK_DEBIAN_COMPONENT_LIBS_NAME "libopencv")

modules/core/include/opencv2/core/version.hpp

@@ -50,7 +50,7 @@
 #define CV_VERSION_EPOCH    2
 #define CV_VERSION_MAJOR    4
 #define CV_VERSION_MINOR    10
-#define CV_VERSION_REVISION 1
+#define CV_VERSION_REVISION 4
 
 #define CVAUX_STR_EXP(__A)  #__A
 #define CVAUX_STR(__A)      CVAUX_STR_EXP(__A)

modules/gpu/perf/perf_denoising.cpp

@@ -99,7 +99,8 @@ PERF_TEST_P(Sz_Depth_Cn_KernelSz, Denoising_BilateralFilter,
 
 DEF_PARAM_TEST(Sz_Depth_Cn_WinSz_BlockSz, cv::Size, MatDepth, MatCn, int, int);
 
-PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, Denoising_NonLocalMeans,
+// disabled, since it takes too much time
+PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, DISABLED_Denoising_NonLocalMeans,
             Combine(GPU_DENOISING_IMAGE_SIZES,
                     Values<MatDepth>(CV_8U),
                     GPU_CHANNELS_1_3,
@@ -143,7 +144,8 @@ PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, Denoising_NonLocalMeans,
 
 DEF_PARAM_TEST(Sz_Depth_Cn_WinSz_BlockSz, cv::Size, MatDepth, MatCn, int, int);
 
-PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, Denoising_FastNonLocalMeans,
+// disabled, since it takes too much time
+PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, DISABLED_Denoising_FastNonLocalMeans,
             Combine(GPU_DENOISING_IMAGE_SIZES,
                     Values<MatDepth>(CV_8U),
                     GPU_CHANNELS_1_3,

modules/gpu/perf/perf_video.cpp

@@ -307,7 +307,8 @@ PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, Video_PyrLKOpticalFlowSparse
 
 DEF_PARAM_TEST(ImagePair_WinSz_Levels_Iters, pair_string, int, int, int);
 
-PERF_TEST_P(ImagePair_WinSz_Levels_Iters, Video_PyrLKOpticalFlowDense,
+// Sanity test fails on Maxwell and CUDA 7.0
+PERF_TEST_P(ImagePair_WinSz_Levels_Iters, DISABLED_Video_PyrLKOpticalFlowDense,
             Combine(Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")),
                     Values(3, 5, 7, 9, 13, 17, 21),
                     Values(1, 3),
@@ -463,7 +464,8 @@ void calcOpticalFlowBM(const cv::Mat& prev, const cv::Mat& curr,
     cvCalcOpticalFlowBM(&cvprev, &cvcurr, bSize, shiftSize, maxRange, usePrevious, &cvvelx, &cvvely);
 }
 
-PERF_TEST_P(ImagePair, Video_OpticalFlowBM,
+// disabled, since it takes too much time
+PERF_TEST_P(ImagePair, DISABLED_Video_OpticalFlowBM,
             Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
 {
     declare.time(400);
@@ -541,7 +543,8 @@ PERF_TEST_P(ImagePair, DISABLED_Video_FastOpticalFlowBM,
 
 DEF_PARAM_TEST_1(Video, string);
 
-PERF_TEST_P(Video, Video_FGDStatModel,
+// disabled, since it takes too much time
+PERF_TEST_P(Video, DISABLED_Video_FGDStatModel,
             Values(string("gpu/video/768x576.avi")))
 {
     const int numIters = 10;

modules/gpu/src/calib3d.cpp

@@ -138,14 +138,16 @@ namespace
             bool was;
             do
             {
-                subset[i] = rand() % num_points;
+                subset[i] = cv::theRNG().uniform(0, num_points);
                 was = false;
                 for (int j = 0; j < i; ++j)
+                {
                     if (subset[j] == subset[i])
                     {
                         was = true;
                         break;
                     }
+                }
             } while (was);
         }
     }

modules/gpu/src/cuda/element_operations.cu

@@ -1896,53 +1896,53 @@ namespace cv { namespace gpu { namespace device
 
 namespace arithm
 {
-    template <typename T> void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream)
+    template <typename T> void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream)
     {
         if (mask.data)
-            transform((PtrStepSz<T>) src, (PtrStepSz<T>) dst, bit_not<T>(), mask, stream);
+            transform((PtrStepSz<T>) src, (PtrStepSz<T>) dst, bit_not<T>(), SingleMaskChannels(mask, num_channels), stream);
         else
             transform((PtrStepSz<T>) src, (PtrStepSz<T>) dst, bit_not<T>(), WithOutMask(), stream);
     }
 
-    template <typename T> void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream)
+    template <typename T> void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream)
     {
         if (mask.data)
-            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_and<T>(), mask, stream);
+            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_and<T>(), SingleMaskChannels(mask, num_channels), stream);
         else
             transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_and<T>(), WithOutMask(), stream);
     }
 
-    template <typename T> void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream)
+    template <typename T> void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream)
     {
         if (mask.data)
-            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_or<T>(), mask, stream);
+            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_or<T>(), SingleMaskChannels(mask, num_channels), stream);
         else
             transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_or<T>(), WithOutMask(), stream);
     }
 
-    template <typename T> void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream)
+    template <typename T> void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream)
     {
         if (mask.data)
-            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_xor<T>(), mask, stream);
+            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_xor<T>(), SingleMaskChannels(mask, num_channels), stream);
         else
             transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_xor<T>(), WithOutMask(), stream);
     }
 
-    template void bitMatNot<uchar>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatNot<uchar>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template void bitMatNot<ushort>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatNot<ushort>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template void bitMatNot<uint>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatNot<uint>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
 
-    template void bitMatAnd<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatAnd<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template void bitMatAnd<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatAnd<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template void bitMatAnd<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatAnd<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
 
-    template void bitMatOr<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatOr<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template void bitMatOr<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatOr<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template void bitMatOr<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatOr<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
 
-    template void bitMatXor<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatXor<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template void bitMatXor<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatXor<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template void bitMatXor<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatXor<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
 }
 
 //////////////////////////////////////////////////////////////////////////////////////

modules/gpu/src/cuda/pyr_down.cu

@@ -114,7 +114,7 @@ namespace cv { namespace gpu { namespace device
 
                     sum =       0.0625f * src(b.idx_row_low (src_y - 2), b.idx_col_high(x));
                     sum = sum + 0.25f   * src(b.idx_row_low (src_y - 1), b.idx_col_high(x));
-                    sum = sum + 0.375f  * src(src_y                    , b.idx_col_high(x));
+                    sum = sum + 0.375f  * src(b.idx_row_high(src_y    ), b.idx_col_high(x));
                     sum = sum + 0.25f   * src(b.idx_row_high(src_y + 1), b.idx_col_high(x));
                     sum = sum + 0.0625f * src(b.idx_row_high(src_y + 2), b.idx_col_high(x));
 
@@ -129,7 +129,7 @@ namespace cv { namespace gpu { namespace device
 
                     sum =       0.0625f * src(b.idx_row_low (src_y - 2), b.idx_col(left_x));
                     sum = sum + 0.25f   * src(b.idx_row_low (src_y - 1), b.idx_col(left_x));
-                    sum = sum + 0.375f  * src(src_y                    , b.idx_col(left_x));
+                    sum = sum + 0.375f  * src(b.idx_row_high(src_y    ), b.idx_col(left_x));
                     sum = sum + 0.25f   * src(b.idx_row_high(src_y + 1), b.idx_col(left_x));
                     sum = sum + 0.0625f * src(b.idx_row_high(src_y + 2), b.idx_col(left_x));
 
@@ -144,7 +144,7 @@ namespace cv { namespace gpu { namespace device
 
                     sum =       0.0625f * src(b.idx_row_low (src_y - 2), b.idx_col_high(right_x));
                     sum = sum + 0.25f   * src(b.idx_row_low (src_y - 1), b.idx_col_high(right_x));
-                    sum = sum + 0.375f  * src(src_y                    , b.idx_col_high(right_x));
+                    sum = sum + 0.375f  * src(b.idx_row_high(src_y    ), b.idx_col_high(right_x));
                     sum = sum + 0.25f   * src(b.idx_row_high(src_y + 1), b.idx_col_high(right_x));
                     sum = sum + 0.0625f * src(b.idx_row_high(src_y + 2), b.idx_col_high(right_x));
 

modules/gpu/src/element_operations.cpp

@@ -1955,7 +1955,7 @@ void cv::gpu::compare(const GpuMat& src, Scalar sc, GpuMat& dst, int cmpop, Stre
 
 namespace arithm
 {
-    template <typename T> void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template <typename T> void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
 }
 
 void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask, Stream& s)
@@ -1964,39 +1964,73 @@ void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask, St
 
     const int depth = src.depth();
 
-    CV_Assert( depth <= CV_64F );
+    CV_Assert( depth < CV_32F );
     CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src.size()) );
 
     dst.create(src.size(), src.type());
 
     cudaStream_t stream = StreamAccessor::getStream(s);
 
-    const int bcols = (int) (src.cols * src.elemSize());
+    if (mask.empty())
-
-    if ((bcols & 3) == 0)
     {
-        const int vcols = bcols >> 2;
+        const int bcols = (int) (src.cols * src.elemSize());
+        bool aligned =
+                isAligned(src.data, sizeof(unsigned int)) &&
+                isAligned(dst.data, sizeof(unsigned int));
 
-        bitMatNot<unsigned int>(
+        if (aligned && (bcols & 3) == 0)
-                    PtrStepSzb(src.rows, vcols, src.data, src.step),
+        {
-                    PtrStepSzb(src.rows, vcols, dst.data, dst.step),
+            const int vcols = bcols >> 2;
-                    mask, stream);
-    }
-    else if ((bcols & 1) == 0)
-    {
-        const int vcols = bcols >> 1;
 
-        bitMatNot<unsigned short>(
+            bitMatNot<unsigned int>(
-                    PtrStepSzb(src.rows, vcols, src.data, src.step),
+                        PtrStepSzb(src.rows, vcols, src.data, src.step),
-                    PtrStepSzb(src.rows, vcols, dst.data, dst.step),
+                        PtrStepSzb(src.rows, vcols, dst.data, dst.step),
-                    mask, stream);
+                        PtrStepb(), 1, stream);
+        }
+        else if (aligned && (bcols & 1) == 0)
+        {
+            const int vcols = bcols >> 1;
+
+            bitMatNot<unsigned short>(
+                        PtrStepSzb(src.rows, vcols, src.data, src.step),
+                        PtrStepSzb(src.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else
+        {
+            bitMatNot<unsigned char>(
+                        PtrStepSzb(src.rows, bcols, src.data, src.step),
+                        PtrStepSzb(src.rows, bcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
     }
     else
     {
-        bitMatNot<unsigned char>(
+        const int elem_size = src.elemSize1();
-                    PtrStepSzb(src.rows, bcols, src.data, src.step),
+        const int num_channels = src.channels();
-                    PtrStepSzb(src.rows, bcols, dst.data, dst.step),
+        const int bcols = src.cols * num_channels;
-                    mask, stream);
+
+        if (elem_size == 1)
+        {
+            bitMatNot<unsigned char>(
+                        PtrStepSzb(src.rows, bcols, src.data, src.step),
+                        PtrStepSzb(src.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 2)
+        {
+            bitMatNot<unsigned short>(
+                        PtrStepSzb(src.rows, bcols, src.data, src.step),
+                        PtrStepSzb(src.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 4)
+        {
+            bitMatNot<unsigned int>(
+                        PtrStepSzb(src.rows, bcols, src.data, src.step),
+                        PtrStepSzb(src.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
     }
 }
 
@@ -2005,9 +2039,9 @@ void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask, St
 
 namespace arithm
 {
-    template <typename T> void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template <typename T> void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template <typename T> void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template <typename T> void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
-    template <typename T> void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template <typename T> void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
 }
 
 void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, Stream& s)
@@ -2016,7 +2050,7 @@ void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, c
 
     const int depth = src1.depth();
 
-    CV_Assert( depth <= CV_64F );
+    CV_Assert( depth < CV_32F );
     CV_Assert( src2.size() == src1.size() && src2.type() == src1.type() );
     CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src1.size()) );
 
@@ -2024,36 +2058,73 @@ void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, c
 
     cudaStream_t stream = StreamAccessor::getStream(s);
 
-    const int bcols = (int) (src1.cols * src1.elemSize());
+    if (mask.empty())
-
-    if ((bcols & 3) == 0)
     {
-        const int vcols = bcols >> 2;
+        const int bcols = (int) (src1.cols * src1.elemSize());
+        bool aligned =
+                isAligned(src1.data, sizeof(unsigned int)) &&
+                isAligned(src2.data, sizeof(unsigned int)) &&
+                isAligned(dst.data, sizeof(unsigned int));
 
-        bitMatAnd<unsigned int>(
+        if (aligned && (bcols & 3) == 0)
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+        {
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+            const int vcols = bcols >> 2;
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
-    }
-    else if ((bcols & 1) == 0)
-    {
-        const int vcols = bcols >> 1;
 
-        bitMatAnd<unsigned short>(
+            bitMatAnd<unsigned int>(
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
+                        PtrStepb(), 1, stream);
+        }
+        else if (aligned && (bcols & 1) == 0)
+        {
+            const int vcols = bcols >> 1;
+
+            bitMatAnd<unsigned short>(
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else
+        {
+            bitMatAnd<unsigned char>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
     }
     else
     {
+        const int elem_size = src1.elemSize1();
+        const int num_channels = src1.channels();
+        const int bcols = src1.cols * num_channels;
 
-        bitMatAnd<unsigned char>(
+        if (elem_size == 1)
-                    PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+        {
-                    PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+            bitMatAnd<unsigned char>(
-                    PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
-                    mask, stream);
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 2)
+        {
+            bitMatAnd<unsigned short>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 4)
+        {
+            bitMatAnd<unsigned int>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
     }
 }
 
@@ -2063,7 +2134,7 @@ void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, co
 
     const int depth = src1.depth();
 
-    CV_Assert( depth <= CV_64F );
+    CV_Assert( depth < CV_32F );
     CV_Assert( src2.size() == src1.size() && src2.type() == src1.type() );
     CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src1.size()) );
 
@@ -2071,36 +2142,73 @@ void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, co
 
     cudaStream_t stream = StreamAccessor::getStream(s);
 
-    const int bcols = (int) (src1.cols * src1.elemSize());
+    if (mask.empty())
-
-    if ((bcols & 3) == 0)
     {
-        const int vcols = bcols >> 2;
+        const int bcols = (int) (src1.cols * src1.elemSize());
+        bool aligned =
+                isAligned(src1.data, sizeof(unsigned int)) &&
+                isAligned(src2.data, sizeof(unsigned int)) &&
+                isAligned(dst.data, sizeof(unsigned int));
 
-        bitMatOr<unsigned int>(
+        if (aligned && (bcols & 3) == 0)
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+        {
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+            const int vcols = bcols >> 2;
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
-    }
-    else if ((bcols & 1) == 0)
-    {
-        const int vcols = bcols >> 1;
 
-        bitMatOr<unsigned short>(
+            bitMatOr<unsigned int>(
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
+                        PtrStepb(), 1, stream);
+        }
+        else if (aligned && (bcols & 1) == 0)
+        {
+            const int vcols = bcols >> 1;
+
+            bitMatOr<unsigned short>(
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else
+        {
+            bitMatOr<unsigned char>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
     }
     else
     {
+        const int elem_size = src1.elemSize1();
+        const int num_channels = src1.channels();
+        const int bcols = src1.cols * num_channels;
 
-        bitMatOr<unsigned char>(
+        if (elem_size == 1)
-                    PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+        {
-                    PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+            bitMatOr<unsigned char>(
-                    PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
-                    mask, stream);
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 2)
+        {
+            bitMatOr<unsigned short>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 4)
+        {
+            bitMatOr<unsigned int>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
     }
 }
 
@@ -2110,7 +2218,7 @@ void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, c
 
     const int depth = src1.depth();
 
-    CV_Assert( depth <= CV_64F );
+    CV_Assert( depth < CV_32F );
     CV_Assert( src2.size() == src1.size() && src2.type() == src1.type() );
     CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src1.size()) );
 
@@ -2118,36 +2226,73 @@ void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, c
 
     cudaStream_t stream = StreamAccessor::getStream(s);
 
-    const int bcols = (int) (src1.cols * src1.elemSize());
+    if (mask.empty())
-
-    if ((bcols & 3) == 0)
     {
-        const int vcols = bcols >> 2;
+        const int bcols = (int) (src1.cols * src1.elemSize());
+        bool aligned =
+                isAligned(src1.data, sizeof(unsigned int)) &&
+                isAligned(src2.data, sizeof(unsigned int)) &&
+                isAligned(dst.data, sizeof(unsigned int));
 
-        bitMatXor<unsigned int>(
+        if (aligned && (bcols & 3) == 0)
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+        {
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+            const int vcols = bcols >> 2;
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
-    }
-    else if ((bcols & 1) == 0)
-    {
-        const int vcols = bcols >> 1;
 
-        bitMatXor<unsigned short>(
+            bitMatXor<unsigned int>(
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
+                        PtrStepb(), 1, stream);
+        }
+        else if (aligned && (bcols & 1) == 0)
+        {
+            const int vcols = bcols >> 1;
+
+            bitMatXor<unsigned short>(
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else
+        {
+            bitMatXor<unsigned char>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
     }
     else
     {
+        const int elem_size = src1.elemSize1();
+        const int num_channels = src1.channels();
+        const int bcols = src1.cols * num_channels;
 
-        bitMatXor<unsigned char>(
+        if (elem_size == 1)
-                    PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+        {
-                    PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+            bitMatXor<unsigned char>(
-                    PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
-                    mask, stream);
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 2)
+        {
+            bitMatXor<unsigned short>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 4)
+        {
+            bitMatXor<unsigned int>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
     }
 }
 

modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu

@@ -280,7 +280,8 @@ __global__ void scanRows(T_in *d_src, Ncv32u texOffs, Ncv32u srcWidth, Ncv32u sr
 
     __shared__ T_out shmem[NUM_SCAN_THREADS * 2];
     __shared__ T_out carryElem;
-    carryElem = 0;
+    if (threadIdx.x == 0)
+        carryElem = 0;
     __syncthreads();
 
     while (numBuckets--)

modules/gpu/test/nvidia/main_nvidia.cpp

@@ -340,7 +340,8 @@ bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel
 bool nvidia_NPPST_Squared_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerSII("NPPST Squared Integral Image", outputLevel);
 
@@ -354,7 +355,8 @@ bool nvidia_NPPST_Squared_Integral_Image(const std::string& test_data_path, Outp
 bool nvidia_NPPST_RectStdDev(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerRStdDev("NPPST RectStdDev", outputLevel);
 
@@ -368,7 +370,8 @@ bool nvidia_NPPST_RectStdDev(const std::string& test_data_path, OutputLevel outp
 bool nvidia_NPPST_Resize(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerResize("NPPST Resize", outputLevel);
 
@@ -384,7 +387,8 @@ bool nvidia_NPPST_Resize(const std::string& test_data_path, OutputLevel outputLe
 bool nvidia_NPPST_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerNPPSTVectorOperations("NPPST Vector Operations", outputLevel);
 
@@ -398,7 +402,8 @@ bool nvidia_NPPST_Vector_Operations(const std::string& test_data_path, OutputLev
 bool nvidia_NPPST_Transpose(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerTranspose("NPPST Transpose", outputLevel);
 
@@ -414,7 +419,8 @@ bool nvidia_NPPST_Transpose(const std::string& test_data_path, OutputLevel outpu
 bool nvidia_NCV_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerVectorOperations("Vector Operations", outputLevel);
 
@@ -429,7 +435,8 @@ bool nvidia_NCV_Vector_Operations(const std::string& test_data_path, OutputLevel
 bool nvidia_NCV_Haar_Cascade_Loader(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerHaarLoader("Haar Cascade Loader", outputLevel);
 
@@ -441,7 +448,8 @@ bool nvidia_NCV_Haar_Cascade_Loader(const std::string& test_data_path, OutputLev
 bool nvidia_NCV_Haar_Cascade_Application(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerHaarAppl("Haar Cascade Application", outputLevel);
 
@@ -455,7 +463,8 @@ bool nvidia_NCV_Haar_Cascade_Application(const std::string& test_data_path, Outp
 bool nvidia_NCV_Hypotheses_Filtration(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerHypFiltration("Hypotheses Filtration", outputLevel);
 
@@ -469,7 +478,8 @@ bool nvidia_NCV_Hypotheses_Filtration(const std::string& test_data_path, OutputL
 bool nvidia_NCV_Visualization(const std::string& test_data_path, OutputLevel outputLevel)
 {
     path = test_data_path;
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);
 
     NCVAutoTestLister testListerVisualize("Visualization", outputLevel);
 

modules/gpu/test/test_calib3d.cpp

@@ -278,18 +278,20 @@ struct SolvePnPRansac : testing::TestWithParam<cv::gpu::DeviceInfo>
 
 GPU_TEST_P(SolvePnPRansac, Accuracy)
 {
-    cv::Mat object = randomMat(cv::Size(5000, 1), CV_32FC3, 0, 100);
+    // Use RNG with fixed seed to be reproducable
-    cv::Mat camera_mat = randomMat(cv::Size(3, 3), CV_32F, 0.5, 1);
+    cv::RNG rng(123456789);
+    cv::theRNG() = rng;
+
+    cv::Mat object = cvtest::randomMat(rng, cv::Size(5000, 1), CV_32FC3, 0, 100, false);
+    cv::Mat camera_mat = cvtest::randomMat(rng, cv::Size(3, 3), CV_32FC1, 0.5, 1, false);
     camera_mat.at<float>(0, 1) = 0.f;
     camera_mat.at<float>(1, 0) = 0.f;
     camera_mat.at<float>(2, 0) = 0.f;
     camera_mat.at<float>(2, 1) = 0.f;
 
     std::vector<cv::Point2f> image_vec;
-    cv::Mat rvec_gold;
+    cv::Mat rvec_gold = cvtest::randomMat(rng, cv::Size(3, 1), CV_32F, 0, 1, false);
-    cv::Mat tvec_gold;
+    cv::Mat tvec_gold = cvtest::randomMat(rng, cv::Size(3, 1), CV_32F, 0, 1, false);
-    rvec_gold = randomMat(cv::Size(3, 1), CV_32F, 0, 1);
-    tvec_gold = randomMat(cv::Size(3, 1), CV_32F, 0, 1);
     cv::projectPoints(object, rvec_gold, tvec_gold, camera_mat, cv::Mat(1, 8, CV_32F, cv::Scalar::all(0)), image_vec);
 
     cv::Mat rvec, tvec;
@@ -298,8 +300,8 @@ GPU_TEST_P(SolvePnPRansac, Accuracy)
                             camera_mat, cv::Mat(1, 8, CV_32F, cv::Scalar::all(0)),
                             rvec, tvec, false, 200, 2.f, 100, &inliers);
 
-    ASSERT_LE(cv::norm(rvec - rvec_gold), 1e-3);
+    ASSERT_LE(cv::norm(rvec - rvec_gold), 2e-3);
-    ASSERT_LE(cv::norm(tvec - tvec_gold), 1e-3);
+    ASSERT_LE(cv::norm(tvec - tvec_gold), 2e-3);
 }
 
 INSTANTIATE_TEST_CASE_P(GPU_Calib3D, SolvePnPRansac, ALL_DEVICES);

modules/gpu/test/test_core.cpp

@@ -1785,72 +1785,95 @@ INSTANTIATE_TEST_CASE_P(GPU_Core, Compare_Scalar, testing::Combine(
 //////////////////////////////////////////////////////////////////////////////
 // Bitwise_Array
 
-PARAM_TEST_CASE(Bitwise_Array, cv::gpu::DeviceInfo, cv::Size, MatType)
+PARAM_TEST_CASE(Bitwise_Array, cv::gpu::DeviceInfo, cv::Size, MatType, UseRoi)
 {
     cv::gpu::DeviceInfo devInfo;
     cv::Size size;
     int type;
+    bool useRoi;
 
     cv::Mat src1;
     cv::Mat src2;
 
+    cv::Mat mask;
+
     virtual void SetUp()
     {
         devInfo = GET_PARAM(0);
         size = GET_PARAM(1);
         type = GET_PARAM(2);
+        useRoi = GET_PARAM(3);
 
         cv::gpu::setDevice(devInfo.deviceID());
 
         src1 = randomMat(size, type, 0.0, std::numeric_limits<int>::max());
         src2 = randomMat(size, type, 0.0, std::numeric_limits<int>::max());
+
+        mask = randomMat(size, CV_8UC1, 0.0, 2.0);
     }
 };
 
 GPU_TEST_P(Bitwise_Array, Not)
 {
-    cv::gpu::GpuMat dst;
+    cv::gpu::GpuMat dst_nomask, dst_mask(src1.size(), src1.type(), cv::Scalar::all(0));
-    cv::gpu::bitwise_not(loadMat(src1), dst);
+    cv::gpu::bitwise_not(loadMat(src1, useRoi), dst_nomask);
+    cv::gpu::bitwise_not(loadMat(src1, useRoi), dst_mask, loadMat(mask, useRoi));
 
-    cv::Mat dst_gold = ~src1;
+    cv::Mat dst_gold_nomask, dst_gold_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::bitwise_not(src1, dst_gold_nomask);
+    cv::bitwise_not(src1, dst_gold_mask, mask);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_nomask, dst_nomask, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_mask, dst_mask, 0.0);
 }
 
 GPU_TEST_P(Bitwise_Array, Or)
 {
-    cv::gpu::GpuMat dst;
+    cv::gpu::GpuMat dst_nomask, dst_mask(src1.size(), src1.type(), cv::Scalar::all(0));
-    cv::gpu::bitwise_or(loadMat(src1), loadMat(src2), dst);
+    cv::gpu::bitwise_or(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_nomask);
+    cv::gpu::bitwise_or(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_mask, loadMat(mask, useRoi));
 
-    cv::Mat dst_gold = src1 | src2;
+    cv::Mat dst_gold_nomask, dst_gold_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::bitwise_or(src1, src2, dst_gold_nomask);
+    cv::bitwise_or(src1, src2, dst_gold_mask, mask);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_nomask, dst_nomask, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_mask, dst_mask, 0.0);
 }
 
 GPU_TEST_P(Bitwise_Array, And)
 {
-    cv::gpu::GpuMat dst;
+    cv::gpu::GpuMat dst_nomask, dst_mask(src1.size(), src1.type(), cv::Scalar::all(0));
-    cv::gpu::bitwise_and(loadMat(src1), loadMat(src2), dst);
+    cv::gpu::bitwise_and(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_nomask);
+    cv::gpu::bitwise_and(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_mask, loadMat(mask, useRoi));
 
-    cv::Mat dst_gold = src1 & src2;
+    cv::Mat dst_gold_nomask, dst_gold_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::bitwise_and(src1, src2, dst_gold_nomask);
+    cv::bitwise_and(src1, src2, dst_gold_mask, mask);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_nomask, dst_nomask, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_mask, dst_mask, 0.0);
 }
 
 GPU_TEST_P(Bitwise_Array, Xor)
 {
-    cv::gpu::GpuMat dst;
+    cv::gpu::GpuMat dst_nomask, dst_mask(src1.size(), src1.type(), cv::Scalar::all(0));
-    cv::gpu::bitwise_xor(loadMat(src1), loadMat(src2), dst);
+    cv::gpu::bitwise_xor(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_nomask);
+    cv::gpu::bitwise_xor(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_mask, loadMat(mask, useRoi));
 
-    cv::Mat dst_gold = src1 ^ src2;
+    cv::Mat dst_gold_nomask, dst_gold_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::bitwise_xor(src1, src2, dst_gold_nomask);
+    cv::bitwise_xor(src1, src2, dst_gold_mask, mask);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_nomask, dst_nomask, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_mask, dst_mask, 0.0);
 }
 
 INSTANTIATE_TEST_CASE_P(GPU_Core, Bitwise_Array, testing::Combine(
     ALL_DEVICES,
     DIFFERENT_SIZES,
-    TYPES(CV_8U, CV_32S, 1, 4)));
+    TYPES(CV_8U, CV_32S, 1, 4),
+    WHOLE_SUBMAT));
 
 //////////////////////////////////////////////////////////////////////////////
 // Bitwise_Scalar

modules/objdetect/src/hog.cpp

@@ -858,6 +858,7 @@ void HOGDescriptor::detect(const Mat& img,
     Size winStride, Size padding, const vector<Point>& locations) const
 {
     hits.clear();
+    weights.clear();
     if( svmDetector.empty() )
         return;