diff --git a/CMakeLists.txt b/CMakeLists.txt index 8cec98fc9..73def9547 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 b8945c257..4b80b1f78 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 881cd371a..362841b21 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 7bfc7bc5a..a36b70e94 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 5dbe957dd..ed5acc76b 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 818ec9b4b..13461e82c 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") +# 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) + if(fname MATCHES "\\.lib$") + string(REGEX REPLACE "\\.lib$" ".pdb" fname "${fname}") + install(FILES ${fname} DESTINATION ${__dst} CONFIGURATIONS Release) + endif() endif() - - get_target_property(fname ${__target} LOCATION_DEBUG) - string(REPLACE ".lib" ".pdb" fname "${fname}") - install(FILES ${fname} DESTINATION ${DST} CONFIGURATIONS Debug) - - get_target_property(fname ${__target} LOCATION_RELEASE) - string(REPLACE ".lib" ".pdb" fname "${fname}") - install(FILES ${fname} DESTINATION ${DST} CONFIGURATIONS Release) endif() endif() endfunction() diff --git a/cmake/templates/OpenCVConfig.cmake.in b/cmake/templates/OpenCVConfig.cmake.in index c1d802172..2dfc78169 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 c4cde1990..3eb8913be 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 893c0e9a9..b93b3f7eb 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 a15ef8947..5e6252a61 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 fdd1931bd..cc48aa28c 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 aace143d5..02f77c224 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 6aa71d8ea..20367ab98 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 d6a877bc8..c52b6e366 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 1361367fc..33748c0a5 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 78ebc1a26..836f8ee9b 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 880bdff5e..025221b4e 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 b3d02f27a..95099640f 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 018ec6315..a5e057ffc 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 33c42c72d..86c0a3c70 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 997e84856..12fee549b 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 259684092..33723daa3 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 b66fc2b0a..e56db2495 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 4cba47e96..49da65936 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 7f2758877..b6dcd2a08 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 803e1f91b..4dcd5d4d6 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 8a29adc0c..af24c3489 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 9c258fe25..5332a2d93 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 fe5d9d190..2a6731117 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 48dbf4dde..7a9eab5d3 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 017a8a70d..946444ad9 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 869e01e60..68192cf80 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 8266f0930..3035c97f0 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 fac471ed4..db45eceb8 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 6ecc76651..a36e1a13e 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 35ac00639..ff49eb4ed 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 861307526..bc1761b49 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 0235b3d04..f4b59a5fd 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 19c728bb7..c799853db 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 3821a8e16..f2f7c4115 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 cff087574..6bfa7333a 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 6ea520e37..c6883661b 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 5372702a5..c9bba13c9 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 59a274fae..3c2988dc8 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 fda3b1c5f..bc24f5e38 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 d321d91c5..3f5de5274 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 72ffce485..2b3129d05 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 5f04561e9..f83220dae 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 ba279b794..1254e3063 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 e96bd57b6..ebf3e7676 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 9b91802d8..eec103a6d 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 33f4b0047..ccbd960bc 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 9689e0a30..a1e2d807d 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 0fbcecc5c..e10d67152 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 984f3a9f9..816988db7 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"; + size_t localThreads[3] = { 16, 16, 1 }; + size_t globalThreads[3] = { dst.cols, dst.rows, 1 }; - char btype[30]; - - 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 504b697a7..92119df5e 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 28af01dc2..b07286553 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 15dcbf9b6..811649636 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 7427620f6..70fe99187 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 26c5052ea..ed39868c7 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 78669bae1..47d6c837a 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 6f8243457..5a133a7b1 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 c5a03bacd..5486aa495 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 9b77f7a58..143e7aa7a 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 9720c7421..b822aaafc 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 f3dc7b56f..287cfa686 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 866dbbef6..f6efbf7ad 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 a5f684415..6372364dd 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 0b31f16ff..163e017c4 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 ff21e9a31..658e1f4bc 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 020880606..fcf38749d 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 2f34bbbff..a73b65da6 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 159a970db..8272806e2 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 c3958bf1f..ea96d8a8a 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 7f4e41327..b82eff595 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 b93de0c6b..0762b19b1 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 a4fa205c1..56cd745d2 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 d244e572d..328ccd91a 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 5a7d5938c..434bd5eca 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 714220ddf..e5b46c936 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 a2f65e0b7..6c779ead9 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 d0842db18..005d3c73f 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 b2143ba14..835bc95c3 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 49242d07c..7c2a04d74 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 68e26a439..4a2038275 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 ef8c4dd04..fe1b3046a 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 921367b3d..085386f5c 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 a30eba431..b6bc7b42b 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 dd9250c7f..1704f6b42 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 921026b40..fb684c367 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 39bcf949a..6eb6e4832 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 5328d1f1b..bd06a5208 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 2e269999a..8621ff31b 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 a05c98ee0..d6a89f205 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 07cccee1a..4402e8c38 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 1908f92a2..b3e699dc4 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 2b1cfccd0..fcbf67ca7 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 60ce51360..8dd77d5a9 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) +#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 0) +#else +#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 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) -#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)) +#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 0) +#else +#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 1) #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)) -#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 d163ebe76..030c13cc5 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 7c5b0c321..ea22967df 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 db88aca65..c402ff721 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 22a7fe7cb..5fa353305 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 c12ab5950..17e95b4e4 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 5402759e3..ca670b6db 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 55692ae3b..16c68fd47 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 1609d7c55..6b596a322 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 db7a7dfc3..fb9596e5d 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 ff7509ffd..d97f66068 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 +#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_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) -#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)) + int src_x = x - left; + int src_y = y - top; + int dst_index = mad24(y, dst_step, x + dst_offset); + + if (NEED_EXTRAPOLATION(src_x, src_y)) + EXTRAPOLATE(src_x, src_y, dst[dst_index]) + else { - 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; + int src_index = mad24(src_y, src_step, src_x + src_offset); + dst[dst_index] = src[src_index]; } - 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 d2e136301..8b4c865c3 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 8e6e26939..63693749d 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 5b80449e2..d31b3715b 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 532e61d13..9fd0b2329 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 b7d2d6d44..82286031f 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 886f3c1e5..cc7843db6 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 d9cc7b159..20d88c21f 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 2e54570e7..3caea04fb 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 68dab0ac0..c2a44842c 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 fa7a70f4d..eb983fb17 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 f02df6af7..045cd9815 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 5e4026694..94263d8f7 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 0c2e9bd4e..edbc36a3f 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 27a587260..c7ceef453 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 684a2a937..6ee3e35a7 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 76940fff0..a06407089 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 da38fe471..7296a6b7e 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 3b91f12ba..2d861b627 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 5a7d4a300..b25914968 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 b090045c7..6148e95cb 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 717bbc7a2..bfe5b638f 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 6e519991d..8521dff82 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 e578effca..0f0fac313 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 b5df8b5fb..44edc815e 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 472062323..3c0cff85b 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) + } + else { - 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 - { - 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 bee73dbea..a9d4d7c42 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 12a061c68..67d0532a6 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 f67c54109..c274d903a 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 a91c2f11a..19d87ef03 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 0b2643773..894a41483 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 }"; - CommandLineParser cmd(argc, argv, keys); - string infile = cmd.get("i"); - outfile = cmd.get("o"); + cv::CommandLineParser cmd(argc, argv, keys); + if (cmd.has("help")) + { + cout << "Usage : clahe [options]" << endl; + cout << "Available options:" << endl; + cmd.printMessage(); + return EXIT_SUCCESS; + } + + 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 66ef05230..866971950 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 e29f78bb2..a3c5c9936 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 858a740af..89137d96e 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 797775f6b..b53648f3f 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 932f99699..880ad51c0 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 49321269e..f88678b7b 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 b067476c4..046e0cba7 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; }