temporary disable performance test for alphaComp function(cherry picked from commit 1ce5165cb7)

(cherry picked from commit dbe7634286)

3rdparty/libjasper/CMakeLists.txt

@@ -46,5 +46,5 @@ if(ENABLE_SOLUTION_FOLDERS)
 endif()
 
 if(NOT BUILD_SHARED_LIBS)
-  ocv_install_target(${JASPER_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main)
+  ocv_install_target(${JASPER_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev)
 endif()

3rdparty/libjpeg/CMakeLists.txt

@@ -39,5 +39,5 @@ if(ENABLE_SOLUTION_FOLDERS)
 endif()
 
 if(NOT BUILD_SHARED_LIBS)
-  ocv_install_target(${JPEG_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main)
+  ocv_install_target(${JPEG_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev)
 endif()

3rdparty/libpng/CMakeLists.txt

@@ -55,5 +55,5 @@ if(ENABLE_SOLUTION_FOLDERS)
 endif()
 
 if(NOT BUILD_SHARED_LIBS)
-  ocv_install_target(${PNG_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main)
+  ocv_install_target(${PNG_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev)
 endif()

3rdparty/libtiff/CMakeLists.txt

@@ -115,5 +115,5 @@ if(ENABLE_SOLUTION_FOLDERS)
 endif()
 
 if(NOT BUILD_SHARED_LIBS)
-  ocv_install_target(${TIFF_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main)
+  ocv_install_target(${TIFF_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev)
 endif()

3rdparty/openexr/CMakeLists.txt

@@ -62,7 +62,7 @@ if(ENABLE_SOLUTION_FOLDERS)
 endif()
 
 if(NOT BUILD_SHARED_LIBS)
-  ocv_install_target(IlmImf EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main)
+  ocv_install_target(IlmImf EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev)
 endif()
 
 set(OPENEXR_INCLUDE_PATHS ${OPENEXR_INCLUDE_PATHS} PARENT_SCOPE)

3rdparty/tbb/CMakeLists.txt

@@ -248,9 +248,9 @@ if(ENABLE_SOLUTION_FOLDERS)
 endif()
 
 ocv_install_target(tbb EXPORT OpenCVModules
-    RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main
-    LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main
-    ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main
+    RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT libs
+    LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs
+    ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev
     )
 
 # get TBB version

3rdparty/zlib/CMakeLists.txt

@@ -95,5 +95,5 @@ if(ENABLE_SOLUTION_FOLDERS)
 endif()
 
 if(NOT BUILD_SHARED_LIBS)
-  ocv_install_target(${ZLIB_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main)
+  ocv_install_target(${ZLIB_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev)
 endif()

CMakeLists.txt

@@ -197,7 +197,7 @@ OCV_OPTION(INSTALL_C_EXAMPLES       "Install C examples"        OFF )
 OCV_OPTION(INSTALL_PYTHON_EXAMPLES  "Install Python examples"   OFF )
 OCV_OPTION(INSTALL_ANDROID_EXAMPLES "Install Android examples"  OFF IF ANDROID )
 OCV_OPTION(INSTALL_TO_MANGLED_PATHS "Enables mangled install paths, that help with side by side installs." OFF IF (UNIX AND NOT ANDROID AND NOT IOS AND BUILD_SHARED_LIBS) )
-
+OCV_OPTION(INSTALL_TESTS            "Install accuracy and performance test binaries and test data" OFF)
 
 # OpenCV build options
 # ===================================================
@@ -268,13 +268,27 @@ if(WIN32)
     message(STATUS "Can't detect runtime and/or arch")
     set(OpenCV_INSTALL_BINARIES_PREFIX "")
   endif()
+elseif(ANDROID)
+  set(OpenCV_INSTALL_BINARIES_PREFIX "sdk/native/")
 else()
   set(OpenCV_INSTALL_BINARIES_PREFIX "")
 endif()
 
-set(OPENCV_SAMPLES_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}samples")
+if(ANDROID)
+  set(OPENCV_SAMPLES_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}samples/${ANDROID_NDK_ABI_NAME}")
+else()
+  set(OPENCV_SAMPLES_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}samples")
+endif()
 
-set(OPENCV_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}bin")
+if(ANDROID)
+  set(OPENCV_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}bin/${ANDROID_NDK_ABI_NAME}")
+else()
+  set(OPENCV_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}bin")
+endif()
+
+if(NOT OPENCV_TEST_INSTALL_PATH)
+  set(OPENCV_TEST_INSTALL_PATH "${OPENCV_BIN_INSTALL_PATH}")
+endif()
 
 if(ANDROID)
   set(LIBRARY_OUTPUT_PATH         "${OpenCV_BINARY_DIR}/lib/${ANDROID_NDK_ABI_NAME}")
@@ -283,6 +297,7 @@ if(ANDROID)
   set(OPENCV_3P_LIB_INSTALL_PATH  sdk/native/3rdparty/libs/${ANDROID_NDK_ABI_NAME})
   set(OPENCV_CONFIG_INSTALL_PATH  sdk/native/jni)
   set(OPENCV_INCLUDE_INSTALL_PATH sdk/native/jni/include)
+  set(OPENCV_SAMPLES_SRC_INSTALL_PATH samples/native)
 else()
   set(LIBRARY_OUTPUT_PATH         "${OpenCV_BINARY_DIR}/lib")
   set(3P_LIBRARY_OUTPUT_PATH      "${OpenCV_BINARY_DIR}/3rdparty/lib${LIB_SUFFIX}")
@@ -293,9 +308,11 @@ else()
       set(OPENCV_LIB_INSTALL_PATH   "${OpenCV_INSTALL_BINARIES_PREFIX}lib${LIB_SUFFIX}")
     endif()
     set(OPENCV_3P_LIB_INSTALL_PATH  "${OpenCV_INSTALL_BINARIES_PREFIX}staticlib${LIB_SUFFIX}")
+    set(OPENCV_SAMPLES_SRC_INSTALL_PATH    samples/native)
   else()
     set(OPENCV_LIB_INSTALL_PATH     lib${LIB_SUFFIX})
     set(OPENCV_3P_LIB_INSTALL_PATH  share/OpenCV/3rdparty/${OPENCV_LIB_INSTALL_PATH})
+    set(OPENCV_SAMPLES_SRC_INSTALL_PATH    share/OpenCV/samples)
   endif()
   set(OPENCV_INCLUDE_INSTALL_PATH "include")
 
@@ -558,6 +575,49 @@ include(cmake/OpenCVGenConfig.cmake)
 # Generate Info.plist for the IOS framework
 include(cmake/OpenCVGenInfoPlist.cmake)
 
+# Generate environment setup file
+if(INSTALL_TESTS AND OPENCV_TEST_DATA_PATH AND UNIX)
+  if(ANDROID)
+    get_filename_component(TEST_PATH ${OPENCV_TEST_INSTALL_PATH} DIRECTORY)
+    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)
+  else()
+    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"
+            DESTINATION ${OPENCV_TEST_INSTALL_PATH} COMPONENT tests)
+
+  endif()
+endif()
+
+if(NOT OPENCV_README_FILE)
+  if(ANDROID)
+    set(OPENCV_README_FILE ${CMAKE_CURRENT_SOURCE_DIR}/platforms/android/README.android)
+  endif()
+endif()
+
+if(NOT OPENCV_LICENSE_FILE)
+  set(OPENCV_LICENSE_FILE ${CMAKE_CURRENT_SOURCE_DIR}/LICENSE)
+endif()
+
+# for UNIX it does not make sense as LICENSE and readme will be part of the package automatically
+if(ANDROID OR NOT UNIX)
+  install(FILES ${OPENCV_LICENSE_FILE}
+        PERMISSIONS OWNER_READ GROUP_READ WORLD_READ
+        DESTINATION ${CMAKE_INSTALL_PREFIX} COMPONENT libs)
+  if(OPENCV_README_FILE)
+    install(FILES ${OPENCV_README_FILE}
+            PERMISSIONS OWNER_READ GROUP_READ WORLD_READ
+            DESTINATION ${CMAKE_INSTALL_PREFIX} COMPONENT libs)
+  endif()
+endif()
+
 # ----------------------------------------------------------------------------
 # Summary:
 # ----------------------------------------------------------------------------
@@ -975,3 +1035,9 @@ ocv_finalize_status()
 if("${CMAKE_CURRENT_SOURCE_DIR}" STREQUAL "${CMAKE_CURRENT_BINARY_DIR}")
   message(WARNING "The source directory is the same as binary directory. \"make clean\" may damage the source tree")
 endif()
+
+# ----------------------------------------------------------------------------
+# CPack stuff
+# ----------------------------------------------------------------------------
+
+include(cmake/OpenCVPackaging.cmake)

LICENSE

@@ -1,16 +1,11 @@
-IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
-
- By downloading, copying, installing or using the software you agree to this license.
- If you do not agree to this license, do not download, install,
- copy or use the software.
+By downloading, copying, installing or using the software you agree to this license.
+If you do not agree to this license, do not download, install,
+copy or use the software.
 
 
                           License Agreement
                For Open Source Computer Vision Library
-
-Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
-Copyright (C) 2008-2011, Willow Garage Inc., all rights reserved.
-Third party copyrights are property of their respective owners.
+                       (3-clause BSD License)
 
 Redistribution and use in source and binary forms, with or without modification,
 are permitted provided that the following conditions are met:
@@ -22,13 +17,14 @@ are permitted provided that the following conditions are met:
     this list of conditions and the following disclaimer in the documentation
     and/or other materials provided with the distribution.
 
-  * The name of the copyright holders may not be used to endorse or promote products
-    derived from this software without specific prior written permission.
+  * Neither the names of the copyright holders nor the names of the contributors
+    may be used to endorse or promote products derived from this software
+    without specific prior written permission.
 
 This software is provided by the copyright holders and contributors "as is" and
 any express or implied warranties, including, but not limited to, the implied
 warranties of merchantability and fitness for a particular purpose are disclaimed.
-In no event shall the Intel Corporation or contributors be liable for any direct,
+In no event shall copyright holders or contributors be liable for any direct,
 indirect, incidental, special, exemplary, or consequential damages
 (including, but not limited to, procurement of substitute goods or services;
 loss of use, data, or profits; or business interruption) however caused

apps/haartraining/CMakeLists.txt

@@ -71,14 +71,14 @@ set_target_properties(opencv_performance PROPERTIES
 
 if(INSTALL_CREATE_DISTRIB)
   if(BUILD_SHARED_LIBS)
-    install(TARGETS opencv_haartraining RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT main)
-    install(TARGETS opencv_createsamples RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT main)
-    install(TARGETS opencv_performance RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT main)
+    install(TARGETS opencv_haartraining RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT dev)
+    install(TARGETS opencv_createsamples RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT dev)
+    install(TARGETS opencv_performance RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT dev)
   endif()
 else()
-  install(TARGETS opencv_haartraining RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main)
-  install(TARGETS opencv_createsamples RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main)
-  install(TARGETS opencv_performance RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main)
+  install(TARGETS opencv_haartraining RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev)
+  install(TARGETS opencv_createsamples RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev)
+  install(TARGETS opencv_performance RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev)
 endif()
 
 if(ENABLE_SOLUTION_FOLDERS)

apps/traincascade/CMakeLists.txt

@@ -35,8 +35,8 @@ endif()
 
 if(INSTALL_CREATE_DISTRIB)
   if(BUILD_SHARED_LIBS)
-    install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT main)
+    install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT dev)
   endif()
 else()
-  install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main)
+  install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev)
 endif()

cmake/OpenCVDetectAndroidSDK.cmake

@@ -180,7 +180,7 @@ unset(__android_project_chain CACHE)
 # add_android_project(target_name ${path} NATIVE_DEPS opencv_core LIBRARY_DEPS ${OpenCV_BINARY_DIR} SDK_TARGET 11)
 macro(add_android_project target path)
   # parse arguments
-  set(android_proj_arglist NATIVE_DEPS LIBRARY_DEPS SDK_TARGET IGNORE_JAVA IGNORE_MANIFEST)
+  set(android_proj_arglist NATIVE_DEPS LIBRARY_DEPS SDK_TARGET IGNORE_JAVA IGNORE_MANIFEST EMBED_CUDA FORCE_EMBED_OPENCV)
   set(__varname "android_proj_")
   foreach(v ${android_proj_arglist})
     set(${__varname}${v} "")
@@ -303,6 +303,46 @@ macro(add_android_project target path)
             add_custom_command(TARGET ${JNI_LIB_NAME} POST_BUILD COMMAND ${CMAKE_STRIP} --strip-unneeded "${android_proj_jni_location}")
         endif()
       endif()
+
+      # copy opencv_java, tbb if it is shared and dynamicuda if present if FORCE_EMBED_OPENCV flag is set
+      if(android_proj_FORCE_EMBED_OPENCV)
+        set(native_deps ${android_proj_NATIVE_DEPS})
+        # filter out gpu module as it is always static library on Android
+        list(REMOVE_ITEM native_deps "opencv_gpu")
+        if(ENABLE_DYNAMIC_CUDA)
+          list(APPEND native_deps "opencv_dynamicuda")
+        endif()
+        foreach(lib ${native_deps})
+          get_property(f TARGET ${lib} PROPERTY LOCATION)
+          get_filename_component(f_name ${f} NAME)
+          add_custom_command(
+            OUTPUT "${android_proj_bin_dir}/libs/${ANDROID_NDK_ABI_NAME}/${f_name}"
+            COMMAND ${CMAKE_COMMAND} -E copy "${f}" "${android_proj_bin_dir}/libs/${ANDROID_NDK_ABI_NAME}/${f_name}"
+            DEPENDS "${lib}" VERBATIM
+            COMMENT "Embedding ${f}")
+            list(APPEND android_proj_file_deps "${android_proj_bin_dir}/libs/${ANDROID_NDK_ABI_NAME}/${f_name}")
+        endforeach()
+      endif()
+
+      # copy all needed CUDA libs to project if EMBED_CUDA flag is present
+      if(android_proj_EMBED_CUDA)
+        set(android_proj_culibs ${CUDA_npp_LIBRARY} ${CUDA_LIBRARIES})
+        if(HAVE_CUFFT)
+          list(INSERT android_proj_culibs 0 ${CUDA_cufft_LIBRARY})
+        endif()
+        if(HAVE_CUBLAS)
+          list(INSERT android_proj_culibs 0 ${CUDA_cublas_LIBRARY})
+        endif()
+        foreach(lib ${android_proj_culibs})
+          get_filename_component(f "${lib}" NAME)
+          add_custom_command(
+            OUTPUT "${android_proj_bin_dir}/libs/${ANDROID_NDK_ABI_NAME}/${f}"
+            COMMAND ${CMAKE_COMMAND} -E copy "${lib}" "${android_proj_bin_dir}/libs/${ANDROID_NDK_ABI_NAME}/${f}"
+            DEPENDS "${lib}" VERBATIM
+            COMMENT "Embedding ${f}")
+          list(APPEND android_proj_file_deps "${android_proj_bin_dir}/libs/${ANDROID_NDK_ABI_NAME}/${f}")
+        endforeach()
+      endif()
     endif()
 
     # build java part
@@ -344,20 +384,20 @@ macro(add_android_project target path)
     add_custom_command(TARGET ${target} POST_BUILD COMMAND ${CMAKE_COMMAND} -E copy "${android_proj_bin_dir}/bin/${target}-debug.apk" "${OpenCV_BINARY_DIR}/bin/${target}.apk")
     if(INSTALL_ANDROID_EXAMPLES AND "${target}" MATCHES "^example-")
       #apk
-      install(FILES "${OpenCV_BINARY_DIR}/bin/${target}.apk" DESTINATION "samples" COMPONENT main)
+      install(FILES "${OpenCV_BINARY_DIR}/bin/${target}.apk" DESTINATION "samples" COMPONENT samples)
       get_filename_component(sample_dir "${path}" NAME)
       #java part
       list(REMOVE_ITEM android_proj_files ${ANDROID_MANIFEST_FILE})
       foreach(f ${android_proj_files} ${ANDROID_MANIFEST_FILE})
         get_filename_component(install_subdir "${f}" PATH)
-        install(FILES "${android_proj_bin_dir}/${f}" DESTINATION "samples/${sample_dir}/${install_subdir}" COMPONENT main)
+        install(FILES "${android_proj_bin_dir}/${f}" DESTINATION "samples/${sample_dir}/${install_subdir}" COMPONENT samples)
       endforeach()
       #jni part + eclipse files
       file(GLOB_RECURSE jni_files RELATIVE "${path}" "${path}/jni/*" "${path}/.cproject")
       ocv_list_filterout(jni_files "\\\\.svn")
       foreach(f ${jni_files} ".classpath" ".project" ".settings/org.eclipse.jdt.core.prefs")
         get_filename_component(install_subdir "${f}" PATH)
-        install(FILES "${path}/${f}" DESTINATION "samples/${sample_dir}/${install_subdir}" COMPONENT main)
+        install(FILES "${path}/${f}" DESTINATION "samples/${sample_dir}/${install_subdir}" COMPONENT samples)
       endforeach()
       #update proj
       if(android_proj_lib_deps_commands)
@@ -365,9 +405,9 @@ macro(add_android_project target path)
       endif()
       install(CODE "EXECUTE_PROCESS(COMMAND ${ANDROID_EXECUTABLE} --silent update project --path . --target \"${android_proj_sdk_target}\" --name \"${target}\" ${inst_lib_opt}
                                     WORKING_DIRECTORY \"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/samples/${sample_dir}\"
-                                   )"  COMPONENT main)
+                                   )"  COMPONENT samples)
       #empty 'gen'
-      install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/samples/${sample_dir}/gen\")" COMPONENT main)
+      install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/samples/${sample_dir}/gen\")" COMPONENT samples)
     endif()
   endif()
 endmacro()

cmake/OpenCVGenAndroidMK.cmake

@@ -59,6 +59,24 @@ if(ANDROID)
   ocv_list_filterout(OPENCV_EXTRA_COMPONENTS_CONFIGMAKE "libcu")
   ocv_list_filterout(OPENCV_EXTRA_COMPONENTS_CONFIGMAKE "libnpp")
 
+  if(HAVE_CUDA)
+    # CUDA runtime libraries and are required always
+    set(culibs ${CUDA_LIBRARIES})
+
+    # right now NPP is requared always too
+    list(INSERT culibs 0 ${CUDA_npp_LIBRARY})
+
+    if(HAVE_CUFFT)
+      list(INSERT culibs 0 ${CUDA_cufft_LIBRARY})
+    endif()
+
+    if(HAVE_CUBLAS)
+      list(INSERT culibs 0 ${CUDA_cublas_LIBRARY})
+    endif()
+  endif()
+
+  ocv_convert_to_lib_name(CUDA_RUNTIME_LIBS_CONFIGMAKE ${culibs})
+
   # split 3rdparty libs and modules
   foreach(mod ${OPENCV_MODULES_CONFIGMAKE})
     if(NOT mod MATCHES "^opencv_.+$")
@@ -69,6 +87,10 @@ if(ANDROID)
     list(REMOVE_ITEM OPENCV_MODULES_CONFIGMAKE ${OPENCV_3RDPARTY_COMPONENTS_CONFIGMAKE})
   endif()
 
+  if(ENABLE_DYNAMIC_CUDA)
+    set(OPENCV_DYNAMICUDA_MODULE_CONFIGMAKE "dynamicuda")
+  endif()
+
   # GPU module enabled separately
   list(REMOVE_ITEM OPENCV_MODULES_CONFIGMAKE "opencv_gpu")
   list(REMOVE_ITEM OPENCV_MODULES_CONFIGMAKE "opencv_dynamicuda")
@@ -84,6 +106,7 @@ if(ANDROID)
     string(REPLACE ";" " " ${lst} "${${lst}}")
   endforeach()
   string(REPLACE "opencv_" "" OPENCV_MODULES_CONFIGMAKE "${OPENCV_MODULES_CONFIGMAKE}")
+  string(REPLACE ";" " " CUDA_RUNTIME_LIBS_CONFIGMAKE "${CUDA_RUNTIME_LIBS_CONFIGMAKE}")
 
   # prepare 3rd-party component list without TBB for armeabi and mips platforms. TBB is useless there.
   set(OPENCV_3RDPARTY_COMPONENTS_CONFIGMAKE_NO_TBB ${OPENCV_3RDPARTY_COMPONENTS_CONFIGMAKE})
@@ -116,5 +139,5 @@ if(ANDROID)
   set(OPENCV_3RDPARTY_LIBS_DIR_CONFIGCMAKE "\$(OPENCV_THIS_DIR)/../3rdparty/libs/\$(OPENCV_TARGET_ARCH_ABI)")
 
   configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCV.mk.in" "${CMAKE_BINARY_DIR}/unix-install/OpenCV.mk" IMMEDIATE @ONLY)
-  install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCV.mk DESTINATION ${OPENCV_CONFIG_INSTALL_PATH})
+  install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCV.mk DESTINATION ${OPENCV_CONFIG_INSTALL_PATH} COMPONENT dev)
 endif(ANDROID)

cmake/OpenCVGenConfig.cmake

@@ -109,18 +109,18 @@ if(UNIX) # ANDROID configuration is created here also
   #                <prefix>/(share|lib)/<name>*/                           (U)
   #                <prefix>/(share|lib)/<name>*/(cmake|CMake)/             (U)
   if(INSTALL_TO_MANGLED_PATHS)
-    install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/)
-    install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/)
-    install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ FILE OpenCVModules${modules_file_suffix}.cmake)
+    install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ COMPONENT dev)
+    install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ COMPONENT dev)
+    install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ FILE OpenCVModules${modules_file_suffix}.cmake COMPONENT dev)
   else()
-    install(FILES "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/)
-    install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/)
-    install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ FILE OpenCVModules${modules_file_suffix}.cmake)
+    install(FILES "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ COMPONENT dev)
+    install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ COMPONENT dev)
+    install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ FILE OpenCVModules${modules_file_suffix}.cmake COMPONENT dev)
   endif()
 endif()
 
 if(ANDROID)
-  install(FILES "${OpenCV_SOURCE_DIR}/platforms/android/android.toolchain.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/)
+  install(FILES "${OpenCV_SOURCE_DIR}/platforms/android/android.toolchain.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ COMPONENT dev)
 endif()
 
 # --------------------------------------------------------------------------------------------
@@ -134,12 +134,12 @@ if(WIN32)
   configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig.cmake.in" "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" IMMEDIATE @ONLY)
   configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig-version.cmake.in" "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" IMMEDIATE @ONLY)
   if(BUILD_SHARED_LIBS)
-    install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}lib")
-    install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}lib" FILE OpenCVModules${modules_file_suffix}.cmake)
+    install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}lib" COMPONENT dev)
+    install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}lib" FILE OpenCVModules${modules_file_suffix}.cmake COMPONENT dev)
   else()
-    install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}staticlib")
-    install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}staticlib" FILE OpenCVModules${modules_file_suffix}.cmake)
+    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}")
-  install(FILES "${OpenCV_SOURCE_DIR}/cmake/OpenCVConfig.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}/")
+  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)
 endif()

cmake/OpenCVGenHeaders.cmake

@@ -23,4 +23,4 @@ set(OPENCV_MODULE_DEFINITIONS_CONFIGMAKE "${OPENCV_MODULE_DEFINITIONS_CONFIGMAKE
 #endforeach()
 
 configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/opencv_modules.hpp.in" "${OPENCV_CONFIG_FILE_INCLUDE_DIR}/opencv2/opencv_modules.hpp")
-install(FILES "${OPENCV_CONFIG_FILE_INCLUDE_DIR}/opencv2/opencv_modules.hpp" DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2 COMPONENT main)
+install(FILES "${OPENCV_CONFIG_FILE_INCLUDE_DIR}/opencv2/opencv_modules.hpp" DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2 COMPONENT dev)

cmake/OpenCVGenPkgconfig.cmake

@@ -81,5 +81,5 @@ configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/opencv-XXX.pc.in"
                @ONLY IMMEDIATE)
 
 if(UNIX AND NOT ANDROID)
-  install(FILES ${CMAKE_BINARY_DIR}/unix-install/${OPENCV_PC_FILE_NAME} DESTINATION ${OPENCV_LIB_INSTALL_PATH}/pkgconfig)
+  install(FILES ${CMAKE_BINARY_DIR}/unix-install/${OPENCV_PC_FILE_NAME} DESTINATION ${OPENCV_LIB_INSTALL_PATH}/pkgconfig COMPONENT dev)
 endif()

cmake/OpenCVModule.cmake

@@ -577,9 +577,9 @@ macro(ocv_create_module)
   endif()
 
   ocv_install_target(${the_module} EXPORT OpenCVModules
-    RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main
-    LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main
-    ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main
+    RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT libs
+    LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs
+    ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT dev
     )
 
   # only "public" headers need to be installed
@@ -587,7 +587,7 @@ macro(ocv_create_module)
     foreach(hdr ${OPENCV_MODULE_${the_module}_HEADERS})
       string(REGEX REPLACE "^.*opencv2/" "opencv2/" hdr2 "${hdr}")
       if(hdr2 MATCHES "^(opencv2/.*)/[^/]+.h(..)?$")
-        install(FILES ${hdr} DESTINATION "${OPENCV_INCLUDE_INSTALL_PATH}/${CMAKE_MATCH_1}" COMPONENT main)
+        install(FILES ${hdr} DESTINATION "${OPENCV_INCLUDE_INSTALL_PATH}/${CMAKE_MATCH_1}" COMPONENT dev)
       endif()
     endforeach()
   endif()
@@ -711,6 +711,9 @@ function(ocv_add_perf_tests)
     else(OCV_DEPENDENCIES_FOUND)
       # TODO: warn about unsatisfied dependencies
     endif(OCV_DEPENDENCIES_FOUND)
+    if(INSTALL_TESTS)
+      install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_TEST_INSTALL_PATH} COMPONENT tests)
+    endif()
   endif()
 endfunction()
 
@@ -764,6 +767,10 @@ function(ocv_add_accuracy_tests)
     else(OCV_DEPENDENCIES_FOUND)
       # TODO: warn about unsatisfied dependencies
     endif(OCV_DEPENDENCIES_FOUND)
+
+    if(INSTALL_TESTS)
+      install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_TEST_INSTALL_PATH} COMPONENT tests)
+    endif()
   endif()
 endfunction()
 
@@ -795,7 +802,7 @@ function(ocv_add_samples)
         endif()
 
         if(WIN32)
-          install(TARGETS ${the_target} RUNTIME DESTINATION "samples/${module_id}" COMPONENT main)
+          install(TARGETS ${the_target} RUNTIME DESTINATION "samples/${module_id}" COMPONENT samples)
         endif()
       endforeach()
     endif()
@@ -804,8 +811,8 @@ function(ocv_add_samples)
   if(INSTALL_C_EXAMPLES AND NOT WIN32 AND EXISTS "${samples_path}")
     file(GLOB sample_files "${samples_path}/*")
     install(FILES ${sample_files}
-            DESTINATION share/OpenCV/samples/${module_id}
-            PERMISSIONS OWNER_READ GROUP_READ WORLD_READ)
+            DESTINATION ${OPENCV_SAMPLES_SRC_INSTALL_PATH}/${module_id}
+            PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples)
   endif()
 endfunction()
 

cmake/OpenCVPackaging.cmake

@@ -0,0 +1,110 @@
+if(EXISTS "${CMAKE_ROOT}/Modules/CPack.cmake")
+set(CPACK_set_DESTDIR "on")
+
+if(NOT OPENCV_CUSTOM_PACKAGE_INFO)
+  set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "Open Computer Vision Library")
+  set(CPACK_PACKAGE_DESCRIPTION
+"OpenCV (Open Source Computer Vision Library) is an open source computer vision
+and machine learning software library. OpenCV was built to provide a common
+infrastructure for computer vision applications and to accelerate the use of
+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_VERSION_MAJOR "${OPENCV_VERSION_MAJOR}")
+  set(CPACK_PACKAGE_VERSION_MINOR "${OPENCV_VERSION_MINOR}")
+  set(CPACK_PACKAGE_VERSION_PATCH "${OPENCV_VERSION_PATCH}")
+  set(CPACK_PACKAGE_VERSION "${OPENCV_VCSVERSION}")
+endif(NOT OPENCV_CUSTOM_PACKAGE_INFO)
+
+#arch
+if(X86)
+  set(CPACK_DEBIAN_ARCHITECTURE "i386")
+  set(CPACK_RPM_PACKAGE_ARCHITECTURE "i686")
+elseif(X86_64)
+  set(CPACK_DEBIAN_ARCHITECTURE "amd64")
+  set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64")
+elseif(ARM)
+  set(CPACK_DEBIAN_ARCHITECTURE "armhf")
+  set(CPACK_RPM_PACKAGE_ARCHITECTURE "armhf")
+else()
+  set(CPACK_DEBIAN_ARCHITECTURE ${CMAKE_SYSTEM_PROCESSOR})
+  set(CPACK_RPM_PACKAGE_ARCHITECTURE ${CMAKE_SYSTEM_PROCESSOR})
+endif()
+
+if(CPACK_GENERATOR STREQUAL "DEB")
+  set(OPENCV_PACKAGE_ARCH_SUFFIX ${CPACK_DEBIAN_ARCHITECTURE})
+elseif(CPACK_GENERATOR STREQUAL "RPM")
+  set(OPENCV_PACKAGE_ARCH_SUFFIX ${CPACK_RPM_PACKAGE_ARCHITECTURE})
+else()
+  set(OPENCV_PACKAGE_ARCH_SUFFIX ${CMAKE_SYSTEM_PROCESSOR})
+endif()
+
+set(CPACK_PACKAGE_FILE_NAME "${CMAKE_PROJECT_NAME}-${OPENCV_VCSVERSION}-${OPENCV_PACKAGE_ARCH_SUFFIX}")
+set(CPACK_SOURCE_PACKAGE_FILE_NAME "${CMAKE_PROJECT_NAME}-${OPENCV_VCSVERSION}-${OPENCV_PACKAGE_ARCH_SUFFIX}")
+
+#rpm options
+set(CPACK_RPM_COMPONENT_INSTALL TRUE)
+set(CPACK_RPM_PACKAGE_SUMMARY ${CPACK_PACKAGE_DESCRIPTION_SUMMARY})
+set(CPACK_RPM_PACKAGE_DESCRIPTION ${CPACK_PACKAGE_DESCRIPTION})
+set(CPACK_RPM_PACKAGE_URL "http://opencv.org")
+set(CPACK_RPM_PACKAGE_LICENSE "BSD")
+
+#deb options
+set(CPACK_DEB_COMPONENT_INSTALL TRUE)
+set(CPACK_DEBIAN_PACKAGE_PRIORITY "optional")
+set(CPACK_DEBIAN_PACKAGE_SECTION "libs")
+set(CPACK_DEBIAN_PACKAGE_HOMEPAGE "http://opencv.org")
+
+#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_COMPONENT_tests_DEPENDS libs)
+
+if(HAVE_CUDA)
+  string(REPLACE "." "-" cuda_version_suffix ${CUDA_VERSION})
+  set(CPACK_DEB_libs_PACKAGE_DEPENDS "cuda-core-libs-${cuda_version_suffix}, cuda-extra-libs-${cuda_version_suffix}")
+  set(CPACK_COMPONENT_dev_DEPENDS libs)
+  set(CPACK_DEB_dev_PACKAGE_DEPENDS "cuda-headers-${cuda_version_suffix}")
+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_python_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-python")
+  set(CPACK_COMPONENT_python_DESCRIPTION "Python bindings for Open Source Computer Vision Library")
+
+  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_dev_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-dev")
+  set(CPACK_COMPONENT_dev_DESCRIPTION "Development files for Open Source Computer Vision Library")
+
+  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_samples_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-samples")
+  set(CPACK_COMPONENT_samples_DESCRIPTION "Samples for Open Source Computer Vision Library")
+
+  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")
+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)
+
+include(CPack)
+
+ENDif(EXISTS "${CMAKE_ROOT}/Modules/CPack.cmake")

cmake/OpenCVUtils.cmake

@@ -448,6 +448,20 @@ macro(ocv_convert_to_full_paths VAR)
 endmacro()
 
 
+# convert list of paths to libraries names without lib prefix
+macro(ocv_convert_to_lib_name var)
+  set(__tmp "")
+  foreach(path ${ARGN})
+    get_filename_component(__tmp_name "${path}" NAME_WE)
+    string(REGEX REPLACE "^lib" "" __tmp_name ${__tmp_name})
+    list(APPEND __tmp "${__tmp_name}")
+  endforeach()
+  set(${var} ${__tmp})
+  unset(__tmp)
+  unset(__tmp_name)
+endmacro()
+
+
 # add install command
 function(ocv_install_target)
   install(TARGETS ${ARGN})

cmake/templates/OpenCV.mk.in

@@ -2,6 +2,13 @@
 # you might need to define NDK_USE_CYGPATH=1 before calling the ndk-build
 
 USER_LOCAL_PATH:=$(LOCAL_PATH)
+
+USER_LOCAL_C_INCLUDES:=$(LOCAL_C_INCLUDES)
+USER_LOCAL_CFLAGS:=$(LOCAL_CFLAGS)
+USER_LOCAL_STATIC_LIBRARIES:=$(LOCAL_STATIC_LIBRARIES)
+USER_LOCAL_SHARED_LIBRARIES:=$(LOCAL_SHARED_LIBRARIES)
+USER_LOCAL_LDLIBS:=$(LOCAL_LDLIBS)
+
 LOCAL_PATH:=$(subst ?,,$(firstword ?$(subst \, ,$(subst /, ,$(call my-dir)))))
 
 OPENCV_TARGET_ARCH_ABI:=$(TARGET_ARCH_ABI)
@@ -13,7 +20,7 @@ OPENCV_BASEDIR:=@OPENCV_BASE_INCLUDE_DIR_CONFIGCMAKE@
 OPENCV_LOCAL_C_INCLUDES:=@OPENCV_INCLUDE_DIRS_CONFIGCMAKE@
 OPENCV_MODULES:=@OPENCV_MODULES_CONFIGMAKE@
 
-OPENCV_HAVE_GPU_MODULE=@OPENCV_HAVE_GPU_MODULE_CONFIGMAKE@
+OPENCV_HAVE_GPU_MODULE:=@OPENCV_HAVE_GPU_MODULE_CONFIGMAKE@
 OPENCV_USE_GPU_MODULE:=
 
 ifeq ($(TARGET_ARCH_ABI),armeabi-v7a)
@@ -22,9 +29,12 @@ ifeq ($(TARGET_ARCH_ABI),armeabi-v7a)
             OPENCV_USE_GPU_MODULE:=on
         endif
     endif
+    OPENCV_DYNAMICUDA_MODULE:=@OPENCV_DYNAMICUDA_MODULE_CONFIGMAKE@
+else
+    OPENCV_DYNAMICUDA_MODULE:=
 endif
 
-CUDA_RUNTIME_LIBS:=cufft npps nppi nppc cudart
+CUDA_RUNTIME_LIBS:=@CUDA_RUNTIME_LIBS_CONFIGMAKE@
 
 ifeq ($(OPENCV_LIB_TYPE),)
     OPENCV_LIB_TYPE:=@OPENCV_LIBTYPE_CONFIGMAKE@
@@ -60,7 +70,7 @@ else
     endif
 endif
 
-ifeq (${OPENCV_CAMERA_MODULES},on)
+ifeq ($(OPENCV_CAMERA_MODULES),on)
     ifeq ($(TARGET_ARCH_ABI),armeabi)
         OPENCV_CAMERA_MODULES:=@OPENCV_CAMERA_LIBS_ARMEABI_CONFIGCMAKE@
     endif
@@ -91,6 +101,13 @@ define add_opencv_module
     include $(PREBUILT_$(OPENCV_LIB_TYPE)_LIBRARY)
 endef
 
+define add_cuda_module
+    include $(CLEAR_VARS)
+    LOCAL_MODULE:=$1
+    LOCAL_SRC_FILES:=$(CUDA_TOOLKIT_DIR)/targets/armv7-linux-androideabi/lib/lib$1.so
+    include $(PREBUILT_SHARED_LIBRARY)
+endef
+
 define add_opencv_3rdparty_component
     include $(CLEAR_VARS)
     LOCAL_MODULE:=$1
@@ -108,6 +125,17 @@ endef
 ifeq ($(OPENCV_MK_$(OPENCV_TARGET_ARCH_ABI)_ALREADY_INCLUDED),)
     ifeq ($(OPENCV_INSTALL_MODULES),on)
         $(foreach module,$(OPENCV_LIBS),$(eval $(call add_opencv_module,$(module))))
+        ifneq ($(OPENCV_DYNAMICUDA_MODULE),)
+            ifeq ($(OPENCV_LIB_TYPE),SHARED)
+              $(eval $(call add_opencv_module,$(OPENCV_DYNAMICUDA_MODULE)))
+            endif
+        endif
+    endif
+
+    ifeq ($(OPENCV_USE_GPU_MODULE),on)
+        ifeq ($(INSTALL_CUDA_LIBRARIES),on)
+            $(foreach module,$(CUDA_RUNTIME_LIBS),$(eval $(call add_cuda_module,$(module))))
+        endif
     endif
 
     $(foreach module,$(OPENCV_3RDPARTY_COMPONENTS),$(eval $(call add_opencv_3rdparty_component,$(module))))
@@ -136,6 +164,13 @@ ifeq ($(OPENCV_LOCAL_CFLAGS),)
 endif
 
 include $(CLEAR_VARS)
+
+LOCAL_C_INCLUDES:=$(USER_LOCAL_C_INCLUDES)
+LOCAL_CFLAGS:=$(USER_LOCAL_CFLAGS)
+LOCAL_STATIC_LIBRARIES:=$(USER_LOCAL_STATIC_LIBRARIES)
+LOCAL_SHARED_LIBRARIES:=$(USER_LOCAL_SHARED_LIBRARIES)
+LOCAL_LDLIBS:=$(USER_LOCAL_LDLIBS)
+
 LOCAL_C_INCLUDES += $(OPENCV_LOCAL_C_INCLUDES)
 LOCAL_CFLAGS     += $(OPENCV_LOCAL_CFLAGS)
 
@@ -145,6 +180,11 @@ endif
 
 ifeq ($(OPENCV_INSTALL_MODULES),on)
     LOCAL_$(OPENCV_LIB_TYPE)_LIBRARIES += $(foreach mod, $(OPENCV_LIBS), opencv_$(mod))
+    ifeq ($(OPENCV_LIB_TYPE),SHARED)
+        ifneq ($(OPENCV_DYNAMICUDA_MODULE),)
+            LOCAL_$(OPENCV_LIB_TYPE)_LIBRARIES += $(OPENCV_DYNAMICUDA_MODULE)
+        endif
+    endif
 else
     LOCAL_LDLIBS += -L$(call host-path,$(LOCAL_PATH)/$(OPENCV_LIBS_DIR)) $(foreach lib, $(OPENCV_LIBS), -lopencv_$(lib))
 endif
@@ -156,8 +196,12 @@ endif
 LOCAL_LDLIBS += $(foreach lib,$(OPENCV_EXTRA_COMPONENTS), -l$(lib))
 
 ifeq ($(OPENCV_USE_GPU_MODULE),on)
+    ifeq ($(INSTALL_CUDA_LIBRARIES),on)
+        LOCAL_SHARED_LIBRARIES += $(foreach mod, $(CUDA_RUNTIME_LIBS), $(mod))
+    else
+        LOCAL_LDLIBS += -L$(CUDA_TOOLKIT_DIR)/targets/armv7-linux-androideabi/lib $(foreach lib, $(CUDA_RUNTIME_LIBS), -l$(lib))
+    endif
     LOCAL_STATIC_LIBRARIES+=libopencv_gpu
-    LOCAL_LDLIBS += -L$(CUDA_TOOLKIT_DIR)/lib $(foreach lib, $(CUDA_RUNTIME_LIBS), -l$(lib))
 endif
 
 #restore the LOCAL_PATH

cmake/templates/opencv_run_all_tests_android.sh.in

@@ -0,0 +1,51 @@
+#!/bin/sh
+
+BASE_DIR=`dirname $0`
+OPENCV_TEST_PATH=$BASE_DIR/@TEST_PATH@
+OPENCV_TEST_DATA_PATH=$BASE_DIR/sdk/etc/testdata/
+
+if [ $# -ne 1 ]; then
+  echo "Device architecture is not preset in command line"
+  echo "Tests are available for architectures: `ls -m ${OPENCV_TEST_PATH}`"
+  echo "Usage: $0 <target_device_arch>"
+  return 1
+else
+  TARGET_ARCH=$1
+fi
+
+if [ -z `which adb` ]; then
+  echo "adb command was not found in PATH"
+  return 1
+fi
+
+adb push $OPENCV_TEST_DATA_PATH /sdcard/opencv_testdata
+
+adb shell "mkdir -p /data/local/tmp/opencv_test"
+SUMMARY_STATUS=0
+for t in "$OPENCV_TEST_PATH/$TARGET_ARCH/"opencv_test_* "$OPENCV_TEST_PATH/$TARGET_ARCH/"opencv_perf_*;
+do
+  test_name=`basename "$t"`
+  report="$test_name-`date --rfc-3339=date`.xml"
+  adb push $t /data/local/tmp/opencv_test/
+  adb shell "export OPENCV_TEST_DATA_PATH=/sdcard/opencv_testdata && /data/local/tmp/opencv_test/$test_name --perf_min_samples=1 --perf_force_samples=1 --gtest_output=xml:/data/local/tmp/opencv_test/$report"
+  adb pull "/data/local/tmp/opencv_test/$report" $report
+  TEST_STATUS=0
+  if [ -e $report ]; then
+    if [ `grep -c "<fail" $report` -ne 0 ]; then
+      TEST_STATUS=2
+    fi
+  else
+    TEST_STATUS=3
+  fi
+  if [ $TEST_STATUS -ne 0 ]; then
+    SUMMARY_STATUS=$TEST_STATUS
+  fi
+done
+
+if [ $SUMMARY_STATUS -eq 0 ]; then
+  echo "All OpenCV tests finished successfully"
+else
+  echo "OpenCV tests finished with status $SUMMARY_STATUS"
+fi
+
+return $SUMMARY_STATUS

cmake/templates/opencv_run_all_tests_unix.sh.in

@@ -0,0 +1,25 @@
+#!/bin/sh
+
+OPENCV_TEST_PATH=@CMAKE_INSTALL_PREFIX@/@OPENCV_TEST_INSTALL_PATH@
+export OPENCV_TEST_DATA_PATH=@CMAKE_INSTALL_PREFIX@/share/OpenCV/testdata
+
+SUMMARY_STATUS=0
+for t in "$OPENCV_TEST_PATH/"opencv_test_* "$OPENCV_TEST_PATH/"opencv_perf_*;
+do
+  report="`basename "$t"`-`date --rfc-3339=date`.xml"
+  "$t" --perf_min_samples=1 --perf_force_samples=1 --gtest_output=xml:"$report"
+  TEST_STATUS=$?
+  if [ $TEST_STATUS -ne 0 ]; then
+    SUMMARY_STATUS=$TEST_STATUS
+  fi
+done
+
+rm -f /tmp/__opencv_temp.*
+
+if [ $SUMMARY_STATUS -eq 0 ]; then
+  echo "All OpenCV tests finished successfully"
+else
+  echo "OpenCV tests finished with status $SUMMARY_STATUS"
+fi
+
+return $SUMMARY_STATUS

cmake/templates/opencv_testing.sh.in

@@ -0,0 +1,2 @@
+# Environment setup for OpenCV testing
+export OPENCV_TEST_DATA_PATH=@CMAKE_INSTALL_PREFIX@/share/OpenCV/testdata

data/CMakeLists.txt

@@ -2,9 +2,21 @@ file(GLOB HAAR_CASCADES haarcascades/*.xml)
 file(GLOB LBP_CASCADES lbpcascades/*.xml)
 
 if(ANDROID)
-  install(FILES ${HAAR_CASCADES} DESTINATION sdk/etc/haarcascades COMPONENT main)
-  install(FILES ${LBP_CASCADES}  DESTINATION sdk/etc/lbpcascades  COMPONENT main)
+  install(FILES ${HAAR_CASCADES} DESTINATION sdk/etc/haarcascades COMPONENT libs)
+  install(FILES ${LBP_CASCADES}  DESTINATION sdk/etc/lbpcascades  COMPONENT libs)
 elseif(NOT WIN32)
-  install(FILES ${HAAR_CASCADES} DESTINATION share/OpenCV/haarcascades COMPONENT main)
-  install(FILES ${LBP_CASCADES}  DESTINATION share/OpenCV/lbpcascades  COMPONENT main)
+  install(FILES ${HAAR_CASCADES} DESTINATION share/OpenCV/haarcascades COMPONENT libs)
+  install(FILES ${LBP_CASCADES}  DESTINATION share/OpenCV/lbpcascades  COMPONENT libs)
 endif()
+
+if(INSTALL_TESTS AND OPENCV_TEST_DATA_PATH)
+  if(ANDROID)
+    install(DIRECTORY ${OPENCV_TEST_DATA_PATH} DESTINATION sdk/etc/testdata COMPONENT tests)
+  elseif(NOT WIN32)
+    # CPack does not set correct permissions by default, so we do it explicitly.
+    install(DIRECTORY ${OPENCV_TEST_DATA_PATH}
+            DIRECTORY_PERMISSIONS OWNER_WRITE OWNER_READ OWNER_EXECUTE
+              GROUP_READ GROUP_EXECUTE WORLD_READ WORLD_EXECUTE
+            DESTINATION share/OpenCV/testdata COMPONENT tests)
+  endif()
+endif()

doc/CMakeLists.txt

@@ -143,11 +143,11 @@ if(BUILD_DOCS AND HAVE_SPHINX)
   endif()
 
   foreach(f ${DOC_LIST})
-    install(FILES "${f}" DESTINATION "${OPENCV_DOC_INSTALL_PATH}" COMPONENT main)
+    install(FILES "${f}" DESTINATION "${OPENCV_DOC_INSTALL_PATH}" COMPONENT docs)
   endforeach()
 
   foreach(f ${OPTIONAL_DOC_LIST})
-    install(FILES "${f}" DESTINATION "${OPENCV_DOC_INSTALL_PATH}" OPTIONAL)
+    install(FILES "${f}" DESTINATION "${OPENCV_DOC_INSTALL_PATH}" OPTIONAL COMPONENT docs)
   endforeach()
 
 endif()

doc/tutorials/introduction/android_binary_package/O4A_SDK.rst

@@ -48,10 +48,10 @@ The structure of package contents looks as follows:
 
 ::
 
-    OpenCV-2.4.8-android-sdk
+    OpenCV-2.4.8.2-android-sdk
     |_ apk
-    |   |_ OpenCV_2.4.8_binary_pack_armv7a.apk
-    |   |_ OpenCV_2.4.8_Manager_2.16_XXX.apk
+    |   |_ OpenCV_2.4.8.2_binary_pack_armv7a.apk
+    |   |_ OpenCV_2.4.8.2_Manager_2.17_XXX.apk
     |
     |_ doc
     |_ samples
@@ -66,7 +66,7 @@ The structure of package contents looks as follows:
     |               |_ armeabi-v7a
     |               |_ x86
     |
-    |_ license.txt
+    |_ LICENSE
     |_ README.android
 
 * :file:`sdk` folder contains OpenCV API and libraries for Android:
@@ -157,10 +157,10 @@ Get the OpenCV4Android SDK
 
    .. code-block:: bash
 
-      unzip ~/Downloads/OpenCV-2.4.8-android-sdk.zip
+      unzip ~/Downloads/OpenCV-2.4.8.2-android-sdk.zip
 
-.. |opencv_android_bin_pack| replace:: :file:`OpenCV-2.4.8-android-sdk.zip`
-.. _opencv_android_bin_pack_url: http://sourceforge.net/projects/opencvlibrary/files/opencv-android/2.4.8/OpenCV-2.4.8-android-sdk.zip/download
+.. |opencv_android_bin_pack| replace:: :file:`OpenCV-2.4.8.2-android-sdk.zip`
+.. _opencv_android_bin_pack_url: http://sourceforge.net/projects/opencvlibrary/files/opencv-android/2.4.8.2/OpenCV-2.4.8.2-android-sdk.zip/download
 .. |opencv_android_bin_pack_url| replace:: |opencv_android_bin_pack|
 .. |seven_zip| replace:: 7-Zip
 .. _seven_zip: http://www.7-zip.org/
@@ -295,7 +295,7 @@ Well, running samples from Eclipse is very simple:
   .. code-block:: sh
     :linenos:
 
-    <Android SDK path>/platform-tools/adb install <OpenCV4Android SDK path>/apk/OpenCV_2.4.8_Manager_2.16_armv7a-neon.apk
+    <Android SDK path>/platform-tools/adb install <OpenCV4Android SDK path>/apk/OpenCV_2.4.8.2_Manager_2.17_armv7a-neon.apk
 
   .. note:: ``armeabi``, ``armv7a-neon``, ``arm7a-neon-android8``, ``mips`` and ``x86`` stand for
             platform targets:

doc/tutorials/introduction/android_binary_package/dev_with_OCV_on_Android.rst

@@ -55,14 +55,14 @@ Manager to access OpenCV libraries externally installed in the target system.
    :guilabel:`File -> Import -> Existing project in your workspace`.
 
    Press :guilabel:`Browse`  button and locate OpenCV4Android SDK
-   (:file:`OpenCV-2.4.8-android-sdk/sdk`).
+   (:file:`OpenCV-2.4.8.2-android-sdk/sdk`).
 
    .. image:: images/eclipse_opencv_dependency0.png
         :alt: Add dependency from OpenCV library
         :align: center
 
 #. In application project add a reference to the OpenCV Java SDK in
-   :guilabel:`Project -> Properties -> Android -> Library -> Add` select ``OpenCV Library - 2.4.8``.
+   :guilabel:`Project -> Properties -> Android -> Library -> Add` select ``OpenCV Library - 2.4.8.2``.
 
    .. image:: images/eclipse_opencv_dependency1.png
         :alt: Add dependency from OpenCV library
@@ -128,27 +128,27 @@ described above.
 #. Add the OpenCV library project to your workspace the same way as for the async initialization
    above. Use menu :guilabel:`File -> Import -> Existing project in your workspace`,
    press :guilabel:`Browse` button and select OpenCV SDK path
-   (:file:`OpenCV-2.4.8-android-sdk/sdk`).
+   (:file:`OpenCV-2.4.8.2-android-sdk/sdk`).
 
    .. image:: images/eclipse_opencv_dependency0.png
         :alt: Add dependency from OpenCV library
         :align: center
 
 #. In the application project add a reference to the OpenCV4Android SDK in
-   :guilabel:`Project -> Properties -> Android -> Library -> Add` select ``OpenCV Library - 2.4.8``;
+   :guilabel:`Project -> Properties -> Android -> Library -> Add` select ``OpenCV Library - 2.4.8.2``;
 
    .. image:: images/eclipse_opencv_dependency1.png
        :alt: Add dependency from OpenCV library
        :align: center
 
 #. If your application project **doesn't have a JNI part**, just copy the corresponding OpenCV
-   native libs from :file:`<OpenCV-2.4.8-android-sdk>/sdk/native/libs/<target_arch>` to your
+   native libs from :file:`<OpenCV-2.4.8.2-android-sdk>/sdk/native/libs/<target_arch>` to your
    project directory to folder :file:`libs/<target_arch>`.
 
    In case of the application project **with a JNI part**, instead of manual libraries copying you
    need to modify your ``Android.mk`` file:
    add the following two code lines after the ``"include $(CLEAR_VARS)"`` and before
-   ``"include path_to_OpenCV-2.4.8-android-sdk/sdk/native/jni/OpenCV.mk"``
+   ``"include path_to_OpenCV-2.4.8.2-android-sdk/sdk/native/jni/OpenCV.mk"``
 
    .. code-block:: make
       :linenos:
@@ -221,7 +221,7 @@ taken:
 
    .. code-block:: make
 
-      include C:\Work\OpenCV4Android\OpenCV-2.4.8-android-sdk\sdk\native\jni\OpenCV.mk
+      include C:\Work\OpenCV4Android\OpenCV-2.4.8.2-android-sdk\sdk\native\jni\OpenCV.mk
 
    Should be inserted into the :file:`jni/Android.mk` file **after** this line:
 

doc/tutorials/introduction/linux_install/linux_install.rst

@@ -16,7 +16,7 @@ Required Packages
   * CMake 2.6 or higher;
   * Git;
   * GTK+2.x or higher, including headers (libgtk2.0-dev);
-  * pkgconfig;
+  * pkg-config;
   * Python 2.6 or later and Numpy 1.5 or later with developer packages (python-dev, python-numpy);
   * ffmpeg or libav development packages: libavcodec-dev, libavformat-dev, libswscale-dev;
   * [optional] libdc1394 2.x;

include/CMakeLists.txt

@@ -1,7 +1,7 @@
 file(GLOB old_hdrs "opencv/*.h*")
 install(FILES ${old_hdrs}
     DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv
-    COMPONENT main)
+    COMPONENT dev)
 install(FILES "opencv2/opencv.hpp"
     DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2
-    COMPONENT main)
+    COMPONENT dev)

modules/androidcamera/CMakeLists.txt

@@ -40,6 +40,6 @@ else()
     get_filename_component(wrapper_name "${wrapper}" NAME)
     install(FILES "${LIBRARY_OUTPUT_PATH}/${wrapper_name}"
             DESTINATION ${OPENCV_LIB_INSTALL_PATH}
-            COMPONENT main)
+            COMPONENT libs)
   endforeach()
 endif()

modules/androidcamera/camera_wrapper/CMakeLists.txt

@@ -63,4 +63,4 @@ if (NOT (CMAKE_BUILD_TYPE MATCHES "debug"))
 endif()
 
 
-install(TARGETS ${the_target} LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main)
+install(TARGETS ${the_target} LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs)

modules/core/include/opencv2/core/mat.hpp

@@ -2401,7 +2401,7 @@ template<typename _Tp> inline SparseMat_<_Tp>::SparseMat_(const SparseMat& m)
     if( m.type() == DataType<_Tp>::type )
         *this = (const SparseMat_<_Tp>&)m;
     else
-        m.convertTo(this, DataType<_Tp>::type);
+        m.convertTo(*this, DataType<_Tp>::type);
 }
 
 template<typename _Tp> inline SparseMat_<_Tp>::SparseMat_(const SparseMat_<_Tp>& m)

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

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

modules/core/src/algorithm.cpp

@@ -647,7 +647,7 @@ void AlgorithmInfo::set(Algorithm* algo, const char* parameter, int argType, con
             || argType == Param::FLOAT || argType == Param::UNSIGNED_INT || argType == Param::UINT64 || argType == Param::UCHAR)
     {
         if ( !( p->type == Param::INT || p->type == Param::REAL || p->type == Param::BOOLEAN
-                || p->type == Param::UNSIGNED_INT || p->type == Param::UINT64 || p->type == Param::FLOAT || argType == Param::UCHAR
+                || p->type == Param::UNSIGNED_INT || p->type == Param::UINT64 || p->type == Param::FLOAT || p->type == Param::UCHAR
                 || (p->type == Param::SHORT && argType == Param::INT)) )
         {
             string message = getErrorMessageForWrongArgumentInSetter(algo->name(), parameter, p->type, argType);

modules/core/src/arithm.cpp

@@ -1272,8 +1272,8 @@ static void arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
     bool haveScalar = false, swapped12 = false;
     int depth2 = src2.depth();
     if( src1.size != src2.size || src1.channels() != src2.channels() ||
-        ((kind1 == _InputArray::MATX || kind2 == _InputArray::MATX) &&
-         src1.cols == 1 && src2.rows == 4) )
+        (kind1 == _InputArray::MATX && (src1.size() == Size(1,4) || src1.size() == Size(1,1))) ||
+        (kind2 == _InputArray::MATX && (src2.size() == Size(1,4) || src2.size() == Size(1,1))) )
     {
         if( checkScalar(src1, src2.type(), kind1, kind2) )
         {

modules/core/src/gpumat.cpp

@@ -620,11 +620,19 @@ void cv::gpu::GpuMat::copyTo(GpuMat& m) const
 void cv::gpu::GpuMat::copyTo(GpuMat& mat, const GpuMat& mask) const
 {
     if (mask.empty())
+    {
         copyTo(mat);
+    }
     else
     {
+        uchar* data0 = mat.data;
+
         mat.create(size(), type());
 
+        // do not leave dst uninitialized
+        if (mat.data != data0)
+            mat.setTo(Scalar::all(0));
+
         gpuFuncTable()->copyWithMask(*this, mat, mask);
     }
 }

modules/core/src/precomp.hpp

@@ -129,12 +129,14 @@ template<typename T> struct OpMax
 
 inline Size getContinuousSize( const Mat& m1, int widthScale=1 )
 {
+    CV_Assert(m1.dims <= 2);
     return m1.isContinuous() ? Size(m1.cols*m1.rows*widthScale, 1) :
         Size(m1.cols*widthScale, m1.rows);
 }
 
 inline Size getContinuousSize( const Mat& m1, const Mat& m2, int widthScale=1 )
 {
+    CV_Assert(m1.dims <= 2 && m1.size() == m2.size());
     return (m1.flags & m2.flags & Mat::CONTINUOUS_FLAG) != 0 ?
         Size(m1.cols*m1.rows*widthScale, 1) : Size(m1.cols*widthScale, m1.rows);
 }
@@ -142,6 +144,7 @@ inline Size getContinuousSize( const Mat& m1, const Mat& m2, int widthScale=1 )
 inline Size getContinuousSize( const Mat& m1, const Mat& m2,
                                const Mat& m3, int widthScale=1 )
 {
+    CV_Assert(m1.dims <= 2 && m1.size() == m2.size() && m1.size() == m3.size());
     return (m1.flags & m2.flags & m3.flags & Mat::CONTINUOUS_FLAG) != 0 ?
         Size(m1.cols*m1.rows*widthScale, 1) : Size(m1.cols*widthScale, m1.rows);
 }
@@ -150,6 +153,7 @@ inline Size getContinuousSize( const Mat& m1, const Mat& m2,
                                const Mat& m3, const Mat& m4,
                                int widthScale=1 )
 {
+    CV_Assert(m1.dims <= 2 && m1.size() == m2.size() && m1.size() == m3.size() && m1.size() == m4.size());
     return (m1.flags & m2.flags & m3.flags & m4.flags & Mat::CONTINUOUS_FLAG) != 0 ?
         Size(m1.cols*m1.rows*widthScale, 1) : Size(m1.cols*widthScale, m1.rows);
 }
@@ -158,6 +162,7 @@ inline Size getContinuousSize( const Mat& m1, const Mat& m2,
                                const Mat& m3, const Mat& m4,
                                const Mat& m5, int widthScale=1 )
 {
+    CV_Assert(m1.dims <= 2 && m1.size() == m2.size() && m1.size() == m3.size() && m1.size() == m4.size() && m1.size() == m5.size());
     return (m1.flags & m2.flags & m3.flags & m4.flags & m5.flags & Mat::CONTINUOUS_FLAG) != 0 ?
         Size(m1.cols*m1.rows*widthScale, 1) : Size(m1.cols*widthScale, m1.rows);
 }

modules/core/src/system.cpp

@@ -1129,17 +1129,24 @@ public:
         }
     }
 };
-static TLSContainerStorage tlsContainerStorage;
+
+// This is a wrapper function that will ensure 'tlsContainerStorage' is constructed on first use.
+// For more information: http://www.parashift.com/c++-faq/static-init-order-on-first-use.html
+static TLSContainerStorage& getTLSContainerStorage()
+{
+    static TLSContainerStorage *tlsContainerStorage = new TLSContainerStorage();
+    return *tlsContainerStorage;
+}
 
 TLSDataContainer::TLSDataContainer()
     : key_(-1)
 {
-    key_ = tlsContainerStorage.allocateKey(this);
+    key_ = getTLSContainerStorage().allocateKey(this);
 }
 
 TLSDataContainer::~TLSDataContainer()
 {
-    tlsContainerStorage.releaseKey(key_, this);
+    getTLSContainerStorage().releaseKey(key_, this);
     key_ = -1;
 }
 
@@ -1164,7 +1171,7 @@ TLSStorage::~TLSStorage()
         void*& data = tlsData_[i];
         if (data)
         {
-            tlsContainerStorage.destroyData(i, data);
+            getTLSContainerStorage().destroyData(i, data);
             data = NULL;
         }
     }

modules/core/test/test_arithm.cpp

@@ -1564,3 +1564,19 @@ TEST(Core_round, CvRound)
     ASSERT_EQ(-2, cvRound(-2.5));
     ASSERT_EQ(-4, cvRound(-3.5));
 }
+
+
+typedef testing::TestWithParam<Size> Mul1;
+
+TEST_P(Mul1, One)
+{
+    Size size = GetParam();
+    cv::Mat src(size, CV_32FC1, cv::Scalar::all(2)), dst,
+            ref_dst(size, CV_32FC1, cv::Scalar::all(6));
+
+    cv::multiply(3, src, dst);
+
+    ASSERT_EQ(0, cv::norm(dst, ref_dst, cv::NORM_INF));
+}
+
+INSTANTIATE_TEST_CASE_P(Arithm, Mul1, testing::Values(Size(2, 2), Size(1, 1)));

modules/core/test/test_math.cpp

@@ -2392,16 +2392,14 @@ TYPED_TEST_P(Core_CheckRange, Negative)
     double min_bound = 4.5;
     double max_bound = 16.0;
 
-    TypeParam data[] = {5, 10, 15, 4, 10 ,2, 8, 12, 14};
+    TypeParam data[] = {5, 10, 15, 4, 10, 2, 8, 12, 14};
     cv::Mat src = cv::Mat(3,3, cv::DataDepth<TypeParam>::value, data);
 
-    cv::Point* bad_pt = new cv::Point(0, 0);
+    cv::Point bad_pt(0, 0);
 
-    ASSERT_FALSE(checkRange(src, true, bad_pt, min_bound, max_bound));
-    ASSERT_EQ(bad_pt->x,0);
-    ASSERT_EQ(bad_pt->y,1);
-
-    delete bad_pt;
+    ASSERT_FALSE(checkRange(src, true, &bad_pt, min_bound, max_bound));
+    ASSERT_EQ(bad_pt.x, 0);
+    ASSERT_EQ(bad_pt.y, 1);
 }
 
 TYPED_TEST_P(Core_CheckRange, Positive)
@@ -2409,16 +2407,14 @@ TYPED_TEST_P(Core_CheckRange, Positive)
     double min_bound = -1;
     double max_bound = 16.0;
 
-    TypeParam data[] = {5, 10, 15, 4, 10 ,2, 8, 12, 14};
+    TypeParam data[] = {5, 10, 15, 4, 10, 2, 8, 12, 14};
     cv::Mat src = cv::Mat(3,3, cv::DataDepth<TypeParam>::value, data);
 
-    cv::Point* bad_pt = new cv::Point(0, 0);
+    cv::Point bad_pt(0, 0);
 
-    ASSERT_TRUE(checkRange(src, true, bad_pt, min_bound, max_bound));
-    ASSERT_EQ(bad_pt->x,0);
-    ASSERT_EQ(bad_pt->y,0);
-
-    delete bad_pt;
+    ASSERT_TRUE(checkRange(src, true, &bad_pt, min_bound, max_bound));
+    ASSERT_EQ(bad_pt.x, 0);
+    ASSERT_EQ(bad_pt.y, 0);
 }
 
 TYPED_TEST_P(Core_CheckRange, Bounds)
@@ -2426,16 +2422,14 @@ TYPED_TEST_P(Core_CheckRange, Bounds)
     double min_bound = 24.5;
     double max_bound = 1.0;
 
-    TypeParam data[] = {5, 10, 15, 4, 10 ,2, 8, 12, 14};
+    TypeParam data[] = {5, 10, 15, 4, 10, 2, 8, 12, 14};
     cv::Mat src = cv::Mat(3,3, cv::DataDepth<TypeParam>::value, data);
 
-    cv::Point* bad_pt = new cv::Point(0, 0);
+    cv::Point bad_pt(0, 0);
 
-    ASSERT_FALSE(checkRange(src, true, bad_pt, min_bound, max_bound));
-    ASSERT_EQ(bad_pt->x,0);
-    ASSERT_EQ(bad_pt->y,0);
-
-    delete bad_pt;
+    ASSERT_FALSE(checkRange(src, true, &bad_pt, min_bound, max_bound));
+    ASSERT_EQ(bad_pt.x, 0);
+    ASSERT_EQ(bad_pt.y, 0);
 }
 
 TYPED_TEST_P(Core_CheckRange, Zero)

modules/dynamicuda/include/opencv2/dynamicuda/dynamicuda.hpp

@@ -129,15 +129,20 @@ public:
 
 #if defined(USE_CUDA)
 
-#define cudaSafeCall(expr)  ___cudaSafeCall(expr, __FILE__, __LINE__, CV_Func)
-#define nppSafeCall(expr)  ___nppSafeCall(expr, __FILE__, __LINE__, CV_Func)
+// Disable NPP for this file
+//#define USE_NPP
+#undef USE_NPP
 
+#define cudaSafeCall(expr)  ___cudaSafeCall(expr, __FILE__, __LINE__, CV_Func)
 inline void ___cudaSafeCall(cudaError_t err, const char *file, const int line, const char *func = "")
 {
     if (cudaSuccess != err)
         cv::gpu::error(cudaGetErrorString(err), file, line, func);
 }
 
+#ifdef USE_NPP
+
+#define nppSafeCall(expr)  ___nppSafeCall(expr, __FILE__, __LINE__, CV_Func)
 inline void ___nppSafeCall(int err, const char *file, const int line, const char *func = "")
 {
     if (err < 0)
@@ -148,6 +153,8 @@ inline void ___nppSafeCall(int err, const char *file, const int line, const char
     }
 }
 
+#endif
+
 namespace cv { namespace gpu { namespace device
 {
     void copyToWithMask_gpu(PtrStepSzb src, PtrStepSzb dst, size_t elemSize1, int cn, PtrStepSzb mask, bool colorMask, cudaStream_t stream);
@@ -173,6 +180,8 @@ template <typename T> void kernelSetCaller(GpuMat& src, Scalar s, const GpuMat&
     cv::gpu::device::set_to_gpu(src, sf.val, mask, src.channels(), stream);
 }
 
+#ifdef USE_NPP
+
 template<int n> struct NPPTypeTraits;
 template<> struct NPPTypeTraits<CV_8U>  { typedef Npp8u npp_type; };
 template<> struct NPPTypeTraits<CV_8S>  { typedef Npp8s npp_type; };
@@ -182,9 +191,13 @@ template<> struct NPPTypeTraits<CV_32S> { typedef Npp32s npp_type; };
 template<> struct NPPTypeTraits<CV_32F> { typedef Npp32f npp_type; };
 template<> struct NPPTypeTraits<CV_64F> { typedef Npp64f npp_type; };
 
+#endif
+
 //////////////////////////////////////////////////////////////////////////
 // Convert
 
+#ifdef USE_NPP
+
 template<int SDEPTH, int DDEPTH> struct NppConvertFunc
 {
     typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
@@ -232,9 +245,13 @@ template<int DDEPTH, typename NppConvertFunc<CV_32F, DDEPTH>::func_ptr func> str
     }
 };
 
+#endif
+
 //////////////////////////////////////////////////////////////////////////
 // Set
 
+#ifdef USE_NPP
+
 template<int SDEPTH, int SCN> struct NppSetFunc
 {
     typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
@@ -339,9 +356,13 @@ template<int SDEPTH, typename NppSetMaskFunc<SDEPTH, 1>::func_ptr func> struct N
     }
 };
 
+#endif
+
 //////////////////////////////////////////////////////////////////////////
 // CopyMasked
 
+#ifdef USE_NPP
+
 template<int SDEPTH> struct NppCopyMaskedFunc
 {
     typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
@@ -365,6 +386,8 @@ template<int SDEPTH, typename NppCopyMaskedFunc<SDEPTH>::func_ptr func> struct N
     }
 };
 
+#endif
+
 template <typename T> static inline bool isAligned(const T* ptr, size_t size)
 {
     return reinterpret_cast<size_t>(ptr) % size == 0;
@@ -877,6 +900,8 @@ public:
         }
 
         typedef void (*func_t)(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream);
+
+#ifdef USE_NPP
         static const func_t funcs[7][4] =
         {
             /*  8U */ {NppCopyMasked<CV_8U , nppiCopy_8u_C1MR >::call, cv::gpu::device::copyWithMask, NppCopyMasked<CV_8U , nppiCopy_8u_C3MR >::call, NppCopyMasked<CV_8U , nppiCopy_8u_C4MR >::call},
@@ -889,6 +914,9 @@ public:
          };
 
          const func_t func =  mask.channels() == src.channels() ? funcs[src.depth()][src.channels() - 1] : cv::gpu::device::copyWithMask;
+#else
+        const func_t func = cv::gpu::device::copyWithMask;
+#endif
 
          func(src, dst, mask, 0);
     }
@@ -896,6 +924,8 @@ public:
     void convert(const GpuMat& src, GpuMat& dst) const
     {
         typedef void (*func_t)(const GpuMat& src, GpuMat& dst);
+
+#ifdef USE_NPP
         static const func_t funcs[7][7][4] =
         {
             {
@@ -962,6 +992,7 @@ public:
                 /* 64F -> 64F */ {0,0,0,0}
             }
         };
+#endif
 
         CV_Assert(src.depth() <= CV_64F && src.channels() <= 4);
         CV_Assert(dst.depth() <= CV_64F);
@@ -980,8 +1011,12 @@ public:
             return;
         }
 
+#ifdef USE_NPP
         const func_t func = funcs[src.depth()][dst.depth()][src.channels() - 1];
         CV_DbgAssert(func != 0);
+#else
+        const func_t func = cv::gpu::device::convertTo;
+#endif
 
         func(src, dst);
     }
@@ -1023,6 +1058,8 @@ public:
             }
 
             typedef void (*func_t)(GpuMat& src, Scalar s);
+
+#ifdef USE_NPP
             static const func_t funcs[7][4] =
             {
                 {NppSet<CV_8U , 1, nppiSet_8u_C1R >::call, cv::gpu::device::setTo                  , cv::gpu::device::setTo                        , NppSet<CV_8U , 4, nppiSet_8u_C4R >::call},
@@ -1033,6 +1070,7 @@ public:
                 {NppSet<CV_32F, 1, nppiSet_32f_C1R>::call, cv::gpu::device::setTo                  , cv::gpu::device::setTo                        , NppSet<CV_32F, 4, nppiSet_32f_C4R>::call},
                 {cv::gpu::device::setTo                  , cv::gpu::device::setTo                  , cv::gpu::device::setTo                        , cv::gpu::device::setTo                          }
             };
+#endif
 
             CV_Assert(m.depth() <= CV_64F && m.channels() <= 4);
 
@@ -1042,14 +1080,22 @@ public:
                     CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");
             }
 
+#ifdef USE_NPP
+        const func_t func = funcs[m.depth()][m.channels() - 1];
+#else
+        const func_t func = cv::gpu::device::setTo;
+#endif
+
             if (stream)
                 cv::gpu::device::setTo(m, s, stream);
             else
-                funcs[m.depth()][m.channels() - 1](m, s);
+                func(m, s);
         }
         else
         {
             typedef void (*func_t)(GpuMat& src, Scalar s, const GpuMat& mask);
+
+#ifdef USE_NPP
             static const func_t funcs[7][4] =
             {
                 {NppSetMask<CV_8U , 1, nppiSet_8u_C1MR >::call, cv::gpu::device::setTo, cv::gpu::device::setTo, NppSetMask<CV_8U , 4, nppiSet_8u_C4MR >::call},
@@ -1060,6 +1106,7 @@ public:
                 {NppSetMask<CV_32F, 1, nppiSet_32f_C1MR>::call, cv::gpu::device::setTo, cv::gpu::device::setTo, NppSetMask<CV_32F, 4, nppiSet_32f_C4MR>::call},
                 {cv::gpu::device::setTo                       , cv::gpu::device::setTo, cv::gpu::device::setTo, cv::gpu::device::setTo                               }
             };
+#endif
 
             CV_Assert(m.depth() <= CV_64F && m.channels() <= 4);
 
@@ -1069,10 +1116,16 @@ public:
                     CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");
             }
 
+#ifdef USE_NPP
+        const func_t func = funcs[m.depth()][m.channels() - 1];
+#else
+        const func_t func = cv::gpu::device::setTo;
+#endif
+
             if (stream)
                 cv::gpu::device::setTo(m, s, mask, stream);
             else
-                funcs[m.depth()][m.channels() - 1](m, s, mask);
+                func(m, s, mask);
         }
     }
 

modules/flann/include/opencv2/flann/hierarchical_clustering_index.h

@@ -414,12 +414,6 @@ public:
 
     void loadIndex(FILE* stream)
     {
-        load_value(stream, branching_);
-        load_value(stream, trees_);
-        load_value(stream, centers_init_);
-        load_value(stream, leaf_size_);
-        load_value(stream, memoryCounter);
-
         free_elements();
 
         if (root!=NULL) {
@@ -430,6 +424,12 @@ public:
             delete[] indices;
         }
 
+        load_value(stream, branching_);
+        load_value(stream, trees_);
+        load_value(stream, centers_init_);
+        load_value(stream, leaf_size_);
+        load_value(stream, memoryCounter);
+
         indices = new int*[trees_];
         root = new NodePtr[trees_];
         for (int i=0; i<trees_; ++i) {

modules/gpu/CMakeLists.txt

@@ -82,7 +82,7 @@ ocv_create_module(${cuda_link_libs})
 if(HAVE_CUDA)
   install(FILES src/nvidia/NPP_staging/NPP_staging.hpp  src/nvidia/core/NCV.hpp
     DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2/${name}
-    COMPONENT main)
+    COMPONENT dev)
 endif()
 
 ocv_add_precompiled_headers(${the_module})

modules/gpu/perf/perf_calib3d.cpp

@@ -93,7 +93,7 @@ PERF_TEST_P(ImagePair, Calib3D_StereoBM,
 //////////////////////////////////////////////////////////////////////
 // StereoBeliefPropagation
 
-PERF_TEST_P(ImagePair, Calib3D_StereoBeliefPropagation,
+PERF_TEST_P(ImagePair, DISABLED_Calib3D_StereoBeliefPropagation,
             Values(pair_string("gpu/stereobp/aloe-L.png", "gpu/stereobp/aloe-R.png")))
 {
     declare.time(300.0);

modules/gpu/perf/perf_imgproc.cpp

@@ -851,6 +851,8 @@ PERF_TEST_P(Sz_Depth_Cn, ImgProc_BlendLinear,
     }
 }
 
+#ifdef HAVE_CUFFT
+
 //////////////////////////////////////////////////////////////////////
 // Convolve
 
@@ -1085,6 +1087,8 @@ PERF_TEST_P(Sz_Flags, ImgProc_Dft,
     }
 }
 
+#endif
+
 //////////////////////////////////////////////////////////////////////
 // CornerHarris
 
@@ -1538,7 +1542,7 @@ CV_ENUM(AlphaOp, ALPHA_OVER, ALPHA_IN, ALPHA_OUT, ALPHA_ATOP, ALPHA_XOR, ALPHA_P
 
 DEF_PARAM_TEST(Sz_Type_Op, cv::Size, MatType, AlphaOp);
 
-PERF_TEST_P(Sz_Type_Op, ImgProc_AlphaComp,
+PERF_TEST_P(Sz_Type_Op, DISABLED_ImgProc_AlphaComp,
             Combine(GPU_TYPICAL_MAT_SIZES,
                     Values(CV_8UC4, CV_16UC4, CV_32SC4, CV_32FC4),
                     AlphaOp::all()))
@@ -1559,7 +1563,14 @@ PERF_TEST_P(Sz_Type_Op, ImgProc_AlphaComp,
 
         TEST_CYCLE() cv::gpu::alphaComp(d_img1, d_img2, dst, alpha_op);
 
-        GPU_SANITY_CHECK(dst, 1e-3, ERROR_RELATIVE);
+        if (CV_MAT_DEPTH(type) < CV_32F)
+        {
+            GPU_SANITY_CHECK(dst, 1);
+        }
+        else
+        {
+            GPU_SANITY_CHECK(dst, 1e-3, ERROR_RELATIVE);
+        }
     }
     else
     {
@@ -1825,7 +1836,7 @@ CV_FLAGS(GHMethod, GHT_POSITION, GHT_SCALE, GHT_ROTATION);
 
 DEF_PARAM_TEST(Method_Sz, GHMethod, cv::Size);
 
-PERF_TEST_P(Method_Sz, ImgProc_GeneralizedHough,
+PERF_TEST_P(Method_Sz, DISABLED_ImgProc_GeneralizedHough,
             Combine(Values(GHMethod(cv::GHT_POSITION), GHMethod(cv::GHT_POSITION | cv::GHT_SCALE), GHMethod(cv::GHT_POSITION | cv::GHT_ROTATION), GHMethod(cv::GHT_POSITION | cv::GHT_SCALE | cv::GHT_ROTATION)),
                     GPU_TYPICAL_MAT_SIZES))
 {

modules/gpu/src/cuda/build_point_list.cu

@@ -0,0 +1,138 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#if !defined CUDA_DISABLER
+
+#include "opencv2/gpu/device/common.hpp"
+#include "opencv2/gpu/device/emulation.hpp"
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        __device__ static int g_counter;
+
+        template <int PIXELS_PER_THREAD>
+        __global__ void buildPointList(const PtrStepSzb src, unsigned int* list)
+        {
+            __shared__ unsigned int s_queues[4][32 * PIXELS_PER_THREAD];
+            __shared__ int s_qsize[4];
+            __shared__ int s_globStart[4];
+
+            const int x = blockIdx.x * blockDim.x * PIXELS_PER_THREAD + threadIdx.x;
+            const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+            if (threadIdx.x == 0)
+                s_qsize[threadIdx.y] = 0;
+            __syncthreads();
+
+            if (y < src.rows)
+            {
+                // fill the queue
+                const uchar* srcRow = src.ptr(y);
+                for (int i = 0, xx = x; i < PIXELS_PER_THREAD && xx < src.cols; ++i, xx += blockDim.x)
+                {
+                    if (srcRow[xx])
+                    {
+                        const unsigned int val = (y << 16) | xx;
+                        const int qidx = Emulation::smem::atomicAdd(&s_qsize[threadIdx.y], 1);
+                        s_queues[threadIdx.y][qidx] = val;
+                    }
+                }
+            }
+
+            __syncthreads();
+
+            // let one thread reserve the space required in the global list
+            if (threadIdx.x == 0 && threadIdx.y == 0)
+            {
+                // find how many items are stored in each list
+                int totalSize = 0;
+                for (int i = 0; i < blockDim.y; ++i)
+                {
+                    s_globStart[i] = totalSize;
+                    totalSize += s_qsize[i];
+                }
+
+                // calculate the offset in the global list
+                const int globalOffset = atomicAdd(&g_counter, totalSize);
+                for (int i = 0; i < blockDim.y; ++i)
+                    s_globStart[i] += globalOffset;
+            }
+
+            __syncthreads();
+
+            // copy local queues to global queue
+            const int qsize = s_qsize[threadIdx.y];
+            int gidx = s_globStart[threadIdx.y] + threadIdx.x;
+            for(int i = threadIdx.x; i < qsize; i += blockDim.x, gidx += blockDim.x)
+                list[gidx] = s_queues[threadIdx.y][i];
+        }
+
+        int buildPointList_gpu(PtrStepSzb src, unsigned int* list)
+        {
+            const int PIXELS_PER_THREAD = 16;
+
+            void* counterPtr;
+            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
+
+            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
+
+            const dim3 block(32, 4);
+            const dim3 grid(divUp(src.cols, block.x * PIXELS_PER_THREAD), divUp(src.rows, block.y));
+
+            cudaSafeCall( cudaFuncSetCacheConfig(buildPointList<PIXELS_PER_THREAD>, cudaFuncCachePreferShared) );
+
+            buildPointList<PIXELS_PER_THREAD><<<grid, block>>>(src, list);
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+
+            int totalCount;
+            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
+
+            return totalCount;
+        }
+    }
+}}}
+
+#endif /* CUDA_DISABLER */

modules/gpu/src/cuda/canny.cu

@@ -239,30 +239,35 @@ namespace canny
 {
     __device__ int counter = 0;
 
-    __global__ void edgesHysteresisLocalKernel(PtrStepSzi map, ushort2* st)
+    __device__ __forceinline__ bool checkIdx(int y, int x, int rows, int cols)
+    {
+        return (y >= 0) && (y < rows) && (x >= 0) && (x < cols);
+    }
+
+    __global__ void edgesHysteresisLocalKernel(PtrStepSzi map, short2* st)
     {
         __shared__ volatile int smem[18][18];
 
         const int x = blockIdx.x * blockDim.x + threadIdx.x;
         const int y = blockIdx.y * blockDim.y + threadIdx.y;
 
-        smem[threadIdx.y + 1][threadIdx.x + 1] = x < map.cols && y < map.rows ? map(y, x) : 0;
+        smem[threadIdx.y + 1][threadIdx.x + 1] = checkIdx(y, x, map.rows, map.cols) ? map(y, x) : 0;
         if (threadIdx.y == 0)
-            smem[0][threadIdx.x + 1] = y > 0 ? map(y - 1, x) : 0;
+            smem[0][threadIdx.x + 1] = checkIdx(y - 1, x, map.rows, map.cols) ? map(y - 1, x) : 0;
         if (threadIdx.y == blockDim.y - 1)
-            smem[blockDim.y + 1][threadIdx.x + 1] = y + 1 < map.rows ? map(y + 1, x) : 0;
+            smem[blockDim.y + 1][threadIdx.x + 1] = checkIdx(y + 1, x, map.rows, map.cols) ? map(y + 1, x) : 0;
         if (threadIdx.x == 0)
-            smem[threadIdx.y + 1][0] = x > 0 ? map(y, x - 1) : 0;
+            smem[threadIdx.y + 1][0] = checkIdx(y, x - 1, map.rows, map.cols) ? map(y, x - 1) : 0;
         if (threadIdx.x == blockDim.x - 1)
-            smem[threadIdx.y + 1][blockDim.x + 1] = x + 1 < map.cols ? map(y, x + 1) : 0;
+            smem[threadIdx.y + 1][blockDim.x + 1] = checkIdx(y, x + 1, map.rows, map.cols) ? map(y, x + 1) : 0;
         if (threadIdx.x == 0 && threadIdx.y == 0)
-            smem[0][0] = y > 0 && x > 0 ? map(y - 1, x - 1) : 0;
+            smem[0][0] = checkIdx(y - 1, x - 1, map.rows, map.cols) ? map(y - 1, x - 1) : 0;
         if (threadIdx.x == blockDim.x - 1 && threadIdx.y == 0)
-            smem[0][blockDim.x + 1] = y > 0 && x + 1 < map.cols ? map(y - 1, x + 1) : 0;
+            smem[0][blockDim.x + 1] = checkIdx(y - 1, x + 1, map.rows, map.cols) ? map(y - 1, x + 1) : 0;
         if (threadIdx.x == 0 && threadIdx.y == blockDim.y - 1)
-            smem[blockDim.y + 1][0] = y + 1 < map.rows && x > 0 ? map(y + 1, x - 1) : 0;
+            smem[blockDim.y + 1][0] = checkIdx(y + 1, x - 1, map.rows, map.cols) ? map(y + 1, x - 1) : 0;
         if (threadIdx.x == blockDim.x - 1 && threadIdx.y == blockDim.y - 1)
-            smem[blockDim.y + 1][blockDim.x + 1] = y + 1 < map.rows && x + 1 < map.cols ? map(y + 1, x + 1) : 0;
+            smem[blockDim.y + 1][blockDim.x + 1] = checkIdx(y + 1, x + 1, map.rows, map.cols) ? map(y + 1, x + 1) : 0;
 
         __syncthreads();
 
@@ -317,11 +322,11 @@ namespace canny
         if (n > 0)
         {
             const int ind =  ::atomicAdd(&counter, 1);
-            st[ind] = make_ushort2(x, y);
+            st[ind] = make_short2(x, y);
         }
     }
 
-    void edgesHysteresisLocal(PtrStepSzi map, ushort2* st1)
+    void edgesHysteresisLocal(PtrStepSzi map, short2* st1)
     {
         void* counter_ptr;
         cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, counter) );
@@ -345,13 +350,13 @@ namespace canny
     __constant__ int c_dx[8] = {-1,  0,  1, -1, 1, -1, 0, 1};
     __constant__ int c_dy[8] = {-1, -1, -1,  0, 0,  1, 1, 1};
 
-    __global__ void edgesHysteresisGlobalKernel(PtrStepSzi map, ushort2* st1, ushort2* st2, const int count)
+    __global__ void edgesHysteresisGlobalKernel(PtrStepSzi map, short2* st1, short2* st2, const int count)
     {
         const int stack_size = 512;
 
         __shared__ int s_counter;
         __shared__ int s_ind;
-        __shared__ ushort2 s_st[stack_size];
+        __shared__ short2 s_st[stack_size];
 
         if (threadIdx.x == 0)
             s_counter = 0;
@@ -363,14 +368,14 @@ namespace canny
         if (ind >= count)
             return;
 
-        ushort2 pos = st1[ind];
+        short2 pos = st1[ind];
 
         if (threadIdx.x < 8)
         {
             pos.x += c_dx[threadIdx.x];
             pos.y += c_dy[threadIdx.x];
 
-            if (pos.x > 0 && pos.x < map.cols && pos.y > 0 && pos.y < map.rows && map(pos.y, pos.x) == 1)
+            if (pos.x > 0 && pos.x < map.cols - 1 && pos.y > 0 && pos.y < map.rows - 1 && map(pos.y, pos.x) == 1)
             {
                 map(pos.y, pos.x) = 2;
 
@@ -402,7 +407,7 @@ namespace canny
                 pos.x += c_dx[threadIdx.x & 7];
                 pos.y += c_dy[threadIdx.x & 7];
 
-                if (pos.x > 0 && pos.x < map.cols && pos.y > 0 && pos.y < map.rows && map(pos.y, pos.x) == 1)
+                if (pos.x > 0 && pos.x < map.cols - 1 && pos.y > 0 && pos.y < map.rows - 1 && map(pos.y, pos.x) == 1)
                 {
                     map(pos.y, pos.x) = 2;
 
@@ -419,8 +424,10 @@ namespace canny
         {
             if (threadIdx.x == 0)
             {
-                ind = ::atomicAdd(&counter, s_counter);
-                s_ind = ind - s_counter;
+                s_ind = ::atomicAdd(&counter, s_counter);
+
+                if (s_ind + s_counter > map.cols * map.rows)
+                    s_counter = 0;
             }
 
             __syncthreads();
@@ -432,7 +439,7 @@ namespace canny
         }
     }
 
-    void edgesHysteresisGlobal(PtrStepSzi map, ushort2* st1, ushort2* st2)
+    void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2)
     {
         void* counter_ptr;
         cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, canny::counter) );
@@ -454,6 +461,8 @@ namespace canny
 
             cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) );
 
+            count = min(count, map.cols * map.rows);
+
             std::swap(st1, st2);
         }
     }

modules/gpu/src/cuda/generalized_hough.cu

@@ -43,651 +43,18 @@
 #if !defined CUDA_DISABLER
 
 #include <thrust/device_ptr.h>
-#include <thrust/sort.h>
+#include <thrust/transform.h>
 
 #include "opencv2/gpu/device/common.hpp"
 #include "opencv2/gpu/device/emulation.hpp"
 #include "opencv2/gpu/device/vec_math.hpp"
 #include "opencv2/gpu/device/functional.hpp"
-#include "opencv2/gpu/device/limits.hpp"
-#include "opencv2/gpu/device/dynamic_smem.hpp"
 
 namespace cv { namespace gpu { namespace device
 {
     namespace hough
     {
-        __device__ int g_counter;
-
-        ////////////////////////////////////////////////////////////////////////
-        // buildPointList
-
-        template <int PIXELS_PER_THREAD>
-        __global__ void buildPointList(const PtrStepSzb src, unsigned int* list)
-        {
-            __shared__ unsigned int s_queues[4][32 * PIXELS_PER_THREAD];
-            __shared__ int s_qsize[4];
-            __shared__ int s_globStart[4];
-
-            const int x = blockIdx.x * blockDim.x * PIXELS_PER_THREAD + threadIdx.x;
-            const int y = blockIdx.y * blockDim.y + threadIdx.y;
-
-            if (threadIdx.x == 0)
-                s_qsize[threadIdx.y] = 0;
-            __syncthreads();
-
-            if (y < src.rows)
-            {
-                // fill the queue
-                const uchar* srcRow = src.ptr(y);
-                for (int i = 0, xx = x; i < PIXELS_PER_THREAD && xx < src.cols; ++i, xx += blockDim.x)
-                {
-                    if (srcRow[xx])
-                    {
-                        const unsigned int val = (y << 16) | xx;
-                        const int qidx = Emulation::smem::atomicAdd(&s_qsize[threadIdx.y], 1);
-                        s_queues[threadIdx.y][qidx] = val;
-                    }
-                }
-            }
-
-            __syncthreads();
-
-            // let one thread reserve the space required in the global list
-            if (threadIdx.x == 0 && threadIdx.y == 0)
-            {
-                // find how many items are stored in each list
-                int totalSize = 0;
-                for (int i = 0; i < blockDim.y; ++i)
-                {
-                    s_globStart[i] = totalSize;
-                    totalSize += s_qsize[i];
-                }
-
-                // calculate the offset in the global list
-                const int globalOffset = atomicAdd(&g_counter, totalSize);
-                for (int i = 0; i < blockDim.y; ++i)
-                    s_globStart[i] += globalOffset;
-            }
-
-            __syncthreads();
-
-            // copy local queues to global queue
-            const int qsize = s_qsize[threadIdx.y];
-            int gidx = s_globStart[threadIdx.y] + threadIdx.x;
-            for(int i = threadIdx.x; i < qsize; i += blockDim.x, gidx += blockDim.x)
-                list[gidx] = s_queues[threadIdx.y][i];
-        }
-
-        int buildPointList_gpu(PtrStepSzb src, unsigned int* list)
-        {
-            const int PIXELS_PER_THREAD = 16;
-
-            void* counterPtr;
-            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
-
-            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
-
-            const dim3 block(32, 4);
-            const dim3 grid(divUp(src.cols, block.x * PIXELS_PER_THREAD), divUp(src.rows, block.y));
-
-            cudaSafeCall( cudaFuncSetCacheConfig(buildPointList<PIXELS_PER_THREAD>, cudaFuncCachePreferShared) );
-
-            buildPointList<PIXELS_PER_THREAD><<<grid, block>>>(src, list);
-            cudaSafeCall( cudaGetLastError() );
-
-            cudaSafeCall( cudaDeviceSynchronize() );
-
-            int totalCount;
-            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
-
-            return totalCount;
-        }
-
-        ////////////////////////////////////////////////////////////////////////
-        // linesAccum
-
-        __global__ void linesAccumGlobal(const unsigned int* list, const int count, PtrStepi accum, const float irho, const float theta, const int numrho)
-        {
-            const int n = blockIdx.x;
-            const float ang = n * theta;
-
-            float sinVal;
-            float cosVal;
-            sincosf(ang, &sinVal, &cosVal);
-            sinVal *= irho;
-            cosVal *= irho;
-
-            const int shift = (numrho - 1) / 2;
-
-            int* accumRow = accum.ptr(n + 1);
-            for (int i = threadIdx.x; i < count; i += blockDim.x)
-            {
-                const unsigned int val = list[i];
-
-                const int x = (val & 0xFFFF);
-                const int y = (val >> 16) & 0xFFFF;
-
-                int r = __float2int_rn(x * cosVal + y * sinVal);
-                r += shift;
-
-                ::atomicAdd(accumRow + r + 1, 1);
-            }
-        }
-
-        __global__ void linesAccumShared(const unsigned int* list, const int count, PtrStepi accum, const float irho, const float theta, const int numrho)
-        {
-            int* smem = DynamicSharedMem<int>();
-
-            for (int i = threadIdx.x; i < numrho + 1; i += blockDim.x)
-                smem[i] = 0;
-
-            __syncthreads();
-
-            const int n = blockIdx.x;
-            const float ang = n * theta;
-
-            float sinVal;
-            float cosVal;
-            sincosf(ang, &sinVal, &cosVal);
-            sinVal *= irho;
-            cosVal *= irho;
-
-            const int shift = (numrho - 1) / 2;
-
-            for (int i = threadIdx.x; i < count; i += blockDim.x)
-            {
-                const unsigned int val = list[i];
-
-                const int x = (val & 0xFFFF);
-                const int y = (val >> 16) & 0xFFFF;
-
-                int r = __float2int_rn(x * cosVal + y * sinVal);
-                r += shift;
-
-                Emulation::smem::atomicAdd(&smem[r + 1], 1);
-            }
-
-            __syncthreads();
-
-            int* accumRow = accum.ptr(n + 1);
-            for (int i = threadIdx.x; i < numrho + 1; i += blockDim.x)
-                accumRow[i] = smem[i];
-        }
-
-        void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20)
-        {
-            const dim3 block(has20 ? 1024 : 512);
-            const dim3 grid(accum.rows - 2);
-
-            size_t smemSize = (accum.cols - 1) * sizeof(int);
-
-            if (smemSize < sharedMemPerBlock - 1000)
-                linesAccumShared<<<grid, block, smemSize>>>(list, count, accum, 1.0f / rho, theta, accum.cols - 2);
-            else
-                linesAccumGlobal<<<grid, block>>>(list, count, accum, 1.0f / rho, theta, accum.cols - 2);
-
-            cudaSafeCall( cudaGetLastError() );
-
-            cudaSafeCall( cudaDeviceSynchronize() );
-        }
-
-        ////////////////////////////////////////////////////////////////////////
-        // linesGetResult
-
-        __global__ void linesGetResult(const PtrStepSzi accum, float2* out, int* votes, const int maxSize, const float rho, const float theta, const int threshold, const int numrho)
-        {
-            const int r = blockIdx.x * blockDim.x + threadIdx.x;
-            const int n = blockIdx.y * blockDim.y + threadIdx.y;
-
-            if (r >= accum.cols - 2 || n >= accum.rows - 2)
-                return;
-
-            const int curVotes = accum(n + 1, r + 1);
-
-            if (curVotes > threshold &&
-                curVotes >  accum(n + 1, r) &&
-                curVotes >= accum(n + 1, r + 2) &&
-                curVotes >  accum(n, r + 1) &&
-                curVotes >= accum(n + 2, r + 1))
-            {
-                const float radius = (r - (numrho - 1) * 0.5f) * rho;
-                const float angle = n * theta;
-
-                const int ind = ::atomicAdd(&g_counter, 1);
-                if (ind < maxSize)
-                {
-                    out[ind] = make_float2(radius, angle);
-                    votes[ind] = curVotes;
-                }
-            }
-        }
-
-        int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort)
-        {
-            void* counterPtr;
-            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
-
-            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
-
-            const dim3 block(32, 8);
-            const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y));
-
-            cudaSafeCall( cudaFuncSetCacheConfig(linesGetResult, cudaFuncCachePreferL1) );
-
-            linesGetResult<<<grid, block>>>(accum, out, votes, maxSize, rho, theta, threshold, accum.cols - 2);
-            cudaSafeCall( cudaGetLastError() );
-
-            cudaSafeCall( cudaDeviceSynchronize() );
-
-            int totalCount;
-            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
-
-            totalCount = ::min(totalCount, maxSize);
-
-            if (doSort && totalCount > 0)
-            {
-                thrust::device_ptr<float2> outPtr(out);
-                thrust::device_ptr<int> votesPtr(votes);
-                thrust::sort_by_key(votesPtr, votesPtr + totalCount, outPtr, thrust::greater<int>());
-            }
-
-            return totalCount;
-        }
-
-        ////////////////////////////////////////////////////////////////////////
-        // houghLinesProbabilistic
-
-        texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex_mask(false, cudaFilterModePoint, cudaAddressModeClamp);
-
-        __global__ void houghLinesProbabilistic(const PtrStepSzi accum,
-                                                int4* out, const int maxSize,
-                                                const float rho, const float theta,
-                                                const int lineGap, const int lineLength,
-                                                const int rows, const int cols)
-        {
-            const int r = blockIdx.x * blockDim.x + threadIdx.x;
-            const int n = blockIdx.y * blockDim.y + threadIdx.y;
-
-            if (r >= accum.cols - 2 || n >= accum.rows - 2)
-                return;
-
-            const int curVotes = accum(n + 1, r + 1);
-
-            if (curVotes >= lineLength &&
-                curVotes > accum(n, r) &&
-                curVotes > accum(n, r + 1) &&
-                curVotes > accum(n, r + 2) &&
-                curVotes > accum(n + 1, r) &&
-                curVotes > accum(n + 1, r + 2) &&
-                curVotes > accum(n + 2, r) &&
-                curVotes > accum(n + 2, r + 1) &&
-                curVotes > accum(n + 2, r + 2))
-            {
-                const float radius = (r - (accum.cols - 2 - 1) * 0.5f) * rho;
-                const float angle = n * theta;
-
-                float cosa;
-                float sina;
-                sincosf(angle, &sina, &cosa);
-
-                float2 p0 = make_float2(cosa * radius, sina * radius);
-                float2 dir = make_float2(-sina, cosa);
-
-                float2 pb[4] = {make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1)};
-                float a;
-
-                if (dir.x != 0)
-                {
-                    a = -p0.x / dir.x;
-                    pb[0].x = 0;
-                    pb[0].y = p0.y + a * dir.y;
-
-                    a = (cols - 1 - p0.x) / dir.x;
-                    pb[1].x = cols - 1;
-                    pb[1].y = p0.y + a * dir.y;
-                }
-                if (dir.y != 0)
-                {
-                    a = -p0.y / dir.y;
-                    pb[2].x = p0.x + a * dir.x;
-                    pb[2].y = 0;
-
-                    a = (rows - 1 - p0.y) / dir.y;
-                    pb[3].x = p0.x + a * dir.x;
-                    pb[3].y = rows - 1;
-                }
-
-                if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < rows))
-                {
-                    p0 = pb[0];
-                    if (dir.x < 0)
-                        dir = -dir;
-                }
-                else if (pb[1].x == cols - 1 && (pb[0].y >= 0 && pb[0].y < rows))
-                {
-                    p0 = pb[1];
-                    if (dir.x > 0)
-                        dir = -dir;
-                }
-                else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < cols))
-                {
-                    p0 = pb[2];
-                    if (dir.y < 0)
-                        dir = -dir;
-                }
-                else if (pb[3].y == rows - 1 && (pb[3].x >= 0 && pb[3].x < cols))
-                {
-                    p0 = pb[3];
-                    if (dir.y > 0)
-                        dir = -dir;
-                }
-
-                float2 d;
-                if (::fabsf(dir.x) > ::fabsf(dir.y))
-                {
-                    d.x = dir.x > 0 ? 1 : -1;
-                    d.y = dir.y / ::fabsf(dir.x);
-                }
-                else
-                {
-                    d.x = dir.x / ::fabsf(dir.y);
-                    d.y = dir.y > 0 ? 1 : -1;
-                }
-
-                float2 line_end[2];
-                int gap;
-                bool inLine = false;
-
-                float2 p1 = p0;
-                if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows)
-                    return;
-
-                for (;;)
-                {
-                    if (tex2D(tex_mask, p1.x, p1.y))
-                    {
-                        gap = 0;
-
-                        if (!inLine)
-                        {
-                            line_end[0] = p1;
-                            line_end[1] = p1;
-                            inLine = true;
-                        }
-                        else
-                        {
-                            line_end[1] = p1;
-                        }
-                    }
-                    else if (inLine)
-                    {
-                        if (++gap > lineGap)
-                        {
-                            bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength ||
-                                             ::abs(line_end[1].y - line_end[0].y) >= lineLength;
-
-                            if (good_line)
-                            {
-                                const int ind = ::atomicAdd(&g_counter, 1);
-                                if (ind < maxSize)
-                                    out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
-                            }
-
-                            gap = 0;
-                            inLine = false;
-                        }
-                    }
-
-                    p1 = p1 + d;
-                    if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows)
-                    {
-                        if (inLine)
-                        {
-                            bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength ||
-                                             ::abs(line_end[1].y - line_end[0].y) >= lineLength;
-
-                            if (good_line)
-                            {
-                                const int ind = ::atomicAdd(&g_counter, 1);
-                                if (ind < maxSize)
-                                    out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
-                            }
-
-                        }
-                        break;
-                    }
-                }
-            }
-        }
-
-        int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength)
-        {
-            void* counterPtr;
-            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
-
-            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
-
-            const dim3 block(32, 8);
-            const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y));
-
-            bindTexture(&tex_mask, mask);
-
-            houghLinesProbabilistic<<<grid, block>>>(accum,
-                                                     out, maxSize,
-                                                     rho, theta,
-                                                     lineGap, lineLength,
-                                                     mask.rows, mask.cols);
-            cudaSafeCall( cudaGetLastError() );
-
-            cudaSafeCall( cudaDeviceSynchronize() );
-
-            int totalCount;
-            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
-
-            totalCount = ::min(totalCount, maxSize);
-
-            return totalCount;
-        }
-
-        ////////////////////////////////////////////////////////////////////////
-        // circlesAccumCenters
-
-        __global__ void circlesAccumCenters(const unsigned int* list, const int count, const PtrStepi dx, const PtrStepi dy,
-                                            PtrStepi accum, const int width, const int height, const int minRadius, const int maxRadius, const float idp)
-        {
-            const int SHIFT = 10;
-            const int ONE = 1 << SHIFT;
-
-            const int tid = blockIdx.x * blockDim.x + threadIdx.x;
-
-            if (tid >= count)
-                return;
-
-            const unsigned int val = list[tid];
-
-            const int x = (val & 0xFFFF);
-            const int y = (val >> 16) & 0xFFFF;
-
-            const int vx = dx(y, x);
-            const int vy = dy(y, x);
-
-            if (vx == 0 && vy == 0)
-                return;
-
-            const float mag = ::sqrtf(vx * vx + vy * vy);
-
-            const int x0 = __float2int_rn((x * idp) * ONE);
-            const int y0 = __float2int_rn((y * idp) * ONE);
-
-            int sx = __float2int_rn((vx * idp) * ONE / mag);
-            int sy = __float2int_rn((vy * idp) * ONE / mag);
-
-            // Step from minRadius to maxRadius in both directions of the gradient
-            for (int k1 = 0; k1 < 2; ++k1)
-            {
-                int x1 = x0 + minRadius * sx;
-                int y1 = y0 + minRadius * sy;
-
-                for (int r = minRadius; r <= maxRadius; x1 += sx, y1 += sy, ++r)
-                {
-                    const int x2 = x1 >> SHIFT;
-                    const int y2 = y1 >> SHIFT;
-
-                    if (x2 < 0 || x2 >= width || y2 < 0 || y2 >= height)
-                        break;
-
-                    ::atomicAdd(accum.ptr(y2 + 1) + x2 + 1, 1);
-                }
-
-                sx = -sx;
-                sy = -sy;
-            }
-        }
-
-        void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp)
-        {
-            const dim3 block(256);
-            const dim3 grid(divUp(count, block.x));
-
-            cudaSafeCall( cudaFuncSetCacheConfig(circlesAccumCenters, cudaFuncCachePreferL1) );
-
-            circlesAccumCenters<<<grid, block>>>(list, count, dx, dy, accum, accum.cols - 2, accum.rows - 2, minRadius, maxRadius, idp);
-            cudaSafeCall( cudaGetLastError() );
-
-            cudaSafeCall( cudaDeviceSynchronize() );
-        }
-
-        ////////////////////////////////////////////////////////////////////////
-        // buildCentersList
-
-        __global__ void buildCentersList(const PtrStepSzi accum, unsigned int* centers, const int threshold)
-        {
-            const int x = blockIdx.x * blockDim.x + threadIdx.x;
-            const int y = blockIdx.y * blockDim.y + threadIdx.y;
-
-            if (x < accum.cols - 2 && y < accum.rows - 2)
-            {
-                const int top = accum(y, x + 1);
-
-                const int left = accum(y + 1, x);
-                const int cur = accum(y + 1, x + 1);
-                const int right = accum(y + 1, x + 2);
-
-                const int bottom = accum(y + 2, x + 1);
-
-                if (cur > threshold && cur > top && cur >= bottom && cur >  left && cur >= right)
-                {
-                    const unsigned int val = (y << 16) | x;
-                    const int idx = ::atomicAdd(&g_counter, 1);
-                    centers[idx] = val;
-                }
-            }
-        }
-
-        int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold)
-        {
-            void* counterPtr;
-            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
-
-            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
-
-            const dim3 block(32, 8);
-            const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y));
-
-            cudaSafeCall( cudaFuncSetCacheConfig(buildCentersList, cudaFuncCachePreferL1) );
-
-            buildCentersList<<<grid, block>>>(accum, centers, threshold);
-            cudaSafeCall( cudaGetLastError() );
-
-            cudaSafeCall( cudaDeviceSynchronize() );
-
-            int totalCount;
-            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
-
-            return totalCount;
-        }
-
-        ////////////////////////////////////////////////////////////////////////
-        // circlesAccumRadius
-
-        __global__ void circlesAccumRadius(const unsigned int* centers, const unsigned int* list, const int count,
-                                           float3* circles, const int maxCircles, const float dp,
-                                           const int minRadius, const int maxRadius, const int histSize, const int threshold)
-        {
-            int* smem = DynamicSharedMem<int>();
-
-            for (int i = threadIdx.x; i < histSize + 2; i += blockDim.x)
-                smem[i] = 0;
-            __syncthreads();
-
-            unsigned int val = centers[blockIdx.x];
-
-            float cx = (val & 0xFFFF);
-            float cy = (val >> 16) & 0xFFFF;
-
-            cx = (cx + 0.5f) * dp;
-            cy = (cy + 0.5f) * dp;
-
-            for (int i = threadIdx.x; i < count; i += blockDim.x)
-            {
-                val = list[i];
-
-                const int x = (val & 0xFFFF);
-                const int y = (val >> 16) & 0xFFFF;
-
-                const float rad = ::sqrtf((cx - x) * (cx - x) + (cy - y) * (cy - y));
-                if (rad >= minRadius && rad <= maxRadius)
-                {
-                    const int r = __float2int_rn(rad - minRadius);
-
-                    Emulation::smem::atomicAdd(&smem[r + 1], 1);
-                }
-            }
-
-            __syncthreads();
-
-            for (int i = threadIdx.x; i < histSize; i += blockDim.x)
-            {
-                const int curVotes = smem[i + 1];
-
-                if (curVotes >= threshold && curVotes > smem[i] && curVotes >= smem[i + 2])
-                {
-                    const int ind = ::atomicAdd(&g_counter, 1);
-                    if (ind < maxCircles)
-                        circles[ind] = make_float3(cx, cy, i + minRadius);
-                }
-            }
-        }
-
-        int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count,
-                                   float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20)
-        {
-            void* counterPtr;
-            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
-
-            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
-
-            const dim3 block(has20 ? 1024 : 512);
-            const dim3 grid(centersCount);
-
-            const int histSize = maxRadius - minRadius + 1;
-            size_t smemSize = (histSize + 2) * sizeof(int);
-
-            circlesAccumRadius<<<grid, block, smemSize>>>(centers, list, count, circles, maxCircles, dp, minRadius, maxRadius, histSize, threshold);
-            cudaSafeCall( cudaGetLastError() );
-
-            cudaSafeCall( cudaDeviceSynchronize() );
-
-            int totalCount;
-            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
-
-            totalCount = ::min(totalCount, maxCircles);
-
-            return totalCount;
-        }
-
-        ////////////////////////////////////////////////////////////////////////
-        // Generalized Hough
+        __device__ static int g_counter;
 
         template <typename T, int PIXELS_PER_THREAD>
         __global__ void buildEdgePointList(const PtrStepSzb edges, const PtrStep<T> dx, const PtrStep<T> dy, unsigned int* coordList, float* thetaList)
@@ -1706,5 +1073,4 @@ namespace cv { namespace gpu { namespace device
     }
 }}}
 
-
 #endif /* CUDA_DISABLER */

modules/gpu/src/cuda/hough_circles.cu

@@ -0,0 +1,254 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#if !defined CUDA_DISABLER
+
+#include "opencv2/gpu/device/common.hpp"
+#include "opencv2/gpu/device/emulation.hpp"
+#include "opencv2/gpu/device/dynamic_smem.hpp"
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        __device__ static int g_counter;
+
+        ////////////////////////////////////////////////////////////////////////
+        // circlesAccumCenters
+
+        __global__ void circlesAccumCenters(const unsigned int* list, const int count, const PtrStepi dx, const PtrStepi dy,
+                                            PtrStepi accum, const int width, const int height, const int minRadius, const int maxRadius, const float idp)
+        {
+            const int SHIFT = 10;
+            const int ONE = 1 << SHIFT;
+
+            const int tid = blockIdx.x * blockDim.x + threadIdx.x;
+
+            if (tid >= count)
+                return;
+
+            const unsigned int val = list[tid];
+
+            const int x = (val & 0xFFFF);
+            const int y = (val >> 16) & 0xFFFF;
+
+            const int vx = dx(y, x);
+            const int vy = dy(y, x);
+
+            if (vx == 0 && vy == 0)
+                return;
+
+            const float mag = ::sqrtf(vx * vx + vy * vy);
+
+            const int x0 = __float2int_rn((x * idp) * ONE);
+            const int y0 = __float2int_rn((y * idp) * ONE);
+
+            int sx = __float2int_rn((vx * idp) * ONE / mag);
+            int sy = __float2int_rn((vy * idp) * ONE / mag);
+
+            // Step from minRadius to maxRadius in both directions of the gradient
+            for (int k1 = 0; k1 < 2; ++k1)
+            {
+                int x1 = x0 + minRadius * sx;
+                int y1 = y0 + minRadius * sy;
+
+                for (int r = minRadius; r <= maxRadius; x1 += sx, y1 += sy, ++r)
+                {
+                    const int x2 = x1 >> SHIFT;
+                    const int y2 = y1 >> SHIFT;
+
+                    if (x2 < 0 || x2 >= width || y2 < 0 || y2 >= height)
+                        break;
+
+                    ::atomicAdd(accum.ptr(y2 + 1) + x2 + 1, 1);
+                }
+
+                sx = -sx;
+                sy = -sy;
+            }
+        }
+
+        void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp)
+        {
+            const dim3 block(256);
+            const dim3 grid(divUp(count, block.x));
+
+            cudaSafeCall( cudaFuncSetCacheConfig(circlesAccumCenters, cudaFuncCachePreferL1) );
+
+            circlesAccumCenters<<<grid, block>>>(list, count, dx, dy, accum, accum.cols - 2, accum.rows - 2, minRadius, maxRadius, idp);
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+        }
+
+        ////////////////////////////////////////////////////////////////////////
+        // buildCentersList
+
+        __global__ void buildCentersList(const PtrStepSzi accum, unsigned int* centers, const int threshold)
+        {
+            const int x = blockIdx.x * blockDim.x + threadIdx.x;
+            const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+            if (x < accum.cols - 2 && y < accum.rows - 2)
+            {
+                const int top = accum(y, x + 1);
+
+                const int left = accum(y + 1, x);
+                const int cur = accum(y + 1, x + 1);
+                const int right = accum(y + 1, x + 2);
+
+                const int bottom = accum(y + 2, x + 1);
+
+                if (cur > threshold && cur > top && cur >= bottom && cur >  left && cur >= right)
+                {
+                    const unsigned int val = (y << 16) | x;
+                    const int idx = ::atomicAdd(&g_counter, 1);
+                    centers[idx] = val;
+                }
+            }
+        }
+
+        int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold)
+        {
+            void* counterPtr;
+            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
+
+            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
+
+            const dim3 block(32, 8);
+            const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y));
+
+            cudaSafeCall( cudaFuncSetCacheConfig(buildCentersList, cudaFuncCachePreferL1) );
+
+            buildCentersList<<<grid, block>>>(accum, centers, threshold);
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+
+            int totalCount;
+            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
+
+            return totalCount;
+        }
+
+        ////////////////////////////////////////////////////////////////////////
+        // circlesAccumRadius
+
+        __global__ void circlesAccumRadius(const unsigned int* centers, const unsigned int* list, const int count,
+                                           float3* circles, const int maxCircles, const float dp,
+                                           const int minRadius, const int maxRadius, const int histSize, const int threshold)
+        {
+            int* smem = DynamicSharedMem<int>();
+
+            for (int i = threadIdx.x; i < histSize + 2; i += blockDim.x)
+                smem[i] = 0;
+            __syncthreads();
+
+            unsigned int val = centers[blockIdx.x];
+
+            float cx = (val & 0xFFFF);
+            float cy = (val >> 16) & 0xFFFF;
+
+            cx = (cx + 0.5f) * dp;
+            cy = (cy + 0.5f) * dp;
+
+            for (int i = threadIdx.x; i < count; i += blockDim.x)
+            {
+                val = list[i];
+
+                const int x = (val & 0xFFFF);
+                const int y = (val >> 16) & 0xFFFF;
+
+                const float rad = ::sqrtf((cx - x) * (cx - x) + (cy - y) * (cy - y));
+                if (rad >= minRadius && rad <= maxRadius)
+                {
+                    const int r = __float2int_rn(rad - minRadius);
+
+                    Emulation::smem::atomicAdd(&smem[r + 1], 1);
+                }
+            }
+
+            __syncthreads();
+
+            for (int i = threadIdx.x; i < histSize; i += blockDim.x)
+            {
+                const int curVotes = smem[i + 1];
+
+                if (curVotes >= threshold && curVotes > smem[i] && curVotes >= smem[i + 2])
+                {
+                    const int ind = ::atomicAdd(&g_counter, 1);
+                    if (ind < maxCircles)
+                        circles[ind] = make_float3(cx, cy, i + minRadius);
+                }
+            }
+        }
+
+        int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count,
+                                   float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20)
+        {
+            void* counterPtr;
+            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
+
+            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
+
+            const dim3 block(has20 ? 1024 : 512);
+            const dim3 grid(centersCount);
+
+            const int histSize = maxRadius - minRadius + 1;
+            size_t smemSize = (histSize + 2) * sizeof(int);
+
+            circlesAccumRadius<<<grid, block, smemSize>>>(centers, list, count, circles, maxCircles, dp, minRadius, maxRadius, histSize, threshold);
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+
+            int totalCount;
+            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
+
+            totalCount = ::min(totalCount, maxCircles);
+
+            return totalCount;
+        }
+    }
+}}}
+
+#endif /* CUDA_DISABLER */

modules/gpu/src/cuda/hough_lines.cu

@@ -0,0 +1,212 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#if !defined CUDA_DISABLER
+
+#include <thrust/device_ptr.h>
+#include <thrust/sort.h>
+
+#include "opencv2/gpu/device/common.hpp"
+#include "opencv2/gpu/device/emulation.hpp"
+#include "opencv2/gpu/device/dynamic_smem.hpp"
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        __device__ static int g_counter;
+
+        ////////////////////////////////////////////////////////////////////////
+        // linesAccum
+
+        __global__ void linesAccumGlobal(const unsigned int* list, const int count, PtrStepi accum, const float irho, const float theta, const int numrho)
+        {
+            const int n = blockIdx.x;
+            const float ang = n * theta;
+
+            float sinVal;
+            float cosVal;
+            sincosf(ang, &sinVal, &cosVal);
+            sinVal *= irho;
+            cosVal *= irho;
+
+            const int shift = (numrho - 1) / 2;
+
+            int* accumRow = accum.ptr(n + 1);
+            for (int i = threadIdx.x; i < count; i += blockDim.x)
+            {
+                const unsigned int val = list[i];
+
+                const int x = (val & 0xFFFF);
+                const int y = (val >> 16) & 0xFFFF;
+
+                int r = __float2int_rn(x * cosVal + y * sinVal);
+                r += shift;
+
+                ::atomicAdd(accumRow + r + 1, 1);
+            }
+        }
+
+        __global__ void linesAccumShared(const unsigned int* list, const int count, PtrStepi accum, const float irho, const float theta, const int numrho)
+        {
+            int* smem = DynamicSharedMem<int>();
+
+            for (int i = threadIdx.x; i < numrho + 1; i += blockDim.x)
+                smem[i] = 0;
+
+            __syncthreads();
+
+            const int n = blockIdx.x;
+            const float ang = n * theta;
+
+            float sinVal;
+            float cosVal;
+            sincosf(ang, &sinVal, &cosVal);
+            sinVal *= irho;
+            cosVal *= irho;
+
+            const int shift = (numrho - 1) / 2;
+
+            for (int i = threadIdx.x; i < count; i += blockDim.x)
+            {
+                const unsigned int val = list[i];
+
+                const int x = (val & 0xFFFF);
+                const int y = (val >> 16) & 0xFFFF;
+
+                int r = __float2int_rn(x * cosVal + y * sinVal);
+                r += shift;
+
+                Emulation::smem::atomicAdd(&smem[r + 1], 1);
+            }
+
+            __syncthreads();
+
+            int* accumRow = accum.ptr(n + 1);
+            for (int i = threadIdx.x; i < numrho + 1; i += blockDim.x)
+                accumRow[i] = smem[i];
+        }
+
+        void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20)
+        {
+            const dim3 block(has20 ? 1024 : 512);
+            const dim3 grid(accum.rows - 2);
+
+            size_t smemSize = (accum.cols - 1) * sizeof(int);
+
+            if (smemSize < sharedMemPerBlock - 1000)
+                linesAccumShared<<<grid, block, smemSize>>>(list, count, accum, 1.0f / rho, theta, accum.cols - 2);
+            else
+                linesAccumGlobal<<<grid, block>>>(list, count, accum, 1.0f / rho, theta, accum.cols - 2);
+
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+        }
+
+        ////////////////////////////////////////////////////////////////////////
+        // linesGetResult
+
+        __global__ void linesGetResult(const PtrStepSzi accum, float2* out, int* votes, const int maxSize, const float rho, const float theta, const int threshold, const int numrho)
+        {
+            const int r = blockIdx.x * blockDim.x + threadIdx.x;
+            const int n = blockIdx.y * blockDim.y + threadIdx.y;
+
+            if (r >= accum.cols - 2 || n >= accum.rows - 2)
+                return;
+
+            const int curVotes = accum(n + 1, r + 1);
+
+            if (curVotes > threshold &&
+                curVotes >  accum(n + 1, r) &&
+                curVotes >= accum(n + 1, r + 2) &&
+                curVotes >  accum(n, r + 1) &&
+                curVotes >= accum(n + 2, r + 1))
+            {
+                const float radius = (r - (numrho - 1) * 0.5f) * rho;
+                const float angle = n * theta;
+
+                const int ind = ::atomicAdd(&g_counter, 1);
+                if (ind < maxSize)
+                {
+                    out[ind] = make_float2(radius, angle);
+                    votes[ind] = curVotes;
+                }
+            }
+        }
+
+        int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort)
+        {
+            void* counterPtr;
+            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
+
+            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
+
+            const dim3 block(32, 8);
+            const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y));
+
+            cudaSafeCall( cudaFuncSetCacheConfig(linesGetResult, cudaFuncCachePreferL1) );
+
+            linesGetResult<<<grid, block>>>(accum, out, votes, maxSize, rho, theta, threshold, accum.cols - 2);
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+
+            int totalCount;
+            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
+
+            totalCount = ::min(totalCount, maxSize);
+
+            if (doSort && totalCount > 0)
+            {
+                thrust::device_ptr<float2> outPtr(out);
+                thrust::device_ptr<int> votesPtr(votes);
+                thrust::sort_by_key(votesPtr, votesPtr + totalCount, outPtr, thrust::greater<int>());
+            }
+
+            return totalCount;
+        }
+    }
+}}}
+
+
+#endif /* CUDA_DISABLER */

modules/gpu/src/cuda/hough_segments.cu

@@ -0,0 +1,249 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#if !defined CUDA_DISABLER
+
+#include "opencv2/gpu/device/common.hpp"
+#include "opencv2/gpu/device/vec_math.hpp"
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        __device__ int g_counter;
+
+        texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex_mask(false, cudaFilterModePoint, cudaAddressModeClamp);
+
+        __global__ void houghLinesProbabilistic(const PtrStepSzi accum,
+                                                int4* out, const int maxSize,
+                                                const float rho, const float theta,
+                                                const int lineGap, const int lineLength,
+                                                const int rows, const int cols)
+        {
+            const int r = blockIdx.x * blockDim.x + threadIdx.x;
+            const int n = blockIdx.y * blockDim.y + threadIdx.y;
+
+            if (r >= accum.cols - 2 || n >= accum.rows - 2)
+                return;
+
+            const int curVotes = accum(n + 1, r + 1);
+
+            if (curVotes >= lineLength &&
+                curVotes > accum(n, r) &&
+                curVotes > accum(n, r + 1) &&
+                curVotes > accum(n, r + 2) &&
+                curVotes > accum(n + 1, r) &&
+                curVotes > accum(n + 1, r + 2) &&
+                curVotes > accum(n + 2, r) &&
+                curVotes > accum(n + 2, r + 1) &&
+                curVotes > accum(n + 2, r + 2))
+            {
+                const float radius = (r - (accum.cols - 2 - 1) * 0.5f) * rho;
+                const float angle = n * theta;
+
+                float cosa;
+                float sina;
+                sincosf(angle, &sina, &cosa);
+
+                float2 p0 = make_float2(cosa * radius, sina * radius);
+                float2 dir = make_float2(-sina, cosa);
+
+                float2 pb[4] = {make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1)};
+                float a;
+
+                if (dir.x != 0)
+                {
+                    a = -p0.x / dir.x;
+                    pb[0].x = 0;
+                    pb[0].y = p0.y + a * dir.y;
+
+                    a = (cols - 1 - p0.x) / dir.x;
+                    pb[1].x = cols - 1;
+                    pb[1].y = p0.y + a * dir.y;
+                }
+                if (dir.y != 0)
+                {
+                    a = -p0.y / dir.y;
+                    pb[2].x = p0.x + a * dir.x;
+                    pb[2].y = 0;
+
+                    a = (rows - 1 - p0.y) / dir.y;
+                    pb[3].x = p0.x + a * dir.x;
+                    pb[3].y = rows - 1;
+                }
+
+                if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < rows))
+                {
+                    p0 = pb[0];
+                    if (dir.x < 0)
+                        dir = -dir;
+                }
+                else if (pb[1].x == cols - 1 && (pb[0].y >= 0 && pb[0].y < rows))
+                {
+                    p0 = pb[1];
+                    if (dir.x > 0)
+                        dir = -dir;
+                }
+                else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < cols))
+                {
+                    p0 = pb[2];
+                    if (dir.y < 0)
+                        dir = -dir;
+                }
+                else if (pb[3].y == rows - 1 && (pb[3].x >= 0 && pb[3].x < cols))
+                {
+                    p0 = pb[3];
+                    if (dir.y > 0)
+                        dir = -dir;
+                }
+
+                float2 d;
+                if (::fabsf(dir.x) > ::fabsf(dir.y))
+                {
+                    d.x = dir.x > 0 ? 1 : -1;
+                    d.y = dir.y / ::fabsf(dir.x);
+                }
+                else
+                {
+                    d.x = dir.x / ::fabsf(dir.y);
+                    d.y = dir.y > 0 ? 1 : -1;
+                }
+
+                float2 line_end[2];
+                int gap;
+                bool inLine = false;
+
+                float2 p1 = p0;
+                if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows)
+                    return;
+
+                for (;;)
+                {
+                    if (tex2D(tex_mask, p1.x, p1.y))
+                    {
+                        gap = 0;
+
+                        if (!inLine)
+                        {
+                            line_end[0] = p1;
+                            line_end[1] = p1;
+                            inLine = true;
+                        }
+                        else
+                        {
+                            line_end[1] = p1;
+                        }
+                    }
+                    else if (inLine)
+                    {
+                        if (++gap > lineGap)
+                        {
+                            bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength ||
+                                             ::abs(line_end[1].y - line_end[0].y) >= lineLength;
+
+                            if (good_line)
+                            {
+                                const int ind = ::atomicAdd(&g_counter, 1);
+                                if (ind < maxSize)
+                                    out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
+                            }
+
+                            gap = 0;
+                            inLine = false;
+                        }
+                    }
+
+                    p1 = p1 + d;
+                    if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows)
+                    {
+                        if (inLine)
+                        {
+                            bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength ||
+                                             ::abs(line_end[1].y - line_end[0].y) >= lineLength;
+
+                            if (good_line)
+                            {
+                                const int ind = ::atomicAdd(&g_counter, 1);
+                                if (ind < maxSize)
+                                    out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
+                            }
+
+                        }
+                        break;
+                    }
+                }
+            }
+        }
+
+        int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength)
+        {
+            void* counterPtr;
+            cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
+
+            cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
+
+            const dim3 block(32, 8);
+            const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y));
+
+            bindTexture(&tex_mask, mask);
+
+            houghLinesProbabilistic<<<grid, block>>>(accum,
+                                                     out, maxSize,
+                                                     rho, theta,
+                                                     lineGap, lineLength,
+                                                     mask.rows, mask.cols);
+            cudaSafeCall( cudaGetLastError() );
+
+            cudaSafeCall( cudaDeviceSynchronize() );
+
+            int totalCount;
+            cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
+
+            totalCount = ::min(totalCount, maxSize);
+
+            return totalCount;
+        }
+    }
+}}}
+
+
+#endif /* CUDA_DISABLER */

modules/gpu/src/generalized_hough.cpp

@@ -48,16 +48,6 @@ using namespace cv::gpu;
 
 #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
 
-void cv::gpu::HoughLines(const GpuMat&, GpuMat&, float, float, int, bool, int) { throw_nogpu(); }
-void cv::gpu::HoughLines(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, bool, int) { throw_nogpu(); }
-void cv::gpu::HoughLinesDownload(const GpuMat&, OutputArray, OutputArray) { throw_nogpu(); }
-
-void cv::gpu::HoughLinesP(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, int, int) { throw_nogpu(); }
-
-void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, int, float, float, int, int, int, int, int) { throw_nogpu(); }
-void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, HoughCirclesBuf&, int, float, float, int, int, int, int, int) { throw_nogpu(); }
-void cv::gpu::HoughCirclesDownload(const GpuMat&, OutputArray) { throw_nogpu(); }
-
 Ptr<GeneralizedHough_GPU> cv::gpu::GeneralizedHough_GPU::create(int) { throw_nogpu(); return Ptr<GeneralizedHough_GPU>(); }
 cv::gpu::GeneralizedHough_GPU::~GeneralizedHough_GPU() {}
 void cv::gpu::GeneralizedHough_GPU::setTemplate(const GpuMat&, int, Point) { throw_nogpu(); }
@@ -77,299 +67,6 @@ namespace cv { namespace gpu { namespace device
     }
 }}}
 
-//////////////////////////////////////////////////////////
-// HoughLines
-
-namespace cv { namespace gpu { namespace device
-{
-    namespace hough
-    {
-        void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20);
-        int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort);
-    }
-}}}
-
-void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, float rho, float theta, int threshold, bool doSort, int maxLines)
-{
-    HoughLinesBuf buf;
-    HoughLines(src, lines, buf, rho, theta, threshold, doSort, maxLines);
-}
-
-void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int threshold, bool doSort, int maxLines)
-{
-    using namespace cv::gpu::device::hough;
-
-    CV_Assert(src.type() == CV_8UC1);
-    CV_Assert(src.cols < std::numeric_limits<unsigned short>::max());
-    CV_Assert(src.rows < std::numeric_limits<unsigned short>::max());
-
-    ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list);
-    unsigned int* srcPoints = buf.list.ptr<unsigned int>();
-
-    const int pointsCount = buildPointList_gpu(src, srcPoints);
-    if (pointsCount == 0)
-    {
-        lines.release();
-        return;
-    }
-
-    const int numangle = cvRound(CV_PI / theta);
-    const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho);
-    CV_Assert(numangle > 0 && numrho > 0);
-
-    ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum);
-    buf.accum.setTo(Scalar::all(0));
-
-    DeviceInfo devInfo;
-    linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, devInfo.sharedMemPerBlock(), devInfo.supports(FEATURE_SET_COMPUTE_20));
-
-    ensureSizeIsEnough(2, maxLines, CV_32FC2, lines);
-
-    int linesCount = linesGetResult_gpu(buf.accum, lines.ptr<float2>(0), lines.ptr<int>(1), maxLines, rho, theta, threshold, doSort);
-    if (linesCount > 0)
-        lines.cols = linesCount;
-    else
-        lines.release();
-}
-
-void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, OutputArray h_votes_)
-{
-    if (d_lines.empty())
-    {
-        h_lines_.release();
-        if (h_votes_.needed())
-            h_votes_.release();
-        return;
-    }
-
-    CV_Assert(d_lines.rows == 2 && d_lines.type() == CV_32FC2);
-
-    h_lines_.create(1, d_lines.cols, CV_32FC2);
-    Mat h_lines = h_lines_.getMat();
-    d_lines.row(0).download(h_lines);
-
-    if (h_votes_.needed())
-    {
-        h_votes_.create(1, d_lines.cols, CV_32SC1);
-        Mat h_votes = h_votes_.getMat();
-        GpuMat d_votes(1, d_lines.cols, CV_32SC1, const_cast<int*>(d_lines.ptr<int>(1)));
-        d_votes.download(h_votes);
-    }
-}
-
-//////////////////////////////////////////////////////////
-// HoughLinesP
-
-namespace cv { namespace gpu { namespace device
-{
-    namespace hough
-    {
-        int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength);
-    }
-}}}
-
-void cv::gpu::HoughLinesP(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int minLineLength, int maxLineGap, int maxLines)
-{
-    using namespace cv::gpu::device::hough;
-
-    CV_Assert( src.type() == CV_8UC1 );
-    CV_Assert( src.cols < std::numeric_limits<unsigned short>::max() );
-    CV_Assert( src.rows < std::numeric_limits<unsigned short>::max() );
-
-    ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list);
-    unsigned int* srcPoints = buf.list.ptr<unsigned int>();
-
-    const int pointsCount = buildPointList_gpu(src, srcPoints);
-    if (pointsCount == 0)
-    {
-        lines.release();
-        return;
-    }
-
-    const int numangle = cvRound(CV_PI / theta);
-    const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho);
-    CV_Assert( numangle > 0 && numrho > 0 );
-
-    ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum);
-    buf.accum.setTo(Scalar::all(0));
-
-    DeviceInfo devInfo;
-    linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, devInfo.sharedMemPerBlock(), devInfo.supports(FEATURE_SET_COMPUTE_20));
-
-    ensureSizeIsEnough(1, maxLines, CV_32SC4, lines);
-
-    int linesCount = houghLinesProbabilistic_gpu(src, buf.accum, lines.ptr<int4>(), maxLines, rho, theta, maxLineGap, minLineLength);
-
-    if (linesCount > 0)
-        lines.cols = linesCount;
-    else
-        lines.release();
-}
-
-//////////////////////////////////////////////////////////
-// HoughCircles
-
-namespace cv { namespace gpu { namespace device
-{
-    namespace hough
-    {
-        void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp);
-        int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold);
-        int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count,
-                                   float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20);
-    }
-}}}
-
-void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, int method, float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles)
-{
-    HoughCirclesBuf buf;
-    HoughCircles(src, circles, buf, method, dp, minDist, cannyThreshold, votesThreshold, minRadius, maxRadius, maxCircles);
-}
-
-void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, HoughCirclesBuf& buf, int method,
-                           float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles)
-{
-    using namespace cv::gpu::device::hough;
-
-    CV_Assert(src.type() == CV_8UC1);
-    CV_Assert(src.cols < std::numeric_limits<unsigned short>::max());
-    CV_Assert(src.rows < std::numeric_limits<unsigned short>::max());
-    CV_Assert(method == CV_HOUGH_GRADIENT);
-    CV_Assert(dp > 0);
-    CV_Assert(minRadius > 0 && maxRadius > minRadius);
-    CV_Assert(cannyThreshold > 0);
-    CV_Assert(votesThreshold > 0);
-    CV_Assert(maxCircles > 0);
-
-    const float idp = 1.0f / dp;
-
-    cv::gpu::Canny(src, buf.cannyBuf, buf.edges, std::max(cannyThreshold / 2, 1), cannyThreshold);
-
-    ensureSizeIsEnough(2, src.size().area(), CV_32SC1, buf.list);
-    unsigned int* srcPoints = buf.list.ptr<unsigned int>(0);
-    unsigned int* centers = buf.list.ptr<unsigned int>(1);
-
-    const int pointsCount = buildPointList_gpu(buf.edges, srcPoints);
-    if (pointsCount == 0)
-    {
-        circles.release();
-        return;
-    }
-
-    ensureSizeIsEnough(cvCeil(src.rows * idp) + 2, cvCeil(src.cols * idp) + 2, CV_32SC1, buf.accum);
-    buf.accum.setTo(Scalar::all(0));
-
-    circlesAccumCenters_gpu(srcPoints, pointsCount, buf.cannyBuf.dx, buf.cannyBuf.dy, buf.accum, minRadius, maxRadius, idp);
-
-    int centersCount = buildCentersList_gpu(buf.accum, centers, votesThreshold);
-    if (centersCount == 0)
-    {
-        circles.release();
-        return;
-    }
-
-    if (minDist > 1)
-    {
-        cv::AutoBuffer<ushort2> oldBuf_(centersCount);
-        cv::AutoBuffer<ushort2> newBuf_(centersCount);
-        int newCount = 0;
-
-        ushort2* oldBuf = oldBuf_;
-        ushort2* newBuf = newBuf_;
-
-        cudaSafeCall( cudaMemcpy(oldBuf, centers, centersCount * sizeof(ushort2), cudaMemcpyDeviceToHost) );
-
-        const int cellSize = cvRound(minDist);
-        const int gridWidth = (src.cols + cellSize - 1) / cellSize;
-        const int gridHeight = (src.rows + cellSize - 1) / cellSize;
-
-        std::vector< std::vector<ushort2> > grid(gridWidth * gridHeight);
-
-        const float minDist2 = minDist * minDist;
-
-        for (int i = 0; i < centersCount; ++i)
-        {
-            ushort2 p = oldBuf[i];
-
-            bool good = true;
-
-            int xCell = static_cast<int>(p.x / cellSize);
-            int yCell = static_cast<int>(p.y / cellSize);
-
-            int x1 = xCell - 1;
-            int y1 = yCell - 1;
-            int x2 = xCell + 1;
-            int y2 = yCell + 1;
-
-            // boundary check
-            x1 = std::max(0, x1);
-            y1 = std::max(0, y1);
-            x2 = std::min(gridWidth - 1, x2);
-            y2 = std::min(gridHeight - 1, y2);
-
-            for (int yy = y1; yy <= y2; ++yy)
-            {
-                for (int xx = x1; xx <= x2; ++xx)
-                {
-                    vector<ushort2>& m = grid[yy * gridWidth + xx];
-
-                    for(size_t j = 0; j < m.size(); ++j)
-                    {
-                        float dx = (float)(p.x - m[j].x);
-                        float dy = (float)(p.y - m[j].y);
-
-                        if (dx * dx + dy * dy < minDist2)
-                        {
-                            good = false;
-                            goto break_out;
-                        }
-                    }
-                }
-            }
-
-            break_out:
-
-            if(good)
-            {
-                grid[yCell * gridWidth + xCell].push_back(p);
-
-                newBuf[newCount++] = p;
-            }
-        }
-
-        cudaSafeCall( cudaMemcpy(centers, newBuf, newCount * sizeof(unsigned int), cudaMemcpyHostToDevice) );
-        centersCount = newCount;
-    }
-
-    ensureSizeIsEnough(1, maxCircles, CV_32FC3, circles);
-
-    const int circlesCount = circlesAccumRadius_gpu(centers, centersCount, srcPoints, pointsCount, circles.ptr<float3>(), maxCircles,
-                                                    dp, minRadius, maxRadius, votesThreshold, deviceSupports(FEATURE_SET_COMPUTE_20));
-
-    if (circlesCount > 0)
-        circles.cols = circlesCount;
-    else
-        circles.release();
-}
-
-void cv::gpu::HoughCirclesDownload(const GpuMat& d_circles, cv::OutputArray h_circles_)
-{
-    if (d_circles.empty())
-    {
-        h_circles_.release();
-        return;
-    }
-
-    CV_Assert(d_circles.rows == 1 && d_circles.type() == CV_32FC3);
-
-    h_circles_.create(1, d_circles.cols, CV_32FC3);
-    Mat h_circles = h_circles_.getMat();
-    d_circles.download(h_circles);
-}
-
-//////////////////////////////////////////////////////////
-// GeneralizedHough
-
 namespace cv { namespace gpu { namespace device
 {
     namespace hough

modules/gpu/src/hough_circles.cpp

@@ -0,0 +1,223 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "precomp.hpp"
+
+using namespace std;
+using namespace cv;
+using namespace cv::gpu;
+
+#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
+
+void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, int, float, float, int, int, int, int, int) { throw_nogpu(); }
+void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, HoughCirclesBuf&, int, float, float, int, int, int, int, int) { throw_nogpu(); }
+void cv::gpu::HoughCirclesDownload(const GpuMat&, OutputArray) { throw_nogpu(); }
+
+#else /* !defined (HAVE_CUDA) */
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        int buildPointList_gpu(PtrStepSzb src, unsigned int* list);
+    }
+}}}
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp);
+        int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold);
+        int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count,
+                                   float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20);
+    }
+}}}
+
+void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, int method, float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles)
+{
+    HoughCirclesBuf buf;
+    HoughCircles(src, circles, buf, method, dp, minDist, cannyThreshold, votesThreshold, minRadius, maxRadius, maxCircles);
+}
+
+void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, HoughCirclesBuf& buf, int method,
+                           float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles)
+{
+    using namespace cv::gpu::device::hough;
+
+    CV_Assert(src.type() == CV_8UC1);
+    CV_Assert(src.cols < std::numeric_limits<unsigned short>::max());
+    CV_Assert(src.rows < std::numeric_limits<unsigned short>::max());
+    CV_Assert(method == CV_HOUGH_GRADIENT);
+    CV_Assert(dp > 0);
+    CV_Assert(minRadius > 0 && maxRadius > minRadius);
+    CV_Assert(cannyThreshold > 0);
+    CV_Assert(votesThreshold > 0);
+    CV_Assert(maxCircles > 0);
+
+    const float idp = 1.0f / dp;
+
+    cv::gpu::Canny(src, buf.cannyBuf, buf.edges, std::max(cannyThreshold / 2, 1), cannyThreshold);
+
+    ensureSizeIsEnough(2, src.size().area(), CV_32SC1, buf.list);
+    unsigned int* srcPoints = buf.list.ptr<unsigned int>(0);
+    unsigned int* centers = buf.list.ptr<unsigned int>(1);
+
+    const int pointsCount = buildPointList_gpu(buf.edges, srcPoints);
+    if (pointsCount == 0)
+    {
+        circles.release();
+        return;
+    }
+
+    ensureSizeIsEnough(cvCeil(src.rows * idp) + 2, cvCeil(src.cols * idp) + 2, CV_32SC1, buf.accum);
+    buf.accum.setTo(Scalar::all(0));
+
+    circlesAccumCenters_gpu(srcPoints, pointsCount, buf.cannyBuf.dx, buf.cannyBuf.dy, buf.accum, minRadius, maxRadius, idp);
+
+    int centersCount = buildCentersList_gpu(buf.accum, centers, votesThreshold);
+    if (centersCount == 0)
+    {
+        circles.release();
+        return;
+    }
+
+    if (minDist > 1)
+    {
+        cv::AutoBuffer<ushort2> oldBuf_(centersCount);
+        cv::AutoBuffer<ushort2> newBuf_(centersCount);
+        int newCount = 0;
+
+        ushort2* oldBuf = oldBuf_;
+        ushort2* newBuf = newBuf_;
+
+        cudaSafeCall( cudaMemcpy(oldBuf, centers, centersCount * sizeof(ushort2), cudaMemcpyDeviceToHost) );
+
+        const int cellSize = cvRound(minDist);
+        const int gridWidth = (src.cols + cellSize - 1) / cellSize;
+        const int gridHeight = (src.rows + cellSize - 1) / cellSize;
+
+        std::vector< std::vector<ushort2> > grid(gridWidth * gridHeight);
+
+        const float minDist2 = minDist * minDist;
+
+        for (int i = 0; i < centersCount; ++i)
+        {
+            ushort2 p = oldBuf[i];
+
+            bool good = true;
+
+            int xCell = static_cast<int>(p.x / cellSize);
+            int yCell = static_cast<int>(p.y / cellSize);
+
+            int x1 = xCell - 1;
+            int y1 = yCell - 1;
+            int x2 = xCell + 1;
+            int y2 = yCell + 1;
+
+            // boundary check
+            x1 = std::max(0, x1);
+            y1 = std::max(0, y1);
+            x2 = std::min(gridWidth - 1, x2);
+            y2 = std::min(gridHeight - 1, y2);
+
+            for (int yy = y1; yy <= y2; ++yy)
+            {
+                for (int xx = x1; xx <= x2; ++xx)
+                {
+                    vector<ushort2>& m = grid[yy * gridWidth + xx];
+
+                    for(size_t j = 0; j < m.size(); ++j)
+                    {
+                        float dx = (float)(p.x - m[j].x);
+                        float dy = (float)(p.y - m[j].y);
+
+                        if (dx * dx + dy * dy < minDist2)
+                        {
+                            good = false;
+                            goto break_out;
+                        }
+                    }
+                }
+            }
+
+            break_out:
+
+            if(good)
+            {
+                grid[yCell * gridWidth + xCell].push_back(p);
+
+                newBuf[newCount++] = p;
+            }
+        }
+
+        cudaSafeCall( cudaMemcpy(centers, newBuf, newCount * sizeof(unsigned int), cudaMemcpyHostToDevice) );
+        centersCount = newCount;
+    }
+
+    ensureSizeIsEnough(1, maxCircles, CV_32FC3, circles);
+
+    const int circlesCount = circlesAccumRadius_gpu(centers, centersCount, srcPoints, pointsCount, circles.ptr<float3>(), maxCircles,
+                                                    dp, minRadius, maxRadius, votesThreshold, deviceSupports(FEATURE_SET_COMPUTE_20));
+
+    if (circlesCount > 0)
+        circles.cols = circlesCount;
+    else
+        circles.release();
+}
+
+void cv::gpu::HoughCirclesDownload(const GpuMat& d_circles, cv::OutputArray h_circles_)
+{
+    if (d_circles.empty())
+    {
+        h_circles_.release();
+        return;
+    }
+
+    CV_Assert(d_circles.rows == 1 && d_circles.type() == CV_32FC3);
+
+    h_circles_.create(1, d_circles.cols, CV_32FC3);
+    Mat h_circles = h_circles_.getMat();
+    d_circles.download(h_circles);
+}
+
+#endif /* !defined (HAVE_CUDA) */

modules/gpu/src/hough_lines.cpp

@@ -0,0 +1,142 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "precomp.hpp"
+
+using namespace std;
+using namespace cv;
+using namespace cv::gpu;
+
+#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
+
+void cv::gpu::HoughLines(const GpuMat&, GpuMat&, float, float, int, bool, int) { throw_nogpu(); }
+void cv::gpu::HoughLines(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, bool, int) { throw_nogpu(); }
+void cv::gpu::HoughLinesDownload(const GpuMat&, OutputArray, OutputArray) { throw_nogpu(); }
+
+#else /* !defined (HAVE_CUDA) */
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        int buildPointList_gpu(PtrStepSzb src, unsigned int* list);
+    }
+}}}
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20);
+        int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort);
+    }
+}}}
+
+void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, float rho, float theta, int threshold, bool doSort, int maxLines)
+{
+    HoughLinesBuf buf;
+    HoughLines(src, lines, buf, rho, theta, threshold, doSort, maxLines);
+}
+
+void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int threshold, bool doSort, int maxLines)
+{
+    using namespace cv::gpu::device::hough;
+
+    CV_Assert(src.type() == CV_8UC1);
+    CV_Assert(src.cols < std::numeric_limits<unsigned short>::max());
+    CV_Assert(src.rows < std::numeric_limits<unsigned short>::max());
+
+    ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list);
+    unsigned int* srcPoints = buf.list.ptr<unsigned int>();
+
+    const int pointsCount = buildPointList_gpu(src, srcPoints);
+    if (pointsCount == 0)
+    {
+        lines.release();
+        return;
+    }
+
+    const int numangle = cvRound(CV_PI / theta);
+    const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho);
+    CV_Assert(numangle > 0 && numrho > 0);
+
+    ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum);
+    buf.accum.setTo(Scalar::all(0));
+
+    DeviceInfo devInfo;
+    linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, devInfo.sharedMemPerBlock(), devInfo.supports(FEATURE_SET_COMPUTE_20));
+
+    ensureSizeIsEnough(2, maxLines, CV_32FC2, lines);
+
+    int linesCount = linesGetResult_gpu(buf.accum, lines.ptr<float2>(0), lines.ptr<int>(1), maxLines, rho, theta, threshold, doSort);
+    if (linesCount > 0)
+        lines.cols = linesCount;
+    else
+        lines.release();
+}
+
+void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, OutputArray h_votes_)
+{
+    if (d_lines.empty())
+    {
+        h_lines_.release();
+        if (h_votes_.needed())
+            h_votes_.release();
+        return;
+    }
+
+    CV_Assert(d_lines.rows == 2 && d_lines.type() == CV_32FC2);
+
+    h_lines_.create(1, d_lines.cols, CV_32FC2);
+    Mat h_lines = h_lines_.getMat();
+    d_lines.row(0).download(h_lines);
+
+    if (h_votes_.needed())
+    {
+        h_votes_.create(1, d_lines.cols, CV_32SC1);
+        Mat h_votes = h_votes_.getMat();
+        GpuMat d_votes(1, d_lines.cols, CV_32SC1, const_cast<int*>(d_lines.ptr<int>(1)));
+        d_votes.download(h_votes);
+    }
+}
+
+#endif /* !defined (HAVE_CUDA) */

modules/gpu/src/hough_segments.cpp

@@ -0,0 +1,110 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "precomp.hpp"
+
+using namespace std;
+using namespace cv;
+using namespace cv::gpu;
+
+#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
+
+void cv::gpu::HoughLinesP(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, int, int) { throw_nogpu(); }
+
+#else /* !defined (HAVE_CUDA) */
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        int buildPointList_gpu(PtrStepSzb src, unsigned int* list);
+    }
+}}}
+
+namespace cv { namespace gpu { namespace device
+{
+    namespace hough
+    {
+        void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20);
+        int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength);
+    }
+}}}
+
+void cv::gpu::HoughLinesP(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int minLineLength, int maxLineGap, int maxLines)
+{
+    using namespace cv::gpu::device::hough;
+
+    CV_Assert( src.type() == CV_8UC1 );
+    CV_Assert( src.cols < std::numeric_limits<unsigned short>::max() );
+    CV_Assert( src.rows < std::numeric_limits<unsigned short>::max() );
+
+    ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list);
+    unsigned int* srcPoints = buf.list.ptr<unsigned int>();
+
+    const int pointsCount = buildPointList_gpu(src, srcPoints);
+    if (pointsCount == 0)
+    {
+        lines.release();
+        return;
+    }
+
+    const int numangle = cvRound(CV_PI / theta);
+    const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho);
+    CV_Assert( numangle > 0 && numrho > 0 );
+
+    ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum);
+    buf.accum.setTo(Scalar::all(0));
+
+    DeviceInfo devInfo;
+    linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, devInfo.sharedMemPerBlock(), devInfo.supports(FEATURE_SET_COMPUTE_20));
+
+    ensureSizeIsEnough(1, maxLines, CV_32SC4, lines);
+
+    int linesCount = houghLinesProbabilistic_gpu(src, buf.accum, lines.ptr<int4>(), maxLines, rho, theta, maxLineGap, minLineLength);
+
+    if (linesCount > 0)
+        lines.cols = linesCount;
+    else
+        lines.release();
+}
+
+#endif /* !defined (HAVE_CUDA) */

modules/gpu/src/imgproc.cpp

@@ -244,6 +244,10 @@ void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyz, const Mat& Q,
 ////////////////////////////////////////////////////////////////////////
 // copyMakeBorder
 
+// Disable NPP for this file
+//#define USE_NPP
+#undef USE_NPP
+
 namespace cv { namespace gpu { namespace device
 {
     namespace imgproc
@@ -279,6 +283,7 @@ void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom
 
     cudaStream_t stream = StreamAccessor::getStream(s);
 
+#ifdef USE_NPP
     if (borderType == BORDER_CONSTANT && (src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1))
     {
         NppiSize srcsz;
@@ -328,6 +333,7 @@ void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom
             cudaSafeCall( cudaDeviceSynchronize() );
     }
     else
+#endif
     {
         typedef void (*caller_t)(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream);
         static const caller_t callers[6][4] =
@@ -1485,6 +1491,8 @@ void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result,
 
 void cv::gpu::CannyBuf::create(const Size& image_size, int apperture_size)
 {
+    CV_Assert(image_size.width < std::numeric_limits<short>::max() && image_size.height < std::numeric_limits<short>::max());
+
     if (apperture_size > 0)
     {
         ensureSizeIsEnough(image_size, CV_32SC1, dx);
@@ -1500,8 +1508,8 @@ void cv::gpu::CannyBuf::create(const Size& image_size, int apperture_size)
     ensureSizeIsEnough(image_size, CV_32FC1, mag);
     ensureSizeIsEnough(image_size, CV_32SC1, map);
 
-    ensureSizeIsEnough(1, image_size.area(), CV_16UC2, st1);
-    ensureSizeIsEnough(1, image_size.area(), CV_16UC2, st2);
+    ensureSizeIsEnough(1, image_size.area(), CV_16SC2, st1);
+    ensureSizeIsEnough(1, image_size.area(), CV_16SC2, st2);
 }
 
 void cv::gpu::CannyBuf::release()
@@ -1521,9 +1529,9 @@ namespace canny
 
     void calcMap(PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, PtrStepSzi map, float low_thresh, float high_thresh);
 
-    void edgesHysteresisLocal(PtrStepSzi map, ushort2* st1);
+    void edgesHysteresisLocal(PtrStepSzi map, short2* st1);
 
-    void edgesHysteresisGlobal(PtrStepSzi map, ushort2* st1, ushort2* st2);
+    void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2);
 
     void getEdges(PtrStepSzi map, PtrStepSzb dst);
 }
@@ -1537,9 +1545,9 @@ namespace
         buf.map.setTo(Scalar::all(0));
         calcMap(dx, dy, buf.mag, buf.map, low_thresh, high_thresh);
 
-        edgesHysteresisLocal(buf.map, buf.st1.ptr<ushort2>());
+        edgesHysteresisLocal(buf.map, buf.st1.ptr<short2>());
 
-        edgesHysteresisGlobal(buf.map, buf.st1.ptr<ushort2>(), buf.st2.ptr<ushort2>());
+        edgesHysteresisGlobal(buf.map, buf.st1.ptr<short2>(), buf.st2.ptr<short2>());
 
         getEdges(buf.map, dst);
     }

modules/gpu/test/nvidia/TestHaarCascadeLoader.cpp

@@ -98,7 +98,7 @@ bool TestHaarCascadeLoader::process()
     NCV_SET_SKIP_COND(this->allocatorGPU.get()->isCounting());
     NCV_SKIP_COND_BEGIN
 
-    const std::string testNvbinName = "test.nvbin";
+    const std::string testNvbinName = cv::tempfile("test.nvbin");
     ncvStat = ncvHaarLoadFromFile_host(this->cascadeName, haar, h_HaarStages, h_HaarNodes, h_HaarFeatures);
     ncvAssertReturn(ncvStat == NCV_SUCCESS, false);
 

modules/gpu/test/test_color.cpp

@@ -715,7 +715,7 @@ GPU_TEST_P(CvtColor, BGR2YCrCb)
     cv::Mat dst_gold;
     cv::cvtColor(src, dst_gold, cv::COLOR_BGR2YCrCb);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_32F ? 1e-2 : 1);
 }
 
 GPU_TEST_P(CvtColor, RGB2YCrCb)
@@ -728,7 +728,7 @@ GPU_TEST_P(CvtColor, RGB2YCrCb)
     cv::Mat dst_gold;
     cv::cvtColor(src, dst_gold, cv::COLOR_RGB2YCrCb);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_32F ? 1e-2 : 1);
 }
 
 GPU_TEST_P(CvtColor, BGR2YCrCb4)
@@ -749,7 +749,7 @@ GPU_TEST_P(CvtColor, BGR2YCrCb4)
     cv::split(h_dst, channels);
     cv::merge(channels, 3, h_dst);
 
-    EXPECT_MAT_NEAR(dst_gold, h_dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, h_dst, depth == CV_32F ? 1e-2 : 1);
 }
 
 GPU_TEST_P(CvtColor, RGBA2YCrCb4)
@@ -771,7 +771,7 @@ GPU_TEST_P(CvtColor, RGBA2YCrCb4)
     cv::split(h_dst, channels);
     cv::merge(channels, 3, h_dst);
 
-    EXPECT_MAT_NEAR(dst_gold, h_dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, h_dst, depth == CV_32F ? 1e-2 : 1);
 }
 
 GPU_TEST_P(CvtColor, YCrCb2BGR)

modules/gpu/test/test_core.cpp

@@ -3582,7 +3582,7 @@ GPU_TEST_P(Normalize, WithOutMask)
     cv::Mat dst_gold;
     cv::normalize(src, dst_gold, alpha, beta, norm_type, type);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, dst, type < CV_32F ? 1.0 : 1e-4);
 }
 
 GPU_TEST_P(Normalize, WithMask)
@@ -3598,7 +3598,7 @@ GPU_TEST_P(Normalize, WithMask)
     dst_gold.setTo(cv::Scalar::all(0));
     cv::normalize(src, dst_gold, alpha, beta, norm_type, type, mask);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, dst, type < CV_32F ? 1.0 : 1e-4);
 }
 
 INSTANTIATE_TEST_CASE_P(GPU_Core, Normalize, testing::Combine(

modules/gpu/test/test_denoising.cpp

@@ -114,6 +114,7 @@ GPU_TEST_P(BruteForceNonLocalMeans, Regression)
 
     cv::Mat bgr  = readImage("denoising/lena_noised_gaussian_sigma=20_multi_0.png", cv::IMREAD_COLOR);
     ASSERT_FALSE(bgr.empty());
+    cv::resize(bgr, bgr, cv::Size(256, 256));
 
     cv::Mat gray;
     cv::cvtColor(bgr, gray, CV_BGR2GRAY);
@@ -130,6 +131,8 @@ GPU_TEST_P(BruteForceNonLocalMeans, Regression)
     cv::Mat bgr_gold  = readImage("denoising/nlm_denoised_lena_bgr.png", cv::IMREAD_COLOR);
     cv::Mat gray_gold  = readImage("denoising/nlm_denoised_lena_gray.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(bgr_gold.empty() || gray_gold.empty());
+    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);

modules/gpu/test/test_gpumat.cpp

@@ -281,7 +281,7 @@ GPU_TEST_P(ConvertTo, WithOutScaling)
         cv::Mat dst_gold;
         src.convertTo(dst_gold, depth2);
 
-        EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
+        EXPECT_MAT_NEAR(dst_gold, dst, depth2 < CV_32F ? 1.0 : 1e-4);
     }
 }
 

modules/gpu/test/test_imgproc.cpp

@@ -563,6 +563,8 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, Blend, testing::Combine(
     testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)),
     WHOLE_SUBMAT));
 
+#ifdef HAVE_CUFFT
+
 ////////////////////////////////////////////////////////
 // Convolve
 
@@ -1090,6 +1092,8 @@ GPU_TEST_P(Dft, R2CThenC2R)
 
 INSTANTIATE_TEST_CASE_P(GPU_ImgProc, Dft, ALL_DEVICES);
 
+#endif
+
 ///////////////////////////////////////////////////////////////////////////////////////////////////////
 // CornerHarris
 

modules/gpu/test/test_optflow.cpp

@@ -483,13 +483,15 @@ GPU_TEST_P(OpticalFlowBM, Accuracy)
 
     cv::Mat frame0 = readImage("opticalflow/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame0.empty());
+    cv::resize(frame0, frame0, cv::Size(), 0.5, 0.5);
 
     cv::Mat frame1 = readImage("opticalflow/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame1.empty());
+    cv::resize(frame1, frame1, cv::Size(), 0.5, 0.5);
 
-    cv::Size block_size(16, 16);
+    cv::Size block_size(8, 8);
     cv::Size shift_size(1, 1);
-    cv::Size max_range(16, 16);
+    cv::Size max_range(8, 8);
 
     cv::gpu::GpuMat d_velx, d_vely, buf;
     cv::gpu::calcOpticalFlowBM(loadMat(frame0), loadMat(frame1),

modules/highgui/CMakeLists.txt

@@ -315,7 +315,7 @@ if(WIN32 AND WITH_FFMPEG)
                        COMMENT "Copying ${ffmpeg_path} to the output directory")
   endif()
 
-  install(FILES "${ffmpeg_path}" DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main RENAME "${ffmpeg_bare_name_ver}")
+  install(FILES "${ffmpeg_path}" DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT libs RENAME "${ffmpeg_bare_name_ver}")
 endif()
 
 ocv_add_accuracy_tests()

modules/highgui/src/precomp.hpp

@@ -57,6 +57,14 @@
 #include <assert.h>
 
 #if defined WIN32 || defined WINCE
+    #if !defined _WIN32_WINNT
+        #ifdef HAVE_MSMF
+            #define _WIN32_WINNT 0x0600 // Windows Vista
+        #else
+            #define _WIN32_WINNT 0x0500 // Windows 2000
+        #endif
+    #endif
+
     #include <windows.h>
     #undef small
     #undef min

modules/highgui/src/window_w32.cpp

@@ -43,27 +43,11 @@
 
 #if defined WIN32 || defined _WIN32
 
-#define COMPILE_MULTIMON_STUBS // Required for multi-monitor support
-#ifndef _MULTIMON_USE_SECURE_CRT
-#  define _MULTIMON_USE_SECURE_CRT 0 // some MinGW platforms have no strncpy_s
-#endif
-
-#if defined SM_CMONITORS && !defined MONITOR_DEFAULTTONEAREST
-#  define MONITOR_DEFAULTTONULL       0x00000000
-#  define MONITOR_DEFAULTTOPRIMARY    0x00000001
-#  define MONITOR_DEFAULTTONEAREST    0x00000002
-#  define MONITORINFOF_PRIMARY        0x00000001
-#endif
-#ifndef __inout
-#  define __inout
-#endif
-
 #ifdef __GNUC__
 #  pragma GCC diagnostic ignored "-Wmissing-declarations"
 #endif
 
 #include <commctrl.h>
-#include <winuser.h>
 #include <stdlib.h>
 #include <string.h>
 #include <stdio.h>

modules/highgui/test/test_ffmpeg.cpp

@@ -88,7 +88,7 @@ public:
             stringstream s;
             s << tag;
 
-            const string filename = "output_"+s.str()+".avi";
+            const string filename = tempfile((s.str()+".avi").c_str());
 
             try
             {

modules/highgui/test/test_video_io.cpp

@@ -332,9 +332,7 @@ void CV_HighGuiTest::VideoTest(const string& dir, const cvtest::VideoFormat& fmt
         }
     }
 
-    printf("Before saved release for %s\n", tmp_name.c_str());
     cvReleaseCapture( &saved );
-    printf("After release\n");
 
     ts->printf(ts->LOG, "end test function : ImagesVideo \n");
 }

modules/imgproc/src/smooth.cpp

@@ -718,7 +718,7 @@ void cv::boxFilter( InputArray _src, OutputArray _dst, int ddepth,
         ddepth = sdepth;
     _dst.create( src.size(), CV_MAKETYPE(ddepth, cn) );
     Mat dst = _dst.getMat();
-    if( borderType != BORDER_CONSTANT && normalize )
+    if( borderType != BORDER_CONSTANT && normalize && (borderType & BORDER_ISOLATED) != 0 )
     {
         if( src.rows == 1 )
             ksize.height = 1;

modules/imgproc/test/test_filter.cpp

@@ -1886,3 +1886,35 @@ protected:
 };
 
 TEST(Imgproc_Filtering, supportedFormats) { CV_FilterSupportedFormatsTest test; test.safe_run(); }
+
+TEST(Imgproc_Blur, borderTypes)
+{
+    Size kernelSize(3, 3);
+
+    /// ksize > src_roi.size()
+    Mat src(3, 3, CV_8UC1, cv::Scalar::all(255)), dst;
+    Mat src_roi = src(Rect(1, 1, 1, 1));
+    src_roi.setTo(cv::Scalar::all(0));
+
+    // should work like !BORDER_ISOLATED
+    blur(src_roi, dst, kernelSize, Point(-1, -1), BORDER_REPLICATE);
+    EXPECT_EQ(227, dst.at<uchar>(0, 0));
+
+    // should work like BORDER_ISOLATED
+    blur(src_roi, dst, kernelSize, Point(-1, -1), BORDER_REPLICATE | BORDER_ISOLATED);
+    EXPECT_EQ(0, dst.at<uchar>(0, 0));
+
+    /// ksize <= src_roi.size()
+    src = Mat(5, 5, CV_8UC1, cv::Scalar(255));
+    src_roi = src(Rect(1, 1, 3, 3));
+    src_roi.setTo(0);
+    src.at<uchar>(2, 2) = 255;
+
+    // should work like !BORDER_ISOLATED
+    blur(src_roi, dst, kernelSize, Point(-1, -1), BORDER_REPLICATE);
+    Mat expected_dst =
+            (Mat_<uchar>(3, 3) << 170, 113, 170, 113, 28, 113, 170, 113, 170);
+    EXPECT_EQ(expected_dst.type(), dst.type());
+    EXPECT_EQ(expected_dst.size(), dst.size());
+    EXPECT_DOUBLE_EQ(0.0, cvtest::norm(expected_dst, dst, NORM_INF));
+}

modules/java/CMakeLists.txt

@@ -175,7 +175,7 @@ foreach(java_file ${step3_input_files})
 
   if(ANDROID)
     get_filename_component(install_subdir "${java_file_name}" PATH)
-    install(FILES "${output_name}" DESTINATION "${JAVA_INSTALL_ROOT}/src/org/opencv/${install_subdir}" COMPONENT main)
+    install(FILES "${output_name}" DESTINATION "${JAVA_INSTALL_ROOT}/src/org/opencv/${install_subdir}" COMPONENT java)
   endif()
 endforeach()
 
@@ -189,7 +189,7 @@ if(ANDROID)
 
     if(NOT file MATCHES "jni/.+")
       get_filename_component(install_subdir "${file}" PATH)
-      install(FILES "${OpenCV_BINARY_DIR}/${file}" DESTINATION "${JAVA_INSTALL_ROOT}/${install_subdir}" COMPONENT main)
+      install(FILES "${OpenCV_BINARY_DIR}/${file}" DESTINATION "${JAVA_INSTALL_ROOT}/${install_subdir}" COMPONENT java)
     endif()
   endforeach()
 
@@ -225,11 +225,11 @@ if(ANDROID AND ANDROID_EXECUTABLE)
   list(APPEND copied_files ${lib_target_files} "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}")
   list(APPEND step3_input_files "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}")
 
-  install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_PROJECT_PROPERTIES_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT main)
-  install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT main)
+  install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_PROJECT_PROPERTIES_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT java)
+  install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT java)
   # creating empty 'gen' and 'res' folders
-  install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/gen\")" COMPONENT main)
-  install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/res\")" COMPONENT main)
+  install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/gen\")" COMPONENT java)
+  install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/res\")" COMPONENT java)
 endif(ANDROID AND ANDROID_EXECUTABLE)
 
 set(step3_depends ${step2_depends} ${step3_input_files} ${copied_files})
@@ -282,7 +282,7 @@ else(ANDROID)
   else(WIN32)
     set(JAR_INSTALL_DIR share/OpenCV/java)
   endif(WIN32)
-  install(FILES ${JAR_FILE} DESTINATION ${JAR_INSTALL_DIR} COMPONENT main)
+  install(FILES ${JAR_FILE} DESTINATION ${JAR_INSTALL_DIR} COMPONENT java)
 endif(ANDROID)
 
 # step 5: build native part
@@ -353,17 +353,17 @@ endif()
 
 if(ANDROID)
   ocv_install_target(${the_module} EXPORT OpenCVModules
-          LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main
-          ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main)
+          LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT java
+          ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT java)
 else()
   if(NOT INSTALL_CREATE_DISTRIB)
     ocv_install_target(${the_module} EXPORT OpenCVModules
-            RUNTIME DESTINATION ${JAR_INSTALL_DIR} COMPONENT main
-            LIBRARY DESTINATION ${JAR_INSTALL_DIR} COMPONENT main)
+            RUNTIME DESTINATION ${JAR_INSTALL_DIR} COMPONENT java
+            LIBRARY DESTINATION ${JAR_INSTALL_DIR} COMPONENT java)
   else()
     ocv_install_target(${the_module} EXPORT OpenCVModules
-            RUNTIME DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT main
-            LIBRARY DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT main)
+            RUNTIME DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT java
+            LIBRARY DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT java)
   endif()
 endif()
 

modules/java/generator/src/cpp/VideoCapture.cpp

@@ -329,7 +329,10 @@ JNIEXPORT jstring JNICALL Java_org_opencv_highgui_VideoCapture_n_1getSupportedPr
         VideoCapture* me = (VideoCapture*) self; //TODO: check for NULL
         union {double prop; const char* name;} u;
         u.prop = me->get(CV_CAP_PROP_SUPPORTED_PREVIEW_SIZES_STRING);
-        return env->NewStringUTF(u.name);
+        // VideoCapture::get can return 0.0 or -1.0 if it doesn't support
+        // CV_CAP_PROP_SUPPORTED_PREVIEW_SIZES_STRING
+        if (u.prop != 0.0 && u.prop != -1.0)
+            return env->NewStringUTF(u.name);
     } catch(const std::exception &e) {
         throwJavaException(env, &e, method_name);
     } catch (...) {

modules/java/generator/src/java/android+OpenCVLoader.java

@@ -48,7 +48,17 @@ public class OpenCVLoader
      */
     public static boolean initDebug()
     {
-        return StaticHelper.initOpenCV();
+        return StaticHelper.initOpenCV(false);
+    }
+
+    /**
+     * Loads and initializes OpenCV library from current application package. Roughly, it's an analog of system.loadLibrary("opencv_java").
+     * @param InitCuda load and initialize CUDA runtime libraries.
+     * @return Returns true is initialization of OpenCV was successful.
+     */
+    public static boolean initDebug(boolean InitCuda)
+    {
+        return StaticHelper.initOpenCV(InitCuda);
     }
 
     /**

modules/java/generator/src/java/android+StaticHelper.java

@@ -7,11 +7,21 @@ import android.util.Log;
 
 class StaticHelper {
 
-    public static boolean initOpenCV()
+    public static boolean initOpenCV(boolean InitCuda)
     {
         boolean result;
         String libs = "";
 
+        if(InitCuda)
+        {
+            loadLibrary("cudart");
+            loadLibrary("nppc");
+            loadLibrary("nppi");
+            loadLibrary("npps");
+            loadLibrary("cufft");
+            loadLibrary("cublas");
+        }
+
         Log.d(TAG, "Trying to get library list");
 
         try
@@ -52,7 +62,7 @@ class StaticHelper {
         try
         {
             System.loadLibrary(Name);
-            Log.d(TAG, "OpenCV libs init was ok!");
+            Log.d(TAG, "Library " + Name + " loaded");
         }
         catch(UnsatisfiedLinkError e)
         {

modules/java/generator/src/java/core+TermCriteria.java

@@ -87,7 +87,6 @@ public class TermCriteria {
 
     @Override
     public String toString() {
-        if (this == null) return "null";
         return "{ type: " + type + ", maxCount: " + maxCount + ", epsilon: " + epsilon + "}";
     }
 }

modules/nonfree/CMakeLists.txt

@@ -4,9 +4,9 @@ endif()
 
 set(the_description "Functionality with possible limitations on the use")
 ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef)
-if (ENABLE_DYNAMIC_CUDA)
-  set(HAVE_CUDA FALSE)
+if(ENABLE_DYNAMIC_CUDA)
+  add_definitions(-DDYNAMIC_CUDA_SUPPORT)
   ocv_define_module(nonfree opencv_imgproc opencv_features2d opencv_calib3d OPTIONAL opencv_ocl)
 else()
   ocv_define_module(nonfree opencv_imgproc opencv_features2d opencv_calib3d OPTIONAL opencv_gpu opencv_ocl)
-endif()
+endif()

modules/nonfree/include/opencv2/nonfree/gpu.hpp

@@ -43,11 +43,7 @@
 #ifndef __OPENCV_NONFREE_GPU_HPP__
 #define __OPENCV_NONFREE_GPU_HPP__
 
-#include "opencv2/opencv_modules.hpp"
-
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
-
-#include "opencv2/gpu/gpu.hpp"
+#include "opencv2/core/gpumat.hpp"
 
 namespace cv { namespace gpu {
 
@@ -129,6 +125,4 @@ public:
 
 } // namespace cv
 
-#endif // defined(HAVE_OPENCV_GPU)
-
 #endif // __OPENCV_NONFREE_GPU_HPP__

modules/nonfree/perf/perf_gpu.cpp

@@ -42,7 +42,9 @@
 
 #include "perf_precomp.hpp"
 
-#if defined(HAVE_OPENCV_GPU) && defined(HAVE_CUDA)
+#include "cvconfig.h"
+
+#if defined(HAVE_OPENCV_GPU) && defined(HAVE_CUDA) && !defined(DYNAMIC_CUDA_SUPPORT)
 
 #include "opencv2/ts/gpu_perf.hpp"
 

modules/nonfree/src/cuda/surf.cu

@@ -42,7 +42,7 @@
 
 #include "opencv2/opencv_modules.hpp"
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
 
 #include "opencv2/gpu/device/common.hpp"
 #include "opencv2/gpu/device/limits.hpp"

modules/nonfree/src/precomp.hpp

@@ -51,17 +51,14 @@
 #include "opencv2/imgproc/imgproc.hpp"
 #include "opencv2/core/internal.hpp"
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
-    #include "opencv2/nonfree/gpu.hpp"
+#include "opencv2/nonfree/gpu.hpp"
 
-    #if defined(HAVE_CUDA)
-        #include "opencv2/gpu/stream_accessor.hpp"
-        #include "opencv2/gpu/device/common.hpp"
-
-        static inline void throw_nogpu() { CV_Error(CV_StsNotImplemented, "The called functionality is disabled for current build or platform"); }
-    #else
-        static inline void throw_nogpu() { CV_Error(CV_GpuNotSupported, "The library is compiled without GPU support"); }
-    #endif
+#if defined(HAVE_CUDA) && defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
+    #include "opencv2/gpu/stream_accessor.hpp"
+    #include "opencv2/gpu/device/common.hpp"
+    static inline void throw_nogpu() { CV_Error(CV_StsNotImplemented, "The called functionality is disabled for current build or platform"); }
+#else
+    static inline void throw_nogpu() { CV_Error(CV_GpuNotSupported, "The library is compiled without GPU support"); }
 #endif
 
 #ifdef HAVE_OPENCV_OCL

modules/nonfree/src/surf_gpu.cpp

@@ -42,12 +42,10 @@
 
 #include "precomp.hpp"
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
-
 using namespace cv;
 using namespace cv::gpu;
 
-#if !defined (HAVE_CUDA)
+#if !defined (HAVE_CUDA) || !defined(HAVE_OPENCV_GPU) || defined(DYNAMIC_CUDA_SUPPORT)
 
 cv::gpu::SURF_GPU::SURF_GPU() { throw_nogpu(); }
 cv::gpu::SURF_GPU::SURF_GPU(double, int, int, bool, float, bool) { throw_nogpu(); }
@@ -421,5 +419,3 @@ void cv::gpu::SURF_GPU::releaseMemory()
 }
 
 #endif // !defined (HAVE_CUDA)
-
-#endif // defined(HAVE_OPENCV_GPU) && !defined(ANDROID)

modules/nonfree/test/test_gpu.cpp

@@ -42,7 +42,9 @@
 
 #include "test_precomp.hpp"
 
-#if defined(HAVE_OPENCV_GPU) && defined(HAVE_CUDA)
+#include "cvconfig.h"
+
+#if defined(HAVE_OPENCV_GPU) && defined(HAVE_CUDA) && !defined(DYNAMIC_CUDA_SUPPORT)
 
 using namespace cvtest;
 

modules/python/CMakeLists.txt

@@ -108,17 +108,17 @@ endif()
 if(WIN32)
   set(PYTHON_INSTALL_ARCHIVE "")
 else()
-  set(PYTHON_INSTALL_ARCHIVE ARCHIVE DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT main)
+  set(PYTHON_INSTALL_ARCHIVE ARCHIVE DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT python)
 endif()
 
 if(NOT INSTALL_CREATE_DISTRIB)
   install(TARGETS ${the_module}
           ${PYTHON_INSTALL_CONFIGURATIONS}
-          RUNTIME DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT main
-          LIBRARY DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT main
+          RUNTIME DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT python
+          LIBRARY DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT python
           ${PYTHON_INSTALL_ARCHIVE}
           )
-  install(FILES src2/cv.py ${PYTHON_INSTALL_CONFIGURATIONS} DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT main)
+  install(FILES src2/cv.py ${PYTHON_INSTALL_CONFIGURATIONS} DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT python)
 else()
   if(DEFINED PYTHON_VERSION_MAJOR)
     set(__ver "${PYTHON_VERSION_MAJOR}.${PYTHON_VERSION_MINOR}")
@@ -127,7 +127,7 @@ else()
   endif()
   install(TARGETS ${the_module}
           CONFIGURATIONS Release
-          RUNTIME DESTINATION python/${__ver}/${OpenCV_ARCH} COMPONENT main
-          LIBRARY DESTINATION python/${__ver}/${OpenCV_ARCH} COMPONENT main
+          RUNTIME DESTINATION python/${__ver}/${OpenCV_ARCH} COMPONENT python
+          LIBRARY DESTINATION python/${__ver}/${OpenCV_ARCH} COMPONENT python
           )
 endif()

modules/stitching/CMakeLists.txt

@@ -1,6 +1,7 @@
 set(the_description "Images stitching")
-if (ENABLE_DYNAMIC_CUDA)
+if(ENABLE_DYNAMIC_CUDA)
+  add_definitions(-DDYNAMIC_CUDA_SUPPORT)
   ocv_define_module(stitching opencv_imgproc opencv_features2d opencv_calib3d opencv_objdetect OPTIONAL opencv_nonfree)
 else()
   ocv_define_module(stitching opencv_imgproc opencv_features2d opencv_calib3d opencv_objdetect OPTIONAL opencv_gpu opencv_nonfree)
-endif()
+endif()

modules/stitching/include/opencv2/stitching/detail/matchers.hpp

@@ -44,11 +44,12 @@
 #define __OPENCV_STITCHING_MATCHERS_HPP__
 
 #include "opencv2/core/core.hpp"
+#include "opencv2/core/gpumat.hpp"
 #include "opencv2/features2d/features2d.hpp"
 
 #include "opencv2/opencv_modules.hpp"
 
-#if defined(HAVE_OPENCV_NONFREE) && defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_NONFREE)
     #include "opencv2/nonfree/gpu.hpp"
 #endif
 
@@ -104,7 +105,7 @@ private:
 };
 
 
-#if defined(HAVE_OPENCV_NONFREE) && defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_NONFREE)
 class CV_EXPORTS SurfFeaturesFinderGpu : public FeaturesFinder
 {
 public:

modules/stitching/include/opencv2/stitching/detail/seam_finders.hpp

@@ -45,7 +45,7 @@
 
 #include <set>
 #include "opencv2/core/core.hpp"
-#include "opencv2/opencv_modules.hpp"
+#include "opencv2/core/gpumat.hpp"
 
 namespace cv {
 namespace detail {
@@ -227,7 +227,6 @@ private:
 };
 
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
 class CV_EXPORTS GraphCutSeamFinderGpu : public GraphCutSeamFinderBase, public PairwiseSeamFinder
 {
 public:
@@ -251,7 +250,6 @@ private:
     float terminal_cost_;
     float bad_region_penalty_;
 };
-#endif
 
 } // namespace detail
 } // namespace cv

modules/stitching/include/opencv2/stitching/detail/warpers.hpp

@@ -44,11 +44,8 @@
 #define __OPENCV_STITCHING_WARPERS_HPP__
 
 #include "opencv2/core/core.hpp"
+#include "opencv2/core/gpumat.hpp"
 #include "opencv2/imgproc/imgproc.hpp"
-#include "opencv2/opencv_modules.hpp"
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
-# include "opencv2/gpu/gpu.hpp"
-#endif
 
 namespace cv {
 namespace detail {
@@ -331,7 +328,6 @@ public:
 };
 
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
 class CV_EXPORTS PlaneWarperGpu : public PlaneWarper
 {
 public:
@@ -448,7 +444,6 @@ public:
 private:
     gpu::GpuMat d_xmap_, d_ymap_, d_src_, d_dst_;
 };
-#endif
 
 
 struct SphericalPortraitProjector : ProjectorBase

modules/stitching/include/opencv2/stitching/warpers.hpp

@@ -145,7 +145,6 @@ public:
 
 
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
 class PlaneWarperGpu: public WarperCreator
 {
 public:
@@ -165,7 +164,6 @@ class SphericalWarperGpu: public WarperCreator
 public:
     Ptr<detail::RotationWarper> create(float scale) const { return new detail::SphericalWarperGpu(scale); }
 };
-#endif
 
 } // namespace cv
 

modules/stitching/src/blenders.cpp

@@ -189,7 +189,7 @@ Rect FeatherBlender::createWeightMaps(const vector<Mat> &masks, const vector<Poi
 MultiBandBlender::MultiBandBlender(int try_gpu, int num_bands, int weight_type)
 {
     setNumBands(num_bands);
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
     can_use_gpu_ = try_gpu && gpu::getCudaEnabledDeviceCount();
 #else
     (void)try_gpu;
@@ -491,7 +491,7 @@ void createLaplacePyr(const Mat &img, int num_levels, vector<Mat> &pyr)
 
 void createLaplacePyrGpu(const Mat &img, int num_levels, vector<Mat> &pyr)
 {
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
     pyr.resize(num_levels + 1);
 
     vector<gpu::GpuMat> gpu_pyr(num_levels + 1);
@@ -512,6 +512,7 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, vector<Mat> &pyr)
     (void)img;
     (void)num_levels;
     (void)pyr;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
 #endif
 }
 
@@ -531,7 +532,7 @@ void restoreImageFromLaplacePyr(vector<Mat> &pyr)
 
 void restoreImageFromLaplacePyrGpu(vector<Mat> &pyr)
 {
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
     if (pyr.empty())
         return;
 
@@ -549,6 +550,7 @@ void restoreImageFromLaplacePyrGpu(vector<Mat> &pyr)
     gpu_pyr[0].download(pyr[0]);
 #else
     (void)pyr;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
 #endif
 }
 

modules/stitching/src/matchers.cpp

@@ -45,10 +45,7 @@
 using namespace std;
 using namespace cv;
 using namespace cv::detail;
-
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
 using namespace cv::gpu;
-#endif
 
 #ifdef HAVE_OPENCV_NONFREE
 #include "opencv2/nonfree/nonfree.hpp"
@@ -129,7 +126,7 @@ private:
     float match_conf_;
 };
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
 class GpuMatcher : public FeaturesMatcher
 {
 public:
@@ -204,7 +201,7 @@ void CpuMatcher::match(const ImageFeatures &features1, const ImageFeatures &feat
     LOG("1->2 & 2->1 matches: " << matches_info.matches.size() << endl);
 }
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
 void GpuMatcher::match(const ImageFeatures &features1, const ImageFeatures &features2, MatchesInfo& matches_info)
 {
     matches_info.matches.clear();
@@ -432,7 +429,7 @@ void OrbFeaturesFinder::find(const Mat &image, ImageFeatures &features)
     }
 }
 
-#if defined(HAVE_OPENCV_NONFREE) && defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_NONFREE) && defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
 SurfFeaturesFinderGpu::SurfFeaturesFinderGpu(double hess_thresh, int num_octaves, int num_layers,
                                              int num_octaves_descr, int num_layers_descr)
 {
@@ -478,6 +475,29 @@ void SurfFeaturesFinderGpu::collectGarbage()
     keypoints_.release();
     descriptors_.release();
 }
+#elif defined(HAVE_OPENCV_NONFREE)
+SurfFeaturesFinderGpu::SurfFeaturesFinderGpu(double hess_thresh, int num_octaves, int num_layers,
+                                             int num_octaves_descr, int num_layers_descr)
+{
+    (void)hess_thresh;
+    (void)num_octaves;
+    (void)num_layers;
+    (void)num_octaves_descr;
+    (void)num_layers_descr;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+}
+
+
+void SurfFeaturesFinderGpu::find(const Mat &image, ImageFeatures &features)
+{
+    (void)image;
+    (void)features;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+}
+
+void SurfFeaturesFinderGpu::collectGarbage()
+{
+}
 #endif
 
 
@@ -533,7 +553,7 @@ void FeaturesMatcher::operator ()(const vector<ImageFeatures> &features, vector<
 
 BestOf2NearestMatcher::BestOf2NearestMatcher(bool try_use_gpu, float match_conf, int num_matches_thresh1, int num_matches_thresh2)
 {
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
     if (try_use_gpu && getCudaEnabledDeviceCount() > 0)
         impl_ = new GpuMatcher(match_conf);
     else

modules/stitching/src/precomp.hpp

@@ -68,7 +68,7 @@
 #include "opencv2/imgproc/imgproc.hpp"
 #include "opencv2/features2d/features2d.hpp"
 #include "opencv2/calib3d/calib3d.hpp"
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
     #include "opencv2/gpu/gpu.hpp"
 
     #ifdef HAVE_OPENCV_NONFREE

modules/stitching/src/seam_finders.cpp

@@ -1318,7 +1318,7 @@ void GraphCutSeamFinder::find(const vector<Mat> &src, const vector<Point> &corne
 }
 
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
 void GraphCutSeamFinderGpu::find(const vector<Mat> &src, const vector<Point> &corners,
                                  vector<Mat> &masks)
 {
@@ -1642,6 +1642,62 @@ void GraphCutSeamFinderGpu::setGraphWeightsColorGrad(
         }
     }
 }
+#else
+void GraphCutSeamFinderGpu::find(const vector<Mat> &src, const vector<Point> &corners,
+                                 vector<Mat> &masks)
+{
+    (void)src;
+    (void)corners;
+    (void)masks;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+}
+
+
+void GraphCutSeamFinderGpu::findInPair(size_t first, size_t second, Rect roi)
+{
+    (void)first;
+    (void)second;
+    (void)roi;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+}
+
+
+void GraphCutSeamFinderGpu::setGraphWeightsColor(const Mat &img1, const Mat &img2, const Mat &mask1, const Mat &mask2,
+                                                 Mat &terminals, Mat &leftT, Mat &rightT, Mat &top, Mat &bottom)
+{
+    (void)img1;
+    (void)img2;
+    (void)mask1;
+    (void)mask2;
+    (void)terminals;
+    (void)leftT;
+    (void)rightT;
+    (void)top;
+    (void)bottom;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+}
+
+
+void GraphCutSeamFinderGpu::setGraphWeightsColorGrad(
+        const Mat &img1, const Mat &img2, const Mat &dx1, const Mat &dx2,
+        const Mat &dy1, const Mat &dy2, const Mat &mask1, const Mat &mask2,
+        Mat &terminals, Mat &leftT, Mat &rightT, Mat &top, Mat &bottom)
+{
+    (void)img1;
+    (void)img2;
+    (void)dx1;
+    (void)dx2;
+    (void)dy1;
+    (void)dy2;
+    (void)mask1;
+    (void)mask2;
+    (void)terminals;
+    (void)leftT;
+    (void)rightT;
+    (void)top;
+    (void)bottom;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+}
 #endif
 
 } // namespace detail

modules/stitching/src/stitcher.cpp

@@ -58,7 +58,7 @@ Stitcher Stitcher::createDefault(bool try_use_gpu)
     stitcher.setFeaturesMatcher(new detail::BestOf2NearestMatcher(try_use_gpu));
     stitcher.setBundleAdjuster(new detail::BundleAdjusterRay());
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
     if (try_use_gpu && gpu::getCudaEnabledDeviceCount() > 0)
     {
 #if defined(HAVE_OPENCV_NONFREE)

modules/stitching/src/warpers.cpp

@@ -212,7 +212,7 @@ void SphericalWarper::detectResultRoi(Size src_size, Point &dst_tl, Point &dst_b
 }
 
 
-#if defined(HAVE_OPENCV_GPU) && !defined(ANDROID)
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
 Rect PlaneWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, gpu::GpuMat &xmap, gpu::GpuMat &ymap)
 {
     return buildMaps(src_size, K, R, Mat::zeros(3, 1, CV_32F), xmap, ymap);
@@ -294,6 +294,96 @@ Point CylindricalWarperGpu::warp(const gpu::GpuMat &src, const Mat &K, const Mat
     gpu::remap(src, dst, d_xmap_, d_ymap_, interp_mode, border_mode);
     return dst_roi.tl();
 }
+#else
+Rect PlaneWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, gpu::GpuMat &xmap, gpu::GpuMat &ymap)
+{
+    return buildMaps(src_size, K, R, Mat::zeros(3, 1, CV_32F), xmap, ymap);
+}
+
+Rect PlaneWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, gpu::GpuMat &xmap, gpu::GpuMat &ymap)
+{
+    (void)src_size;
+    (void)K;
+    (void)R;
+    (void)T;
+    (void)xmap;
+    (void)ymap;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+    return Rect();
+}
+
+Point PlaneWarperGpu::warp(const gpu::GpuMat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode,
+                           gpu::GpuMat &dst)
+{
+    return warp(src, K, R, Mat::zeros(3, 1, CV_32F), interp_mode, border_mode, dst);
+}
+
+
+Point PlaneWarperGpu::warp(const gpu::GpuMat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode,
+                           gpu::GpuMat &dst)
+{
+    (void)src;
+    (void)K;
+    (void)R;
+    (void)T;
+    (void)interp_mode;
+    (void)border_mode;
+    (void)dst;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+    return Point();
+}
+
+
+Rect SphericalWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, gpu::GpuMat &xmap, gpu::GpuMat &ymap)
+{
+    (void)src_size;
+    (void)K;
+    (void)R;
+    (void)xmap;
+    (void)ymap;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+    return Rect();
+}
+
+
+Point SphericalWarperGpu::warp(const gpu::GpuMat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode,
+                               gpu::GpuMat &dst)
+{
+    (void)src;
+    (void)K;
+    (void)R;
+    (void)interp_mode;
+    (void)border_mode;
+    (void)dst;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+    return Point();
+}
+
+
+Rect CylindricalWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, gpu::GpuMat &xmap, gpu::GpuMat &ymap)
+{
+    (void)src_size;
+    (void)K;
+    (void)R;
+    (void)xmap;
+    (void)ymap;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+    return Rect();
+}
+
+
+Point CylindricalWarperGpu::warp(const gpu::GpuMat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode,
+                                 gpu::GpuMat &dst)
+{
+    (void)src;
+    (void)K;
+    (void)R;
+    (void)interp_mode;
+    (void)border_mode;
+    (void)dst;
+    CV_Error(CV_StsNotImplemented, "CUDA optimization is unavailable");
+    return Point();
+}
 #endif
 
 void SphericalPortraitWarper::detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br)

platforms/android/android.toolchain.cmake

@@ -1,5 +1,5 @@
 # Copyright (c) 2010-2011, Ethan Rublee
-# Copyright (c) 2011-2013, Andrey Kamaev
+# Copyright (c) 2011-2014, Andrey Kamaev
 # All rights reserved.
 #
 # Redistribution and use in source and binary forms, with or without
@@ -12,9 +12,9 @@
 #     this list of conditions and the following disclaimer in the documentation
 #     and/or other materials provided with the distribution.
 #
-# 3.  The name of the copyright holders may be used to endorse or promote
-#     products derived from this software without specific prior written
-#     permission.
+# 3.  Neither the name of the copyright holder nor the names of its
+#     contributors may be used to endorse or promote products derived from this
+#     software without specific prior written permission.
 #
 # THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
 # AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
@@ -29,12 +29,12 @@
 # POSSIBILITY OF SUCH DAMAGE.
 
 # ------------------------------------------------------------------------------
-#  Android CMake toolchain file, for use with the Android NDK r5-r8
+#  Android CMake toolchain file, for use with the Android NDK r5-r9
 #  Requires cmake 2.6.3 or newer (2.8.5 or newer is recommended).
 #  See home page: https://github.com/taka-no-me/android-cmake
 #
 #  The file is mantained by the OpenCV project. The latest version can be get at
-#  https://github.com/Itseez/opencv/tree/master/platforms/android/android.toolchain.cmake
+#  http://code.opencv.org/projects/opencv/repository/revisions/master/changes/android/android.toolchain.cmake
 #
 #  Usage Linux:
 #   $ export ANDROID_NDK=/absolute/path/to/the/android-ndk
@@ -87,8 +87,7 @@
 #        "armeabi-v6 with VFP" - tuned for ARMv6 processors having VFP.
 #        "x86" - matches to the NDK ABI with the same name.
 #            See ${ANDROID_NDK}/docs/CPU-ARCH-ABIS.html for the documentation.
-#        "mips" - matches to the NDK ABI with the same name
-#            (It is not tested on real devices by the authos of this toolchain)
+#        "mips" - matches to the NDK ABI with the same name.
 #            See ${ANDROID_NDK}/docs/CPU-ARCH-ABIS.html for the documentation.
 #
 #    ANDROID_NATIVE_API_LEVEL=android-8 - level of Android API compile for.
@@ -292,6 +291,16 @@
 #   - April 2013
 #     [+] support non-release NDK layouts (from Linaro git and Android git)
 #     [~] automatically detect if explicit link to crtbegin_*.o is needed
+#   - June 2013
+#     [~] fixed stl include path for standalone toolchain made by NDK >= r8c
+#   - July 2013
+#     [+] updated for NDK r9
+#   - November 2013
+#     [+] updated for NDK r9b
+#   - December 2013
+#     [+] updated for NDK r9c
+#   - January 2014
+#     [~] fix copying of shared STL
 # ------------------------------------------------------------------------------
 
 cmake_minimum_required( VERSION 2.6.3 )
@@ -318,7 +327,7 @@ set( CMAKE_SYSTEM_VERSION 1 )
 # rpath makes low sence for Android
 set( CMAKE_SKIP_RPATH TRUE CACHE BOOL "If set, runtime paths are not added when using shared libraries." )
 
-set( ANDROID_SUPPORTED_NDK_VERSIONS ${ANDROID_EXTRA_NDK_VERSIONS} -r9b -r9 -r8e -r8d -r8c -r8b -r8 -r7c -r7b -r7 -r6b -r6 -r5c -r5b -r5 "" )
+set( ANDROID_SUPPORTED_NDK_VERSIONS ${ANDROID_EXTRA_NDK_VERSIONS} -r9c -r9b -r9 -r8e -r8d -r8c -r8b -r8 -r7c -r7b -r7 -r6b -r6 -r5c -r5b -r5 "" )
 if(NOT DEFINED ANDROID_NDK_SEARCH_PATHS)
  if( CMAKE_HOST_WIN32 )
   file( TO_CMAKE_PATH "$ENV{PROGRAMFILES}" ANDROID_NDK_SEARCH_PATHS )
@@ -464,7 +473,7 @@ endif()
 
 
 # detect current host platform
-if( NOT DEFINED ANDROID_NDK_HOST_X64 AND CMAKE_HOST_SYSTEM_PROCESSOR MATCHES "amd64|x86_64|AMD64")
+if( NOT DEFINED ANDROID_NDK_HOST_X64 AND (CMAKE_HOST_SYSTEM_PROCESSOR MATCHES "amd64|x86_64|AMD64" OR CMAKE_HOST_APPLE) )
  set( ANDROID_NDK_HOST_X64 1 CACHE BOOL "Try to use 64-bit compiler toolchain" )
  mark_as_advanced( ANDROID_NDK_HOST_X64 )
 endif()
@@ -484,9 +493,7 @@ else()
  message( FATAL_ERROR "Cross-compilation on your platform is not supported by this cmake toolchain" )
 endif()
 
-# CMAKE_HOST_SYSTEM_PROCESSOR on MacOS X always says i386 on Intel platform
-# So we do not trust ANDROID_NDK_HOST_X64 on Apple hosts
-if( NOT ANDROID_NDK_HOST_X64 AND NOT CMAKE_HOST_APPLE)
+if( NOT ANDROID_NDK_HOST_X64 )
  set( ANDROID_NDK_HOST_SYSTEM_NAME ${ANDROID_NDK_HOST_SYSTEM_NAME2} )
 endif()
 
@@ -634,30 +641,27 @@ endif()
 
 macro( __GLOB_NDK_TOOLCHAINS __availableToolchainsVar __availableToolchainsLst __toolchain_subpath )
  foreach( __toolchain ${${__availableToolchainsLst}} )
-  # Skip renderscript folder. It's not C++ toolchain
-  if (NOT ${__toolchain} STREQUAL "renderscript")
-   if( "${__toolchain}" MATCHES "-clang3[.][0-9]$" AND NOT EXISTS "${ANDROID_NDK_TOOLCHAINS_PATH}/${__toolchain}${__toolchain_subpath}" )
-     string( REGEX REPLACE "-clang3[.][0-9]$" "-4.6" __gcc_toolchain "${__toolchain}" )
-   else()
-     set( __gcc_toolchain "${__toolchain}" )
-   endif()
-   __DETECT_TOOLCHAIN_MACHINE_NAME( __machine "${ANDROID_NDK_TOOLCHAINS_PATH}/${__gcc_toolchain}${__toolchain_subpath}" )
-   if( __machine )
-    string( REGEX MATCH "[0-9]+[.][0-9]+([.][0-9x]+)?$" __version "${__gcc_toolchain}" )
-    if( __machine MATCHES i686 )
-     set( __arch "x86" )
-    elseif( __machine MATCHES arm )
-     set( __arch "arm" )
-    elseif( __machine MATCHES mipsel )
-     set( __arch "mipsel" )
-    endif()
-    list( APPEND __availableToolchainMachines "${__machine}" )
-    list( APPEND __availableToolchainArchs "${__arch}" )
-    list( APPEND __availableToolchainCompilerVersions "${__version}" )
-    list( APPEND ${__availableToolchainsVar} "${__toolchain}" )
-   endif()
-   unset( __gcc_toolchain )
+  if( "${__toolchain}" MATCHES "-clang3[.][0-9]$" AND NOT EXISTS "${ANDROID_NDK_TOOLCHAINS_PATH}/${__toolchain}${__toolchain_subpath}" )
+   string( REGEX REPLACE "-clang3[.][0-9]$" "-4.6" __gcc_toolchain "${__toolchain}" )
+  else()
+   set( __gcc_toolchain "${__toolchain}" )
   endif()
+  __DETECT_TOOLCHAIN_MACHINE_NAME( __machine "${ANDROID_NDK_TOOLCHAINS_PATH}/${__gcc_toolchain}${__toolchain_subpath}" )
+  if( __machine )
+   string( REGEX MATCH "[0-9]+[.][0-9]+([.][0-9x]+)?$" __version "${__gcc_toolchain}" )
+   if( __machine MATCHES i686 )
+    set( __arch "x86" )
+   elseif( __machine MATCHES arm )
+    set( __arch "arm" )
+   elseif( __machine MATCHES mipsel )
+    set( __arch "mipsel" )
+   endif()
+   list( APPEND __availableToolchainMachines "${__machine}" )
+   list( APPEND __availableToolchainArchs "${__arch}" )
+   list( APPEND __availableToolchainCompilerVersions "${__version}" )
+   list( APPEND ${__availableToolchainsVar} "${__toolchain}" )
+  endif()
+  unset( __gcc_toolchain )
  endforeach()
 endmacro()
 
@@ -687,6 +691,7 @@ if( BUILD_WITH_ANDROID_NDK )
   endif()
   __LIST_FILTER( __availableToolchainsLst "^[.]" )
   __LIST_FILTER( __availableToolchainsLst "llvm" )
+  __LIST_FILTER( __availableToolchainsLst "renderscript" )
   __GLOB_NDK_TOOLCHAINS( __availableToolchains __availableToolchainsLst "${ANDROID_NDK_TOOLCHAINS_SUBPATH}" )
   if( NOT __availableToolchains AND NOT ANDROID_NDK_TOOLCHAINS_SUBPATH STREQUAL ANDROID_NDK_TOOLCHAINS_SUBPATH2 )
    __GLOB_NDK_TOOLCHAINS( __availableToolchains __availableToolchainsLst "${ANDROID_NDK_TOOLCHAINS_SUBPATH2}" )
@@ -975,7 +980,11 @@ if( BUILD_WITH_STANDALONE_TOOLCHAIN )
  set( ANDROID_SYSROOT "${ANDROID_STANDALONE_TOOLCHAIN}/sysroot" )
 
  if( NOT ANDROID_STL STREQUAL "none" )
-  set( ANDROID_STL_INCLUDE_DIRS "${ANDROID_STANDALONE_TOOLCHAIN}/${ANDROID_TOOLCHAIN_MACHINE_NAME}/include/c++/${ANDROID_COMPILER_VERSION}" )
+  set( ANDROID_STL_INCLUDE_DIRS "${ANDROID_STANDALONE_TOOLCHAIN}/include/c++/${ANDROID_COMPILER_VERSION}" )
+  if( NOT EXISTS "${ANDROID_STL_INCLUDE_DIRS}" )
+   # old location ( pre r8c )
+   set( ANDROID_STL_INCLUDE_DIRS "${ANDROID_STANDALONE_TOOLCHAIN}/${ANDROID_TOOLCHAIN_MACHINE_NAME}/include/c++/${ANDROID_COMPILER_VERSION}" )
+  endif()
   if( ARMEABI_V7A AND EXISTS "${ANDROID_STL_INCLUDE_DIRS}/${ANDROID_TOOLCHAIN_MACHINE_NAME}/${CMAKE_SYSTEM_PROCESSOR}/bits" )
    list( APPEND ANDROID_STL_INCLUDE_DIRS "${ANDROID_STL_INCLUDE_DIRS}/${ANDROID_TOOLCHAIN_MACHINE_NAME}/${CMAKE_SYSTEM_PROCESSOR}" )
   elseif( ARMEABI AND NOT ANDROID_FORCE_ARM_BUILD AND EXISTS "${ANDROID_STL_INCLUDE_DIRS}/${ANDROID_TOOLCHAIN_MACHINE_NAME}/thumb/bits" )
@@ -1130,15 +1139,7 @@ endif()
 # case of shared STL linkage
 if( ANDROID_STL MATCHES "shared" AND DEFINED __libstl )
  string( REPLACE "_static.a" "_shared.so" __libstl "${__libstl}" )
- if( NOT _CMAKE_IN_TRY_COMPILE AND __libstl MATCHES "[.]so$" )
-  get_filename_component( __libstlname "${__libstl}" NAME )
-  execute_process( COMMAND "${CMAKE_COMMAND}" -E copy_if_different "${__libstl}" "${LIBRARY_OUTPUT_PATH}/${__libstlname}" RESULT_VARIABLE __fileCopyProcess )
-  if( NOT __fileCopyProcess EQUAL 0 OR NOT EXISTS "${LIBRARY_OUTPUT_PATH}/${__libstlname}")
-   message( SEND_ERROR "Failed copying of ${__libstl} to the ${LIBRARY_OUTPUT_PATH}/${__libstlname}" )
-  endif()
-  unset( __fileCopyProcess )
-  unset( __libstlname )
- endif()
+ # TODO: check if .so file exists before the renaming
 endif()
 
 
@@ -1503,7 +1504,8 @@ endif()
 
 # global includes and link directories
 include_directories( SYSTEM "${ANDROID_SYSROOT}/usr/include" ${ANDROID_STL_INCLUDE_DIRS} )
-link_directories( "${CMAKE_INSTALL_PREFIX}/libs/${ANDROID_NDK_ABI_NAME}" )
+get_filename_component(__android_install_path "${CMAKE_INSTALL_PREFIX}/libs/${ANDROID_NDK_ABI_NAME}" ABSOLUTE) # avoid CMP0015 policy warning
+link_directories( "${__android_install_path}" )
 
 # detect if need link crtbegin_so.o explicitly
 if( NOT DEFINED ANDROID_EXPLICIT_CRT_LINK )
@@ -1555,6 +1557,18 @@ if(NOT _CMAKE_IN_TRY_COMPILE)
  set( LIBRARY_OUTPUT_PATH "${LIBRARY_OUTPUT_PATH_ROOT}/libs/${ANDROID_NDK_ABI_NAME}" CACHE PATH "path for android libs" )
 endif()
 
+# copy shaed stl library to build directory
+if( NOT _CMAKE_IN_TRY_COMPILE AND __libstl MATCHES "[.]so$" )
+ get_filename_component( __libstlname "${__libstl}" NAME )
+ execute_process( COMMAND "${CMAKE_COMMAND}" -E copy_if_different "${__libstl}" "${LIBRARY_OUTPUT_PATH}/${__libstlname}" RESULT_VARIABLE __fileCopyProcess )
+ if( NOT __fileCopyProcess EQUAL 0 OR NOT EXISTS "${LIBRARY_OUTPUT_PATH}/${__libstlname}")
+  message( SEND_ERROR "Failed copying of ${__libstl} to the ${LIBRARY_OUTPUT_PATH}/${__libstlname}" )
+ endif()
+ unset( __fileCopyProcess )
+ unset( __libstlname )
+endif()
+
+
 # set these global flags for cmake client scripts to change behavior
 set( ANDROID True )
 set( BUILD_ANDROID True )
@@ -1663,6 +1677,19 @@ if( NOT PROJECT_NAME STREQUAL "CMAKE_TRY_COMPILE" )
 endif()
 
 
+# force cmake to produce / instead of \ in build commands for Ninja generator
+if( CMAKE_GENERATOR MATCHES "Ninja" AND CMAKE_HOST_WIN32 )
+ # it is a bad hack after all
+ # CMake generates Ninja makefiles with UNIX paths only if it thinks that we are going to build with MinGW
+ set( CMAKE_COMPILER_IS_MINGW TRUE ) # tell CMake that we are MinGW
+ set( CMAKE_CROSSCOMPILING TRUE )    # stop recursion
+ enable_language( C )
+ enable_language( CXX )
+ # unset( CMAKE_COMPILER_IS_MINGW ) # can't unset because CMake does not convert back-slashes in response files without it
+ unset( MINGW )
+endif()
+
+
 # set some obsolete variables for backward compatibility
 set( ANDROID_SET_OBSOLETE_VARIABLES ON CACHE BOOL "Define obsolete Andrid-specific cmake variables" )
 mark_as_advanced( ANDROID_SET_OBSOLETE_VARIABLES )
@@ -1717,7 +1744,7 @@ endif()
 #   BUILD_WITH_STANDALONE_TOOLCHAIN : TRUE if standalone toolchain is used
 #   ANDROID_NDK_HOST_SYSTEM_NAME : "windows", "linux-x86" or "darwin-x86" depending on host platform
 #   ANDROID_NDK_ABI_NAME : "armeabi", "armeabi-v7a", "x86" or "mips" depending on ANDROID_ABI
-#   ANDROID_NDK_RELEASE : one of r5, r5b, r5c, r6, r6b, r7, r7b, r7c, r8, r8b, r8c, r8d, r8e; set only for NDK
+#   ANDROID_NDK_RELEASE : one of r5, r5b, r5c, r6, r6b, r7, r7b, r7c, r8, r8b, r8c, r8d, r8e, r9, r9b, r9c; set only for NDK
 #   ANDROID_ARCH_NAME : "arm" or "x86" or "mips" depending on ANDROID_ABI
 #   ANDROID_SYSROOT : path to the compiler sysroot
 #   TOOL_OS_SUFFIX : "" or ".exe" depending on host platform

platforms/android/libinfo/CMakeLists.txt

@@ -36,4 +36,4 @@ set_target_properties(${the_module} PROPERTIES
   )
 
 get_filename_component(lib_name "libopencv_info.so" NAME)
-install(FILES "${LIBRARY_OUTPUT_PATH}/${lib_name}" DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main)
+install(FILES "${LIBRARY_OUTPUT_PATH}/${lib_name}" DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs)

platforms/android/package/CMakeLists.txt

@@ -89,6 +89,6 @@ add_custom_command(
          DEPENDS "${OpenCV_BINARY_DIR}/bin/classes.jar.dephelper" "${PACKAGE_DIR}/res/values/strings.xml" "${PACKAGE_DIR}/res/drawable/icon.png" ${camera_wrappers} opencv_java
        )
 
-install(FILES "${APK_NAME}" DESTINATION "apk/" COMPONENT main)
+install(FILES "${APK_NAME}" DESTINATION "apk/" COMPONENT libs)
 add_custom_target(android_package ALL SOURCES "${APK_NAME}" )
 add_dependencies(android_package opencv_java)

platforms/android/service/CMakeLists.txt

@@ -3,4 +3,4 @@ if(BUILD_ANDROID_SERVICE)
   #add_subdirectory(engine_test)
 endif()
 
-install(FILES "readme.txt" DESTINATION "apk/" COMPONENT main)
+install(FILES "readme.txt" DESTINATION "apk/" COMPONENT libs)

platforms/android/service/engine/AndroidManifest.xml

@@ -1,8 +1,8 @@
 <?xml version="1.0" encoding="utf-8"?>
 <manifest xmlns:android="http://schemas.android.com/apk/res/android"
     package="org.opencv.engine"
-    android:versionCode="216@ANDROID_PLATFORM_VERSION_CODE@"
-    android:versionName="2.16" >
+    android:versionCode="217@ANDROID_PLATFORM_VERSION_CODE@"
+    android:versionName="2.17" >
 
     <uses-sdk android:minSdkVersion="@ANDROID_NATIVE_API_LEVEL@" />
     <uses-feature android:name="android.hardware.touchscreen" android:required="false"/>

samples/android/native-activity/src/org/opencv/samples/NativeActivity/CvNativeActivity.java

@@ -21,6 +21,7 @@ public class CvNativeActivity extends Activity {
                     System.loadLibrary("native_activity");
                     Intent intent = new Intent(CvNativeActivity.this, android.app.NativeActivity.class);
                     CvNativeActivity.this.startActivity(intent);
+                    CvNativeActivity.this.finish();
                 } break;
                 default:
                 {
@@ -34,7 +35,7 @@ public class CvNativeActivity extends Activity {
         Log.i(TAG, "Instantiated new " + this.getClass());
     }
 
-   @Override
+    @Override
     public void onResume()
     {
         super.onResume();