diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r2.2.0.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r2.2.0.so index 5b618a874..aac6634b4 100755 Binary files a/3rdparty/lib/armeabi-v7a/libnative_camera_r2.2.0.so and b/3rdparty/lib/armeabi-v7a/libnative_camera_r2.2.0.so differ diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r2.3.3.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r2.3.3.so index 846fc88bd..d523f69de 100755 Binary files a/3rdparty/lib/armeabi-v7a/libnative_camera_r2.3.3.so and b/3rdparty/lib/armeabi-v7a/libnative_camera_r2.3.3.so differ diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r3.0.1.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r3.0.1.so index 80bf459cc..e386bf4f9 100755 Binary files a/3rdparty/lib/armeabi-v7a/libnative_camera_r3.0.1.so and b/3rdparty/lib/armeabi-v7a/libnative_camera_r3.0.1.so differ diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.0.0.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.0.0.so index e5cc7d296..028ab7d1e 100755 Binary files a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.0.0.so and b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.0.0.so differ diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.0.3.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.0.3.so index d3cf3b124..48cbdd096 100755 Binary files a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.0.3.so and b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.0.3.so differ diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.1.1.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.1.1.so index 6498151ba..7fe50875c 100755 Binary files a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.1.1.so and b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.1.1.so differ diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.2.0.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.2.0.so index 58bef3455..15827d818 100755 Binary files a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.2.0.so and b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.2.0.so differ diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.3.0.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.3.0.so index ce69b52ea..ec1edfb04 100755 Binary files a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.3.0.so and b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.3.0.so differ diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.4.0.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.4.0.so index 3e65fb171..4d777edf8 100755 Binary files a/3rdparty/lib/armeabi-v7a/libnative_camera_r4.4.0.so and b/3rdparty/lib/armeabi-v7a/libnative_camera_r4.4.0.so differ diff --git a/3rdparty/lib/armeabi/libnative_camera_r2.2.0.so b/3rdparty/lib/armeabi/libnative_camera_r2.2.0.so index 68805b589..1707a8850 100755 Binary files a/3rdparty/lib/armeabi/libnative_camera_r2.2.0.so and b/3rdparty/lib/armeabi/libnative_camera_r2.2.0.so differ diff --git a/3rdparty/lib/armeabi/libnative_camera_r2.3.3.so b/3rdparty/lib/armeabi/libnative_camera_r2.3.3.so index 88ac3f7e3..fb4b125fd 100755 Binary files a/3rdparty/lib/armeabi/libnative_camera_r2.3.3.so and b/3rdparty/lib/armeabi/libnative_camera_r2.3.3.so differ diff --git a/3rdparty/lib/armeabi/libnative_camera_r3.0.1.so b/3rdparty/lib/armeabi/libnative_camera_r3.0.1.so index fa41cb250..96b264d0e 100755 Binary files a/3rdparty/lib/armeabi/libnative_camera_r3.0.1.so and b/3rdparty/lib/armeabi/libnative_camera_r3.0.1.so differ diff --git a/3rdparty/lib/armeabi/libnative_camera_r4.0.0.so b/3rdparty/lib/armeabi/libnative_camera_r4.0.0.so index a305c2b00..179eef9a9 100755 Binary files a/3rdparty/lib/armeabi/libnative_camera_r4.0.0.so and b/3rdparty/lib/armeabi/libnative_camera_r4.0.0.so differ diff --git a/3rdparty/lib/armeabi/libnative_camera_r4.0.3.so b/3rdparty/lib/armeabi/libnative_camera_r4.0.3.so index 8c34357cc..165dc463c 100755 Binary files a/3rdparty/lib/armeabi/libnative_camera_r4.0.3.so and b/3rdparty/lib/armeabi/libnative_camera_r4.0.3.so differ diff --git a/3rdparty/lib/armeabi/libnative_camera_r4.1.1.so b/3rdparty/lib/armeabi/libnative_camera_r4.1.1.so index a01ee15e2..a9a5d7da7 100755 Binary files a/3rdparty/lib/armeabi/libnative_camera_r4.1.1.so and b/3rdparty/lib/armeabi/libnative_camera_r4.1.1.so differ diff --git a/3rdparty/lib/armeabi/libnative_camera_r4.2.0.so b/3rdparty/lib/armeabi/libnative_camera_r4.2.0.so index a8ff89465..9037c6860 100755 Binary files a/3rdparty/lib/armeabi/libnative_camera_r4.2.0.so and b/3rdparty/lib/armeabi/libnative_camera_r4.2.0.so differ diff --git a/3rdparty/lib/armeabi/libnative_camera_r4.3.0.so b/3rdparty/lib/armeabi/libnative_camera_r4.3.0.so index aa1cfd844..026f0b48b 100755 Binary files a/3rdparty/lib/armeabi/libnative_camera_r4.3.0.so and b/3rdparty/lib/armeabi/libnative_camera_r4.3.0.so differ diff --git a/3rdparty/lib/armeabi/libnative_camera_r4.4.0.so b/3rdparty/lib/armeabi/libnative_camera_r4.4.0.so index 264f6f217..6aebec923 100755 Binary files a/3rdparty/lib/armeabi/libnative_camera_r4.4.0.so and b/3rdparty/lib/armeabi/libnative_camera_r4.4.0.so differ diff --git a/3rdparty/lib/mips/libnative_camera_r4.0.3.so b/3rdparty/lib/mips/libnative_camera_r4.0.3.so index 14dfaf23b..6dee89780 100755 Binary files a/3rdparty/lib/mips/libnative_camera_r4.0.3.so and b/3rdparty/lib/mips/libnative_camera_r4.0.3.so differ diff --git a/3rdparty/lib/mips/libnative_camera_r4.1.1.so b/3rdparty/lib/mips/libnative_camera_r4.1.1.so index a37474256..71a6354ac 100755 Binary files a/3rdparty/lib/mips/libnative_camera_r4.1.1.so and b/3rdparty/lib/mips/libnative_camera_r4.1.1.so differ diff --git a/3rdparty/lib/mips/libnative_camera_r4.2.0.so b/3rdparty/lib/mips/libnative_camera_r4.2.0.so index 31cbb3a99..21bcffb4a 100755 Binary files a/3rdparty/lib/mips/libnative_camera_r4.2.0.so and b/3rdparty/lib/mips/libnative_camera_r4.2.0.so differ diff --git a/3rdparty/lib/mips/libnative_camera_r4.3.0.so b/3rdparty/lib/mips/libnative_camera_r4.3.0.so index 379fc7003..653c2f1ca 100755 Binary files a/3rdparty/lib/mips/libnative_camera_r4.3.0.so and b/3rdparty/lib/mips/libnative_camera_r4.3.0.so differ diff --git a/3rdparty/lib/mips/libnative_camera_r4.4.0.so b/3rdparty/lib/mips/libnative_camera_r4.4.0.so index 0f6c83713..8d6fdf2bc 100755 Binary files a/3rdparty/lib/mips/libnative_camera_r4.4.0.so and b/3rdparty/lib/mips/libnative_camera_r4.4.0.so differ diff --git a/3rdparty/lib/x86/libnative_camera_r2.3.3.so b/3rdparty/lib/x86/libnative_camera_r2.3.3.so index 5c46b1607..a47b8b2ce 100755 Binary files a/3rdparty/lib/x86/libnative_camera_r2.3.3.so and b/3rdparty/lib/x86/libnative_camera_r2.3.3.so differ diff --git a/3rdparty/lib/x86/libnative_camera_r3.0.1.so b/3rdparty/lib/x86/libnative_camera_r3.0.1.so index 77512e5de..faa13461f 100755 Binary files a/3rdparty/lib/x86/libnative_camera_r3.0.1.so and b/3rdparty/lib/x86/libnative_camera_r3.0.1.so differ diff --git a/3rdparty/lib/x86/libnative_camera_r4.0.3.so b/3rdparty/lib/x86/libnative_camera_r4.0.3.so index b5de08299..2d2fb8eb1 100755 Binary files a/3rdparty/lib/x86/libnative_camera_r4.0.3.so and b/3rdparty/lib/x86/libnative_camera_r4.0.3.so differ diff --git a/3rdparty/lib/x86/libnative_camera_r4.1.1.so b/3rdparty/lib/x86/libnative_camera_r4.1.1.so index 867137410..f40da0d9d 100755 Binary files a/3rdparty/lib/x86/libnative_camera_r4.1.1.so and b/3rdparty/lib/x86/libnative_camera_r4.1.1.so differ diff --git a/3rdparty/lib/x86/libnative_camera_r4.2.0.so b/3rdparty/lib/x86/libnative_camera_r4.2.0.so index 52e9a5792..0d4ac03b5 100755 Binary files a/3rdparty/lib/x86/libnative_camera_r4.2.0.so and b/3rdparty/lib/x86/libnative_camera_r4.2.0.so differ diff --git a/3rdparty/lib/x86/libnative_camera_r4.3.0.so b/3rdparty/lib/x86/libnative_camera_r4.3.0.so index af898ccad..7e1c5803a 100755 Binary files a/3rdparty/lib/x86/libnative_camera_r4.3.0.so and b/3rdparty/lib/x86/libnative_camera_r4.3.0.so differ diff --git a/3rdparty/lib/x86/libnative_camera_r4.4.0.so b/3rdparty/lib/x86/libnative_camera_r4.4.0.so index 108862f56..37ab6d080 100755 Binary files a/3rdparty/lib/x86/libnative_camera_r4.4.0.so and b/3rdparty/lib/x86/libnative_camera_r4.4.0.so differ diff --git a/CMakeLists.txt b/CMakeLists.txt index 2d4704779..442edf32c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -438,7 +438,6 @@ include(cmake/OpenCVFindLibsGUI.cmake) include(cmake/OpenCVFindLibsVideo.cmake) include(cmake/OpenCVFindLibsPerf.cmake) - # ---------------------------------------------------------------------------- # Detect other 3rd-party libraries/tools # ---------------------------------------------------------------------------- diff --git a/cmake/OpenCVDetectCUDA.cmake b/cmake/OpenCVDetectCUDA.cmake index 89602acaa..2685171bb 100644 --- a/cmake/OpenCVDetectCUDA.cmake +++ b/cmake/OpenCVDetectCUDA.cmake @@ -217,3 +217,42 @@ else() unset(CUDA_ARCH_BIN CACHE) unset(CUDA_ARCH_PTX CACHE) endif() + +if(HAVE_CUDA) + set(CUDA_LIBS_PATH "") + foreach(p ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY}) + get_filename_component(_tmp ${p} PATH) + list(APPEND CUDA_LIBS_PATH ${_tmp}) + endforeach() + + if(HAVE_CUBLAS) + foreach(p ${CUDA_cublas_LIBRARY}) + get_filename_component(_tmp ${p} PATH) + list(APPEND CUDA_LIBS_PATH ${_tmp}) + endforeach() + endif() + + if(HAVE_CUFFT) + foreach(p ${CUDA_cufft_LIBRARY}) + get_filename_component(_tmp ${p} PATH) + list(APPEND CUDA_LIBS_PATH ${_tmp}) + endforeach() + endif() + + list(REMOVE_DUPLICATES CUDA_LIBS_PATH) + link_directories(${CUDA_LIBS_PATH}) + + set(CUDA_LIBRARIES_ABS ${CUDA_LIBRARIES}) + ocv_convert_to_lib_name(CUDA_LIBRARIES ${CUDA_LIBRARIES}) + set(CUDA_npp_LIBRARY_ABS ${CUDA_npp_LIBRARY}) + ocv_convert_to_lib_name(CUDA_npp_LIBRARY ${CUDA_npp_LIBRARY}) + if(HAVE_CUBLAS) + set(CUDA_cublas_LIBRARY_ABS ${CUDA_cublas_LIBRARY}) + ocv_convert_to_lib_name(CUDA_cublas_LIBRARY ${CUDA_cublas_LIBRARY}) + endif() + + if(HAVE_CUFFT) + set(CUDA_cufft_LIBRARY_ABS ${CUDA_cufft_LIBRARY}) + ocv_convert_to_lib_name(CUDA_cufft_LIBRARY ${CUDA_cufft_LIBRARY}) + endif() +endif() diff --git a/cmake/OpenCVModule.cmake b/cmake/OpenCVModule.cmake index 372d4504c..e6fa19911 100644 --- a/cmake/OpenCVModule.cmake +++ b/cmake/OpenCVModule.cmake @@ -27,7 +27,8 @@ # The verbose template for OpenCV module: # # ocv_add_module(modname ) -# ocv_glob_module_sources() or glob them manually and ocv_set_module_sources(...) +# ocv_glob_module_sources(([EXCLUDE_CUDA] ) +# or glob them manually and ocv_set_module_sources(...) # ocv_module_include_directories() # ocv_create_module() # @@ -478,9 +479,15 @@ endmacro() # finds and sets headers and sources for the standard OpenCV module # Usage: -# ocv_glob_module_sources() +# ocv_glob_module_sources([EXCLUDE_CUDA] ) macro(ocv_glob_module_sources) - file(GLOB_RECURSE lib_srcs "src/*.cpp") + set(_argn ${ARGN}) + list(FIND _argn "EXCLUDE_CUDA" exclude_cuda) + if(NOT exclude_cuda EQUAL -1) + list(REMOVE_AT _argn ${exclude_cuda}) + endif() + + file(GLOB_RECURSE lib_srcs "src/*.cpp") file(GLOB_RECURSE lib_int_hdrs "src/*.hpp" "src/*.h") file(GLOB lib_hdrs "include/opencv2/*.hpp" "include/opencv2/${name}/*.hpp" "include/opencv2/${name}/*.h") file(GLOB lib_hdrs_detail "include/opencv2/${name}/detail/*.hpp" "include/opencv2/${name}/detail/*.h") @@ -492,15 +499,21 @@ macro(ocv_glob_module_sources) ocv_source_group("Src" DIRBASE "${CMAKE_CURRENT_SOURCE_DIR}/src" FILES ${lib_srcs} ${lib_int_hdrs}) ocv_source_group("Include" DIRBASE "${CMAKE_CURRENT_SOURCE_DIR}/include" FILES ${lib_hdrs} ${lib_hdrs_detail}) - file(GLOB lib_cuda_srcs "src/cuda/*.cu") - set(cuda_objs "") - set(lib_cuda_hdrs "") - if(HAVE_CUDA AND lib_cuda_srcs) - ocv_include_directories(${CUDA_INCLUDE_DIRS}) - file(GLOB lib_cuda_hdrs "src/cuda/*.hpp") + if (exclude_cuda EQUAL -1) + file(GLOB lib_cuda_srcs "src/cuda/*.cu") + set(cuda_objs "") + set(lib_cuda_hdrs "") + if(HAVE_CUDA) + ocv_include_directories(${CUDA_INCLUDE_DIRS}) + file(GLOB lib_cuda_hdrs "src/cuda/*.hpp") - ocv_cuda_compile(cuda_objs ${lib_cuda_srcs} ${lib_cuda_hdrs}) - source_group("Src\\Cuda" FILES ${lib_cuda_srcs} ${lib_cuda_hdrs}) + ocv_cuda_compile(cuda_objs ${lib_cuda_srcs} ${lib_cuda_hdrs}) + source_group("Src\\Cuda" FILES ${lib_cuda_srcs} ${lib_cuda_hdrs}) + endif() + else() + set(cuda_objs "") + set(lib_cuda_srcs "") + set(lib_cuda_hdrs "") endif() file(GLOB cl_kernels "src/opencl/*.cl") @@ -516,8 +529,8 @@ macro(ocv_glob_module_sources) list(APPEND lib_srcs ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp") endif() - ocv_set_module_sources(${ARGN} HEADERS ${lib_hdrs} ${lib_hdrs_detail} - SOURCES ${lib_srcs} ${lib_int_hdrs} ${cuda_objs} ${lib_cuda_srcs} ${lib_cuda_hdrs}) + ocv_set_module_sources(${_argn} HEADERS ${lib_hdrs} ${lib_hdrs_detail} + SOURCES ${lib_srcs} ${lib_int_hdrs} ${cuda_objs} ${lib_cuda_srcs} ${lib_cuda_hdrs}) endmacro() # creates OpenCV module in current folder @@ -622,11 +635,20 @@ endmacro() # short command for adding simple OpenCV module # see ocv_add_module for argument details # Usage: -# ocv_define_module(module_name [INTERNAL] [REQUIRED] [] [OPTIONAL ]) +# ocv_define_module(module_name [INTERNAL] [EXCLUDE_CUDA] [REQUIRED] [] [OPTIONAL ]) macro(ocv_define_module module_name) - ocv_add_module(${module_name} ${ARGN}) + set(_argn ${ARGN}) + set(exclude_cuda "") + foreach(arg ${_argn}) + if("${arg}" STREQUAL "EXCLUDE_CUDA") + set(exclude_cuda "${arg}") + list(REMOVE_ITEM _argn ${arg}) + endif() + endforeach() + + ocv_add_module(${module_name} ${_argn}) ocv_module_include_directories() - ocv_glob_module_sources() + ocv_glob_module_sources(${exclude_cuda}) ocv_create_module() ocv_add_precompiled_headers(${the_module}) diff --git a/cmake/templates/OpenCVConfig.cmake.in b/cmake/templates/OpenCVConfig.cmake.in index 88eed8ee0..2d5a14ce3 100644 --- a/cmake/templates/OpenCVConfig.cmake.in +++ b/cmake/templates/OpenCVConfig.cmake.in @@ -19,8 +19,8 @@ # This file will define the following variables: # - OpenCV_LIBS : The list of all imported targets for OpenCV modules. # - OpenCV_INCLUDE_DIRS : The OpenCV include directories. -# - OpenCV_COMPUTE_CAPABILITIES : The version of compute capability -# - OpenCV_ANDROID_NATIVE_API_LEVEL : Minimum required level of Android API +# - OpenCV_COMPUTE_CAPABILITIES : The version of compute capability. +# - OpenCV_ANDROID_NATIVE_API_LEVEL : Minimum required level of Android API. # - OpenCV_VERSION : The version of this OpenCV build: "@OPENCV_VERSION_PLAIN@" # - OpenCV_VERSION_MAJOR : Major version part of OpenCV_VERSION: "@OPENCV_VERSION_MAJOR@" # - OpenCV_VERSION_MINOR : Minor version part of OpenCV_VERSION: "@OPENCV_VERSION_MINOR@" @@ -28,25 +28,29 @@ # - OpenCV_VERSION_STATUS : Development status of this build: "@OPENCV_VERSION_STATUS@" # # Advanced variables: -# - OpenCV_SHARED -# - OpenCV_CONFIG_PATH -# - OpenCV_INSTALL_PATH (not set on Windows) -# - OpenCV_LIB_COMPONENTS -# - OpenCV_USE_MANGLED_PATHS -# - OpenCV_HAVE_ANDROID_CAMERA +# - OpenCV_SHARED : Use OpenCV as shared library +# - OpenCV_CONFIG_PATH : Path to this OpenCVConfig.cmake +# - OpenCV_INSTALL_PATH : OpenCV location (not set on Windows) +# - OpenCV_LIB_COMPONENTS : Present OpenCV modules list +# - OpenCV_USE_MANGLED_PATHS : Mangled OpenCV path flag +# - OpenCV_MODULES_SUFFIX : The suffix for OpenCVModules-XXX.cmake file +# - OpenCV_HAVE_ANDROID_CAMERA : Presence of Android native camera wrappers # # Deprecated variables: # - OpenCV_VERSION_TWEAK : Always "0" # # =================================================================================== -set(modules_file_suffix "") -if(ANDROID) - string(REPLACE - _ modules_file_suffix "_${ANDROID_NDK_ABI_NAME}") +if(NOT DEFINED OpenCV_MODULES_SUFFIX) + if(ANDROID) + string(REPLACE - _ OpenCV_MODULES_SUFFIX "_${ANDROID_NDK_ABI_NAME}") + else() + set(OpenCV_MODULES_SUFFIX "") + endif() endif() if(NOT TARGET opencv_core) - include(${CMAKE_CURRENT_LIST_DIR}/OpenCVModules${modules_file_suffix}.cmake) + include(${CMAKE_CURRENT_LIST_DIR}/OpenCVModules${OpenCV_MODULES_SUFFIX}.cmake) endif() # TODO All things below should be reviewed. What is about of moving this code into related modules (special vars/hooks/files) diff --git a/doc/py_tutorials/py_imgproc/py_contours/py_contour_features/py_contour_features.rst b/doc/py_tutorials/py_imgproc/py_contours/py_contour_features/py_contour_features.rst index 6b7c661cc..8220fb501 100644 --- a/doc/py_tutorials/py_imgproc/py_contours/py_contour_features/py_contour_features.rst +++ b/doc/py_tutorials/py_imgproc/py_contours/py_contour_features/py_contour_features.rst @@ -123,7 +123,7 @@ Let (x,y) be the top-left coordinate of the rectangle and (w,h) be its width and 7.b. Rotated Rectangle ----------------------- -Here, bounding rectangle is drawn with minimum area, so it considers the rotation also. The function used is **cv2.minAreaRect()**. It returns a Box2D structure which contains following detals - ( top-left corner(x,y), (width, height), angle of rotation ). But to draw this rectangle, we need 4 corners of the rectangle. It is obtained by the function **cv2.boxPoints()** +Here, bounding rectangle is drawn with minimum area, so it considers the rotation also. The function used is **cv2.minAreaRect()**. It returns a Box2D structure which contains following detals - ( center (x,y), (width, height), angle of rotation ). But to draw this rectangle, we need 4 corners of the rectangle. It is obtained by the function **cv2.boxPoints()** :: rect = cv2.minAreaRect(cnt) diff --git a/doc/py_tutorials/py_video/py_meanshift/py_meanshift.rst b/doc/py_tutorials/py_video/py_meanshift/py_meanshift.rst index a111311af..87ece6935 100644 --- a/doc/py_tutorials/py_video/py_meanshift/py_meanshift.rst +++ b/doc/py_tutorials/py_video/py_meanshift/py_meanshift.rst @@ -52,7 +52,7 @@ To use meanshift in OpenCV, first we need to setup the target, find its histogra # set up the ROI for tracking roi = frame[r:r+h, c:c+w] - hsv_roi = cv2.cvtColor(frame, cv2.COLOR_BGR2HSV) + hsv_roi = cv2.cvtColor(roi, cv2.COLOR_BGR2HSV) mask = cv2.inRange(hsv_roi, np.array((0., 60.,32.)), np.array((180.,255.,255.))) roi_hist = cv2.calcHist([hsv_roi],[0],mask,[180],[0,180]) cv2.normalize(roi_hist,roi_hist,0,255,cv2.NORM_MINMAX) @@ -127,7 +127,7 @@ It is almost same as meanshift, but it returns a rotated rectangle (that is our # set up the ROI for tracking roi = frame[r:r+h, c:c+w] - hsv_roi = cv2.cvtColor(frame, cv2.COLOR_BGR2HSV) + hsv_roi = cv2.cvtColor(roi, cv2.COLOR_BGR2HSV) mask = cv2.inRange(hsv_roi, np.array((0., 60.,32.)), np.array((180.,255.,255.))) roi_hist = cv2.calcHist([hsv_roi],[0],mask,[180],[0,180]) cv2.normalize(roi_hist,roi_hist,0,255,cv2.NORM_MINMAX) diff --git a/doc/tutorials/core/adding_images/adding_images.rst b/doc/tutorials/core/adding_images/adding_images.rst index 3559132dc..a6d201e43 100644 --- a/doc/tutorials/core/adding_images/adding_images.rst +++ b/doc/tutorials/core/adding_images/adding_images.rst @@ -6,12 +6,12 @@ Adding (blending) two images using OpenCV Goal ===== -In this tutorial you will learn how to: +In this tutorial you will learn: .. container:: enumeratevisibleitemswithsquare - * What is *linear blending* and why it is useful. - * Add two images using :add_weighted:`addWeighted <>` + * what is *linear blending* and why it is useful; + * how to add two images using :add_weighted:`addWeighted <>` Theory ======= diff --git a/doc/tutorials/core/how_to_scan_images/how_to_scan_images.rst b/doc/tutorials/core/how_to_scan_images/how_to_scan_images.rst index ef0f8640c..b6a18fee8 100644 --- a/doc/tutorials/core/how_to_scan_images/how_to_scan_images.rst +++ b/doc/tutorials/core/how_to_scan_images/how_to_scan_images.rst @@ -18,7 +18,7 @@ We'll seek answers for the following questions: Our test case ============= -Let us consider a simple color reduction method. Using the unsigned char C and C++ type for matrix item storing a channel of pixel may have up to 256 different values. For a three channel image this can allow the formation of way too many colors (16 million to be exact). Working with so many color shades may give a heavy blow to our algorithm performance. However, sometimes it is enough to work with a lot less of them to get the same final result. +Let us consider a simple color reduction method. By using the unsigned char C and C++ type for matrix item storing, a channel of pixel may have up to 256 different values. For a three channel image this can allow the formation of way too many colors (16 million to be exact). Working with so many color shades may give a heavy blow to our algorithm performance. However, sometimes it is enough to work with a lot less of them to get the same final result. In this cases it's common that we make a *color space reduction*. This means that we divide the color space current value with a new input value to end up with fewer colors. For instance every value between zero and nine takes the new value zero, every value between ten and nineteen the value ten and so on. diff --git a/doc/tutorials/imgproc/histograms/histogram_comparison/histogram_comparison.rst b/doc/tutorials/imgproc/histograms/histogram_comparison/histogram_comparison.rst index 6a8345d69..f5f636d08 100644 --- a/doc/tutorials/imgproc/histograms/histogram_comparison/histogram_comparison.rst +++ b/doc/tutorials/imgproc/histograms/histogram_comparison/histogram_comparison.rst @@ -84,88 +84,10 @@ Code * **Code at glance:** -.. code-block:: cpp +.. literalinclude:: ../../../../../samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp + :language: cpp + :tab-width: 4 - #include "opencv2/highgui.hpp" - #include "opencv2/imgproc.hpp" - #include - #include - - using namespace std; - using namespace cv; - - /** @function main */ - int main( int argc, char** argv ) - { - Mat src_base, hsv_base; - Mat src_test1, hsv_test1; - Mat src_test2, hsv_test2; - Mat hsv_half_down; - - /// Load three images with different environment settings - if( argc < 4 ) - { printf("** Error. Usage: ./compareHist_Demo \n"); - return -1; - } - - src_base = imread( argv[1], 1 ); - src_test1 = imread( argv[2], 1 ); - src_test2 = imread( argv[3], 1 ); - - /// Convert to HSV - cvtColor( src_base, hsv_base, CV_BGR2HSV ); - cvtColor( src_test1, hsv_test1, CV_BGR2HSV ); - cvtColor( src_test2, hsv_test2, CV_BGR2HSV ); - - hsv_half_down = hsv_base( Range( hsv_base.rows/2, hsv_base.rows - 1 ), Range( 0, hsv_base.cols - 1 ) ); - - /// Using 30 bins for hue and 32 for saturation - int h_bins = 50; int s_bins = 60; - int histSize[] = { h_bins, s_bins }; - - // hue varies from 0 to 256, saturation from 0 to 180 - float h_ranges[] = { 0, 256 }; - float s_ranges[] = { 0, 180 }; - - const float* ranges[] = { h_ranges, s_ranges }; - - // Use the o-th and 1-st channels - int channels[] = { 0, 1 }; - - /// Histograms - MatND hist_base; - MatND hist_half_down; - MatND hist_test1; - MatND hist_test2; - - /// Calculate the histograms for the HSV images - calcHist( &hsv_base, 1, channels, Mat(), hist_base, 2, histSize, ranges, true, false ); - normalize( hist_base, hist_base, 0, 1, NORM_MINMAX, -1, Mat() ); - - calcHist( &hsv_half_down, 1, channels, Mat(), hist_half_down, 2, histSize, ranges, true, false ); - normalize( hist_half_down, hist_half_down, 0, 1, NORM_MINMAX, -1, Mat() ); - - calcHist( &hsv_test1, 1, channels, Mat(), hist_test1, 2, histSize, ranges, true, false ); - normalize( hist_test1, hist_test1, 0, 1, NORM_MINMAX, -1, Mat() ); - - calcHist( &hsv_test2, 1, channels, Mat(), hist_test2, 2, histSize, ranges, true, false ); - normalize( hist_test2, hist_test2, 0, 1, NORM_MINMAX, -1, Mat() ); - - /// Apply the histogram comparison methods - for( int i = 0; i < 4; i++ ) - { int compare_method = i; - double base_base = compareHist( hist_base, hist_base, compare_method ); - double base_half = compareHist( hist_base, hist_half_down, compare_method ); - double base_test1 = compareHist( hist_base, hist_test1, compare_method ); - double base_test2 = compareHist( hist_base, hist_test2, compare_method ); - - printf( " Method [%d] Perfect, Base-Half, Base-Test(1), Base-Test(2) : %f, %f, %f, %f \n", i, base_base, base_half , base_test1, base_test2 ); - } - - printf( "Done \n" ); - - return 0; - } Explanation @@ -211,11 +133,11 @@ Explanation .. code-block:: cpp - int h_bins = 50; int s_bins = 32; + int h_bins = 50; int s_bins = 60; int histSize[] = { h_bins, s_bins }; - float h_ranges[] = { 0, 256 }; - float s_ranges[] = { 0, 180 }; + float h_ranges[] = { 0, 180 }; + float s_ranges[] = { 0, 256 }; const float* ranges[] = { h_ranges, s_ranges }; diff --git a/modules/androidcamera/camera_wrapper/CMakeLists.txt b/modules/androidcamera/camera_wrapper/CMakeLists.txt index bc5585a7a..d08e2c469 100644 --- a/modules/androidcamera/camera_wrapper/CMakeLists.txt +++ b/modules/androidcamera/camera_wrapper/CMakeLists.txt @@ -58,7 +58,7 @@ SET_TARGET_PROPERTIES(${the_target} PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${EXECUTABLE_OUTPUT_PATH} ) -if (NOT (CMAKE_BUILD_TYPE MATCHES "debug")) +if (NOT (CMAKE_BUILD_TYPE MATCHES "Debug")) ADD_CUSTOM_COMMAND( TARGET ${the_target} POST_BUILD COMMAND ${CMAKE_STRIP} --strip-unneeded "${LIBRARY_OUTPUT_PATH}/lib${the_target}.so" ) endif() diff --git a/modules/androidcamera/camera_wrapper/camera_wrapper.cpp b/modules/androidcamera/camera_wrapper/camera_wrapper.cpp index 5ca1778a5..0ed301323 100644 --- a/modules/androidcamera/camera_wrapper/camera_wrapper.cpp +++ b/modules/androidcamera/camera_wrapper/camera_wrapper.cpp @@ -61,6 +61,12 @@ using namespace android; +// non-public camera related classes are not binary compatible +// objects of these classes have different sizeof on different platforms +// additional memory tail to all system objects to overcome sizeof issue +#define MAGIC_TAIL 4096 + + void debugShowFPS(); #if defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0) @@ -90,6 +96,7 @@ public: }; #endif + std::string getProcessName() { std::string result; @@ -142,12 +149,22 @@ class CameraHandler: public CameraListener protected: int cameraId; sp camera; - CameraParameters params; +#if defined(ANDROID_r3_0_1) || defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) + sp surface; +#endif +#if defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0) || defined(ANDROID_r4_4_0) + sp queue; + sp listener; +#endif + CameraParameters* params; CameraCallback cameraCallback; void* userData; int emptyCameraCallbackReported; + int width; + int height; + static const char* flashModesNames[ANDROID_CAMERA_FLASH_MODES_NUM]; static const char* focusModesNames[ANDROID_CAMERA_FOCUS_MODES_NUM]; static const char* whiteBalanceModesNames[ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM]; @@ -258,7 +275,7 @@ protected: int is_supported(const char* supp_modes_key, const char* mode) { - const char* supported_modes = params.get(supp_modes_key); + const char* supported_modes = params->get(supp_modes_key); return (supported_modes && mode && (strstr(supported_modes, mode) > 0)); } @@ -268,7 +285,7 @@ protected: if (focus_distance_type >= 0 && focus_distance_type < 3) { float focus_distances[3]; - const char* output = params.get(CameraParameters::KEY_FOCUS_DISTANCES); + const char* output = params->get(CameraParameters::KEY_FOCUS_DISTANCES); int val_num = CameraHandler::split_float(output, focus_distances, ',', 3); if(val_num == 3) { @@ -300,10 +317,15 @@ public: emptyCameraCallbackReported(0) { LOGD("Instantiated new CameraHandler (%p, %p)", callback, _userData); + void* params_buffer = operator new(sizeof(CameraParameters) + MAGIC_TAIL); + params = new(params_buffer) CameraParameters(); } virtual ~CameraHandler() { + if (params) + params->~CameraParameters(); + operator delete(params); LOGD("CameraHandler destructor is called"); } @@ -371,10 +393,18 @@ const char* CameraHandler::focusModesNames[ANDROID_CAMERA_FOCUS_MODES_NUM] = CameraParameters::FOCUS_MODE_AUTO, #if !defined(ANDROID_r2_2_0) CameraParameters::FOCUS_MODE_CONTINUOUS_VIDEO, +#else + CameraParameters::FOCUS_MODE_AUTO, #endif CameraParameters::FOCUS_MODE_EDOF, CameraParameters::FOCUS_MODE_FIXED, - CameraParameters::FOCUS_MODE_INFINITY + CameraParameters::FOCUS_MODE_INFINITY, + CameraParameters::FOCUS_MODE_MACRO, +#if !defined(ANDROID_r2_2_0) && !defined(ANDROID_r2_3_3) && !defined(ANDROID_r3_0_1) + CameraParameters::FOCUS_MODE_CONTINUOUS_PICTURE +#else + CameraParameters::FOCUS_MODE_AUTO +#endif }; const char* CameraHandler::whiteBalanceModesNames[ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM] = @@ -534,39 +564,39 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback, { LOGI("initCameraConnect: Setting paramers from previous camera handler"); camera->setParameters(prevCameraParameters->flatten()); - handler->params.unflatten(prevCameraParameters->flatten()); + handler->params->unflatten(prevCameraParameters->flatten()); } else { android::String8 params_str = camera->getParameters(); LOGI("initCameraConnect: [%s]", params_str.string()); - handler->params.unflatten(params_str); + handler->params->unflatten(params_str); - LOGD("Supported Cameras: %s", handler->params.get("camera-indexes")); - LOGD("Supported Picture Sizes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PICTURE_SIZES)); - LOGD("Supported Picture Formats: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PICTURE_FORMATS)); - LOGD("Supported Preview Sizes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_SIZES)); - LOGD("Supported Preview Formats: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_FORMATS)); - LOGD("Supported Preview Frame Rates: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_FRAME_RATES)); - LOGD("Supported Thumbnail Sizes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_JPEG_THUMBNAIL_SIZES)); - LOGD("Supported Whitebalance Modes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_WHITE_BALANCE)); - LOGD("Supported Effects: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_EFFECTS)); - LOGD("Supported Scene Modes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_SCENE_MODES)); - LOGD("Supported Focus Modes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_FOCUS_MODES)); - LOGD("Supported Antibanding Options: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_ANTIBANDING)); - LOGD("Supported Flash Modes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_FLASH_MODES)); + LOGD("Supported Cameras: %s", handler->params->get("camera-indexes")); + LOGD("Supported Picture Sizes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PICTURE_SIZES)); + LOGD("Supported Picture Formats: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PICTURE_FORMATS)); + LOGD("Supported Preview Sizes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_SIZES)); + LOGD("Supported Preview Formats: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_FORMATS)); + LOGD("Supported Preview Frame Rates: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_FRAME_RATES)); + LOGD("Supported Thumbnail Sizes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_JPEG_THUMBNAIL_SIZES)); + LOGD("Supported Whitebalance Modes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_WHITE_BALANCE)); + LOGD("Supported Effects: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_EFFECTS)); + LOGD("Supported Scene Modes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_SCENE_MODES)); + LOGD("Supported Focus Modes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_FOCUS_MODES)); + LOGD("Supported Antibanding Options: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_ANTIBANDING)); + LOGD("Supported Flash Modes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_FLASH_MODES)); #if !defined(ANDROID_r2_2_0) // Set focus mode to continuous-video if supported - const char* available_focus_modes = handler->params.get(CameraParameters::KEY_SUPPORTED_FOCUS_MODES); + const char* available_focus_modes = handler->params->get(CameraParameters::KEY_SUPPORTED_FOCUS_MODES); if (available_focus_modes != 0) { if (strstr(available_focus_modes, "continuous-video") != NULL) { - handler->params.set(CameraParameters::KEY_FOCUS_MODE, CameraParameters::FOCUS_MODE_CONTINUOUS_VIDEO); + handler->params->set(CameraParameters::KEY_FOCUS_MODE, CameraParameters::FOCUS_MODE_CONTINUOUS_VIDEO); - status_t resParams = handler->camera->setParameters(handler->params.flatten()); + status_t resParams = handler->camera->setParameters(handler->params->flatten()); if (resParams != 0) { @@ -581,7 +611,7 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback, #endif //check if yuv420sp format available. Set this format as preview format. - const char* available_formats = handler->params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_FORMATS); + const char* available_formats = handler->params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_FORMATS); if (available_formats != 0) { const char* format_to_set = 0; @@ -607,9 +637,9 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback, if (0 != format_to_set) { - handler->params.setPreviewFormat(format_to_set); + handler->params->setPreviewFormat(format_to_set); - status_t resParams = handler->camera->setParameters(handler->params.flatten()); + status_t resParams = handler->camera->setParameters(handler->params->flatten()); if (resParams != 0) LOGE("initCameraConnect: failed to set preview format to %s", format_to_set); @@ -617,6 +647,13 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback, LOGD("initCameraConnect: preview format is set to %s", format_to_set); } } + + handler->params->setPreviewSize(640, 480); + status_t resParams = handler->camera->setParameters(handler->params->flatten()); + if (resParams != 0) + LOGE("initCameraConnect: failed to set preview resolution to 640x480"); + else + LOGD("initCameraConnect: preview format is set to 640x480"); } status_t bufferStatus; @@ -627,22 +664,27 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback, #elif defined(ANDROID_r2_3_3) /* Do nothing in case of 2.3 for now */ #elif defined(ANDROID_r3_0_1) || defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) - sp surfaceTexture = new SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID); - bufferStatus = camera->setPreviewTexture(surfaceTexture); + void* surface_texture_obj = operator new(sizeof(SurfaceTexture) + MAGIC_TAIL); + handler->surface = new(surface_texture_obj) SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID); + bufferStatus = camera->setPreviewTexture(handler->surface); if (bufferStatus != 0) LOGE("initCameraConnect: failed setPreviewTexture call (status %d); camera might not work correctly", bufferStatus); #elif defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0) - sp bufferQueue = new BufferQueue(); - sp queueListener = new ConsumerListenerStub(); - bufferQueue->consumerConnect(queueListener); - bufferStatus = camera->setPreviewTexture(bufferQueue); + void* buffer_queue_obj = operator new(sizeof(BufferQueue) + MAGIC_TAIL); + handler->queue = new(buffer_queue_obj) BufferQueue(); + void* consumer_listener_obj = operator new(sizeof(ConsumerListenerStub) + MAGIC_TAIL); + handler->listener = new(consumer_listener_obj) ConsumerListenerStub(); + handler->queue->consumerConnect(handler->listener); + bufferStatus = camera->setPreviewTexture(handler->queue); if (bufferStatus != 0) LOGE("initCameraConnect: failed setPreviewTexture call; camera might not work correctly"); # elif defined(ANDROID_r4_4_0) - sp bufferQueue = new BufferQueue(); - sp queueListener = new ConsumerListenerStub(); - bufferQueue->consumerConnect(queueListener, true); - bufferStatus = handler->camera->setPreviewTarget(bufferQueue); + void* buffer_queue_obj = operator new(sizeof(BufferQueue) + MAGIC_TAIL); + handler->queue = new(buffer_queue_obj) BufferQueue(); + void* consumer_listener_obj = operator new(sizeof(ConsumerListenerStub) + MAGIC_TAIL); + handler->listener = new(consumer_listener_obj) ConsumerListenerStub(); + handler->queue->consumerConnect(handler->listener, true); + bufferStatus = handler->camera->setPreviewTarget(handler->queue); if (bufferStatus != 0) LOGE("applyProperties: failed setPreviewTexture call; camera might not work correctly"); # endif @@ -723,18 +765,18 @@ double CameraHandler::getProperty(int propIdx) case ANDROID_CAMERA_PROPERTY_FRAMEWIDTH: { int w,h; - params.getPreviewSize(&w, &h); + params->getPreviewSize(&w, &h); return w; } case ANDROID_CAMERA_PROPERTY_FRAMEHEIGHT: { int w,h; - params.getPreviewSize(&w, &h); + params->getPreviewSize(&w, &h); return h; } case ANDROID_CAMERA_PROPERTY_SUPPORTED_PREVIEW_SIZES_STRING: { - cameraPropertySupportedPreviewSizesString = params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_SIZES); + cameraPropertySupportedPreviewSizesString = params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_SIZES); union {const char* str;double res;} u; memset(&u.res, 0, sizeof(u.res)); u.str = cameraPropertySupportedPreviewSizesString.c_str(); @@ -742,7 +784,7 @@ double CameraHandler::getProperty(int propIdx) } case ANDROID_CAMERA_PROPERTY_PREVIEW_FORMAT_STRING: { - const char* fmt = params.get(CameraParameters::KEY_PREVIEW_FORMAT); + const char* fmt = params->get(CameraParameters::KEY_PREVIEW_FORMAT); if (fmt == CameraParameters::PIXEL_FORMAT_YUV422SP) fmt = "yuv422sp"; else if (fmt == CameraParameters::PIXEL_FORMAT_YUV420SP) @@ -762,44 +804,44 @@ double CameraHandler::getProperty(int propIdx) } case ANDROID_CAMERA_PROPERTY_EXPOSURE: { - int exposure = params.getInt(CameraParameters::KEY_EXPOSURE_COMPENSATION); + int exposure = params->getInt(CameraParameters::KEY_EXPOSURE_COMPENSATION); return exposure; } case ANDROID_CAMERA_PROPERTY_FPS: { - return params.getPreviewFrameRate(); + return params->getPreviewFrameRate(); } case ANDROID_CAMERA_PROPERTY_FLASH_MODE: { int flash_mode = getModeNum(CameraHandler::flashModesNames, ANDROID_CAMERA_FLASH_MODES_NUM, - params.get(CameraParameters::KEY_FLASH_MODE)); + params->get(CameraParameters::KEY_FLASH_MODE)); return flash_mode; } case ANDROID_CAMERA_PROPERTY_FOCUS_MODE: { int focus_mode = getModeNum(CameraHandler::focusModesNames, ANDROID_CAMERA_FOCUS_MODES_NUM, - params.get(CameraParameters::KEY_FOCUS_MODE)); + params->get(CameraParameters::KEY_FOCUS_MODE)); return focus_mode; } case ANDROID_CAMERA_PROPERTY_WHITE_BALANCE: { int white_balance = getModeNum(CameraHandler::whiteBalanceModesNames, ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM, - params.get(CameraParameters::KEY_WHITE_BALANCE)); + params->get(CameraParameters::KEY_WHITE_BALANCE)); return white_balance; } case ANDROID_CAMERA_PROPERTY_ANTIBANDING: { int antibanding = getModeNum(CameraHandler::antibandingModesNames, ANDROID_CAMERA_ANTIBANDING_MODES_NUM, - params.get(CameraParameters::KEY_ANTIBANDING)); + params->get(CameraParameters::KEY_ANTIBANDING)); return antibanding; } case ANDROID_CAMERA_PROPERTY_FOCAL_LENGTH: { - float focal_length = params.getFloat(CameraParameters::KEY_FOCAL_LENGTH); + float focal_length = params->getFloat(CameraParameters::KEY_FOCAL_LENGTH); return focal_length; } case ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_NEAR: @@ -814,6 +856,24 @@ double CameraHandler::getProperty(int propIdx) { return getFocusDistance(ANDROID_CAMERA_FOCUS_DISTANCE_FAR_INDEX); } +#if !defined(ANDROID_r2_2_0) && !defined(ANDROID_r2_3_3) && !defined(ANDROID_r3_0_1) + case ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK: + { + const char* status = params->get(CameraParameters::KEY_AUTO_WHITEBALANCE_LOCK); + if (status == CameraParameters::TRUE) + return 1.; + else + return 0.; + } + case ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK: + { + const char* status = params->get(CameraParameters::KEY_AUTO_EXPOSURE_LOCK); + if (status == CameraParameters::TRUE) + return 1.; + else + return 0.; + } +#endif default: LOGW("CameraHandler::getProperty - Unsupported property."); }; @@ -824,99 +884,151 @@ void CameraHandler::setProperty(int propIdx, double value) { LOGD("CameraHandler::setProperty(%d, %f)", propIdx, value); + android::String8 params_str; + params_str = camera->getParameters(); + LOGI("Params before set: [%s]", params_str.string()); + switch (propIdx) { case ANDROID_CAMERA_PROPERTY_FRAMEWIDTH: { int w,h; - params.getPreviewSize(&w, &h); - w = (int)value; - params.setPreviewSize(w, h); + params->getPreviewSize(&w, &h); + width = (int)value; } break; case ANDROID_CAMERA_PROPERTY_FRAMEHEIGHT: { int w,h; - params.getPreviewSize(&w, &h); - h = (int)value; - params.setPreviewSize(w, h); + params->getPreviewSize(&w, &h); + height = (int)value; } break; case ANDROID_CAMERA_PROPERTY_EXPOSURE: { - int max_exposure = params.getInt("max-exposure-compensation"); - int min_exposure = params.getInt("min-exposure-compensation"); - if(max_exposure && min_exposure){ + int max_exposure = params->getInt("max-exposure-compensation"); + int min_exposure = params->getInt("min-exposure-compensation"); + if(max_exposure && min_exposure) + { int exposure = (int)value; - if(exposure >= min_exposure && exposure <= max_exposure){ - params.set("exposure-compensation", exposure); - } else { + if(exposure >= min_exposure && exposure <= max_exposure) + params->set("exposure-compensation", exposure); + else LOGE("Exposure compensation not in valid range (%i,%i).", min_exposure, max_exposure); - } - } else { + } else LOGE("Exposure compensation adjust is not supported."); - } + + camera->setParameters(params->flatten()); } break; case ANDROID_CAMERA_PROPERTY_FLASH_MODE: { int new_val = (int)value; - if(new_val >= 0 && new_val < ANDROID_CAMERA_FLASH_MODES_NUM){ + if(new_val >= 0 && new_val < ANDROID_CAMERA_FLASH_MODES_NUM) + { const char* mode_name = flashModesNames[new_val]; if(is_supported(CameraParameters::KEY_SUPPORTED_FLASH_MODES, mode_name)) - params.set(CameraParameters::KEY_FLASH_MODE, mode_name); + params->set(CameraParameters::KEY_FLASH_MODE, mode_name); else LOGE("Flash mode %s is not supported.", mode_name); - } else { - LOGE("Flash mode value not in valid range."); } + else + LOGE("Flash mode value not in valid range."); + + camera->setParameters(params->flatten()); } break; case ANDROID_CAMERA_PROPERTY_FOCUS_MODE: { int new_val = (int)value; - if(new_val >= 0 && new_val < ANDROID_CAMERA_FOCUS_MODES_NUM){ + if(new_val >= 0 && new_val < ANDROID_CAMERA_FOCUS_MODES_NUM) + { const char* mode_name = focusModesNames[new_val]; if(is_supported(CameraParameters::KEY_SUPPORTED_FOCUS_MODES, mode_name)) - params.set(CameraParameters::KEY_FOCUS_MODE, mode_name); + params->set(CameraParameters::KEY_FOCUS_MODE, mode_name); else LOGE("Focus mode %s is not supported.", mode_name); - } else { - LOGE("Focus mode value not in valid range."); } + else + LOGE("Focus mode value not in valid range."); + + camera->setParameters(params->flatten()); } break; case ANDROID_CAMERA_PROPERTY_WHITE_BALANCE: { int new_val = (int)value; - if(new_val >= 0 && new_val < ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM){ + if(new_val >= 0 && new_val < ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM) + { const char* mode_name = whiteBalanceModesNames[new_val]; if(is_supported(CameraParameters::KEY_SUPPORTED_WHITE_BALANCE, mode_name)) - params.set(CameraParameters::KEY_WHITE_BALANCE, mode_name); + params->set(CameraParameters::KEY_WHITE_BALANCE, mode_name); else LOGE("White balance mode %s is not supported.", mode_name); - } else { - LOGE("White balance mode value not in valid range."); } + else + LOGE("White balance mode value not in valid range."); + + camera->setParameters(params->flatten()); } break; case ANDROID_CAMERA_PROPERTY_ANTIBANDING: { int new_val = (int)value; - if(new_val >= 0 && new_val < ANDROID_CAMERA_ANTIBANDING_MODES_NUM){ + if(new_val >= 0 && new_val < ANDROID_CAMERA_ANTIBANDING_MODES_NUM) + { const char* mode_name = antibandingModesNames[new_val]; if(is_supported(CameraParameters::KEY_SUPPORTED_ANTIBANDING, mode_name)) - params.set(CameraParameters::KEY_ANTIBANDING, mode_name); + params->set(CameraParameters::KEY_ANTIBANDING, mode_name); else LOGE("Antibanding mode %s is not supported.", mode_name); - } else { - LOGE("Antibanding mode value not in valid range."); } + else + LOGE("Antibanding mode value not in valid range."); + + camera->setParameters(params->flatten()); } break; +#if !defined(ANDROID_r2_2_0) && !defined(ANDROID_r2_3_3) && !defined(ANDROID_r3_0_1) + case ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK: + { + if (is_supported(CameraParameters::KEY_AUTO_EXPOSURE_LOCK_SUPPORTED, "true")) + { + if (value != 0) + params->set(CameraParameters::KEY_AUTO_EXPOSURE_LOCK, CameraParameters::TRUE); + else + params->set(CameraParameters::KEY_AUTO_EXPOSURE_LOCK, CameraParameters::FALSE); + LOGE("Expose lock is set"); + } + else + LOGE("Expose lock is not supported"); + + camera->setParameters(params->flatten()); + } + break; + case ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK: + { + if (is_supported(CameraParameters::KEY_AUTO_WHITEBALANCE_LOCK_SUPPORTED, "true")) + { + if (value != 0) + params->set(CameraParameters::KEY_AUTO_WHITEBALANCE_LOCK, CameraParameters::TRUE); + else + params->set(CameraParameters::KEY_AUTO_WHITEBALANCE_LOCK, CameraParameters::FALSE); + LOGE("White balance lock is set"); + } + else + LOGE("White balance lock is not supported"); + + camera->setParameters(params->flatten()); + } + break; +#endif default: LOGW("CameraHandler::setProperty - Unsupported property."); }; + + params_str = camera->getParameters(); + LOGI("Params after set: [%s]", params_str.string()); } void CameraHandler::applyProperties(CameraHandler** ppcameraHandler) @@ -935,7 +1047,10 @@ void CameraHandler::applyProperties(CameraHandler** ppcameraHandler) return; } - CameraParameters curCameraParameters((*ppcameraHandler)->params.flatten()); + // delayed resolution setup to exclude errors during other parameres setup on the fly + // without camera restart + if (((*ppcameraHandler)->width != 0) && ((*ppcameraHandler)->height != 0)) + (*ppcameraHandler)->params->setPreviewSize((*ppcameraHandler)->width, (*ppcameraHandler)->height); #if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) || defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) \ || defined(ANDROID_r4_3_0) || defined(ANDROID_r4_4_0) @@ -951,27 +1066,27 @@ void CameraHandler::applyProperties(CameraHandler** ppcameraHandler) return; } - handler->camera->setParameters(curCameraParameters.flatten()); - handler->params.unflatten(curCameraParameters.flatten()); + handler->camera->setParameters((*ppcameraHandler)->params->flatten()); status_t bufferStatus; # if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) - sp surfaceTexture = new SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID); - bufferStatus = handler->camera->setPreviewTexture(surfaceTexture); + void* surface_texture_obj = operator new(sizeof(SurfaceTexture) + MAGIC_TAIL); + handler->surface = new(surface_texture_obj) SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID); + bufferStatus = handler->camera->setPreviewTexture(handler->surface); if (bufferStatus != 0) LOGE("applyProperties: failed setPreviewTexture call (status %d); camera might not work correctly", bufferStatus); # elif defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0) - sp bufferQueue = new BufferQueue(); - sp queueListener = new ConsumerListenerStub(); - bufferQueue->consumerConnect(queueListener); - bufferStatus = handler->camera->setPreviewTexture(bufferQueue); + void* buffer_queue_obj = operator new(sizeof(BufferQueue) + MAGIC_TAIL); + handler->queue = new(buffer_queue_obj) BufferQueue(); + handler->queue->consumerConnect(handler->listener); + bufferStatus = handler->camera->setPreviewTexture(handler->queue); if (bufferStatus != 0) LOGE("applyProperties: failed setPreviewTexture call; camera might not work correctly"); # elif defined(ANDROID_r4_4_0) - sp bufferQueue = new BufferQueue(); - sp queueListener = new ConsumerListenerStub(); - bufferQueue->consumerConnect(queueListener, true); - bufferStatus = handler->camera->setPreviewTarget(bufferQueue); + void* buffer_queue_obj = operator new(sizeof(BufferQueue) + MAGIC_TAIL); + handler->queue = new(buffer_queue_obj) BufferQueue(); + handler->queue->consumerConnect(handler->listener, true); + bufferStatus = handler->camera->setPreviewTarget(handler->queue); if (bufferStatus != 0) LOGE("applyProperties: failed setPreviewTexture call; camera might not work correctly"); # endif @@ -1002,7 +1117,7 @@ void CameraHandler::applyProperties(CameraHandler** ppcameraHandler) LOGD("CameraHandler::applyProperties(): after previousCameraHandler->closeCameraConnect"); LOGD("CameraHandler::applyProperties(): before initCameraConnect"); - CameraHandler* handler=initCameraConnect(cameraCallback, cameraId, userData, &curCameraParameters); + CameraHandler* handler=initCameraConnect(cameraCallback, cameraId, userData, (*ppcameraHandler)->params); LOGD("CameraHandler::applyProperties(): after initCameraConnect, handler=0x%x", (int)handler); if (handler == NULL) { LOGE("ERROR in applyProperties --- cannot reinit camera"); diff --git a/modules/androidcamera/include/camera_properties.h b/modules/androidcamera/include/camera_properties.h index 2fec745fa..65499be2d 100644 --- a/modules/androidcamera/include/camera_properties.h +++ b/modules/androidcamera/include/camera_properties.h @@ -15,7 +15,9 @@ enum { ANDROID_CAMERA_PROPERTY_FOCAL_LENGTH = 105, ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_NEAR = 106, ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_OPTIMAL = 107, - ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_FAR = 108 + ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_FAR = 108, + ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK = 109, + ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK = 110 }; @@ -30,12 +32,12 @@ enum { enum { ANDROID_CAMERA_FOCUS_MODE_AUTO = 0, - ANDROID_CAMERA_FOCUS_MODE_CONTINUOUS_PICTURE, ANDROID_CAMERA_FOCUS_MODE_CONTINUOUS_VIDEO, ANDROID_CAMERA_FOCUS_MODE_EDOF, ANDROID_CAMERA_FOCUS_MODE_FIXED, ANDROID_CAMERA_FOCUS_MODE_INFINITY, ANDROID_CAMERA_FOCUS_MODE_MACRO, + ANDROID_CAMERA_FOCUS_MODE_CONTINUOUS_PICTURE, ANDROID_CAMERA_FOCUS_MODES_NUM }; diff --git a/modules/core/include/opencv2/core/mat.hpp b/modules/core/include/opencv2/core/mat.hpp index 33167fab2..6b8368fd5 100644 --- a/modules/core/include/opencv2/core/mat.hpp +++ b/modules/core/include/opencv2/core/mat.hpp @@ -118,6 +118,8 @@ public: virtual int kind() const; virtual int dims(int i=-1) const; + virtual int cols(int i=-1) const; + virtual int rows(int i=-1) const; virtual Size size(int i=-1) const; virtual int sizend(int* sz, int i=-1) const; virtual bool sameSize(const _InputArray& arr) const; diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index fb9ec24c5..fdb6f9a0a 100644 --- a/modules/core/include/opencv2/core/ocl.hpp +++ b/modules/core/include/opencv2/core/ocl.hpp @@ -592,7 +592,7 @@ protected: CV_EXPORTS const char* convertTypeStr(int sdepth, int ddepth, int cn, char* buf); CV_EXPORTS const char* typeToStr(int t); CV_EXPORTS const char* memopTypeToStr(int t); -CV_EXPORTS String kernelToStr(InputArray _kernel, int ddepth = -1); +CV_EXPORTS String kernelToStr(InputArray _kernel, int ddepth = -1, const char * name = NULL); CV_EXPORTS void getPlatfomsInfo(std::vector& platform_info); CV_EXPORTS int predictOptimalVectorWidth(InputArray src1, InputArray src2 = noArray(), InputArray src3 = noArray(), InputArray src4 = noArray(), InputArray src5 = noArray(), InputArray src6 = noArray(), diff --git a/modules/core/src/copy.cpp b/modules/core/src/copy.cpp index e3e959c95..5ac5f22c5 100644 --- a/modules/core/src/copy.cpp +++ b/modules/core/src/copy.cpp @@ -482,9 +482,9 @@ enum { FLIP_COLS = 1 << 0, FLIP_ROWS = 1 << 1, FLIP_BOTH = FLIP_ROWS | FLIP_COLS static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode ) { CV_Assert(flipCode >= - 1 && flipCode <= 1); - int type = _src.type(), cn = CV_MAT_CN(type), flipType; + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), flipType; - if (cn > 4 || cn == 3) + if (cn > 4) return false; const char * kernelName; @@ -506,7 +506,8 @@ static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode ) } ocl::Kernel k(kernelName, ocl::core::flip_oclsrc, - format( "-D type=%s", ocl::memopTypeToStr(type))); + format( "-D T=%s -D T1=%s -D cn=%d", ocl::memopTypeToStr(type), + ocl::memopTypeToStr(depth), cn)); if (k.empty()) return false; diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index db1ce760f..45ae3d512 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -1416,6 +1416,16 @@ int _InputArray::kind() const return flags & KIND_MASK; } +int _InputArray::rows(int i) const +{ + return size(i).height; +} + +int _InputArray::cols(int i) const +{ + return size(i).width; +} + Size _InputArray::size(int i) const { int k = kind(); @@ -2078,45 +2088,45 @@ void _OutputArray::create(Size _sz, int mtype, int i, bool allowTransposed, int create(2, sizes, mtype, i, allowTransposed, fixedDepthMask); } -void _OutputArray::create(int rows, int cols, int mtype, int i, bool allowTransposed, int fixedDepthMask) const +void _OutputArray::create(int _rows, int _cols, int mtype, int i, bool allowTransposed, int fixedDepthMask) const { int k = kind(); if( k == MAT && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((Mat*)obj)->size.operator()() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((Mat*)obj)->size.operator()() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((Mat*)obj)->type() == mtype); - ((Mat*)obj)->create(rows, cols, mtype); + ((Mat*)obj)->create(_rows, _cols, mtype); return; } if( k == UMAT && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((UMat*)obj)->size.operator()() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((UMat*)obj)->size.operator()() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((UMat*)obj)->type() == mtype); - ((UMat*)obj)->create(rows, cols, mtype); + ((UMat*)obj)->create(_rows, _cols, mtype); return; } if( k == GPU_MAT && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((cuda::GpuMat*)obj)->size() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((cuda::GpuMat*)obj)->size() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((cuda::GpuMat*)obj)->type() == mtype); - ((cuda::GpuMat*)obj)->create(rows, cols, mtype); + ((cuda::GpuMat*)obj)->create(_rows, _cols, mtype); return; } if( k == OPENGL_BUFFER && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((ogl::Buffer*)obj)->size() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((ogl::Buffer*)obj)->size() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((ogl::Buffer*)obj)->type() == mtype); - ((ogl::Buffer*)obj)->create(rows, cols, mtype); + ((ogl::Buffer*)obj)->create(_rows, _cols, mtype); return; } if( k == CUDA_MEM && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((cuda::CudaMem*)obj)->size() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((cuda::CudaMem*)obj)->size() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((cuda::CudaMem*)obj)->type() == mtype); - ((cuda::CudaMem*)obj)->create(rows, cols, mtype); + ((cuda::CudaMem*)obj)->create(_rows, _cols, mtype); return; } - int sizes[] = {rows, cols}; + int sizes[] = {_rows, _cols}; create(2, sizes, mtype, i, allowTransposed, fixedDepthMask); } @@ -2679,17 +2689,17 @@ namespace cv { static bool ocl_setIdentity( InputOutputArray _m, const Scalar& s ) { - int type = _m.type(), cn = CV_MAT_CN(type); - if (cn == 3) - return false; + int type = _m.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), + sctype = CV_MAKE_TYPE(depth, cn == 3 ? 4 : cn); ocl::Kernel k("setIdentity", ocl::core::set_identity_oclsrc, - format("-D T=%s", ocl::memopTypeToStr(type))); + format("-D T=%s -D T1=%s -D cn=%d -D ST=%s", ocl::memopTypeToStr(type), + ocl::memopTypeToStr(depth), cn, ocl::memopTypeToStr(sctype))); if (k.empty()) return false; UMat m = _m.getUMat(); - k.args(ocl::KernelArg::WriteOnly(m), ocl::KernelArg::Constant(Mat(1, 1, type, s))); + k.args(ocl::KernelArg::WriteOnly(m), ocl::KernelArg::Constant(Mat(1, 1, sctype, s))); size_t globalsize[2] = { m.cols, m.rows }; return k.run(2, globalsize, NULL, false); diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index c6e0c4039..ffea804ed 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -2179,7 +2179,6 @@ static cl_device_id selectOpenCLDevice() goto not_found; } } - if (deviceTypes.size() == 0) { if (!isID) @@ -2193,13 +2192,16 @@ static cl_device_id selectOpenCLDevice() for (size_t t = 0; t < deviceTypes.size(); t++) { int deviceType = 0; - if (deviceTypes[t] == "GPU") + std::string tempStrDeviceType = deviceTypes[t]; + std::transform( tempStrDeviceType.begin(), tempStrDeviceType.end(), tempStrDeviceType.begin(), tolower ); + + if (tempStrDeviceType == "gpu" || tempStrDeviceType == "dgpu" || tempStrDeviceType == "igpu") deviceType = Device::TYPE_GPU; - else if (deviceTypes[t] == "CPU") + else if (tempStrDeviceType == "cpu") deviceType = Device::TYPE_CPU; - else if (deviceTypes[t] == "ACCELERATOR") + else if (tempStrDeviceType == "accelerator") deviceType = Device::TYPE_ACCELERATOR; - else if (deviceTypes[t] == "ALL") + else if (tempStrDeviceType == "all") deviceType = Device::TYPE_ALL; else { @@ -2229,7 +2231,14 @@ static cl_device_id selectOpenCLDevice() { std::string name; CV_OclDbgAssert(getStringInfo(clGetDeviceInfo, devices[i], CL_DEVICE_NAME, name) == CL_SUCCESS); - if (isID || name.find(deviceName) != std::string::npos) + cl_bool useGPU = true; + if(tempStrDeviceType == "dgpu" || tempStrDeviceType == "igpu") + { + cl_bool isIGPU = CL_FALSE; + clGetDeviceInfo(devices[i], CL_DEVICE_HOST_UNIFIED_MEMORY, sizeof(isIGPU), &isIGPU, NULL); + useGPU = tempStrDeviceType == "dgpu" ? !isIGPU : isIGPU; + } + if ( (isID || name.find(deviceName) != std::string::npos) && useGPU) { // TODO check for OpenCL 1.1 return devices[i]; @@ -4307,7 +4316,7 @@ static std::string kerToStr(const Mat & k) return stream.str(); } -String kernelToStr(InputArray _kernel, int ddepth) +String kernelToStr(InputArray _kernel, int ddepth, const char * name) { Mat kernel = _kernel.getMat().reshape(1, 1); @@ -4318,13 +4327,13 @@ String kernelToStr(InputArray _kernel, int ddepth) if (ddepth != depth) kernel.convertTo(kernel, ddepth); - typedef std::string (*func_t)(const Mat &); - static const func_t funcs[] = { kerToStr, kerToStr, kerToStr,kerToStr, + typedef std::string (* func_t)(const Mat &); + static const func_t funcs[] = { kerToStr, kerToStr, kerToStr, kerToStr, kerToStr, kerToStr, kerToStr, 0 }; const func_t func = funcs[depth]; CV_Assert(func != 0); - return cv::format(" -D COEFF=%s", func(kernel).c_str()); + return cv::format(" -D %s=%s", name ? name : "COEFF", func(kernel).c_str()); } #define PROCESS_SRC(src) \ diff --git a/modules/core/src/opencl/flip.cl b/modules/core/src/opencl/flip.cl index 0c874dbe6..bacfe7adf 100644 --- a/modules/core/src/opencl/flip.cl +++ b/modules/core/src/opencl/flip.cl @@ -39,10 +39,18 @@ // //M*/ -#define sizeoftype ((int)sizeof(type)) +#if cn != 3 +#define loadpix(addr) *(__global const T *)(addr) +#define storepix(val, addr) *(__global T *)(addr) = val +#define TSIZE (int)sizeof(T) +#else +#define loadpix(addr) vload3(0, (__global const T1 *)(addr)) +#define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr)) +#define TSIZE ((int)sizeof(T1)*3) +#endif -__kernel void arithm_flip_rows(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +__kernel void arithm_flip_rows(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols, int thread_rows, int thread_cols) { int x = get_global_id(0); @@ -50,19 +58,16 @@ __kernel void arithm_flip_rows(__global const uchar* srcptr, int srcstep, int sr if (x < cols && y < thread_rows) { - __global const type* src0 = (__global const type*)(srcptr + mad24(y, srcstep, mad24(x, sizeoftype, srcoffset))); - __global const type* src1 = (__global const type*)(srcptr + mad24(rows - y - 1, srcstep, mad24(x, sizeoftype, srcoffset))); + T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset))); + T src1 = loadpix(srcptr + mad24(rows - y - 1, src_step, mad24(x, TSIZE, src_offset))); - __global type* dst0 = (__global type*)(dstptr + mad24(y, dststep, mad24(x, sizeoftype, dstoffset))); - __global type* dst1 = (__global type*)(dstptr + mad24(rows - y - 1, dststep, mad24(x, sizeoftype, dstoffset))); - - dst0[0] = src1[0]; - dst1[0] = src0[0]; + storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset))); + storepix(src0, dstptr + mad24(rows - y - 1, dst_step, mad24(x, TSIZE, dst_offset))); } } -__kernel void arithm_flip_rows_cols(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +__kernel void arithm_flip_rows_cols(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols, int thread_rows, int thread_cols) { int x = get_global_id(0); @@ -71,19 +76,16 @@ __kernel void arithm_flip_rows_cols(__global const uchar* srcptr, int srcstep, i if (x < cols && y < thread_rows) { int x1 = cols - x - 1; - __global const type* src0 = (__global const type*)(srcptr + mad24(y, srcstep, mad24(x, sizeoftype, srcoffset))); - __global const type* src1 = (__global const type*)(srcptr + mad24(rows - y - 1, srcstep, mad24(x1, sizeoftype, srcoffset))); + T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset))); + T src1 = loadpix(srcptr + mad24(rows - y - 1, src_step, mad24(x1, TSIZE, src_offset))); - __global type* dst0 = (__global type*)(dstptr + mad24(rows - y - 1, dststep, mad24(x1, sizeoftype, dstoffset))); - __global type* dst1 = (__global type*)(dstptr + mad24(y, dststep, mad24(x, sizeoftype, dstoffset))); - - dst0[0] = src0[0]; - dst1[0] = src1[0]; + storepix(src0, dstptr + mad24(rows - y - 1, dst_step, mad24(x1, TSIZE, dst_offset))); + storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset))); } } -__kernel void arithm_flip_cols(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +__kernel void arithm_flip_cols(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols, int thread_rows, int thread_cols) { int x = get_global_id(0); @@ -92,13 +94,10 @@ __kernel void arithm_flip_cols(__global const uchar* srcptr, int srcstep, int sr if (x < thread_cols && y < rows) { int x1 = cols - x - 1; - __global const type* src0 = (__global const type*)(srcptr + mad24(y, srcstep, mad24(x, sizeoftype, srcoffset))); - __global const type* src1 = (__global const type*)(srcptr + mad24(y, srcstep, mad24(x1, sizeoftype, srcoffset))); + T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset))); + T src1 = loadpix(srcptr + mad24(y, src_step, mad24(x1, TSIZE, src_offset))); - __global type* dst0 = (__global type*)(dstptr + mad24(y, dststep, mad24(x1, sizeoftype, dstoffset))); - __global type* dst1 = (__global type*)(dstptr + mad24(y, dststep, mad24(x, sizeoftype, dstoffset))); - - dst1[0] = src1[0]; - dst0[0] = src0[0]; + storepix(src0, dstptr + mad24(y, dst_step, mad24(x1, TSIZE, dst_offset))); + storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset))); } } diff --git a/modules/core/src/opencl/set_identity.cl b/modules/core/src/opencl/set_identity.cl index d63ce793d..0e8f1424f 100644 --- a/modules/core/src/opencl/set_identity.cl +++ b/modules/core/src/opencl/set_identity.cl @@ -43,17 +43,28 @@ // //M*/ +#if cn != 3 +#define loadpix(addr) *(__global const T *)(addr) +#define storepix(val, addr) *(__global T *)(addr) = val +#define TSIZE (int)sizeof(T) +#define scalar scalar_ +#else +#define loadpix(addr) vload3(0, (__global const T1 *)(addr)) +#define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr)) +#define TSIZE ((int)sizeof(T1)*3) +#define scalar (T)(scalar_.x, scalar_.y, scalar_.z) +#endif + __kernel void setIdentity(__global uchar * srcptr, int src_step, int src_offset, int rows, int cols, - T scalar) + ST scalar_) { int x = get_global_id(0); int y = get_global_id(1); if (x < cols && y < rows) { - int src_index = mad24(y, src_step, mad24(x, (int)sizeof(T), src_offset)); - __global T * src = (__global T *)(srcptr + src_index); + int src_index = mad24(y, src_step, mad24(x, TSIZE, src_offset)); - src[0] = x == y ? scalar : (T)(0); + storepix(x == y ? scalar : (T)(0), srcptr + src_index); } } diff --git a/modules/core/src/persistence.cpp b/modules/core/src/persistence.cpp index 3755eccf7..88c7d6e8a 100644 --- a/modules/core/src/persistence.cpp +++ b/modules/core/src/persistence.cpp @@ -4824,7 +4824,7 @@ cvRegisterType( const CvTypeInfo* _info ) "Type name should contain only letters, digits, - and _" ); } - info = (CvTypeInfo*)malloc( sizeof(*info) + len + 1 ); + info = (CvTypeInfo*)cvAlloc( sizeof(*info) + len + 1 ); *info = *_info; info->type_name = (char*)(info + 1); @@ -4862,7 +4862,7 @@ cvUnregisterType( const char* type_name ) if( !CvType::first || !CvType::last ) CvType::first = CvType::last = 0; - free( info ); + cvFree( &info ); } } diff --git a/modules/cudabgsegm/src/cuda/mog2.cu b/modules/cudabgsegm/src/cuda/mog2.cu index de8df6c94..789afa47a 100644 --- a/modules/cudabgsegm/src/cuda/mog2.cu +++ b/modules/cudabgsegm/src/cuda/mog2.cu @@ -163,7 +163,7 @@ namespace cv { namespace cuda { namespace device { //need only weight if fit is found float weight = alpha1 * gmm_weight(mode * frame.rows + y, x) + prune; - + int swap_count = 0; //fit not found yet if (!fitsPDF) { @@ -214,6 +214,7 @@ namespace cv { namespace cuda { namespace device if (weight < gmm_weight((i - 1) * frame.rows + y, x)) break; + swap_count++; //swap one up swap(gmm_weight, x, y, i - 1, frame.rows); swap(gmm_variance, x, y, i - 1, frame.rows); @@ -231,7 +232,7 @@ namespace cv { namespace cuda { namespace device nmodes--; } - gmm_weight(mode * frame.rows + y, x) = weight; //update weight by the calculated value + gmm_weight((mode - swap_count) * frame.rows + y, x) = weight; //update weight by the calculated value totalWeight += weight; } diff --git a/modules/flann/include/opencv2/flann/kmeans_index.h b/modules/flann/include/opencv2/flann/kmeans_index.h index c3b1fd5b8..dc1781cc5 100644 --- a/modules/flann/include/opencv2/flann/kmeans_index.h +++ b/modules/flann/include/opencv2/flann/kmeans_index.h @@ -758,10 +758,13 @@ private: for (int k=0; k`_ page. + :param fourcc: 4-character code of codec used to compress the frames. For example, ``CV_FOURCC('P','I','M','1')`` is a MPEG-1 codec, ``CV_FOURCC('M','J','P','G')`` is a motion-jpeg codec etc. List of codes can be obtained at `Video Codecs by FOURCC `_ page. :param fps: Framerate of the created video stream. diff --git a/modules/highgui/include/opencv2/highgui/highgui_c.h b/modules/highgui/include/opencv2/highgui/highgui_c.h index ed8e2df0a..ec9b7853a 100644 --- a/modules/highgui/include/opencv2/highgui/highgui_c.h +++ b/modules/highgui/include/opencv2/highgui/highgui_c.h @@ -463,6 +463,8 @@ enum CV_CAP_PROP_ANDROID_FOCUS_DISTANCE_NEAR = 8006, CV_CAP_PROP_ANDROID_FOCUS_DISTANCE_OPTIMAL = 8007, CV_CAP_PROP_ANDROID_FOCUS_DISTANCE_FAR = 8008, + CV_CAP_PROP_ANDROID_EXPOSE_LOCK = 8009, + CV_CAP_PROP_ANDROID_WHITEBALANCE_LOCK = 8010, // Properties of cameras available through AVFOUNDATION interface CV_CAP_PROP_IOS_DEVICE_FOCUS = 9001, @@ -543,6 +545,7 @@ enum enum { CV_CAP_ANDROID_FOCUS_MODE_AUTO = 0, + CV_CAP_ANDROID_FOCUS_MODE_CONTINUOUS_PICTURE, CV_CAP_ANDROID_FOCUS_MODE_CONTINUOUS_VIDEO, CV_CAP_ANDROID_FOCUS_MODE_EDOF, CV_CAP_ANDROID_FOCUS_MODE_FIXED, diff --git a/modules/highgui/src/cap_android.cpp b/modules/highgui/src/cap_android.cpp index 082a12f00..dac245d2b 100644 --- a/modules/highgui/src/cap_android.cpp +++ b/modules/highgui/src/cap_android.cpp @@ -289,6 +289,10 @@ double CvCapture_Android::getProperty( int propIdx ) return (double)m_activity->getProperty(ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_OPTIMAL); case CV_CAP_PROP_ANDROID_FOCUS_DISTANCE_FAR: return (double)m_activity->getProperty(ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_FAR); + case CV_CAP_PROP_ANDROID_EXPOSE_LOCK: + return (double)m_activity->getProperty(ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK); + case CV_CAP_PROP_ANDROID_WHITEBALANCE_LOCK: + return (double)m_activity->getProperty(ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK); default: CV_Error( CV_StsOutOfRange, "Failed attempt to GET unsupported camera property." ); break; @@ -327,14 +331,23 @@ bool CvCapture_Android::setProperty( int propIdx, double propValue ) case CV_CAP_PROP_ANDROID_ANTIBANDING: m_activity->setProperty(ANDROID_CAMERA_PROPERTY_ANTIBANDING, propValue); break; + case CV_CAP_PROP_ANDROID_EXPOSE_LOCK: + m_activity->setProperty(ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK, propValue); + break; + case CV_CAP_PROP_ANDROID_WHITEBALANCE_LOCK: + m_activity->setProperty(ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK, propValue); + break; default: CV_Error( CV_StsOutOfRange, "Failed attempt to SET unsupported camera property." ); return false; } - if (propIdx != CV_CAP_PROP_AUTOGRAB) {// property for highgui class CvCapture_Android only + // Only changes in frame size require camera restart + if ((propIdx == CV_CAP_PROP_FRAME_WIDTH) || (propIdx == CV_CAP_PROP_FRAME_HEIGHT)) + { // property for highgui class CvCapture_Android only m_CameraParamsChanged = true; } + res = true; } diff --git a/modules/imgproc/perf/opencl/perf_filters.cpp b/modules/imgproc/perf/opencl/perf_filters.cpp index 57b928c28..f7329e319 100644 --- a/modules/imgproc/perf/opencl/perf_filters.cpp +++ b/modules/imgproc/perf/opencl/perf_filters.cpp @@ -211,7 +211,7 @@ OCL_PERF_TEST_P(SobelFixture, Sobel, OCL_TEST_CYCLE() cv::Sobel(src, dst, -1, dx, dy); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 1e-6); } ///////////// Scharr //////////////////////// diff --git a/modules/imgproc/src/canny.cpp b/modules/imgproc/src/canny.cpp index 766cb309c..fbc92dde2 100644 --- a/modules/imgproc/src/canny.cpp +++ b/modules/imgproc/src/canny.cpp @@ -100,19 +100,29 @@ static bool ocl_Canny(InputArray _src, OutputArray _dst, float low_thresh, float low_thresh = std::min(32767.0f, low_thresh); high_thresh = std::min(32767.0f, high_thresh); - if (low_thresh > 0) low_thresh *= low_thresh; - if (high_thresh > 0) high_thresh *= high_thresh; + if (low_thresh > 0) + low_thresh *= low_thresh; + if (high_thresh > 0) + high_thresh *= high_thresh; } int low = cvFloor(low_thresh), high = cvFloor(high_thresh); Size esize(size.width + 2, size.height + 2); UMat mag; - size_t globalsize[2] = { size.width * cn, size.height }, localsize[2] = { 16, 16 }; + size_t globalsize[2] = { size.width, size.height }, localsize[2] = { 16, 16 }; if (aperture_size == 3 && !_src.isSubmatrix()) { // Sobel calculation - ocl::Kernel calcSobelRowPassKernel("calcSobelRowPass", ocl::imgproc::canny_oclsrc); + char cvt[2][40]; + ocl::Kernel calcSobelRowPassKernel("calcSobelRowPass", ocl::imgproc::canny_oclsrc, + format("-D OP_SOBEL -D cn=%d -D shortT=%s -D ucharT=%s" + " -D convertToIntT=%s -D intT=%s -D convertToShortT=%s", cn, + ocl::typeToStr(CV_16SC(cn)), + ocl::typeToStr(CV_8UC(cn)), + ocl::convertTypeStr(CV_8U, CV_32S, cn, cvt[0]), + ocl::typeToStr(CV_32SC(cn)), + ocl::convertTypeStr(CV_32S, CV_16S, cn, cvt[1]))); if (calcSobelRowPassKernel.empty()) return false; @@ -126,58 +136,62 @@ static bool ocl_Canny(InputArray _src, OutputArray _dst, float low_thresh, float // magnitude calculation ocl::Kernel magnitudeKernel("calcMagnitude_buf", ocl::imgproc::canny_oclsrc, - L2gradient ? " -D L2GRAD" : ""); + format("-D cn=%d%s -D OP_MAG_BUF -D shortT=%s -D convertToIntT=%s -D intT=%s", + cn, L2gradient ? " -D L2GRAD" : "", + ocl::typeToStr(CV_16SC(cn)), + ocl::convertTypeStr(CV_16S, CV_32S, cn, cvt[0]), + ocl::typeToStr(CV_32SC(cn)))); if (magnitudeKernel.empty()) return false; - mag = UMat(esize, CV_32SC(cn), Scalar::all(0)); + mag = UMat(esize, CV_32SC1, Scalar::all(0)); dx.create(size, CV_16SC(cn)); dy.create(size, CV_16SC(cn)); magnitudeKernel.args(ocl::KernelArg::ReadOnlyNoSize(dxBuf), ocl::KernelArg::ReadOnlyNoSize(dyBuf), ocl::KernelArg::WriteOnlyNoSize(dx), ocl::KernelArg::WriteOnlyNoSize(dy), - ocl::KernelArg::WriteOnlyNoSize(mag, cn), size.height, size.width); + ocl::KernelArg::WriteOnlyNoSize(mag), size.height, size.width); if (!magnitudeKernel.run(2, globalsize, localsize, false)) return false; } else { - dx.create(size, CV_16SC(cn)); - dy.create(size, CV_16SC(cn)); - - Sobel(_src, dx, CV_16SC1, 1, 0, aperture_size, 1, 0, BORDER_REPLICATE); - Sobel(_src, dy, CV_16SC1, 0, 1, aperture_size, 1, 0, BORDER_REPLICATE); + Sobel(_src, dx, CV_16S, 1, 0, aperture_size, 1, 0, BORDER_REPLICATE); + Sobel(_src, dy, CV_16S, 0, 1, aperture_size, 1, 0, BORDER_REPLICATE); // magnitude calculation ocl::Kernel magnitudeKernel("calcMagnitude", ocl::imgproc::canny_oclsrc, - L2gradient ? " -D L2GRAD" : ""); + format("-D OP_MAG -D cn=%d%s -D intT=int -D shortT=short -D convertToIntT=convert_int_sat", + cn, L2gradient ? " -D L2GRAD" : "")); if (magnitudeKernel.empty()) return false; - mag = UMat(esize, CV_32SC(cn), Scalar::all(0)); + mag = UMat(esize, CV_32SC1, Scalar::all(0)); magnitudeKernel.args(ocl::KernelArg::ReadOnlyNoSize(dx), ocl::KernelArg::ReadOnlyNoSize(dy), - ocl::KernelArg::WriteOnlyNoSize(mag, cn), size.height, size.width); + ocl::KernelArg::WriteOnlyNoSize(mag), size.height, size.width); if (!magnitudeKernel.run(2, globalsize, NULL, false)) return false; } // map calculation - ocl::Kernel calcMapKernel("calcMap", ocl::imgproc::canny_oclsrc); + ocl::Kernel calcMapKernel("calcMap", ocl::imgproc::canny_oclsrc, + format("-D OP_MAP -D cn=%d", cn)); if (calcMapKernel.empty()) return false; - UMat map(esize, CV_32SC(cn)); + UMat map(esize, CV_32SC1); calcMapKernel.args(ocl::KernelArg::ReadOnlyNoSize(dx), ocl::KernelArg::ReadOnlyNoSize(dy), - ocl::KernelArg::ReadOnlyNoSize(mag), ocl::KernelArg::WriteOnlyNoSize(map, cn), + ocl::KernelArg::ReadOnlyNoSize(mag), ocl::KernelArg::WriteOnlyNoSize(map), size.height, size.width, low, high); if (!calcMapKernel.run(2, globalsize, localsize, false)) return false; // local hysteresis thresholding - ocl::Kernel edgesHysteresisLocalKernel("edgesHysteresisLocal", ocl::imgproc::canny_oclsrc); + ocl::Kernel edgesHysteresisLocalKernel("edgesHysteresisLocal", ocl::imgproc::canny_oclsrc, + "-D OP_HYST_LOCAL"); if (edgesHysteresisLocalKernel.empty()) return false; @@ -193,7 +207,8 @@ static bool ocl_Canny(InputArray _src, OutputArray _dst, float low_thresh, float for ( ; ; ) { - ocl::Kernel edgesHysteresisGlobalKernel("edgesHysteresisGlobal", ocl::imgproc::canny_oclsrc); + ocl::Kernel edgesHysteresisGlobalKernel("edgesHysteresisGlobal", ocl::imgproc::canny_oclsrc, + "-D OP_HYST_GLOBAL"); if (edgesHysteresisGlobalKernel.empty()) return false; @@ -221,14 +236,15 @@ static bool ocl_Canny(InputArray _src, OutputArray _dst, float low_thresh, float } // get edges - ocl::Kernel getEdgesKernel("getEdges", ocl::imgproc::canny_oclsrc); + ocl::Kernel getEdgesKernel("getEdges", ocl::imgproc::canny_oclsrc, "-D OP_EDGES"); if (getEdgesKernel.empty()) return false; - _dst.create(size, CV_8UC(cn)); + _dst.create(size, CV_8UC1); UMat dst = _dst.getUMat(); getEdgesKernel.args(ocl::KernelArg::ReadOnlyNoSize(map), ocl::KernelArg::WriteOnly(dst)); + return getEdgesKernel.run(2, globalsize, NULL, false); } @@ -254,12 +270,12 @@ void cv::Canny( InputArray _src, OutputArray _dst, } if ((aperture_size & 1) == 0 || (aperture_size != -1 && (aperture_size < 3 || aperture_size > 7))) - CV_Error(CV_StsBadFlag, ""); + CV_Error(CV_StsBadFlag, "Aperture size should be odd"); if (low_thresh > high_thresh) std::swap(low_thresh, high_thresh); - CV_OCL_RUN(_dst.isUMat() && cn == 1, + CV_OCL_RUN(_dst.isUMat() && (cn == 1 || cn == 3), ocl_Canny(_src, _dst, (float)low_thresh, (float)high_thresh, aperture_size, L2gradient, cn, size)) Mat src = _src.getMat(), dst = _dst.getMat(); diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index ea0baf6b0..c49f23efa 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -3134,7 +3134,7 @@ template struct Filter2D : public BaseFi // b e h b e h 0 0 // c f i c f i 0 0 template -static int _prepareKernelFilter2D(std::vector& data, const Mat &kernel) +static int _prepareKernelFilter2D(std::vector & data, const Mat & kernel) { Mat _kernel; kernel.convertTo(_kernel, DataDepth::value); int size_y_aligned = ROUNDUP(kernel.rows * 2, 4); @@ -3317,200 +3317,224 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, return kernel.run(2, globalsize, localsize, true); } -static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, int borderType, bool sync) +static bool ocl_sepRowFilter2D(const UMat & src, UMat & buf, const Mat & kernelX, int anchor, + int borderType, int ddepth, bool fast8uc1) { - int type = src.type(); - int cn = CV_MAT_CN(type); - int sdepth = CV_MAT_DEPTH(type); + int type = src.type(), cn = CV_MAT_CN(type), sdepth = CV_MAT_DEPTH(type); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; Size bufSize = buf.size(); + if (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F)) + return false; + #ifdef ANDROID size_t localsize[2] = {16, 10}; #else size_t localsize[2] = {16, 16}; #endif + size_t globalsize[2] = {DIVUP(bufSize.width, localsize[0]) * localsize[0], DIVUP(bufSize.height, localsize[1]) * localsize[1]}; - if (CV_8U == sdepth) - { - switch (cn) - { - case 1: - globalsize[0] = DIVUP((bufSize.width + 3) >> 2, localsize[0]) * localsize[0]; - break; - case 2: - globalsize[0] = DIVUP((bufSize.width + 1) >> 1, localsize[0]) * localsize[0]; - break; - case 4: - globalsize[0] = DIVUP(bufSize.width, localsize[0]) * localsize[0]; - break; - } - } + if (fast8uc1) + globalsize[0] = DIVUP((bufSize.width + 3) >> 2, localsize[0]) * localsize[0]; - int radiusX = anchor; - int radiusY = (int)((buf.rows - src.rows) >> 1); + int radiusX = anchor, radiusY = (buf.rows - src.rows) >> 1; - bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0; - const char* btype = NULL; - switch (borderType & ~BORDER_ISOLATED) - { - case BORDER_CONSTANT: - btype = "BORDER_CONSTANT"; - break; - case BORDER_REPLICATE: - btype = "BORDER_REPLICATE"; - break; - case BORDER_REFLECT: - btype = "BORDER_REFLECT"; - break; - case BORDER_WRAP: - btype = "BORDER_WRAP"; - break; - case BORDER_REFLECT101: - btype = "BORDER_REFLECT_101"; - break; - default: - return false; - } + bool isolated = (borderType & BORDER_ISOLATED) != 0; + const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" }, + * const btype = borderMap[borderType & ~BORDER_ISOLATED]; bool extra_extrapolation = src.rows < (int)((-radiusY + globalsize[1]) >> 1) + 1; extra_extrapolation |= src.rows < radiusY; extra_extrapolation |= src.cols < (int)((-radiusX + globalsize[0] + 8 * localsize[0] + 3) >> 1) + 1; extra_extrapolation |= src.cols < radiusX; - cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s", - radiusX, (int)localsize[0], (int)localsize[1], cn, - btype, - extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION", - isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED"); + char cvt[40]; + cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s" + " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s%s", + radiusX, (int)localsize[0], (int)localsize[1], cn, btype, + extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION", + isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED", + ocl::typeToStr(type), ocl::typeToStr(CV_32FC(cn)), + ocl::convertTypeStr(sdepth, CV_32F, cn, cvt), + ocl::typeToStr(sdepth), ocl::typeToStr(CV_32F), + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); build_options += ocl::kernelToStr(kernelX, CV_32F); Size srcWholeSize; Point srcOffset; src.locateROI(srcWholeSize, srcOffset); - std::stringstream strKernel; - strKernel << "row_filter"; - if (-1 != cn) - strKernel << "_C" << cn; - if (-1 != sdepth) - strKernel << "_D" << sdepth; + String kernelName("row_filter"); + if (fast8uc1) + kernelName += "_C1_D0"; - ocl::Kernel kernelRow; - if (!kernelRow.create(strKernel.str().c_str(), cv::ocl::imgproc::filterSepRow_oclsrc, - build_options)) + ocl::Kernel k(kernelName.c_str(), cv::ocl::imgproc::filterSepRow_oclsrc, + build_options); + if (k.empty()) return false; - int idxArg = 0; - idxArg = kernelRow.set(idxArg, ocl::KernelArg::PtrReadOnly(src)); - idxArg = kernelRow.set(idxArg, (int)(src.step / src.elemSize())); + if (fast8uc1) + k.args(ocl::KernelArg::PtrReadOnly(src), (int)(src.step / src.elemSize()), srcOffset.x, + srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height, + ocl::KernelArg::PtrWriteOnly(buf), (int)(buf.step / buf.elemSize()), + buf.cols, buf.rows, radiusY); + else + k.args(ocl::KernelArg::PtrReadOnly(src), (int)src.step, srcOffset.x, + srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height, + ocl::KernelArg::PtrWriteOnly(buf), (int)buf.step, buf.cols, buf.rows, radiusY); - idxArg = kernelRow.set(idxArg, srcOffset.x); - idxArg = kernelRow.set(idxArg, srcOffset.y); - idxArg = kernelRow.set(idxArg, src.cols); - idxArg = kernelRow.set(idxArg, src.rows); - idxArg = kernelRow.set(idxArg, srcWholeSize.width); - idxArg = kernelRow.set(idxArg, srcWholeSize.height); - - idxArg = kernelRow.set(idxArg, ocl::KernelArg::PtrWriteOnly(buf)); - idxArg = kernelRow.set(idxArg, (int)(buf.step / buf.elemSize())); - idxArg = kernelRow.set(idxArg, buf.cols); - idxArg = kernelRow.set(idxArg, buf.rows); - idxArg = kernelRow.set(idxArg, radiusY); - - return kernelRow.run(2, globalsize, localsize, sync); + return k.run(2, globalsize, localsize, false); } -static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anchor, bool sync) +static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY, double delta, int anchor) { + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + if (dst.depth() == CV_64F && !doubleSupport) + return false; + #ifdef ANDROID - size_t localsize[2] = {16, 10}; + size_t localsize[2] = { 16, 10 }; #else - size_t localsize[2] = {16, 16}; + size_t localsize[2] = { 16, 16 }; #endif - size_t globalsize[2] = {0, 0}; + size_t globalsize[2] = { 0, 0 }; int dtype = dst.type(), cn = CV_MAT_CN(dtype), ddepth = CV_MAT_DEPTH(dtype); Size sz = dst.size(); globalsize[1] = DIVUP(sz.height, localsize[1]) * localsize[1]; - - if (dtype == CV_8UC2) - globalsize[0] = DIVUP((sz.width + 1) / 2, localsize[0]) * localsize[0]; - else - globalsize[0] = DIVUP(sz.width, localsize[0]) * localsize[0]; + globalsize[0] = DIVUP(sz.width, localsize[0]) * localsize[0]; char cvt[40]; - cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D GENTYPE_SRC=%s -D GENTYPE_DST=%s -D convert_to_DST=%s", - anchor, (int)localsize[0], (int)localsize[1], cn, ocl::typeToStr(buf.type()), - ocl::typeToStr(dtype), ocl::convertTypeStr(CV_32F, ddepth, cn, cvt)); + cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d" + " -D srcT=%s -D dstT=%s -D convertToDstT=%s" + " -D srcT1=%s -D dstT1=%s%s", + anchor, (int)localsize[0], (int)localsize[1], cn, + ocl::typeToStr(buf.type()), ocl::typeToStr(dtype), + ocl::convertTypeStr(CV_32F, ddepth, cn, cvt), + ocl::typeToStr(CV_32F), ocl::typeToStr(ddepth), + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); build_options += ocl::kernelToStr(kernelY, CV_32F); - ocl::Kernel kernelCol; - if (!kernelCol.create("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc, build_options)) + ocl::Kernel k("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc, + build_options); + if (k.empty()) return false; - int idxArg = 0; - idxArg = kernelCol.set(idxArg, ocl::KernelArg::PtrReadOnly(buf)); - idxArg = kernelCol.set(idxArg, (int)(buf.step / buf.elemSize())); - idxArg = kernelCol.set(idxArg, buf.cols); - idxArg = kernelCol.set(idxArg, buf.rows); + k.args(ocl::KernelArg::ReadOnly(buf), ocl::KernelArg::WriteOnly(dst), + static_cast(delta)); - idxArg = kernelCol.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst)); - idxArg = kernelCol.set(idxArg, (int)(dst.offset / dst.elemSize())); - idxArg = kernelCol.set(idxArg, (int)(dst.step / dst.elemSize())); - idxArg = kernelCol.set(idxArg, dst.cols); - idxArg = kernelCol.set(idxArg, dst.rows); + return k.run(2, globalsize, localsize, false); +} - return kernelCol.run(2, globalsize, localsize, sync); +const int optimizedSepFilterLocalSize = 16; + +static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, + Mat row_kernel, Mat col_kernel, + double delta, int borderType, int ddepth) +{ + Size size = _src.size(), wholeSize; + Point origin; + int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype), + esz = CV_ELEM_SIZE(stype), wdepth = std::max(std::max(sdepth, ddepth), CV_32F), + dtype = CV_MAKE_TYPE(ddepth, cn); + size_t src_step = _src.step(), src_offset = _src.offset(); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + + if ((src_offset % src_step) % esz != 0 || (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F)) || + !(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE || + borderType == BORDER_REFLECT || borderType == BORDER_WRAP || + borderType == BORDER_REFLECT_101)) + return false; + + size_t lt2[2] = { optimizedSepFilterLocalSize, optimizedSepFilterLocalSize }; + size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), lt2[1] * (1 + (size.height - 1) / lt2[1]) }; + + char cvt[2][40]; + const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", + "BORDER_REFLECT_101" }; + + String opts = cv::format("-D BLK_X=%d -D BLK_Y=%d -D RADIUSX=%d -D RADIUSY=%d%s%s" + " -D srcT=%s -D convertToWT=%s -D WT=%s -D dstT=%s -D convertToDstT=%s" + " -D %s -D srcT1=%s -D dstT1=%s -D CN=%d", (int)lt2[0], (int)lt2[1], + row_kernel.cols / 2, col_kernel.cols / 2, + ocl::kernelToStr(row_kernel, CV_32F, "KERNEL_MATRIX_X").c_str(), + ocl::kernelToStr(col_kernel, CV_32F, "KERNEL_MATRIX_Y").c_str(), + ocl::typeToStr(stype), ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]), + ocl::typeToStr(CV_MAKE_TYPE(wdepth, cn)), ocl::typeToStr(dtype), + ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]), borderMap[borderType], + ocl::typeToStr(sdepth), ocl::typeToStr(ddepth), cn); + + ocl::Kernel k("sep_filter", ocl::imgproc::filterSep_singlePass_oclsrc, opts); + if (k.empty()) + return false; + + UMat src = _src.getUMat(); + _dst.create(size, dtype); + UMat dst = _dst.getUMat(); + + int src_offset_x = static_cast((src_offset % src_step) / esz); + int src_offset_y = static_cast(src_offset / src_step); + + src.locateROI(wholeSize, origin); + + k.args(ocl::KernelArg::PtrReadOnly(src), (int)src_step, src_offset_x, src_offset_y, + wholeSize.height, wholeSize.width, ocl::KernelArg::WriteOnly(dst), + static_cast(delta)); + + return k.run(2, gt2, lt2, false); } static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, InputArray _kernelX, InputArray _kernelY, Point anchor, double delta, int borderType ) { - if (abs(delta)> FLT_MIN) - return false; + const ocl::Device & d = ocl::Device::getDefault(); + Size imgSize = _src.size(); - int type = _src.type(); - if ( !( (type == CV_8UC1 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC4) && - (ddepth == CV_32F || ddepth == CV_16S || ddepth == CV_8U || ddepth < 0) ) ) + int type = _src.type(), sdepth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + if (cn > 4) return false; - int cn = CV_MAT_CN(type); - Mat kernelX = _kernelX.getMat().reshape(1, 1); - if (1 != (kernelX.cols % 2)) + if (kernelX.cols % 2 != 1) return false; Mat kernelY = _kernelY.getMat().reshape(1, 1); - if (1 != (kernelY.cols % 2)) + if (kernelY.cols % 2 != 1) return false; - int sdepth = CV_MAT_DEPTH(type); - if( anchor.x < 0 ) - anchor.x = kernelX.cols >> 1; - if( anchor.y < 0 ) - anchor.y = kernelY.cols >> 1; - - if( ddepth < 0 ) + if (ddepth < 0) ddepth = sdepth; + CV_OCL_RUN_(kernelY.cols <= 21 && kernelX.cols <= 21 && + imgSize.width > optimizedSepFilterLocalSize + (kernelX.cols >> 1) && + imgSize.height > optimizedSepFilterLocalSize + (kernelY.cols >> 1) && + (!(borderType & BORDER_ISOLATED) || _src.offset() == 0) && anchor == Point(-1, -1) && + (d.isIntel() || (d.isAMD() && !d.hostUnifiedMemory())), + ocl_sepFilter2D_SinglePass(_src, _dst, kernelX, kernelY, delta, + borderType & ~BORDER_ISOLATED, ddepth), true) + + if (anchor.x < 0) + anchor.x = kernelX.cols >> 1; + if (anchor.y < 0) + anchor.y = kernelY.cols >> 1; + UMat src = _src.getUMat(); Size srcWholeSize; Point srcOffset; src.locateROI(srcWholeSize, srcOffset); - if ( (0 != (srcOffset.x % 4)) || - (0 != (src.cols % 4)) || - (0 != ((src.step / src.elemSize()) % 4)) - ) - return false; + + bool fast8uc1 = type == CV_8UC1 && srcOffset.x % 4 == 0 && + src.cols % 4 == 0 && src.step % 4 == 0; Size srcSize = src.size(); Size bufSize(srcSize.width, srcSize.height + kernelY.cols - 1); - UMat buf; buf.create(bufSize, CV_MAKETYPE(CV_32F, cn)); - if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType, false)) + UMat buf(bufSize, CV_32FC(cn)); + if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType, ddepth, fast8uc1)) return false; _dst.create(srcSize, CV_MAKETYPE(ddepth, cn)); UMat dst = _dst.getUMat(); - return ocl_sepColFilter2D(buf, dst, kernelY, anchor.y, false); + + return ocl_sepColFilter2D(buf, dst, kernelY, delta, anchor.y); } #endif diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index ac958fc69..1dd0a252e 100644 --- a/modules/imgproc/src/morph.cpp +++ b/modules/imgproc/src/morph.cpp @@ -42,7 +42,6 @@ #include "precomp.hpp" #include -#include #include "opencl_kernels.hpp" /****************************************************************************************\ @@ -1291,9 +1290,10 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, { CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE); + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; - if (_src.depth() == CV_64F && !doubleSupport) + if (depth == CV_64F && !doubleSupport) return false; UMat kernel8U; @@ -1324,13 +1324,14 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, return false; static const char * const op2str[] = { "ERODE", "DILATE" }; - String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s -D GENTYPE=%s -D DEPTH_%d", - anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], op2str[op], + String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s" + " -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s", anchor.x, anchor.y, + (int)localThreads[0], (int)localThreads[1], op2str[op], doubleSupport ? " -D DOUBLE_SUPPORT" : "", rectKernel ? " -D RECTKERNEL" : "", - ocl::typeToStr(_src.type()), _src.depth() ); + ocl::typeToStr(_src.type()), _src.depth(), cn, ocl::typeToStr(depth)); std::vector kernels; - for (int i = 0; i= 0 ? x : -x) + (y >= 0 ? y : -y); + return convertToIntT( (x >= (shortT)(0) ? x : -x) + (y >= (shortT)(0) ? y : -y) ); #endif } +#ifdef OP_MAG + // calculate the magnitude of the filter pass combining both x and y directions // This is the non-buffered version(non-3x3 sobel) // @@ -112,18 +127,43 @@ __kernel void calcMagnitude(__global const uchar * dxptr, int dx_step, int dx_of if (y < rows && x < cols) { - int dx_index = mad24(dx_step, y, x * (int)sizeof(short) + dx_offset); - int dy_index = mad24(dy_step, y, x * (int)sizeof(short) + dy_offset); - int mag_index = mad24(mag_step, y + 1, (x + 1) * (int)sizeof(int) + mag_offset); + int dx_index = mad24(dx_step, y, mad24(x, (int)sizeof(short) * cn, dx_offset)); + int dy_index = mad24(dy_step, y, mad24(x, (int)sizeof(short) * cn, dy_offset)); + int mag_index = mad24(mag_step, y + 1, mad24(x + 1, (int)sizeof(int), mag_offset)); - __global const short * dx = (__global const short *)(dxptr + dx_index); - __global const short * dy = (__global const short *)(dyptr + dy_index); + __global short * dx = (__global short *)(dxptr + dx_index); + __global short * dy = (__global short *)(dyptr + dy_index); __global int * mag = (__global int *)(magptr + mag_index); - mag[0] = calc(dx[0], dy[0]); + int cmag = calc(dx[0], dy[0]); +#if cn > 1 + short cx = dx[0], cy = dy[0]; + int pmag; + + #pragma unroll + for (int i = 1; i < cn; ++i) + { + pmag = calc(dx[i], dy[i]); + if (pmag > cmag) + cmag = pmag, cx = dx[i], cy = dy[i]; + } + + dx[0] = cx, dy[0] = cy; +#endif + mag[0] = cmag; } } +#elif defined OP_MAG_BUF + +#if cn != 3 +#define loadpix(addr) *(__global const shortT *)(addr) +#define shortSize (int)sizeof(shortT) +#else +#define loadpix(addr) vload3(0, (__global const short *)(addr)) +#define shortSize (int)sizeof(short)*cn +#endif + // calculate the magnitude of the filter pass combining both x and y directions // This is the buffered version(3x3 sobel) // @@ -132,59 +172,64 @@ __kernel void calcMagnitude(__global const uchar * dxptr, int dx_step, int dx_of // dx direvitive in x direction output // dy direvitive in y direction output // mag magnitude direvitive of xy output -__kernel void __attribute__((reqd_work_group_size(16, 16, 1))) -calcMagnitude_buf - (__global const short * dx_buf, int dx_buf_step, int dx_buf_offset, - __global const short * dy_buf, int dy_buf_step, int dy_buf_offset, - __global short * dx, int dx_step, int dx_offset, - __global short * dy, int dy_step, int dy_offset, - __global int * mag, int mag_step, int mag_offset, - int rows, int cols) +__kernel void calcMagnitude_buf(__global const uchar * dx_buf, int dx_buf_step, int dx_buf_offset, + __global const uchar * dy_buf, int dy_buf_step, int dy_buf_offset, + __global uchar * dx, int dx_step, int dx_offset, + __global uchar * dy, int dy_step, int dy_offset, + __global uchar * mag, int mag_step, int mag_offset, int rows, int cols) { - dx_buf_step /= sizeof(*dx_buf); - dx_buf_offset /= sizeof(*dx_buf); - dy_buf_step /= sizeof(*dy_buf); - dy_buf_offset /= sizeof(*dy_buf); - dx_step /= sizeof(*dx); - dx_offset /= sizeof(*dx); - dy_step /= sizeof(*dy); - dy_offset /= sizeof(*dy); - mag_step /= sizeof(*mag); - mag_offset /= sizeof(*mag); - int gidx = get_global_id(0); int gidy = get_global_id(1); int lidx = get_local_id(0); int lidy = get_local_id(1); - __local short sdx[18][16]; - __local short sdy[18][16]; + __local shortT sdx[18][16]; + __local shortT sdy[18][16]; - sdx[lidy + 1][lidx] = dx_buf[gidx + min(gidy, rows - 1) * dx_buf_step + dx_buf_offset]; - sdy[lidy + 1][lidx] = dy_buf[gidx + min(gidy, rows - 1) * dy_buf_step + dy_buf_offset]; + sdx[lidy + 1][lidx] = loadpix(dx_buf + mad24(min(gidy, rows - 1), dx_buf_step, mad24(gidx, shortSize, dx_buf_offset))); + sdy[lidy + 1][lidx] = loadpix(dy_buf + mad24(min(gidy, rows - 1), dy_buf_step, mad24(gidx, shortSize, dy_buf_offset))); if (lidy == 0) { - sdx[0][lidx] = dx_buf[gidx + min(max(gidy - 1, 0), rows - 1) * dx_buf_step + dx_buf_offset]; - sdx[17][lidx] = dx_buf[gidx + min(gidy + 16, rows - 1) * dx_buf_step + dx_buf_offset]; + sdx[0][lidx] = loadpix(dx_buf + mad24(clamp(gidy - 1, 0, rows - 1), dx_buf_step, mad24(gidx, shortSize, dx_buf_offset))); + sdx[17][lidx] = loadpix(dx_buf + mad24(min(gidy + 16, rows - 1), dx_buf_step, mad24(gidx, shortSize, dx_buf_offset))); - sdy[0][lidx] = dy_buf[gidx + min(max(gidy - 1, 0), rows - 1) * dy_buf_step + dy_buf_offset]; - sdy[17][lidx] = dy_buf[gidx + min(gidy + 16, rows - 1) * dy_buf_step + dy_buf_offset]; + sdy[0][lidx] = loadpix(dy_buf + mad24(clamp(gidy - 1, 0, rows - 1), dy_buf_step, mad24(gidx, shortSize, dy_buf_offset))); + sdy[17][lidx] = loadpix(dy_buf + mad24(min(gidy + 16, rows - 1), dy_buf_step, mad24(gidx, shortSize, dy_buf_offset))); } barrier(CLK_LOCAL_MEM_FENCE); if (gidx < cols && gidy < rows) { - short x = sdx[lidy][lidx] + 2 * sdx[lidy + 1][lidx] + sdx[lidy + 2][lidx]; - short y = -sdy[lidy][lidx] + sdy[lidy + 2][lidx]; + shortT x = sdx[lidy + 1][lidx] * (shortT)(2) + sdx[lidy][lidx] + sdx[lidy + 2][lidx]; + shortT y = -sdy[lidy][lidx] + sdy[lidy + 2][lidx]; - dx[gidx + gidy * dx_step + dx_offset] = x; - dy[gidx + gidy * dy_step + dy_offset] = y; +#if cn == 1 + *(__global short *)(dx + mad24(gidy, dx_step, mad24(gidx, shortSize, dx_offset))) = x; + *(__global short *)(dy + mad24(gidy, dy_step, mad24(gidx, shortSize, dy_offset))) = y; - mag[(gidx + 1) + (gidy + 1) * mag_step + mag_offset] = calc(x, y); + *(__global int *)(mag + mad24(gidy + 1, mag_step, mad24(gidx + 1, (int)sizeof(int), mag_offset))) = calc(x, y); +#elif cn == 3 + intT magv = calc(x, y); + short cx = x.x, cy = y.x; + int cmag = magv.x; + + if (cmag < magv.y) + cx = x.y, cy = y.y, cmag = magv.y; + if (cmag < magv.z) + cx = x.z, cy = y.z, cmag = magv.z; + + *(__global short *)(dx + mad24(gidy, dx_step, mad24(gidx, shortSize, dx_offset))) = cx; + *(__global short *)(dy + mad24(gidy, dy_step, mad24(gidx, shortSize, dy_offset))) = cy; + + *(__global int *)(mag + mad24(gidy + 1, mag_step, mad24(gidx + 1, (int)sizeof(int), mag_offset))) = cmag; +#endif } } +#endif + +#elif defined OP_MAP ////////////////////////////////////////////////////////////////////////////////////////// // 0.4142135623730950488016887242097 is tan(22.5) @@ -208,13 +253,11 @@ calcMagnitude_buf // mag magnitudes calculated from calcMagnitude function // map output containing raw edge types -__kernel void __attribute__((reqd_work_group_size(16,16,1))) -calcMap( - __global const uchar * dx, int dx_step, int dx_offset, - __global const uchar * dy, int dy_step, int dy_offset, - __global const uchar * mag, int mag_step, int mag_offset, - __global uchar * map, int map_step, int map_offset, - int rows, int cols, int low_thresh, int high_thresh) +__kernel void calcMap(__global const uchar * dx, int dx_step, int dx_offset, + __global const uchar * dy, int dy_step, int dy_offset, + __global const uchar * mag, int mag_step, int mag_offset, + __global uchar * map, int map_step, int map_offset, + int rows, int cols, int low_thresh, int high_thresh) { __local int smem[18][18]; @@ -227,7 +270,7 @@ calcMap( int grp_idx = get_global_id(0) & 0xFFFFF0; int grp_idy = get_global_id(1) & 0xFFFFF0; - int tid = lidx + lidy * 16; + int tid = mad24(lidy, 16, lidx); int lx = tid % 18; int ly = tid / 18; @@ -250,8 +293,8 @@ calcMap( if (m > low_thresh) { - short xs = *(__global const short *)(dx + mad24(gidy, dx_step, dx_offset + (int)sizeof(short) * gidx)); - short ys = *(__global const short *)(dy + mad24(gidy, dy_step, dy_offset + (int)sizeof(short) * gidx)); + short xs = *(__global const short *)(dx + mad24(gidy, dx_step, mad24(gidx, (int)sizeof(short) * cn, dx_offset))); + short ys = *(__global const short *)(dy + mad24(gidy, dy_step, mad24(gidx, (int)sizeof(short) * cn, dy_offset))); int x = abs(xs), y = abs(ys); int tg22x = x * TG22; @@ -278,13 +321,15 @@ calcMap( } } } - *(__global int *)(map + mad24(map_step, gidy + 1, (gidx + 1) * (int)sizeof(int) + map_offset)) = edge_type; + *(__global int *)(map + mad24(map_step, gidy + 1, mad24(gidx + 1, (int)sizeof(int), + map_offset))) = edge_type; } } #undef CANNY_SHIFT #undef TG22 +#elif defined OP_HYST_LOCAL + struct PtrStepSz { __global uchar * ptr; @@ -312,11 +357,9 @@ inline void set(struct PtrStepSz data, int y, int x, int value) // stack the potiential edge points found in this kernel call // counter the number of potiential edge points -__kernel void __attribute__((reqd_work_group_size(16,16,1))) -edgesHysteresisLocal - (__global uchar * map_ptr, int map_step, int map_offset, - __global ushort2 * st, __global unsigned int * counter, - int rows, int cols) +__kernel void edgesHysteresisLocal(__global uchar * map_ptr, int map_step, int map_offset, + __global ushort2 * st, __global unsigned int * counter, + int rows, int cols) { struct PtrStepSz map = { map_ptr + map_offset, map_step, rows + 1, cols + 1 }; @@ -402,6 +445,8 @@ edgesHysteresisLocal } } +#elif defined OP_HYST_GLOBAL + __constant int c_dx[8] = {-1, 0, 1, -1, 1, -1, 0, 1}; __constant int c_dy[8] = {-1, -1, -1, 0, 0, 1, 1, 1}; @@ -409,10 +454,9 @@ __constant int c_dy[8] = {-1, -1, -1, 0, 0, 1, 1, 1}; #define stack_size 512 #define map_index mad24(map_step, pos.y, pos.x * (int)sizeof(int)) -__kernel void __attribute__((reqd_work_group_size(128, 1, 1))) -edgesHysteresisGlobal(__global uchar * map, int map_step, int map_offset, - __global ushort2 * st1, __global ushort2 * st2, __global int * counter, - int rows, int cols, int count) +__kernel void edgesHysteresisGlobal(__global uchar * map, int map_step, int map_offset, + __global ushort2 * st1, __global ushort2 * st2, __global int * counter, + int rows, int cols, int count) { map += map_offset; @@ -492,6 +536,8 @@ edgesHysteresisGlobal(__global uchar * map, int map_step, int map_offset, #undef map_index #undef stack_size +#elif defined OP_EDGES + // Get the edge result. egde type of value 2 will be marked as an edge point and set to 255. Otherwise 0. // map edge type mappings // dst edge output @@ -504,7 +550,7 @@ __kernel void getEdges(__global const uchar * mapptr, int map_step, int map_offs if (y < rows && x < cols) { - int map_index = mad24(map_step, y + 1, (x + 1) * (int)sizeof(int) + map_offset); + int map_index = mad24(map_step, y + 1, mad24(x + 1, (int)sizeof(int), map_offset)); int dst_index = mad24(dst_step, y, x + dst_offset); __global const int * map = (__global const int *)(mapptr + map_index); @@ -512,3 +558,5 @@ __kernel void getEdges(__global const uchar * mapptr, int map_step, int map_offs dst[dst_index] = (uchar)(-(map[0] >> 1)); } } + +#endif diff --git a/modules/imgproc/src/opencl/filterSepCol.cl b/modules/imgproc/src/opencl/filterSepCol.cl index 30a2221cf..29514cc21 100644 --- a/modules/imgproc/src/opencl/filterSepCol.cl +++ b/modules/imgproc/src/opencl/filterSepCol.cl @@ -34,47 +34,36 @@ // // +#ifdef DOUBLE_SUPPORT +#ifdef cl_amd_fp64 +#pragma OPENCL EXTENSION cl_amd_fp64:enable +#elif defined (cl_khr_fp64) +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#endif +#endif + #define READ_TIMES_COL ((2*(RADIUSY+LSIZE1)-1)/LSIZE1) #define RADIUS 1 -#if CN ==1 -#define ALIGN (((RADIUS)+3)>>2<<2) -#elif CN==2 -#define ALIGN (((RADIUS)+1)>>1<<1) -#elif CN==3 -#define ALIGN (((RADIUS)+3)>>2<<2) -#elif CN==4 -#define ALIGN (RADIUS) -#define READ_TIMES_ROW ((2*(RADIUS+LSIZE0)-1)/LSIZE0) -#endif #define noconvert -/********************************************************************************** -These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur. -Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle -kernel must be in the center. ROI is not supported either. -Each kernels read 4 elements(not 4 pixels), save them to LDS and read the data needed -from LDS to calculate the result. -The length of the convovle kernel supported is only related to the MAX size of LDS, -which is HW related. -Niko -6/29/2011 -The info above maybe obsolete. -***********************************************************************************/ +#if CN != 3 +#define loadpix(addr) *(__global const srcT *)(addr) +#define storepix(val, addr) *(__global dstT *)(addr) = val +#define SRCSIZE (int)sizeof(srcT) +#define DSTSIZE (int)sizeof(dstT) +#else +#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr)) +#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr)) +#define SRCSIZE (int)sizeof(srcT1)*3 +#define DSTSIZE (int)sizeof(dstT1)*3 +#endif #define DIG(a) a, __constant float mat_kernel[] = { COEFF }; -__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter - (__global const GENTYPE_SRC * restrict src, - const int src_step_in_pixel, - const int src_whole_cols, - const int src_whole_rows, - __global GENTYPE_DST * dst, - const int dst_offset_in_pixel, - const int dst_step_in_pixel, - const int dst_cols, - const int dst_rows) +__kernel void col_filter(__global const uchar * src, int src_step, int src_offset, int src_whole_rows, int src_whole_cols, + __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, float delta) { int x = get_global_id(0); int y = get_global_id(1); @@ -82,38 +71,38 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter int l_x = get_local_id(0); int l_y = get_local_id(1); - int start_addr = mad24(y, src_step_in_pixel, x); - int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols); + int start_addr = mad24(y, src_step, x * SRCSIZE); + int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * SRCSIZE); - int i; - GENTYPE_SRC sum, temp[READ_TIMES_COL]; - __local GENTYPE_SRC LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1]; + srcT sum, temp[READ_TIMES_COL]; + __local srcT LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1]; - //read pixels from src - for(i = 0;i>2<<2) -#elif CN==2 -#define ALIGN (((RADIUS)+1)>>1<<1) -#elif CN==3 -#define ALIGN (((RADIUS)+3)>>2<<2) -#elif CN==4 -#define ALIGN (RADIUS) +#ifdef DOUBLE_SUPPORT +#ifdef cl_amd_fp64 +#pragma OPENCL EXTENSION cl_amd_fp64:enable +#elif defined (cl_khr_fp64) +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#endif #endif +#define READ_TIMES_ROW ((2*(RADIUSX+LSIZE0)-1)/LSIZE0) //for c4 only +#define RADIUS 1 + #ifdef BORDER_REPLICATE -//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh +// BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh #define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i)) #define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr)) #endif #ifdef BORDER_REFLECT -//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb +// BORDER_REFLECT: fedcba|abcdefgh|hgfedcb #define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i)-1 : (i)) #define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)) #endif #ifdef BORDER_REFLECT_101 -//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba +// BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba #define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i)) #define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)) #endif -//blur function does not support BORDER_WRAP #ifdef BORDER_WRAP -//BORDER_WRAP: cdefgh|abcdefgh|abcdefg +// BORDER_WRAP: cdefgh|abcdefgh|abcdefg #define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (i)+(r_edge) : (i)) #define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (i)-(r_edge) : (addr)) #endif @@ -127,65 +123,56 @@ #endif //BORDER_CONSTANT #endif //EXTRA_EXTRAPOLATION -/********************************************************************************** -These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur. -Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle -kernel must be in the center. ROI is not supported either. -For channels =1,2,4, each kernels read 4 elements(not 4 pixels), and for channels =3, -the kernel read 4 pixels, save them to LDS and read the data needed from LDS to -calculate the result. -The length of the convovle kernel supported is related to the LSIZE0 and the MAX size -of LDS, which is HW related. -For channels = 1,3 the RADIUS is no more than LSIZE0*2 -For channels = 2, the RADIUS is no more than LSIZE0 -For channels = 4, arbitary RADIUS is supported unless the LDS is not enough -Niko -6/29/2011 -The info above maybe obsolete. -***********************************************************************************/ +#define noconvert + +#if CN != 3 +#define loadpix(addr) *(__global const srcT *)(addr) +#define storepix(val, addr) *(__global dstT *)(addr) = val +#define SRCSIZE (int)sizeof(srcT) +#define DSTSIZE (int)sizeof(dstT) +#else +#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr)) +#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr)) +#define SRCSIZE (int)sizeof(srcT1)*3 +#define DSTSIZE (int)sizeof(dstT1)*3 +#endif #define DIG(a) a, __constant float mat_kernel[] = { COEFF }; -__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C1_D0 - (__global uchar * restrict src, - int src_step_in_pixel, - int src_offset_x, int src_offset_y, - int src_cols, int src_rows, - int src_whole_cols, int src_whole_rows, - __global float * dst, - int dst_step_in_pixel, - int dst_cols, int dst_rows, - int radiusy) +__kernel void row_filter_C1_D0(__global const uchar * src, int src_step_in_pixel, int src_offset_x, int src_offset_y, + int src_cols, int src_rows, int src_whole_cols, int src_whole_rows, + __global float * dst, int dst_step_in_pixel, int dst_cols, int dst_rows, + int radiusy) { int x = get_global_id(0)<<2; int y = get_global_id(1); int l_x = get_local_id(0); int l_y = get_local_id(1); - int start_x = x+src_offset_x - RADIUSX & 0xfffffffc; + int start_x = x + src_offset_x - RADIUSX & 0xfffffffc; int offset = src_offset_x - RADIUSX & 3; int start_y = y + src_offset_y - radiusy; int start_addr = mad24(start_y, src_step_in_pixel, start_x); - int i; + float4 sum; uchar4 temp[READ_TIMES_ROW]; - __local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1]; + __local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW * LSIZE0 + 1]; #ifdef BORDER_CONSTANT int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols); // read pixels from src - for (i = 0; i < READ_TIMES_ROW; i++) + for (int i = 0; i < READ_TIMES_ROW; ++i) { - int current_addr = start_addr+i*LSIZE0*4; - current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0; - temp[i] = *(__global uchar4*)&src[current_addr]; + int current_addr = mad24(i, LSIZE0 << 2, start_addr); + current_addr = current_addr < end_addr && current_addr > 0 ? current_addr : 0; + temp[i] = *(__global const uchar4 *)&src[current_addr]; } // judge if read out of boundary #ifdef BORDER_ISOLATED - for (i = 0; isrc_whole_cols)| (start_y<0) | (start_y >= src_whole_rows); #endif - int4 index[READ_TIMES_ROW]; - int4 addr; + int4 index[READ_TIMES_ROW], addr; int s_y; if (not_all_in_range) { // judge if read out of boundary - for (i = 0; i < READ_TIMES_ROW; i++) + for (int i = 0; i < READ_TIMES_ROW; ++i) { - index[i] = (int4)(start_x+i*LSIZE0*4) + (int4)(0, 1, 2, 3); + index[i] = (int4)(mad24(i, LSIZE0 << 2, start_x)) + (int4)(0, 1, 2, 3); #ifdef BORDER_ISOLATED EXTRAPOLATE(index[i].x, src_offset_x, src_offset_x + src_cols); EXTRAPOLATE(index[i].y, src_offset_x, src_offset_x + src_cols); @@ -231,6 +217,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_ EXTRAPOLATE(index[i].w, 0, src_whole_cols); #endif } + s_y = start_y; #ifdef BORDER_ISOLATED EXTRAPOLATE(s_y, src_offset_y, src_offset_y + src_rows); @@ -239,9 +226,9 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_ #endif // read pixels from src - for (i = 0; i 0)) ? current_addr : 0; - temp[i] = src[current_addr]; + int current_addr = mad24(i, LSIZE0 * SRCSIZE, start_addr); + current_addr = current_addr < end_addr && current_addr >= 0 ? current_addr : 0; + temp[i] = loadpix(src + current_addr); } - //judge if read out of boundary + // judge if read out of boundary #ifdef BORDER_ISOLATED - for (i = 0; i 0)) ? current_addr : 0; - temp[i] = src[current_addr]; - } - - // judge if read out of boundary -#ifdef BORDER_ISOLATED - for (i = 0; i 0)) ? current_addr : 0; - temp[i] = src[current_addr]; - } - - // judge if read out of boundary -#ifdef BORDER_ISOLATED - for (i = 0; i= (r_edge) | (_y)<0 | (_y) >= (t_edge) ? (const_v) : SRC((_x),(_y)) +#else +#define ELEM(_x,_y,r_edge,t_edge,const_v) SRC((_x),(_y)) +#endif + +#define noconvert + +// horizontal and vertical filter kernels +// should be defined on host during compile time to avoid overhead +#define DIG(a) a, +__constant float mat_kernelX[] = { KERNEL_MATRIX_X }; +__constant float mat_kernelY[] = { KERNEL_MATRIX_Y }; + +__kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int srcOffsetY, int height, int width, + __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, float delta) +{ + // RADIUSX, RADIUSY are filter dimensions + // BLK_X, BLK_Y are local wrogroup sizes + // all these should be defined on host during compile time + // first lsmem array for source pixels used in first pass, + // second lsmemDy for storing first pass results + __local WT lsmem[BLK_Y + 2 * RADIUSY][BLK_X + 2 * RADIUSX]; + __local WT lsmemDy[BLK_Y][BLK_X + 2 * RADIUSX]; + + // get local and global ids - used as image and local memory array indexes + int lix = get_local_id(0); + int liy = get_local_id(1); + + int x = get_global_id(0); + int y = get_global_id(1); + + // calculate pixel position in source image taking image offset into account + int srcX = x + srcOffsetX - RADIUSX; + int srcY = y + srcOffsetY - RADIUSY; + int xb = srcX; + int yb = srcY; + + // extrapolate coordinates, if needed + // and read my own source pixel into local memory + // with account for extra border pixels, which will be read by starting workitems + int clocY = liy; + int cSrcY = srcY; + do + { + int yb = cSrcY; + EXTRAPOLATE(yb, (height)); + + int clocX = lix; + int cSrcX = srcX; + do + { + int xb = cSrcX; + EXTRAPOLATE(xb,(width)); + lsmem[clocY][clocX] = ELEM(xb, yb, (width), (height), 0 ); + + clocX += BLK_X; + cSrcX += BLK_X; + } + while(clocX < BLK_X+(RADIUSX*2)); + + clocY += BLK_Y; + cSrcY += BLK_Y; + } + while (clocY < BLK_Y+(RADIUSY*2)); + barrier(CLK_LOCAL_MEM_FENCE); + + // do vertical filter pass + // and store intermediate results to second local memory array + int i, clocX = lix; + WT sum = 0.0f; + do + { + sum = 0.0f; + for (i=0; i<=2*RADIUSY; i++) + sum = mad(lsmem[liy+i][clocX], mat_kernelY[i], sum); + lsmemDy[liy][clocX] = sum; + clocX += BLK_X; + } + while(clocX < BLK_X+(RADIUSX*2)); + barrier(CLK_LOCAL_MEM_FENCE); + + // if this pixel happened to be out of image borders because of global size rounding, + // then just return + if( x >= dst_cols || y >=dst_rows ) + return; + + // do second horizontal filter pass + // and calculate final result + sum = 0.0f; + for (i=0; i<=2*RADIUSX; i++) + sum = mad(lsmemDy[liy][lix+i], mat_kernelX[i], sum); + + // store result into destination image + storepix(convertToDstT(sum + (WT)(delta)), Dst + mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset))); +} diff --git a/modules/imgproc/src/opencl/morph.cl b/modules/imgproc/src/opencl/morph.cl index cb6e733ed..fe11b4994 100644 --- a/modules/imgproc/src/opencl/morph.cl +++ b/modules/imgproc/src/opencl/morph.cl @@ -43,6 +43,16 @@ #endif #endif +#if cn != 3 +#define loadpix(addr) *(__global const T *)(addr) +#define storepix(val, addr) *(__global T *)(addr) = val +#define TSIZE (int)sizeof(T) +#else +#define loadpix(addr) vload3(0, (__global const T1 *)(addr)) +#define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr)) +#define TSIZE ((int)sizeof(T1)*3) +#endif + #ifdef DEPTH_0 #ifdef ERODE #define VAL 255 @@ -50,16 +60,14 @@ #ifdef DILATE #define VAL 0 #endif -#endif -#ifdef DEPTH_5 +#elif defined DEPTH_5 #ifdef ERODE #define VAL FLT_MAX #endif #ifdef DILATE #define VAL -FLT_MAX #endif -#endif -#ifdef DEPTH_6 +#elif defined DEPTH_6 #ifdef ERODE #define VAL DBL_MAX #endif @@ -69,84 +77,80 @@ #endif #ifdef ERODE +#ifdef INTEL_DEVICE +// workaround for bug in Intel HD graphics drivers (10.18.10.3496 or older) +#define MORPH_OP(A,B) ((A) < (B) ? (A) : (B)) +#else #define MORPH_OP(A,B) min((A),(B)) #endif +#endif #ifdef DILATE #define MORPH_OP(A,B) max((A),(B)) #endif -//BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii -#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2) -__kernel void morph(__global const uchar * restrict srcptr, int src_step, int src_offset, +// BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii +#define ELEM(i, l_edge, r_edge, elem1, elem2) (i) < (l_edge) | (i) >= (r_edge) ? (elem1) : (elem2) + +__kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, __global uchar * dstptr, int dst_step, int dst_offset, - int src_offset_x, int src_offset_y, - int cols, int rows, - __constant uchar * mat_kernel, - int src_whole_cols, int src_whole_rows) + int src_offset_x, int src_offset_y, int cols, int rows, + __constant uchar * mat_kernel, int src_whole_cols, int src_whole_rows) { - int l_x = get_local_id(0); - int l_y = get_local_id(1); - int x = get_group_id(0)*LSIZE0; - int y = get_group_id(1)*LSIZE1; - int start_x = x+src_offset_x-RADIUSX; - int end_x = x + src_offset_x+LSIZE0+RADIUSX; - int width = end_x -(x+src_offset_x-RADIUSX)+1; - int start_y = y+src_offset_y-RADIUSY; - int point1 = mad24(l_y,LSIZE0,l_x); - int point2 = point1 + LSIZE0*LSIZE1; - int tl_x = point1 % width; - int tl_y = point1 / width; - int tl_x2 = point2 % width; - int tl_y2 = point2 / width; - int cur_x = start_x + tl_x; - int cur_y = start_y + tl_y; - int cur_x2 = start_x + tl_x2; - int cur_y2 = start_y + tl_y2; - int start_addr = mad24(cur_y,src_step, cur_x*(int)sizeof(GENTYPE)); - int start_addr2 = mad24(cur_y2,src_step, cur_x2*(int)sizeof(GENTYPE)); - GENTYPE temp0,temp1; - __local GENTYPE LDS_DAT[2*LSIZE1*LSIZE0]; + int gidx = get_global_id(0), gidy = get_global_id(1); + int l_x = get_local_id(0), l_y = get_local_id(1); + int x = get_group_id(0) * LSIZE0, y = get_group_id(1) * LSIZE1; + int start_x = x + src_offset_x - RADIUSX; + int end_x = x + src_offset_x + LSIZE0 + RADIUSX; + int width = end_x - (x + src_offset_x - RADIUSX) + 1; + int start_y = y + src_offset_y - RADIUSY; + int point1 = mad24(l_y, LSIZE0, l_x); + int point2 = point1 + LSIZE0 * LSIZE1; + int tl_x = point1 % width, tl_y = point1 / width; + int tl_x2 = point2 % width, tl_y2 = point2 / width; + int cur_x = start_x + tl_x, cur_y = start_y + tl_y; + int cur_x2 = start_x + tl_x2, cur_y2 = start_y + tl_y2; + int start_addr = mad24(cur_y, src_step, cur_x * TSIZE); + int start_addr2 = mad24(cur_y2, src_step, cur_x2 * TSIZE); - int end_addr = mad24(src_whole_rows - 1,src_step,src_whole_cols*(int)sizeof(GENTYPE)); - //read pixels from src - start_addr = ((start_addr < end_addr) && (start_addr > 0)) ? start_addr : 0; - start_addr2 = ((start_addr2 < end_addr) && (start_addr2 > 0)) ? start_addr2 : 0; - __global const GENTYPE * src; - src = (__global const GENTYPE *)(srcptr+start_addr); - temp0 = src[0]; - src = (__global const GENTYPE *)(srcptr+start_addr2); - temp1 = src[0]; - //judge if read out of boundary - temp0= ELEM(cur_x,0,src_whole_cols,(GENTYPE)VAL,temp0); - temp0= ELEM(cur_y,0,src_whole_rows,(GENTYPE)VAL,temp0); + __local T LDS_DAT[2*LSIZE1*LSIZE0]; - temp1= ELEM(cur_x2,0,src_whole_cols,(GENTYPE)VAL,temp1); - temp1= ELEM(cur_y2,0,src_whole_rows,(GENTYPE)VAL,temp1); + // read pixels from src + int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * TSIZE); + start_addr = start_addr < end_addr && start_addr > 0 ? start_addr : 0; + start_addr2 = start_addr2 < end_addr && start_addr2 > 0 ? start_addr2 : 0; + + T temp0 = loadpix(srcptr + start_addr); + T temp1 = loadpix(srcptr + start_addr2); + + // judge if read out of boundary + temp0 = ELEM(cur_x, 0, src_whole_cols, (T)(VAL),temp0); + temp0 = ELEM(cur_y, 0, src_whole_rows, (T)(VAL),temp0); + + temp1 = ELEM(cur_x2, 0, src_whole_cols, (T)(VAL), temp1); + temp1 = ELEM(cur_y2, 0, src_whole_rows, (T)(VAL), temp1); LDS_DAT[point1] = temp0; LDS_DAT[point2] = temp1; barrier(CLK_LOCAL_MEM_FENCE); - GENTYPE res = (GENTYPE)VAL; - for(int i=0; i<2*RADIUSY+1; i++) - for(int j=0; j<2*RADIUSX+1; j++) + + T res = (T)(VAL); + for (int i = 0, sizey = 2 * RADIUSY + 1; i < sizey; i++) + for (int j = 0, sizex = 2 * RADIUSX + 1; j < sizex; j++) { res = #ifndef RECTKERNEL mat_kernel[i*(2*RADIUSX+1)+j] ? #endif - MORPH_OP(res,LDS_DAT[mad24(l_y+i,width,l_x+j)]) + MORPH_OP(res, LDS_DAT[mad24(l_y + i, width, l_x + j)]) #ifndef RECTKERNEL - :res + : res #endif ; } - int gidx = get_global_id(0); - int gidy = get_global_id(1); - if(gidx> 3; - int qangle_step_shift = 0; - int qangle_step = (int)qangle.step >> (1 + qangle_step_shift); + int qangle_elem_size = CV_ELEM_SIZE1(qangle.type()); + int qangle_step = (int)qangle.step / (2 * qangle_elem_size); int idx = 0; idx = k.set(idx, height); @@ -1137,9 +1137,9 @@ static bool ocl_compute_hists(int nbins, int block_stride_x, int block_stride_y, int img_block_height = (height - CELLS_PER_BLOCK_Y * CELL_HEIGHT + block_stride_y)/block_stride_y; int blocks_total = img_block_width * img_block_height; - int qangle_step_shift = 0; + int qangle_elem_size = CV_ELEM_SIZE1(qangle.type()); int grad_quadstep = (int)grad.step >> 2; - int qangle_step = (int)qangle.step >> qangle_step_shift; + int qangle_step = (int)qangle.step / qangle_elem_size; int blocks_in_group = 4; size_t localThreads[3] = { blocks_in_group * 24, 2, 1 }; @@ -1316,11 +1316,12 @@ static bool ocl_extract_descrs_by_cols(int win_height, int win_width, int block_ static bool ocl_compute(InputArray _img, Size win_stride, std::vector& _descriptors, int descr_format, Size blockSize, Size cellSize, int nbins, Size blockStride, Size winSize, float sigma, bool gammaCorrection, double L2HysThreshold) { - Size imgSize = _img.size(); + Size imgSize = _img.size(); Size effect_size = imgSize; UMat grad(imgSize, CV_32FC2); - UMat qangle(imgSize, CV_8UC2); + int qangle_type = ocl::Device::getDefault().isIntel() ? CV_32SC2 : CV_8UC2; + UMat qangle(imgSize, qangle_type); const size_t block_hist_size = getBlockHistogramSize(blockSize, cellSize, nbins); const Size blocks_per_img = numPartsWithin(imgSize, blockSize, blockStride); @@ -1720,7 +1721,8 @@ static bool ocl_detect(InputArray img, std::vector &hits, double hit_thre Size imgSize = img.size(); Size effect_size = imgSize; UMat grad(imgSize, CV_32FC2); - UMat qangle(imgSize, CV_8UC2); + int qangle_type = ocl::Device::getDefault().isIntel() ? CV_32SC2 : CV_8UC2; + UMat qangle(imgSize, qangle_type); const size_t block_hist_size = getBlockHistogramSize(blockSize, cellSize, nbins); const Size blocks_per_img = numPartsWithin(imgSize, blockSize, blockStride); diff --git a/modules/objdetect/src/opencl/objdetect_hog.cl b/modules/objdetect/src/opencl/objdetect_hog.cl index 082f9ab7f..704dec444 100644 --- a/modules/objdetect/src/opencl/objdetect_hog.cl +++ b/modules/objdetect/src/opencl/objdetect_hog.cl @@ -50,6 +50,14 @@ #define NTHREADS 256 #define CV_PI_F 3.1415926535897932384626433832795f +#ifdef INTEL_DEVICE +#define QANGLE_TYPE int +#define QANGLE_TYPE2 int2 +#else +#define QANGLE_TYPE uchar +#define QANGLE_TYPE2 uchar2 +#endif + //---------------------------------------------------------------------------- // Histogram computation // 12 threads for a cell, 12x4 threads per block @@ -59,7 +67,7 @@ __kernel void compute_hists_lut_kernel( const int cnbins, const int cblock_hist_size, const int img_block_width, const int blocks_in_group, const int blocks_total, const int grad_quadstep, const int qangle_step, - __global const float* grad, __global const uchar* qangle, + __global const float* grad, __global const QANGLE_TYPE* qangle, __global const float* gauss_w_lut, __global float* block_hists, __local float* smem) { @@ -86,7 +94,7 @@ __kernel void compute_hists_lut_kernel( __global const float* grad_ptr = (gid < blocks_total) ? grad + offset_y * grad_quadstep + (offset_x << 1) : grad; - __global const uchar* qangle_ptr = (gid < blocks_total) ? + __global const QANGLE_TYPE* qangle_ptr = (gid < blocks_total) ? qangle + offset_y * qangle_step + (offset_x << 1) : qangle; __local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) + @@ -101,7 +109,7 @@ __kernel void compute_hists_lut_kernel( for (int dist_y = dist_y_begin; dist_y < dist_y_begin + 12; ++dist_y) { float2 vote = (float2) (grad_ptr[0], grad_ptr[1]); - uchar2 bin = (uchar2) (qangle_ptr[0], qangle_ptr[1]); + QANGLE_TYPE2 bin = (QANGLE_TYPE2) (qangle_ptr[0], qangle_ptr[1]); grad_ptr += grad_quadstep; qangle_ptr += qangle_step; @@ -133,9 +141,8 @@ __kernel void compute_hists_lut_kernel( final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] = hist_[0] + hist_[1] + hist_[2]; } -#ifdef CPU + barrier(CLK_LOCAL_MEM_FENCE); -#endif int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 12 + cell_thread_x; if ((tid < cblock_hist_size) && (gid < blocks_total)) @@ -558,7 +565,7 @@ __kernel void extract_descrs_by_cols_kernel( __kernel void compute_gradients_8UC4_kernel( const int height, const int width, const int img_step, const int grad_quadstep, const int qangle_step, - const __global uchar4 * img, __global float * grad, __global uchar * qangle, + const __global uchar4 * img, __global float * grad, __global QANGLE_TYPE * qangle, const float angle_scale, const char correct_gamma, const int cnbins) { const int x = get_global_id(0); @@ -660,7 +667,7 @@ __kernel void compute_gradients_8UC4_kernel( __kernel void compute_gradients_8UC1_kernel( const int height, const int width, const int img_step, const int grad_quadstep, const int qangle_step, - __global const uchar * img, __global float * grad, __global uchar * qangle, + __global const uchar * img, __global float * grad, __global QANGLE_TYPE * qangle, const float angle_scale, const char correct_gamma, const int cnbins) { const int x = get_global_id(0); diff --git a/modules/objdetect/test/opencl/test_hogdetector.cpp b/modules/objdetect/test/opencl/test_hogdetector.cpp index 8568352b6..b3ef6b48f 100644 --- a/modules/objdetect/test/opencl/test_hogdetector.cpp +++ b/modules/objdetect/test/opencl/test_hogdetector.cpp @@ -110,7 +110,7 @@ OCL_TEST_P(HOG, Detect) OCL_OFF(hog.detectMultiScale(img, cpu_found, 0, Size(8, 8), Size(0, 0), 1.05, 6)); OCL_ON(hog.detectMultiScale(uimg, gpu_found, 0, Size(8, 8), Size(0, 0), 1.05, 6)); - EXPECT_LT(checkRectSimilarity(img.size(), cpu_found, gpu_found), 1.0); + EXPECT_LT(checkRectSimilarity(img.size(), cpu_found, gpu_found), 0.05); } INSTANTIATE_TEST_CASE_P(OCL_ObjDetect, HOG, testing::Combine( diff --git a/modules/photo/src/merge.cpp b/modules/photo/src/merge.cpp index 7adfb5ec6..295e03c95 100644 --- a/modules/photo/src/merge.cpp +++ b/modules/photo/src/merge.cpp @@ -208,7 +208,7 @@ public: if(channels == 3) { weights[i] = weights[i].mul(saturation); } - weights[i] = weights[i].mul(wellexp); + weights[i] = weights[i].mul(wellexp) + 1e-12f; weight_sum += weights[i]; } int maxlevel = static_cast(logf(static_cast(min(size.width, size.height))) / logf(2.0f)); diff --git a/modules/photo/test/test_hdr.cpp b/modules/photo/test/test_hdr.cpp index 82ae25f52..27773fb38 100644 --- a/modules/photo/test/test_hdr.cpp +++ b/modules/photo/test/test_hdr.cpp @@ -166,6 +166,16 @@ TEST(Photo_MergeMertens, regression) merge->process(images, result); result.convertTo(result, CV_8UC3, 255); checkEqual(expected, result, 3, "Mertens"); + + Mat uniform(100, 100, CV_8UC3); + uniform = Scalar(0, 255, 0); + + images.clear(); + images.push_back(uniform); + + merge->process(images, result); + result.convertTo(result, CV_8UC3, 255); + checkEqual(uniform, result, 1e-2f, "Mertens"); } TEST(Photo_MergeDebevec, regression) diff --git a/modules/superres/CMakeLists.txt b/modules/superres/CMakeLists.txt index 092b1cd40..c360303f6 100644 --- a/modules/superres/CMakeLists.txt +++ b/modules/superres/CMakeLists.txt @@ -1,4 +1,4 @@ -if(ANDROID OR IOS) +if(IOS) ocv_module_disable(superres) endif() diff --git a/modules/superres/src/btv_l1.cpp b/modules/superres/src/btv_l1.cpp index 1e4aa48a7..d54b4b398 100644 --- a/modules/superres/src/btv_l1.cpp +++ b/modules/superres/src/btv_l1.cpp @@ -1014,10 +1014,8 @@ namespace return; #ifdef HAVE_OPENCL - if (isUmat_ && curFrame_.channels() == 1) + if (isUmat_) curFrame_.copyTo(ucurFrame_); - else - isUmat_ = false; #endif ++storePos_; diff --git a/modules/ts/src/ts_perf.cpp b/modules/ts/src/ts_perf.cpp index f1403b9c5..25981a721 100644 --- a/modules/ts/src/ts_perf.cpp +++ b/modules/ts/src/ts_perf.cpp @@ -20,7 +20,7 @@ static std::vector available_impls; static std::string param_impl; static enum PERF_STRATEGY strategyForce = PERF_STRATEGY_DEFAULT; -static enum PERF_STRATEGY strategyModule = PERF_STRATEGY_BASE; +static enum PERF_STRATEGY strategyModule = PERF_STRATEGY_SIMPLE; static double param_max_outliers; static double param_max_deviation; diff --git a/modules/video/src/bgfg_gaussmix2.cpp b/modules/video/src/bgfg_gaussmix2.cpp index 1e6ee0d88..098310f6b 100644 --- a/modules/video/src/bgfg_gaussmix2.cpp +++ b/modules/video/src/bgfg_gaussmix2.cpp @@ -578,7 +578,7 @@ public: for( int mode = 0; mode < nmodes; mode++, mean_m += nchannels ) { float weight = alpha1*gmm[mode].weight + prune;//need only weight if fit is found - + int swap_count = 0; //// //fit not found yet if( !fitsPDF ) @@ -643,6 +643,7 @@ public: if( weight < gmm[i-1].weight ) break; + swap_count++; //swap one up std::swap(gmm[i], gmm[i-1]); for( int c = 0; c < nchannels; c++ ) @@ -660,7 +661,7 @@ public: nmodes--; } - gmm[mode].weight = weight;//update weight by the calculated value + gmm[mode-swap_count].weight = weight;//update weight by the calculated value totalWeight += weight; } //go through all modes @@ -918,4 +919,4 @@ Ptr createBackgroundSubtractorMOG2(int _history, doubl } -/* End of file. */ \ No newline at end of file +/* End of file. */ diff --git a/modules/video/src/lkpyramid.cpp b/modules/video/src/lkpyramid.cpp index cd5758565..a33e47664 100644 --- a/modules/video/src/lkpyramid.cpp +++ b/modules/video/src/lkpyramid.cpp @@ -975,9 +975,7 @@ namespace cv idxArg = kernel.set(idxArg, imageI); //image2d_t I idxArg = kernel.set(idxArg, imageJ); //image2d_t J idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(prevPts)); // __global const float2* prevPts - idxArg = kernel.set(idxArg, (int)prevPts.step); // int prevPtsStep idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadWrite(nextPts)); // __global const float2* nextPts - idxArg = kernel.set(idxArg, (int)nextPts.step); // int nextPtsStep idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadWrite(status)); // __global uchar* status idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadWrite(err)); // __global float* err idxArg = kernel.set(idxArg, (int)level); // const int level diff --git a/modules/video/src/opencl/bgfg_mog2.cl b/modules/video/src/opencl/bgfg_mog2.cl index f895b5be7..9bc18b215 100644 --- a/modules/video/src/opencl/bgfg_mog2.cl +++ b/modules/video/src/opencl/bgfg_mog2.cl @@ -102,7 +102,7 @@ __kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame { float c_weight = alpha1 * _weight[(mode * frame_row + y) * weight_step + x] + prune; - + int swap_count = 0; if (!fitsPDF) { float c_var = _variance[(mode * frame_row + y) * var_step + x]; @@ -132,6 +132,7 @@ __kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame { if (c_weight < _weight[((i - 1) * frame_row + y) * weight_step + x]) break; + swap_count++; swap(_weight, x, y, i - 1, frame_row, weight_step); swap(_variance, x, y, i - 1, frame_row, var_step); #if (CN==1) @@ -149,7 +150,7 @@ __kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame nmodes--; } - _weight[(mode * frame_row + y) * weight_step + x] = c_weight; //update weight by the calculated value + _weight[((mode - swap_count) * frame_row + y) * weight_step + x] = c_weight; //update weight by the calculated value totalWeight += c_weight; } diff --git a/modules/video/src/opencl/pyrlk.cl b/modules/video/src/opencl/pyrlk.cl index c01855490..822e628f2 100644 --- a/modules/video/src/opencl/pyrlk.cl +++ b/modules/video/src/opencl/pyrlk.cl @@ -262,50 +262,9 @@ inline void GetError(image2d_t J, const float x, const float y, const float* Pch *errval += fabs(diff); } -inline void SetPatch4(image2d_t I, const float x, const float y, - float4* Pch, float4* Dx, float4* Dy, - float* A11, float* A12, float* A22) -{ - *Pch = read_imagef(I, sampler, (float2)(x, y)); - - float4 dIdx = 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x + 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) - - (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x - 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1))); - - float4 dIdy = 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y + 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) - - (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y - 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1))); - - - *Dx = dIdx; - *Dy = dIdy; - float4 sqIdx = dIdx * dIdx; - *A11 += sqIdx.x + sqIdx.y + sqIdx.z; - sqIdx = dIdx * dIdy; - *A12 += sqIdx.x + sqIdx.y + sqIdx.z; - sqIdx = dIdy * dIdy; - *A22 += sqIdx.x + sqIdx.y + sqIdx.z; -} - -inline void GetPatch4(image2d_t J, const float x, const float y, - const float4* Pch, const float4* Dx, const float4* Dy, - float* b1, float* b2) -{ - float4 J_val = read_imagef(J, sampler, (float2)(x, y)); - float4 diff = (J_val - *Pch) * 32.0f; - float4 xdiff = diff* *Dx; - *b1 += xdiff.x + xdiff.y + xdiff.z; - xdiff = diff* *Dy; - *b2 += xdiff.x + xdiff.y + xdiff.z; -} - -inline void GetError4(image2d_t J, const float x, const float y, const float4* Pch, float* errval) -{ - float4 diff = read_imagef(J, sampler, (float2)(x,y))-*Pch; - *errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z); -} - #define GRIDSIZE 3 __kernel void lkSparse(image2d_t I, image2d_t J, - __global const float2* prevPts, int prevPtsStep, __global float2* nextPts, int nextPtsStep, __global uchar* status, __global float* err, + __global const float2* prevPts, __global float2* nextPts, __global uchar* status, __global float* err, const int level, const int rows, const int cols, int PATCH_X, int PATCH_Y, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr) { __local float smem1[BUFFER]; @@ -434,9 +393,8 @@ __kernel void lkSparse(image2d_t I, image2d_t J, { if (tid == 0 && level == 0) status[gid] = 0; - return; + break; } - float b1 = 0; float b2 = 0; diff --git a/modules/video/test/test_tvl1optflow.cpp b/modules/video/test/test_tvl1optflow.cpp index 804eae8b6..274c13e65 100644 --- a/modules/video/test/test_tvl1optflow.cpp +++ b/modules/video/test/test_tvl1optflow.cpp @@ -133,14 +133,13 @@ namespace } } } - return sqrt(sum / (1e-9 + counter)); } } TEST(Video_calcOpticalFlowDual_TVL1, Regression) { - const double MAX_RMSE = 0.02; + const double MAX_RMSE = 0.03; const string frame1_path = TS::ptr()->get_data_path() + "optflow/RubberWhale1.png"; const string frame2_path = TS::ptr()->get_data_path() + "optflow/RubberWhale2.png"; diff --git a/modules/viz/doc/widget.rst b/modules/viz/doc/widget.rst index a8d9ac6fd..906adf9ba 100644 --- a/modules/viz/doc/widget.rst +++ b/modules/viz/doc/widget.rst @@ -1050,7 +1050,7 @@ viz::WWidgetMerger::WWidgetMerger --------------------------------------- Constructs a WWidgetMerger. -.. ocv:WWidgetMerger:: WWidgetMerger() +.. ocv:function:: WWidgetMerger() viz::WWidgetMerger::addCloud ------------------------------- diff --git a/samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp b/samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp index f4dd4e5e4..424a38e93 100644 --- a/samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp +++ b/samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp @@ -40,13 +40,13 @@ int main( int argc, char** argv ) hsv_half_down = hsv_base( Range( hsv_base.rows/2, hsv_base.rows - 1 ), Range( 0, hsv_base.cols - 1 ) ); - /// Using 30 bins for hue and 32 for saturation + /// Using 50 bins for hue and 60 for saturation int h_bins = 50; int s_bins = 60; int histSize[] = { h_bins, s_bins }; - // hue varies from 0 to 256, saturation from 0 to 180 - float s_ranges[] = { 0, 256 }; + // hue varies from 0 to 179, saturation from 0 to 255 float h_ranges[] = { 0, 180 }; + float s_ranges[] = { 0, 256 }; const float* ranges[] = { h_ranges, s_ranges }; diff --git a/samples/gpu/CMakeLists.txt b/samples/gpu/CMakeLists.txt index 46b465a87..ca5243aa0 100644 --- a/samples/gpu/CMakeLists.txt +++ b/samples/gpu/CMakeLists.txt @@ -49,8 +49,8 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND) target_link_libraries(${the_target} ${OPENCV_LINKER_LIBS} ${OPENCV_CUDA_SAMPLES_REQUIRED_DEPS}) - if(HAVE_CUDA) - target_link_libraries(${the_target} ${CUDA_CUDA_LIBRARY}) + if(HAVE_CUDA AND NOT ANDROID) + target_link_libraries(${the_target} ${CUDA_CUDA_LIBRARY}) endif() if(HAVE_opencv_nonfree) diff --git a/samples/gpu/brox_optical_flow.cpp b/samples/gpu/brox_optical_flow.cpp index 08973861f..638aade45 100644 --- a/samples/gpu/brox_optical_flow.cpp +++ b/samples/gpu/brox_optical_flow.cpp @@ -1,7 +1,7 @@ #include #include #include -#include +#include #include "opencv2/core.hpp" #include "opencv2/core/utility.hpp" diff --git a/samples/gpu/opticalflow_nvidia_api.cpp b/samples/gpu/opticalflow_nvidia_api.cpp index 29aaa0645..63eebfdaf 100644 --- a/samples/gpu/opticalflow_nvidia_api.cpp +++ b/samples/gpu/opticalflow_nvidia_api.cpp @@ -7,6 +7,7 @@ #include #include #include +#include #include "cvconfig.h" #include diff --git a/samples/gpu/super_resolution.cpp b/samples/gpu/super_resolution.cpp index 3066e8f74..4e3de21db 100644 --- a/samples/gpu/super_resolution.cpp +++ b/samples/gpu/super_resolution.cpp @@ -1,6 +1,8 @@ #include #include #include +#include + #include "opencv2/core.hpp" #include "opencv2/core/utility.hpp" #include "opencv2/highgui.hpp"