OpenCV patch verison++.

Bug #4315 : fix CUDA bitwise operations with mask
(cherry picked from commit d87c30dc84)
2026-07-25 05:13:04 +04:00 · 2015-05-06 16:29:51 +03:00 · 2015-05-06 16:27:12 +03:00 · 2015-05-06 16:26:50 +03:00 · 2015-05-06 16:24:39 +03:00 · 2015-05-06 16:24:07 +03:00
49 changed files with 1048 additions and 541 deletions
@@ -605,7 +605,7 @@ if(INSTALL_TESTS AND OPENCV_TEST_DATA_PATH AND UNIX)
    configure_file("${CMAKE_CURRENT_SOURCE_DIR}/cmake/templates/opencv_run_all_tests_android.sh.in"
                   "${CMAKE_BINARY_DIR}/unix-install/opencv_run_all_tests.sh" @ONLY)
    install(PROGRAMS "${CMAKE_BINARY_DIR}/unix-install/opencv_run_all_tests.sh"
-            DESTINATION ${CMAKE_INSTALL_PREFIX} COMPONENT tests)
+            DESTINATION . COMPONENT tests)
  else()
    set(OPENCV_PYTHON_TESTS_LIST "")
    if(BUILD_opencv_python)
@@ -613,10 +613,6 @@ if(INSTALL_TESTS AND OPENCV_TEST_DATA_PATH AND UNIX)
      install(PROGRAMS ${py_tests} DESTINATION ${OPENCV_TEST_INSTALL_PATH} COMPONENT tests)
      set(OPENCV_PYTHON_TESTS_LIST "test2.py")
    endif()
-    configure_file("${CMAKE_CURRENT_SOURCE_DIR}/cmake/templates/opencv_testing.sh.in"
-                   "${CMAKE_BINARY_DIR}/unix-install/opencv_testing.sh" @ONLY)
-    install(FILES "${CMAKE_BINARY_DIR}/unix-install/opencv_testing.sh"
-            DESTINATION /etc/profile.d/ COMPONENT tests)
    configure_file("${CMAKE_CURRENT_SOURCE_DIR}/cmake/templates/opencv_run_all_tests_unix.sh.in"
                   "${CMAKE_BINARY_DIR}/unix-install/opencv_run_all_tests.sh" @ONLY)
    install(PROGRAMS "${CMAKE_BINARY_DIR}/unix-install/opencv_run_all_tests.sh"
@@ -638,11 +634,11 @@ endif()
 if(ANDROID OR NOT UNIX)
  install(FILES ${OPENCV_LICENSE_FILE}
        PERMISSIONS OWNER_READ GROUP_READ WORLD_READ
-        DESTINATION ${CMAKE_INSTALL_PREFIX} COMPONENT libs)
+        DESTINATION . COMPONENT libs)
  if(OPENCV_README_FILE)
    install(FILES ${OPENCV_README_FILE}
            PERMISSIONS OWNER_READ GROUP_READ WORLD_READ
-            DESTINATION ${CMAKE_INSTALL_PREFIX} COMPONENT libs)
+            DESTINATION . COMPONENT libs)
  endif()
 endif()

@@ -253,6 +253,11 @@ if(MSVC)
  endif()
 endif()

+if(MSVC12 AND NOT CMAKE_GENERATOR MATCHES "Visual Studio")
+  set(OPENCV_EXTRA_C_FLAGS "${OPENCV_EXTRA_C_FLAGS} /FS")
+  set(OPENCV_EXTRA_CXX_FLAGS "${OPENCV_EXTRA_CXX_FLAGS} /FS")
+endif()
+
 # Extra link libs if the user selects building static libs:
 if(NOT BUILD_SHARED_LIBS AND CMAKE_COMPILER_IS_GNUCXX AND NOT ANDROID)
  # Android does not need these settings because they are already set by toolchain file
@@ -140,6 +140,6 @@ if(WIN32)
    install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}staticlib" COMPONENT dev)
    install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}staticlib" FILE OpenCVModules${modules_file_suffix}.cmake COMPONENT dev)
  endif()
-  install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}" COMPONENT dev)
-  install(FILES "${OpenCV_SOURCE_DIR}/cmake/OpenCVConfig.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}/" COMPONENT dev)
+  install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" DESTINATION "." COMPONENT dev)
+  install(FILES "${OpenCV_SOURCE_DIR}/cmake/OpenCVConfig.cmake" DESTINATION "." COMPONENT dev)
 endif()
@@ -617,9 +617,14 @@ macro(ocv_create_module)

  ocv_install_target(${the_module} EXPORT OpenCVModules
    RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT libs
-    LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs
+    LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs NAMELINK_SKIP
    ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT dev
    )
+  get_target_property(_target_type ${the_module} TYPE)
+  if("${_target_type}" STREQUAL "SHARED_LIBRARY")
+    install(TARGETS ${the_module}
+      LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT dev NAMELINK_ONLY)
+  endif()

  # only "public" headers need to be installed
  if(OPENCV_MODULE_${the_module}_HEADERS AND ";${OPENCV_MODULES_PUBLIC};" MATCHES ";${the_module};")
@@ -860,7 +865,7 @@ function(ocv_add_samples)
    file(GLOB sample_files "${samples_path}/*")
    install(FILES ${sample_files}
            DESTINATION ${OPENCV_SAMPLES_SRC_INSTALL_PATH}/${module_id}
-            PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples)
+            PERMISSIONS OWNER_WRITE OWNER_READ GROUP_READ WORLD_READ COMPONENT samples)
  endif()
 endfunction()

@@ -14,7 +14,7 @@ machine perception in the commercial products. Being a BSD-licensed product,
 OpenCV makes it easy for businesses to utilize and modify the code.")
  set(CPACK_PACKAGE_VENDOR "OpenCV Foundation")
  set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE")
-  set(CPACK_PACKAGE_CONTACT "admin@opencv.org")
+  set(CPACK_PACKAGE_CONTACT "OpenCV Developers <admin@opencv.org>")
  set(CPACK_PACKAGE_VERSION_MAJOR "${OPENCV_VERSION_MAJOR}")
  set(CPACK_PACKAGE_VERSION_MINOR "${OPENCV_VERSION_MINOR}")
  set(CPACK_PACKAGE_VERSION_PATCH "${OPENCV_VERSION_PATCH}")
@@ -62,79 +62,115 @@ set(CPACK_DEBIAN_PACKAGE_PRIORITY "optional")
 set(CPACK_DEBIAN_PACKAGE_SECTION "libs")
 set(CPACK_DEBIAN_PACKAGE_HOMEPAGE "http://opencv.org")

+#display names
+set(CPACK_COMPONENT_DEV_DISPLAY_NAME     "Development files")
+set(CPACK_COMPONENT_DOCS_DISPLAY_NAME    "Documentation")
+set(CPACK_COMPONENT_JAVA_DISPLAY_NAME    "Java bindings")
+set(CPACK_COMPONENT_LIBS_DISPLAY_NAME    "Libraries and data")
+set(CPACK_COMPONENT_PYTHON_DISPLAY_NAME  "Python bindings")
+set(CPACK_COMPONENT_SAMPLES_DISPLAY_NAME "Samples")
+set(CPACK_COMPONENT_TESTS_DISPLAY_NAME   "Tests")
+
 #depencencies
 set(CPACK_DEBIAN_PACKAGE_SHLIBDEPS TRUE)
-set(CPACK_COMPONENT_samples_DEPENDS libs)
-set(CPACK_COMPONENT_dev_DEPENDS libs)
-set(CPACK_COMPONENT_docs_DEPENDS libs)
-set(CPACK_COMPONENT_java_DEPENDS libs)
-set(CPACK_COMPONENT_python_DEPENDS libs)
-set(CPACK_DEB_python_PACKAGE_DEPENDS "python${PYTHON_VERSION_MAJOR_MINOR}")
-set(CPACK_COMPONENT_tests_DEPENDS libs)
+set(CPACK_COMPONENT_LIBS_REQUIRED TRUE)
+set(CPACK_COMPONENT_SAMPLES_DEPENDS libs dev)
+set(CPACK_COMPONENT_DEV_DEPENDS libs)
+set(CPACK_COMPONENT_DOCS_DEPENDS libs)
+set(CPACK_COMPONENT_JAVA_DEPENDS libs)
+set(CPACK_COMPONENT_PYTHON_DEPENDS libs)
+set(CPACK_DEB_PYTHON_PACKAGE_DEPENDS "python-numpy (>=${PYTHON_NUMPY_VERSION}), python${PYTHON_VERSION_MAJOR_MINOR}")
+set(CPACK_COMPONENT_TESTS_DEPENDS libs)
 if (HAVE_opencv_python)
-  set(CPACK_DEB_tests_PACKAGE_DEPENDS "python${PYTHON_VERSION_MAJOR_MINOR}, python-py | python-pytest")
+  set(CPACK_DEB_TESTS_PACKAGE_DEPENDS "python-numpy (>=${PYTHON_NUMPY_VERSION}), python${PYTHON_VERSION_MAJOR_MINOR}, python-py | python-pytest")
 endif()

 if(HAVE_CUDA)
  string(REPLACE "." "-" cuda_version_suffix ${CUDA_VERSION})
  if(CUDA_VERSION VERSION_LESS "6.5")
-    set(CPACK_DEB_libs_PACKAGE_DEPENDS "cuda-core-libs-${cuda_version_suffix}, cuda-extra-libs-${cuda_version_suffix}")
-    set(CPACK_DEB_dev_PACKAGE_DEPENDS "cuda-headers-${cuda_version_suffix}")
+    set(CPACK_DEB_LIBS_PACKAGE_DEPENDS "cuda-core-libs-${cuda_version_suffix}, cuda-extra-libs-${cuda_version_suffix}")
+    set(CPACK_DEB_DEV_PACKAGE_DEPENDS "cuda-headers-${cuda_version_suffix}")
  else()
-    set(CPACK_DEB_libs_PACKAGE_DEPENDS "cuda-cudart-${cuda_version_suffix}, cuda-npp-${cuda_version_suffix}")
-    set(CPACK_DEB_dev_PACKAGE_DEPENDS "cuda-cudart-dev-${cuda_version_suffix}, cuda-npp-dev-${cuda_version_suffix}")
+    set(CPACK_DEB_LIBS_PACKAGE_DEPENDS "cuda-cudart-${cuda_version_suffix}, cuda-npp-${cuda_version_suffix}")
+    set(CPACK_DEB_DEV_PACKAGE_DEPENDS "cuda-cudart-dev-${cuda_version_suffix}, cuda-npp-dev-${cuda_version_suffix}")
    if(HAVE_CUFFT)
-      set(CPACK_DEB_libs_PACKAGE_DEPENDS "${CPACK_DEB_libs_PACKAGE_DEPENDS}, cuda-cufft-${cuda_version_suffix}")
-      set(CPACK_DEB_dev_PACKAGE_DEPENDS "${CPACK_DEB_dev_PACKAGE_DEPENDS}, cuda-cufft-dev-${cuda_version_suffix}")
+      set(CPACK_DEB_LIBS_PACKAGE_DEPENDS "${CPACK_DEB_LIBS_PACKAGE_DEPENDS}, cuda-cufft-${cuda_version_suffix}")
+      set(CPACK_DEB_DEV_PACKAGE_DEPENDS "${CPACK_DEB_DEV_PACKAGE_DEPENDS}, cuda-cufft-dev-${cuda_version_suffix}")
    endif()
    if(HAVE_HAVE_CUBLAS)
-      set(CPACK_DEB_libs_PACKAGE_DEPENDS "${CPACK_DEB_libs_PACKAGE_DEPENDS}, cuda-cublas-${cuda_version_suffix}")
-      set(CPACK_DEB_dev_PACKAGE_DEPENDS "${CPACK_DEB_dev_PACKAGE_DEPENDS}, cuda-cublas-dev-${cuda_version_suffix}")
+      set(CPACK_DEB_LIBS_PACKAGE_DEPENDS "${CPACK_DEB_LIBS_PACKAGE_DEPENDS}, cuda-cublas-${cuda_version_suffix}")
+      set(CPACK_DEB_DEV_PACKAGE_DEPENDS "${CPACK_DEB_DEV_PACKAGE_DEPENDS}, cuda-cublas-dev-${cuda_version_suffix}")
    endif()
  endif()
+endif()

-  set(CPACK_COMPONENT_dev_DEPENDS libs)
+if(HAVE_TBB AND NOT BUILD_TBB)
+  if(CPACK_DEB_DEV_PACKAGE_DEPENDS)
+    set(CPACK_DEB_DEV_PACKAGE_DEPENDS "${CPACK_DEB_DEV_PACKAGE_DEPENDS}, libtbb-dev")
+  else()
+    set(CPACK_DEB_DEV_PACKAGE_DEPENDS "libtbb-dev")
+  endif()
 endif()

 if(NOT OPENCV_CUSTOM_PACKAGE_INFO)
-  set(CPACK_COMPONENT_libs_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}")
-  set(CPACK_COMPONENT_libs_DESCRIPTION "Open Computer Vision Library")
-  set(CPACK_COMPONENT_libs_SECTION "libs")
+  set(CPACK_COMPONENT_LIBS_DESCRIPTION "Open Computer Vision Library")
+  set(CPACK_DEBIAN_COMPONENT_LIBS_NAME "libopencv")
+  set(CPACK_DEBIAN_COMPONENT_LIBS_SECTION "libs")

-  set(CPACK_COMPONENT_python_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-python")
-  set(CPACK_COMPONENT_python_DESCRIPTION "Python bindings for Open Source Computer Vision Library")
-  set(CPACK_COMPONENT_python_SECTION "python")
+  set(CPACK_COMPONENT_PYTHON_DESCRIPTION "Python bindings for Open Source Computer Vision Library")
+  set(CPACK_DEBIAN_COMPONENT_PYTHON_NAME "libopencv-python")
+  set(CPACK_DEBIAN_COMPONENT_PYTHON_SECTION "python")

-  set(CPACK_COMPONENT_java_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-java")
-  set(CPACK_COMPONENT_java_DESCRIPTION "Java bindings for Open Source Computer Vision Library")
-  set(CPACK_COMPONENT_java_SECTION "java")
+  set(CPACK_COMPONENT_JAVA_DESCRIPTION "Java bindings for Open Source Computer Vision Library")
+  set(CPACK_DEBIAN_COMPONENT_JAVA_NAME "libopencv-java")
+  set(CPACK_DEBIAN_COMPONENT_JAVA_SECTION "java")

-  set(CPACK_COMPONENT_dev_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-dev")
-  set(CPACK_COMPONENT_dev_DESCRIPTION "Development files for Open Source Computer Vision Library")
-  set(CPACK_COMPONENT_dev_SECTION "libdevel")
+  set(CPACK_COMPONENT_DEV_DESCRIPTION "Development files for Open Source Computer Vision Library")
+  set(CPACK_DEBIAN_COMPONENT_DEV_NAME "libopencv-dev")
+  set(CPACK_DEBIAN_COMPONENT_DEV_SECTION "libdevel")

-  set(CPACK_COMPONENT_docs_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-docs")
-  set(CPACK_COMPONENT_docs_DESCRIPTION "Documentation for Open Source Computer Vision Library")
-  set(CPACK_COMPONENT_docs_SECTION "doc")
+  set(CPACK_COMPONENT_DOCS_DESCRIPTION "Documentation for Open Source Computer Vision Library")
+  set(CPACK_DEBIAN_COMPONENT_DOCS_NAME "libopencv-docs")
+  set(CPACK_DEBIAN_COMPONENT_DOCS_SECTION "doc")

-  set(CPACK_COMPONENT_samples_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-samples")
-  set(CPACK_COMPONENT_samples_DESCRIPTION "Samples for Open Source Computer Vision Library")
-  set(CPACK_COMPONENT_samples_SECTION "devel")
+  set(CPACK_COMPONENT_SAMPLES_DESCRIPTION "Samples for Open Source Computer Vision Library")
+  set(CPACK_DEBIAN_COMPONENT_SAMPLES_NAME "libopencv-samples")
+  set(CPACK_DEBIAN_COMPONENT_SAMPLES_SECTION "devel")

-  set(CPACK_COMPONENT_tests_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-tests")
-  set(CPACK_COMPONENT_tests_DESCRIPTION "Accuracy and performance tests for Open Source Computer Vision Library")
-  set(CPACK_COMPONENT_tests_SECTION "misc")
+  set(CPACK_COMPONENT_TESTS_DESCRIPTION "Accuracy and performance tests for Open Source Computer Vision Library")
+  set(CPACK_DEBIAN_COMPONENT_TESTS_NAME "libopencv-tests")
+  set(CPACK_DEBIAN_COMPONENT_TESTS_SECTION "misc")
 endif(NOT OPENCV_CUSTOM_PACKAGE_INFO)

-if(NOT OPENCV_CUSTOM_PACKAGE_LAYOUT)
-  set(CPACK_libs_COMPONENT_INSTALL TRUE)
-  set(CPACK_dev_COMPONENT_INSTALL TRUE)
-  set(CPACK_docs_COMPONENT_INSTALL TRUE)
-  set(CPACK_python_COMPONENT_INSTALL TRUE)
-  set(CPACK_java_COMPONENT_INSTALL TRUE)
-  set(CPACK_samples_COMPONENT_INSTALL TRUE)
-endif(NOT OPENCV_CUSTOM_PACKAGE_LAYOUT)
+if(CPACK_GENERATOR STREQUAL "DEB")
+  find_program(GZIP_TOOL NAMES "gzip" PATHS "/bin" "/usr/bin" "/usr/local/bin")
+  if(NOT GZIP_TOOL)
+    message(FATAL_ERROR "Unable to find 'gzip' program")
+  endif()
+
+  execute_process(COMMAND "date" "-R"
+                  OUTPUT_VARIABLE CHANGELOG_PACKAGE_DATE
+                  OUTPUT_STRIP_TRAILING_WHITESPACE)
+
+  set(CHANGELOG_PACKAGE_VERSION "${CPACK_PACKAGE_VERSION}")
+  set(ALL_COMPONENTS "libs" "dev" "docs" "python" "java" "samples" "tests")
+  foreach (comp ${ALL_COMPONENTS})
+    string(TOUPPER "${comp}" comp_upcase)
+    set(DEBIAN_CHANGELOG_OUT_FILE    "${CMAKE_BINARY_DIR}/deb-packages-gen/${comp}/changelog.Debian")
+    set(DEBIAN_CHANGELOG_OUT_FILE_GZ "${CMAKE_BINARY_DIR}/deb-packages-gen/${comp}/changelog.Debian.gz")
+    set(CHANGELOG_PACKAGE_NAME "${CPACK_DEBIAN_COMPONENT_${comp_upcase}_NAME}")
+    configure_file("${CMAKE_CURRENT_SOURCE_DIR}/cmake/templates/changelog.Debian.in" "${DEBIAN_CHANGELOG_OUT_FILE}" @ONLY)
+
+    execute_process(COMMAND "${GZIP_TOOL}" "-cf9" "${DEBIAN_CHANGELOG_OUT_FILE}"
+                    OUTPUT_FILE "${DEBIAN_CHANGELOG_OUT_FILE_GZ}"
+                    WORKING_DIRECTORY "${CMAKE_BINARY_DIR}")
+
+    install(FILES "${DEBIAN_CHANGELOG_OUT_FILE_GZ}"
+            DESTINATION "share/doc/${CPACK_DEBIAN_COMPONENT_${comp_upcase}_NAME}"
+            COMPONENT "${comp}")
+  endforeach()
+endif()

 include(CPack)

-ENDif(EXISTS "${CMAKE_ROOT}/Modules/CPack.cmake")
+ENDif(EXISTS "${CMAKE_ROOT}/Modules/CPack.cmake")
@@ -0,0 +1,4 @@
+@CHANGELOG_PACKAGE_NAME@ (@CPACK_PACKAGE_VERSION@) unstable; urgency=low
+  * Debian changelog stub. See upstream changelog or release notes in user
+     documentation for more details.
+ -- @CPACK_PACKAGE_CONTACT@  @CHANGELOG_PACKAGE_DATE@
@@ -1,2 +0,0 @@
-# Environment setup for OpenCV testing
-export OPENCV_TEST_DATA_PATH=@CMAKE_INSTALL_PREFIX@/share/OpenCV/testdata
@@ -196,18 +196,16 @@ namespace cv
                }
            }

+            resultsMutex.lock();
            if (localInliers.size() > inliers.size())
            {
-                resultsMutex.lock();
-
                inliers.clear();
                inliers.resize(localInliers.size());
                memcpy(&inliers[0], &localInliers[0], sizeof(int) * localInliers.size());
                localRvec.copyTo(rvec);
                localTvec.copyTo(tvec);
-
-                resultsMutex.unlock();
            }
+            resultsMutex.unlock();
        }

        static void pnpTask(const vector<char>& pointsMask, const Mat& objectPoints, const Mat& imagePoints,
@@ -50,7 +50,7 @@
 #define CV_VERSION_EPOCH    2
 #define CV_VERSION_MAJOR    4
 #define CV_VERSION_MINOR    10
-#define CV_VERSION_REVISION 0
+#define CV_VERSION_REVISION 3

 #define CVAUX_STR_EXP(__A)  #__A
 #define CVAUX_STR(__A)      CVAUX_STR_EXP(__A)
@@ -2683,7 +2683,6 @@ CV_IMPL CvFileStorage*
 cvOpenFileStorage( const char* filename, CvMemStorage* dststorage, int flags, const char* encoding )
 {
    CvFileStorage* fs = 0;
-    char* xml_buf = 0;
    int default_block_size = 1 << 18;
    bool append = (flags & 3) == CV_STORAGE_APPEND;
    bool mem = (flags & CV_STORAGE_MEMORY) != 0;
@@ -2815,7 +2814,7 @@ cvOpenFileStorage( const char* filename, CvMemStorage* dststorage, int flags, co
                int last_occurence = -1;
                xml_buf_size = MIN(xml_buf_size, int(file_size));
                fseek( fs->file, -xml_buf_size, SEEK_END );
-                xml_buf = (char*)cvAlloc( xml_buf_size+2 );
+                char* xml_buf = (char*)cvAlloc( xml_buf_size+2 );
                // find the last occurence of </opencv_storage>
                for(;;)
                {
@@ -2833,6 +2832,7 @@ cvOpenFileStorage( const char* filename, CvMemStorage* dststorage, int flags, co
                        ptr += strlen(substr);
                    }
                }
+                cvFree( &xml_buf );
                if( last_occurence < 0 )
                    CV_Error( CV_StsError, "Could not find </opencv_storage> in the end of file.\n" );
                icvCloseFile( fs );
@@ -2936,7 +2936,6 @@ _exit_:
        }
    }

-    cvFree( &xml_buf );
    return  fs;
 }

@@ -44,6 +44,7 @@
 #include "opencv2/gpu/device/transform.hpp"
 #include "opencv2/gpu/device/functional.hpp"
 #include "opencv2/gpu/device/type_traits.hpp"
+#include "opencv2/gpu/device/vec_traits.hpp"

 namespace cv { namespace gpu { namespace device
 {
@@ -105,87 +106,59 @@ namespace cv { namespace gpu { namespace device
    ////////////////////////////////// SetTo //////////////////////////////////
    ///////////////////////////////////////////////////////////////////////////

-    __constant__ uchar scalar_8u[4];
-    __constant__ schar scalar_8s[4];
-    __constant__ ushort scalar_16u[4];
-    __constant__ short scalar_16s[4];
-    __constant__ int scalar_32s[4];
-    __constant__ float scalar_32f[4];
-    __constant__ double scalar_64f[4];
-
-    template <typename T> __device__ __forceinline__ T readScalar(int i);
-    template <> __device__ __forceinline__ uchar readScalar<uchar>(int i) {return scalar_8u[i];}
-    template <> __device__ __forceinline__ schar readScalar<schar>(int i) {return scalar_8s[i];}
-    template <> __device__ __forceinline__ ushort readScalar<ushort>(int i) {return scalar_16u[i];}
-    template <> __device__ __forceinline__ short readScalar<short>(int i) {return scalar_16s[i];}
-    template <> __device__ __forceinline__ int readScalar<int>(int i) {return scalar_32s[i];}
-    template <> __device__ __forceinline__ float readScalar<float>(int i) {return scalar_32f[i];}
-    template <> __device__ __forceinline__ double readScalar<double>(int i) {return scalar_64f[i];}
-
-    void writeScalar(const uchar* vals)
-    {
-        cudaSafeCall( cudaMemcpyToSymbol(scalar_8u, vals, sizeof(uchar) * 4) );
-    }
-    void writeScalar(const schar* vals)
-    {
-        cudaSafeCall( cudaMemcpyToSymbol(scalar_8s, vals, sizeof(schar) * 4) );
-    }
-    void writeScalar(const ushort* vals)
-    {
-        cudaSafeCall( cudaMemcpyToSymbol(scalar_16u, vals, sizeof(ushort) * 4) );
-    }
-    void writeScalar(const short* vals)
-    {
-        cudaSafeCall( cudaMemcpyToSymbol(scalar_16s, vals, sizeof(short) * 4) );
-    }
-    void writeScalar(const int* vals)
-    {
-        cudaSafeCall( cudaMemcpyToSymbol(scalar_32s, vals, sizeof(int) * 4) );
-    }
-    void writeScalar(const float* vals)
-    {
-        cudaSafeCall( cudaMemcpyToSymbol(scalar_32f, vals, sizeof(float) * 4) );
-    }
-    void writeScalar(const double* vals)
-    {
-        cudaSafeCall( cudaMemcpyToSymbol(scalar_64f, vals, sizeof(double) * 4) );
-    }
-
    template<typename T>
-    __global__ void set_to_without_mask(T* mat, int cols, int rows, size_t step, int channels)
+    __global__ void set_to_without_mask(PtrStepSz<T> mat, typename TypeVec<T, 4>::vec_type val, int channels)
    {
-        size_t x = blockIdx.x * blockDim.x + threadIdx.x;
-        size_t y = blockIdx.y * blockDim.y + threadIdx.y;
+        const int y = blockIdx.x * blockDim.y + threadIdx.y;

-        if ((x < cols * channels ) && (y < rows))
+        if (y < mat.rows)
        {
-            size_t idx = y * ( step >> shift_and_sizeof<T>::shift ) + x;
-            mat[idx] = readScalar<T>(x % channels);
+            const T vals[] = {
+                val.x, val.y, val.z, val.w
+            };
+
+            T* row = mat.ptr(y);
+
+            for (int x = threadIdx.x; x < mat.cols * channels; x += 32)
+            {
+                row[x] = vals[x % channels];
+            }
        }
    }

    template<typename T>
-    __global__ void set_to_with_mask(T* mat, const uchar* mask, int cols, int rows, size_t step, int channels, size_t step_mask)
+    __global__ void set_to_with_mask(PtrStepSz<T> mat, const PtrStepb mask, typename TypeVec<T, 4>::vec_type val, int channels)
    {
-        size_t x = blockIdx.x * blockDim.x + threadIdx.x;
-        size_t y = blockIdx.y * blockDim.y + threadIdx.y;
+        const int y = blockIdx.x * blockDim.y + threadIdx.y;

-        if ((x < cols * channels ) && (y < rows))
-            if (mask[y * step_mask + x / channels] != 0)
+        if (y < mat.rows)
+        {
+            const T vals[] = {
+                val.x, val.y, val.z, val.w
+            };
+
+            T* row = mat.ptr(y);
+            const uchar* mask_row = mask.ptr(y);
+
+            for (int x = threadIdx.x; x < mat.cols * channels; x += 32)
            {
-                size_t idx = y * ( step >> shift_and_sizeof<T>::shift ) + x;
-                mat[idx] = readScalar<T>(x % channels);
+                if (mask_row[x / channels])
+                {
+                    row[x] = vals[x % channels];
+                }
            }
+        }
    }
+
    template <typename T>
    void set_to_gpu(PtrStepSzb mat, const T* scalar, PtrStepSzb mask, int channels, cudaStream_t stream)
    {
-        writeScalar(scalar);
+        typedef typename TypeVec<T, 4>::vec_type vec_type;

-        dim3 threadsPerBlock(32, 8, 1);
-        dim3 numBlocks (mat.cols * channels / threadsPerBlock.x + 1, mat.rows / threadsPerBlock.y + 1, 1);
+        dim3 block(32, 8);
+        dim3 grid(divUp(mat.rows, block.y));

-        set_to_with_mask<T><<<numBlocks, threadsPerBlock, 0, stream>>>((T*)mat.data, (uchar*)mask.data, mat.cols, mat.rows, mat.step, channels, mask.step);
+        set_to_with_mask<T><<<grid, block, 0, stream>>>(PtrStepSz<T>(mat), mask, VecTraits<vec_type>::make(scalar), channels);
        cudaSafeCall( cudaGetLastError() );

        if (stream == 0)
@@ -203,12 +176,12 @@ namespace cv { namespace gpu { namespace device
    template <typename T>
    void set_to_gpu(PtrStepSzb mat, const T* scalar, int channels, cudaStream_t stream)
    {
-        writeScalar(scalar);
+        typedef typename TypeVec<T, 4>::vec_type vec_type;

-        dim3 threadsPerBlock(32, 8, 1);
-        dim3 numBlocks (mat.cols * channels / threadsPerBlock.x + 1, mat.rows / threadsPerBlock.y + 1, 1);
+        dim3 block(32, 8);
+        dim3 grid(divUp(mat.rows, block.y));

-        set_to_without_mask<T><<<numBlocks, threadsPerBlock, 0, stream>>>((T*)mat.data, mat.cols, mat.rows, mat.step, channels);
+        set_to_without_mask<T><<<grid, block, 0, stream>>>(PtrStepSz<T>(mat), VecTraits<vec_type>::make(scalar), channels);
        cudaSafeCall( cudaGetLastError() );

        if (stream == 0)
@@ -626,12 +626,12 @@ namespace cv { namespace gpu { namespace device

        __device__ __forceinline__ int idx_row_low(int y) const
        {
-            return (y >= 0) * y + (y < 0) * (y - ((y - height + 1) / height) * height);
+            return (y >= 0) ? y : (y - ((y - height + 1) / height) * height);
        }

        __device__ __forceinline__ int idx_row_high(int y) const
        {
-            return (y < height) * y + (y >= height) * (y % height);
+            return (y < height) ? y : (y % height);
        }

        __device__ __forceinline__ int idx_row(int y) const
@@ -641,12 +641,12 @@ namespace cv { namespace gpu { namespace device

        __device__ __forceinline__ int idx_col_low(int x) const
        {
-            return (x >= 0) * x + (x < 0) * (x - ((x - width + 1) / width) * width);
+            return (x >= 0) ? x : (x - ((x - width + 1) / width) * width);
        }

        __device__ __forceinline__ int idx_col_high(int x) const
        {
-            return (x < width) * x + (x >= width) * (x % width);
+            return (x < width) ? x : (x % width);
        }

        __device__ __forceinline__ int idx_col(int x) const
@@ -99,7 +99,8 @@ PERF_TEST_P(Sz_Depth_Cn_KernelSz, Denoising_BilateralFilter,

 DEF_PARAM_TEST(Sz_Depth_Cn_WinSz_BlockSz, cv::Size, MatDepth, MatCn, int, int);

-PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, Denoising_NonLocalMeans,
+// disabled, since it takes too much time
+PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, DISABLED_Denoising_NonLocalMeans,
            Combine(GPU_DENOISING_IMAGE_SIZES,
                    Values<MatDepth>(CV_8U),
                    GPU_CHANNELS_1_3,
@@ -143,7 +144,8 @@ PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, Denoising_NonLocalMeans,

 DEF_PARAM_TEST(Sz_Depth_Cn_WinSz_BlockSz, cv::Size, MatDepth, MatCn, int, int);

-PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, Denoising_FastNonLocalMeans,
+// disabled, since it takes too much time
+PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, DISABLED_Denoising_FastNonLocalMeans,
            Combine(GPU_DENOISING_IMAGE_SIZES,
                    Values<MatDepth>(CV_8U),
                    GPU_CHANNELS_1_3,
@@ -123,7 +123,7 @@ PERF_TEST_P(Image_NFeatures, Features2D_ORB,

        sortKeyPoints(gpu_keypoints, gpu_descriptors);

-        SANITY_CHECK_KEYPOINTS(gpu_keypoints);
+        SANITY_CHECK_KEYPOINTS(gpu_keypoints, 1e-10);
        SANITY_CHECK(gpu_descriptors);
    }
    else
@@ -908,7 +908,7 @@ CV_ENUM(TemplateMethod, TM_SQDIFF, TM_SQDIFF_NORMED, TM_CCORR, TM_CCORR_NORMED,

 DEF_PARAM_TEST(Sz_TemplateSz_Cn_Method, cv::Size, cv::Size, MatCn, TemplateMethod);

-PERF_TEST_P(Sz_TemplateSz_Cn_Method, ImgProc_MatchTemplate8U,
+PERF_TEST_P(Sz_TemplateSz_Cn_Method, DISABLED_ImgProc_MatchTemplate8U,
            Combine(GPU_TYPICAL_MAT_SIZES,
                    Values(cv::Size(5, 5), cv::Size(16, 16), cv::Size(30, 30)),
                    GPU_CHANNELS_1_3_4,
@@ -948,7 +948,7 @@ PERF_TEST_P(Sz_TemplateSz_Cn_Method, ImgProc_MatchTemplate8U,
 ////////////////////////////////////////////////////////////////////////////////
 // MatchTemplate32F

-PERF_TEST_P(Sz_TemplateSz_Cn_Method, ImgProc_MatchTemplate32F,
+PERF_TEST_P(Sz_TemplateSz_Cn_Method, DISABLED_ImgProc_MatchTemplate32F,
            Combine(GPU_TYPICAL_MAT_SIZES,
                    Values(cv::Size(5, 5), cv::Size(16, 16), cv::Size(30, 30)),
                    GPU_CHANNELS_1_3_4,
@@ -1011,7 +1011,7 @@ PERF_TEST_P(Sz_Flags, ImgProc_MulSpectrums,

        TEST_CYCLE() cv::gpu::mulSpectrums(d_a, d_b, dst, flag);

-        GPU_SANITY_CHECK(dst);
+        GPU_SANITY_CHECK(dst, 2);
    }
    else
    {
@@ -1045,7 +1045,7 @@ PERF_TEST_P(Sz, ImgProc_MulAndScaleSpectrums,

        TEST_CYCLE() cv::gpu::mulAndScaleSpectrums(d_src1, d_src2, dst, cv::DFT_ROWS, scale, false);

-        GPU_SANITY_CHECK(dst);
+        GPU_SANITY_CHECK(dst, 1e-5);
    }
    else
    {
@@ -1413,7 +1413,7 @@ PERF_TEST_P(Sz_Depth_Code, ImgProc_CvtColor,

        TEST_CYCLE() cv::gpu::cvtColor(d_src, dst, info.code, info.dcn);

-        GPU_SANITY_CHECK(dst, 1e-4);
+        GPU_SANITY_CHECK(dst, 1e-2);
    }
    else
    {
@@ -1609,7 +1609,7 @@ PERF_TEST_P(Sz_Depth_Cn, ImgProc_ImagePyramidBuild,
        cv::gpu::GpuMat dst;
        d_pyr.getLayer(dst, dstSize);

-        GPU_SANITY_CHECK(dst);
+        GPU_SANITY_CHECK(dst, 1e-3);
    }
    else
    {
@@ -1646,7 +1646,7 @@ PERF_TEST_P(Sz_Depth_Cn, ImgProc_ImagePyramidGetLayer,

        TEST_CYCLE() d_pyr.getLayer(dst, dstSize);

-        GPU_SANITY_CHECK(dst);
+        GPU_SANITY_CHECK(dst, 1e-3);
    }
    else
    {
@@ -143,7 +143,7 @@ PERF_TEST_P(ImagePair, Video_CreateOpticalFlowNeedleMap,

        TEST_CYCLE() cv::gpu::createOpticalFlowNeedleMap(u, v, vertex, colors);

-        GPU_SANITY_CHECK(vertex, 1e-6);
+        GPU_SANITY_CHECK(vertex, 1e-5);
        GPU_SANITY_CHECK(colors);
    }
    else
@@ -307,7 +307,8 @@ PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, Video_PyrLKOpticalFlowSparse

 DEF_PARAM_TEST(ImagePair_WinSz_Levels_Iters, pair_string, int, int, int);

-PERF_TEST_P(ImagePair_WinSz_Levels_Iters, Video_PyrLKOpticalFlowDense,
+// Sanity test fails on Maxwell and CUDA 7.0
+PERF_TEST_P(ImagePair_WinSz_Levels_Iters, DISABLED_Video_PyrLKOpticalFlowDense,
            Combine(Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")),
                    Values(3, 5, 7, 9, 13, 17, 21),
                    Values(1, 3),
@@ -340,8 +341,8 @@ PERF_TEST_P(ImagePair_WinSz_Levels_Iters, Video_PyrLKOpticalFlowDense,

        TEST_CYCLE() d_pyrLK.dense(d_frame0, d_frame1, u, v);

-        GPU_SANITY_CHECK(u);
-        GPU_SANITY_CHECK(v);
+        GPU_SANITY_CHECK(u, 0.5);
+        GPU_SANITY_CHECK(v, 0.5);
    }
    else
    {
@@ -463,7 +464,8 @@ void calcOpticalFlowBM(const cv::Mat& prev, const cv::Mat& curr,
    cvCalcOpticalFlowBM(&cvprev, &cvcurr, bSize, shiftSize, maxRange, usePrevious, &cvvelx, &cvvely);
 }

-PERF_TEST_P(ImagePair, Video_OpticalFlowBM,
+// disabled, since it takes too much time
+PERF_TEST_P(ImagePair, DISABLED_Video_OpticalFlowBM,
            Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
 {
    declare.time(400);
@@ -541,7 +543,8 @@ PERF_TEST_P(ImagePair, DISABLED_Video_FastOpticalFlowBM,

 DEF_PARAM_TEST_1(Video, string);

-PERF_TEST_P(Video, Video_FGDStatModel,
+// disabled, since it takes too much time
+PERF_TEST_P(Video, DISABLED_Video_FGDStatModel,
            Values(string("gpu/video/768x576.avi")))
 {
    const int numIters = 10;
@@ -1577,7 +1577,7 @@ namespace

    void rgba_to_mbgra(const GpuMat& src, GpuMat& dst, int, Stream& st)
    {
-    #if (CUDA_VERSION < 5000)
+    #if (CUDART_VERSION < 5000)
        (void)src;
        (void)dst;
        (void)st;
@@ -1947,7 +1947,7 @@ void cv::gpu::swapChannels(GpuMat& image, const int dstOrder[4], Stream& s)

 void cv::gpu::gammaCorrection(const GpuMat& src, GpuMat& dst, bool forward, Stream& stream)
 {
-#if (CUDA_VERSION < 5000)
+#if (CUDART_VERSION < 5000)
    (void)src;
    (void)dst;
    (void)forward;
@@ -42,7 +42,7 @@

 #include "cu_safe_call.h"

-#ifdef HAVE_CUDA
+#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)

 namespace
 {
@@ -45,7 +45,7 @@

 #include "precomp.hpp"

-#ifdef HAVE_CUDA
+#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)

 namespace cv { namespace gpu {
    namespace detail
@@ -374,6 +374,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void matchUnrolledCached(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, int2* bestTrainIdx, float2* bestDistance)
        {
            extern __shared__ int smem[];
@@ -424,6 +425,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void matchUnrolledCached(const PtrStepSz<T> query, const PtrStepSz<T>* trains, int n, const Mask mask, int2* bestTrainIdx, int2* bestImgIdx, float2* bestDistance)
        {
            extern __shared__ int smem[];
@@ -553,6 +555,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void matchUnrolled(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, int2* bestTrainIdx, float2* bestDistance)
        {
            extern __shared__ int smem[];
@@ -601,6 +604,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void matchUnrolled(const PtrStepSz<T> query, const PtrStepSz<T>* trains, int n, const Mask mask, int2* bestTrainIdx, int2* bestImgIdx, float2* bestDistance)
        {
            extern __shared__ int smem[];
@@ -727,6 +731,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void match(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, int2* bestTrainIdx, float2* bestDistance)
        {
            extern __shared__ int smem[];
@@ -775,6 +780,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void match(const PtrStepSz<T> query, const PtrStepSz<T>* trains, int n, const Mask mask, int2* bestTrainIdx, int2* bestImgIdx, float2* bestDistance)
        {
            extern __shared__ int smem[];
@@ -902,6 +908,7 @@ namespace cv { namespace gpu { namespace device
        // Calc distance kernel

        template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void calcDistanceUnrolled(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, PtrStepf allDist)
        {
            extern __shared__ int smem[];
@@ -966,6 +973,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void calcDistance(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, PtrStepf allDist)
        {
            extern __shared__ int smem[];
@@ -1066,6 +1074,7 @@ namespace cv { namespace gpu { namespace device
        // find knn match kernel

        template <int BLOCK_SIZE>
+        __launch_bounds__(BLOCK_SIZE)
        __global__ void findBestMatch(PtrStepSzf allDist, int i, PtrStepi trainIdx, PtrStepf distance)
        {
            const int SMEM_SIZE = BLOCK_SIZE > 64 ? BLOCK_SIZE : 64;
@@ -136,6 +136,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void matchUnrolledCached(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, int* bestTrainIdx, float* bestDistance)
        {
            extern __shared__ int smem[];
@@ -184,6 +185,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void matchUnrolledCached(const PtrStepSz<T> query, const PtrStepSz<T>* trains, int n, const Mask mask,
                                            int* bestTrainIdx, int* bestImgIdx, float* bestDistance)
        {
@@ -296,6 +298,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void matchUnrolled(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, int* bestTrainIdx, float* bestDistance)
        {
            extern __shared__ int smem[];
@@ -342,6 +345,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void matchUnrolled(const PtrStepSz<T> query, const PtrStepSz<T>* trains, int n, const Mask mask,
                                      int* bestTrainIdx, int* bestImgIdx, float* bestDistance)
        {
@@ -451,6 +455,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void match(const PtrStepSz<T> query, const PtrStepSz<T> train, const Mask mask, int* bestTrainIdx, float* bestDistance)
        {
            extern __shared__ int smem[];
@@ -497,6 +502,7 @@ namespace cv { namespace gpu { namespace device
        }

        template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void match(const PtrStepSz<T> query, const PtrStepSz<T>* trains, int n, const Mask mask,
                              int* bestTrainIdx, int* bestImgIdx, float* bestDistance)
        {
@@ -56,6 +56,7 @@ namespace cv { namespace gpu { namespace device
        // Match Unrolled

        template <int BLOCK_SIZE, int MAX_DESC_LEN, bool SAVE_IMG_IDX, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void matchUnrolled(const PtrStepSz<T> query, int imgIdx, const PtrStepSz<T> train, float maxDistance, const Mask mask,
            PtrStepi bestTrainIdx, PtrStepi bestImgIdx, PtrStepf bestDistance, unsigned int* nMatches, int maxCount)
        {
@@ -164,6 +165,7 @@ namespace cv { namespace gpu { namespace device
        // Match

        template <int BLOCK_SIZE, bool SAVE_IMG_IDX, typename Dist, typename T, typename Mask>
+        __launch_bounds__(BLOCK_SIZE * BLOCK_SIZE)
        __global__ void match(const PtrStepSz<T> query, int imgIdx, const PtrStepSz<T> train, float maxDistance, const Mask mask,
            PtrStepi bestTrainIdx, PtrStepi bestImgIdx, PtrStepf bestDistance, unsigned int* nMatches, int maxCount)
        {
@@ -1896,53 +1896,53 @@ namespace cv { namespace gpu { namespace device

 namespace arithm
 {
-    template <typename T> void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream)
+    template <typename T> void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream)
    {
        if (mask.data)
-            transform((PtrStepSz<T>) src, (PtrStepSz<T>) dst, bit_not<T>(), mask, stream);
+            transform((PtrStepSz<T>) src, (PtrStepSz<T>) dst, bit_not<T>(), SingleMaskChannels(mask, num_channels), stream);
        else
            transform((PtrStepSz<T>) src, (PtrStepSz<T>) dst, bit_not<T>(), WithOutMask(), stream);
    }

-    template <typename T> void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream)
+    template <typename T> void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream)
    {
        if (mask.data)
-            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_and<T>(), mask, stream);
+            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_and<T>(), SingleMaskChannels(mask, num_channels), stream);
        else
            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_and<T>(), WithOutMask(), stream);
    }

-    template <typename T> void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream)
+    template <typename T> void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream)
    {
        if (mask.data)
-            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_or<T>(), mask, stream);
+            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_or<T>(), SingleMaskChannels(mask, num_channels), stream);
        else
            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_or<T>(), WithOutMask(), stream);
    }

-    template <typename T> void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream)
+    template <typename T> void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream)
    {
        if (mask.data)
-            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_xor<T>(), mask, stream);
+            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_xor<T>(), SingleMaskChannels(mask, num_channels), stream);
        else
            transform((PtrStepSz<T>) src1, (PtrStepSz<T>) src2, (PtrStepSz<T>) dst, bit_xor<T>(), WithOutMask(), stream);
    }

-    template void bitMatNot<uchar>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template void bitMatNot<ushort>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template void bitMatNot<uint>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatNot<uchar>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template void bitMatNot<ushort>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template void bitMatNot<uint>(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);

-    template void bitMatAnd<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template void bitMatAnd<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template void bitMatAnd<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatAnd<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template void bitMatAnd<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template void bitMatAnd<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);

-    template void bitMatOr<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template void bitMatOr<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template void bitMatOr<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatOr<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template void bitMatOr<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template void bitMatOr<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);

-    template void bitMatXor<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template void bitMatXor<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template void bitMatXor<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template void bitMatXor<uchar>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template void bitMatXor<ushort>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template void bitMatXor<uint>(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
 }

 //////////////////////////////////////////////////////////////////////////////////////
@@ -114,7 +114,7 @@ namespace cv { namespace gpu { namespace device

                    sum =       0.0625f * src(b.idx_row_low (src_y - 2), b.idx_col_high(x));
                    sum = sum + 0.25f   * src(b.idx_row_low (src_y - 1), b.idx_col_high(x));
-                    sum = sum + 0.375f  * src(src_y                    , b.idx_col_high(x));
+                    sum = sum + 0.375f  * src(b.idx_row_high(src_y    ), b.idx_col_high(x));
                    sum = sum + 0.25f   * src(b.idx_row_high(src_y + 1), b.idx_col_high(x));
                    sum = sum + 0.0625f * src(b.idx_row_high(src_y + 2), b.idx_col_high(x));

@@ -129,7 +129,7 @@ namespace cv { namespace gpu { namespace device

                    sum =       0.0625f * src(b.idx_row_low (src_y - 2), b.idx_col(left_x));
                    sum = sum + 0.25f   * src(b.idx_row_low (src_y - 1), b.idx_col(left_x));
-                    sum = sum + 0.375f  * src(src_y                    , b.idx_col(left_x));
+                    sum = sum + 0.375f  * src(b.idx_row_high(src_y    ), b.idx_col(left_x));
                    sum = sum + 0.25f   * src(b.idx_row_high(src_y + 1), b.idx_col(left_x));
                    sum = sum + 0.0625f * src(b.idx_row_high(src_y + 2), b.idx_col(left_x));

@@ -144,7 +144,7 @@ namespace cv { namespace gpu { namespace device

                    sum =       0.0625f * src(b.idx_row_low (src_y - 2), b.idx_col_high(right_x));
                    sum = sum + 0.25f   * src(b.idx_row_low (src_y - 1), b.idx_col_high(right_x));
-                    sum = sum + 0.375f  * src(src_y                    , b.idx_col_high(right_x));
+                    sum = sum + 0.375f  * src(b.idx_row_high(src_y    ), b.idx_col_high(right_x));
                    sum = sum + 0.25f   * src(b.idx_row_high(src_y + 1), b.idx_col_high(right_x));
                    sum = sum + 0.0625f * src(b.idx_row_high(src_y + 2), b.idx_col_high(right_x));

@@ -103,16 +103,22 @@ namespace cv { namespace gpu { namespace device
        {
            static __device__ __forceinline__ float compute(const uchar* left, const uchar* right)
            {
-                return fmin(cdata_weight * ::abs((int)*left - *right), cdata_weight * cmax_data_term);
+                int l = *(left);
+                int r = *(right);
+
+                return fmin(cdata_weight * ::abs(l - r), cdata_weight * cmax_data_term);
            }
        };
        template <> struct DataCostPerPixel<3>
        {
            static __device__ __forceinline__ float compute(const uchar* left, const uchar* right)
            {
-                float tb = 0.114f * ::abs((int)left[0] - right[0]);
-                float tg = 0.587f * ::abs((int)left[1] - right[1]);
-                float tr = 0.299f * ::abs((int)left[2] - right[2]);
+                uchar3 l = *((const uchar3*)left);
+                uchar3 r = *((const uchar3*)right);
+
+                float tb = 0.114f * ::abs((int)l.x - r.x);
+                float tg = 0.587f * ::abs((int)l.y - r.y);
+                float tr = 0.299f * ::abs((int)l.z - r.z);

                return fmin(cdata_weight * (tr + tg + tb), cdata_weight * cmax_data_term);
            }
@@ -270,7 +270,7 @@ void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int dtype,
    convertTo(src, dst, alpha, beta, stream);
 }

-#if CUDA_VERSION >= 5000
+#if CUDART_VERSION >= 5000

 namespace
 {
@@ -293,7 +293,7 @@ namespace

 void cv::gpu::Stream::enqueueHostCallback(StreamCallback callback, void* userData)
 {
-#if CUDA_VERSION >= 5000
+#if CUDART_VERSION >= 5000
    CallbackData* data = new CallbackData;
    data->callback = callback;
    data->userData = userData;
@@ -1955,7 +1955,7 @@ void cv::gpu::compare(const GpuMat& src, Scalar sc, GpuMat& dst, int cmpop, Stre

 namespace arithm
 {
-    template <typename T> void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template <typename T> void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
 }

 void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask, Stream& s)
@@ -1964,39 +1964,73 @@ void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask, St

    const int depth = src.depth();

-    CV_Assert( depth <= CV_64F );
+    CV_Assert( depth < CV_32F );
    CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src.size()) );

    dst.create(src.size(), src.type());

    cudaStream_t stream = StreamAccessor::getStream(s);

-    const int bcols = (int) (src.cols * src.elemSize());
-
-    if ((bcols & 3) == 0)
+    if (mask.empty())
    {
-        const int vcols = bcols >> 2;
+        const int bcols = (int) (src.cols * src.elemSize());
+        bool aligned =
+                isAligned(src.data, sizeof(unsigned int)) &&
+                isAligned(dst.data, sizeof(unsigned int));

-        bitMatNot<unsigned int>(
-                    PtrStepSzb(src.rows, vcols, src.data, src.step),
-                    PtrStepSzb(src.rows, vcols, dst.data, dst.step),
-                    mask, stream);
-    }
-    else if ((bcols & 1) == 0)
-    {
-        const int vcols = bcols >> 1;
+        if (aligned && (bcols & 3) == 0)
+        {
+            const int vcols = bcols >> 2;

-        bitMatNot<unsigned short>(
-                    PtrStepSzb(src.rows, vcols, src.data, src.step),
-                    PtrStepSzb(src.rows, vcols, dst.data, dst.step),
-                    mask, stream);
+            bitMatNot<unsigned int>(
+                        PtrStepSzb(src.rows, vcols, src.data, src.step),
+                        PtrStepSzb(src.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else if (aligned && (bcols & 1) == 0)
+        {
+            const int vcols = bcols >> 1;
+
+            bitMatNot<unsigned short>(
+                        PtrStepSzb(src.rows, vcols, src.data, src.step),
+                        PtrStepSzb(src.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else
+        {
+            bitMatNot<unsigned char>(
+                        PtrStepSzb(src.rows, bcols, src.data, src.step),
+                        PtrStepSzb(src.rows, bcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
    }
    else
    {
-        bitMatNot<unsigned char>(
-                    PtrStepSzb(src.rows, bcols, src.data, src.step),
-                    PtrStepSzb(src.rows, bcols, dst.data, dst.step),
-                    mask, stream);
+        const int elem_size = src.elemSize1();
+        const int num_channels = src.channels();
+        const int bcols = src.cols * num_channels;
+
+        if (elem_size == 1)
+        {
+            bitMatNot<unsigned char>(
+                        PtrStepSzb(src.rows, bcols, src.data, src.step),
+                        PtrStepSzb(src.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 2)
+        {
+            bitMatNot<unsigned short>(
+                        PtrStepSzb(src.rows, bcols, src.data, src.step),
+                        PtrStepSzb(src.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 4)
+        {
+            bitMatNot<unsigned int>(
+                        PtrStepSzb(src.rows, bcols, src.data, src.step),
+                        PtrStepSzb(src.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
    }
 }

@@ -2005,9 +2039,9 @@ void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask, St

 namespace arithm
 {
-    template <typename T> void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template <typename T> void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
-    template <typename T> void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream);
+    template <typename T> void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template <typename T> void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
+    template <typename T> void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, int num_channels, cudaStream_t stream);
 }

 void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, Stream& s)
@@ -2016,7 +2050,7 @@ void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, c

    const int depth = src1.depth();

-    CV_Assert( depth <= CV_64F );
+    CV_Assert( depth < CV_32F );
    CV_Assert( src2.size() == src1.size() && src2.type() == src1.type() );
    CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src1.size()) );

@@ -2024,36 +2058,73 @@ void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, c

    cudaStream_t stream = StreamAccessor::getStream(s);

-    const int bcols = (int) (src1.cols * src1.elemSize());
-
-    if ((bcols & 3) == 0)
+    if (mask.empty())
    {
-        const int vcols = bcols >> 2;
+        const int bcols = (int) (src1.cols * src1.elemSize());
+        bool aligned =
+                isAligned(src1.data, sizeof(unsigned int)) &&
+                isAligned(src2.data, sizeof(unsigned int)) &&
+                isAligned(dst.data, sizeof(unsigned int));

-        bitMatAnd<unsigned int>(
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
-    }
-    else if ((bcols & 1) == 0)
-    {
-        const int vcols = bcols >> 1;
+        if (aligned && (bcols & 3) == 0)
+        {
+            const int vcols = bcols >> 2;

-        bitMatAnd<unsigned short>(
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
+            bitMatAnd<unsigned int>(
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else if (aligned && (bcols & 1) == 0)
+        {
+            const int vcols = bcols >> 1;
+
+            bitMatAnd<unsigned short>(
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else
+        {
+            bitMatAnd<unsigned char>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
    }
    else
    {
+        const int elem_size = src1.elemSize1();
+        const int num_channels = src1.channels();
+        const int bcols = src1.cols * num_channels;

-        bitMatAnd<unsigned char>(
-                    PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
-                    mask, stream);
+        if (elem_size == 1)
+        {
+            bitMatAnd<unsigned char>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 2)
+        {
+            bitMatAnd<unsigned short>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 4)
+        {
+            bitMatAnd<unsigned int>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
    }
 }

@@ -2063,7 +2134,7 @@ void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, co

    const int depth = src1.depth();

-    CV_Assert( depth <= CV_64F );
+    CV_Assert( depth < CV_32F );
    CV_Assert( src2.size() == src1.size() && src2.type() == src1.type() );
    CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src1.size()) );

@@ -2071,36 +2142,73 @@ void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, co

    cudaStream_t stream = StreamAccessor::getStream(s);

-    const int bcols = (int) (src1.cols * src1.elemSize());
-
-    if ((bcols & 3) == 0)
+    if (mask.empty())
    {
-        const int vcols = bcols >> 2;
+        const int bcols = (int) (src1.cols * src1.elemSize());
+        bool aligned =
+                isAligned(src1.data, sizeof(unsigned int)) &&
+                isAligned(src2.data, sizeof(unsigned int)) &&
+                isAligned(dst.data, sizeof(unsigned int));

-        bitMatOr<unsigned int>(
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
-    }
-    else if ((bcols & 1) == 0)
-    {
-        const int vcols = bcols >> 1;
+        if (aligned && (bcols & 3) == 0)
+        {
+            const int vcols = bcols >> 2;

-        bitMatOr<unsigned short>(
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
+            bitMatOr<unsigned int>(
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else if (aligned && (bcols & 1) == 0)
+        {
+            const int vcols = bcols >> 1;
+
+            bitMatOr<unsigned short>(
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else
+        {
+            bitMatOr<unsigned char>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
    }
    else
    {
+        const int elem_size = src1.elemSize1();
+        const int num_channels = src1.channels();
+        const int bcols = src1.cols * num_channels;

-        bitMatOr<unsigned char>(
-                    PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
-                    mask, stream);
+        if (elem_size == 1)
+        {
+            bitMatOr<unsigned char>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 2)
+        {
+            bitMatOr<unsigned short>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 4)
+        {
+            bitMatOr<unsigned int>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
    }
 }

@@ -2110,7 +2218,7 @@ void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, c

    const int depth = src1.depth();

-    CV_Assert( depth <= CV_64F );
+    CV_Assert( depth < CV_32F );
    CV_Assert( src2.size() == src1.size() && src2.type() == src1.type() );
    CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src1.size()) );

@@ -2118,36 +2226,73 @@ void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, c

    cudaStream_t stream = StreamAccessor::getStream(s);

-    const int bcols = (int) (src1.cols * src1.elemSize());
-
-    if ((bcols & 3) == 0)
+    if (mask.empty())
    {
-        const int vcols = bcols >> 2;
+        const int bcols = (int) (src1.cols * src1.elemSize());
+        bool aligned =
+                isAligned(src1.data, sizeof(unsigned int)) &&
+                isAligned(src2.data, sizeof(unsigned int)) &&
+                isAligned(dst.data, sizeof(unsigned int));

-        bitMatXor<unsigned int>(
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
-    }
-    else if ((bcols & 1) == 0)
-    {
-        const int vcols = bcols >> 1;
+        if (aligned && (bcols & 3) == 0)
+        {
+            const int vcols = bcols >> 2;

-        bitMatXor<unsigned short>(
-                    PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
-                    mask, stream);
+            bitMatXor<unsigned int>(
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else if (aligned && (bcols & 1) == 0)
+        {
+            const int vcols = bcols >> 1;
+
+            bitMatXor<unsigned short>(
+                        PtrStepSzb(src1.rows, vcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, vcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, vcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
+        else
+        {
+            bitMatXor<unsigned char>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        PtrStepb(), 1, stream);
+        }
    }
    else
    {
+        const int elem_size = src1.elemSize1();
+        const int num_channels = src1.channels();
+        const int bcols = src1.cols * num_channels;

-        bitMatXor<unsigned char>(
-                    PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
-                    PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
-                    PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
-                    mask, stream);
+        if (elem_size == 1)
+        {
+            bitMatXor<unsigned char>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 2)
+        {
+            bitMatXor<unsigned short>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
+        else if (elem_size == 4)
+        {
+            bitMatXor<unsigned int>(
+                        PtrStepSzb(src1.rows, bcols, src1.data, src1.step),
+                        PtrStepSzb(src1.rows, bcols, src2.data, src2.step),
+                        PtrStepSzb(src1.rows, bcols, dst.data, dst.step),
+                        mask, num_channels, stream);
+        }
    }
 }

@@ -144,7 +144,7 @@ namespace

 void cv::gpu::graphcut(GpuMat& terminals, GpuMat& leftTransp, GpuMat& rightTransp, GpuMat& top, GpuMat& bottom, GpuMat& labels, GpuMat& buf, Stream& s)
 {
-#if (CUDA_VERSION < 5000)
+#if (CUDART_VERSION < 5000)
    CV_Assert(terminals.type() == CV_32S);
 #else
    CV_Assert(terminals.type() == CV_32S || terminals.type() == CV_32F);
@@ -181,7 +181,7 @@ void cv::gpu::graphcut(GpuMat& terminals, GpuMat& leftTransp, GpuMat& rightTrans

    NppiGraphcutStateHandler state(sznpp, buf.ptr<Npp8u>(), nppiGraphcutInitAlloc);

-#if (CUDA_VERSION < 5000)
+#if (CUDART_VERSION < 5000)
    nppSafeCall( nppiGraphcut_32s8u(terminals.ptr<Npp32s>(), leftTransp.ptr<Npp32s>(), rightTransp.ptr<Npp32s>(), top.ptr<Npp32s>(), bottom.ptr<Npp32s>(),
        static_cast<int>(terminals.step), static_cast<int>(leftTransp.step), sznpp, labels.ptr<Npp8u>(), static_cast<int>(labels.step), state) );
 #else
@@ -204,7 +204,7 @@ void cv::gpu::graphcut(GpuMat& terminals, GpuMat& leftTransp, GpuMat& rightTrans
 void cv::gpu::graphcut(GpuMat& terminals, GpuMat& leftTransp, GpuMat& rightTransp, GpuMat& top, GpuMat& topLeft, GpuMat& topRight,
              GpuMat& bottom, GpuMat& bottomLeft, GpuMat& bottomRight, GpuMat& labels, GpuMat& buf, Stream& s)
 {
-#if (CUDA_VERSION < 5000)
+#if (CUDART_VERSION < 5000)
    CV_Assert(terminals.type() == CV_32S);
 #else
    CV_Assert(terminals.type() == CV_32S || terminals.type() == CV_32F);
@@ -253,7 +253,7 @@ void cv::gpu::graphcut(GpuMat& terminals, GpuMat& leftTransp, GpuMat& rightTrans

    NppiGraphcutStateHandler state(sznpp, buf.ptr<Npp8u>(), nppiGraphcut8InitAlloc);

-#if (CUDA_VERSION < 5000)
+#if (CUDART_VERSION < 5000)
    nppSafeCall( nppiGraphcut8_32s8u(terminals.ptr<Npp32s>(), leftTransp.ptr<Npp32s>(), rightTransp.ptr<Npp32s>(),
        top.ptr<Npp32s>(), topLeft.ptr<Npp32s>(), topRight.ptr<Npp32s>(),
        bottom.ptr<Npp32s>(), bottomLeft.ptr<Npp32s>(), bottomRight.ptr<Npp32s>(),
@@ -127,9 +127,6 @@ cv::gpu::HOGDescriptor::HOGDescriptor(Size win_size_, Size block_size_, Size blo

    Size cells_per_block = Size(block_size.width / cell_size.width, block_size.height / cell_size.height);
    CV_Assert(cells_per_block == Size(2, 2));
-
-    cv::Size blocks_per_win = numPartsWithin(win_size, block_size, block_stride);
-    hog::set_up_constants(nbins, block_stride.width, block_stride.height, blocks_per_win.width, blocks_per_win.height);
 }

 size_t cv::gpu::HOGDescriptor::getDescriptorSize() const
@@ -221,6 +218,9 @@ void cv::gpu::HOGDescriptor::computeGradient(const GpuMat& img, GpuMat& _grad, G

 void cv::gpu::HOGDescriptor::computeBlockHistograms(const GpuMat& img)
 {
+    cv::Size blocks_per_win = numPartsWithin(win_size, block_size, block_stride);
+    hog::set_up_constants(nbins, block_stride.width, block_stride.height, blocks_per_win.width, blocks_per_win.height);
+
    computeGradient(img, grad, qangle);

    size_t block_hist_size = getBlockHistogramSize();
@@ -132,7 +132,7 @@ void cv::gpu::meanStdDev(const GpuMat& src, Scalar& mean, Scalar& stddev, GpuMat
    DeviceBuffer dbuf(2);

    int bufSize;
-#if (CUDA_VERSION <= 4020)
+#if (CUDART_VERSION <= 4020)
    nppSafeCall( nppiMeanStdDev8uC1RGetBufferHostSize(sz, &bufSize) );
 #else
    nppSafeCall( nppiMeanStdDevGetBufferHostSize_8u_C1R(sz, &bufSize) );
@@ -187,7 +187,7 @@ double cv::gpu::norm(const GpuMat& src1, const GpuMat& src2, int normType)
    CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());
    CV_Assert(normType == NORM_INF || normType == NORM_L1 || normType == NORM_L2);

-#if CUDA_VERSION < 5050
+#if CUDART_VERSION < 5050
    typedef NppStatus (*func_t)(const Npp8u* pSrc1, int nSrcStep1, const Npp8u* pSrc2, int nSrcStep2, NppiSize oSizeROI, Npp64f* pRetVal);

    static const func_t funcs[] = {nppiNormDiff_Inf_8u_C1R, nppiNormDiff_L1_8u_C1R, nppiNormDiff_L2_8u_C1R};
@@ -212,7 +212,7 @@ double cv::gpu::norm(const GpuMat& src1, const GpuMat& src2, int normType)

    DeviceBuffer dbuf;

-#if CUDA_VERSION < 5050
+#if CUDART_VERSION < 5050
    nppSafeCall( funcs[funcIdx](src1.ptr<Npp8u>(), static_cast<int>(src1.step), src2.ptr<Npp8u>(), static_cast<int>(src2.step), sz, dbuf) );
 #else
    int bufSize;
@@ -280,7 +280,8 @@ __global__ void scanRows(T_in *d_src, Ncv32u texOffs, Ncv32u srcWidth, Ncv32u sr

    __shared__ T_out shmem[NUM_SCAN_THREADS * 2];
    __shared__ T_out carryElem;
-    carryElem = 0;
+    if (threadIdx.x == 0)
+        carryElem = 0;
    __syncthreads();

    while (numBuckets--)
@@ -288,7 +289,7 @@ __global__ void scanRows(T_in *d_src, Ncv32u texOffs, Ncv32u srcWidth, Ncv32u sr
        Ncv32u curElemOffs = offsetX + threadIdx.x;
        T_out curScanElem;

-        T_in curElem;
+        T_in curElem = 0;
        T_out curElemMod;

        if (curElemOffs < srcWidth)
@@ -82,8 +82,6 @@
 #define OPENCV_GPU_UNUSED(x) (void)x

 #ifdef HAVE_CUDA
-
-    #include <cuda.h>
    #include <cuda_runtime.h>
    #include <npp.h>

@@ -96,6 +94,7 @@
    #endif

    #ifdef HAVE_NVCUVID
+        #include <cuda.h>
        #include <nvcuvid.h>

        #ifdef WIN32
@@ -58,7 +58,7 @@ int main(int argc, char** argv)
                "{ h | help ?            | false | Print help}"
                "{ i | info              | false | Print information about system and exit }"
                "{ d | device            | -1   | Device on which tests will be executed (-1 means all devices) }"
-                "{ nvtest_output_level | nvtest_output_level | none | NVidia test verbosity level (none, compact, full) }"
+                "{ nvtest_output_level | nvtest_output_level | full | NVidia test verbosity level (none, compact, full) }"
                ;

        CommandLineParser cmd(argc, (const char**)argv, keys);
@@ -323,7 +323,8 @@ static void devNullOutput(const std::string& msg)
 bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel)
 {
    path = test_data_path.c_str();
-    ncvSetDebugOutputHandler(devNullOutput);
+    if (outputLevel != OutputLevelFull)
+        ncvSetDebugOutputHandler(devNullOutput);

    NCVAutoTestLister testListerII("NPPST Integral Image", outputLevel);

@@ -158,7 +158,7 @@ GPU_TEST_P(StereoConstantSpaceBP, Regression)
    cv::Mat h_disp(disp);
    h_disp.convertTo(h_disp, disp_gold.depth());

-    EXPECT_MAT_NEAR(disp_gold, h_disp, 1.0);
+    EXPECT_MAT_SIMILAR(disp_gold, h_disp, 1e-4);
 }

 INSTANTIATE_TEST_CASE_P(GPU_Calib3D, StereoConstantSpaceBP, ALL_DEVICES);
@@ -298,8 +298,8 @@ GPU_TEST_P(SolvePnPRansac, Accuracy)
                            camera_mat, cv::Mat(1, 8, CV_32F, cv::Scalar::all(0)),
                            rvec, tvec, false, 200, 2.f, 100, &inliers);

-    ASSERT_LE(cv::norm(rvec - rvec_gold), 1e-3);
-    ASSERT_LE(cv::norm(tvec - tvec_gold), 1e-3);
+    ASSERT_LE(cv::norm(rvec - rvec_gold), 2e-3);
+    ASSERT_LE(cv::norm(tvec - tvec_gold), 2e-3);
 }

 INSTANTIATE_TEST_CASE_P(GPU_Calib3D, SolvePnPRansac, ALL_DEVICES);
@@ -2048,7 +2048,7 @@ GPU_TEST_P(CvtColor, Luv2LRGBA)
    EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_8U ? 1 : 1e-4);
 }

-#if defined (CUDA_VERSION) && (CUDA_VERSION >= 5000)
+#if defined (CUDART_VERSION) && (CUDART_VERSION >= 5000)

 GPU_TEST_P(CvtColor, RGBA2mRGBA)
 {
@@ -2066,7 +2066,7 @@ GPU_TEST_P(CvtColor, RGBA2mRGBA)
    EXPECT_MAT_NEAR(dst_gold, dst, 1);
 }

-#endif // defined (CUDA_VERSION) && (CUDA_VERSION >= 5000)
+#endif // defined (CUDART_VERSION) && (CUDART_VERSION >= 5000)

 GPU_TEST_P(CvtColor, BayerBG2BGR)
 {
@@ -1785,72 +1785,95 @@ INSTANTIATE_TEST_CASE_P(GPU_Core, Compare_Scalar, testing::Combine(
 //////////////////////////////////////////////////////////////////////////////
 // Bitwise_Array

-PARAM_TEST_CASE(Bitwise_Array, cv::gpu::DeviceInfo, cv::Size, MatType)
+PARAM_TEST_CASE(Bitwise_Array, cv::gpu::DeviceInfo, cv::Size, MatType, UseRoi)
 {
    cv::gpu::DeviceInfo devInfo;
    cv::Size size;
    int type;
+    bool useRoi;

    cv::Mat src1;
    cv::Mat src2;

+    cv::Mat mask;
+
    virtual void SetUp()
    {
        devInfo = GET_PARAM(0);
        size = GET_PARAM(1);
        type = GET_PARAM(2);
+        useRoi = GET_PARAM(3);

        cv::gpu::setDevice(devInfo.deviceID());

        src1 = randomMat(size, type, 0.0, std::numeric_limits<int>::max());
        src2 = randomMat(size, type, 0.0, std::numeric_limits<int>::max());
+
+        mask = randomMat(size, CV_8UC1, 0.0, 2.0);
    }
 };

 GPU_TEST_P(Bitwise_Array, Not)
 {
-    cv::gpu::GpuMat dst;
-    cv::gpu::bitwise_not(loadMat(src1), dst);
+    cv::gpu::GpuMat dst_nomask, dst_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::gpu::bitwise_not(loadMat(src1, useRoi), dst_nomask);
+    cv::gpu::bitwise_not(loadMat(src1, useRoi), dst_mask, loadMat(mask, useRoi));

-    cv::Mat dst_gold = ~src1;
+    cv::Mat dst_gold_nomask, dst_gold_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::bitwise_not(src1, dst_gold_nomask);
+    cv::bitwise_not(src1, dst_gold_mask, mask);

-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_nomask, dst_nomask, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_mask, dst_mask, 0.0);
 }

 GPU_TEST_P(Bitwise_Array, Or)
 {
-    cv::gpu::GpuMat dst;
-    cv::gpu::bitwise_or(loadMat(src1), loadMat(src2), dst);
+    cv::gpu::GpuMat dst_nomask, dst_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::gpu::bitwise_or(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_nomask);
+    cv::gpu::bitwise_or(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_mask, loadMat(mask, useRoi));

-    cv::Mat dst_gold = src1 | src2;
+    cv::Mat dst_gold_nomask, dst_gold_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::bitwise_or(src1, src2, dst_gold_nomask);
+    cv::bitwise_or(src1, src2, dst_gold_mask, mask);

-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_nomask, dst_nomask, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_mask, dst_mask, 0.0);
 }

 GPU_TEST_P(Bitwise_Array, And)
 {
-    cv::gpu::GpuMat dst;
-    cv::gpu::bitwise_and(loadMat(src1), loadMat(src2), dst);
+    cv::gpu::GpuMat dst_nomask, dst_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::gpu::bitwise_and(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_nomask);
+    cv::gpu::bitwise_and(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_mask, loadMat(mask, useRoi));

-    cv::Mat dst_gold = src1 & src2;
+    cv::Mat dst_gold_nomask, dst_gold_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::bitwise_and(src1, src2, dst_gold_nomask);
+    cv::bitwise_and(src1, src2, dst_gold_mask, mask);

-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_nomask, dst_nomask, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_mask, dst_mask, 0.0);
 }

 GPU_TEST_P(Bitwise_Array, Xor)
 {
-    cv::gpu::GpuMat dst;
-    cv::gpu::bitwise_xor(loadMat(src1), loadMat(src2), dst);
+    cv::gpu::GpuMat dst_nomask, dst_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::gpu::bitwise_xor(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_nomask);
+    cv::gpu::bitwise_xor(loadMat(src1, useRoi), loadMat(src2, useRoi), dst_mask, loadMat(mask, useRoi));

-    cv::Mat dst_gold = src1 ^ src2;
+    cv::Mat dst_gold_nomask, dst_gold_mask(src1.size(), src1.type(), cv::Scalar::all(0));
+    cv::bitwise_xor(src1, src2, dst_gold_nomask);
+    cv::bitwise_xor(src1, src2, dst_gold_mask, mask);

-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_nomask, dst_nomask, 0.0);
+    EXPECT_MAT_NEAR(dst_gold_mask, dst_mask, 0.0);
 }

 INSTANTIATE_TEST_CASE_P(GPU_Core, Bitwise_Array, testing::Combine(
    ALL_DEVICES,
    DIFFERENT_SIZES,
-    TYPES(CV_8U, CV_32S, 1, 4)));
+    TYPES(CV_8U, CV_32S, 1, 4),
+    WHOLE_SUBMAT));

 //////////////////////////////////////////////////////////////////////////////
 // Bitwise_Scalar
@@ -134,8 +134,8 @@ GPU_TEST_P(BruteForceNonLocalMeans, Regression)
    cv::resize(bgr_gold, bgr_gold, cv::Size(256, 256));
    cv::resize(gray_gold, gray_gold, cv::Size(256, 256));

-    EXPECT_MAT_NEAR(bgr_gold, dbgr, 1e-4);
-    EXPECT_MAT_NEAR(gray_gold, dgray, 1e-4);
+    EXPECT_MAT_NEAR(bgr_gold, dbgr, 1);
+    EXPECT_MAT_NEAR(gray_gold, dgray, 1);
 }

 INSTANTIATE_TEST_CASE_P(GPU_Denoising, BruteForceNonLocalMeans, ALL_DEVICES);
@@ -681,7 +681,7 @@ PARAM_TEST_CASE(MatchTemplate8U, cv::gpu::DeviceInfo, cv::Size, TemplateSize, Ch
    }
 };

-GPU_TEST_P(MatchTemplate8U, Accuracy)
+GPU_TEST_P(MatchTemplate8U, DISABLED_Accuracy)
 {
    cv::Mat image = randomMat(size, CV_MAKETYPE(CV_8U, cn));
    cv::Mat templ = randomMat(templ_size, CV_MAKETYPE(CV_8U, cn));
@@ -692,7 +692,18 @@ GPU_TEST_P(MatchTemplate8U, Accuracy)
    cv::Mat dst_gold;
    cv::matchTemplate(image, templ, dst_gold, method);

-    EXPECT_MAT_NEAR(dst_gold, dst, templ_size.area() * 1e-1);
+    cv::Mat h_dst(dst);
+    ASSERT_EQ(dst_gold.size(), h_dst.size());
+    ASSERT_EQ(dst_gold.type(), h_dst.type());
+    for (int y = 0; y < h_dst.rows; ++y)
+    {
+        for (int x = 0; x < h_dst.cols; ++x)
+        {
+            float gold_val = dst_gold.at<float>(y, x);
+            float actual_val = dst_gold.at<float>(y, x);
+            ASSERT_FLOAT_EQ(gold_val, actual_val) << y << ", " << x;
+        }
+    }
 }

 INSTANTIATE_TEST_CASE_P(GPU_ImgProc, MatchTemplate8U, testing::Combine(
@@ -727,7 +738,7 @@ PARAM_TEST_CASE(MatchTemplate32F, cv::gpu::DeviceInfo, cv::Size, TemplateSize, C
    }
 };

-GPU_TEST_P(MatchTemplate32F, Regression)
+GPU_TEST_P(MatchTemplate32F, DISABLED_Regression)
 {
    cv::Mat image = randomMat(size, CV_MAKETYPE(CV_32F, cn));
    cv::Mat templ = randomMat(templ_size, CV_MAKETYPE(CV_32F, cn));
@@ -738,7 +749,18 @@ GPU_TEST_P(MatchTemplate32F, Regression)
    cv::Mat dst_gold;
    cv::matchTemplate(image, templ, dst_gold, method);

-    EXPECT_MAT_NEAR(dst_gold, dst, templ_size.area() * 1e-1);
+    cv::Mat h_dst(dst);
+    ASSERT_EQ(dst_gold.size(), h_dst.size());
+    ASSERT_EQ(dst_gold.type(), h_dst.type());
+    for (int y = 0; y < h_dst.rows; ++y)
+    {
+        for (int x = 0; x < h_dst.cols; ++x)
+        {
+            float gold_val = dst_gold.at<float>(y, x);
+            float actual_val = dst_gold.at<float>(y, x);
+            ASSERT_FLOAT_EQ(gold_val, actual_val) << y << ", " << x;
+        }
+    }
 }

 INSTANTIATE_TEST_CASE_P(GPU_ImgProc, MatchTemplate32F, testing::Combine(
@@ -765,7 +787,7 @@ PARAM_TEST_CASE(MatchTemplateBlackSource, cv::gpu::DeviceInfo, TemplateMethod)
    }
 };

-GPU_TEST_P(MatchTemplateBlackSource, Accuracy)
+GPU_TEST_P(MatchTemplateBlackSource, DISABLED_Accuracy)
 {
    cv::Mat image = readImage("matchtemplate/black.png");
    ASSERT_FALSE(image.empty());
@@ -810,7 +832,7 @@ PARAM_TEST_CASE(MatchTemplate_CCOEF_NORMED, cv::gpu::DeviceInfo, std::pair<std::
    }
 };

-GPU_TEST_P(MatchTemplate_CCOEF_NORMED, Accuracy)
+GPU_TEST_P(MatchTemplate_CCOEF_NORMED, DISABLED_Accuracy)
 {
    cv::Mat image = readImage(imageName);
    ASSERT_FALSE(image.empty());
@@ -859,7 +881,7 @@ struct MatchTemplate_CanFindBigTemplate : testing::TestWithParam<cv::gpu::Device
    }
 };

-GPU_TEST_P(MatchTemplate_CanFindBigTemplate, SQDIFF_NORMED)
+GPU_TEST_P(MatchTemplate_CanFindBigTemplate, DISABLED_SQDIFF_NORMED)
 {
    cv::Mat scene = readImage("matchtemplate/scene.png");
    ASSERT_FALSE(scene.empty());
@@ -882,7 +904,7 @@ GPU_TEST_P(MatchTemplate_CanFindBigTemplate, SQDIFF_NORMED)
    ASSERT_EQ(0, minLoc.y);
 }

-GPU_TEST_P(MatchTemplate_CanFindBigTemplate, SQDIFF)
+GPU_TEST_P(MatchTemplate_CanFindBigTemplate, DISABLED_SQDIFF)
 {
    cv::Mat scene = readImage("matchtemplate/scene.png");
    ASSERT_FALSE(scene.empty());
@@ -44,7 +44,7 @@

 #ifdef HAVE_CUDA

-OutputLevel nvidiaTestOutputLevel = OutputLevelNone;
+OutputLevel nvidiaTestOutputLevel = OutputLevelFull;

 using namespace cvtest;
 using namespace testing;
@@ -48,9 +48,24 @@ using namespace cvtest;

 //#define DUMP

-struct HOG : testing::TestWithParam<cv::gpu::DeviceInfo>, cv::gpu::HOGDescriptor
+class HogForTest : public cv::gpu::HOGDescriptor
+{
+public:
+    cv::gpu::GpuMat getBlockHists() const
+    {
+        return block_hists;
+    }
+
+    void computeBlockHistograms(const cv::gpu::GpuMat& img)
+    {
+        cv::gpu::HOGDescriptor::computeBlockHistograms(img);
+    }
+};
+
+struct HOG : testing::TestWithParam<cv::gpu::DeviceInfo>
 {
    cv::gpu::DeviceInfo devInfo;
+    cv::Ptr<HogForTest> hog;

 #ifdef DUMP
    std::ofstream f;
@@ -69,6 +84,8 @@ struct HOG : testing::TestWithParam<cv::gpu::DeviceInfo>, cv::gpu::HOGDescriptor
        devInfo = GetParam();

        cv::gpu::setDevice(devInfo.deviceID());
+
+        hog = new HogForTest;
    }

 #ifdef DUMP
@@ -126,39 +143,39 @@ struct HOG : testing::TestWithParam<cv::gpu::DeviceInfo>, cv::gpu::HOGDescriptor

    void testDetect(const cv::Mat& img)
    {
-        gamma_correction = false;
-        setSVMDetector(cv::gpu::HOGDescriptor::getDefaultPeopleDetector());
+        hog->gamma_correction = false;
+        hog->setSVMDetector(cv::gpu::HOGDescriptor::getDefaultPeopleDetector());

        std::vector<cv::Point> locations;

        // Test detect
-        detect(loadMat(img), locations, 0);
+        hog->detect(loadMat(img), locations, 0);

 #ifdef DUMP
-        dump(cv::Mat(block_hists), locations);
+        dump(cv::Mat(hog->getBlockHist()), locations);
 #else
-        compare(cv::Mat(block_hists), locations);
+        compare(cv::Mat(hog->getBlockHists()), locations);
 #endif

        // Test detect on smaller image
        cv::Mat img2;
        cv::resize(img, img2, cv::Size(img.cols / 2, img.rows / 2));
-        detect(loadMat(img2), locations, 0);
+        hog->detect(loadMat(img2), locations, 0);

 #ifdef DUMP
-        dump(cv::Mat(block_hists), locations);
+        dump(cv::Mat(hog->getBlockHist()), locations);
 #else
-        compare(cv::Mat(block_hists), locations);
+        compare(cv::Mat(hog->getBlockHists()), locations);
 #endif

        // Test detect on greater image
        cv::resize(img, img2, cv::Size(img.cols * 2, img.rows * 2));
-        detect(loadMat(img2), locations, 0);
+        hog->detect(loadMat(img2), locations, 0);

 #ifdef DUMP
-        dump(cv::Mat(block_hists), locations);
+        dump(cv::Mat(hog->getBlockHist()), locations);
 #else
-        compare(cv::Mat(block_hists), locations);
+        compare(cv::Mat(hog->getBlockHists()), locations);
 #endif
    }

@@ -216,8 +233,8 @@ GPU_TEST_P(HOG, GetDescriptors)

    // Convert train images into feature vectors (train table)
    cv::gpu::GpuMat descriptors, descriptors_by_cols;
-    getDescriptors(d_img, win_size, descriptors, DESCR_FORMAT_ROW_BY_ROW);
-    getDescriptors(d_img, win_size, descriptors_by_cols, DESCR_FORMAT_COL_BY_COL);
+    hog->getDescriptors(d_img, hog->win_size, descriptors, cv::gpu::HOGDescriptor::DESCR_FORMAT_ROW_BY_ROW);
+    hog->getDescriptors(d_img, hog->win_size, descriptors_by_cols, cv::gpu::HOGDescriptor::DESCR_FORMAT_COL_BY_COL);

    // Check size of the result train table
    wins_per_img_x = 3;
@@ -251,39 +268,39 @@ GPU_TEST_P(HOG, GetDescriptors)
    img_rgb = readImage("hog/positive1.png");
    ASSERT_TRUE(!img_rgb.empty());
    cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
-    computeBlockHistograms(cv::gpu::GpuMat(img));
+    hog->computeBlockHistograms(cv::gpu::GpuMat(img));
    // Everything is fine with interpolation for left top subimage
-    ASSERT_EQ(0.0, cv::norm((cv::Mat)block_hists, (cv::Mat)descriptors.rowRange(0, 1)));
+    ASSERT_EQ(0.0, cv::norm(cv::Mat(hog->getBlockHists()), (cv::Mat)descriptors.rowRange(0, 1)));

    img_rgb = readImage("hog/positive2.png");
    ASSERT_TRUE(!img_rgb.empty());
    cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
-    computeBlockHistograms(cv::gpu::GpuMat(img));
-    compare_inner_parts(cv::Mat(block_hists), cv::Mat(descriptors.rowRange(1, 2)));
+    hog->computeBlockHistograms(cv::gpu::GpuMat(img));
+    compare_inner_parts(cv::Mat(hog->getBlockHists()), cv::Mat(descriptors.rowRange(1, 2)));

    img_rgb = readImage("hog/negative1.png");
    ASSERT_TRUE(!img_rgb.empty());
    cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
-    computeBlockHistograms(cv::gpu::GpuMat(img));
-    compare_inner_parts(cv::Mat(block_hists), cv::Mat(descriptors.rowRange(2, 3)));
+    hog->computeBlockHistograms(cv::gpu::GpuMat(img));
+    compare_inner_parts(cv::Mat(hog->getBlockHists()), cv::Mat(descriptors.rowRange(2, 3)));

    img_rgb = readImage("hog/negative2.png");
    ASSERT_TRUE(!img_rgb.empty());
    cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
-    computeBlockHistograms(cv::gpu::GpuMat(img));
-    compare_inner_parts(cv::Mat(block_hists), cv::Mat(descriptors.rowRange(3, 4)));
+    hog->computeBlockHistograms(cv::gpu::GpuMat(img));
+    compare_inner_parts(cv::Mat(hog->getBlockHists()), cv::Mat(descriptors.rowRange(3, 4)));

    img_rgb = readImage("hog/positive3.png");
    ASSERT_TRUE(!img_rgb.empty());
    cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
-    computeBlockHistograms(cv::gpu::GpuMat(img));
-    compare_inner_parts(cv::Mat(block_hists), cv::Mat(descriptors.rowRange(4, 5)));
+    hog->computeBlockHistograms(cv::gpu::GpuMat(img));
+    compare_inner_parts(cv::Mat(hog->getBlockHists()), cv::Mat(descriptors.rowRange(4, 5)));

    img_rgb = readImage("hog/negative3.png");
    ASSERT_TRUE(!img_rgb.empty());
    cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
-    computeBlockHistograms(cv::gpu::GpuMat(img));
-    compare_inner_parts(cv::Mat(block_hists), cv::Mat(descriptors.rowRange(5, 6)));
+    hog->computeBlockHistograms(cv::gpu::GpuMat(img));
+    compare_inner_parts(cv::Mat(hog->getBlockHists()), cv::Mat(descriptors.rowRange(5, 6)));
 }

 INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, HOG, ALL_DEVICES);
@@ -501,8 +501,8 @@ GPU_TEST_P(OpticalFlowBM, Accuracy)
    cv::Mat velx, vely;
    calcOpticalFlowBM(frame0, frame1, block_size, shift_size, max_range, false, velx, vely);

-    EXPECT_MAT_NEAR(velx, d_velx, 0);
-    EXPECT_MAT_NEAR(vely, d_vely, 0);
+    EXPECT_MAT_NEAR(velx, d_velx, 1e-6);
+    EXPECT_MAT_NEAR(vely, d_vely, 1e-6);
 }

 INSTANTIATE_TEST_CASE_P(GPU_Video, OpticalFlowBM, ALL_DEVICES);
@@ -67,7 +67,6 @@
 #include "cvconfig.h"

 #ifdef HAVE_CUDA
-    #include <cuda.h>
    #include <cuda_runtime.h>

    #include "opencv2/ts/ts.hpp"
@@ -46,7 +46,7 @@

 using namespace cvtest;

-#if CUDA_VERSION >= 5000
+#if CUDART_VERSION >= 5000

 struct Async : testing::TestWithParam<cv::gpu::DeviceInfo>
 {
@@ -292,7 +292,15 @@ void CvCapture_FFMPEG::close()
    }

    if( picture )
+    {
+        // FFmpeg and Libav added avcodec_free_frame in different versions.
+#if LIBAVCODEC_BUILD >= (LIBAVCODEC_VERSION_MICRO >= 100 \
+    ? CALC_FFMPEG_VERSION(54, 59, 100) : CALC_FFMPEG_VERSION(54, 28, 0))
+        avcodec_free_frame(&picture);
+#else
        av_free(picture);
+#endif
+    }

    if( video_st )
    {
@@ -157,7 +157,10 @@ static int64 icvGCMaxFlow( GCVtx* vtx, int nvtx, GCEdge* edges, GCVtx**& _orphan
            v->t = v->weight < 0;
        }
        else
+        {
+            v->next = NULL;
            v->parent = 0;
+        }
    }

    first = first->next;
@@ -396,6 +399,8 @@ void cvReleaseStereoGCState( CvStereoGCState** _state )
    cvReleaseMat( &state->right );
    cvReleaseMat( &state->ptrLeft );
    cvReleaseMat( &state->ptrRight );
+    cvReleaseMat( &state->dispLeft );
+    cvReleaseMat( &state->dispRight );
    cvReleaseMat( &state->vtxBuf );
    cvReleaseMat( &state->edgeBuf );
    cvFree( _state );
@@ -858,6 +858,7 @@ void HOGDescriptor::detect(const Mat& img,
    Size winStride, Size padding, const vector<Point>& locations) const
 {
    hits.clear();
+    weights.clear();
    if( svmDetector.empty() )
        return;

@@ -260,7 +260,7 @@ namespace

 Ptr<FrameSource> cv::superres::createFrameSource_Video_GPU(const string& fileName)
 {
-    return new VideoFrameSource(fileName);
+    return new VideoFrameSource_GPU(fileName);
 }

 #endif // HAVE_OPENCV_GPU
Author	SHA1	Message	Date
Alexander Smorkalov	e49e75da06	OpenCV patch verison++.	2015-05-06 16:29:51 +03:00
Vladislav Vinogradov	1d40946959	Bug #4315 : fix CUDA bitwise operations with mask (cherry picked from commit `d87c30dc84`)	2015-05-06 16:27:12 +03:00
Alexander Smorkalov	2598392295	Added explicit deb package dependency from libtbb-dev if TBB is enabled. (cherry picked from commit `63d6cc5ca6`)	2015-05-06 16:26:50 +03:00
lujia	17cc5e2c40	bugfix_for_hog_detectMultiScale_with_weights (cherry picked from commit `7ce116695d`)	2015-05-06 16:24:39 +03:00
Vladislav Vinogradov	4b14400976	disable several heavy performance tests (cherry picked from commit `cbdddb473c`)	2015-05-06 16:24:07 +03:00
Vladislav Vinogradov	4704a254f7	disable Video_PyrLKOpticalFlowDense performance test sanity fails on Maxwell and CUDA 7.0 due to unknow reason (cherry picked from commit `b4c2891ef3`)	2015-05-06 16:23:41 +03:00
Vladislav Vinogradov	7708e9882e	use border extrapolation for central pixel in pyrDown in case if image has odd dimension (cherry picked from commit `1325213303`)	2015-05-06 16:23:28 +03:00
Vladislav Vinogradov	fe1bd0cc2f	fix racecheck warning in scanRows kernel (cherry picked from commit `fb15bdfb21`)	2015-05-06 16:23:09 +03:00
Vladislav Vinogradov	a752f25944	increase epsilon for solvePnPRansac function (cherry picked from commit `9d2d173485`)	2015-05-06 16:22:57 +03:00
Vladislav Vinogradov	a984da911b	increase epsilons for some sanity tests (cherry picked from commit `6a6619ec1e`)	2015-04-19 14:06:47 +03:00
Vladislav Vinogradov	df55be3c3d	fix BruteForceMatcher resource distribution added launch bounds attributes for all CUDA kernels (cherry picked from commit `d22516872c`)	2015-04-19 14:06:36 +03:00
Vladislav Vinogradov	55339de684	make NVIDIA tests verbose by default (cherry picked from commit `17608f7ade`)	2015-04-19 14:06:25 +03:00
Vladislav Vinogradov	f08dd510fa	fixed a bug in scanRows CUDA kernel (part of nppStIntegral) uninitialized value (cherry picked from commit `81ebe28c24`)	2015-04-19 14:06:11 +03:00
Vladislav Vinogradov	d308ed3712	fix GPU WARP border mode in CUDA 7.0 and Maxwell architecture (cherry picked from commit `27302c367c`)	2015-04-19 14:02:54 +03:00
Vladislav Vinogradov	6d0f8aa893	fix tests for gpu HOG initialize HOG after CUDA device switch (cherry picked from commit `c849492dfa`)	2015-04-19 14:01:16 +03:00
Vladimir Kolesnikov	bea98bd22a	Not block PDB file in multithreaded build on Windows. If used cl compiler, but generator is not Visual Studio e.g. Ninja, enable FS option to prevent blocking PDB file in multithreaded build. (cherry picked from commit `58c9135626`)	2015-04-19 14:00:24 +03:00
Vladislav Vinogradov	4539607ab1	fix gpu HOG implementation: move hog::set_up_constants from constructor to compute method if user changed CUDA device between constructor and computation, some variables were uninitialized (cherry picked from commit `21bbed7baf`)	2015-04-19 13:57:37 +03:00
Vladislav Vinogradov	b320138dba	fix GpuMat::setTo implementation previous implementation was not thread/stream safe, since it used constant memory new implementation doesn't use any global objects, so it is thread/stream safe (cherry picked from commit `4f5d30a865`)	2015-04-19 13:57:23 +03:00
Vladislav Vinogradov	95ea12588e	set epsilon for gpu OpticalFlowBM to 1e-6, since it uses floating point arithmetic (cherry picked from commit `c147ab1e85`)	2015-04-19 13:57:03 +03:00
Roman Donchenko	bcd08b5d18	Mark the libs component required Everything else depends on it, after all. (cherry picked from commit `cf54e3b97e`)	2015-04-19 13:55:35 +03:00
Roman Donchenko	19298ae3cb	Add component display names (cherry picked from commit `6d52ea8984`)	2015-04-19 13:55:22 +03:00
Roman Donchenko	5e06da3050	Update the CPack variables to match the changes in asmorkalov/CMake#1 Which also happens to align the non-Debian specific variables with the ones used by upstream CMake. (cherry picked from commit `b8c60234c3`) Conflicts: cmake/OpenCVPackaging.cmake	2015-04-19 13:54:32 +03:00
Roman Donchenko	17ac18a7b9	Remove useless CPACK_*_COMPONENT_INSTALL variables They don't actually do anything. And even if they did, all components are enabled by default, anyway. (cherry picked from commit `49fe496914`)	2015-04-19 13:42:14 +03:00
Roman Donchenko	c259590b26	Fix a memory leak in CvCapture_FFMPEG::close FFmpeg now requires that frames allocated with avcodec_alloc_frame are freed with avcodec_free_frame. (cherry picked from commit `77578d415f`)	2015-04-19 13:41:22 +03:00
Roman Donchenko	3f3ca85103	Don't use ${CMAKE_INSTALL_PREFIX} as an install destination Absolute destinations are not necessary, and prevent CPack from working. (cherry picked from commit `0387f8ad56`)	2015-04-19 13:39:55 +03:00
Alexander Smorkalov	05ddc16eaa	Added Debian changelog to -tests package. (cherry picked from commit `a87ccb9ac0`)	2015-04-19 13:39:01 +03:00
Alexander Smorkalov	2ba77614aa	Debian package names replaced by lower case version to satisfy lintian. (cherry picked from commit `e6ac64008b`)	2015-04-19 13:38:46 +03:00
Vladislav Vinogradov	ef347aa6a4	disable gpu::matchTemplate tests (cherry picked from commit `7bb8c50080`)	2015-04-19 13:38:20 +03:00
Alexander Smorkalov	b1cdb91139	Fixed samples install permissions for Debian packaging. (cherry picked from commit `cf852972d1`)	2015-04-19 13:37:50 +03:00
Alexander Smorkalov	a09ad35d98	opencv_testing.sh script installation is removed as run tests script does the same thing. (cherry picked from commit `9206ec30a2`)	2015-04-19 13:37:29 +03:00
Alexander Smorkalov	1b5835cd35	Added dependency from numpy to debian package with python bindings. (cherry picked from commit `be6b847675`)	2015-04-19 13:37:04 +03:00
Alexander Smorkalov	9932bbad3a	Added Debian changelog installation step for Debian package generation. (cherry picked from commit `ddc1b965b6`)	2015-04-19 13:36:47 +03:00
Vijay Pradeep	39ac84ff04	Fixing race condition by expanding resultsMutex lock section (cherry picked from commit `042ff210d5`)	2015-04-19 13:36:17 +03:00
Vladislav Vinogradov	fcbefaff86	fix tests for gpu::matchTemplate: use ASSERT_FLOAT_EQ to compare float values, it is more robust for large values (cherry picked from commit `d00f36ec75`)	2015-04-19 13:32:24 +03:00
Roman Donchenko	bf2256fb89	cvOpenFileStorage: reduce the scope of xml_buf and make sure it's freed... ... before any exceptions occur. (cherry picked from commit `08da247a87`)	2015-04-19 13:31:52 +03:00
Vladislav Vinogradov	b0b2fc9e3f	get rid of cuda.h usage (cherry picked from commit `eeb997261d`)	2015-04-19 13:29:45 +03:00
Roman Donchenko	1ccd64e102	Fix uninitialized memory reads and memory leaks in StereoGC (cherry picked from commit `7d8e5f623a`)	2015-04-19 13:27:55 +03:00
Vicente Olivert Riera	23bf3e337a	superres: Fix return value VideoFrameSource_GPU superres module fails to compile with the following error messages: [100%] Building CXX object modules/superres/CMakeFiles/opencv_superres.dir/src/super_resolution.cpp.o /opencv-2.4.10/modules/superres/src/frame_source.cpp: In function 'cv::Ptr<cv::superres::FrameSource> cv::superres::createFrameSource_Video_GPU(const string&)': /opencv-2.4.10/modules/superres/src/frame_source.cpp:263:16: error: expected type-specifier before 'VideoFrameSource' /opencv-2.4.10/modules/superres/src/frame_source.cpp:263:16: error: could not convert '(int)operator new(4ul)' from 'int' to 'cv::Ptr<cv::superres::FrameSource>' /opencv-2.4.10/modules/superres/src/frame_source.cpp:263:16: error: expected ';' before 'VideoFrameSource' /opencv-2.4.10/modules/superres/src/frame_source.cpp:263:41: error: 'VideoFrameSource' was not declared in this scope /opencv-2.4.10/modules/superres/src/frame_source.cpp:264:1: error: control reaches end of non-void function [-Werror=return-type] cc1plus: some warnings being treated as errors make[3]: * [modules/superres/CMakeFiles/opencv_superres.dir/src/frame_source.cpp.o] Error 1 make[3]: * Waiting for unfinished jobs.... This is caused because the return value of the createFrameSource_Video_GPU function should be a VideoFrameSource_GPU object. (cherry picked from commit `2e393ab833`)	2015-04-19 13:25:16 +03:00
Vladislav Vinogradov	3cf265992f	fix installation layout for debian packages: Install symlinks to shared libraries as a part of development package, not runtime package. It is default behavior for debian packages. (cherry picked from commit `f55c1cc0fb`)	2015-04-19 13:23:47 +03:00
Alexander Smorkalov	67635c6d65	Version++.	2014-11-05 08:46:55 +03:00
Vladislav Vinogradov	a26e496d00	minor fix for StereoCSBP data cost compute kernel and test (cherry picked from commit `84f33d0578`)	2014-11-04 10:19:38 +03:00
Vladislav Vinogradov	d579d3e596	increase epsilons for some tests, which functions use floating point arithm (cherry picked from commit `5c07e0b6d3`)	2014-11-04 10:19:18 +03:00
Vladislav Vinogradov	5a77176654	avoid pointer arithmetic on register memory in color conversion (cherry picked from commit `e0827069c1`)	2014-11-04 10:18:58 +03:00