diff --git a/CMakeLists.txt b/CMakeLists.txt index 8cec98fc94..73def95474 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -32,6 +32,10 @@ else(NOT CMAKE_TOOLCHAIN_FILE) endif(NOT CMAKE_TOOLCHAIN_FILE) +if(POLICY CMP0022) + cmake_policy(SET CMP0022 OLD) +endif() + # must go before the project command set(CMAKE_CONFIGURATION_TYPES "Debug;Release" CACHE STRING "Configs" FORCE) if(DEFINED CMAKE_BUILD_TYPE) @@ -136,6 +140,7 @@ OCV_OPTION(WITH_QT "Build with Qt Backend support" OFF OCV_OPTION(WITH_WIN32UI "Build with Win32 UI Backend support" ON IF WIN32 ) OCV_OPTION(WITH_QUICKTIME "Use QuickTime for Video I/O insted of QTKit" OFF IF APPLE ) OCV_OPTION(WITH_TBB "Include Intel TBB support" OFF IF (NOT IOS) ) +OCV_OPTION(WITH_OPENMP "Include OpenMP support" OFF) OCV_OPTION(WITH_CSTRIPES "Include C= support" OFF IF WIN32 ) OCV_OPTION(WITH_TIFF "Include TIFF support" ON IF (NOT IOS) ) OCV_OPTION(WITH_UNICAP "Include Unicap support (GPL)" OFF IF (UNIX AND NOT APPLE AND NOT ANDROID) ) diff --git a/cmake/OpenCVFindLibsPerf.cmake b/cmake/OpenCVFindLibsPerf.cmake index b8945c257b..4b80b1f786 100644 --- a/cmake/OpenCVFindLibsPerf.cmake +++ b/cmake/OpenCVFindLibsPerf.cmake @@ -86,13 +86,13 @@ else() endif() # --- OpenMP --- -if(NOT HAVE_TBB AND NOT HAVE_CSTRIPES) - set(_fname "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/omptest.cpp") - file(WRITE "${_fname}" "#ifndef _OPENMP\n#error\n#endif\nint main() { return 0; }\n") - try_compile(HAVE_OPENMP "${CMAKE_BINARY_DIR}" "${_fname}") - file(REMOVE "${_fname}") -else() - set(HAVE_OPENMP 0) +if(WITH_OPENMP AND NOT HAVE_TBB AND NOT HAVE_CSTRIPES) + find_package(OpenMP) + if(OPENMP_FOUND) + set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${OpenMP_C_FLAGS}") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}") + endif() + set(HAVE_OPENMP "${OPENMP_FOUND}") endif() # --- GCD --- diff --git a/cmake/OpenCVGenConfig.cmake b/cmake/OpenCVGenConfig.cmake index 881cd371a9..362841b217 100644 --- a/cmake/OpenCVGenConfig.cmake +++ b/cmake/OpenCVGenConfig.cmake @@ -74,7 +74,14 @@ if(ANDROID AND NOT BUILD_SHARED_LIBS AND HAVE_TBB) list(APPEND OpenCV2_INCLUDE_DIRS_CONFIGCMAKE ${TBB_INCLUDE_DIRS}) endif() -export(TARGETS ${OpenCVModules_TARGETS} FILE "${CMAKE_BINARY_DIR}/OpenCVModules.cmake") +set(modules_file_suffix "") +if(ANDROID) + # the REPLACE here is needed, because OpenCVModules_armeabi.cmake includes + # OpenCVModules_armeabi-*.cmake, which would match OpenCVModules_armeabi-v7a*.cmake. + string(REPLACE - _ modules_file_suffix "_${ANDROID_NDK_ABI_NAME}") +endif() + +export(TARGETS ${OpenCVModules_TARGETS} FILE "${CMAKE_BINARY_DIR}/OpenCVModules${modules_file_suffix}.cmake") configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig.cmake.in" "${CMAKE_BINARY_DIR}/OpenCVConfig.cmake" IMMEDIATE @ONLY) #support for version checking when finding opencv. find_package(OpenCV 2.3.1 EXACT) should now work. @@ -94,7 +101,7 @@ endif() configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig.cmake.in" "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake" IMMEDIATE @ONLY) configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig-version.cmake.in" "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake" IMMEDIATE @ONLY) -if(UNIX) +if(UNIX) # ANDROID configuration is created here also #http://www.vtk.org/Wiki/CMake/Tutorials/Packaging reference # For a command "find_package( [major[.minor]] [EXACT] [REQUIRED|QUIET])" # cmake will look in the following dir on unix: @@ -104,11 +111,11 @@ if(UNIX) if(INSTALL_TO_MANGLED_PATHS) install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/) install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/) - install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/) + install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ FILE OpenCVModules${modules_file_suffix}.cmake) else() install(FILES "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/) install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/) - install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/) + install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ FILE OpenCVModules${modules_file_suffix}.cmake) endif() endif() @@ -128,10 +135,10 @@ if(WIN32) configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig-version.cmake.in" "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" IMMEDIATE @ONLY) if(BUILD_SHARED_LIBS) install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/lib") - install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/lib") + install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/lib" FILE OpenCVModules${modules_file_suffix}.cmake) else() install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/staticlib") - install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/staticlib") + install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}/staticlib" FILE OpenCVModules${modules_file_suffix}.cmake) endif() install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}") install(FILES "${OpenCV_SOURCE_DIR}/cmake/OpenCVConfig.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}/") diff --git a/cmake/OpenCVGenPkgconfig.cmake b/cmake/OpenCVGenPkgconfig.cmake index 7bfc7bc5af..a36b70e941 100644 --- a/cmake/OpenCVGenPkgconfig.cmake +++ b/cmake/OpenCVGenPkgconfig.cmake @@ -57,8 +57,17 @@ endforeach() # add extra dependencies required for OpenCV set(OpenCV_LIB_COMPONENTS ${OpenCV_LIB_COMPONENTS_}) if(OpenCV_EXTRA_COMPONENTS) - string(REPLACE ";" " " OpenCV_EXTRA_COMPONENTS "${OpenCV_EXTRA_COMPONENTS}") - set(OpenCV_LIB_COMPONENTS "${OpenCV_LIB_COMPONENTS} ${OpenCV_EXTRA_COMPONENTS}") + foreach(extra_component ${OpenCV_EXTRA_COMPONENTS}) + + if(extra_component MATCHES "^-[lL]" OR extra_component MATCHES "[\\/]") + set(maybe_l_prefix "") + else() + set(maybe_l_prefix "-l") + endif() + + set(OpenCV_LIB_COMPONENTS "${OpenCV_LIB_COMPONENTS} ${maybe_l_prefix}${extra_component}") + + endforeach() endif() #generate the .pc file diff --git a/cmake/OpenCVModule.cmake b/cmake/OpenCVModule.cmake index 5dbe957ddd..ed5acc76bd 100644 --- a/cmake/OpenCVModule.cmake +++ b/cmake/OpenCVModule.cmake @@ -538,9 +538,10 @@ macro(ocv_create_module) if(NOT "${ARGN}" STREQUAL "SKIP_LINK") target_link_libraries(${the_module} ${OPENCV_MODULE_${the_module}_DEPS}) - target_link_libraries(${the_module} LINK_PRIVATE ${OPENCV_MODULE_${the_module}_DEPS_EXT} ${OPENCV_LINKER_LIBS} ${IPP_LIBS} ${ARGN}) + target_link_libraries(${the_module} LINK_INTERFACE_LIBRARIES ${OPENCV_MODULE_${the_module}_DEPS}) + target_link_libraries(${the_module} ${OPENCV_MODULE_${the_module}_DEPS_EXT} ${OPENCV_LINKER_LIBS} ${IPP_LIBS} ${ARGN}) if (HAVE_CUDA) - target_link_libraries(${the_module} LINK_PRIVATE ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY}) + target_link_libraries(${the_module} ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY}) endif() endif() diff --git a/cmake/OpenCVUtils.cmake b/cmake/OpenCVUtils.cmake index 818ec9b4ba..13461e82c1 100644 --- a/cmake/OpenCVUtils.cmake +++ b/cmake/OpenCVUtils.cmake @@ -11,6 +11,18 @@ if(NOT COMMAND find_host_program) endmacro() endif() +# assert macro +# Note: it doesn't support lists in arguments +# Usage samples: +# ocv_assert(MyLib_FOUND) +# ocv_assert(DEFINED MyLib_INCLUDE_DIRS) +macro(ocv_assert) + if(NOT (${ARGN})) + string(REPLACE ";" " " __assert_msg "${ARGN}") + message(AUTHOR_WARNING "Assertion failed: ${__assert_msg}") + endif() +endmacro() + macro(ocv_check_environment_variables) foreach(_var ${ARGN}) if(NOT DEFINED ${_var} AND DEFINED ENV{${_var}}) @@ -467,9 +479,10 @@ function(ocv_install_target) set(isArchive 0) set(isDst 0) + unset(__dst) foreach(e ${ARGN}) if(isDst EQUAL 1) - set(DST "${e}") + set(__dst "${e}") break() endif() if(isArchive EQUAL 1 AND e STREQUAL "DESTINATION") @@ -482,18 +495,20 @@ function(ocv_install_target) endif() endforeach() -# message(STATUS "Process ${__target} dst=${DST}...") - if(NOT DEFINED DST) - set(DST "OPENCV_LIB_INSTALL_PATH") - endif() - - get_target_property(fname ${__target} LOCATION_DEBUG) - string(REPLACE ".lib" ".pdb" fname "${fname}") - install(FILES ${fname} DESTINATION ${DST} CONFIGURATIONS Debug) +# message(STATUS "Process ${__target} dst=${__dst}...") + if(DEFINED __dst) + get_target_property(fname ${__target} LOCATION_DEBUG) + if(fname MATCHES "\\.lib$") + string(REGEX REPLACE "\\.lib$" ".pdb" fname "${fname}") + install(FILES ${fname} DESTINATION ${__dst} CONFIGURATIONS Debug) + endif() - get_target_property(fname ${__target} LOCATION_RELEASE) - string(REPLACE ".lib" ".pdb" fname "${fname}") - install(FILES ${fname} DESTINATION ${DST} CONFIGURATIONS Release) + get_target_property(fname ${__target} LOCATION_RELEASE) + if(fname MATCHES "\\.lib$") + string(REGEX REPLACE "\\.lib$" ".pdb" fname "${fname}") + install(FILES ${fname} DESTINATION ${__dst} CONFIGURATIONS Release) + endif() + endif() endif() endif() endfunction() diff --git a/cmake/templates/OpenCVConfig.cmake.in b/cmake/templates/OpenCVConfig.cmake.in index c1d8021721..2dfc78169b 100644 --- a/cmake/templates/OpenCVConfig.cmake.in +++ b/cmake/templates/OpenCVConfig.cmake.in @@ -37,7 +37,12 @@ # # =================================================================================== -include(${CMAKE_CURRENT_LIST_DIR}/OpenCVModules.cmake) +set(modules_file_suffix "") +if(ANDROID) + string(REPLACE - _ modules_file_suffix "_${ANDROID_NDK_ABI_NAME}") +endif() + +include(${CMAKE_CURRENT_LIST_DIR}/OpenCVModules${modules_file_suffix}.cmake) # 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/tutorials/ios/image_manipulation/image_manipulation.rst b/doc/tutorials/ios/image_manipulation/image_manipulation.rst index c4cde19907..3eb8913be4 100644 --- a/doc/tutorials/ios/image_manipulation/image_manipulation.rst +++ b/doc/tutorials/ios/image_manipulation/image_manipulation.rst @@ -12,7 +12,8 @@ In this tutorial we will learn how to do basic image processing using OpenCV in *Introduction* ============== -In *OpenCV* all the image processing operations are done on *Mat*. iOS uses UIImage object to display image. One of the thing is to convert UIImage object to Mat object. Below is the code to convert UIImage to Mat. +In *OpenCV* all the image processing operations are usually carried out on the *Mat* structure. In iOS however, to render an image on screen it have to be an instance of the *UIImage* class. To convert an *OpenCV Mat* to an *UIImage* we use the *Core Graphics* framework available in iOS. Below is the code needed to covert back and forth between Mat's and UIImage's. + .. code-block:: cpp @@ -22,7 +23,7 @@ In *OpenCV* all the image processing operations are done on *Mat*. iOS uses UIIm CGFloat cols = image.size.width; CGFloat rows = image.size.height; - cv::Mat cvMat(rows, cols, CV_8UC4); // 8 bits per component, 4 channels + cv::Mat cvMat(rows, cols, CV_8UC4); // 8 bits per component, 4 channels (color channels + alpha) CGContextRef contextRef = CGBitmapContextCreate(cvMat.data, // Pointer to data cols, // Width of bitmap @@ -35,7 +36,6 @@ In *OpenCV* all the image processing operations are done on *Mat*. iOS uses UIIm CGContextDrawImage(contextRef, CGRectMake(0, 0, cols, rows), image.CGImage); CGContextRelease(contextRef); - CGColorSpaceRelease(colorSpace); return cvMat; } @@ -61,12 +61,11 @@ In *OpenCV* all the image processing operations are done on *Mat*. iOS uses UIIm CGContextDrawImage(contextRef, CGRectMake(0, 0, cols, rows), image.CGImage); CGContextRelease(contextRef); - CGColorSpaceRelease(colorSpace); return cvMat; } -Once we obtain the Mat Object. We can do all our processing on Mat object, similar to cpp. For example if we want to convert image to gray, we can do it via below code. +After the processing we need to convert it back to UIImage. The code below can handle both gray-scale and color image conversions (determined by the number of channels in the *if* statement). .. code-block:: cpp diff --git a/modules/calib3d/src/calibration.cpp b/modules/calib3d/src/calibration.cpp index 893c0e9a96..b93b3f7eb9 100644 --- a/modules/calib3d/src/calibration.cpp +++ b/modules/calib3d/src/calibration.cpp @@ -1164,8 +1164,8 @@ CV_IMPL void cvInitIntrinsicParams2D( const CvMat* objectPoints, matA.reset(cvCreateMat( 2*nimages, 2, CV_64F )); _b.reset(cvCreateMat( 2*nimages, 1, CV_64F )); - a[2] = (imageSize.width - 1)*0.5; - a[5] = (imageSize.height - 1)*0.5; + a[2] = (!imageSize.width) ? 0.5 : (imageSize.width - 1)*0.5; + a[5] = (!imageSize.height) ? 0.5 : (imageSize.height - 1)*0.5; _allH.reset(cvCreateMat( nimages, 9, CV_64F )); // extract vanishing points in order to obtain initial value for the focal length diff --git a/modules/java/CMakeLists.txt b/modules/java/CMakeLists.txt index a15ef8947f..5e6252a612 100644 --- a/modules/java/CMakeLists.txt +++ b/modules/java/CMakeLists.txt @@ -280,7 +280,12 @@ set_target_properties(${the_module} PROPERTIES COMPILE_DEFINITIONS OPENCV_NOSTL) if(BUILD_FAT_JAVA_LIB) set(__deps ${OPENCV_MODULE_${the_module}_DEPS} ${OPENCV_MODULES_BUILD}) - list(REMOVE_ITEM __deps ${the_module} opencv_ts) + foreach(m ${OPENCV_MODULES_BUILD}) # filterout INTERNAL (like opencv_ts) and BINDINGS (like opencv_python) modules + ocv_assert(DEFINED OPENCV_MODULE_${m}_CLASS) + if(NOT OPENCV_MODULE_${m}_CLASS STREQUAL "PUBLIC") + list(REMOVE_ITEM __deps ${m}) + endif() + endforeach() ocv_list_unique(__deps) set(__extradeps ${__deps}) ocv_list_filterout(__extradeps "^opencv_") diff --git a/modules/nonfree/perf/perf_surf.ocl.cpp b/modules/nonfree/perf/perf_surf.ocl.cpp index fdd1931bd8..cc48aa28c5 100644 --- a/modules/nonfree/perf/perf_surf.ocl.cpp +++ b/modules/nonfree/perf/perf_surf.ocl.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/nonfree/src/opencl/surf.cl b/modules/nonfree/src/opencl/surf.cl index aace143d53..02f77c224d 100644 --- a/modules/nonfree/src/opencl/surf.cl +++ b/modules/nonfree/src/opencl/surf.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/nonfree/src/surf.ocl.cpp b/modules/nonfree/src/surf.ocl.cpp index 6aa71d8ea8..20367ab98f 100644 --- a/modules/nonfree/src/surf.ocl.cpp +++ b/modules/nonfree/src/surf.ocl.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/nonfree/test/test_surf.ocl.cpp b/modules/nonfree/test/test_surf.ocl.cpp index d6a877bc80..c52b6e366e 100644 --- a/modules/nonfree/test/test_surf.ocl.cpp +++ b/modules/nonfree/test/test_surf.ocl.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/include/opencv2/ocl/matrix_operations.hpp b/modules/ocl/include/opencv2/ocl/matrix_operations.hpp index 1361367fca..33748c0a5d 100644 --- a/modules/ocl/include/opencv2/ocl/matrix_operations.hpp +++ b/modules/ocl/include/opencv2/ocl/matrix_operations.hpp @@ -23,7 +23,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/main.cpp b/modules/ocl/perf/main.cpp index 78ebc1a268..836f8ee9bd 100644 --- a/modules/ocl/perf/main.cpp +++ b/modules/ocl/perf/main.cpp @@ -22,7 +22,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_arithm.cpp b/modules/ocl/perf/perf_arithm.cpp index 880bdff5eb..025221b4ee 100644 --- a/modules/ocl/perf/perf_arithm.cpp +++ b/modules/ocl/perf/perf_arithm.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_bgfg.cpp b/modules/ocl/perf/perf_bgfg.cpp index b3d02f27a5..95099640f5 100644 --- a/modules/ocl/perf/perf_bgfg.cpp +++ b/modules/ocl/perf/perf_bgfg.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_blend.cpp b/modules/ocl/perf/perf_blend.cpp index 018ec63154..a5e057ffca 100644 --- a/modules/ocl/perf/perf_blend.cpp +++ b/modules/ocl/perf/perf_blend.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_brute_force_matcher.cpp b/modules/ocl/perf/perf_brute_force_matcher.cpp index 33c42c72dc..86c0a3c70d 100644 --- a/modules/ocl/perf/perf_brute_force_matcher.cpp +++ b/modules/ocl/perf/perf_brute_force_matcher.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_calib3d.cpp b/modules/ocl/perf/perf_calib3d.cpp index 997e848561..12fee549b8 100644 --- a/modules/ocl/perf/perf_calib3d.cpp +++ b/modules/ocl/perf/perf_calib3d.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_canny.cpp b/modules/ocl/perf/perf_canny.cpp index 259684092f..33723daa32 100644 --- a/modules/ocl/perf/perf_canny.cpp +++ b/modules/ocl/perf/perf_canny.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_color.cpp b/modules/ocl/perf/perf_color.cpp index b66fc2b0ae..e56db2495a 100644 --- a/modules/ocl/perf/perf_color.cpp +++ b/modules/ocl/perf/perf_color.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_fft.cpp b/modules/ocl/perf/perf_fft.cpp index 4cba47e960..49da659361 100644 --- a/modules/ocl/perf/perf_fft.cpp +++ b/modules/ocl/perf/perf_fft.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_filters.cpp b/modules/ocl/perf/perf_filters.cpp index 7f2758877b..b6dcd2a08a 100644 --- a/modules/ocl/perf/perf_filters.cpp +++ b/modules/ocl/perf/perf_filters.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_gemm.cpp b/modules/ocl/perf/perf_gemm.cpp index 803e1f91b6..4dcd5d4d6a 100644 --- a/modules/ocl/perf/perf_gemm.cpp +++ b/modules/ocl/perf/perf_gemm.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_gftt.cpp b/modules/ocl/perf/perf_gftt.cpp index 8a29adc0cc..af24c3489a 100644 --- a/modules/ocl/perf/perf_gftt.cpp +++ b/modules/ocl/perf/perf_gftt.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_haar.cpp b/modules/ocl/perf/perf_haar.cpp index 9c258fe259..5332a2d932 100644 --- a/modules/ocl/perf/perf_haar.cpp +++ b/modules/ocl/perf/perf_haar.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_hog.cpp b/modules/ocl/perf/perf_hog.cpp index fe5d9d1904..2a67311170 100644 --- a/modules/ocl/perf/perf_hog.cpp +++ b/modules/ocl/perf/perf_hog.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_imgproc.cpp b/modules/ocl/perf/perf_imgproc.cpp index 48dbf4dde2..7a9eab5d3b 100644 --- a/modules/ocl/perf/perf_imgproc.cpp +++ b/modules/ocl/perf/perf_imgproc.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_kalman.cpp b/modules/ocl/perf/perf_kalman.cpp index 017a8a70dc..946444ad9f 100644 --- a/modules/ocl/perf/perf_kalman.cpp +++ b/modules/ocl/perf/perf_kalman.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_match_template.cpp b/modules/ocl/perf/perf_match_template.cpp index 869e01e602..68192cf803 100644 --- a/modules/ocl/perf/perf_match_template.cpp +++ b/modules/ocl/perf/perf_match_template.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_matrix_operation.cpp b/modules/ocl/perf/perf_matrix_operation.cpp index 8266f0930d..3035c97f04 100644 --- a/modules/ocl/perf/perf_matrix_operation.cpp +++ b/modules/ocl/perf/perf_matrix_operation.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_ml.cpp b/modules/ocl/perf/perf_ml.cpp index fac471ed4c..db45eceb8f 100644 --- a/modules/ocl/perf/perf_ml.cpp +++ b/modules/ocl/perf/perf_ml.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_moments.cpp b/modules/ocl/perf/perf_moments.cpp index 6ecc766517..a36e1a13ed 100644 --- a/modules/ocl/perf/perf_moments.cpp +++ b/modules/ocl/perf/perf_moments.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_norm.cpp b/modules/ocl/perf/perf_norm.cpp index 35ac006394..ff49eb4eda 100644 --- a/modules/ocl/perf/perf_norm.cpp +++ b/modules/ocl/perf/perf_norm.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_opticalflow.cpp b/modules/ocl/perf/perf_opticalflow.cpp index 8613075261..bc1761b497 100644 --- a/modules/ocl/perf/perf_opticalflow.cpp +++ b/modules/ocl/perf/perf_opticalflow.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_precomp.hpp b/modules/ocl/perf/perf_precomp.hpp index 0235b3d045..f4b59a5fde 100644 --- a/modules/ocl/perf/perf_precomp.hpp +++ b/modules/ocl/perf/perf_precomp.hpp @@ -22,7 +22,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_pyramid.cpp b/modules/ocl/perf/perf_pyramid.cpp index 19c728bb77..c799853db4 100644 --- a/modules/ocl/perf/perf_pyramid.cpp +++ b/modules/ocl/perf/perf_pyramid.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/perf/perf_split_merge.cpp b/modules/ocl/perf/perf_split_merge.cpp index 3821a8e163..f2f7c41155 100644 --- a/modules/ocl/perf/perf_split_merge.cpp +++ b/modules/ocl/perf/perf_split_merge.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/arithm.cpp b/modules/ocl/src/arithm.cpp index cff0875743..6bfa7333a4 100644 --- a/modules/ocl/src/arithm.cpp +++ b/modules/ocl/src/arithm.cpp @@ -32,7 +32,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/bgfg_mog.cpp b/modules/ocl/src/bgfg_mog.cpp index 6ea520e37d..c6883661bc 100644 --- a/modules/ocl/src/bgfg_mog.cpp +++ b/modules/ocl/src/bgfg_mog.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/blend.cpp b/modules/ocl/src/blend.cpp index 5372702a51..c9bba13c94 100644 --- a/modules/ocl/src/blend.cpp +++ b/modules/ocl/src/blend.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/brute_force_matcher.cpp b/modules/ocl/src/brute_force_matcher.cpp index 59a274faee..3c2988dc8d 100644 --- a/modules/ocl/src/brute_force_matcher.cpp +++ b/modules/ocl/src/brute_force_matcher.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/build_warps.cpp b/modules/ocl/src/build_warps.cpp index fda3b1c5f5..bc24f5e385 100644 --- a/modules/ocl/src/build_warps.cpp +++ b/modules/ocl/src/build_warps.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/canny.cpp b/modules/ocl/src/canny.cpp index d321d91c52..3f5de52748 100644 --- a/modules/ocl/src/canny.cpp +++ b/modules/ocl/src/canny.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/cl_context.cpp b/modules/ocl/src/cl_context.cpp index 72ffce485f..2b3129d05e 100644 --- a/modules/ocl/src/cl_context.cpp +++ b/modules/ocl/src/cl_context.cpp @@ -27,7 +27,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -757,6 +757,9 @@ __Module::~__Module() #if defined(WIN32) && defined(CVAPI_EXPORTS) +extern "C" +BOOL WINAPI DllMain(HINSTANCE /*hInst*/, DWORD fdwReason, LPVOID lpReserved); + extern "C" BOOL WINAPI DllMain(HINSTANCE /*hInst*/, DWORD fdwReason, LPVOID lpReserved) { diff --git a/modules/ocl/src/cl_operations.cpp b/modules/ocl/src/cl_operations.cpp index 5f04561e94..f83220dae7 100644 --- a/modules/ocl/src/cl_operations.cpp +++ b/modules/ocl/src/cl_operations.cpp @@ -27,7 +27,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/cl_programcache.cpp b/modules/ocl/src/cl_programcache.cpp index ba279b794d..1254e30634 100644 --- a/modules/ocl/src/cl_programcache.cpp +++ b/modules/ocl/src/cl_programcache.cpp @@ -27,7 +27,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -61,12 +61,16 @@ namespace cv { namespace ocl { cv::Mutex ProgramCache::mutexFiles; cv::Mutex ProgramCache::mutexCache; -std::auto_ptr _programCache; +ProgramCache* _programCache = NULL; ProgramCache* ProgramCache::getProgramCache() { - if (NULL == _programCache.get()) - _programCache.reset(new ProgramCache()); - return _programCache.get(); + if (NULL == _programCache) + { + cv::AutoLock lock(getInitializationMutex()); + if (NULL == _programCache) + _programCache = new ProgramCache(); + } + return _programCache; } ProgramCache::ProgramCache() @@ -78,6 +82,12 @@ ProgramCache::ProgramCache() ProgramCache::~ProgramCache() { releaseProgram(); + if (this == _programCache) + { + cv::AutoLock lock(getInitializationMutex()); + if (this == _programCache) + _programCache = NULL; + } } cl_program ProgramCache::progLookup(const String& srcsign) @@ -420,22 +430,17 @@ struct ProgramFileCache { if(status == CL_BUILD_PROGRAM_FAILURE) { - cl_int logStatus; - char *buildLog = NULL; size_t buildLogSize = 0; - logStatus = clGetProgramBuildInfo(program, - getClDeviceID(ctx), CL_PROGRAM_BUILD_LOG, buildLogSize, - buildLog, &buildLogSize); - if(logStatus != CL_SUCCESS) - std::cout << "Failed to build the program and get the build info." << std::endl; - buildLog = new char[buildLogSize]; - CV_DbgAssert(!!buildLog); - memset(buildLog, 0, buildLogSize); openCLSafeCall(clGetProgramBuildInfo(program, getClDeviceID(ctx), - CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, NULL)); - std::cout << "\nBUILD LOG: " << options << "\n"; - std::cout << buildLog << std::endl; - delete [] buildLog; + CL_PROGRAM_BUILD_LOG, 0, NULL, &buildLogSize)); + std::vector buildLog; buildLog.resize(buildLogSize); + memset(&buildLog[0], 0, buildLogSize); + openCLSafeCall(clGetProgramBuildInfo(program, getClDeviceID(ctx), + CL_PROGRAM_BUILD_LOG, buildLogSize, &buildLog[0], NULL)); + std::cout << std::endl << "BUILD LOG: " + << (source->name ? source->name : "dynamic program") << ": " + << options << "\n"; + std::cout << &buildLog[0] << std::endl; } openCLVerifyCall(status); } diff --git a/modules/ocl/src/cl_programcache.hpp b/modules/ocl/src/cl_programcache.hpp index e96bd57b68..ebf3e76760 100644 --- a/modules/ocl/src/cl_programcache.hpp +++ b/modules/ocl/src/cl_programcache.hpp @@ -24,7 +24,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -52,7 +52,6 @@ class ProgramCache protected: ProgramCache(); ~ProgramCache(); - friend class std::auto_ptr; public: static ProgramCache *getProgramCache(); diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 9b91802d85..eec103a6dc 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/columnsum.cpp b/modules/ocl/src/columnsum.cpp index 33f4b00476..ccbd960bc7 100644 --- a/modules/ocl/src/columnsum.cpp +++ b/modules/ocl/src/columnsum.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/error.cpp b/modules/ocl/src/error.cpp index 9689e0a306..a1e2d807de 100644 --- a/modules/ocl/src/error.cpp +++ b/modules/ocl/src/error.cpp @@ -24,7 +24,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/fft.cpp b/modules/ocl/src/fft.cpp index 0fbcecc5c3..e10d671528 100644 --- a/modules/ocl/src/fft.cpp +++ b/modules/ocl/src/fft.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -126,7 +126,8 @@ void cv::ocl::fft_setup() { return; } - pCache.setupData = new clAmdFftSetupData; + if (pCache.setupData == NULL) + pCache.setupData = new clAmdFftSetupData; openCLSafeCall(clAmdFftInitSetupData( pCache.setupData )); pCache.started = true; } diff --git a/modules/ocl/src/filtering.cpp b/modules/ocl/src/filtering.cpp index 984f3a9f92..816988db78 100644 --- a/modules/ocl/src/filtering.cpp +++ b/modules/ocl/src/filtering.cpp @@ -29,7 +29,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -1059,74 +1059,39 @@ template <> struct index_and_sizeof template void linearRowFilter_gpu(const oclMat &src, const oclMat &dst, oclMat mat_kernel, int ksize, int anchor, int bordertype) { - Context *clCxt = src.clCxt; + CV_Assert(bordertype <= BORDER_REFLECT_101); + CV_Assert(ksize == (anchor << 1) + 1); int channels = src.oclchannels(); - size_t localThreads[3] = {16, 16, 1}; - String kernelName = "row_filter"; - - char btype[30]; + size_t localThreads[3] = { 16, 16, 1 }; + size_t globalThreads[3] = { dst.cols, dst.rows, 1 }; - switch (bordertype) - { - case 0: - sprintf(btype, "BORDER_CONSTANT"); - break; - case 1: - sprintf(btype, "BORDER_REPLICATE"); - break; - case 2: - sprintf(btype, "BORDER_REFLECT"); - break; - case 3: - sprintf(btype, "BORDER_WRAP"); - break; - case 4: - sprintf(btype, "BORDER_REFLECT_101"); - break; - } - - char compile_option[128]; - sprintf(compile_option, "-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s", anchor, (int)localThreads[0], (int)localThreads[1], channels, btype); - - size_t globalThreads[3]; - globalThreads[1] = (dst.rows + localThreads[1] - 1) / localThreads[1] * localThreads[1]; - globalThreads[2] = (1 + localThreads[2] - 1) / localThreads[2] * localThreads[2]; + const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" }; + std::string buildOptions = format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s", + anchor, (int)localThreads[0], (int)localThreads[1], channels, borderMap[bordertype]); if (src.depth() == CV_8U) { switch (channels) { case 1: - case 3: - globalThreads[0] = ((dst.cols + 4) / 4 + localThreads[0] - 1) / localThreads[0] * localThreads[0]; + globalThreads[0] = (dst.cols + 3) >> 2; break; case 2: - globalThreads[0] = ((dst.cols + 1) / 2 + localThreads[0] - 1) / localThreads[0] * localThreads[0]; + globalThreads[0] = (dst.cols + 1) >> 1; break; case 4: - globalThreads[0] = (dst.cols + localThreads[0] - 1) / localThreads[0] * localThreads[0]; + globalThreads[0] = dst.cols; break; } } - else - { - globalThreads[0] = (dst.cols + localThreads[0] - 1) / localThreads[0] * localThreads[0]; - } - //sanity checks - CV_Assert(clCxt == dst.clCxt); - CV_Assert(src.cols == dst.cols); - CV_Assert(src.oclchannels() == dst.oclchannels()); - CV_Assert(ksize == (anchor << 1) + 1); - int src_pix_per_row, dst_pix_per_row; - int src_offset_x, src_offset_y;//, dst_offset_in_pixel; - src_pix_per_row = src.step / src.elemSize(); - src_offset_x = (src.offset % src.step) / src.elemSize(); - src_offset_y = src.offset / src.step; - dst_pix_per_row = dst.step / dst.elemSize(); - //dst_offset_in_pixel = dst.offset / dst.elemSize(); + int src_pix_per_row = src.step / src.elemSize(); + int src_offset_x = (src.offset % src.step) / src.elemSize(); + int src_offset_y = src.offset / src.step; + int dst_pix_per_row = dst.step / dst.elemSize(); int ridusy = (dst.rows - src.rows) >> 1; + std::vector > args; args.push_back(std::make_pair(sizeof(cl_mem), &src.data)); args.push_back(std::make_pair(sizeof(cl_mem), &dst.data)); @@ -1141,7 +1106,8 @@ void linearRowFilter_gpu(const oclMat &src, const oclMat &dst, oclMat mat_kernel args.push_back(std::make_pair(sizeof(cl_int), (void *)&ridusy)); args.push_back(std::make_pair(sizeof(cl_mem), (void *)&mat_kernel.data)); - openCLExecuteKernel(clCxt, &filter_sep_row, kernelName, globalThreads, localThreads, args, channels, src.depth(), compile_option); + openCLExecuteKernel(src.clCxt, &filter_sep_row, "row_filter", globalThreads, localThreads, + args, channels, src.depth(), buildOptions.c_str()); } Ptr cv::ocl::getLinearRowFilter_GPU(int srcType, int /*bufType*/, const Mat &rowKernel, int anchor, int bordertype) diff --git a/modules/ocl/src/gemm.cpp b/modules/ocl/src/gemm.cpp index 504b697a76..92119df5e7 100644 --- a/modules/ocl/src/gemm.cpp +++ b/modules/ocl/src/gemm.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/gftt.cpp b/modules/ocl/src/gftt.cpp index 28af01dc2f..b072865539 100644 --- a/modules/ocl/src/gftt.cpp +++ b/modules/ocl/src/gftt.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/haar.cpp b/modules/ocl/src/haar.cpp index 15dcbf9b6a..8116496362 100644 --- a/modules/ocl/src/haar.cpp +++ b/modules/ocl/src/haar.cpp @@ -29,7 +29,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/hog.cpp b/modules/ocl/src/hog.cpp index 7427620f6d..70fe991875 100644 --- a/modules/ocl/src/hog.cpp +++ b/modules/ocl/src/hog.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/imgproc.cpp b/modules/ocl/src/imgproc.cpp index 26c5052eac..ed39868c76 100644 --- a/modules/ocl/src/imgproc.cpp +++ b/modules/ocl/src/imgproc.cpp @@ -36,7 +36,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -437,7 +437,7 @@ namespace cv CV_Assert(top >= 0 && bottom >= 0 && left >= 0 && right >= 0); - if( _src.offset != 0 && (bordertype & BORDER_ISOLATED) == 0 ) + if( (_src.wholecols != _src.cols || _src.wholerows != _src.rows) && (bordertype & BORDER_ISOLATED) == 0 ) { Size wholeSize; Point ofs; @@ -454,34 +454,25 @@ namespace cv } bordertype &= ~cv::BORDER_ISOLATED; - // TODO need to remove this conditions and fix the code - if (bordertype == cv::BORDER_REFLECT || bordertype == cv::BORDER_WRAP) - { - CV_Assert((_src.cols >= left) && (_src.cols >= right) && (_src.rows >= top) && (_src.rows >= bottom)); - } - else if (bordertype == cv::BORDER_REFLECT_101) - { - CV_Assert((_src.cols > left) && (_src.cols > right) && (_src.rows > top) && (_src.rows > bottom)); - } - dst.create(_src.rows + top + bottom, _src.cols + left + right, _src.type()); - int srcStep = _src.step1() / _src.oclchannels(), dstStep = dst.step1() / dst.oclchannels(); + int srcStep = _src.step / _src.elemSize(), dstStep = dst.step / dst.elemSize(); int srcOffset = _src.offset / _src.elemSize(), dstOffset = dst.offset / dst.elemSize(); int depth = _src.depth(), ochannels = _src.oclchannels(); - int __bordertype[] = {cv::BORDER_CONSTANT, cv::BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101}; - const char *borderstr[] = {"BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101"}; - size_t bordertype_index; + int __bordertype[] = { BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101 }; + const char *borderstr[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" }; - for(bordertype_index = 0; bordertype_index < sizeof(__bordertype) / sizeof(int); bordertype_index++) - if (__bordertype[bordertype_index] == bordertype) + int bordertype_index = -1; + for (int i = 0, end = sizeof(__bordertype) / sizeof(int); i < end; i++) + if (__bordertype[i] == bordertype) + { + bordertype_index = i; break; - - if (bordertype_index == sizeof(__bordertype) / sizeof(int)) + } + if (bordertype_index < 0) CV_Error(Error::StsBadArg, "Unsupported border type"); - String kernelName = "copymakeborder"; - size_t localThreads[3] = {16, 16, 1}; + size_t localThreads[3] = { 16, 16, 1 }; size_t globalThreads[3] = { dst.cols, dst.rows, 1 }; std::vector< std::pair > args; @@ -504,12 +495,6 @@ namespace cv typeMap[depth], channelMap[ochannels], borderstr[bordertype_index]); - if (src.type() == CV_8UC1 && (dst.offset & 3) == 0 && (dst.cols & 3) == 0) - { - kernelName = "copymakeborder_C1_D0"; - globalThreads[0] = dst.cols >> 2; - } - int cn = src.channels(), ocn = src.oclchannels(); int bufSize = src.elemSize1() * ocn; AutoBuffer _buf(bufSize); @@ -519,7 +504,7 @@ namespace cv args.push_back( std::make_pair( bufSize , (void *)buf )); - openCLExecuteKernel(src.clCxt, &imgproc_copymakeboder, kernelName, globalThreads, + openCLExecuteKernel(src.clCxt, &imgproc_copymakeboder, "copymakeborder", globalThreads, localThreads, args, -1, -1, buildOptions.c_str()); } diff --git a/modules/ocl/src/interpolate_frames.cpp b/modules/ocl/src/interpolate_frames.cpp index 78669bae10..47d6c837a3 100644 --- a/modules/ocl/src/interpolate_frames.cpp +++ b/modules/ocl/src/interpolate_frames.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/kalman.cpp b/modules/ocl/src/kalman.cpp index 6f8243457c..5a133a7b1c 100644 --- a/modules/ocl/src/kalman.cpp +++ b/modules/ocl/src/kalman.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/kmeans.cpp b/modules/ocl/src/kmeans.cpp index c5a03bacd9..5486aa495a 100644 --- a/modules/ocl/src/kmeans.cpp +++ b/modules/ocl/src/kmeans.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/knearest.cpp b/modules/ocl/src/knearest.cpp index 9b77f7a580..143e7aa7a0 100644 --- a/modules/ocl/src/knearest.cpp +++ b/modules/ocl/src/knearest.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/match_template.cpp b/modules/ocl/src/match_template.cpp index 9720c74214..b822aaafc0 100644 --- a/modules/ocl/src/match_template.cpp +++ b/modules/ocl/src/match_template.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/matrix_operations.cpp b/modules/ocl/src/matrix_operations.cpp index 68e42970af..021c7a3647 100644 --- a/modules/ocl/src/matrix_operations.cpp +++ b/modules/ocl/src/matrix_operations.cpp @@ -27,7 +27,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/mcwutil.cpp b/modules/ocl/src/mcwutil.cpp index 866dbbef6e..f6efbf7adc 100644 --- a/modules/ocl/src/mcwutil.cpp +++ b/modules/ocl/src/mcwutil.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/moments.cpp b/modules/ocl/src/moments.cpp index a5f684415e..6372364dd4 100644 --- a/modules/ocl/src/moments.cpp +++ b/modules/ocl/src/moments.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/mssegmentation.cpp b/modules/ocl/src/mssegmentation.cpp index 0b31f16ff5..163e017c4b 100644 --- a/modules/ocl/src/mssegmentation.cpp +++ b/modules/ocl/src/mssegmentation.cpp @@ -24,7 +24,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_LUT.cl b/modules/ocl/src/opencl/arithm_LUT.cl index ff21e9a315..658e1f4bc1 100644 --- a/modules/ocl/src/opencl/arithm_LUT.cl +++ b/modules/ocl/src/opencl/arithm_LUT.cl @@ -16,7 +16,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl b/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl index 0208806069..fcf38749d8 100644 --- a/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl +++ b/modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_add.cl b/modules/ocl/src/opencl/arithm_add.cl index 2f34bbbffe..a73b65da6f 100644 --- a/modules/ocl/src/opencl/arithm_add.cl +++ b/modules/ocl/src/opencl/arithm_add.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_addWeighted.cl b/modules/ocl/src/opencl/arithm_addWeighted.cl index 159a970db4..8272806e2b 100644 --- a/modules/ocl/src/opencl/arithm_addWeighted.cl +++ b/modules/ocl/src/opencl/arithm_addWeighted.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_add_mask.cl b/modules/ocl/src/opencl/arithm_add_mask.cl index c3958bf1fb..ea96d8a8a2 100644 --- a/modules/ocl/src/opencl/arithm_add_mask.cl +++ b/modules/ocl/src/opencl/arithm_add_mask.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_add_scalar.cl b/modules/ocl/src/opencl/arithm_add_scalar.cl index 7f4e413277..b82eff5954 100644 --- a/modules/ocl/src/opencl/arithm_add_scalar.cl +++ b/modules/ocl/src/opencl/arithm_add_scalar.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_add_scalar_mask.cl b/modules/ocl/src/opencl/arithm_add_scalar_mask.cl index b93de0c6b2..0762b19b10 100644 --- a/modules/ocl/src/opencl/arithm_add_scalar_mask.cl +++ b/modules/ocl/src/opencl/arithm_add_scalar_mask.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_bitwise_binary.cl b/modules/ocl/src/opencl/arithm_bitwise_binary.cl index a4fa205c1f..56cd745d29 100644 --- a/modules/ocl/src/opencl/arithm_bitwise_binary.cl +++ b/modules/ocl/src/opencl/arithm_bitwise_binary.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_bitwise_binary_mask.cl b/modules/ocl/src/opencl/arithm_bitwise_binary_mask.cl index d244e572d9..328ccd91ae 100644 --- a/modules/ocl/src/opencl/arithm_bitwise_binary_mask.cl +++ b/modules/ocl/src/opencl/arithm_bitwise_binary_mask.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_bitwise_binary_scalar.cl b/modules/ocl/src/opencl/arithm_bitwise_binary_scalar.cl index 5a7d5938cb..434bd5eca8 100644 --- a/modules/ocl/src/opencl/arithm_bitwise_binary_scalar.cl +++ b/modules/ocl/src/opencl/arithm_bitwise_binary_scalar.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_bitwise_not.cl b/modules/ocl/src/opencl/arithm_bitwise_not.cl index 714220ddf4..e5b46c9368 100644 --- a/modules/ocl/src/opencl/arithm_bitwise_not.cl +++ b/modules/ocl/src/opencl/arithm_bitwise_not.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_cartToPolar.cl b/modules/ocl/src/opencl/arithm_cartToPolar.cl index a2f65e0b73..6c779ead90 100644 --- a/modules/ocl/src/opencl/arithm_cartToPolar.cl +++ b/modules/ocl/src/opencl/arithm_cartToPolar.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_compare.cl b/modules/ocl/src/opencl/arithm_compare.cl index d0842db180..005d3c73f5 100644 --- a/modules/ocl/src/opencl/arithm_compare.cl +++ b/modules/ocl/src/opencl/arithm_compare.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_exp.cl b/modules/ocl/src/opencl/arithm_exp.cl index b2143ba142..835bc95c3d 100644 --- a/modules/ocl/src/opencl/arithm_exp.cl +++ b/modules/ocl/src/opencl/arithm_exp.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_flip.cl b/modules/ocl/src/opencl/arithm_flip.cl index 49242d07c7..7c2a04d74f 100644 --- a/modules/ocl/src/opencl/arithm_flip.cl +++ b/modules/ocl/src/opencl/arithm_flip.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_flip_rc.cl b/modules/ocl/src/opencl/arithm_flip_rc.cl index 68e26a439e..4a20382755 100644 --- a/modules/ocl/src/opencl/arithm_flip_rc.cl +++ b/modules/ocl/src/opencl/arithm_flip_rc.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_log.cl b/modules/ocl/src/opencl/arithm_log.cl index ef8c4dd04e..fe1b3046a3 100644 --- a/modules/ocl/src/opencl/arithm_log.cl +++ b/modules/ocl/src/opencl/arithm_log.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_nonzero.cl b/modules/ocl/src/opencl/arithm_nonzero.cl index 921367b3df..085386f5c3 100644 --- a/modules/ocl/src/opencl/arithm_nonzero.cl +++ b/modules/ocl/src/opencl/arithm_nonzero.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_phase.cl b/modules/ocl/src/opencl/arithm_phase.cl index a30eba4310..b6bc7b42b4 100644 --- a/modules/ocl/src/opencl/arithm_phase.cl +++ b/modules/ocl/src/opencl/arithm_phase.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_pow.cl b/modules/ocl/src/opencl/arithm_pow.cl index dd9250c7f9..1704f6b42e 100644 --- a/modules/ocl/src/opencl/arithm_pow.cl +++ b/modules/ocl/src/opencl/arithm_pow.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_setidentity.cl b/modules/ocl/src/opencl/arithm_setidentity.cl index 921026b40d..fb684c367f 100644 --- a/modules/ocl/src/opencl/arithm_setidentity.cl +++ b/modules/ocl/src/opencl/arithm_setidentity.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_sum.cl b/modules/ocl/src/opencl/arithm_sum.cl index 39bcf949a0..6eb6e48323 100644 --- a/modules/ocl/src/opencl/arithm_sum.cl +++ b/modules/ocl/src/opencl/arithm_sum.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/arithm_transpose.cl b/modules/ocl/src/opencl/arithm_transpose.cl index 5328d1f1b2..bd06a52083 100644 --- a/modules/ocl/src/opencl/arithm_transpose.cl +++ b/modules/ocl/src/opencl/arithm_transpose.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/bgfg_mog.cl b/modules/ocl/src/opencl/bgfg_mog.cl index 2e269999ac..8621ff31b0 100644 --- a/modules/ocl/src/opencl/bgfg_mog.cl +++ b/modules/ocl/src/opencl/bgfg_mog.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/brute_force_match.cl b/modules/ocl/src/opencl/brute_force_match.cl index a05c98ee03..d6a89f2057 100644 --- a/modules/ocl/src/opencl/brute_force_match.cl +++ b/modules/ocl/src/opencl/brute_force_match.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/build_warps.cl b/modules/ocl/src/opencl/build_warps.cl index 07cccee1a3..4402e8c388 100644 --- a/modules/ocl/src/opencl/build_warps.cl +++ b/modules/ocl/src/opencl/build_warps.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/convertC3C4.cl b/modules/ocl/src/opencl/convertC3C4.cl index 1908f92a2a..b3e699dc47 100644 --- a/modules/ocl/src/opencl/convertC3C4.cl +++ b/modules/ocl/src/opencl/convertC3C4.cl @@ -15,7 +15,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 2b1cfccd03..fcbf67ca7a 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/filter_sep_col.cl b/modules/ocl/src/opencl/filter_sep_col.cl index 60ce51360e..8dd77d5a97 100644 --- a/modules/ocl/src/opencl/filter_sep_col.cl +++ b/modules/ocl/src/opencl/filter_sep_col.cl @@ -47,36 +47,6 @@ #define READ_TIMES_ROW ((2*(RADIUS+LSIZE0)-1)/LSIZE0) #endif -#ifdef BORDER_CONSTANT -//BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii -#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2) -#endif - -#ifdef BORDER_REPLICATE -//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 -#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 -#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 - -#ifdef BORDER_WRAP -//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 - - /********************************************************************************** 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 @@ -107,15 +77,16 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter { int x = get_global_id(0); int y = get_global_id(1); + 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 i; - GENTYPE_SRC sum; - GENTYPE_SRC temp[READ_TIMES_COL]; - __local GENTYPE_SRC LDS_DAT[LSIZE1*READ_TIMES_COL][LSIZE0+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 i; + GENTYPE_SRC sum, temp[READ_TIMES_COL]; + __local GENTYPE_SRC LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1]; //read pixels from src for(i = 0;i= (r_edge) ? (elem1) : (elem2) -#endif - -#ifdef BORDER_REPLICATE -//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh -#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (l_edge) : (addr) -#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (r_edge)-1 : (addr) -#endif - +#elif defined BORDER_REPLICATE +#define EXTRAPOLATE(x, maxV) \ + { \ + x = max(min(x, maxV - 1), 0); \ + } +#elif defined BORDER_WRAP +#define EXTRAPOLATE(x, maxV) \ + { \ + if (x < 0) \ + x -= ((x - maxV + 1) / maxV) * maxV; \ + if (x >= maxV) \ + x %= maxV; \ + } +#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101) +#define EXTRAPOLATE_(x, maxV, delta) \ + { \ + if (maxV == 1) \ + x = 0; \ + else \ + do \ + { \ + if ( x < 0 ) \ + x = -x - 1 + delta; \ + else \ + x = maxV - 1 - (x - maxV) - delta; \ + } \ + while (x >= maxV || x < 0); \ + } #ifdef BORDER_REFLECT -//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb -#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i)-1 : (addr) -#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 -#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i) : (addr) -#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr) +#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 0) +#else +#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 1) #endif - -#ifdef BORDER_WRAP -//BORDER_WRAP: cdefgh|abcdefgh|abcdefg -#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (i)+(r_edge) : (addr) -#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (i)-(r_edge) : (addr) +#else +#error No extrapolation method #endif /********************************************************************************** @@ -96,73 +105,71 @@ The info above maybe obsolete. ***********************************************************************************/ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C1_D0 -(__global const uchar * restrict src, - __global float * dst, - const int dst_cols, - const int dst_rows, - const int src_whole_cols, - const int src_whole_rows, - const int src_step_in_pixel, - const int src_offset_x, - const int src_offset_y, - const int dst_step_in_pixel, - const int radiusy, - __constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1))))) + (__global uchar * restrict src, + __global float * dst, + int dst_cols, int dst_rows, + int src_whole_cols, int src_whole_rows, + int src_step_in_pixel, + int src_offset_x, int src_offset_y, + int dst_step_in_pixel, int radiusy, + __constant float * mat_kernel __attribute__((max_constant_size(4*(2*RADIUSX+1))))) { 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 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 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]; #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 0)) ? current_addr : 0; temp[i] = *(__global uchar4*)&src[current_addr]; } - //judge if read out of boundary - for(i = 0; isrc_whole_cols)| (start_y<0) | (start_y >= src_whole_rows); int4 index[READ_TIMES_ROW]; int4 addr; int s_y; - if(not_all_in_range) + + if (not_all_in_range) { - //judge if read out of boundary - for(i = 0; i 0)) ? current_addr : 0; temp[i] = src[current_addr]; } + //judge if read out of boundary - for(i = 0; i 0)) ? current_addr : 0; temp[i] = src[current_addr]; } - //judge if read out of boundary - for(i = 0; i 0)) ? current_addr : 0; temp[i] = src[current_addr]; } - //judge if read out of boundary - for(i = 0; i= (r_edge) ? (r_edge)-1 : (addr)) -#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (t_edge) :(i)) -#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (b_edge)-1 :(addr)) -#endif - +#ifdef BORDER_CONSTANT +#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2) +#elif defined BORDER_REPLICATE +#define EXTRAPOLATE(x, maxV) \ + { \ + x = max(min(x, maxV - 1), 0); \ + } +#elif defined BORDER_WRAP +#define EXTRAPOLATE(x, maxV) \ + { \ + if (x < 0) \ + x -= ((x - maxV + 1) / maxV) * maxV; \ + if (x >= maxV) \ + x %= maxV; \ + } +#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101) +#define EXTRAPOLATE_(x, maxV, delta) \ + { \ + if (maxV == 1) \ + x = 0; \ + else \ + do \ + { \ + if ( x < 0 ) \ + x = -x - 1 + delta; \ + else \ + x = maxV - 1 - (x - maxV) - delta; \ + } \ + while (x >= maxV || x < 0); \ + } #ifdef BORDER_REFLECT -//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)) -#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i)-1 : (i)) -#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr)) -#endif - -#ifdef BORDER_REFLECT_101 -//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)) -#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i) : (i)) -#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr)) +#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 0) +#else +#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 1) #endif - -//blur function does not support BORDER_WRAP -#ifdef BORDER_WRAP -//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)) -#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (i)+(b_edge) : (i)) -#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (i)-(b_edge) : (addr)) +#else +#error No extrapolation method #endif __kernel void @@ -117,9 +122,7 @@ edgeEnhancingFilter_C4_D0( float4 tmp_sum[1+EXTRA]; for(int tmpint = 0; tmpint < 1+EXTRA; tmpint++) - { tmp_sum[tmpint] = (float4)(0,0,0,0); - } #ifdef BORDER_CONSTANT bool con; @@ -127,25 +130,18 @@ edgeEnhancingFilter_C4_D0( for(int j = 0; j < ksY+EXTRA; j++) { con = (startX+col >= 0 && startX+col < src_whole_cols && startY+j >= 0 && startY+j < src_whole_rows); - int cur_col = clamp(startX + col, 0, src_whole_cols); - if(con) - { + if (con) ss = src[(startY+j)*(src_step>>2) + cur_col]; - } data[j][col] = con ? ss : (uchar4)0; } #else for(int j= 0; j < ksY+EXTRA; j++) { - int selected_row; - int selected_col; - selected_row = ADDR_H(startY+j, 0, src_whole_rows); - selected_row = ADDR_B(startY+j, src_whole_rows, selected_row); - - selected_col = ADDR_L(startX+col, 0, src_whole_cols); - selected_col = ADDR_R(startX+col, src_whole_cols, selected_col); + int selected_row = startY+j, selected_col = startX+col; + EXTRAPOLATE(selected_row, src_whole_rows) + EXTRAPOLATE(selected_col, src_whole_cols) data[j][col] = src[selected_row * (src_step>>2) + selected_col]; } @@ -172,7 +168,6 @@ edgeEnhancingFilter_C4_D0( if(col < (THREADS-(ksX-1))) { int4 currVal; - int howManyAll = (2*anX+1)*(ksY); //find variance of all data @@ -187,15 +182,14 @@ edgeEnhancingFilter_C4_D0( sumVal =0; sumValSqr=0; for(int j = startLMj; j < endLMj; j++) - { for(int i=-anX; i<=anX; i++) { - currVal = convert_int4(data[j][col+anX+i]) ; + currVal = convert_int4(data[j][col+anX+i]); sumVal += currVal; sumValSqr += mul24(currVal, currVal); } - } + var[extraCnt] = convert_float4( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) / ( (float)(howManyAll*howManyAll) ) ; #else var[extraCnt] = (float4)(900.0, 900.0, 900.0, 0.0); @@ -228,17 +222,15 @@ edgeEnhancingFilter_C4_D0( weight = 1.0f; #endif #else - currVal = convert_int4(data[j][col+anX+i]) ; + currVal = convert_int4(data[j][col+anX+i]); currWRTCenter = currVal-currValCenter; #if VAR_PER_CHANNEL - weight = var[extraCnt] / (var[extraCnt] + convert_float4(currWRTCenter * currWRTCenter)) * (float4)(lut[lut_j*lut_step+anX+i]); - //weight.x = var[extraCnt].x / ( var[extraCnt].x + (float) mul24(currWRTCenter.x , currWRTCenter.x) ) ; - //weight.y = var[extraCnt].y / ( var[extraCnt].y + (float) mul24(currWRTCenter.y , currWRTCenter.y) ) ; - //weight.z = var[extraCnt].z / ( var[extraCnt].z + (float) mul24(currWRTCenter.z , currWRTCenter.z) ) ; - //weight.w = 0; + weight = var[extraCnt] / (var[extraCnt] + convert_float4(currWRTCenter * currWRTCenter)) * + (float4)(lut[lut_j*lut_step+anX+i]); #else - weight = 1.0f/(1.0f+( mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) + mul24(currWRTCenter.z, currWRTCenter.z))/(var.x+var.y+var.z)); + weight = 1.0f/(1.0f+( mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) + + mul24(currWRTCenter.z, currWRTCenter.z))/(var.x+var.y+var.z)); #endif #endif tmp_sum[extraCnt] += convert_float4(data[j][col+anX+i]) * weight; @@ -249,9 +241,7 @@ edgeEnhancingFilter_C4_D0( tmp_sum[extraCnt] /= totalWeight; if(posX >= 0 && posX < dst_cols && (posY+extraCnt) >= 0 && (posY+extraCnt) < dst_rows) - { dst[(dst_startY+extraCnt) * (dst_step>>2)+ dst_startX + col] = convert_uchar4(tmp_sum[extraCnt]); - } #if VAR_PER_CHANNEL totalWeight = (float4)(0,0,0,0); @@ -323,13 +313,9 @@ edgeEnhancingFilter_C1_D0( #else for(int j= 0; j < ksY+EXTRA; j++) { - int selected_row; - int selected_col; - selected_row = ADDR_H(startY+j, 0, src_whole_rows); - selected_row = ADDR_B(startY+j, src_whole_rows, selected_row); - - selected_col = ADDR_L(startX+col, 0, src_whole_cols); - selected_col = ADDR_R(startX+col, src_whole_cols, selected_col); + int selected_row = startY+j, selected_col = startX+col; + EXTRAPOLATE(selected_row, src_whole_rows) + EXTRAPOLATE(selected_col, src_whole_cols) data[j][col] = src[selected_row * (src_step) + selected_col]; } diff --git a/modules/ocl/src/opencl/filtering_boxFilter.cl b/modules/ocl/src/opencl/filtering_boxFilter.cl index d163ebe76a..030c13cc57 100644 --- a/modules/ocl/src/opencl/filtering_boxFilter.cl +++ b/modules/ocl/src/opencl/filtering_boxFilter.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/filtering_laplacian.cl b/modules/ocl/src/opencl/filtering_laplacian.cl index 7c5b0c321b..ea22967dff 100644 --- a/modules/ocl/src/opencl/filtering_laplacian.cl +++ b/modules/ocl/src/opencl/filtering_laplacian.cl @@ -27,7 +27,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/filtering_morph.cl b/modules/ocl/src/opencl/filtering_morph.cl index db88aca654..c402ff7210 100644 --- a/modules/ocl/src/opencl/filtering_morph.cl +++ b/modules/ocl/src/opencl/filtering_morph.cl @@ -17,7 +17,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/haarobjectdetect.cl b/modules/ocl/src/opencl/haarobjectdetect.cl index 22a7fe7cbf..5fa3533054 100644 --- a/modules/ocl/src/opencl/haarobjectdetect.cl +++ b/modules/ocl/src/opencl/haarobjectdetect.cl @@ -19,7 +19,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/haarobjectdetect_scaled2.cl b/modules/ocl/src/opencl/haarobjectdetect_scaled2.cl index c12ab59506..17e95b4e4a 100644 --- a/modules/ocl/src/opencl/haarobjectdetect_scaled2.cl +++ b/modules/ocl/src/opencl/haarobjectdetect_scaled2.cl @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/imgproc_canny.cl b/modules/ocl/src/opencl/imgproc_canny.cl index 5402759e3c..ca670b6db7 100644 --- a/modules/ocl/src/opencl/imgproc_canny.cl +++ b/modules/ocl/src/opencl/imgproc_canny.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/imgproc_clahe.cl b/modules/ocl/src/opencl/imgproc_clahe.cl index 55692ae3b7..16c68fd474 100644 --- a/modules/ocl/src/opencl/imgproc_clahe.cl +++ b/modules/ocl/src/opencl/imgproc_clahe.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/imgproc_columnsum.cl b/modules/ocl/src/opencl/imgproc_columnsum.cl index 1609d7c552..6b596a3228 100644 --- a/modules/ocl/src/opencl/imgproc_columnsum.cl +++ b/modules/ocl/src/opencl/imgproc_columnsum.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/imgproc_convolve.cl b/modules/ocl/src/opencl/imgproc_convolve.cl index db7a7dfc3e..fb9596e5d6 100644 --- a/modules/ocl/src/opencl/imgproc_convolve.cl +++ b/modules/ocl/src/opencl/imgproc_convolve.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/src/opencl/imgproc_copymakeboder.cl b/modules/ocl/src/opencl/imgproc_copymakeboder.cl index ff7509ffdc..d97f660688 100644 --- a/modules/ocl/src/opencl/imgproc_copymakeboder.cl +++ b/modules/ocl/src/opencl/imgproc_copymakeboder.cl @@ -16,7 +16,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -35,173 +35,100 @@ // #if defined (DOUBLE_SUPPORT) -#ifdef cl_khr_fp64 -#pragma OPENCL EXTENSION cl_khr_fp64:enable -#elif defined (cl_amd_fp64) +#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 #ifdef BORDER_CONSTANT -//BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii -#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2) -#endif - -#ifdef BORDER_REPLICATE -//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh -#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (l_edge) : (addr) -#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (r_edge)-1 : (addr) -#endif - +#define EXTRAPOLATE(x, y, v) v = scalar; +#elif defined BORDER_REPLICATE +#define EXTRAPOLATE(x, y, v) \ + { \ + x = max(min(x, src_cols - 1), 0); \ + y = max(min(y, src_rows - 1), 0); \ + v = src[mad24(y, src_step, x + src_offset)]; \ + } +#elif defined BORDER_WRAP +#define EXTRAPOLATE(x, y, v) \ + { \ + if (x < 0) \ + x -= ((x - src_cols + 1) / src_cols) * src_cols; \ + if (x >= src_cols) \ + x %= src_cols; \ + \ + if (y < 0) \ + y -= ((y - src_rows + 1) / src_rows) * src_rows; \ + if( y >= src_rows ) \ + y %= src_rows; \ + v = src[mad24(y, src_step, x + src_offset)]; \ + } +#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101) #ifdef BORDER_REFLECT -//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb -#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i)-1 : (addr) -#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr) +#define DELTA int delta = 0 +#else +#define DELTA int delta = 1 #endif - -#ifdef BORDER_REFLECT_101 -//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba -#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i) : (addr) -#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr) +#define EXTRAPOLATE(x, y, v) \ + { \ + DELTA; \ + if (src_cols == 1) \ + x = 0; \ + else \ + do \ + { \ + if( x < 0 ) \ + x = -x - 1 + delta; \ + else \ + x = src_cols - 1 - (x - src_cols) - delta; \ + } \ + while (x >= src_cols || x < 0); \ + \ + if (src_rows == 1) \ + y = 0; \ + else \ + do \ + { \ + if( y < 0 ) \ + y = -y - 1 + delta; \ + else \ + y = src_rows - 1 - (y - src_rows) - delta; \ + } \ + while (y >= src_rows || y < 0); \ + v = src[mad24(y, src_step, x + src_offset)]; \ + } +#else +#error No extrapolation method #endif -#ifdef BORDER_WRAP -//BORDER_WRAP: cdefgh|abcdefgh|abcdefg -#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (i)+(r_edge) : (addr) -#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (i)-(r_edge) : (addr) -#endif +#define NEED_EXTRAPOLATION(gx, gy) (gx >= src_cols || gy >= src_rows || gx < 0 || gy < 0) __kernel void copymakeborder (__global const GENTYPE *src, __global GENTYPE *dst, - const int dst_cols, - const int dst_rows, - const int src_cols, - const int src_rows, - const int src_step_in_pixel, - const int src_offset_in_pixel, - const int dst_step_in_pixel, - const int dst_offset_in_pixel, - const int top, - const int left, - const GENTYPE val - ) + int dst_cols, int dst_rows, + int src_cols, int src_rows, + int src_step, int src_offset, + int dst_step, int dst_offset, + int top, int left, GENTYPE scalar) { int x = get_global_id(0); int y = get_global_id(1); - int src_x = x-left; - int src_y = y-top; - int src_addr = mad24(src_y,src_step_in_pixel,src_x+src_offset_in_pixel); - int dst_addr = mad24(y,dst_step_in_pixel,x+dst_offset_in_pixel); - int con = (src_x >= 0) && (src_x < src_cols) && (src_y >= 0) && (src_y < src_rows); - if(con) - { - dst[dst_addr] = src[src_addr]; - } - else - { - #ifdef BORDER_CONSTANT - //write the result to dst - if((x= 0) && (src_x+3 < src_cols) && (src_y >= 0) && (src_y < src_rows); - if(con) + if (x < dst_cols && y < dst_rows) { - uchar4 tmp = vload4(0,src+src_addr); - *(__global uchar4*)(dst+dst_addr) = tmp; - } - else - { - #ifdef BORDER_CONSTANT - //write the result to dst - if((((src_x<0) && (src_x+3>=0))||(src_x < src_cols) && (src_x+3 >= src_cols)) && (src_y >= 0) && (src_y < src_rows)) - { - int4 addr; - uchar4 tmp; - addr.x = ((src_x < 0) || (src_x>= src_cols)) ? 0 : src_addr; - addr.y = ((src_x+1 < 0) || (src_x+1>= src_cols)) ? 0 : (src_addr+1); - addr.z = ((src_x+2 < 0) || (src_x+2>= src_cols)) ? 0 : (src_addr+2); - addr.w = ((src_x+3 < 0) || (src_x+3>= src_cols)) ? 0 : (src_addr+3); - tmp.x = src[addr.x]; - tmp.y = src[addr.y]; - tmp.z = src[addr.z]; - tmp.w = src[addr.w]; - tmp.x = (src_x >=0)&&(src_x < src_cols) ? tmp.x : val; - tmp.y = (src_x+1 >=0)&&(src_x +1 < src_cols) ? tmp.y : val; - tmp.z = (src_x+2 >=0)&&(src_x +2 < src_cols) ? tmp.z : val; - tmp.w = (src_x+3 >=0)&&(src_x +3 < src_cols) ? tmp.w : val; - *(__global uchar4*)(dst+dst_addr) = tmp; - } - else if((x= CV_32F ? 1e-3 : 1); @@ -1466,7 +1471,7 @@ OCL_TEST_P(Norm, NORM_L1) const double cpuRes = cv::norm(src1_roi, src2_roi, type); const double gpuRes = cv::ocl::norm(gsrc1_roi, gsrc2_roi, type); - EXPECT_NEAR(cpuRes, gpuRes, 0.1); + EXPECT_PRED3(relativeError, cpuRes, gpuRes, 1e-6); } } @@ -1484,7 +1489,7 @@ OCL_TEST_P(Norm, NORM_L2) const double cpuRes = cv::norm(src1_roi, src2_roi, type); const double gpuRes = cv::ocl::norm(gsrc1_roi, gsrc2_roi, type); - EXPECT_NEAR(cpuRes, gpuRes, 0.1); + EXPECT_PRED3(relativeError, cpuRes, gpuRes, 1e-6); } } diff --git a/modules/ocl/test/test_bgfg.cpp b/modules/ocl/test/test_bgfg.cpp index d2e1363017..8b4c865c3f 100644 --- a/modules/ocl/test/test_bgfg.cpp +++ b/modules/ocl/test/test_bgfg.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_blend.cpp b/modules/ocl/test/test_blend.cpp index 8e6e269399..63693749db 100644 --- a/modules/ocl/test/test_blend.cpp +++ b/modules/ocl/test/test_blend.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_brute_force_matcher.cpp b/modules/ocl/test/test_brute_force_matcher.cpp index 5b80449e28..d31b3715b5 100644 --- a/modules/ocl/test/test_brute_force_matcher.cpp +++ b/modules/ocl/test/test_brute_force_matcher.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_calib3d.cpp b/modules/ocl/test/test_calib3d.cpp index 532e61d134..9fd0b23298 100644 --- a/modules/ocl/test/test_calib3d.cpp +++ b/modules/ocl/test/test_calib3d.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_canny.cpp b/modules/ocl/test/test_canny.cpp index b7d2d6d44a..82286031f5 100644 --- a/modules/ocl/test/test_canny.cpp +++ b/modules/ocl/test/test_canny.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 886f3c1e59..cc7843db62 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_fft.cpp b/modules/ocl/test/test_fft.cpp index d9cc7b1595..20d88c21fe 100644 --- a/modules/ocl/test/test_fft.cpp +++ b/modules/ocl/test/test_fft.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_filters.cpp b/modules/ocl/test/test_filters.cpp index 2e54570e73..3caea04fb0 100644 --- a/modules/ocl/test/test_filters.cpp +++ b/modules/ocl/test/test_filters.cpp @@ -30,7 +30,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -403,7 +403,7 @@ INSTANTIATE_TEST_CASE_P(Filter, SobelTest, Combine( Bool())); INSTANTIATE_TEST_CASE_P(Filter, ScharrTest, Combine( - Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4), + Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), Values(0), // not used Values(Size(0, 1), Size(1, 0)), Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101, @@ -432,7 +432,7 @@ INSTANTIATE_TEST_CASE_P(Filter, Bilateral, Combine( Values(Size(0, 0)), // not used Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_WRAP, (int)BORDER_REFLECT_101), - Values(false))); // TODO does not work with ROI + Bool())); INSTANTIATE_TEST_CASE_P(Filter, AdaptiveBilateral, Combine( Values(CV_8UC1, CV_8UC3), diff --git a/modules/ocl/test/test_gemm.cpp b/modules/ocl/test/test_gemm.cpp index 68dab0ac01..c2a44842ca 100644 --- a/modules/ocl/test/test_gemm.cpp +++ b/modules/ocl/test/test_gemm.cpp @@ -24,7 +24,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_imgproc.cpp b/modules/ocl/test/test_imgproc.cpp index fa7a70f4d3..eb983fb17e 100644 --- a/modules/ocl/test/test_imgproc.cpp +++ b/modules/ocl/test/test_imgproc.cpp @@ -33,7 +33,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -144,7 +144,7 @@ PARAM_TEST_CASE(CopyMakeBorder, MatDepth, // depth generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder); generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder); - border = randomBorder(0, 10); + border = randomBorder(0, MAX_VALUE << 2); val = randomScalar(-MAX_VALUE, MAX_VALUE); } diff --git a/modules/ocl/test/test_kalman.cpp b/modules/ocl/test/test_kalman.cpp index f02df6af71..045cd98158 100644 --- a/modules/ocl/test/test_kalman.cpp +++ b/modules/ocl/test/test_kalman.cpp @@ -24,7 +24,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_kmeans.cpp b/modules/ocl/test/test_kmeans.cpp index 5e4026694d..94263d8f76 100644 --- a/modules/ocl/test/test_kmeans.cpp +++ b/modules/ocl/test/test_kmeans.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_match_template.cpp b/modules/ocl/test/test_match_template.cpp index 0c2e9bd4e0..edbc36a3f2 100644 --- a/modules/ocl/test/test_match_template.cpp +++ b/modules/ocl/test/test_match_template.cpp @@ -24,7 +24,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_matrix_operation.cpp b/modules/ocl/test/test_matrix_operation.cpp index 27a587260a..c7ceef4539 100644 --- a/modules/ocl/test/test_matrix_operation.cpp +++ b/modules/ocl/test/test_matrix_operation.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_mean_shift.cpp b/modules/ocl/test/test_mean_shift.cpp index 684a2a9378..6ee3e35a76 100644 --- a/modules/ocl/test/test_mean_shift.cpp +++ b/modules/ocl/test/test_mean_shift.cpp @@ -33,7 +33,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_ml.cpp b/modules/ocl/test/test_ml.cpp index 76940fff05..a064070890 100644 --- a/modules/ocl/test/test_ml.cpp +++ b/modules/ocl/test/test_ml.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_optflow.cpp b/modules/ocl/test/test_optflow.cpp index da38fe4716..7296a6b7ea 100644 --- a/modules/ocl/test/test_optflow.cpp +++ b/modules/ocl/test/test_optflow.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_pyramids.cpp b/modules/ocl/test/test_pyramids.cpp index 3b91f12ba3..2d861b627d 100644 --- a/modules/ocl/test/test_pyramids.cpp +++ b/modules/ocl/test/test_pyramids.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_sort.cpp b/modules/ocl/test/test_sort.cpp index 5a7d4a3004..b259149682 100644 --- a/modules/ocl/test/test_sort.cpp +++ b/modules/ocl/test/test_sort.cpp @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_split_merge.cpp b/modules/ocl/test/test_split_merge.cpp index b090045c7d..6148e95cb4 100644 --- a/modules/ocl/test/test_split_merge.cpp +++ b/modules/ocl/test/test_split_merge.cpp @@ -26,7 +26,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. diff --git a/modules/ocl/test/test_warp.cpp b/modules/ocl/test/test_warp.cpp index 717bbc7a2e..bfe5b638f3 100644 --- a/modules/ocl/test/test_warp.cpp +++ b/modules/ocl/test/test_warp.cpp @@ -33,7 +33,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -283,14 +283,21 @@ PARAM_TEST_CASE(Resize, MatType, double, double, Interpolation, bool) void random_roi() { - Size srcRoiSize = randomSize(1, MAX_VALUE); - Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); - randomSubMat(src, src_roi, srcRoiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE); + CV_Assert(fx > 0 && fy > 0); - Size dstRoiSize; + Size srcRoiSize = randomSize(1, MAX_VALUE), dstRoiSize; dstRoiSize.width = cvRound(srcRoiSize.width * fx); dstRoiSize.height = cvRound(srcRoiSize.height * fy); + if (dstRoiSize.area() == 0) + { + random_roi(); + return; + } + + Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, srcRoiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE); + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); randomSubMat(dst_whole, dst_roi, dstRoiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE); @@ -315,7 +322,7 @@ OCL_TEST_P(Resize, Mat) { random_roi(); - resize(src_roi, dst_roi, Size(), fx, fy, interpolation); + cv::resize(src_roi, dst_roi, Size(), fx, fy, interpolation); ocl::resize(gsrc_roi, gdst_roi, Size(), fx, fy, interpolation); Near(1.0); diff --git a/modules/ocl/test/utility.cpp b/modules/ocl/test/utility.cpp index 6e519991d7..8521dff821 100644 --- a/modules/ocl/test/utility.cpp +++ b/modules/ocl/test/utility.cpp @@ -230,4 +230,21 @@ double checkRectSimilarity(Size sz, std::vector& ob1, std::vector& o return final_test_result; } +void showDiff(const Mat& gold, const Mat& actual, double eps) +{ + Mat diff; + absdiff(gold, actual, diff); + threshold(diff, diff, eps, 255.0, cv::THRESH_BINARY); + + namedWindow("gold", WINDOW_NORMAL); + namedWindow("actual", WINDOW_NORMAL); + namedWindow("diff", WINDOW_NORMAL); + + imshow("gold", gold); + imshow("actual", actual); + imshow("diff", diff); + + waitKey(); +} + } // namespace cvtest diff --git a/modules/ocl/test/utility.hpp b/modules/ocl/test/utility.hpp index e578effca6..0f0fac3138 100644 --- a/modules/ocl/test/utility.hpp +++ b/modules/ocl/test/utility.hpp @@ -54,7 +54,7 @@ extern int LOOP_TIMES; namespace cvtest { -//void showDiff(cv::InputArray gold, cv::InputArray actual, double eps); +void showDiff(const Mat& gold, const Mat& actual, double eps); cv::ocl::oclMat createMat_ocl(cv::RNG& rng, Size size, int type, bool useRoi); cv::ocl::oclMat loadMat_ocl(cv::RNG& rng, const Mat& m, bool useRoi); diff --git a/modules/superres/src/btv_l1_ocl.cpp b/modules/superres/src/btv_l1_ocl.cpp index b5df8b5fb6..44edc815ec 100644 --- a/modules/superres/src/btv_l1_ocl.cpp +++ b/modules/superres/src/btv_l1_ocl.cpp @@ -70,6 +70,7 @@ namespace cv { float* btvWeights_ = NULL; size_t btvWeights_size = 0; + oclMat c_btvRegWeights; } } @@ -82,10 +83,6 @@ namespace btv_l1_device_ocl void upscale(const oclMat& src, oclMat& dst, int scale); - float diffSign(float a, float b); - - Point3f diffSign(Point3f a, Point3f b); - void diffSign(const oclMat& src1, const oclMat& src2, oclMat& dst); void calcBtvRegularization(const oclMat& src, oclMat& dst, int ksize); @@ -165,20 +162,6 @@ void btv_l1_device_ocl::upscale(const oclMat& src, oclMat& dst, int scale) } -float btv_l1_device_ocl::diffSign(float a, float b) -{ - return a > b ? 1.0f : a < b ? -1.0f : 0.0f; -} - -Point3f btv_l1_device_ocl::diffSign(Point3f a, Point3f b) -{ - return Point3f( - a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f, - a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f, - a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f - ); -} - void btv_l1_device_ocl::diffSign(const oclMat& src1, const oclMat& src2, oclMat& dst) { Context* clCxt = Context::getContext(); @@ -228,12 +211,6 @@ void btv_l1_device_ocl::calcBtvRegularization(const oclMat& src, oclMat& dst, in int cn = src.oclchannels(); - cl_mem c_btvRegWeights; - size_t count = btvWeights_size * sizeof(float); - c_btvRegWeights = openCLCreateBuffer(clCxt, CL_MEM_READ_ONLY, count); - int cl_safe_check = clEnqueueWriteBuffer(getClCommandQueue(clCxt), c_btvRegWeights, 1, 0, count, btvWeights_, 0, NULL, NULL); - CV_Assert(cl_safe_check == CL_SUCCESS); - args.push_back(make_pair(sizeof(cl_mem), (void*)&src_.data)); args.push_back(make_pair(sizeof(cl_mem), (void*)&dst_.data)); args.push_back(make_pair(sizeof(cl_int), (void*)&src_step)); @@ -242,11 +219,9 @@ void btv_l1_device_ocl::calcBtvRegularization(const oclMat& src, oclMat& dst, in args.push_back(make_pair(sizeof(cl_int), (void*)&src.cols)); args.push_back(make_pair(sizeof(cl_int), (void*)&ksize)); args.push_back(make_pair(sizeof(cl_int), (void*)&cn)); - args.push_back(make_pair(sizeof(cl_mem), (void*)&c_btvRegWeights)); + args.push_back(make_pair(sizeof(cl_mem), (void*)&c_btvRegWeights.data)); openCLExecuteKernel(clCxt, &superres_btvl1, kernel_name, global_thread, local_thread, args, -1, -1); - cl_safe_check = clReleaseMemObject(c_btvRegWeights); - CV_Assert(cl_safe_check == CL_SUCCESS); } namespace @@ -321,9 +296,6 @@ namespace { CV_Assert( src.channels() == 1 || src.channels() == 3 || src.channels() == 4 ); - dst.create(src.rows * scale, src.cols * scale, src.type()); - dst.setTo(Scalar::all(0)); - btv_l1_device_ocl::upscale(src, dst, scale); } @@ -351,12 +323,13 @@ namespace btvWeights_ = &btvWeights[0]; btvWeights_size = size; + Mat btvWeights_mheader(1, static_cast(size), CV_32FC1, btvWeights_); + c_btvRegWeights = btvWeights_mheader; } void calcBtvRegularization(const oclMat& src, oclMat& dst, int btvKernelSize) { dst.create(src.size(), src.type()); - dst.setTo(Scalar::all(0)); const int ksize = (btvKernelSize - 1) / 2; @@ -407,7 +380,7 @@ namespace oclMat highRes_; vector diffTerms_; - vector a_, b_, c_; + oclMat a_, b_, c_, d_; oclMat regTerm_; }; @@ -421,7 +394,7 @@ namespace btvKernelSize_ = 7; blurKernelSize_ = 5; blurSigma_ = 0.0; - opticalFlow_ = createOptFlow_DualTVL1_OCL(); + opticalFlow_ = createOptFlow_Farneback_OCL(); curBlurKernelSize_ = -1; curBlurSigma_ = -1.0; @@ -487,34 +460,36 @@ namespace // iterations diffTerms_.resize(src.size()); - a_.resize(src.size()); - b_.resize(src.size()); - c_.resize(src.size()); - + bool d_inited = false; + a_.create(highRes_.size(), highRes_.type()); + b_.create(highRes_.size(), highRes_.type()); + c_.create(lowResSize, highRes_.type()); + d_.create(highRes_.rows, highRes_.cols, highRes_.type()); for (int i = 0; i < iterations_; ++i) { + if(!d_inited) + { + d_.setTo(0); + d_inited = true; + } for (size_t k = 0; k < src.size(); ++k) { diffTerms_[k].create(highRes_.size(), highRes_.type()); - a_[k].create(highRes_.size(), highRes_.type()); - b_[k].create(highRes_.size(), highRes_.type()); - c_[k].create(lowResSize, highRes_.type()); - // a = M * Ih - ocl::remap(highRes_, a_[k], backwardMaps_[k].first, backwardMaps_[k].second, INTER_NEAREST, BORDER_CONSTANT, Scalar()); + ocl::remap(highRes_, a_, backwardMaps_[k].first, backwardMaps_[k].second, INTER_NEAREST, BORDER_CONSTANT, Scalar()); // b = HM * Ih - filters_[k]->apply(a_[k], b_[k], Rect(0,0,-1,-1)); + filters_[k]->apply(a_, b_, Rect(0,0,-1,-1)); // c = DHF * Ih - ocl::resize(b_[k], c_[k], lowResSize, 0, 0, INTER_NEAREST); + ocl::resize(b_, c_, lowResSize, 0, 0, INTER_NEAREST); - diffSign(src[k], c_[k], c_[k]); + diffSign(src[k], c_, c_); // a = Dt * diff - upscale(c_[k], a_[k], scale_); + upscale(c_, d_, scale_); // b = HtDt * diff - filters_[k]->apply(a_[k], b_[k], Rect(0,0,-1,-1)); + filters_[k]->apply(d_, b_, Rect(0,0,-1,-1)); // diffTerm = MtHtDt * diff - ocl::remap(b_[k], diffTerms_[k], forwardMaps_[k].first, forwardMaps_[k].second, INTER_NEAREST, BORDER_CONSTANT, Scalar()); + ocl::remap(b_, diffTerms_[k], forwardMaps_[k].first, forwardMaps_[k].second, INTER_NEAREST, BORDER_CONSTANT, Scalar()); } if (lambda_ > 0) @@ -549,10 +524,11 @@ namespace highRes_.release(); diffTerms_.clear(); - a_.clear(); - b_.clear(); - c_.clear(); + a_.release(); + b_.release(); + c_.release(); regTerm_.release(); + c_btvRegWeights.release(); } //////////////////////////////////////////////////////////// diff --git a/modules/superres/src/opencl/superres_btvl1.cl b/modules/superres/src/opencl/superres_btvl1.cl index 4720623234..3c0cff85b9 100644 --- a/modules/superres/src/opencl/superres_btvl1.cl +++ b/modules/superres/src/opencl/superres_btvl1.cl @@ -25,7 +25,7 @@ // // * Redistribution's in binary form must reproduce the above copyright notice, // this list of conditions and the following disclaimer in the documentation -// and/or other oclMaterials provided with the distribution. +// and/or other materials provided with the distribution. // // * The name of the copyright holders may not be used to endorse or promote products // derived from this software without specific prior written permission. @@ -44,24 +44,24 @@ //M*/ __kernel void buildMotionMapsKernel(__global float* forwardMotionX, - __global float* forwardMotionY, - __global float* backwardMotionX, - __global float* backwardMotionY, - __global float* forwardMapX, - __global float* forwardMapY, - __global float* backwardMapX, - __global float* backwardMapY, - int forwardMotionX_row, - int forwardMotionX_col, - int forwardMotionX_step, - int forwardMotionY_step, - int backwardMotionX_step, - int backwardMotionY_step, - int forwardMapX_step, - int forwardMapY_step, - int backwardMapX_step, - int backwardMapY_step - ) + __global float* forwardMotionY, + __global float* backwardMotionX, + __global float* backwardMotionY, + __global float* forwardMapX, + __global float* forwardMapY, + __global float* backwardMapX, + __global float* backwardMapY, + int forwardMotionX_row, + int forwardMotionX_col, + int forwardMotionX_step, + int forwardMotionY_step, + int backwardMotionX_step, + int backwardMotionY_step, + int forwardMapX_step, + int forwardMapY_step, + int backwardMapX_step, + int backwardMapY_step + ) { int x = get_global_id(0); int y = get_global_id(1); @@ -83,14 +83,14 @@ __kernel void buildMotionMapsKernel(__global float* forwardMotionX, } __kernel void upscaleKernel(__global float* src, - __global float* dst, - int src_step, - int dst_step, - int src_row, - int src_col, - int scale, - int channels - ) + __global float* dst, + int src_step, + int dst_step, + int src_row, + int src_col, + int scale, + int channels + ) { int x = get_global_id(0); int y = get_global_id(1); @@ -100,17 +100,10 @@ __kernel void upscaleKernel(__global float* src, if(channels == 1) { dst[y * scale * dst_step + x * scale] = src[y * src_step + x]; - }else if(channels == 3) - { - dst[y * channels * scale * dst_step + 3 * x * scale + 0] = src[y * channels * src_step + 3 * x + 0]; - dst[y * channels * scale * dst_step + 3 * x * scale + 1] = src[y * channels * src_step + 3 * x + 1]; - dst[y * channels * scale * dst_step + 3 * x * scale + 2] = src[y * channels * src_step + 3 * x + 2]; - }else + } + else { - dst[y * channels * scale * dst_step + 4 * x * scale + 0] = src[y * channels * src_step + 4 * x + 0]; - dst[y * channels * scale * dst_step + 4 * x * scale + 1] = src[y * channels * src_step + 4 * x + 1]; - dst[y * channels * scale * dst_step + 4 * x * scale + 2] = src[y * channels * src_step + 4 * x + 2]; - dst[y * channels * scale * dst_step + 4 * x * scale + 3] = src[y * channels * src_step + 4 * x + 3]; + vstore4(vload4(0, src + y * channels * src_step + 4 * x), 0, dst + y * channels * scale * dst_step + 4 * x * scale); } } } @@ -121,15 +114,6 @@ float diffSign(float a, float b) return a > b ? 1.0f : a < b ? -1.0f : 0.0f; } -float3 diffSign3(float3 a, float3 b) -{ - float3 pos; - pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f; - pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f; - pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f; - return pos; -} - float4 diffSign4(float4 a, float4 b) { float4 pos; @@ -141,13 +125,13 @@ float4 diffSign4(float4 a, float4 b) } __kernel void diffSignKernel(__global float* src1, - __global float* src2, - __global float* dst, - int src1_row, - int src1_col, - int dst_step, - int src1_step, - int src2_step) + __global float* src2, + __global float* dst, + int src1_row, + int src1_col, + int dst_step, + int src1_step, + int src2_step) { int x = get_global_id(0); int y = get_global_id(1); @@ -156,19 +140,18 @@ __kernel void diffSignKernel(__global float* src1, { dst[y * dst_step + x] = diffSign(src1[y * src1_step + x], src2[y * src2_step + x]); } - barrier(CLK_LOCAL_MEM_FENCE); } __kernel void calcBtvRegularizationKernel(__global float* src, - __global float* dst, - int src_step, - int dst_step, - int src_row, - int src_col, - int ksize, - int channels, - __global float* c_btvRegWeights - ) + __global float* dst, + int src_step, + int dst_step, + int src_row, + int src_col, + int ksize, + int channels, + __constant float* c_btvRegWeights + ) { int x = get_global_id(0) + ksize; int y = get_global_id(1) + ksize; @@ -180,57 +163,19 @@ __kernel void calcBtvRegularizationKernel(__global float* src, const float srcVal = src[y * src_step + x]; float dstVal = 0.0f; - for (int m = 0, count = 0; m <= ksize; ++m) - { - for (int l = ksize; l + m >= 0; --l, ++count) - dstVal = dstVal + c_btvRegWeights[count] * (diffSign(srcVal, src[(y + m) * src_step + (x + l)]) - diffSign(src[(y - m) * src_step + (x - l)], srcVal)); - } - dst[y * dst_step + x] = dstVal; - }else if(channels == 3) - { - float3 srcVal; - srcVal.x = src[y * src_step + 3 * x + 0]; - srcVal.y = src[y * src_step + 3 * x + 1]; - srcVal.z = src[y * src_step + 3 * x + 2]; - - float3 dstVal; - dstVal.x = 0.0f; - dstVal.y = 0.0f; - dstVal.z = 0.0f; - for (int m = 0, count = 0; m <= ksize; ++m) { for (int l = ksize; l + m >= 0; --l, ++count) { - float3 src1; - src1.x = src[(y + m) * src_step + 3 * (x + l) + 0]; - src1.y = src[(y + m) * src_step + 3 * (x + l) + 1]; - src1.z = src[(y + m) * src_step + 3 * (x + l) + 2]; - - float3 src2; - src2.x = src[(y - m) * src_step + 3 * (x - l) + 0]; - src2.y = src[(y - m) * src_step + 3 * (x - l) + 1]; - src2.z = src[(y - m) * src_step + 3 * (x - l) + 2]; - - dstVal = dstVal + c_btvRegWeights[count] * (diffSign3(srcVal, src1) - diffSign3(src2, srcVal)); + dstVal = dstVal + c_btvRegWeights[count] * (diffSign(srcVal, src[(y + m) * src_step + (x + l)]) - diffSign(src[(y - m) * src_step + (x - l)], srcVal)); } } - dst[y * dst_step + 3 * x + 0] = dstVal.x; - dst[y * dst_step + 3 * x + 1] = dstVal.y; - dst[y * dst_step + 3 * x + 2] = dstVal.z; - }else + dst[y * dst_step + x] = dstVal; + } + else { - float4 srcVal; - srcVal.x = src[y * src_step + 4 * x + 0];//r type =float - srcVal.y = src[y * src_step + 4 * x + 1];//g - srcVal.z = src[y * src_step + 4 * x + 2];//b - srcVal.w = src[y * src_step + 4 * x + 3];//a - - float4 dstVal; - dstVal.x = 0.0f; - dstVal.y = 0.0f; - dstVal.z = 0.0f; - dstVal.w = 0.0f; + float4 srcVal = vload4(0, src + y * src_step + 4 * x); + float4 dstVal = 0.f; for (int m = 0, count = 0; m <= ksize; ++m) { @@ -249,13 +194,9 @@ __kernel void calcBtvRegularizationKernel(__global float* src, src2.w = src[(y - m) * src_step + 4 * (x - l) + 3]; dstVal = dstVal + c_btvRegWeights[count] * (diffSign4(srcVal, src1) - diffSign4(src2, srcVal)); - } } - dst[y * dst_step + 4 * x + 0] = dstVal.x; - dst[y * dst_step + 4 * x + 1] = dstVal.y; - dst[y * dst_step + 4 * x + 2] = dstVal.z; - dst[y * dst_step + 4 * x + 3] = dstVal.w; + vstore4(dstVal, 0, dst + y * dst_step + 4 * x); } } } diff --git a/platforms/android/android.toolchain.cmake b/platforms/android/android.toolchain.cmake index bee73dbea1..a9d4d7c425 100644 --- a/platforms/android/android.toolchain.cmake +++ b/platforms/android/android.toolchain.cmake @@ -715,7 +715,7 @@ __INIT_VARIABLE( ANDROID_ABI OBSOLETE_ARM_TARGET OBSOLETE_ARM_TARGETS VALUES ${A # verify that target ABI is supported list( FIND ANDROID_SUPPORTED_ABIS "${ANDROID_ABI}" __androidAbiIdx ) if( __androidAbiIdx EQUAL -1 ) - string( REPLACE ";" "\", \"", PRINTABLE_ANDROID_SUPPORTED_ABIS "${ANDROID_SUPPORTED_ABIS}" ) + string( REPLACE ";" "\", \"" PRINTABLE_ANDROID_SUPPORTED_ABIS "${ANDROID_SUPPORTED_ABIS}" ) message( FATAL_ERROR "Specified ANDROID_ABI = \"${ANDROID_ABI}\" is not supported by this cmake toolchain or your NDK/toolchain. Supported values are: \"${PRINTABLE_ANDROID_SUPPORTED_ABIS}\" " ) diff --git a/samples/gpu/super_resolution.cpp b/samples/gpu/super_resolution.cpp index 12a061c686..67d0532a69 100644 --- a/samples/gpu/super_resolution.cpp +++ b/samples/gpu/super_resolution.cpp @@ -222,7 +222,11 @@ int main(int argc, const char* argv[]) if(useOcl) { - MEASURE_TIME(superRes->nextFrame(result_)); + MEASURE_TIME( + { + superRes->nextFrame(result_); + ocl::finish(); + }); } else #endif diff --git a/samples/ocl/adaptive_bilateral_filter.cpp b/samples/ocl/adaptive_bilateral_filter.cpp index f67c541098..c274d903a1 100644 --- a/samples/ocl/adaptive_bilateral_filter.cpp +++ b/samples/ocl/adaptive_bilateral_filter.cpp @@ -13,17 +13,27 @@ int main( int argc, const char** argv ) { const char* keys = "{ i input | | specify input image }" - "{ k ksize | 5 | specify kernel size }"; + "{ k ksize | 5 | specify kernel size }" + "{ h help | false | print help message }"; + CommandLineParser cmd(argc, argv, keys); + if (cmd.has("help")) + { + cout << "Usage : adaptive_bilateral_filter [options]" << endl; + cout << "Available options:" << endl; + cmd.printMessage(); + return EXIT_SUCCESS; + } + string src_path = cmd.get("i"); int ks = cmd.get("k"); const char * winName[] = {"input", "adaptive bilateral CPU", "adaptive bilateral OpenCL", "bilateralFilter OpenCL"}; - Mat src = imread(src_path); - Mat abFilterCPU; - if(src.empty()){ - //cout << "error read image: " << src_path << endl; - return -1; + Mat src = imread(src_path), abFilterCPU; + if (src.empty()) + { + cout << "error read image: " << src_path << endl; + return EXIT_FAILURE; } ocl::oclMat dsrc(src), dABFilter, dBFilter; @@ -33,17 +43,12 @@ int main( int argc, const char** argv ) ocl::adaptiveBilateralFilter(dsrc, dABFilter, ksize, 10); ocl::bilateralFilter(dsrc, dBFilter, ks, 30, 9); - Mat abFilter = dABFilter; - Mat bFilter = dBFilter; + Mat abFilter = dABFilter, bFilter = dBFilter; imshow(winName[0], src); - imshow(winName[1], abFilterCPU); - imshow(winName[2], abFilter); - imshow(winName[3], bFilter); - waitKey(); - return 0; + return EXIT_SUCCESS; } diff --git a/samples/ocl/bgfg_segm.cpp b/samples/ocl/bgfg_segm.cpp index a91c2f11a0..19d87ef03d 100644 --- a/samples/ocl/bgfg_segm.cpp +++ b/samples/ocl/bgfg_segm.cpp @@ -15,19 +15,18 @@ using namespace cv::ocl; int main(int argc, const char** argv) { - cv::CommandLineParser cmd(argc, argv, "{ c camera | false | use camera }" "{ f file | 768x576.avi | input video file }" "{ m method | mog | method (mog, mog2) }" "{ h help | false | print help message }"); - if (cmd.get("help")) + if (cmd.has("help")) { cout << "Usage : bgfg_segm [options]" << endl; cout << "Available options:" << endl; cmd.printMessage(); - return 0; + return EXIT_SUCCESS; } bool useCamera = cmd.get("camera"); @@ -37,13 +36,12 @@ int main(int argc, const char** argv) if (method != "mog" && method != "mog2") { cerr << "Incorrect method" << endl; - return -1; + return EXIT_FAILURE; } int m = method == "mog" ? M_MOG : M_MOG2; VideoCapture cap; - if (useCamera) cap.open(0); else @@ -51,8 +49,8 @@ int main(int argc, const char** argv) if (!cap.isOpened()) { - cerr << "can not open camera or video file" << endl; - return -1; + cout << "can not open camera or video file" << endl; + return EXIT_FAILURE; } Mat frame; @@ -63,15 +61,11 @@ int main(int argc, const char** argv) cv::ocl::MOG mog; cv::ocl::MOG2 mog2; - oclMat d_fgmask; - oclMat d_fgimg; - oclMat d_bgimg; + oclMat d_fgmask, d_fgimg, d_bgimg; d_fgimg.create(d_frame.size(), d_frame.type()); - Mat fgmask; - Mat fgimg; - Mat bgimg; + Mat fgmask, fgimg, bgimg; switch (m) { @@ -84,7 +78,7 @@ int main(int argc, const char** argv) break; } - for(;;) + for (;;) { cap >> frame; if (frame.empty()) @@ -124,10 +118,9 @@ int main(int argc, const char** argv) if (!bgimg.empty()) imshow("mean background image", bgimg); - int key = waitKey(30); - if (key == 27) + if (27 == waitKey(30)) break; } - return 0; + return EXIT_SUCCESS; } diff --git a/samples/ocl/clahe.cpp b/samples/ocl/clahe.cpp index 0b26437733..894a414831 100644 --- a/samples/ocl/clahe.cpp +++ b/samples/ocl/clahe.cpp @@ -10,15 +10,13 @@ using namespace std; Ptr pFilter; int tilesize; int cliplimit; -string outfile; static void TSize_Callback(int pos) { if(pos==0) - { pFilter->setTilesGridSize(Size(1,1)); - } - pFilter->setTilesGridSize(Size(tilesize,tilesize)); + else + pFilter->setTilesGridSize(Size(tilesize,tilesize)); } static void Clip_Callback(int) @@ -32,63 +30,64 @@ int main(int argc, char** argv) "{ i input | | specify input image }" "{ c camera | 0 | specify camera id }" "{ s use_cpu | false | use cpu algorithm }" - "{ o output | clahe_output.jpg | specify output save path}"; + "{ o output | clahe_output.jpg | specify output save path}" + "{ h help | false | print help message }"; + + cv::CommandLineParser cmd(argc, argv, keys); + if (cmd.has("help")) + { + cout << "Usage : clahe [options]" << endl; + cout << "Available options:" << endl; + cmd.printMessage(); + return EXIT_SUCCESS; + } - CommandLineParser cmd(argc, argv, keys); - string infile = cmd.get("i"); - outfile = cmd.get("o"); + string infile = cmd.get("i"), outfile = cmd.get("o"); int camid = cmd.get("c"); bool use_cpu = cmd.get("s"); VideoCapture capture; - bool running = true; namedWindow("CLAHE"); createTrackbar("Tile Size", "CLAHE", &tilesize, 32, (TrackbarCallback)TSize_Callback); createTrackbar("Clip Limit", "CLAHE", &cliplimit, 20, (TrackbarCallback)Clip_Callback); Mat frame, outframe; - ocl::oclMat d_outframe; + ocl::oclMat d_outframe, d_frame; int cur_clip; Size cur_tilesize; - if(use_cpu) - { - pFilter = createCLAHE(); - } - else - { - pFilter = ocl::createCLAHE(); - } + pFilter = use_cpu ? createCLAHE() : ocl::createCLAHE(); + cur_clip = (int)pFilter->getClipLimit(); cur_tilesize = pFilter->getTilesGridSize(); setTrackbarPos("Tile Size", "CLAHE", cur_tilesize.width); setTrackbarPos("Clip Limit", "CLAHE", cur_clip); + if(infile != "") { frame = imread(infile); if(frame.empty()) { cout << "error read image: " << infile << endl; - return -1; + return EXIT_FAILURE; } } else - { capture.open(camid); - } + cout << "\nControls:\n" << "\to - save output image\n" << "\tESC - exit\n"; - while(running) + + for (;;) { if(capture.isOpened()) capture.read(frame); else frame = imread(infile); if(frame.empty()) - { continue; - } + if(use_cpu) { cvtColor(frame, frame, COLOR_BGR2GRAY); @@ -96,15 +95,18 @@ int main(int argc, char** argv) } else { - ocl::oclMat d_frame(frame); - ocl::cvtColor(d_frame, d_outframe, COLOR_BGR2GRAY); + ocl::cvtColor(d_frame = frame, d_outframe, COLOR_BGR2GRAY); pFilter->apply(d_outframe, d_outframe); d_outframe.download(outframe); } + imshow("CLAHE", outframe); + char key = (char)waitKey(3); - if(key == 'o') imwrite(outfile, outframe); - else if(key == 27) running = false; + if(key == 'o') + imwrite(outfile, outframe); + else if(key == 27) + break; } - return 0; + return EXIT_SUCCESS; } diff --git a/samples/ocl/facedetect.cpp b/samples/ocl/facedetect.cpp index 66ef052304..8669719504 100644 --- a/samples/ocl/facedetect.cpp +++ b/samples/ocl/facedetect.cpp @@ -32,6 +32,7 @@ static void workBegin() { work_begin = getTickCount(); } + static void workEnd() { work_end += (getTickCount() - work_begin); @@ -42,16 +43,17 @@ static double getTime() return work_end /((double)cvGetTickFrequency() * 1000.); } -void detect( Mat& img, vector& faces, + +static void detect( Mat& img, vector& faces, ocl::OclCascadeClassifier& cascade, double scale, bool calTime); -void detectCPU( Mat& img, vector& faces, +static void detectCPU( Mat& img, vector& faces, CascadeClassifier& cascade, double scale, bool calTime); -void Draw(Mat& img, vector& faces, double scale); +static void Draw(Mat& img, vector& faces, double scale); // This function test if gpu_rst matches cpu_rst. @@ -59,7 +61,6 @@ void Draw(Mat& img, vector& faces, double scale); // Else if will return (total diff of each cpu and gpu rects covered pixels)/(total cpu rects covered pixels) double checkRectSimilarity(Size sz, vector& cpu_rst, vector& gpu_rst); - int main( int argc, const char** argv ) { const char* keys = @@ -75,10 +76,12 @@ int main( int argc, const char** argv ) CommandLineParser cmd(argc, argv, keys); if (cmd.get("help")) { + cout << "Usage : facedetect [options]" << endl; cout << "Available options:" << endl; cmd.printMessage(); - return 0; + return EXIT_SUCCESS; } + CvCapture* capture = 0; Mat frame, frameCopy, image; @@ -92,8 +95,8 @@ int main( int argc, const char** argv ) if( !cascade.load( cascadeName ) || !cpu_cascade.load(cascadeName) ) { - cerr << "ERROR: Could not load classifier cascade" << endl; - return -1; + cout << "ERROR: Could not load classifier cascade" << endl; + return EXIT_FAILURE; } if( inputName.empty() ) @@ -102,25 +105,17 @@ int main( int argc, const char** argv ) if(!capture) cout << "Capture from CAM 0 didn't work" << endl; } - else if( inputName.size() ) + else { - image = imread( inputName, 1 ); + image = imread( inputName, CV_LOAD_IMAGE_COLOR ); if( image.empty() ) { capture = cvCaptureFromAVI( inputName.c_str() ); if(!capture) cout << "Capture from AVI didn't work" << endl; - return -1; + return EXIT_FAILURE; } } - else - { - image = imread( "lena.jpg", 1 ); - if(image.empty()) - cout << "Couldn't read lena.jpg" << endl; - return -1; - } - cvNamedWindow( "result", 1 ); if( capture ) @@ -137,24 +132,16 @@ int main( int argc, const char** argv ) frame.copyTo( frameCopy ); else flip( frame, frameCopy, 0 ); + if(useCPU) - { detectCPU(frameCopy, faces, cpu_cascade, scale, false); - } else - { detect(frameCopy, faces, cascade, scale, false); - } + Draw(frameCopy, faces, scale); if( waitKey( 10 ) >= 0 ) - goto _cleanup_; + break; } - - - waitKey(0); - - -_cleanup_: cvReleaseCapture( &capture ); } else @@ -167,9 +154,7 @@ _cleanup_: { cout << "loop" << i << endl; if(useCPU) - { detectCPU(image, faces, cpu_cascade, scale, i==0?false:true); - } else { detect(image, faces, cascade, scale, i==0?false:true); diff --git a/samples/ocl/hog.cpp b/samples/ocl/hog.cpp index e29f78bb27..a3c5c99369 100644 --- a/samples/ocl/hog.cpp +++ b/samples/ocl/hog.cpp @@ -73,6 +73,14 @@ int main(int argc, char** argv) "{ l |larger_win| false | use 64x128 window}" "{ o | output | | specify output path when input is images}"; CommandLineParser cmd(argc, argv, keys); + if (cmd.has("help")) + { + cout << "Usage : hog [options]" << endl; + cout << "Available options:" << endl; + cmd.printMessage(); + return EXIT_SUCCESS; + } + App app(cmd); try { @@ -90,7 +98,7 @@ int main(int argc, char** argv) { return cout << "unknown exception" << endl, 1; } - return 0; + return EXIT_SUCCESS; } App::App(CommandLineParser& cmd) diff --git a/samples/ocl/pyrlk_optical_flow.cpp b/samples/ocl/pyrlk_optical_flow.cpp index 858a740af2..89137d96ee 100644 --- a/samples/ocl/pyrlk_optical_flow.cpp +++ b/samples/ocl/pyrlk_optical_flow.cpp @@ -45,7 +45,8 @@ static void download(const oclMat& d_mat, vector& vec) d_mat.download(mat); } -static void drawArrows(Mat& frame, const vector& prevPts, const vector& nextPts, const vector& status, Scalar line_color = Scalar(0, 0, 255)) +static void drawArrows(Mat& frame, const vector& prevPts, const vector& nextPts, const vector& status, + Scalar line_color = Scalar(0, 0, 255)) { for (size_t i = 0; i < prevPts.size(); ++i) { @@ -105,7 +106,7 @@ int main(int argc, const char* argv[]) cout << "Usage: pyrlk_optical_flow [options]" << endl; cout << "Available options:" << endl; cmd.printMessage(); - return 0; + return EXIT_SUCCESS; } bool defaultPicturesFail = false; @@ -137,7 +138,7 @@ int main(int argc, const char* argv[]) Mat frame0Gray, frame1Gray; Mat ptr0, ptr1; - if(vdofile == "") + if(vdofile.empty()) capture.open( inputName ); else capture.open(vdofile.c_str()); @@ -145,14 +146,12 @@ int main(int argc, const char* argv[]) int c = inputName ; if(!capture.isOpened()) { - if(vdofile == "") + if(vdofile.empty()) cout << "Capture from CAM " << c << " didn't work" << endl; else cout << "Capture from file " << vdofile << " failed" <= 0 ) - goto _cleanup_; + break; } - waitKey(0); - -_cleanup_: capture.release(); } else @@ -264,5 +260,5 @@ nocamera: waitKey(); - return 0; + return EXIT_SUCCESS; } diff --git a/samples/ocl/squares.cpp b/samples/ocl/squares.cpp index 797775f6ba..b53648f3f7 100644 --- a/samples/ocl/squares.cpp +++ b/samples/ocl/squares.cpp @@ -14,9 +14,9 @@ using namespace cv; using namespace std; -#define ACCURACY_CHECK 1 +#define ACCURACY_CHECK -#if ACCURACY_CHECK +#ifdef ACCURACY_CHECK // check if two vectors of vector of points are near or not // prior assumption is that they are in correct order static bool checkPoints( @@ -279,27 +279,31 @@ int main(int argc, char** argv) { const char* keys = "{ i | input | | specify input image }" - "{ o | output | squares_output.jpg | specify output save path}"; + "{ o | output | squares_output.jpg | specify output save path}" + "{ h | help | false | print help message }"; CommandLineParser cmd(argc, argv, keys); string inputName = cmd.get("i"); string outfile = cmd.get("o"); - if(inputName.empty()) + + if(cmd.get("help")) { + cout << "Usage : squares [options]" << endl; cout << "Available options:" << endl; cmd.printMessage(); - return 0; + return EXIT_SUCCESS; } int iterations = 10; - namedWindow( wndname, 1 ); + namedWindow( wndname, WINDOW_AUTOSIZE ); vector > squares_cpu, squares_ocl; Mat image = imread(inputName, 1); if( image.empty() ) { cout << "Couldn't load " << inputName << endl; - return -1; + return EXIT_FAILURE; } + int j = iterations; int64 t_ocl = 0, t_cpp = 0; //warm-ups @@ -308,7 +312,7 @@ int main(int argc, char** argv) findSquares_ocl(image, squares_ocl); -#if ACCURACY_CHECK +#ifdef ACCURACY_CHECK cout << "Checking ocl accuracy ... " << endl; cout << (checkPoints(squares_cpu, squares_ocl) ? "Pass" : "Failed") << endl; #endif @@ -333,5 +337,5 @@ int main(int argc, char** argv) imwrite(outfile, result); waitKey(0); - return 0; + return EXIT_SUCCESS; } diff --git a/samples/ocl/stereo_match.cpp b/samples/ocl/stereo_match.cpp index 932f996990..880ad51c0b 100644 --- a/samples/ocl/stereo_match.cpp +++ b/samples/ocl/stereo_match.cpp @@ -78,8 +78,9 @@ int main(int argc, char** argv) "{ l | left | | specify left image }" "{ r | right | | specify right image }" "{ m | method | BM | specify match method(BM/BP/CSBP) }" - "{ n | ndisp | 64 | specify number of disparity levels }" + "{ n | ndisp | 64 | specify number of disparity levels }" "{ o | output | stereo_match_output.jpg | specify output path when input is images}"; + CommandLineParser cmd(argc, argv, keys); if (cmd.get("help")) { @@ -87,6 +88,7 @@ int main(int argc, char** argv) cmd.printMessage(); return 0; } + try { App app(cmd); @@ -98,7 +100,8 @@ int main(int argc, char** argv) { cout << "error: " << e.what() << endl; } - return 0; + + return EXIT_SUCCESS; } App::App(CommandLineParser& cmd) @@ -116,6 +119,7 @@ App::App(CommandLineParser& cmd) << "\t2/w - increase/decrease window size (for BM only)\n" << "\t3/e - increase/decrease iteration count (for BP and CSBP only)\n" << "\t4/r - increase/decrease level count (for BP and CSBP only)\n"; + l_img = cmd.get("l"); r_img = cmd.get("r"); string mstr = cmd.get("m"); diff --git a/samples/ocl/surf_matcher.cpp b/samples/ocl/surf_matcher.cpp index 49321269ee..f88678b7b4 100644 --- a/samples/ocl/surf_matcher.cpp +++ b/samples/ocl/surf_matcher.cpp @@ -15,21 +15,20 @@ const int LOOP_NUM = 10; const int GOOD_PTS_MAX = 50; const float GOOD_PORTION = 0.15f; -namespace -{ - int64 work_begin = 0; int64 work_end = 0; -void workBegin() +static void workBegin() { work_begin = getTickCount(); } -void workEnd() + +static void workEnd() { work_end = getTickCount() - work_begin; } -double getTime() + +static double getTime() { return work_end /((double)getTickFrequency() * 1000.); } @@ -60,7 +59,7 @@ struct SURFMatcher } }; -Mat drawGoodMatches( +static Mat drawGoodMatches( const Mat& cpu_img1, const Mat& cpu_img2, const std::vector& keypoints1, @@ -130,7 +129,6 @@ Mat drawGoodMatches( return img_matches; } -} //////////////////////////////////////////////////// // This program demonstrates the usage of SURF_OCL. // use cpu findHomography interface to calculate the transformation matrix @@ -143,12 +141,14 @@ int main(int argc, char* argv[]) "{ output o | SURF_output.jpg | specify output save path (only works in CPU or GPU only mode) }" "{ use_cpu c | false | use CPU algorithms }" "{ use_all a | false | use both CPU and GPU algorithms}"; + CommandLineParser cmd(argc, argv, keys); if (cmd.get("help")) { + std::cout << "Usage: surf_matcher [options]" << std::endl; std::cout << "Available options:" << std::endl; cmd.printMessage(); - return 0; + return EXIT_SUCCESS; } Mat cpu_img1, cpu_img2, cpu_img1_grey, cpu_img2_grey; @@ -169,23 +169,17 @@ int main(int argc, char* argv[]) cvtColor(cpu_img2, cpu_img2_grey, COLOR_BGR2GRAY); img2 = cpu_img2_grey; - if(useALL) - { - useCPU = false; - useGPU = false; - } - else if(useCPU==false && useALL==false) - { + if (useALL) + useCPU = useGPU = false; + else if(!useCPU && !useALL) useGPU = true; - } if(!useCPU) - { std::cout << "Device name:" << cv::ocl::Context::getContext()->getDeviceInfo().deviceName << std::endl; - } + double surf_time = 0.; //declare input/output @@ -331,5 +325,5 @@ int main(int argc, char* argv[]) imshow("ocl surf matches", ocl_img_matches); } waitKey(0); - return 0; + return EXIT_SUCCESS; } diff --git a/samples/ocl/tvl1_optical_flow.cpp b/samples/ocl/tvl1_optical_flow.cpp index b067476c4b..046e0cba70 100644 --- a/samples/ocl/tvl1_optical_flow.cpp +++ b/samples/ocl/tvl1_optical_flow.cpp @@ -97,10 +97,9 @@ int main(int argc, const char* argv[]) cout << "Usage: pyrlk_optical_flow [options]" << endl; cout << "Available options:" << endl; cmd.printMessage(); - return 0; + return EXIT_SUCCESS; } - bool defaultPicturesFail = false; string fname0 = cmd.get("l"); string fname1 = cmd.get("r"); string vdofile = cmd.get("v"); @@ -114,21 +113,10 @@ int main(int argc, const char* argv[]) cv::Ptr alg = cv::createOptFlow_DualTVL1(); cv::ocl::OpticalFlowDual_TVL1_OCL d_alg; - Mat flow, show_flow; Mat flow_vec[2]; if (frame0.empty() || frame1.empty()) - { useCamera = true; - defaultPicturesFail = true; - VideoCapture capture( inputName ); - if (!capture.isOpened()) - { - cout << "Can't load input images" << endl; - return -1; - } - } - if (useCamera) { @@ -137,22 +125,17 @@ int main(int argc, const char* argv[]) Mat frame0Gray, frame1Gray; Mat ptr0, ptr1; - if(vdofile == "") + if(vdofile.empty()) capture.open( inputName ); else capture.open(vdofile.c_str()); - int c = inputName ; if(!capture.isOpened()) { - if(vdofile == "") - cout << "Capture from CAM " << c << " didn't work" << endl; + if(vdofile.empty()) + cout << "Capture from CAM " << inputName << " didn't work" << endl; else cout << "Capture from file " << vdofile << " failed" <= 0 ) - goto _cleanup_; + break; } - waitKey(0); - -_cleanup_: capture.release(); } else @@ -253,5 +233,5 @@ nocamera: waitKey(); - return 0; + return EXIT_SUCCESS; }