Merge pull request #14676 from 103yiran:103yiran-clang

make it compatible with clang
Merge pull request #13669 from alalek:test_eigen_relax_eps_2.4
2026-07-22 03:43:03 +04:00 · 2019-05-30 14:19:14 +03:00 · 2019-05-30 14:04:40 +08:00 · 2019-01-21 18:32:22 +03:00 · 2019-01-21 15:33:37 +03:00 · 2019-01-21 15:18:10 +03:00
60 changed files with 1548 additions and 136 deletions
@@ -1,10 +1,10 @@
+# ignore dot files/directories
+.*
+!.gitignore
+
 *.autosave
 *.pyc
 *.user
 *~
-.*.swp
-.DS_Store
-.sw[a-z]
 tags
 tegra/
-.cache
@@ -25,7 +25,7 @@ endif(WIN32 AND NOT MINGW)

 ocv_warnings_disable(CMAKE_C_FLAGS -Wno-implicit-function-declaration -Wno-uninitialized -Wmissing-prototypes
                                   -Wno-unused-but-set-parameter -Wmissing-declarations -Wunused -Wshadow -Wsign-compare)
-ocv_warnings_disable(CMAKE_C_FLAGS -Wunused-parameter) # clang
+ocv_warnings_disable(CMAKE_C_FLAGS -Wunused-parameter -Wstrict-prototypes) # clang
 ocv_warnings_disable(CMAKE_C_FLAGS /wd4013 /wd4018 /wd4101 /wd4244 /wd4267 /wd4715) # vs2005

 if(UNIX)
@@ -3,7 +3,7 @@
 #
 # ----------------------------------------------------------------------------

-if(ARM AND ENABLE_NEON AND NOT AARCH64)
+if(ARM AND ENABLE_NEON AND NOT AARCH64 AND NOT IOS)
  project(${PNG_LIBRARY} ASM)
 else()
  project(${PNG_LIBRARY})
@@ -14,7 +14,7 @@ ocv_include_directories("${CMAKE_CURRENT_SOURCE_DIR}" ${ZLIB_INCLUDE_DIR})
 file(GLOB lib_srcs *.c)
 file(GLOB lib_hdrs *.h)

-if(ARM AND ENABLE_NEON AND NOT AARCH64)
+if(ARM AND ENABLE_NEON AND NOT AARCH64 AND NOT IOS)
  list(APPEND lib_srcs arm/filter_neon.S arm/arm_init.c)
  add_definitions(-DPNG_ARM_NEON_OPT=2)
 else()
@@ -683,6 +683,16 @@ endif()
 # ========================== build platform ==========================
 status("")
 status("  Platform:")
+if(NOT DEFINED OPENCV_TIMESTAMP
+    AND NOT CMAKE_VERSION VERSION_LESS 2.8.11
+    AND NOT BUILD_INFO_SKIP_TIMESTAMP
+)
+  string(TIMESTAMP OPENCV_TIMESTAMP "" UTC)
+  set(OPENCV_TIMESTAMP "${OPENCV_TIMESTAMP}" CACHE STRING "Timestamp of OpenCV build configuration" FORCE)
+endif()
+if(OPENCV_TIMESTAMP)
+  status("    Timestamp:"      ${OPENCV_TIMESTAMP})
+endif()
 status("    Host:"             ${CMAKE_HOST_SYSTEM_NAME} ${CMAKE_HOST_SYSTEM_VERSION} ${CMAKE_HOST_SYSTEM_PROCESSOR})
 if(CMAKE_CROSSCOMPILING)
  status("    Target:"         ${CMAKE_SYSTEM_NAME} ${CMAKE_SYSTEM_VERSION} ${CMAKE_SYSTEM_PROCESSOR})
@@ -164,7 +164,7 @@ if(CMAKE_COMPILER_IS_GNUCXX)
  # Other optimizations
  if(ENABLE_OMIT_FRAME_POINTER)
    add_extra_compiler_option(-fomit-frame-pointer)
-  else()
+  elseif(DEFINED ENABLE_OMIT_FRAME_POINTER)
    add_extra_compiler_option(-fno-omit-frame-pointer)
  endif()
  if(ENABLE_FAST_MATH)
@@ -61,7 +61,9 @@ if(NOT DEFINED OpenCV_CUDA)
  endif()
 endif()

-if(MSVC)
+if(DEFINED OpenCV_ARCH AND DEFINED OpenCV_RUNTIME)
+  # custom overrided values
+elseif(MSVC)
  if(CMAKE_CL_64)
    set(OpenCV_ARCH x64)
    set(OpenCV_TBB_ARCH intel64)
@@ -81,7 +83,7 @@ if(MSVC)
    set(OpenCV_RUNTIME vc12)
  elseif(MSVC_VERSION EQUAL 1900)
    set(OpenCV_RUNTIME vc14)
-  elseif(MSVC_VERSION EQUAL 1910 OR MSVC_VERSION EQUAL 1911)
+  elseif(MSVC_VERSION MATCHES "^191[0-9]$")
    set(OpenCV_RUNTIME vc15)
  endif()
 elseif(MINGW)
@@ -82,7 +82,9 @@ if(NOT DEFINED OpenCV_STATIC)
  endif()
 endif()

-if(MSVC)
+if(DEFINED OpenCV_ARCH AND DEFINED OpenCV_RUNTIME)
+  # custom overrided values
+elseif(MSVC)
  if(CMAKE_CL_64)
    set(OpenCV_ARCH x64)
  else()
@@ -100,7 +102,7 @@ if(MSVC)
    set(OpenCV_RUNTIME vc12)
  elseif(MSVC_VERSION EQUAL 1900)
    set(OpenCV_RUNTIME vc14)
-  elseif(MSVC_VERSION EQUAL 1910 OR MSVC_VERSION EQUAL 1911)
+  elseif(MSVC_VERSION MATCHES "^191[0-9]$")
    set(OpenCV_RUNTIME vc15)
  endif()
 elseif(MINGW)
@@ -751,7 +751,10 @@ endfunction()

 macro(ocv_get_libname var_name)
  get_filename_component(__libname "${ARGN}" NAME)
-  string(REGEX REPLACE "^lib(.+).(a|so)(.[.0-9]+)?$" "\\1" __libname "${__libname}")
+  # libopencv_core.so.3.3 -> opencv_core
+  string(REGEX REPLACE "^lib(.+)\\.(a|so)(\\.[.0-9]+)?$" "\\1" __libname "${__libname}")
+  # MacOSX: libopencv_core.3.3.1.dylib -> opencv_core
+  string(REGEX REPLACE "^lib(.+[^.0-9])\\.([.0-9]+\\.)?dylib$" "\\1" __libname "${__libname}")
  set(${var_name} "${__libname}")
 endmacro()

@@ -440,5 +440,5 @@ You may observe a runtime instance of this on the `YouTube here <https://www.you
 .. raw:: html

  <div align="center">
-  <iframe title=" Camera calibration With OpenCV - Chessboard or asymmetrical circle pattern." width="560" height="349" src="http://www.youtube.com/embed/ViPN810E0SU?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title=" Camera calibration With OpenCV - Chessboard or asymmetrical circle pattern." width="560" height="349" src="https://www.youtube.com/embed/ViPN810E0SU?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>
@@ -275,5 +275,5 @@ You may observe a runtime instance of this on the `YouTube here <https://www.you
 .. raw:: html

  <div align="center">
-  <iframe title="File Input and Output using XML and YAML files in OpenCV" width="560" height="349" src="http://www.youtube.com/embed/A4yqVnByMMM?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="File Input and Output using XML and YAML files in OpenCV" width="560" height="349" src="https://www.youtube.com/embed/A4yqVnByMMM?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>
@@ -179,5 +179,5 @@ Finally, you may watch a sample run of the program on the `video posted <https:/
 .. raw:: html

  <div align="center">
-  <iframe title="How to scan images in OpenCV?" width="560" height="349" src="http://www.youtube.com/embed/fB3AN5fjgwc?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="How to scan images in OpenCV?" width="560" height="349" src="https://www.youtube.com/embed/fB3AN5fjgwc?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>
@@ -125,5 +125,5 @@ You may observe a runtime instance of this on the `YouTube here <https://www.you
 .. raw:: html

  <div align="center">
-  <iframe title="Interoperability with OpenCV 1" width="560" height="349" src="http://www.youtube.com/embed/qckm-zvo31w?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="Interoperability with OpenCV 1" width="560" height="349" src="https://www.youtube.com/embed/qckm-zvo31w?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>
@@ -128,7 +128,7 @@ For example:

 You can download this source code from :download:`here <../../../../samples/cpp/tutorial_code/core/mat_mask_operations/mat_mask_operations.cpp>` or look in the OpenCV source code libraries sample directory at :file:`samples/cpp/tutorial_code/core/mat_mask_operations/mat_mask_operations.cpp`.

-Check out an instance of running the program on our `YouTube channel <http://www.youtube.com/watch?v=7PF1tAU9se4>`_ .
+Check out an instance of running the program on our `YouTube channel <https://www.youtube.com/watch?v=7PF1tAU9se4>`_ .

 .. raw:: html

@@ -307,5 +307,5 @@ You can also find a quick video demonstration of this on `YouTube <https://www.y
 .. raw:: html

  <div align="center">
-  <iframe title="Install OpenCV by using its source files - Part 1" width="560" height="349" src="http://www.youtube.com/embed/1tibU7vGWpk?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="Install OpenCV by using its source files - Part 1" width="560" height="349" src="https://www.youtube.com/embed/1tibU7vGWpk?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>
@@ -144,5 +144,5 @@ In both cases we managed a performance increase of almost 100% compared to the C
 .. raw:: html

  <div align="center">
-  <iframe title="Similarity check (PNSR and SSIM) on the GPU" width="560" height="349" src="http://www.youtube.com/embed/3_ESXmFlnvY?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="Similarity check (PNSR and SSIM) on the GPU" width="560" height="349" src="https://www.youtube.com/embed/3_ESXmFlnvY?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>
@@ -212,5 +212,5 @@ You may observe a runtime instance of this on the `YouTube here <https://www.you
 .. raw:: html

  <div align="center">
-  <iframe title="Video Input with OpenCV (Plus PSNR and MSSIM)" width="560" height="349" src="http://www.youtube.com/embed/iOcNljutOgg?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="Video Input with OpenCV (Plus PSNR and MSSIM)" width="560" height="349" src="https://www.youtube.com/embed/iOcNljutOgg?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>
@@ -132,5 +132,5 @@ You may observe a runtime instance of this on the `YouTube here <https://www.you
 .. raw:: html

  <div align="center">
-  <iframe title="Creating a video with OpenCV" width="560" height="349" src="http://www.youtube.com/embed/jpBwHxsl1_0?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="Creating a video with OpenCV" width="560" height="349" src="https://www.youtube.com/embed/jpBwHxsl1_0?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>
@@ -81,7 +81,7 @@ Gaussian Filter

      G_{0}(x, y) = A  e^{ \dfrac{ -(x - \mu_{x})^{2} }{ 2\sigma^{2}_{x} } +  \dfrac{ -(y - \mu_{y})^{2} }{ 2\sigma^{2}_{y} } }

-   where :math:`\mu` is the mean (the peak) and :math:`\sigma` represents the variance (per each of the variables :math:`x` and :math:`y`)
+   where :math:`\mu` is the mean (the peak) and :math:`\sigma` represents the standard deviation (per each of the variables :math:`x` and :math:`y`)


 Median Filter
@@ -129,5 +129,5 @@ Result
   .. raw:: html

     <div align="center">
-     <iframe title="Introduction - Display an Image" width="560" height="349" src="http://www.youtube.com/embed/1OJEqpuaGc4?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+     <iframe title="Introduction - Display an Image" width="560" height="349" src="https://www.youtube.com/embed/1OJEqpuaGc4?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
     </div>
@@ -358,7 +358,7 @@ Now here's our recommendation for the structure of the tutorial (although, remem
          .. raw:: html

             <div align="center">
-             <iframe title="Creating a video with OpenCV" width="560" height="349" src="http://www.youtube.com/embed/jpBwHxsl1_0?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+             <iframe title="Creating a video with OpenCV" width="560" height="349" src="https://www.youtube.com/embed/jpBwHxsl1_0?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
             </div>

       This results in the text and video: You may observe a runtime instance of this on the `YouTube here <https://www.youtube.com/watch?v=jpBwHxsl1_0>`_.
@@ -366,7 +366,7 @@ Now here's our recommendation for the structure of the tutorial (although, remem
       .. raw:: html

          <div align="center">
-          <iframe title="Creating a video with OpenCV" width="560" height="349" src="http://www.youtube.com/embed/jpBwHxsl1_0?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+          <iframe title="Creating a video with OpenCV" width="560" height="349" src="https://www.youtube.com/embed/jpBwHxsl1_0?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
          </div>

     When these aren't self-explanatory make sure to throw in a few guiding lines about what and why we can see.
@@ -40,8 +40,8 @@ You may find the content of this tutorial also inside the following videos: `Par
 .. raw:: html

  <div align="center">
-  <iframe title="Install OpenCV by using its source files - Part 1" width="560" height="349" src="http://www.youtube.com/embed/NnovZ1cTlMs?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
-  <iframe title="Install OpenCV by using its source files - Part 2" width="560" height="349" src="http://www.youtube.com/embed/qGNWMcfWwPU?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="Install OpenCV by using its source files - Part 1" width="560" height="349" src="https://www.youtube.com/embed/NnovZ1cTlMs?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="Install OpenCV by using its source files - Part 2" width="560" height="349" src="https://www.youtube.com/embed/qGNWMcfWwPU?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>

 .. warning:: These videos above are long-obsolete and contain inaccurate information. Be careful, since solutions described in those videos are no longer supported and may even break your install.
@@ -120,10 +120,10 @@ After the processing we need to convert it back to UIImage.
     :alt: header
     :align: center

-Check out an instance of running code with more Image Effects on `YouTube <http://www.youtube.com/watch?v=Ko3K_xdhJ1I>`_ .
+Check out an instance of running code with more Image Effects on `YouTube <https://www.youtube.com/watch?v=Ko3K_xdhJ1I>`_ .

 .. raw:: html

  <div align="center">
- <iframe width="560" height="350" src="http://www.youtube.com/embed/Ko3K_xdhJ1I" frameborder="0" allowfullscreen></iframe>
+ <iframe width="560" height="350" src="https://www.youtube.com/embed/Ko3K_xdhJ1I" frameborder="0" allowfullscreen></iframe>
  </div>
@@ -215,5 +215,5 @@ You may observe a runtime instance of this on the `YouTube here <https://www.you
 .. raw:: html

  <div align="center">
-  <iframe title="Support Vector Machines for Non-Linearly Separable Data" width="560" height="349" src="http://www.youtube.com/embed/vFv2yPcSo-Q?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
+  <iframe title="Support Vector Machines for Non-Linearly Separable Data" width="560" height="349" src="https://www.youtube.com/embed/vFv2yPcSo-Q?rel=0&loop=1" frameborder="0" allowfullscreen align="middle"></iframe>
  </div>
@@ -5,8 +5,10 @@
 #include <android/log.h>
 #include <cctype>
 #include <string>
+#include <cstring>
 #include <vector>
 #include <algorithm>
+#include <functional>
 #include <opencv2/core/version.hpp>
 #include "camera_activity.hpp"
 #include "camera_wrapper.h"
@@ -864,10 +864,10 @@ Rect getValidDisparityROI( Rect roi1, Rect roi2,
                          int SADWindowSize )
 {
    int SW2 = SADWindowSize/2;
-    int minD = minDisparity, maxD = minDisparity + numberOfDisparities - 1;
+    int maxD = minDisparity + numberOfDisparities - 1;

    int xmin = max(roi1.x, roi2.x + maxD) + SW2;
-    int xmax = min(roi1.x + roi1.width, roi2.x + roi2.width - minD) - SW2;
+    int xmax = min(roi1.x + roi1.width, roi2.x + roi2.width) - SW2;
    int ymin = max(roi1.y, roi2.y) + SW2;
    int ymax = min(roi1.y + roi1.height, roi2.y + roi2.height) - SW2;

@@ -60,7 +60,7 @@ protected:

 protected:
    std::string combine(const std::string& _item1, const std::string& _item2);
-    cv::Mat mergeRectification(const cv::Mat& l, const cv::Mat& r);
+    static void merge4(const cv::Mat& tl, const cv::Mat& tr, const cv::Mat& bl, const cv::Mat& br, cv::Mat& merged);
 };

 ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -385,11 +385,7 @@ TEST_F(fisheyeTest, EstimateUncertainties)
    CV_Assert(errors.alpha == 0);
 }

-#ifdef HAVE_TEGRA_OPTIMIZATION
-TEST_F(fisheyeTest, DISABLED_rectify)
-#else
-TEST_F(fisheyeTest, rectify)
-#endif
+TEST_F(fisheyeTest, stereoRectify)
 {
    const std::string folder =combine(datasets_repository_path, "calib-3_stereo_from_JY");

@@ -405,20 +401,65 @@ TEST_F(fisheyeTest, rectify)
    cv::fisheye::stereoRectify(K1, D1, K2, D2, calibration_size, theR, theT, R1, R2, P1, P2, Q,
                      cv::CALIB_ZERO_DISPARITY, requested_size, balance, fov_scale);

+    // Collected with these CMake flags: -DWITH_IPP=OFF -DCMAKE_BUILD_TYPE=Debug
+    cv::Matx33d R1_ref(
+        0.9992853269091279, 0.03779164101000276, -0.0007920188690205426,
+        -0.03778569762983931, 0.9992646472015868, 0.006511981857667881,
+        0.001037534936357442, -0.006477400933964018, 0.9999784831677112
+    );
+    cv::Matx33d R2_ref(
+        0.9994868963898833, -0.03197579751378937, -0.001868774538573449,
+        0.03196298186616116, 0.9994677442608699, -0.0065265589947392,
+        0.002076471801477729, 0.006463478587068991, 0.9999769555891836
+    );
+    cv::Matx34d P1_ref(
+        420.8551870450913, 0, 586.501617798451, 0,
+        0, 420.8551870450913, 374.7667511986098, 0,
+        0, 0, 1, 0
+    );
+    cv::Matx34d P2_ref(
+        420.8551870450913, 0, 586.501617798451, -41.77758076597302,
+        0, 420.8551870450913, 374.7667511986098, 0,
+        0, 0, 1, 0
+    );
+    cv::Matx44d Q_ref(
+        1, 0, 0, -586.501617798451,
+        0, 1, 0, -374.7667511986098,
+        0, 0, 0, 420.8551870450913,
+        0, 0, 10.07370889670733, -0
+    );
+
+    const double eps = 1e-10;
+    EXPECT_MAT_NEAR(R1_ref, R1, eps);
+    EXPECT_MAT_NEAR(R2_ref, R2, eps);
+    EXPECT_MAT_NEAR(P1_ref, P1, eps);
+    EXPECT_MAT_NEAR(P2_ref, P2, eps);
+    EXPECT_MAT_NEAR(Q_ref, Q, eps);
+
+    if (::testing::Test::HasFailure())
+    {
+        std::cout << "Actual values are:" << std::endl
+            << "R1 =" << std::endl << R1 << std::endl
+            << "R2 =" << std::endl << R2 << std::endl
+            << "P1 =" << std::endl << P1 << std::endl
+            << "P2 =" << std::endl << P2 << std::endl
+            << "Q =" << std::endl << Q << std::endl;
+    }
+
+#if 1 // Debug code
    cv::Mat lmapx, lmapy, rmapx, rmapy;
    //rewrite for fisheye
    cv::fisheye::initUndistortRectifyMap(K1, D1, R1, P1, requested_size, CV_32F, lmapx, lmapy);
    cv::fisheye::initUndistortRectifyMap(K2, D2, R2, P2, requested_size, CV_32F, rmapx, rmapy);

    cv::Mat l, r, lundist, rundist;
-    cv::VideoCapture lcap(combine(folder, "left/stereo_pair_%03d.jpg")),
-                     rcap(combine(folder, "right/stereo_pair_%03d.jpg"));
-
-    for(int i = 0;; ++i)
+    for (int i = 0; i < 34; ++i)
    {
-        lcap >> l; rcap >> r;
-        if (l.empty() || r.empty())
-            break;
+        SCOPED_TRACE(cv::format("image %d", i));
+        l = imread(combine(folder, cv::format("left/stereo_pair_%03d.jpg", i)), cv::IMREAD_COLOR);
+        r = imread(combine(folder, cv::format("right/stereo_pair_%03d.jpg", i)), cv::IMREAD_COLOR);
+        ASSERT_FALSE(l.empty());
+        ASSERT_FALSE(r.empty());

        int ndisp = 128;
        cv::rectangle(l, cv::Rect(255,       0, 829,       l.rows-1), CV_RGB(255, 0, 0));
@@ -427,15 +468,18 @@ TEST_F(fisheyeTest, rectify)
        cv::remap(l, lundist, lmapx, lmapy, cv::INTER_LINEAR);
        cv::remap(r, rundist, rmapx, rmapy, cv::INTER_LINEAR);

-        cv::Mat rectification = mergeRectification(lundist, rundist);
+        for (int ii = 0; ii < lundist.rows; ii += 20)
+        {
+            cv::line(lundist, cv::Point(0, ii), cv::Point(lundist.cols, ii), cv::Scalar(0, 255, 0));
+            cv::line(rundist, cv::Point(0, ii), cv::Point(lundist.cols, ii), cv::Scalar(0, 255, 0));
+        }

-        cv::Mat correct = cv::imread(combine(datasets_repository_path, cv::format("rectification_AB_%03d.png", i)));
+        cv::Mat rectification;
+        merge4(l, r, lundist, rundist, rectification);

-        if (correct.empty())
-            cv::imwrite(combine(datasets_repository_path, cv::format("rectification_AB_%03d.png", i)), rectification);
-         else
-             EXPECT_MAT_NEAR(correct, rectification, 1e-10);
-     }
+        cv::imwrite(cv::format("fisheye_rectification_AB_%03d.png", i), rectification);
+    }
+#endif
 }

 TEST_F(fisheyeTest, stereoCalibrate)
@@ -601,17 +645,17 @@ std::string fisheyeTest::combine(const std::string& _item1, const std::string& _
    return item1 + (last != '/' ? "/" : "") + item2;
 }

-cv::Mat fisheyeTest::mergeRectification(const cv::Mat& l, const cv::Mat& r)
+void fisheyeTest::merge4(const cv::Mat& tl, const cv::Mat& tr, const cv::Mat& bl, const cv::Mat& br, cv::Mat& merged)
 {
-    CV_Assert(l.type() == r.type() && l.size() == r.size());
-    cv::Mat merged(l.rows, l.cols * 2, l.type());
-    cv::Mat lpart = merged.colRange(0, l.cols);
-    cv::Mat rpart = merged.colRange(l.cols, merged.cols);
-    l.copyTo(lpart);
-    r.copyTo(rpart);
+    int type = tl.type();
+    cv::Size sz = tl.size();
+    ASSERT_EQ(type, tr.type()); ASSERT_EQ(type, bl.type()); ASSERT_EQ(type, br.type());
+    ASSERT_EQ(sz.width, tr.cols); ASSERT_EQ(sz.width, bl.cols); ASSERT_EQ(sz.width, br.cols);
+    ASSERT_EQ(sz.height, tr.rows); ASSERT_EQ(sz.height, bl.rows); ASSERT_EQ(sz.height, br.rows);

-    for(int i = 0; i < l.rows; i+=20)
-        cv::line(merged, cv::Point(0, i), cv::Point(merged.cols, i), CV_RGB(0, 255, 0));
-
-    return merged;
+    merged.create(cv::Size(sz.width * 2, sz.height * 2), type);
+    tl.copyTo(merged(cv::Rect(0, 0, sz.width, sz.height)));
+    tr.copyTo(merged(cv::Rect(sz.width, 0, sz.width, sz.height)));
+    bl.copyTo(merged(cv::Rect(0, sz.height, sz.width, sz.height)));
+    bl.copyTo(merged(cv::Rect(sz.width, sz.height, sz.width, sz.height)));
 }
@@ -326,9 +326,15 @@ string ERROR_PREFIXES[] = { "borderedAll",
                            "borderedTextureless",
                            "borderedDepthDiscont" }; // size of ERROR_KINDS_COUNT

+string ROI_PREFIXES[] = {   "roiX",
+                            "roiY",
+                            "roiWidth",
+                            "roiHeight" };
+

 const string RMS_STR = "RMS";
 const string BAD_PXLS_FRACTION_STR = "BadPxlsFraction";
+const string ROI_STR = "ValidDisparityROI";

 class QualityEvalParams
 {
@@ -366,16 +372,20 @@ public:
 protected:
    // assumed that left image is a reference image
    virtual int runStereoMatchingAlgorithm( const Mat& leftImg, const Mat& rightImg,
-                   Mat& leftDisp, Mat& rightDisp, int caseIdx ) = 0; // return ignored border width
+                   Rect& calcROI, Mat& leftDisp, Mat& rightDisp, int caseIdx ) = 0; // return ignored border width

    int readDatasetsParams( FileStorage& fs );
    virtual int readRunParams( FileStorage& fs );
    void writeErrors( const string& errName, const vector<float>& errors, FileStorage* fs = 0 );
+    void writeROI( const Rect& calcROI, FileStorage* fs = 0 );
    void readErrors( FileNode& fn, const string& errName, vector<float>& errors );
+    void readROI( FileNode& fn, Rect& trueROI );
    int compareErrors( const vector<float>& calcErrors, const vector<float>& validErrors,
                       const vector<float>& eps, const string& errName );
+    int compareROI( const Rect& calcROI, const Rect& validROI );
    int processStereoMatchingResults( FileStorage& fs, int caseIdx, bool isWrite,
                  const Mat& leftImg, const Mat& rightImg,
+                  const Rect& calcROI,
                  const Mat& trueLeftDisp, const Mat& trueRightDisp,
                  const Mat& leftDisp, const Mat& rightDisp,
                  const QualityEvalParams& qualityEvalParams  );
@@ -451,6 +461,7 @@ void CV_StereoMatchingTest::run(int)
        Mat rightImg = imread(datasetFullDirName + RIGHT_IMG_NAME);
        Mat trueLeftDisp = imread(datasetFullDirName + TRUE_LEFT_DISP_NAME, 0);
        Mat trueRightDisp = imread(datasetFullDirName + TRUE_RIGHT_DISP_NAME, 0);
+        Rect calcROI;

        if( leftImg.empty() || rightImg.empty() || trueLeftDisp.empty() )
        {
@@ -473,7 +484,7 @@ void CV_StereoMatchingTest::run(int)
        }

        Mat leftDisp, rightDisp;
-        int ignBorder = max(runStereoMatchingAlgorithm(leftImg, rightImg, leftDisp, rightDisp, ci), EVAL_IGNORE_BORDER);
+        int ignBorder = max(runStereoMatchingAlgorithm(leftImg, rightImg, calcROI, leftDisp, rightDisp, ci), EVAL_IGNORE_BORDER);

        leftDisp.convertTo( tmp, CV_32FC1 );
        leftDisp = tmp;
@@ -484,7 +495,7 @@ void CV_StereoMatchingTest::run(int)
        tmp.release();

        int tempCode = processStereoMatchingResults( resFS, ci, isWrite,
-                   leftImg, rightImg, trueLeftDisp, trueRightDisp, leftDisp, rightDisp, QualityEvalParams(ignBorder));
+                   leftImg, rightImg, calcROI, trueLeftDisp, trueRightDisp, leftDisp, rightDisp, QualityEvalParams(ignBorder));
        code = tempCode==cvtest::TS::OK ? code : tempCode;
    }

@@ -538,6 +549,7 @@ void calcErrors( const Mat& leftImg, const Mat& /*rightImg*/,

 int CV_StereoMatchingTest::processStereoMatchingResults( FileStorage& fs, int caseIdx, bool isWrite,
              const Mat& leftImg, const Mat& rightImg,
+              const Rect& calcROI,
              const Mat& trueLeftDisp, const Mat& trueRightDisp,
              const Mat& leftDisp, const Mat& rightDisp,
              const QualityEvalParams& qualityEvalParams )
@@ -574,6 +586,8 @@ int CV_StereoMatchingTest::processStereoMatchingResults( FileStorage& fs, int ca
        writeErrors( RMS_STR, rmss, &fs );
        cvWriteComment( fs.fs, BAD_PXLS_FRACTION_STR.c_str(), 0 );
        writeErrors( BAD_PXLS_FRACTION_STR, badPxlsFractions, &fs );
+        cvWriteComment( fs.fs, ROI_STR.c_str(), 0 );
+        writeROI( calcROI, &fs );
        fs << "}"; // datasetName
    }
    else // compare
@@ -583,16 +597,22 @@ int CV_StereoMatchingTest::processStereoMatchingResults( FileStorage& fs, int ca
        writeErrors( RMS_STR, rmss );
        ts->printf( cvtest::TS::LOG, "%s\n", BAD_PXLS_FRACTION_STR.c_str() );
        writeErrors( BAD_PXLS_FRACTION_STR, badPxlsFractions );
+        ts->printf( cvtest::TS::LOG, "%s\n", ROI_STR.c_str() );
+        writeROI( calcROI );

        FileNode fn = fs.getFirstTopLevelNode()[caseNames[caseIdx]];
        vector<float> validRmss, validBadPxlsFractions;
+        Rect validROI;

        readErrors( fn, RMS_STR, validRmss );
        readErrors( fn, BAD_PXLS_FRACTION_STR, validBadPxlsFractions );
+        readROI( fn, validROI );
        int tempCode = compareErrors( rmss, validRmss, rmsEps, RMS_STR );
        code = tempCode==cvtest::TS::OK ? code : tempCode;
        tempCode = compareErrors( badPxlsFractions, validBadPxlsFractions, fracEps, BAD_PXLS_FRACTION_STR );
        code = tempCode==cvtest::TS::OK ? code : tempCode;
+        tempCode = compareROI( calcROI, validROI );
+        code = tempCode==cvtest::TS::OK ? code : tempCode;
    }
    return code;
 }
@@ -642,6 +662,24 @@ void CV_StereoMatchingTest::writeErrors( const string& errName, const vector<flo
            ts->printf( cvtest::TS::LOG, "%s = %f\n", string(ERROR_PREFIXES[i]+errName).c_str(), *it );
 }

+void CV_StereoMatchingTest::writeROI( const Rect& calcROI, FileStorage* fs )
+{
+    if( fs )
+    {
+        *fs << ROI_PREFIXES[0] << calcROI.x;
+        *fs << ROI_PREFIXES[1] << calcROI.y;
+        *fs << ROI_PREFIXES[2] << calcROI.width;
+        *fs << ROI_PREFIXES[3] << calcROI.height;
+    }
+    else
+    {
+        ts->printf( cvtest::TS::LOG, "%s = %d\n", ROI_PREFIXES[0].c_str(), calcROI.x );
+        ts->printf( cvtest::TS::LOG, "%s = %d\n", ROI_PREFIXES[1].c_str(), calcROI.y );
+        ts->printf( cvtest::TS::LOG, "%s = %d\n", ROI_PREFIXES[2].c_str(), calcROI.width );
+        ts->printf( cvtest::TS::LOG, "%s = %d\n", ROI_PREFIXES[3].c_str(), calcROI.height );
+    }
+}
+
 void CV_StereoMatchingTest::readErrors( FileNode& fn, const string& errName, vector<float>& errors )
 {
    errors.resize( ERROR_KINDS_COUNT );
@@ -650,6 +688,14 @@ void CV_StereoMatchingTest::readErrors( FileNode& fn, const string& errName, vec
        fn[ERROR_PREFIXES[i]+errName] >> *it;
 }

+void CV_StereoMatchingTest::readROI( FileNode& fn, Rect& validROI )
+{
+    fn[ROI_PREFIXES[0]] >> validROI.x;
+    fn[ROI_PREFIXES[1]] >> validROI.y;
+    fn[ROI_PREFIXES[2]] >> validROI.width;
+    fn[ROI_PREFIXES[3]] >> validROI.height;
+}
+
 int CV_StereoMatchingTest::compareErrors( const vector<float>& calcErrors, const vector<float>& validErrors,
                   const vector<float>& eps, const string& errName )
 {
@@ -669,6 +715,26 @@ int CV_StereoMatchingTest::compareErrors( const vector<float>& calcErrors, const
    return ok ? cvtest::TS::OK : cvtest::TS::FAIL_BAD_ACCURACY;
 }

+int CV_StereoMatchingTest::compareROI( const Rect& calcROI, const Rect& validROI )
+{
+    int compare[4][2] = {
+        { calcROI.x, validROI.x },
+        { calcROI.y, validROI.y },
+        { calcROI.width, validROI.width },
+        { calcROI.height, validROI.height },
+    };
+    bool ok = true;
+    for (int i = 0; i < 4; i++)
+    {
+        if (compare[i][0] != compare[i][1])
+        {
+            ts->printf( cvtest::TS::LOG, "bad accuracy of %s (valid=%d; calc=%d)\n", ROI_PREFIXES[i].c_str(), compare[i][1], compare[i][0] );
+            ok = false;
+        }
+    }
+    return ok ? cvtest::TS::OK : cvtest::TS::FAIL_BAD_ACCURACY;
+}
+
 //----------------------------------- StereoBM test -----------------------------------------------------

 class CV_StereoBMTest : public CV_StereoMatchingTest
@@ -685,6 +751,7 @@ protected:
    struct RunParams
    {
        int ndisp;
+        int mindisp;
        int winSize;
    };
    vector<RunParams> caseRunParams;
@@ -694,12 +761,13 @@ protected:
        int code = CV_StereoMatchingTest::readRunParams( fs );
        FileNode fn = fs.getFirstTopLevelNode();
        assert(fn.isSeq());
-        for( int i = 0; i < (int)fn.size(); i+=4 )
+        for( int i = 0; i < (int)fn.size(); i+=5 )
        {
            string caseName = fn[i], datasetName = fn[i+1];
            RunParams params;
-            string ndisp = fn[i+2]; params.ndisp = atoi(ndisp.c_str());
-            string winSize = fn[i+3]; params.winSize = atoi(winSize.c_str());
+            String ndisp = fn[i+2]; params.ndisp = atoi(ndisp.c_str());
+            String mindisp = fn[i+3]; params.mindisp = atoi(mindisp.c_str());
+            String winSize = fn[i+4]; params.winSize = atoi(winSize.c_str());
            caseNames.push_back( caseName );
            caseDatasets.push_back( datasetName );
            caseRunParams.push_back( params );
@@ -708,7 +776,7 @@ protected:
    }

    virtual int runStereoMatchingAlgorithm( const Mat& _leftImg, const Mat& _rightImg,
-                   Mat& leftDisp, Mat& /*rightDisp*/, int caseIdx )
+                   Rect& calcROI, Mat& leftDisp, Mat& /*rightDisp*/, int caseIdx )
    {
        RunParams params = caseRunParams[caseIdx];
        assert( params.ndisp%16 == 0 );
@@ -717,7 +785,21 @@ protected:
        Mat rightImg; cvtColor( _rightImg, rightImg, CV_BGR2GRAY );

        StereoBM bm( StereoBM::BASIC_PRESET, params.ndisp, params.winSize );
+        bm.state->minDisparity = params.mindisp;
+
+        Rect cROI(0, 0, _leftImg.cols, _leftImg.rows);
+        calcROI = getValidDisparityROI(cROI, cROI, params.mindisp, params.ndisp, params.winSize);
+
        bm( leftImg, rightImg, leftDisp, CV_32F );
+
+        if (params.mindisp != 0)
+            for (int y = 0; y < leftDisp.rows; y++)
+                for (int x = 0; x < leftDisp.cols; x++)
+                {
+                    if (leftDisp.at<float>(y, x) < params.mindisp)
+                        leftDisp.at<float>(y, x) = -0.0625; // treat disparity < mindisp as no disparity
+                }
+
        return params.winSize/2;
    }
 };
@@ -754,7 +836,13 @@ protected:
            RunParams params;
            string ndisp = fn[i+2]; params.ndisp = atoi(ndisp.c_str());
            string winSize = fn[i+3]; params.winSize = atoi(winSize.c_str());
-            string fullDP = fn[i+4]; params.fullDP = atoi(fullDP.c_str()) == 0 ? false : true;
+            string fullDP = fn[i+4];
+            int n = atoi(fullDP.c_str());
+            params.fullDP = n == 0 ? false : true;
+
+            if (n > 1)
+                continue;   // OpenCV 3+ compatibility - skip tests for mode > 1
+
            caseNames.push_back( caseName );
            caseDatasets.push_back( datasetName );
            caseRunParams.push_back( params );
@@ -763,12 +851,16 @@ protected:
    }

    virtual int runStereoMatchingAlgorithm( const Mat& leftImg, const Mat& rightImg,
-                   Mat& leftDisp, Mat& /*rightDisp*/, int caseIdx )
+                   Rect& calcROI, Mat& leftDisp, Mat& /*rightDisp*/, int caseIdx )
    {
        RunParams params = caseRunParams[caseIdx];
        assert( params.ndisp%16 == 0 );
        StereoSGBM sgbm( 0, params.ndisp, params.winSize, 10*params.winSize*params.winSize, 40*params.winSize*params.winSize,
                         1, 63, 10, 100, 32, params.fullDP );
+
+        Rect cROI(0, 0, leftImg.cols, leftImg.rows);
+        calcROI = getValidDisparityROI(cROI, cROI, 0, params.ndisp, params.winSize);
+
        sgbm( leftImg, rightImg, leftDisp );
        assert( leftDisp.type() == CV_16SC1 );
        leftDisp/=16;
@@ -50,7 +50,7 @@
 #define CV_VERSION_EPOCH    2
 #define CV_VERSION_MAJOR    4
 #define CV_VERSION_MINOR    13
-#define CV_VERSION_REVISION 4
+#define CV_VERSION_REVISION 7

 #define CVAUX_STR_EXP(__A)  #__A
 #define CVAUX_STR(__A)      CVAUX_STR_EXP(__A)
@@ -2994,6 +2994,7 @@ PCA& PCA::computeVar(InputArray _data, InputArray __mean, int flags, double reta
    {
        CV_Assert( _mean.size() == mean_sz );
        _mean.convertTo(mean, ctype);
+        covar_flags |= CV_COVAR_USE_AVG;
    }

    calcCovarMatrix( data, covar, mean, covar_flags, ctype );
@@ -164,7 +164,7 @@ void Core_EigenTest_32::run(int) { check_full(CV_32FC1); }
 void Core_EigenTest_64::run(int) { check_full(CV_64FC1); }

 Core_EigenTest::Core_EigenTest()
-: eps_val_32(1e-3f), eps_vec_32(1e-2f),
+: eps_val_32(1e-3f), eps_vec_32(2e-2f),
  eps_val_64(1e-4f), eps_vec_64(1e-3f), ntests(100) {}
 Core_EigenTest::~Core_EigenTest() {}

@@ -272,9 +272,8 @@ void SimpleBlobDetector::findBlobs(const cv::Mat &image, const cv::Mat &binaryIm
 #endif
 }

-void SimpleBlobDetector::detectImpl(const cv::Mat& image, std::vector<cv::KeyPoint>& keypoints, const cv::Mat&) const
+void SimpleBlobDetector::detectImpl(const cv::Mat& image, std::vector<cv::KeyPoint>& keypoints, const cv::Mat& mask) const
 {
-    //TODO: support mask
    keypoints.clear();
    Mat grayscaleImage;
    if (image.channels() == 3)
@@ -355,6 +354,11 @@ void SimpleBlobDetector::detectImpl(const cv::Mat& image, std::vector<cv::KeyPoi
        keypoints.push_back(kpt);
    }

+    if (!mask.empty())
+    {
+        KeyPointsFilter::runByPixelsMask(keypoints, mask);
+    }
+
 #ifdef DEBUG_BLOB_DETECTOR
    namedWindow("keypoints", CV_WINDOW_NORMAL);
    Mat outImg = image.clone();
@@ -132,7 +132,12 @@ public:
        /* Construct the randomized trees. */
        for (int i = 0; i < trees_; i++) {
            /* Randomize the order of vectors to allow for unbiased sampling. */
+#ifndef OPENCV_FLANN_USE_STD_RAND
+            cv::randShuffle(vind_);
+#else
            std::random_shuffle(vind_.begin(), vind_.end());
+#endif
+
            tree_roots_[i] = divideTree(&vind_[0], int(size_) );
        }
    }
@@ -136,7 +136,12 @@ public:
              indices.resize( feature_size_ * CHAR_BIT );
              for (size_t j = 0; j < feature_size_ * CHAR_BIT; ++j)
                  indices[j] = j;
+
+#ifndef OPENCV_FLANN_USE_STD_RAND
+              cv::randShuffle(indices);
+#else
              std::random_shuffle(indices.begin(), indices.end());
+#endif
            }

            lsh::LshTable<ElementType>& table = tables_[i];
@@ -40,13 +40,31 @@
 namespace cvflann
 {

+inline int rand()
+{
+#ifndef OPENCV_FLANN_USE_STD_RAND
+#   if INT_MAX == RAND_MAX
+    int v = cv::theRNG().next() & INT_MAX;
+#   else
+    int v = cv::theRNG().uniform(0, RAND_MAX + 1);
+#   endif
+#else
+    int v = std::rand();
+#endif // OPENCV_FLANN_USE_STD_RAND
+    return v;
+}
+
 /**
 * Seeds the random number generator
 *  @param seed Random seed
 */
 inline void seed_random(unsigned int seed)
 {
-    srand(seed);
+#ifndef OPENCV_FLANN_USE_STD_RAND
+    cv::theRNG() = cv::RNG(seed);
+#else
+    std::srand(seed);
+#endif
 }

 /*
@@ -60,7 +78,7 @@ inline void seed_random(unsigned int seed)
 */
 inline double rand_double(double high = 1.0, double low = 0)
 {
-    return low + ((high-low) * (std::rand() / (RAND_MAX + 1.0)));
+    return low + ((high-low) * (rand() / (RAND_MAX + 1.0)));
 }

 /**
@@ -71,7 +89,7 @@ inline double rand_double(double high = 1.0, double low = 0)
 */
 inline int rand_int(int high = RAND_MAX, int low = 0)
 {
-    return low + (int) ( double(high-low) * (std::rand() / (RAND_MAX + 1.0)));
+    return low + (int) ( double(high-low) * (rand() / (RAND_MAX + 1.0)));
 }

 /**
@@ -107,7 +125,11 @@ public:
        for (int i = 0; i < size_; ++i) vals_[i] = i;

        // shuffle the elements in the array
+#ifndef OPENCV_FLANN_USE_STD_RAND
+        cv::randShuffle(vals_);
+#else
        std::random_shuffle(vals_.begin(), vals_.end());
+#endif

        counter_ = 0;
    }
@@ -1824,6 +1824,9 @@ public:
    void sparse(const GpuMat& prevImg, const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts,
        GpuMat& status, GpuMat* err = 0);

+    void sparse_multi(const GpuMat& prevImg, const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts,
+        GpuMat& status, Stream& stream, GpuMat* err = 0);
+
    void dense(const GpuMat& prevImg, const GpuMat& nextImg, GpuMat& u, GpuMat& v, GpuMat* err = 0);

    void releaseMemory();
@@ -303,6 +303,88 @@ PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, Video_PyrLKOpticalFlowSparse
    }
 }

+//////////////////////////////////////////////////////
+// PyrLKOpticalFlowSparseMulti
+
+#if defined(HAVE_TBB) && defined(HAVE_CUDA)
+
+DEF_PARAM_TEST(ImagePair_Gray_NPts_WinSz_Levels_Iters, pair_string, bool, int, int, int, int);
+
+PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, Video_PyrLKOpticalFlowSparseMulti,
+            Combine(Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")),
+                    Bool(),
+                    Values(8000),
+                    Values(21),
+                    Values(1, 3),
+                    Values(1, 30)))
+{
+    declare.time(20.0);
+
+    const pair_string imagePair = GET_PARAM(0);
+    const bool useGray = GET_PARAM(1);
+    const int points = GET_PARAM(2);
+    const int winSize = GET_PARAM(3);
+    const int levels = GET_PARAM(4);
+    const int iters = GET_PARAM(5);
+
+    const cv::Mat frame0 = readImage(imagePair.first, useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+    ASSERT_FALSE(frame0.empty());
+
+    const cv::Mat frame1 = readImage(imagePair.second, useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+    ASSERT_FALSE(frame1.empty());
+
+    cv::Mat gray_frame;
+    if (useGray)
+        gray_frame = frame0;
+    else
+        cv::cvtColor(frame0, gray_frame, cv::COLOR_BGR2GRAY);
+
+    cv::Mat pts;
+    cv::goodFeaturesToTrack(gray_frame, pts, points, 0.01, 0.0);
+
+    if (PERF_RUN_GPU())
+    {
+        const cv::gpu::GpuMat d_pts(pts.reshape(2, 1));
+
+        cv::gpu::PyrLKOpticalFlow d_pyrLK;
+        d_pyrLK.winSize = cv::Size(winSize, winSize);
+        d_pyrLK.maxLevel = levels - 1;
+        d_pyrLK.iters = iters;
+
+        const cv::gpu::GpuMat d_frame0(frame0);
+        const cv::gpu::GpuMat d_frame1(frame1);
+        cv::gpu::GpuMat nextPts;
+        cv::gpu::GpuMat status;
+
+        cv::gpu::Stream stream;
+        TEST_CYCLE()
+        {
+            d_pyrLK.sparse_multi(d_frame0, d_frame1, d_pts, nextPts, status, stream);
+            stream.waitForCompletion();
+        }
+
+        GPU_SANITY_CHECK(nextPts);
+        GPU_SANITY_CHECK(status);
+    }
+    else
+    {
+        cv::Mat nextPts;
+        cv::Mat status;
+
+        TEST_CYCLE()
+        {
+            cv::calcOpticalFlowPyrLK(frame0, frame1, pts, nextPts, status, cv::noArray(),
+                                     cv::Size(winSize, winSize), levels - 1,
+                                     cv::TermCriteria(cv::TermCriteria::COUNT + cv::TermCriteria::EPS, iters, 0.01));
+        }
+
+        CPU_SANITY_CHECK(nextPts);
+        CPU_SANITY_CHECK(status);
+    }
+}
+
+#endif // HAVE_TBB
+
 //////////////////////////////////////////////////////
 // PyrLKOpticalFlowDense

@@ -49,6 +49,10 @@
 #include "opencv2/gpu/device/vec_math.hpp"
 #include "opencv2/gpu/device/reduce.hpp"

+#include "opencv2/core/core.hpp"
+
+#include "cvconfig.h"
+
 using namespace cv::gpu;
 using namespace cv::gpu::device;

@@ -60,12 +64,54 @@ namespace pyrlk
    __constant__ int c_halfWin_y;
    __constant__ int c_iters;

-    texture<float, cudaTextureType2D, cudaReadModeElementType> tex_If(false, cudaFilterModeLinear, cudaAddressModeClamp);
-    texture<float4, cudaTextureType2D, cudaReadModeElementType> tex_If4(false, cudaFilterModeLinear, cudaAddressModeClamp);
-    texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex_Ib(false, cudaFilterModePoint, cudaAddressModeClamp);
+#define CUDA_CONSTANTS(index) \
+    __constant__ int c_winSize_x##index; \
+    __constant__ int c_winSize_y##index; \
+    __constant__ int c_halfWin_x##index; \
+    __constant__ int c_halfWin_y##index; \
+    __constant__ int c_iters##index;

-    texture<float, cudaTextureType2D, cudaReadModeElementType> tex_Jf(false, cudaFilterModeLinear, cudaAddressModeClamp);
-    texture<float4, cudaTextureType2D, cudaReadModeElementType> tex_Jf4(false, cudaFilterModeLinear, cudaAddressModeClamp);
+    CUDA_CONSTANTS(0)
+    CUDA_CONSTANTS(1)
+    CUDA_CONSTANTS(2)
+    CUDA_CONSTANTS(3)
+    CUDA_CONSTANTS(4)
+
+    template <int index> struct c_multi_winSize_x;
+    template <int index> struct c_multi_winSize_y;
+    template <int index> struct c_multi_halfWin_x;
+    template <int index> struct c_multi_halfWin_y;
+    template <int index> struct c_multi_iters;
+
+#define CUDA_CONSTANTS_ACCESSOR(index) \
+    template <> struct c_multi_winSize_x<index> \
+    { static __device__ __forceinline__ int get(void){ return c_winSize_x##index;} }; \
+    template <> struct c_multi_winSize_y<index> \
+    { static __device__ __forceinline__ int get(void){ return c_winSize_y##index;} }; \
+    template <> struct c_multi_halfWin_x<index> \
+    { static __device__ __forceinline__ int get(void){ return c_halfWin_x##index;} }; \
+    template <> struct c_multi_halfWin_y<index> \
+    { static __device__ __forceinline__ int get(void){ return c_halfWin_y##index;} }; \
+    template <> struct c_multi_iters<index> \
+    { static __device__ __forceinline__ int get(void){ return c_iters##index;} };
+
+    CUDA_CONSTANTS_ACCESSOR(0)
+    CUDA_CONSTANTS_ACCESSOR(1)
+    CUDA_CONSTANTS_ACCESSOR(2)
+    CUDA_CONSTANTS_ACCESSOR(3)
+    CUDA_CONSTANTS_ACCESSOR(4)
+
+    texture<float, cudaTextureType2D, cudaReadModeElementType>
+            tex_If(false, cudaFilterModeLinear, cudaAddressModeClamp);
+    texture<float4, cudaTextureType2D, cudaReadModeElementType>
+            tex_If4(false, cudaFilterModeLinear, cudaAddressModeClamp);
+    texture<uchar, cudaTextureType2D, cudaReadModeElementType>
+            tex_Ib(false, cudaFilterModePoint, cudaAddressModeClamp);
+
+    texture<float, cudaTextureType2D, cudaReadModeElementType>
+            tex_Jf(false, cudaFilterModeLinear, cudaAddressModeClamp);
+    texture<float4, cudaTextureType2D, cudaReadModeElementType>
+            tex_Jf4(false, cudaFilterModeLinear, cudaAddressModeClamp);

    template <int cn> struct Tex_I;
    template <> struct Tex_I<1>
@@ -99,6 +145,57 @@ namespace pyrlk
        }
    };

+    //--------------------------------------------------------------------------
+
+#define CUDA_DECL_TEX_MULTI(texname, type, filtermode) \
+    texture<type, cudaTextureType2D, cudaReadModeElementType> \
+            texname##_multi0(false, filtermode, cudaAddressModeClamp); \
+    texture<type, cudaTextureType2D, cudaReadModeElementType> \
+            texname##_multi1(false, filtermode, cudaAddressModeClamp); \
+    texture<type, cudaTextureType2D, cudaReadModeElementType> \
+            texname##_multi2(false, filtermode, cudaAddressModeClamp); \
+    texture<type, cudaTextureType2D, cudaReadModeElementType> \
+            texname##_multi3(false, filtermode, cudaAddressModeClamp); \
+    texture<type, cudaTextureType2D, cudaReadModeElementType> \
+            texname##_multi4(false, filtermode, cudaAddressModeClamp); \
+
+    CUDA_DECL_TEX_MULTI(tex_If1, float, cudaFilterModeLinear)
+    CUDA_DECL_TEX_MULTI(tex_If4, float4, cudaFilterModeLinear)
+    CUDA_DECL_TEX_MULTI(tex_Ib1, uchar, cudaFilterModePoint)
+    CUDA_DECL_TEX_MULTI(tex_Jf1, float, cudaFilterModeLinear)
+    CUDA_DECL_TEX_MULTI(tex_Jf4, float4, cudaFilterModeLinear)
+
+    template <int cn, int index> struct Tex_I_multi;
+    template <int cn, int index> struct Tex_J_multi;
+    template <int cn, int index> struct Tex_B_multi;
+
+#define CUDA_DECL_TEX_MULTI_ACCESS(accessorname, texname, cn, returntype) \
+    template <> struct accessorname##_multi<cn, 0> \
+    { static __device__ __forceinline__ returntype read(float x, float y) \
+        { return tex2D(texname##cn##_multi0, x, y); } }; \
+    template <> struct accessorname##_multi<cn, 1> \
+    { static __device__ __forceinline__ returntype read(float x, float y) \
+        { return tex2D(texname##cn##_multi1, x, y); } }; \
+    template <> struct accessorname##_multi<cn, 2> \
+    { static __device__ __forceinline__ returntype read(float x, float y) \
+        { return tex2D(texname##cn##_multi2, x, y); } }; \
+    template <> struct accessorname##_multi<cn, 3> \
+    { static __device__ __forceinline__ returntype read(float x, float y) \
+        { return tex2D(texname##cn##_multi3, x, y); } }; \
+    template <> struct accessorname##_multi<cn, 4> \
+    { static __device__ __forceinline__ returntype read(float x, float y) \
+        { return tex2D(texname##cn##_multi4, x, y); } };
+
+    CUDA_DECL_TEX_MULTI_ACCESS(Tex_I, tex_If, 1, float)
+    CUDA_DECL_TEX_MULTI_ACCESS(Tex_I, tex_If, 4, float4)
+
+    CUDA_DECL_TEX_MULTI_ACCESS(Tex_B, tex_Ib, 1, uchar)
+
+    CUDA_DECL_TEX_MULTI_ACCESS(Tex_J, tex_Jf, 1, float)
+    CUDA_DECL_TEX_MULTI_ACCESS(Tex_J, tex_Jf, 4, float4)
+
+    //--------------------------------------------------------------------------
+
    __device__ __forceinline__ void accum(float& dst, float val)
    {
        dst += val;
@@ -309,6 +406,200 @@ namespace pyrlk
        }
    }

+#if defined(HAVE_TBB)
+    template <int cn, int index, int PATCH_X, int PATCH_Y, bool calcErr>
+    __global__ void sparseKernel_multi(const float2* prevPts, float2* nextPts, uchar* status, float* err, const int level, const int rows, const int cols)
+    {
+    #if __CUDA_ARCH__ <= 110
+        const int BLOCK_SIZE = 128;
+    #else
+        const int BLOCK_SIZE = 256;
+    #endif
+
+        __shared__ float smem1[BLOCK_SIZE];
+        __shared__ float smem2[BLOCK_SIZE];
+        __shared__ float smem3[BLOCK_SIZE];
+
+        const unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;
+
+        float2 prevPt = prevPts[blockIdx.x];
+        prevPt.x *= (1.0f / (1 << level));
+        prevPt.y *= (1.0f / (1 << level));
+
+        if (prevPt.x < 0 || prevPt.x >= cols || prevPt.y < 0 || prevPt.y >= rows)
+        {
+            if (tid == 0 && level == 0)
+                status[blockIdx.x] = 0;
+
+            return;
+        }
+
+        prevPt.x -= c_multi_halfWin_x<index>::get();
+        prevPt.y -= c_multi_halfWin_y<index>::get();
+
+        // extract the patch from the first image, compute covariation matrix of derivatives
+
+        float A11 = 0;
+        float A12 = 0;
+        float A22 = 0;
+
+        typedef typename TypeVec<float, cn>::vec_type work_type;
+
+        work_type I_patch   [PATCH_Y][PATCH_X];
+        work_type dIdx_patch[PATCH_Y][PATCH_X];
+        work_type dIdy_patch[PATCH_Y][PATCH_X];
+
+        for (int yBase = threadIdx.y, i = 0; yBase < c_multi_winSize_y<index>::get(); yBase += blockDim.y, ++i)
+        {
+            for (int xBase = threadIdx.x, j = 0; xBase < c_multi_winSize_x<index>::get(); xBase += blockDim.x, ++j)
+            {
+                float x = prevPt.x + xBase + 0.5f;
+                float y = prevPt.y + yBase + 0.5f;
+
+                I_patch[i][j] = Tex_I_multi<cn,index>::read(x, y);
+
+                // Sharr Deriv
+
+                work_type dIdx = 3.0f * Tex_I_multi<cn,index>::read(x+1, y-1) + 10.0f * Tex_I_multi<cn,index>::read(x+1, y) + 3.0f * Tex_I_multi<cn,index>::read(x+1, y+1) -
+                                 (3.0f * Tex_I_multi<cn,index>::read(x-1, y-1) + 10.0f * Tex_I_multi<cn,index>::read(x-1, y) + 3.0f * Tex_I_multi<cn,index>::read(x-1, y+1));
+
+                work_type dIdy = 3.0f * Tex_I_multi<cn,index>::read(x-1, y+1) + 10.0f * Tex_I_multi<cn,index>::read(x, y+1) + 3.0f * Tex_I_multi<cn,index>::read(x+1, y+1) -
+                                (3.0f * Tex_I_multi<cn,index>::read(x-1, y-1) + 10.0f * Tex_I_multi<cn,index>::read(x, y-1) + 3.0f * Tex_I_multi<cn,index>::read(x+1, y-1));
+
+                dIdx_patch[i][j] = dIdx;
+                dIdy_patch[i][j] = dIdy;
+
+                accum(A11, dIdx * dIdx);
+                accum(A12, dIdx * dIdy);
+                accum(A22, dIdy * dIdy);
+            }
+        }
+
+        reduce<BLOCK_SIZE>(smem_tuple(smem1, smem2, smem3), thrust::tie(A11, A12, A22), tid, thrust::make_tuple(plus<float>(), plus<float>(), plus<float>()));
+
+    #if __CUDA_ARCH__ >= 300
+        if (tid == 0)
+        {
+            smem1[0] = A11;
+            smem2[0] = A12;
+            smem3[0] = A22;
+        }
+    #endif
+
+        __syncthreads();
+
+        A11 = smem1[0];
+        A12 = smem2[0];
+        A22 = smem3[0];
+
+        float D = A11 * A22 - A12 * A12;
+
+        if (abs_(D) < numeric_limits<float>::epsilon())
+        {
+            if (tid == 0 && level == 0)
+                status[blockIdx.x] = 0;
+
+            return;
+        }
+
+        D = 1.f / D;
+
+        A11 *= D;
+        A12 *= D;
+        A22 *= D;
+
+        float2 nextPt = nextPts[blockIdx.x];
+        nextPt.x *= 2.f;
+        nextPt.y *= 2.f;
+
+        nextPt.x -= c_multi_halfWin_x<index>::get();
+        nextPt.y -= c_multi_halfWin_y<index>::get();
+
+        for (int k = 0; k < c_multi_iters<index>::get(); ++k)
+        {
+            if (nextPt.x < -c_multi_halfWin_x<index>::get() || nextPt.x >= cols || nextPt.y < -c_multi_halfWin_y<index>::get() || nextPt.y >= rows)
+            {
+                if (tid == 0 && level == 0)
+                    status[blockIdx.x] = 0;
+
+                return;
+            }
+
+            float b1 = 0;
+            float b2 = 0;
+
+            for (int y = threadIdx.y, i = 0; y < c_multi_winSize_y<index>::get(); y += blockDim.y, ++i)
+            {
+                for (int x = threadIdx.x, j = 0; x < c_multi_winSize_x<index>::get(); x += blockDim.x, ++j)
+                {
+                    work_type I_val = I_patch[i][j];
+                    work_type J_val = Tex_J_multi<cn,index>::read(nextPt.x + x + 0.5f, nextPt.y + y + 0.5f);
+
+                    work_type diff = (J_val - I_val) * 32.0f;
+
+                    accum(b1, diff * dIdx_patch[i][j]);
+                    accum(b2, diff * dIdy_patch[i][j]);
+                }
+            }
+
+            reduce<BLOCK_SIZE>(smem_tuple(smem1, smem2), thrust::tie(b1, b2), tid, thrust::make_tuple(plus<float>(), plus<float>()));
+
+        #if __CUDA_ARCH__ >= 300
+            if (tid == 0)
+            {
+                smem1[0] = b1;
+                smem2[0] = b2;
+            }
+        #endif
+
+            __syncthreads();
+
+            b1 = smem1[0];
+            b2 = smem2[0];
+
+            float2 delta;
+            delta.x = A12 * b2 - A22 * b1;
+            delta.y = A12 * b1 - A11 * b2;
+
+            nextPt.x += delta.x;
+            nextPt.y += delta.y;
+
+            if (::fabs(delta.x) < 0.01f && ::fabs(delta.y) < 0.01f)
+                break;
+        }
+
+        float errval = 0;
+        if (calcErr)
+        {
+            for (int y = threadIdx.y, i = 0; y < c_multi_winSize_y<index>::get(); y += blockDim.y, ++i)
+            {
+                for (int x = threadIdx.x, j = 0; x < c_multi_winSize_x<index>::get(); x += blockDim.x, ++j)
+                {
+                    work_type I_val = I_patch[i][j];
+                    work_type J_val = Tex_J_multi<cn,index>::read(nextPt.x + x + 0.5f, nextPt.y + y + 0.5f);
+
+                    work_type diff = J_val - I_val;
+
+                    accum(errval, abs_(diff));
+                }
+            }
+
+            reduce<BLOCK_SIZE>(smem1, errval, tid, plus<float>());
+        }
+
+        if (tid == 0)
+        {
+            nextPt.x += c_multi_halfWin_x<index>::get();
+            nextPt.y += c_multi_halfWin_y<index>::get();
+
+            nextPts[blockIdx.x] = nextPt;
+
+            if (calcErr)
+                err[blockIdx.x] = static_cast<float>(errval) / (cn * c_multi_winSize_x<index>::get() * c_multi_winSize_y<index>::get());
+        }
+    }
+#endif // defined(HAVE_TBB)
+
    template <int cn, int PATCH_X, int PATCH_Y>
    void sparse_caller(int rows, int cols, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
                       int level, dim3 block, cudaStream_t stream)
@@ -326,6 +617,26 @@ namespace pyrlk
            cudaSafeCall( cudaDeviceSynchronize() );
    }

+#if defined(HAVE_TBB)
+    template <int cn, int index, int PATCH_X, int PATCH_Y>
+    void sparse_caller_multi(int rows, int cols, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
+                       int level, dim3 block, cudaStream_t stream)
+    {
+        dim3 grid(ptcount);
+
+        if (level == 0 && err)
+            sparseKernel_multi<cn, index, PATCH_X, PATCH_Y, true><<<grid, block>>>(prevPts, nextPts, status, err, level, rows, cols);
+        else
+            sparseKernel_multi<cn, index, PATCH_X, PATCH_Y, false><<<grid, block>>>(prevPts, nextPts, status, err, level, rows, cols);
+
+        cudaSafeCall( cudaGetLastError() );
+
+        if (stream == 0)
+            cudaSafeCall( cudaDeviceSynchronize() );
+    }
+
+#endif // defined(HAVE_TBB)
+
    template <bool calcErr>
    __global__ void denseKernel(PtrStepf u, PtrStepf v, const PtrStepf prevU, const PtrStepf prevV, PtrStepf err, const int rows, const int cols)
    {
@@ -484,6 +795,30 @@ namespace pyrlk
        cudaSafeCall( cudaMemcpyToSymbol(c_iters, &iters, sizeof(int)) );
    }

+#if defined(HAVE_TBB)
+    void loadConstants_multi(int2 winSize, int iters, int index, cudaStream_t stream = 0)
+    {
+        int2 halfWin;
+#define COPY_TO_SYMBOL_CALL(index) \
+        cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_x##index, &winSize.x, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); \
+        cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_y##index, &winSize.y, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); \
+        halfWin = make_int2((winSize.x - 1) / 2, (winSize.y - 1) / 2); \
+        cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_x##index, &halfWin.x, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); \
+        cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_y##index, &halfWin.y, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); \
+        cudaSafeCall( cudaMemcpyToSymbolAsync(c_iters##index, &iters, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
+
+        switch(index)
+        {
+            case 0: COPY_TO_SYMBOL_CALL(0) break;
+            case 1: COPY_TO_SYMBOL_CALL(1) break;
+            case 2: COPY_TO_SYMBOL_CALL(2) break;
+            case 3: COPY_TO_SYMBOL_CALL(3) break;
+            case 4: COPY_TO_SYMBOL_CALL(4) break;
+            default: CV_Error(CV_StsBadArg, "invalid execution line index"); break;
+        }
+    }
+#endif // defined(HAVE_TBB)
+
    void sparse1(PtrStepSzf I, PtrStepSzf J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
                 int level, dim3 block, dim3 patch, cudaStream_t stream)
    {
@@ -528,6 +863,161 @@ namespace pyrlk
            level, block, stream);
    }

+#if defined(HAVE_TBB)
+    void sparse1_multi(PtrStepSzf I, PtrStepSzf J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
+                 int level, dim3 block, dim3 patch, cudaStream_t stream, int index)
+    {
+        typedef void (*func_t)(int rows, int cols, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
+                               int level, dim3 block, cudaStream_t stream);
+
+        static const func_t funcs[5][5][5] =
+        {
+            { // index 0
+                {sparse_caller_multi<1, 0, 1, 1>, sparse_caller_multi<1, 0, 2, 1>, sparse_caller_multi<1, 0, 3, 1>, sparse_caller_multi<1, 0, 4, 1>, sparse_caller_multi<1, 0, 5, 1>},
+                {sparse_caller_multi<1, 0, 1, 2>, sparse_caller_multi<1, 0, 2, 2>, sparse_caller_multi<1, 0, 3, 2>, sparse_caller_multi<1, 0, 4, 2>, sparse_caller_multi<1, 0, 5, 2>},
+                {sparse_caller_multi<1, 0, 1, 3>, sparse_caller_multi<1, 0, 2, 3>, sparse_caller_multi<1, 0, 3, 3>, sparse_caller_multi<1, 0, 4, 3>, sparse_caller_multi<1, 0, 5, 3>},
+                {sparse_caller_multi<1, 0, 1, 4>, sparse_caller_multi<1, 0, 2, 4>, sparse_caller_multi<1, 0, 3, 4>, sparse_caller_multi<1, 0, 4, 4>, sparse_caller_multi<1, 0, 5, 4>},
+                {sparse_caller_multi<1, 0, 1, 5>, sparse_caller_multi<1, 0, 2, 5>, sparse_caller_multi<1, 0, 3, 5>, sparse_caller_multi<1, 0, 4, 5>, sparse_caller_multi<1, 0, 5, 5>}
+            },
+            { // index 1
+                {sparse_caller_multi<1, 1, 1, 1>, sparse_caller_multi<1, 1, 2, 1>, sparse_caller_multi<1, 1, 3, 1>, sparse_caller_multi<1, 1, 4, 1>, sparse_caller_multi<1, 1, 5, 1>},
+                {sparse_caller_multi<1, 1, 1, 2>, sparse_caller_multi<1, 1, 2, 2>, sparse_caller_multi<1, 1, 3, 2>, sparse_caller_multi<1, 1, 4, 2>, sparse_caller_multi<1, 1, 5, 2>},
+                {sparse_caller_multi<1, 1, 1, 3>, sparse_caller_multi<1, 1, 2, 3>, sparse_caller_multi<1, 1, 3, 3>, sparse_caller_multi<1, 1, 4, 3>, sparse_caller_multi<1, 1, 5, 3>},
+                {sparse_caller_multi<1, 1, 1, 4>, sparse_caller_multi<1, 1, 2, 4>, sparse_caller_multi<1, 1, 3, 4>, sparse_caller_multi<1, 1, 4, 4>, sparse_caller_multi<1, 1, 5, 4>},
+                {sparse_caller_multi<1, 1, 1, 5>, sparse_caller_multi<1, 1, 2, 5>, sparse_caller_multi<1, 1, 3, 5>, sparse_caller_multi<1, 1, 4, 5>, sparse_caller_multi<1, 1, 5, 5>}
+            },
+            { // index 2
+                {sparse_caller_multi<1, 2, 1, 1>, sparse_caller_multi<1, 2, 2, 1>, sparse_caller_multi<1, 2, 3, 1>, sparse_caller_multi<1, 2, 4, 1>, sparse_caller_multi<1, 2, 5, 1>},
+                {sparse_caller_multi<1, 2, 1, 2>, sparse_caller_multi<1, 2, 2, 2>, sparse_caller_multi<1, 2, 3, 2>, sparse_caller_multi<1, 2, 4, 2>, sparse_caller_multi<1, 2, 5, 2>},
+                {sparse_caller_multi<1, 2, 1, 3>, sparse_caller_multi<1, 2, 2, 3>, sparse_caller_multi<1, 2, 3, 3>, sparse_caller_multi<1, 2, 4, 3>, sparse_caller_multi<1, 2, 5, 3>},
+                {sparse_caller_multi<1, 2, 1, 4>, sparse_caller_multi<1, 2, 2, 4>, sparse_caller_multi<1, 2, 3, 4>, sparse_caller_multi<1, 2, 4, 4>, sparse_caller_multi<1, 2, 5, 4>},
+                {sparse_caller_multi<1, 2, 1, 5>, sparse_caller_multi<1, 2, 2, 5>, sparse_caller_multi<1, 2, 3, 5>, sparse_caller_multi<1, 2, 4, 5>, sparse_caller_multi<1, 2, 5, 5>}
+            },
+            { // index 3
+                {sparse_caller_multi<1, 3, 1, 1>, sparse_caller_multi<1, 3, 2, 1>, sparse_caller_multi<1, 3, 3, 1>, sparse_caller_multi<1, 3, 4, 1>, sparse_caller_multi<1, 3, 5, 1>},
+                {sparse_caller_multi<1, 3, 1, 2>, sparse_caller_multi<1, 3, 2, 2>, sparse_caller_multi<1, 3, 3, 2>, sparse_caller_multi<1, 3, 4, 2>, sparse_caller_multi<1, 3, 5, 2>},
+                {sparse_caller_multi<1, 3, 1, 3>, sparse_caller_multi<1, 3, 2, 3>, sparse_caller_multi<1, 3, 3, 3>, sparse_caller_multi<1, 3, 4, 3>, sparse_caller_multi<1, 3, 5, 3>},
+                {sparse_caller_multi<1, 3, 1, 4>, sparse_caller_multi<1, 3, 2, 4>, sparse_caller_multi<1, 3, 3, 4>, sparse_caller_multi<1, 3, 4, 4>, sparse_caller_multi<1, 3, 5, 4>},
+                {sparse_caller_multi<1, 3, 1, 5>, sparse_caller_multi<1, 3, 2, 5>, sparse_caller_multi<1, 3, 3, 5>, sparse_caller_multi<1, 3, 4, 5>, sparse_caller_multi<1, 3, 5, 5>}
+            },
+            { // index 4
+                {sparse_caller_multi<1, 4, 1, 1>, sparse_caller_multi<1, 4, 2, 1>, sparse_caller_multi<1, 4, 3, 1>, sparse_caller_multi<1, 4, 4, 1>, sparse_caller_multi<1, 4, 5, 1>},
+                {sparse_caller_multi<1, 4, 1, 2>, sparse_caller_multi<1, 4, 2, 2>, sparse_caller_multi<1, 4, 3, 2>, sparse_caller_multi<1, 4, 4, 2>, sparse_caller_multi<1, 4, 5, 2>},
+                {sparse_caller_multi<1, 4, 1, 3>, sparse_caller_multi<1, 4, 2, 3>, sparse_caller_multi<1, 4, 3, 3>, sparse_caller_multi<1, 4, 4, 3>, sparse_caller_multi<1, 4, 5, 3>},
+                {sparse_caller_multi<1, 4, 1, 4>, sparse_caller_multi<1, 4, 2, 4>, sparse_caller_multi<1, 4, 3, 4>, sparse_caller_multi<1, 4, 4, 4>, sparse_caller_multi<1, 4, 5, 4>},
+                {sparse_caller_multi<1, 4, 1, 5>, sparse_caller_multi<1, 4, 2, 5>, sparse_caller_multi<1, 4, 3, 5>, sparse_caller_multi<1, 4, 4, 5>, sparse_caller_multi<1, 4, 5, 5>}
+            }
+        };
+
+        switch(index)
+        {
+            case 0:
+                bindTexture(&tex_If1_multi0, I);
+                bindTexture(&tex_Jf1_multi0, J);
+                break;
+            case 1:
+                bindTexture(&tex_If1_multi1, I);
+                bindTexture(&tex_Jf1_multi1, J);
+                break;
+            case 2:
+                bindTexture(&tex_If1_multi2, I);
+                bindTexture(&tex_Jf1_multi2, J);
+                break;
+            case 3:
+                bindTexture(&tex_If1_multi3, I);
+                bindTexture(&tex_Jf1_multi3, J);
+                break;
+            case 4:
+                bindTexture(&tex_If1_multi4, I);
+                bindTexture(&tex_Jf1_multi4, J);
+                break;
+            default:
+                CV_Error(CV_StsBadArg, "invalid execution line index");
+                break;
+        }
+
+        funcs[index][patch.y - 1][patch.x - 1](I.rows, I.cols, prevPts, nextPts, status, err, ptcount,
+            level, block, stream);
+    }
+
+    void sparse4_multi(PtrStepSz<float4> I, PtrStepSz<float4> J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
+                 int level, dim3 block, dim3 patch, cudaStream_t stream, int index)
+    {
+        typedef void (*func_t)(int rows, int cols, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
+                               int level, dim3 block, cudaStream_t stream);
+
+        static const func_t funcs[5][5][5] =
+        {
+            { // index 0
+                {sparse_caller_multi<4, 0, 1, 1>, sparse_caller_multi<4, 0, 2, 1>, sparse_caller_multi<4, 0, 3, 1>, sparse_caller_multi<4, 0, 4, 1>, sparse_caller_multi<4, 0, 5, 1>},
+                {sparse_caller_multi<4, 0, 1, 2>, sparse_caller_multi<4, 0, 2, 2>, sparse_caller_multi<4, 0, 3, 2>, sparse_caller_multi<4, 0, 4, 2>, sparse_caller_multi<4, 0, 5, 2>},
+                {sparse_caller_multi<4, 0, 1, 3>, sparse_caller_multi<4, 0, 2, 3>, sparse_caller_multi<4, 0, 3, 3>, sparse_caller_multi<4, 0, 4, 3>, sparse_caller_multi<4, 0, 5, 3>},
+                {sparse_caller_multi<4, 0, 1, 4>, sparse_caller_multi<4, 0, 2, 4>, sparse_caller_multi<4, 0, 3, 4>, sparse_caller_multi<4, 0, 4, 4>, sparse_caller_multi<4, 0, 5, 4>},
+                {sparse_caller_multi<4, 0, 1, 5>, sparse_caller_multi<4, 0, 2, 5>, sparse_caller_multi<4, 0, 3, 5>, sparse_caller_multi<4, 0, 4, 5>, sparse_caller_multi<4, 0, 5, 5>}
+            },
+            { // index 1
+                {sparse_caller_multi<4, 1, 1, 1>, sparse_caller_multi<4, 1, 2, 1>, sparse_caller_multi<4, 1, 3, 1>, sparse_caller_multi<4, 1, 4, 1>, sparse_caller_multi<4, 1, 5, 1>},
+                {sparse_caller_multi<4, 1, 1, 2>, sparse_caller_multi<4, 1, 2, 2>, sparse_caller_multi<4, 1, 3, 2>, sparse_caller_multi<4, 1, 4, 2>, sparse_caller_multi<4, 1, 5, 2>},
+                {sparse_caller_multi<4, 1, 1, 3>, sparse_caller_multi<4, 1, 2, 3>, sparse_caller_multi<4, 1, 3, 3>, sparse_caller_multi<4, 1, 4, 3>, sparse_caller_multi<4, 1, 5, 3>},
+                {sparse_caller_multi<4, 1, 1, 4>, sparse_caller_multi<4, 1, 2, 4>, sparse_caller_multi<4, 1, 3, 4>, sparse_caller_multi<4, 1, 4, 4>, sparse_caller_multi<4, 1, 5, 4>},
+                {sparse_caller_multi<4, 1, 1, 5>, sparse_caller_multi<4, 1, 2, 5>, sparse_caller_multi<4, 1, 3, 5>, sparse_caller_multi<4, 1, 4, 5>, sparse_caller_multi<4, 1, 5, 5>}
+            },
+            { // index 2
+                {sparse_caller_multi<4, 2, 1, 1>, sparse_caller_multi<4, 2, 2, 1>, sparse_caller_multi<4, 2, 3, 1>, sparse_caller_multi<4, 2, 4, 1>, sparse_caller_multi<4, 2, 5, 1>},
+                {sparse_caller_multi<4, 2, 1, 2>, sparse_caller_multi<4, 2, 2, 2>, sparse_caller_multi<4, 2, 3, 2>, sparse_caller_multi<4, 2, 4, 2>, sparse_caller_multi<4, 2, 5, 2>},
+                {sparse_caller_multi<4, 2, 1, 3>, sparse_caller_multi<4, 2, 2, 3>, sparse_caller_multi<4, 2, 3, 3>, sparse_caller_multi<4, 2, 4, 3>, sparse_caller_multi<4, 2, 5, 3>},
+                {sparse_caller_multi<4, 2, 1, 4>, sparse_caller_multi<4, 2, 2, 4>, sparse_caller_multi<4, 2, 3, 4>, sparse_caller_multi<4, 2, 4, 4>, sparse_caller_multi<4, 2, 5, 4>},
+                {sparse_caller_multi<4, 2, 1, 5>, sparse_caller_multi<4, 2, 2, 5>, sparse_caller_multi<4, 2, 3, 5>, sparse_caller_multi<4, 2, 4, 5>, sparse_caller_multi<4, 2, 5, 5>}
+            },
+            { // index 3
+                {sparse_caller_multi<4, 3, 1, 1>, sparse_caller_multi<4, 3, 2, 1>, sparse_caller_multi<4, 3, 3, 1>, sparse_caller_multi<4, 3, 4, 1>, sparse_caller_multi<4, 3, 5, 1>},
+                {sparse_caller_multi<4, 3, 1, 2>, sparse_caller_multi<4, 3, 2, 2>, sparse_caller_multi<4, 3, 3, 2>, sparse_caller_multi<4, 3, 4, 2>, sparse_caller_multi<4, 3, 5, 2>},
+                {sparse_caller_multi<4, 3, 1, 3>, sparse_caller_multi<4, 3, 2, 3>, sparse_caller_multi<4, 3, 3, 3>, sparse_caller_multi<4, 3, 4, 3>, sparse_caller_multi<4, 3, 5, 3>},
+                {sparse_caller_multi<4, 3, 1, 4>, sparse_caller_multi<4, 3, 2, 4>, sparse_caller_multi<4, 3, 3, 4>, sparse_caller_multi<4, 3, 4, 4>, sparse_caller_multi<4, 3, 5, 4>},
+                {sparse_caller_multi<4, 3, 1, 5>, sparse_caller_multi<4, 3, 2, 5>, sparse_caller_multi<4, 3, 3, 5>, sparse_caller_multi<4, 3, 4, 5>, sparse_caller_multi<4, 3, 5, 5>}
+            },
+            { // index 4
+                {sparse_caller_multi<4, 4, 1, 1>, sparse_caller_multi<4, 4, 2, 1>, sparse_caller_multi<4, 4, 3, 1>, sparse_caller_multi<4, 4, 4, 1>, sparse_caller_multi<4, 4, 5, 1>},
+                {sparse_caller_multi<4, 4, 1, 2>, sparse_caller_multi<4, 4, 2, 2>, sparse_caller_multi<4, 4, 3, 2>, sparse_caller_multi<4, 4, 4, 2>, sparse_caller_multi<4, 4, 5, 2>},
+                {sparse_caller_multi<4, 4, 1, 3>, sparse_caller_multi<4, 4, 2, 3>, sparse_caller_multi<4, 4, 3, 3>, sparse_caller_multi<4, 4, 4, 3>, sparse_caller_multi<4, 4, 5, 3>},
+                {sparse_caller_multi<4, 4, 1, 4>, sparse_caller_multi<4, 4, 2, 4>, sparse_caller_multi<4, 4, 3, 4>, sparse_caller_multi<4, 4, 4, 4>, sparse_caller_multi<4, 4, 5, 4>},
+                {sparse_caller_multi<4, 4, 1, 5>, sparse_caller_multi<4, 4, 2, 5>, sparse_caller_multi<4, 4, 3, 5>, sparse_caller_multi<4, 4, 4, 5>, sparse_caller_multi<4, 4, 5, 5>}
+            }
+        };
+
+        switch(index)
+        {
+            case 0:
+                bindTexture(&tex_If4_multi0, I);
+                bindTexture(&tex_Jf4_multi0, J);
+                break;
+            case 1:
+                bindTexture(&tex_If4_multi1, I);
+                bindTexture(&tex_Jf4_multi1, J);
+                break;
+            case 2:
+                bindTexture(&tex_If4_multi2, I);
+                bindTexture(&tex_Jf4_multi2, J);
+                break;
+            case 3:
+                bindTexture(&tex_If4_multi3, I);
+                bindTexture(&tex_Jf4_multi3, J);
+                break;
+            case 4:
+                bindTexture(&tex_If4_multi4, I);
+                bindTexture(&tex_Jf4_multi4, J);
+                break;
+            default:
+                CV_Error(CV_StsBadArg, "invalid execution line index");
+                break;
+        }
+
+        funcs[index][patch.y - 1][patch.x - 1](I.rows, I.cols, prevPts, nextPts, status, err, ptcount,
+            level, block, stream);
+    }
+
+#endif // defined(HAVE_TBB)
+
    void dense(PtrStepSzb I, PtrStepSzf J, PtrStepSzf u, PtrStepSzf v, PtrStepSzf prevU, PtrStepSzf prevV, PtrStepSzf err, int2 winSize, cudaStream_t stream)
    {
        dim3 block(16, 16);
@@ -42,12 +42,18 @@

 #include "precomp.hpp"

+#ifdef HAVE_TBB
+#include <tbb/compat/condition_variable>
+#include <tbb/mutex.h>
+#endif
+
 using namespace std;
 using namespace cv;
 using namespace cv::gpu;

 #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)

+
 cv::gpu::PyrLKOpticalFlow::PyrLKOpticalFlow() { throw_nogpu(); }
 void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat&, const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat*) { throw_nogpu(); }
 void cv::gpu::PyrLKOpticalFlow::dense(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat*) { throw_nogpu(); }
@@ -64,6 +70,23 @@ namespace pyrlk
    void sparse4(PtrStepSz<float4> I, PtrStepSz<float4> J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
                 int level, dim3 block, dim3 patch, cudaStream_t stream = 0);

+
+#if !defined(HAVE_TBB)
+#define throw_notbb() CV_Error(CV_StsNotImplemented, "The library is compiled without TBB support")
+    void loadConstants_multi(int2, int, int, cudaStream_t) { throw_notbb(); }
+    void sparse1_multi(PtrStepSzf, PtrStepSzf, const float2*, float2*, uchar*, float*, int,
+                 int, dim3, dim3, cudaStream_t, int) { throw_notbb(); }
+    void sparse4_multi(PtrStepSz<float4>, PtrStepSz<float4>, const float2*, float2*, uchar*, float*, int,
+                 int, dim3, dim3, cudaStream_t, int) { throw_notbb(); }
+#else
+    void loadConstants_multi(int2 winSize, int iters, int index = 0, cudaStream_t stream = 0);
+
+    void sparse1_multi(PtrStepSzf I, PtrStepSzf J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
+                 int level, dim3 block, dim3 patch, cudaStream_t stream = 0, int index = 0);
+    void sparse4_multi(PtrStepSz<float4> I, PtrStepSz<float4> J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
+                 int level, dim3 block, dim3 patch, cudaStream_t stream = 0, int index = 0);
+#endif
+
    void dense(PtrStepSzb I, PtrStepSzf J, PtrStepSzf u, PtrStepSzf v, PtrStepSzf prevU, PtrStepSzf prevV,
               PtrStepSzf err, int2 winSize, cudaStream_t stream = 0);
 }
@@ -98,7 +121,9 @@ namespace
    }
 }

-void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat& prevImg, const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts, GpuMat& status, GpuMat* err)
+void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat& prevImg,
+                        const GpuMat& nextImg, const GpuMat& prevPts,
+                        GpuMat& nextPts, GpuMat& status, GpuMat* err)
 {
    if (prevPts.empty())
    {
@@ -181,6 +206,137 @@ void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat& prevImg, const GpuMat& next
    }
 }

+#ifdef HAVE_TBB
+//--------------------------------------------------------------------------
+// Multi-threading support
+
+static bool index_vector_use[5] = {true, true, true, true, true}; // all free
+static tbb::mutex s_PyrLKOpticalFlow_Mutex;
+static condition_variable s_PyrLKOpticalFlow_ConditionVariable;
+
+void cv::gpu::PyrLKOpticalFlow::sparse_multi(const GpuMat& prevImg,
+        const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts,
+        GpuMat& status, Stream& stream, GpuMat* err)
+{
+    if (prevPts.empty())
+    {
+        nextPts.release();
+        status.release();
+        if (err) err->release();
+        return;
+    }
+
+    dim3 block, patch;
+    calcPatchSize(winSize, block, patch);
+
+    CV_Assert(prevImg.channels() == 1 || prevImg.channels() == 3 || prevImg.channels() == 4);
+    CV_Assert(prevImg.size() == nextImg.size() && prevImg.type() == nextImg.type());
+    CV_Assert(maxLevel >= 0);
+    CV_Assert(winSize.width > 2 && winSize.height > 2);
+    CV_Assert(patch.x > 0 && patch.x < 6 && patch.y > 0 && patch.y < 6);
+    CV_Assert(prevPts.rows == 1 && prevPts.type() == CV_32FC2);
+
+    if (useInitialFlow)
+        CV_Assert(nextPts.size() == prevPts.size() && nextPts.type() == CV_32FC2);
+    else
+        ensureSizeIsEnough(1, prevPts.cols, prevPts.type(), nextPts);
+
+    GpuMat temp1 = (useInitialFlow ? nextPts : prevPts).reshape(1);
+    GpuMat temp2 = nextPts.reshape(1);
+    multiply(temp1, Scalar::all(1.0 / (1 << maxLevel) / 2.0), temp2);
+
+    ensureSizeIsEnough(1, prevPts.cols, CV_8UC1, status);
+    status.setTo(Scalar::all(1));
+
+    if (err)
+        ensureSizeIsEnough(1, prevPts.cols, CV_32FC1, *err);
+
+    // build the image pyramids.
+
+    prevPyr_.resize(maxLevel + 1);
+    nextPyr_.resize(maxLevel + 1);
+
+    int cn = prevImg.channels();
+
+    if (cn == 1 || cn == 4)
+    {
+        prevImg.convertTo(prevPyr_[0], CV_32F);
+        nextImg.convertTo(nextPyr_[0], CV_32F);
+    }
+    else
+    {
+        buf_.resize(1);
+
+        cvtColor(prevImg, buf_[0], COLOR_BGR2BGRA);
+        buf_[0].convertTo(prevPyr_[0], CV_32F);
+
+        cvtColor(nextImg, buf_[0], COLOR_BGR2BGRA);
+        buf_[0].convertTo(nextPyr_[0], CV_32F);
+    }
+
+    for (int level = 1; level <= maxLevel; ++level)
+    {
+        pyrDown(prevPyr_[level - 1], prevPyr_[level]);
+        pyrDown(nextPyr_[level - 1], nextPyr_[level]);
+    }
+
+    //--------------------------------------------------------------------------
+    // Multithreading support
+    int index = -1;
+
+    do
+    {
+        unique_lock<tbb::mutex> ul(s_PyrLKOpticalFlow_Mutex);
+        for (unsigned int uiI = 0; uiI < 5; ++uiI)
+        {
+            if (index_vector_use[uiI])
+            {
+                index = uiI;
+                index_vector_use[uiI] = false;
+                break;
+            }
+        }
+        if (index < 0)
+            s_PyrLKOpticalFlow_ConditionVariable.wait(ul);
+
+        ul.unlock();
+    }while (index < 0);
+
+    //--------------------------------------------------------------------------
+
+    pyrlk::loadConstants_multi(make_int2(winSize.width, winSize.height), iters, index);
+
+    for (int level = maxLevel; level >= 0; level--)
+    {
+        if (cn == 1)
+        {
+            pyrlk::sparse1_multi(prevPyr_[level], nextPyr_[level],
+                prevPts.ptr<float2>(), nextPts.ptr<float2>(), status.ptr(),
+                level == 0 && err ? err->ptr<float>() : 0, prevPts.cols,
+                level, block, patch, StreamAccessor::getStream(stream), index);
+        }
+        else
+        {
+            pyrlk::sparse4_multi(prevPyr_[level], nextPyr_[level],
+                prevPts.ptr<float2>(), nextPts.ptr<float2>(), status.ptr(),
+                level == 0 && err ? err->ptr<float>() : 0, prevPts.cols,
+                level, block, patch, StreamAccessor::getStream(stream), index);
+        }
+    }
+
+    unique_lock<tbb::mutex> ul(s_PyrLKOpticalFlow_Mutex);
+    index_vector_use[index] = true;
+    s_PyrLKOpticalFlow_ConditionVariable.notify_one();
+}
+#else
+void cv::gpu::PyrLKOpticalFlow::sparse_multi(const GpuMat& /*prevImg*/,
+        const GpuMat& /*nextImg*/, const GpuMat& /*prevPts*/, GpuMat& /*nextPts*/,
+        GpuMat& /*status*/, Stream& /*stream*/, GpuMat* /*err*/)
+{
+    throw_notbb();
+}
+#endif
+
 void cv::gpu::PyrLKOpticalFlow::dense(const GpuMat& prevImg, const GpuMat& nextImg, GpuMat& u, GpuMat& v, GpuMat* err)
 {
    CV_Assert(prevImg.type() == CV_8UC1);
@@ -44,6 +44,10 @@

 #ifdef HAVE_CUDA

+#ifdef HAVE_TBB
+#include <tbb/tbb.h>
+#endif
+
 using namespace cvtest;

 //////////////////////////////////////////////////////
@@ -322,6 +326,134 @@ GPU_TEST_P(PyrLKOpticalFlow, Sparse)
    ASSERT_LE(bad_ratio, 0.01);
 }

+#ifdef HAVE_TBB
+
+struct Sparse_Multi_Functor
+{
+    explicit Sparse_Multi_Functor(const cv::Mat& in_frame0, const cv::Mat& in_frame1,
+                                  const cv::Mat& in_pts_mat,
+                                  cv::gpu::GpuMat* in_d_pts,
+                                  cv::gpu::GpuMat* in_d_nextPts,
+                                  cv::gpu::GpuMat* in_d_status,
+                                  cv::gpu::Stream* in_streams):
+                                          m_frame0(in_frame0), m_frame1(in_frame1),
+                                          m_pts_mat(in_pts_mat),
+                                          m_d_pts(in_d_pts), m_d_nextPts(in_d_nextPts),
+                                          m_d_status(in_d_status), m_streams(in_streams){}
+
+    void operator()( const tbb::blocked_range<size_t>& r ) const
+    {
+        for( size_t i = r.begin(); i != r.end(); ++i )
+        {
+            m_d_pts[i].upload(m_pts_mat);
+            cv::gpu::PyrLKOpticalFlow pyrLK;
+            pyrLK.sparse_multi(loadMat(m_frame0), loadMat(m_frame1), m_d_pts[i],
+                               m_d_nextPts[i], m_d_status[i], m_streams[i]);
+            m_streams[i].waitForCompletion();
+        }
+    }
+
+    const cv::Mat& m_frame0;
+    const cv::Mat& m_frame1;
+    const cv::Mat& m_pts_mat;
+    cv::gpu::GpuMat* m_d_pts;
+    cv::gpu::GpuMat* m_d_nextPts;
+    cv::gpu::GpuMat* m_d_status;
+    cv::gpu::Stream* m_streams;
+};
+
+
+GPU_TEST_P(PyrLKOpticalFlow, Sparse_Multi)
+{
+    cv::Mat frame0 = readImage("opticalflow/frame0.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+    ASSERT_FALSE(frame0.empty());
+
+    cv::Mat frame1 = readImage("opticalflow/frame1.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+    ASSERT_FALSE(frame1.empty());
+
+    cv::Mat gray_frame;
+    if (useGray)
+        gray_frame = frame0;
+    else
+        cv::cvtColor(frame0, gray_frame, cv::COLOR_BGR2GRAY);
+
+    std::vector<cv::Point2f> pts;
+    cv::goodFeaturesToTrack(gray_frame, pts, 1000, 0.01, 0.0);
+
+    //--------------------------------------------------------------------------
+    // GPU
+    const unsigned int NB_EXEC_LINES = 27;
+
+    cv::gpu::GpuMat d_pts[NB_EXEC_LINES];
+    cv::gpu::GpuMat d_nextPts[NB_EXEC_LINES];
+    cv::gpu::GpuMat d_status[NB_EXEC_LINES];
+    cv::gpu::Stream streams[NB_EXEC_LINES];
+
+    cv::Mat pts_mat(1, (int) pts.size(), CV_32FC2, (void*) &pts[0]);
+
+    tbb::parallel_for(tbb::blocked_range<size_t>(0, NB_EXEC_LINES),
+                      Sparse_Multi_Functor(frame0, frame1, pts_mat,
+                                           d_pts, d_nextPts, d_status, streams));
+
+    std::vector<cv::Point2f> nextPts[NB_EXEC_LINES];
+    std::vector<unsigned char> status[NB_EXEC_LINES];
+
+    for (unsigned int i = 0; i < NB_EXEC_LINES; ++i)
+    {
+        nextPts[i].resize(d_nextPts[i].cols);
+        cv::Mat nextPts_mat(1, d_nextPts[i].cols, CV_32FC2, (void*) &(nextPts[i][0]));
+        d_nextPts[i].download(nextPts_mat);
+
+        status[i].resize(d_status[i].cols);
+        cv::Mat status_mat(1, d_status[i].cols, CV_8UC1, (void*) &(status[i][0]));
+        d_status[i].download(status_mat);
+    }
+
+    //--------------------------------------------------------------------------
+    // CPU
+    std::vector<cv::Point2f> nextPts_gold;
+    std::vector<unsigned char> status_gold;
+    cv::calcOpticalFlowPyrLK(frame0, frame1, pts, nextPts_gold, status_gold, cv::noArray());
+
+    //--------------------------------------------------------------------------
+    // CHECKS
+    for (unsigned int uiI = 0; uiI < NB_EXEC_LINES; ++uiI)
+    {
+        ASSERT_EQ(nextPts_gold.size(), nextPts[uiI].size());
+        ASSERT_EQ(status_gold.size(), status[uiI].size());
+    }
+
+    size_t mistmatch = 0;
+    for (unsigned int uiI = 0; uiI < NB_EXEC_LINES; ++uiI)
+    {
+        for (size_t i = 0; i < nextPts[uiI].size(); ++i)
+        {
+            cv::Point2i a = nextPts[uiI][i];
+            cv::Point2i b = nextPts_gold[i];
+
+            if (status[uiI][i] != status_gold[i])
+            {
+                ++mistmatch;
+                continue;
+            }
+
+            if (status[uiI][i])
+            {
+                bool eq = std::abs(a.x - b.x) <= 1 && std::abs(a.y - b.y) <= 1;
+
+                if (!eq)
+                    ++mistmatch;
+            }
+        }
+    }
+
+    double bad_ratio = static_cast<double>(mistmatch) / (nextPts[0].size() * NB_EXEC_LINES);
+
+    ASSERT_LE(bad_ratio, 0.01);
+}
+
+#endif // HAVE_TBB
+
 INSTANTIATE_TEST_CASE_P(GPU_Video, PyrLKOpticalFlow, testing::Combine(
    ALL_DEVICES,
    testing::Values(UseGray(true), UseGray(false))));
@@ -42,6 +42,7 @@

 #include "precomp.hpp"
 #include "bitstrm.hpp"
+#include "utils.hpp"

 namespace cv
 {
@@ -103,7 +104,6 @@ void  RBaseStream::readBlock()
    fseek( m_file, m_block_pos, SEEK_SET );
    size_t readed = fread( m_start, 1, m_block_size, m_file );
    m_end = m_start + readed;
-    m_current = m_start;

    if( readed == 0 || m_current >= m_end )
        throw RBS_THROW_EOS;
@@ -164,7 +164,7 @@ void  RBaseStream::release()

 void  RBaseStream::setPos( int pos )
 {
-    assert( isOpened() && pos >= 0 );
+    CV_Assert(isOpened() && pos >= 0);

    if( !m_file )
    {
@@ -181,14 +181,19 @@ void  RBaseStream::setPos( int pos )

 int  RBaseStream::getPos()
 {
-    assert( isOpened() );
-    return m_block_pos + (int)(m_current - m_start);
+    CV_Assert(isOpened());
+    int pos = validateToInt((m_current - m_start) + m_block_pos);
+    CV_Assert(pos >= m_block_pos); // overflow check
+    CV_Assert(pos >= 0); // overflow check
+    return pos;
 }

 void  RBaseStream::skip( int bytes )
 {
-    assert( bytes >= 0 );
+    CV_Assert(bytes >= 0);
+    uchar* old = m_current;
    m_current += bytes;
+    CV_Assert(m_current >= old);  // overflow check
 }

 /////////////////////////  RLByteStream ////////////////////////////
@@ -220,7 +225,7 @@ int RLByteStream::getBytes( void* buffer, int count )
 {
    uchar*  data = (uchar*)buffer;
    int readed = 0;
-    assert( count >= 0 );
+    CV_Assert(count >= 0);

    while( count > 0 )
    {
@@ -371,7 +376,7 @@ void  WBaseStream::writeBlock()
 {
    int size = (int)(m_current - m_start);

-    assert( isOpened() );
+    CV_Assert(isOpened());
    if( size == 0 )
        return;

@@ -442,7 +447,7 @@ void  WBaseStream::release()

 int  WBaseStream::getPos()
 {
-    assert( isOpened() );
+    CV_Assert(isOpened());
    return m_block_pos + (int)(m_current - m_start);
 }

@@ -465,7 +470,7 @@ void WLByteStream::putBytes( const void* buffer, int count )
 {
    uchar* data = (uchar*)buffer;

-    assert( data && m_current && count >= 0 );
+    CV_Assert(data && m_current && count >= 0);

    while( count )
    {
@@ -192,7 +192,11 @@ CvDC1394::~CvDC1394()
    dc = 0;
 }

-static CvDC1394 dc1394;
+static CvDC1394& getDC1394()
+{
+    static CvDC1394 dc1394;
+    return dc1394;
+}

 class CvCaptureCAM_DC1394_v2_CPP : public CvCapture
 {
@@ -451,7 +455,7 @@ bool CvCaptureCAM_DC1394_v2_CPP::startCapture()
    code = dc1394_capture_setup(dcCam, nDMABufs, DC1394_CAPTURE_FLAGS_DEFAULT);
    if (code >= 0)
    {
-        FD_SET(dc1394_capture_get_fileno(dcCam), &dc1394.camFds);
+        FD_SET(dc1394_capture_get_fileno(dcCam), &getDC1394().camFds);
        dc1394_video_set_transmission(dcCam, DC1394_ON);
        if (cameraId == VIDERE)
        {
@@ -477,15 +481,15 @@ bool CvCaptureCAM_DC1394_v2_CPP::open(int index)

    close();

-    if (!dc1394.dc)
+    if (!getDC1394().dc)
        goto _exit_;

-    err = dc1394_camera_enumerate(dc1394.dc, &cameraList);
+    err = dc1394_camera_enumerate(getDC1394().dc, &cameraList);
    if (err < 0 || !cameraList || (unsigned)index >= (unsigned)cameraList->num)
        goto _exit_;

    guid = cameraList->ids[index].guid;
-    dcCam = dc1394_camera_new(dc1394.dc, guid);
+    dcCam = dc1394_camera_new(getDC1394().dc, guid);
    if (!dcCam)
        goto _exit_;

@@ -510,8 +514,8 @@ void CvCaptureCAM_DC1394_v2_CPP::close()
        // check for fileno valid before using
        int fileno=dc1394_capture_get_fileno(dcCam);

-        if (fileno>=0 && FD_ISSET(fileno, &dc1394.camFds))
-            FD_CLR(fileno, &dc1394.camFds);
+        if (fileno>=0 && FD_ISSET(fileno, &getDC1394().camFds))
+            FD_CLR(fileno, &getDC1394().camFds);
        dc1394_video_set_transmission(dcCam, DC1394_OFF);
        dc1394_capture_stop(dcCam);
        dc1394_camera_free(dcCam);
@@ -173,6 +173,10 @@ extern "C" {
 #define AV_PIX_FMT_GRAY16BE PIX_FMT_GRAY16BE
 #endif

+#ifndef PKT_FLAG_KEY
+#define PKT_FLAG_KEY AV_PKT_FLAG_KEY
+#endif
+
 #if LIBAVUTIL_BUILD >= (LIBAVUTIL_VERSION_MICRO >= 100 \
    ? CALC_FFMPEG_VERSION(52, 38, 100) : CALC_FFMPEG_VERSION(52, 13, 0))
 #define USE_AV_FRAME_GET_BUFFER 1
@@ -1481,13 +1485,20 @@ static AVStream *icv_add_video_stream_FFMPEG(AVFormatContext *oc,
    // some formats want stream headers to be seperate
    if(oc->oformat->flags & AVFMT_GLOBALHEADER)
    {
+#if LIBAVCODEC_BUILD > CALC_FFMPEG_VERSION(56, 35, 0)
+        c->flags |= AV_CODEC_FLAG_GLOBAL_HEADER;
+#else
        c->flags |= CODEC_FLAG_GLOBAL_HEADER;
+#endif
    }
 #endif

 #if LIBAVCODEC_BUILD >= CALC_FFMPEG_VERSION(52, 42, 0)
    st->avg_frame_rate = (AVRational){frame_rate, frame_rate_base};
 #endif
+#if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(55, 20, 0)
+    st->time_base = c->time_base;
+#endif

    return st;
 }
@@ -1509,23 +1520,24 @@ static int icv_av_write_frame_FFMPEG( AVFormatContext * oc, AVStream * video_st,
 #endif
    int ret = OPENCV_NO_FRAMES_WRITTEN_CODE;

-    if (oc->oformat->flags & AVFMT_RAWPICTURE) {
+#if LIBAVFORMAT_BUILD < CALC_FFMPEG_VERSION(57, 0, 0)
+    if (oc->oformat->flags & AVFMT_RAWPICTURE)
+    {
        /* raw video case. The API will change slightly in the near
           futur for that */
        AVPacket pkt;
        av_init_packet(&pkt);

-#ifndef PKT_FLAG_KEY
-#define PKT_FLAG_KEY AV_PKT_FLAG_KEY
-#endif
-
        pkt.flags |= PKT_FLAG_KEY;
        pkt.stream_index= video_st->index;
        pkt.data= (uint8_t *)picture;
        pkt.size= sizeof(AVPicture);

        ret = av_write_frame(oc, &pkt);
-    } else {
+    }
+    else
+#endif
+    {
        /* encode the image */
        AVPacket pkt;
        av_init_packet(&pkt);
@@ -1683,7 +1695,9 @@ void CvVideoWriter_FFMPEG::close()
    /* write the trailer, if any */
    if(ok && oc)
    {
-        if( (oc->oformat->flags & AVFMT_RAWPICTURE) == 0 )
+#if LIBAVFORMAT_BUILD < CALC_FFMPEG_VERSION(57, 0, 0)
+        if (!(oc->oformat->flags & AVFMT_RAWPICTURE))
+#endif
        {
            for(;;)
            {
@@ -1917,7 +1931,10 @@ bool CvVideoWriter_FFMPEG::open( const char * filename, int fourcc,

    outbuf = NULL;

-    if (!(oc->oformat->flags & AVFMT_RAWPICTURE)) {
+#if LIBAVFORMAT_BUILD < CALC_FFMPEG_VERSION(57, 0, 0)
+    if (!(oc->oformat->flags & AVFMT_RAWPICTURE))
+#endif
+    {
        /* allocate output buffer */
        /* assume we will never get codec output with more than 4 bytes per pixel... */
        outbuf_size = width*height*4;
@@ -2211,7 +2228,11 @@ AVStream* OutputMediaStream_FFMPEG::addVideoStream(AVFormatContext *oc, CV_CODEC
        // some formats want stream headers to be seperate
        if (oc->oformat->flags & AVFMT_GLOBALHEADER)
        {
-            c->flags |= CODEC_FLAG_GLOBAL_HEADER;
+            #if LIBAVCODEC_BUILD > CALC_FFMPEG_VERSION(56, 35, 0)
+                c->flags |= AV_CODEC_FLAG_GLOBAL_HEADER;
+            #else
+                c->flags |= CODEC_FLAG_GLOBAL_HEADER;
+            #endif
        }
    #endif

@@ -634,11 +634,11 @@ static int icvOpenCamera_QT (CvCapture_QT_Cam * capture, const int index)
    OPENCV_ASSERT (capture,            "icvOpenCamera_QT", "'capture' is a NULL-pointer");
    OPENCV_ASSERT (index >=0, "icvOpenCamera_QT", "camera index is negative");

-    ComponentDescription	component_description;
-    Component				component = 0;
+    ComponentDescription    component_description;
+    Component               component = 0;
    int                     number_of_inputs = 0;
    Rect                    myRect;
-    ComponentResult			result = noErr;
+    ComponentResult         result = noErr;


    // travers all components and count video digitizer channels
@@ -92,6 +92,7 @@ bool  BmpDecoder::readHeader()
        m_offset = m_strm.getDWord();

        int  size = m_strm.getDWord();
+        CV_Assert(size > 0); // overflow, 2Gb limit

        if( size >= 36 )
        {
@@ -115,7 +116,7 @@ bool  BmpDecoder::readHeader()

                if( m_bpp <= 8 )
                {
-                    CV_Assert(clrused <= 256);
+                    CV_Assert(clrused >= 0 && clrused <= 256);
                    memset(m_palette, 0, sizeof(m_palette));
                    m_strm.getBytes(m_palette, (clrused == 0? 1<<m_bpp : clrused)*4 );
                    iscolor = IsColorPalette( m_palette, m_bpp );
@@ -77,7 +77,8 @@ static JasperInitializer initialize_jasper;

 Jpeg2KDecoder::Jpeg2KDecoder()
 {
-    m_signature = '\0' + string() + '\0' + string() + '\0' + string("\x0cjP  \r\n\x87\n");
+    static const unsigned char signature_[12] = { 0, 0, 0, 0x0c, 'j', 'P', ' ', ' ', 13, 10, 0x87, 10};
+    m_signature = string((const char*)signature_, (const char*)signature_ + sizeof(signature_));
    m_stream = 0;
    m_image = 0;
 }
@@ -121,6 +122,8 @@ bool  Jpeg2KDecoder::readHeader()
        jas_image_t* image = jas_image_decode( stream, -1, 0 );
        m_image = image;
        if( image ) {
+            CV_Assert(0 == (jas_image_tlx(image)) && "not supported");
+            CV_Assert(0 == (jas_image_tly(image)) && "not supported");
            m_width = jas_image_width( image );
            m_height = jas_image_height( image );

@@ -130,14 +133,31 @@ bool  Jpeg2KDecoder::readHeader()
            for( int i = 0; i < numcmpts; i++ )
            {
                int depth_i = jas_image_cmptprec( image, i );
+                CV_Assert(depth == 0 || depth == depth_i); // component data type mismatch
                depth = MAX(depth, depth_i);
                if( jas_image_cmpttype( image, i ) > 2 )
                    continue;
+                int sgnd = jas_image_cmptsgnd(image, i);
+                int xstart = jas_image_cmpttlx(image, i);
+                int xend = jas_image_cmptbrx(image, i);
+                int xstep = jas_image_cmpthstep(image, i);
+                int ystart = jas_image_cmpttly(image, i);
+                int yend = jas_image_cmptbry(image, i);
+                int ystep = jas_image_cmptvstep(image, i);
+                CV_Assert(sgnd == 0 && "not supported");
+                CV_Assert(xstart == 0 && "not supported");
+                CV_Assert(ystart == 0 && "not supported");
+                CV_Assert(xstep == 1 && "not supported");
+                CV_Assert(ystep == 1 && "not supported");
+                CV_Assert(xend == m_width);
+                CV_Assert(yend == m_height);
                cntcmpts++;
            }

            if( cntcmpts )
            {
+                CV_Assert(depth == 8 || depth == 16);
+                CV_Assert(cntcmpts == 1 || cntcmpts == 3);
                m_type = CV_MAKETYPE(depth <= 8 ? CV_8U : CV_16U, cntcmpts > 1 ? 3 : 1);
                result = true;
            }
@@ -150,9 +170,15 @@ bool  Jpeg2KDecoder::readHeader()
    return result;
 }

+static void Jpeg2KDecoder_close(Jpeg2KDecoder* ptr)
+{
+    ptr->close();
+}
+template<> void Ptr<Jpeg2KDecoder>::delete_obj() { Jpeg2KDecoder_close(obj); }

 bool  Jpeg2KDecoder::readData( Mat& img )
 {
+    Ptr<Jpeg2KDecoder> close_this(this);  // auto cleanup: Jpeg2KDecoder_close
    bool result = false;
    int color = img.channels() > 1;
    uchar* data = img.data;
@@ -204,11 +230,16 @@ bool  Jpeg2KDecoder::readData( Mat& img )
                    result = true;
                }
                else
-                    fprintf(stderr, "JPEG 2000 LOADER ERROR: cannot convert colorspace\n");
+                {
+                    jas_cmprof_destroy(clrprof);
+                    CV_Error(CV_StsError, "JPEG 2000 LOADER ERROR: cannot convert colorspace");
+                }
                jas_cmprof_destroy( clrprof );
            }
            else
-                fprintf(stderr, "JPEG 2000 LOADER ERROR: unable to create colorspace\n");
+            {
+                CV_Error(CV_StsError, "JPEG 2000 LOADER ERROR: unable to create colorspace");
+            }
        }
        else
            result = true;
@@ -257,8 +288,8 @@ bool  Jpeg2KDecoder::readData( Mat& img )
                                result = readComponent16u( ((unsigned short *)data) + i, buffer, validateToInt(step / 2), cmptlut[i], maxval, offset, ncmpts );
                            if( !result )
                            {
-                                i = ncmpts;
-                                result = false;
+                                jas_matrix_destroy( buffer );
+                                CV_Error(CV_StsError, "JPEG2000 LOADER ERROR: failed to read component");
                            }
                        }
                        jas_matrix_destroy( buffer );
@@ -267,10 +298,12 @@ bool  Jpeg2KDecoder::readData( Mat& img )
            }
        }
        else
-            fprintf(stderr, "JPEG2000 LOADER ERROR: colorspace conversion failed\n" );
+        {
+            CV_Error(CV_StsError, "JPEG2000 LOADER ERROR: colorspace conversion failed");
+        }
    }

-    close();
+    CV_Assert(result == true);

 #ifndef WIN32
    if (!clr.empty())
@@ -331,7 +331,7 @@ bool  PxMDecoder::readData( Mat& img )
                        }
                    }
                    else
-                        memcpy( data, src, m_width*(bit_depth/8) );
+                        memcpy(data, src, img.elemSize1()*m_width);
                }
                else
                {
@@ -120,7 +120,7 @@ bool  SunRasterDecoder::readHeader()
                    m_type = IsColorPalette( m_palette, m_bpp ) ? CV_8UC3 : CV_8UC1;
                    m_offset = m_strm.getPos();

-                    assert( m_offset == 32 + m_maplength );
+                    CV_Assert(m_offset == 32 + m_maplength);
                    result = true;
                }
            }
@@ -133,7 +133,7 @@ bool  SunRasterDecoder::readHeader()

                m_offset = m_strm.getPos();

-                assert( m_offset == 32 + m_maplength );
+                CV_Assert(m_offset == 32 + m_maplength);
                result = true;
            }
        }
@@ -226,7 +226,7 @@ bool  SunRasterDecoder::readData( Mat& img )
                        code = m_strm.getByte();
                        if( len > line_end - tsrc )
                        {
-                            assert(0);
+                            CV_Error(CV_StsInternal, "");
                            goto bad_decoding_1bpp;
                        }

@@ -367,7 +367,7 @@ bad_decoding_end:
            result = true;
            break;
        default:
-            assert(0);
+            CV_Error(CV_StsInternal, "");
        }
    }
    catch( ... )
@@ -54,7 +54,7 @@
 #include <string.h>
 #include <limits.h>
 #include <ctype.h>
-#include <assert.h>
+#include <assert.h>  // FIX IT: remove this

 #if defined WIN32 || defined WINCE
    #if !defined _WIN32_WINNT
@@ -670,7 +670,7 @@ cvConvertImage( const CvArr* srcarr, CvArr* dstarr, int flags )
            icvCvt_BGR2Gray_8u_C3C1R( s, s_step, d, d_step, size, swap_rb );
            break;
        case 33:
-            assert( swap_rb );
+            CV_Assert(swap_rb);
            icvCvt_RGB2BGR_8u_C3R( s, s_step, d, d_step, size );
            break;
        case 41:
@@ -320,6 +320,7 @@ public:

            for (unsigned int i = 0; i < frameCount && next; ++i)
            {
+                SCOPED_TRACE(cv::format("frame=%d", (int)i));
                Mat actual;
                (*capture) >> actual;

@@ -543,6 +543,7 @@ void FeaturesMatcher::operator ()(const vector<ImageFeatures> &features, vector<
            if (features[i].keypoints.size() > 0 && features[j].keypoints.size() > 0 && mask_(i, j))
                near_pairs.push_back(make_pair(i, j));

+    pairwise_matches.clear(); // clear history values
    pairwise_matches.resize(num_images * num_images);
    MatchPairsBody body(*this, features, pairwise_matches, near_pairs);

@@ -49,6 +49,7 @@ def build_opencv(srcroot, buildroot, target, arch):
    # for some reason, if you do not specify CMAKE_BUILD_TYPE, it puts libs to "RELEASE" rather than "Release"
    cmakeargs = ("-GXcode " +
                "-DCMAKE_BUILD_TYPE=Release " +
+                ("-DIOS_ARCH=%s " % arch) +
                "-DCMAKE_TOOLCHAIN_FILE=%s/platforms/ios/cmake/Toolchains/Toolchain-%s_Xcode.cmake " +
                "-DBUILD_opencv_world=ON " +
                "-DCMAKE_C_FLAGS=\"-Wno-implicit-function-declaration\" " +
@@ -43,7 +43,10 @@ set (no_warn "-Wno-unused-function -Wno-overloaded-virtual")
 set (CMAKE_C_FLAGS "${no_warn}")
 set (CMAKE_CXX_FLAGS "-stdlib=libc++ -fvisibility=hidden -fvisibility-inlines-hidden ${no_warn}")

-set (CMAKE_CXX_FLAGS_RELEASE "-DNDEBUG -O3 -fomit-frame-pointer -ffast-math")
+set(CMAKE_CXX_FLAGS_RELEASE "-DNDEBUG -O3 -ffast-math")
+if(NOT IOS_ARCH STREQUAL "armv7" AND NOT IOS_ARCH STREQUAL "armv7s")
+  set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -fomit-frame-pointer")
+endif()

 if (HAVE_FLAG_SEARCH_PATHS_FIRST)
    set (CMAKE_C_LINK_FLAGS "-Wl,-search_paths_first ${CMAKE_C_LINK_FLAGS}")
@@ -2,7 +2,10 @@ message (STATUS "Setting up iPhoneOS toolchain")
 set (IPHONEOS TRUE)

 # Standard settings
-set (CMAKE_SYSTEM_NAME iOS)
+set(CMAKE_SYSTEM_NAME iOS)
+set(CMAKE_SYSTEM_VERSION 6.0)
+set(CMAKE_SYSTEM_PROCESSOR "${IOS_ARCH}")
+
 # Include extra modules for the iOS platform files
 set (CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/platforms/ios/cmake/Modules")

@@ -2,7 +2,10 @@ message (STATUS "Setting up iPhoneSimulator toolchain")
 set (IPHONESIMULATOR TRUE)

 # Standard settings
-set (CMAKE_SYSTEM_NAME iOS)
+set(CMAKE_SYSTEM_NAME iOS)
+set(CMAKE_SYSTEM_VERSION 6.0)
+set(CMAKE_SYSTEM_PROCESSOR "${IOS_ARCH}")
+
 # Include extra modules for the iOS platform files
 set (CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/platforms/ios/cmake/Modules")

@@ -90,6 +90,12 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
    list(REMOVE_ITEM all_samples "driver_api_multi.cpp")
    list(REMOVE_ITEM all_samples "driver_api_stereo_multi.cpp")
  endif()
+  if(NOT HAVE_CUDA
+      OR NOT HAVE_TBB
+      OR OpenCV_FOUND  # via find_package() there is no access to cvconfig.h
+  )
+    list(REMOVE_ITEM all_samples "pyrlk_optical_flow_multithreading.cpp")
+  endif()

  foreach(sample_filename ${all_samples})
    get_filename_component(sample ${sample_filename} NAME_WE)
@@ -111,6 +117,9 @@ if (OCV_DEPENDENCIES_FOUND AND INSTALL_C_EXAMPLES AND NOT WIN32)
    list_filterout(install_list ".*driver_api_multi.cpp")
    list_filterout(install_list ".*driver_api_stereo_multi.cpp")
  endif()
+  if(NOT HAVE_CUDA OR NOT HAVE_TBB)
+      list(REMOVE_ITEM install_list "pyrlk_optical_flow_multithreading.cpp")
+  endif()
  install(FILES ${install_list}
          DESTINATION "${OPENCV_SAMPLES_SRC_INSTALL_PATH}/gpu"
          PERMISSIONS OWNER_READ OWNER_WRITE GROUP_READ WORLD_READ COMPONENT samples)
@@ -0,0 +1,272 @@
+#include <iostream>
+#include <vector>
+#include <sstream>
+
+#include "opencv2/core/core.hpp"
+#include "cvconfig.h"
+
+#ifdef HAVE_TBB
+#include <tbb/parallel_for_each.h>
+#include <tbb/task_scheduler_init.h>
+
+#include "opencv2/imgproc/imgproc.hpp"
+#include "opencv2/highgui/highgui.hpp"
+#include "opencv2/video/video.hpp"
+#include "opencv2/gpu/gpu.hpp"
+
+using namespace std;
+using namespace cv;
+using namespace cv::gpu;
+
+static void download(const GpuMat& d_mat, vector<Point2f>& vec)
+{
+    vec.resize(d_mat.cols);
+    Mat mat(1, d_mat.cols, CV_32FC2, (void*)&vec[0]);
+    d_mat.download(mat);
+}
+
+static void download(const GpuMat& d_mat, vector<uchar>& vec)
+{
+    vec.resize(d_mat.cols);
+    Mat mat(1, d_mat.cols, CV_8UC1, (void*)&vec[0]);
+    d_mat.download(mat);
+}
+
+static void drawArrows(Mat& frame, const vector<Point2f>& prevPts, const vector<Point2f>& nextPts, const vector<uchar>& status, Scalar line_color = Scalar(0, 0, 255))
+{
+    for (size_t i = 0; i < prevPts.size(); ++i)
+    {
+        if (status[i])
+        {
+            int line_thickness = 1;
+
+            Point p = prevPts[i];
+            Point q = nextPts[i];
+
+            double angle = atan2((double) p.y - q.y, (double) p.x - q.x);
+
+            double hypotenuse = sqrt( (double)(p.y - q.y)*(p.y - q.y) + (double)(p.x - q.x)*(p.x - q.x) );
+
+            if (hypotenuse < 1.0)
+                continue;
+
+            // Here we lengthen the arrow by a factor of three.
+            q.x = (int) (p.x - 3 * hypotenuse * cos(angle));
+            q.y = (int) (p.y - 3 * hypotenuse * sin(angle));
+
+            // Now we draw the main line of the arrow.
+            line(frame, p, q, line_color, line_thickness);
+
+            // Now draw the tips of the arrow. I do some scaling so that the
+            // tips look proportional to the main line of the arrow.
+
+            p.x = (int) (q.x + 9 * cos(angle + CV_PI / 4));
+            p.y = (int) (q.y + 9 * sin(angle + CV_PI / 4));
+            line(frame, p, q, line_color, line_thickness);
+
+            p.x = (int) (q.x + 9 * cos(angle - CV_PI / 4));
+            p.y = (int) (q.y + 9 * sin(angle - CV_PI / 4));
+            line(frame, p, q, line_color, line_thickness);
+        }
+    }
+}
+
+template <typename T> inline T clamp (T x, T a, T b)
+{
+    return ((x) > (a) ? ((x) < (b) ? (x) : (b)) : (a));
+}
+
+template <typename T> inline T mapValue(T x, T a, T b, T c, T d)
+{
+    x = clamp(x, a, b);
+    return c + (d - c) * (x - a) / (b - a);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////
+struct S_Thread_data
+{
+    Size winSize;
+    int maxLevel;
+    int iters;
+
+    Stream stream;
+    Mat frame0;
+    Mat frame1;
+    Mat frame1Gray;
+
+    GpuMat d_frame0Gray;
+    GpuMat d_prevPts;
+    bool useGray;
+};
+
+struct pyrLK_task
+{
+    pyrLK_task(size_t n):
+        _n(n),
+        _thread_data(NULL){}
+
+    void operator()()
+    {
+        // Sparse
+
+        PyrLKOpticalFlow d_pyrLK;
+
+        d_pyrLK.winSize.width = _thread_data->winSize.width;
+        d_pyrLK.winSize.height = _thread_data->winSize.height;
+        d_pyrLK.maxLevel = _thread_data->maxLevel;
+        d_pyrLK.iters = _thread_data->iters;
+
+        GpuMat d_frame0(_thread_data->frame0);
+        GpuMat d_frame1(_thread_data->frame1);
+        GpuMat d_frame1Gray(_thread_data->frame1Gray);
+        GpuMat d_nextPts;
+        GpuMat d_status;
+
+        bool useGray = _thread_data->useGray;
+
+        d_pyrLK.sparse_multi(useGray ? _thread_data->d_frame0Gray : d_frame0,
+                    useGray ? d_frame1Gray : d_frame1,
+                            _thread_data->d_prevPts, d_nextPts,
+                    d_status, _thread_data->stream, NULL);
+
+        // Draw arrows
+
+        vector<Point2f> prevPts(_thread_data->d_prevPts.cols);
+        download(_thread_data->d_prevPts, prevPts);
+
+        vector<Point2f> nextPts(d_nextPts.cols);
+        download(d_nextPts, nextPts);
+
+        vector<uchar> status(d_status.cols);
+        download(d_status, status);
+
+        drawArrows(_thread_data->frame0, prevPts, nextPts, status, Scalar(255, 0, 0));
+    }
+
+    size_t _n;
+    struct S_Thread_data* _thread_data;
+};
+
+template <typename T> struct invoker {
+  void operator()(T& it) const {it();}
+};
+
+#define THREADS_NB    12
+
+int main(int argc, const char* argv[])
+{
+    const char* keys =
+        "{ h            | help           | false | print help message }"
+        "{ l            | left           |       | specify left image }"
+        "{ r            | right          |       | specify right image }"
+        "{ gray         | gray           | false | use grayscale sources [PyrLK Sparse] }"
+        "{ win_size     | win_size       | 21    | specify windows size [PyrLK] }"
+        "{ max_level    | max_level      | 3     | specify max level [PyrLK] }"
+        "{ iters        | iters          | 30    | specify iterations count [PyrLK] }"
+        "{ points       | points         | 4000  | specify points count [GoodFeatureToTrack] }"
+        "{ min_dist     | min_dist       | 0     | specify minimal distance between points [GoodFeatureToTrack] }";
+
+    CommandLineParser cmd(argc, argv, keys);
+
+    if (cmd.get<bool>("help"))
+    {
+        cout << "Usage: pyrlk_optical_flow_multithreading [options]" << endl;
+        cout << "Avaible options:" << endl;
+        cmd.printParams();
+        return 0;
+    }
+
+    string fname0 = cmd.get<string>("left");
+    string fname1 = cmd.get<string>("right");
+
+    if (fname0.empty() || fname1.empty())
+    {
+        cerr << "Missing input file names" << endl;
+        return -1;
+    }
+
+    bool useGray = cmd.get<bool>("gray");
+    int winSize = cmd.get<int>("win_size");
+    int maxLevel = cmd.get<int>("max_level");
+    int iters = cmd.get<int>("iters");
+    int points = cmd.get<int>("points");
+    double minDist = cmd.get<double>("min_dist");
+
+    Mat frame0 = imread(fname0);
+    Mat frame1 = imread(fname1);
+
+    if (frame0.empty() || frame1.empty())
+    {
+        cout << "Can't load input images" << endl;
+        return -1;
+    }
+
+    cout << "Image size : " << frame0.cols << " x " << frame0.rows << endl;
+    cout << "Points count : " << points << endl;
+
+    cout << endl;
+
+    Mat frame0Gray;
+    cvtColor(frame0, frame0Gray, COLOR_BGR2GRAY);
+    Mat frame1Gray;
+    cvtColor(frame1, frame1Gray, COLOR_BGR2GRAY);
+
+    // goodFeaturesToTrack
+
+    GoodFeaturesToTrackDetector_GPU detector(points, 0.01, minDist);
+
+    GpuMat d_frame0Gray(frame0Gray);
+    GpuMat d_prevPts;
+
+    detector(d_frame0Gray, d_prevPts);
+
+    // Sparse
+
+    tbb::task_scheduler_init init(THREADS_NB);
+
+    std::vector<pyrLK_task> tasks;
+
+    S_Thread_data s_thread_data[THREADS_NB];
+
+    for (unsigned int uiI = 0; uiI < THREADS_NB; ++uiI)
+    {
+        s_thread_data[uiI].stream = Stream();
+        s_thread_data[uiI].frame0 = frame0.clone();
+        s_thread_data[uiI].frame1 = frame1.clone();
+        s_thread_data[uiI].frame1Gray = frame0Gray.clone();
+
+        s_thread_data[uiI].iters = iters;
+        s_thread_data[uiI].useGray = useGray;
+        s_thread_data[uiI].maxLevel = maxLevel;
+        s_thread_data[uiI].winSize.height = winSize;
+        s_thread_data[uiI].winSize.width = winSize;
+        s_thread_data[uiI].d_frame0Gray = d_frame0Gray.clone();
+        s_thread_data[uiI].d_prevPts = d_prevPts.clone();
+
+        tasks.push_back(pyrLK_task(uiI));
+        tasks.back()._thread_data = &(s_thread_data[uiI]);
+    }
+
+    tbb::parallel_for_each(tasks.begin(),tasks.end(),invoker<pyrLK_task>());
+
+    for (unsigned int uiI = 0; uiI < THREADS_NB; ++uiI)
+    {
+        stringstream ss;
+        ss << "PyrLK MultiThreading [Sparse] " << uiI;
+
+        imshow(ss.str(), s_thread_data[uiI].frame0);
+        ss.str("");
+    }
+
+    waitKey();
+
+    return 0;
+}
+#else
+int main(int , const char* [])
+{
+    std::cout << "This example pyrlk_optical_flow_multithreading must be compiled with TBB Option" << std::endl;
+    return 0;
+}
+#endif // HAVE_TBB
@@ -8,9 +8,9 @@ Example of using features2d framework for interactive video homography matching.
 ORB features and FLANN matcher are used. The actual tracking is implemented by
 PlaneTracker class in plane_tracker.py

-Inspired by http://www.youtube.com/watch?v=-ZNYoL8rzPY
+Inspired by https://www.youtube.com/watch?v=-ZNYoL8rzPY

-video: http://www.youtube.com/watch?v=FirtmYcC0Vc
+video: https://www.youtube.com/watch?v=FirtmYcC0Vc

 Usage
 -----
@@ -8,7 +8,7 @@ This sample shows an example of augmented reality overlay over a planar object
 tracked by PlaneTracker from plane_tracker.py. solvePnP funciton is used to
 estimate the tracked object location in 3d space.

-video: http://www.youtube.com/watch?v=pzVbhxx6aog
+video: https://www.youtube.com/watch?v=pzVbhxx6aog

 Usage
 -----
@@ -8,7 +8,7 @@ Example of using features2d framework for interactive video homography matching.
 ORB features and FLANN matcher are used. This sample provides PlaneTracker class
 and an example of its usage.

-video: http://www.youtube.com/watch?v=pzVbhxx6aog
+video: https://www.youtube.com/watch?v=pzVbhxx6aog

 Usage
 -----
Author	SHA1	Message	Date
Alexander Alekhin	82f8176b06	Merge pull request #14676 from 103yiran:103yiran-clang	2019-05-30 14:19:14 +03:00
103yiran	a03b7575ba	make it compatible with clang	2019-05-30 14:04:40 +08:00
Alexander Alekhin	8bae39ce8a	Merge pull request #13669 from alalek:test_eigen_relax_eps_2.4	2019-01-21 18:32:22 +03:00
Alexander Alekhin	aa89881321	Merge pull request #13665 from alalek:backport_8294_flann_rand	2019-01-21 15:33:37 +03:00
Alexander Alekhin	30b01a2d29	core(test): relax eigen eps value: 0.01 -> 0.02	2019-01-21 15:18:10 +03:00
Alexander Alekhin	dc328451eb	flann: use OpenCV theRNG() std::rand() has no thread-safe guarantee. backport commit: `147f3ebf0a`	2019-01-21 13:45:08 +03:00
Alexander Alekhin	a49600cb24	Merge pull request #13510 from knsong:2.4	2018-12-25 13:08:02 +03:00
Kangning Song	0157ff0bc3	fix initial values bug	2018-12-23 11:29:42 +08:00
Alexander Alekhin	7e71666a0b	Merge pull request #13448 from alalek:issue_13445_2.4	2018-12-16 00:17:39 +03:00
Alexander Alekhin	7b677bb017	videoio(dc1394): use lazy initialization on demand backport `eb1f3733ee` into 2.4	2018-12-15 08:09:22 +00:00
Alexander Alekhin	9b954de175	Merge pull request #13265 from xusiwei:2.4	2018-11-25 12:12:09 +00:00
Siwei Xu	7f3af2b2d9	androidcamera: add a missing header. fix #11699 Missing std header	2018-11-24 01:31:50 +08:00
Alexander Alekhin	fd63c60418	Merge pull request #12181 from alalek:fix_gpu_sparse_multi_definition	2018-08-08 18:38:30 +03:00
Alexander Alekhin	1c34941537	Merge pull request #12179 from alalek:fix_gpu_samples_2.4	2018-08-08 18:37:12 +03:00
Alexander Alekhin	b0f0194595	gpu(sparse_multi): fix definition without TBB	2018-08-08 16:26:55 +03:00
Alexander Alekhin	8484c8af7c	samples(gpu): fix build (invalid access to cvconfig.h from sample) via find_package(OpenCV)	2018-08-08 16:26:18 +03:00
Alexander Alekhin	51cfa51924	OpenCV version++ 2.4.13.7	2018-07-02 15:41:56 +03:00
Alexander Alekhin	f1c5d8364f	Merge pull request #11640 from alalek:backport_11617	2018-05-31 18:28:49 +00:00
Alexander Alekhin	19f4c4403a	videoio(ffmpeg): specify stream->time_base backport `5128c1ff1f`	2018-05-31 18:34:43 +03:00
Alexander Alekhin	e89405d48f	Merge pull request #11534 from juanecito:2.4 Thread-safe version of sparse function in cv::gpu::PyrLKOpticalFlow 2	2018-05-31 15:29:14 +00:00
juanitov	eaf0b04530	Add sample of sparse pyrlk optical flow thread safe	2018-05-31 16:36:45 +03:00
juanitov	cbae431752	Fix HAVE_TBB is not defined in pyrlk.cu in spite of CMake TBB option is ON	2018-05-31 16:34:49 +03:00
Juan María Gómez López	0239c195d8	Merge pull request #11060 from juanecito:2.4 * Thread-safe version of sparse function in cv::gpu::PyrLKOpticalFlow class. The new function name is sparse_multi * Thread-safe sparse function in cv::gpu::PyrLKOpticalFlow. Tests * Thread-safe sparse function in cv::gpu::PyrLKOpticalFlow class. Add intel_TBB conditional compilation	2018-05-16 14:32:42 +03:00
Alexander Alekhin	a32aec5ba6	Merge pull request #11496 from shengyu7697:tab_to_space tab to space	2018-05-10 19:33:54 +03:00
shengyu	dd131219b2	tab to space	2018-05-10 23:14:04 +08:00
Alexander Alekhin	2cf58febf8	Merge pull request #11369 from ilovezfs:ffmpeg-4.0 Fix build with FFmpeg 4.0	2018-04-21 22:14:40 +03:00
ilovezfs	99091a6246	Fix build with FFmpeg 4.0 Backport of https://github.com/opencv/opencv/pull/10011.	2018-04-21 09:09:12 -07:00
Alexander Alekhin	0354d01e79	OpenCV version++ 2.4.13.6	2018-02-21 18:27:31 +00:00
Alexander Alekhin	45d3aac730	Merge pull request #10913 from alalek:fix_imgcodecs_hang_2.4	2018-02-21 16:39:59 +03:00
Alexander Alekhin	318ac6b8c9	imgcodecs: fix RBaseStream hang on truncated inputs `6e8241b78d`	2018-02-21 13:24:08 +03:00
Alexander Alekhin	7d332100a4	Merge pull request #10901 from alalek:backport_imgcodecs_fixes (2.4) backport imgcodecs fixes	2018-02-20 11:42:04 +03:00
Alexander Alekhin	56072c4406	imgcodecs: add more Jasper checks for supported and tested cases `435a3e337b`	2018-02-19 16:43:47 +03:00
Alexander Alekhin	cd64b504b8	imgcodecs: add overflow checks imgcodecs: remove assert() usage Origin commits: - `be5247921d` - `8a76fadaa3`	2018-02-19 16:38:08 +03:00
Alexander Alekhin	443059e371	imgcodecs(pxm): fix memcpy size `7bbe1a53cf`	2018-02-19 16:37:51 +03:00
Alexander Alekhin	8f9c4d23e0	Merge pull request #10695 from mworchel:blobdetector_mask Add mask support to SimpleBlobDetector	2018-01-30 21:00:46 +03:00
mworchel	8b90db3f25	Add mask support to SimpleBlobDetector	2018-01-25 13:34:12 +01:00
Alexander Alekhin	ec16307632	Merge pull request #10509 from kislinsk:support-msvc-14.1-minor-upgrades	2018-01-04 19:02:31 +03:00
Alexander Alekhin	fdefc4b09d	cmake: allow custom OpenCV_ARCH / OpenCV_RUNTIME values backport `8e21f808be`	2018-01-04 14:58:01 +00:00
Stefan Dinkelacker	61d8292652	cmake: add support for MSVC 14.1 minor upgrades	2018-01-04 15:39:19 +01:00
Alexander Alekhin	43f1b72e92	Merge pull request #10443 from alalek:backport_10435 (2.4 backport) Fixed #10433	2017-12-28 18:27:33 +03:00
Arthur Williams	1f4b8c2785	Fixed #10433 backport #10435	2017-12-28 02:32:43 +00:00
Alexander Alekhin	97eae6b546	OpenCV version++ 2.4.13.5	2017-12-19 23:11:33 +00:00
Alexander Alekhin	572d6d795d	Merge pull request #10301 from alalek:cmake_timestamp_update_2.4 (2.4) cmake: update timestamp status	2017-12-13 17:42:53 +03:00
Alexander Alekhin	7459388dd9	cmake: update timestamp status - avoid unnecessary rebuilding of OpenCV libraries - use timestamp of the first launch of CMake - to return to previous behavior use `-UOPENCV_TIMESTAMP` CMake option original commit: `3e45795232`	2017-12-13 17:05:33 +03:00
Alexander Alekhin	249edb2c99	Merge pull request #10204 from alalek:fix_calib3d_fisheye_rectify_test_2.4	2017-12-01 18:31:21 +03:00
Alexander Alekhin	00cc4aa230	calib3d: fix fisheye stereoRectify test - don't write into testdata directory - check matrices instead of result images	2017-11-30 16:41:37 +03:00
Alexander Alekhin	7c0193bb75	Merge pull request #10198 from alalek:fix_build_xcode_9.1_2.4 (2.4) Fix build with Xcode 9.1	2017-11-30 16:27:26 +03:00
Alexander Alekhin	97507e42a1	build: eliminate Xcode 9.1 warnings - 3rdparty/libjasper/jas_getopt.c:129:25: warning: this function declaration is not a prototype [-Wstrict-prototypes] - clang: warning: optimization flag '-fomit-frame-pointer' is not supported for target 'armv7' [-Wignored-optimization-argument] - clang: warning: optimization flag '-fomit-frame-pointer' is not supported for target 'armv7s' [-Wignored-optimization-argument] - backport IOS_ARCH toolchain option and setup CMAKE_SYSTEM_PROCESSOR: - adds '-mfpu=neon' compiler flag for ARM, SSE flags for i386/x86_64 - fixed proper libpng compilation (no ARM asm)	2017-11-30 14:24:05 +03:00
Alexander Alekhin	91fe01beca	Merge pull request #10052 from alalek:fix_macos_pkg_config_2.4 (2.4) cmake: fix pkg-config generation for MacOSX	2017-11-08 19:03:44 +00:00
Alexander Alekhin	591a08721e	cmake: fix pkg-config generation for MacOSX	2017-11-08 19:40:02 +03:00
Alexander Alekhin	a28733148e	Merge pull request #10038 from alalek:update_gitignore_2.4 (2.4) git: .gitignore update	2017-11-07 20:00:56 +00:00
Alexander Alekhin	e10e628036	git: .gitignore update ignore all "dot" files/directories by default. 'ignored' files can be added via 'git add -f' command if necessary.	2017-11-07 17:35:13 +03:00
Alexander Alekhin	738b388a64	Merge pull request #10030 from hosjiu1702:2.4 change from variance to standard deviation	2017-11-07 10:59:09 +00:00
Alexander Alekhin	9802024f71	Merge pull request #9974 from opalmirror:test_stereo_min_disparity_24 test_stereomatching.cpp: validate min disparity affect on valid ROI	2017-11-07 10:05:21 +00:00
hosjiu1702	1659e96766	change from variance to standard deviation	2017-11-06 11:42:24 +07:00
James Perkins	c1dea8465e	test_stereomatching.cpp: validate min disparity affect on valid ROI	2017-10-30 11:27:11 -07:00
Alexander Alekhin	678d383f26	Merge pull request #9844 from opalmirror:fix_stereobm_mindisp_truncation_24 fix StereoBM disparity map right margin truncation when minDisparities > 0	2017-10-30 18:48:48 +03:00
Alexander Alekhin	e15a56d142	Merge pull request #9915 from alalek:backport_9903 (2.4) Backport #9903: Fix out of bounds write	2017-10-26 12:19:38 +03:00
blendin	c46521ad65	Fix out of bounds write	2017-10-23 12:43:22 +03:00
Alexander Alekhin	d0f3468477	Merge pull request #9873 from alalek:fix_documentation_2.4 (2.4) doc: fix youtube videos handling	2017-10-17 21:40:34 +03:00
Alexander Alekhin	3224b0a75a	doc: fix youtube videos handling	2017-10-17 16:52:00 +03:00
James Perkins	7d8110772a	fix StereoBM disparity map right margin truncation when minDisparities > 0	2017-10-12 17:20:28 -07:00