Merge branch 'android_manager_version_inc4' into 2.4.8.1-prep

disable gpu CvtColor.*2HSV tests:
possible bug in CPU version(cherry picked from commit 3cb8b352e52f85847d98d2b0ece32ba5bdb5c31b)
2026-07-21 19:33:03 +04:00 · 2014-01-31 22:25:36 +04:00 · 2014-01-31 22:24:20 +04:00 · 2014-01-31 22:23:55 +04:00 · 2014-01-31 22:23:39 +04:00 · 2014-01-31 22:23:22 +04:00
61 changed files with 1814 additions and 1116 deletions
@@ -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()
@@ -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()
@@ -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()
@@ -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()
@@ -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)
@@ -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
@@ -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()
@@ -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
 # ===================================================
@@ -275,6 +275,9 @@ endif()
 set(OPENCV_SAMPLES_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}samples")

 set(OPENCV_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}bin")
+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 +286,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 +297,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 +564,19 @@ 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 AND NOT ANDROID)
+  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.sh.in"
+                 "${CMAKE_BINARY_DIR}/unix-install/opencv_run_all_tests.sh" @ONLY)
+  install(FILES "${CMAKE_BINARY_DIR}/unix-install/opencv_run_all_tests.sh"
+          PERMISSIONS OWNER_READ OWNER_WRITE GROUP_READ WORLD_READ OWNER_EXECUTE GROUP_EXECUTE WORLD_EXECUTE
+          DESTINATION ${OPENCV_TEST_INSTALL_PATH} COMPONENT tests)
+endif()
+
 # ----------------------------------------------------------------------------
 # Summary:
 # ----------------------------------------------------------------------------
@@ -975,3 +994,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)
@@ -0,0 +1,33 @@
+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
+                       (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:
+
+  * Redistributions of source code must retain the above copyright notice,
+    this list of conditions and the following disclaimer.
+
+  * Redistributions 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.
+
+  * 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 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
+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.
@@ -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)
@@ -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()
@@ -344,20 +344,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 +365,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 dev)
      #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()
@@ -116,5 +116,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)
@@ -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()
@@ -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)
@@ -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()
@@ -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()

@@ -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")
@@ -0,0 +1,25 @@
+#!/bin/sh
+
+OPENCV_TEST_PATH=@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
@@ -0,0 +1,2 @@
+# Environment setup for OpenCV testing
+export OPENCV_TEST_DATA_PATH=@CMAKE_INSTALL_PREFIX@/share/OpenCV/testdata
@@ -2,9 +2,17 @@ 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)
+    install(DIRECTORY ${OPENCV_TEST_DATA_PATH} DESTINATION share/OpenCV/testdata COMPONENT tests)
+  endif()
+endif()
@@ -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()
@@ -48,10 +48,10 @@ The structure of package contents looks as follows:

 ::

-    OpenCV-2.4.8-android-sdk
+    OpenCV-2.4.8.1-android-sdk
    |_ apk
-    |   |_ OpenCV_2.4.8_binary_pack_armv7a.apk
-    |   |_ OpenCV_2.4.8_Manager_2.16_XXX.apk
+    |   |_ OpenCV_2.4.8.1_binary_pack_armv7a.apk
+    |   |_ OpenCV_2.4.8.1_Manager_2.17_XXX.apk
    |
    |_ doc
    |_ samples
@@ -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.1-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.1-android-sdk.zip`
+.. _opencv_android_bin_pack_url: http://sourceforge.net/projects/opencvlibrary/files/opencv-android/2.4.8.1/OpenCV-2.4.8.1-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.1_Manager_2.17_armv7a-neon.apk

  .. note:: ``armeabi``, ``armv7a-neon``, ``arm7a-neon-android8``, ``mips`` and ``x86`` stand for
            platform targets:
@@ -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.1-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.1``.

   .. 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.1-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.1``;

   .. 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.1-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.1-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.1-android-sdk\sdk\native\jni\OpenCV.mk

   Should be inserted into the :file:`jni/Android.mk` file **after** this line:

@@ -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)
@@ -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()
@@ -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)
@@ -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 1

 #define CVAUX_STR_EXP(__A)  #__A
 #define CVAUX_STR(__A)      CVAUX_STR_EXP(__A)
@@ -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);
        }
    }

@@ -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})
@@ -672,7 +672,7 @@ PERF_TEST_P(Sz, ImgProc_ColumnSum,

 DEF_PARAM_TEST(Image_AppertureSz_L2gradient, string, int, bool);

-PERF_TEST_P(Image_AppertureSz_L2gradient, ImgProc_Canny,
+PERF_TEST_P(Image_AppertureSz_L2gradient, DISABLED_ImgProc_Canny,
            Combine(Values("perf/800x600.png", "perf/1280x1024.png", "perf/1680x1050.png"),
                    Values(3, 5),
                    Bool()))
@@ -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

@@ -1773,7 +1777,7 @@ PERF_TEST_P(Image, ImgProc_HoughLinesP,

 DEF_PARAM_TEST(Sz_Dp_MinDist, cv::Size, float, float);

-PERF_TEST_P(Sz_Dp_MinDist, ImgProc_HoughCircles,
+PERF_TEST_P(Sz_Dp_MinDist, DISABLED_ImgProc_HoughCircles,
            Combine(GPU_TYPICAL_MAT_SIZES,
                    Values(1.0f, 2.0f, 4.0f),
                    Values(1.0f)))
@@ -1825,7 +1829,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))
 {
@@ -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 */
@@ -40,654 +40,23 @@
 //
 //M*/

+#define CUDA_DISABLER
+
 #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 +1075,4 @@ namespace cv { namespace gpu { namespace device
    }
 }}}

-
 #endif /* CUDA_DISABLER */
@@ -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 */
@@ -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 */
@@ -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 */
@@ -40,6 +40,8 @@
 //
 //M*/

+#define CUDA_DISABLER
+
 #include "precomp.hpp"

 using namespace std;
@@ -48,16 +50,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 +69,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
@@ -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) */
@@ -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) */
@@ -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) */
@@ -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] =
@@ -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);

@@ -840,7 +840,7 @@ GPU_TEST_P(CvtColor, YCrCb42RGBA)
    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
 }

-GPU_TEST_P(CvtColor, BGR2HSV)
+GPU_TEST_P(CvtColor, DISABLED_BGR2HSV)
 {
    if (depth == CV_16U)
        return;
@@ -856,7 +856,7 @@ GPU_TEST_P(CvtColor, BGR2HSV)
    EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_32F ? 1e-2 : 1);
 }

-GPU_TEST_P(CvtColor, RGB2HSV)
+GPU_TEST_P(CvtColor, DISABLED_RGB2HSV)
 {
    if (depth == CV_16U)
        return;
@@ -872,7 +872,7 @@ GPU_TEST_P(CvtColor, RGB2HSV)
    EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_32F ? 1e-2 : 1);
 }

-GPU_TEST_P(CvtColor, RGB2HSV4)
+GPU_TEST_P(CvtColor, DISABLED_RGB2HSV4)
 {
    if (depth == CV_16U)
        return;
@@ -896,7 +896,7 @@ GPU_TEST_P(CvtColor, RGB2HSV4)
    EXPECT_MAT_NEAR(dst_gold, h_dst, depth == CV_32F ? 1e-2 : 1);
 }

-GPU_TEST_P(CvtColor, RGBA2HSV4)
+GPU_TEST_P(CvtColor, DISABLED_RGBA2HSV4)
 {
    if (depth == CV_16U)
        return;
@@ -422,7 +422,7 @@ PARAM_TEST_CASE(Subtract_Array, cv::gpu::DeviceInfo, cv::Size, std::pair<MatDept
    }
 };

-GPU_TEST_P(Subtract_Array, Accuracy)
+GPU_TEST_P(Subtract_Array, DISABLED_Accuracy)
 {
    cv::Mat mat1 = randomMat(size, stype);
    cv::Mat mat2 = randomMat(size, stype);
@@ -189,7 +189,7 @@ PARAM_TEST_CASE(GeneralizedHough, cv::gpu::DeviceInfo, UseRoi)
 {
 };

-GPU_TEST_P(GeneralizedHough, POSITION)
+GPU_TEST_P(GeneralizedHough, DISABLED_POSITION)
 {
    const cv::gpu::DeviceInfo devInfo = GET_PARAM(0);
    cv::gpu::setDevice(devInfo.deviceID());
@@ -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

@@ -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()
@@ -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
            {
@@ -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");
 }
@@ -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()

@@ -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()
@@ -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
@@ -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)
@@ -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)
@@ -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)
@@ -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"/>
@@ -39,7 +39,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
        set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG")
      endif()
      install(TARGETS ${the_target}
-              RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/c" COMPONENT main)
+              RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/c" COMPONENT samples)
    endif()
  ENDMACRO()

@@ -54,6 +54,6 @@ endif()
 if (INSTALL_C_EXAMPLES AND NOT WIN32)
  file(GLOB C_SAMPLES *.c *.cpp *.jpg *.png *.data makefile.* build_all.sh *.dsp *.cmd )
  install(FILES ${C_SAMPLES}
-          DESTINATION share/OpenCV/samples/c
-          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ)
+          DESTINATION ${OPENCV_SAMPLES_SRC_INSTALL_PATH}/c
+          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples)
 endif ()
@@ -68,7 +68,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
        set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG")
      endif()
      install(TARGETS ${the_target}
-              RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${sample_subfolder}" COMPONENT main)
+              RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${sample_subfolder}" COMPONENT samples)
    endif()
  ENDMACRO()

@@ -91,6 +91,6 @@ endif()
 if (INSTALL_C_EXAMPLES AND NOT WIN32)
  file(GLOB C_SAMPLES *.c *.cpp *.jpg *.png *.data makefile.* build_all.sh *.dsp *.cmd )
  install(FILES ${C_SAMPLES}
-          DESTINATION share/OpenCV/samples/cpp
-          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ)
+          DESTINATION ${OPENCV_SAMPLES_SRC_INSTALL_PATH}/cpp
+          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples)
 endif()
@@ -65,7 +65,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
      if(MSVC AND NOT BUILD_SHARED_LIBS)
        set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG")
      endif()
-      install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${project}" COMPONENT main)
+      install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${project}" COMPONENT samples)
    endif()
  ENDMACRO()

@@ -83,6 +83,6 @@ endif()
 if (INSTALL_C_EXAMPLES AND NOT WIN32)
  file(GLOB install_list *.c *.cpp *.jpg *.png *.data makefile.* build_all.sh *.dsp *.cmd )
  install(FILES ${install_list}
-          DESTINATION share/OpenCV/samples/${project}
-          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ)
+          DESTINATION ${OPENCV_SAMPLES_SRC_INSTALL_PATH}/gpu
+          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples)
 endif()
@@ -23,12 +23,13 @@ if(ENABLE_SOLUTION_FOLDERS)
 endif()

 if(WIN32)
-  install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/gpu" COMPONENT main)
+  install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/gpu" COMPONENT samples)
 endif()

 if(INSTALL_C_EXAMPLES AND NOT WIN32)
  file(GLOB GPU_FILES performance/*.cpp performance/*.h)
  install(FILES ${GPU_FILES}
-          DESTINATION share/OpenCV/samples/gpu/performance
-          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ)
+          DESTINATION ${OPENCV_SAMPLES_SRC_INSTALL_PATH}/gpu/performance
+          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ
+          COMPONENT samples)
 endif()
@@ -38,7 +38,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
      if(MSVC AND NOT BUILD_SHARED_LIBS)
        set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG")
      endif()
-      install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${project}" COMPONENT main)
+      install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${project}" COMPONENT samples)
    endif()
  ENDMACRO()

@@ -54,6 +54,6 @@ endif()
 if (INSTALL_C_EXAMPLES AND NOT WIN32)
  file(GLOB install_list *.c *.cpp *.jpg *.png *.data makefile.* build_all.sh *.dsp *.cmd )
  install(FILES ${install_list}
-          DESTINATION share/OpenCV/samples/${project}
-          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ)
+          DESTINATION ${OPENCV_SAMPLES_SRC_INSTALL_PATH}/ocl
+          PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples)
 endif()
Author	SHA1	Message	Date
Alexander Smorkalov	215f78eee2	Merge branch 'android_manager_version_inc4' into 2.4.8.1-prep	2014-01-31 22:25:36 +04:00
Vladislav Vinogradov	4425dac7f4	disable gpu CvtColor.*2HSV tests: possible bug in CPU version(cherry picked from commit 3cb8b352e52f85847d98d2b0ece32ba5bdb5c31b)	2014-01-31 22:24:20 +04:00
Vladislav Vinogradov	4f79b9de48	disable gpu Subtract_Array test: possible bug in CPU version(cherry picked from commit 59155c1eefb6f22e362e34a8d73b31841eb3413a)	2014-01-31 22:23:55 +04:00
Vladislav Vinogradov	397ac5e68f	disable gpu Canny and HoughCircles perf tests: it fails because driver terminates CUDA kernels after time out (cherry picked from commit fa5bbb5f8dc1a86cfe004fb258d66e56c90560d6)	2014-01-31 22:23:39 +04:00
Vladislav Vinogradov	43c75c64b5	disable NPP for GpuMat methods and for copyMakeBorder(cherry picked from commit 316d49fc0fb7a609ebb0a65efc207faea6b978a4)	2014-01-31 22:23:22 +04:00
Vladislav Vinogradov	c319625a07	disable some gpu tests if library was built without CUFFT(cherry picked from commit b4b929d27cc25822dd15b5b96b8d335c59d4408c)	2014-01-31 22:23:05 +04:00
Vladislav Vinogradov	a7d0448faa	gpu test output files moved from cwd to temp folder(cherry picked from commit `49731ad530`)	2014-01-31 22:22:49 +04:00
Alexander Smorkalov	ca10e5e8ae	Highgui test output fixes. Useless output to console fixed; Test output files moved from cwd to temp folder. (cherry picked from commit `87935f3560`)	2014-01-31 22:22:33 +04:00
Vladislav Vinogradov	32414afe72	disable performance test for gpu generalized hough(cherry picked from commit `063d8b4211`)	2014-01-31 22:22:14 +04:00
Alexander Smorkalov	ed10f50d25	Reports path fix for opencv_run_all_tests.sh.in script. (cherry picked from commit `3d261e8a01`)	2014-01-31 22:21:57 +04:00
Alexander Smorkalov	48d9be70d5	Android toolchain file sync with original project. Original project: https://github.com/taka-no-me/android-cmake/ Revision: 5db45cfb87fec180b74963d3680dd60d4d8d8c3a (cherry picked from commit `c815043607`)	2014-01-31 22:21:34 +04:00
Alexander Smorkalov	c6b31481b6	OpenCV Manager version++.	2014-01-31 12:15:22 +04:00
Alexander Smorkalov	514b714cc2	opencv_run_all_tests.sh script added to -tests package. (cherry picked from commit `d45350a06a`)	2014-01-29 17:41:59 +04:00
Alexander Smorkalov	7fec87d3f6	Multiple fixes for tests deb package build. Added opencv_testing.sh.in file; opencv_testing.sh installation guarded by OS check. (cherry picked from commit `d9dc5ffa91`)	2014-01-28 17:56:22 +04:00
Alexander Smorkalov	00d555f051	Code review notes fixed. Env setup for testing package implemented using /etc/profile.d; Variable with path for all native samples added; Path for test binaries and test data updated. (cherry picked from commit `39201e68e2`)	2014-01-27 21:37:01 +04:00
Alexander Smorkalov	6cf7d6ef4e	OpenCV C/C++/OCL/CUDA samples install path fixed. Install rools for tests added. (cherry picked from commit `f332cba14b`)	2014-01-27 21:36:43 +04:00
Alexander Smorkalov	3ebdcafdd3	All installed files marked with component names for install customization. (cherry picked from commit `b75cbfde45`) Conflicts: cmake/OpenCVModule.cmake	2014-01-24 17:13:22 +04:00
Alexander Smorkalov	a348f3eeaa	OpenCV version++	2014-01-24 14:49:56 +04:00
Alexander Smorkalov	33f423de04	Improvements in package build. (cherry picked from commit `086792ec06`)	2014-01-24 14:48:51 +04:00
Alexander Smorkalov	d6ba52c3f9	Initial Linux packages build rools for CPack. (cherry picked from commit `7821fe2bde`) Conflicts: cmake/OpenCVModule.cmake	2014-01-24 14:48:19 +04:00
Vladislav Vinogradov	e9638d0997	disable CUDA generalized Hough Transform (cherry picked from commit `33d42b740c`)	2014-01-24 14:46:36 +04:00
Vladislav Vinogradov	5cb0084547	split CUDA Hough sources (cherry picked from commit `d847387694`)	2014-01-24 14:46:16 +04:00