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)
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) )
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 ---
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.
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(<name> [major[.minor]] [EXACT] [REQUIRED|QUIET])"
# cmake will look in the following dir on 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()
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}/")
# 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
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()
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}})
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")
endif()
endforeach()
-# message(STATUS "Process ${__target} dst=${DST}...")
- if(NOT DEFINED DST)
- set(DST "OPENCV_LIB_INSTALL_PATH")
- endif()
-
- get_target_property(fname ${__target} LOCATION_DEBUG)
- string(REPLACE ".lib" ".pdb" fname "${fname}")
- install(FILES ${fname} DESTINATION ${DST} CONFIGURATIONS Debug)
+# message(STATUS "Process ${__target} dst=${__dst}...")
+ if(DEFINED __dst)
+ get_target_property(fname ${__target} LOCATION_DEBUG)
+ if(fname MATCHES "\\.lib$")
+ string(REGEX REPLACE "\\.lib$" ".pdb" fname "${fname}")
+ install(FILES ${fname} DESTINATION ${__dst} CONFIGURATIONS Debug)
+ endif()
- get_target_property(fname ${__target} LOCATION_RELEASE)
- string(REPLACE ".lib" ".pdb" fname "${fname}")
- install(FILES ${fname} DESTINATION ${DST} CONFIGURATIONS Release)
+ get_target_property(fname ${__target} LOCATION_RELEASE)
+ if(fname MATCHES "\\.lib$")
+ string(REGEX REPLACE "\\.lib$" ".pdb" fname "${fname}")
+ install(FILES ${fname} DESTINATION ${__dst} CONFIGURATIONS Release)
+ endif()
+ endif()
endif()
endif()
endfunction()
#
# ===================================================================================
-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)
*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
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
CGContextDrawImage(contextRef, CGRectMake(0, 0, cols, rows), image.CGImage);
CGContextRelease(contextRef);
- CGColorSpaceRelease(colorSpace);
return cvMat;
}
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
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
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_")
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
#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)
{
if (fdwReason == DLL_PROCESS_DETACH)
//
// * 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.
//
// * 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.
cv::Mutex ProgramCache::mutexFiles;
cv::Mutex ProgramCache::mutexCache;
-std::auto_ptr<ProgramCache> _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()
ProgramCache::~ProgramCache()
{
releaseProgram();
+ if (this == _programCache)
+ {
+ cv::AutoLock lock(getInitializationMutex());
+ if (this == _programCache)
+ _programCache = NULL;
+ }
}
cl_program ProgramCache::progLookup(const String& srcsign)
{
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<char> 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);
}
//
// * 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.
protected:
ProgramCache();
~ProgramCache();
- friend class std::auto_ptr<ProgramCache>;
public:
static ProgramCache *getProgramCache();
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
{
return;
}
- pCache.setupData = new clAmdFftSetupData;
+ if (pCache.setupData == NULL)
+ pCache.setupData = new clAmdFftSetupData;
openCLSafeCall(clAmdFftInitSetupData( pCache.setupData ));
pCache.started = true;
}
//
// * 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.
template <typename T>
void linearRowFilter_gpu(const oclMat &src, const oclMat &dst, oclMat mat_kernel, int ksize, int anchor, int bordertype)
{
- Context *clCxt = src.clCxt;
+ CV_Assert(bordertype <= BORDER_REFLECT_101);
+ CV_Assert(ksize == (anchor << 1) + 1);
int channels = src.oclchannels();
- size_t localThreads[3] = {16, 16, 1};
- String kernelName = "row_filter";
-
- char btype[30];
+ size_t localThreads[3] = { 16, 16, 1 };
+ size_t globalThreads[3] = { dst.cols, dst.rows, 1 };
- switch (bordertype)
- {
- case 0:
- sprintf(btype, "BORDER_CONSTANT");
- break;
- case 1:
- sprintf(btype, "BORDER_REPLICATE");
- break;
- case 2:
- sprintf(btype, "BORDER_REFLECT");
- break;
- case 3:
- sprintf(btype, "BORDER_WRAP");
- break;
- case 4:
- sprintf(btype, "BORDER_REFLECT_101");
- break;
- }
-
- char compile_option[128];
- sprintf(compile_option, "-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s", anchor, (int)localThreads[0], (int)localThreads[1], channels, btype);
-
- size_t globalThreads[3];
- globalThreads[1] = (dst.rows + localThreads[1] - 1) / localThreads[1] * localThreads[1];
- globalThreads[2] = (1 + localThreads[2] - 1) / localThreads[2] * localThreads[2];
+ const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" };
+ std::string buildOptions = format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s",
+ anchor, (int)localThreads[0], (int)localThreads[1], channels, borderMap[bordertype]);
if (src.depth() == CV_8U)
{
switch (channels)
{
case 1:
- case 3:
- globalThreads[0] = ((dst.cols + 4) / 4 + localThreads[0] - 1) / localThreads[0] * localThreads[0];
+ globalThreads[0] = (dst.cols + 3) >> 2;
break;
case 2:
- globalThreads[0] = ((dst.cols + 1) / 2 + localThreads[0] - 1) / localThreads[0] * localThreads[0];
+ globalThreads[0] = (dst.cols + 1) >> 1;
break;
case 4:
- globalThreads[0] = (dst.cols + localThreads[0] - 1) / localThreads[0] * localThreads[0];
+ globalThreads[0] = dst.cols;
break;
}
}
- else
- {
- globalThreads[0] = (dst.cols + localThreads[0] - 1) / localThreads[0] * localThreads[0];
- }
- //sanity checks
- CV_Assert(clCxt == dst.clCxt);
- CV_Assert(src.cols == dst.cols);
- CV_Assert(src.oclchannels() == dst.oclchannels());
- CV_Assert(ksize == (anchor << 1) + 1);
- int src_pix_per_row, dst_pix_per_row;
- int src_offset_x, src_offset_y;//, dst_offset_in_pixel;
- src_pix_per_row = src.step / src.elemSize();
- src_offset_x = (src.offset % src.step) / src.elemSize();
- src_offset_y = src.offset / src.step;
- dst_pix_per_row = dst.step / dst.elemSize();
- //dst_offset_in_pixel = dst.offset / dst.elemSize();
+ int src_pix_per_row = src.step / src.elemSize();
+ int src_offset_x = (src.offset % src.step) / src.elemSize();
+ int src_offset_y = src.offset / src.step;
+ int dst_pix_per_row = dst.step / dst.elemSize();
int ridusy = (dst.rows - src.rows) >> 1;
+
std::vector<std::pair<size_t , const void *> > args;
args.push_back(std::make_pair(sizeof(cl_mem), &src.data));
args.push_back(std::make_pair(sizeof(cl_mem), &dst.data));
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<BaseRowFilter_GPU> cv::ocl::getLinearRowFilter_GPU(int srcType, int /*bufType*/, const Mat &rowKernel, int anchor, int bordertype)
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
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;
}
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<size_t, const void *> > args;
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<uchar> _buf(bufSize);
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());
}
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
#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
{
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<READ_TIMES_COL;i++)
//
// * 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.
#define ALIGN (RADIUS)
#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
-
+#elif defined BORDER_REPLICATE
+#define EXTRAPOLATE(x, maxV) \
+ { \
+ x = max(min(x, maxV - 1), 0); \
+ }
+#elif defined BORDER_WRAP
+#define EXTRAPOLATE(x, maxV) \
+ { \
+ if (x < 0) \
+ x -= ((x - maxV + 1) / maxV) * maxV; \
+ if (x >= maxV) \
+ x %= maxV; \
+ }
+#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
+#define EXTRAPOLATE_(x, maxV, delta) \
+ { \
+ if (maxV == 1) \
+ x = 0; \
+ else \
+ do \
+ { \
+ if ( x < 0 ) \
+ x = -x - 1 + delta; \
+ else \
+ x = maxV - 1 - (x - maxV) - delta; \
+ } \
+ while (x >= maxV || x < 0); \
+ }
#ifdef BORDER_REFLECT
-//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
-#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i)-1 : (addr)
-#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)
-#endif
-
-#ifdef BORDER_REFLECT_101
-//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
-#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i) : (addr)
-#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)
+#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 0)
+#else
+#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 1)
#endif
-
-#ifdef BORDER_WRAP
-//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
-#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (i)+(r_edge) : (addr)
-#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (i)-(r_edge) : (addr)
+#else
+#error No extrapolation method
#endif
/**********************************************************************************
***********************************************************************************/
__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<READ_TIMES_ROW; i++)
+ int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols);
+
+ // read pixels from src
+ for (i = 0; i < READ_TIMES_ROW; i++)
{
int current_addr = start_addr+i*LSIZE0*4;
current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
temp[i] = *(__global uchar4*)&src[current_addr];
}
- //judge if read out of boundary
- for(i = 0; i<READ_TIMES_ROW; i++)
+
+ // judge if read out of boundary
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
- temp[i].x= ELEM(start_x+i*LSIZE0*4,0,src_whole_cols,0,temp[i].x);
- temp[i].y= ELEM(start_x+i*LSIZE0*4+1,0,src_whole_cols,0,temp[i].y);
- temp[i].z= ELEM(start_x+i*LSIZE0*4+2,0,src_whole_cols,0,temp[i].z);
- temp[i].w= ELEM(start_x+i*LSIZE0*4+3,0,src_whole_cols,0,temp[i].w);
- temp[i]= ELEM(start_y,0,src_whole_rows,(uchar4)0,temp[i]);
+ temp[i].x = ELEM(start_x+i*LSIZE0*4,0,src_whole_cols,0,temp[i].x);
+ temp[i].y = ELEM(start_x+i*LSIZE0*4+1,0,src_whole_cols,0,temp[i].y);
+ temp[i].z = ELEM(start_x+i*LSIZE0*4+2,0,src_whole_cols,0,temp[i].z);
+ temp[i].w = ELEM(start_x+i*LSIZE0*4+3,0,src_whole_cols,0,temp[i].w);
+ temp[i] = ELEM(start_y,0,src_whole_rows,(uchar4)0,temp[i]);
}
#else
int not_all_in_range = (start_x<0) | (start_x + READ_TIMES_ROW*LSIZE0*4+4>src_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<READ_TIMES_ROW; i++)
+ // judge if read out of boundary
+ for (i = 0; i < READ_TIMES_ROW; i++)
{
- index[i].x= ADDR_L(start_x+i*LSIZE0*4,0,src_whole_cols,start_x+i*LSIZE0*4);
- index[i].x= ADDR_R(start_x+i*LSIZE0*4,src_whole_cols,index[i].x);
- index[i].y= ADDR_L(start_x+i*LSIZE0*4+1,0,src_whole_cols,start_x+i*LSIZE0*4+1);
- index[i].y= ADDR_R(start_x+i*LSIZE0*4+1,src_whole_cols,index[i].y);
- index[i].z= ADDR_L(start_x+i*LSIZE0*4+2,0,src_whole_cols,start_x+i*LSIZE0*4+2);
- index[i].z= ADDR_R(start_x+i*LSIZE0*4+2,src_whole_cols,index[i].z);
- index[i].w= ADDR_L(start_x+i*LSIZE0*4+3,0,src_whole_cols,start_x+i*LSIZE0*4+3);
- index[i].w= ADDR_R(start_x+i*LSIZE0*4+3,src_whole_cols,index[i].w);
+ index[i] = (int4)(start_x+i*LSIZE0*4) + (int4)(0, 1, 2, 3);
+ EXTRAPOLATE(index[i].x, src_whole_cols);
+ EXTRAPOLATE(index[i].y, src_whole_cols);
+ EXTRAPOLATE(index[i].z, src_whole_cols);
+ EXTRAPOLATE(index[i].w, src_whole_cols);
}
- s_y= ADDR_L(start_y,0,src_whole_rows,start_y);
- s_y= ADDR_R(start_y,src_whole_rows,s_y);
- //read pixels from src
- for(i = 0; i<READ_TIMES_ROW; i++)
+ s_y = start_y;
+ EXTRAPOLATE(s_y, src_whole_rows);
+
+ // read pixels from src
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
addr = mad24((int4)s_y,(int4)src_step_in_pixel,index[i]);
temp[i].x = src[addr.x];
}
else
{
- //read pixels from src
- for(i = 0; i<READ_TIMES_ROW; i++)
- {
+ // read pixels from src
+ for (i = 0; i<READ_TIMES_ROW; i++)
temp[i] = *(__global uchar4*)&src[start_addr+i*LSIZE0*4];
- }
}
#endif
- //save pixels to lds
- for(i = 0; i<READ_TIMES_ROW; i++)
- {
+ // save pixels to lds
+ for (i = 0; i<READ_TIMES_ROW; i++)
LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
- }
barrier(CLK_LOCAL_MEM_FENCE);
- //read pixels from lds and calculate the result
+ // read pixels from lds and calculate the result
sum =convert_float4(vload4(0,(__local uchar*)&LDS_DAT[l_y][l_x]+RADIUSX+offset))*mat_kernel[RADIUSX];
- for(i=1; i<=RADIUSX; i++)
+ for (i=1; i<=RADIUSX; i++)
{
- temp[0]=vload4(0,(__local uchar*)&LDS_DAT[l_y][l_x]+RADIUSX+offset-i);
- temp[1]=vload4(0,(__local uchar*)&LDS_DAT[l_y][l_x]+RADIUSX+offset+i);
- sum += convert_float4(temp[0])*mat_kernel[RADIUSX-i]+convert_float4(temp[1])*mat_kernel[RADIUSX+i];
+ temp[0] = vload4(0, (__local uchar*)&LDS_DAT[l_y][l_x] + RADIUSX + offset - i);
+ temp[1] = vload4(0, (__local uchar*)&LDS_DAT[l_y][l_x] + RADIUSX + offset + i);
+ sum += convert_float4(temp[0]) * mat_kernel[RADIUSX-i] + convert_float4(temp[1]) * mat_kernel[RADIUSX+i];
}
+
start_addr = mad24(y,dst_step_in_pixel,x);
- //write the result to dst
- if((x+3<dst_cols) & (y<dst_rows))
- {
+
+ // write the result to dst
+ if ((x+3<dst_cols) & (y<dst_rows))
*(__global float4*)&dst[start_addr] = sum;
- }
- else if((x+2<dst_cols) & (y<dst_rows))
+ else if ((x+2<dst_cols) && (y<dst_rows))
{
dst[start_addr] = sum.x;
dst[start_addr+1] = sum.y;
dst[start_addr+2] = sum.z;
}
- else if((x+1<dst_cols) & (y<dst_rows))
+ else if ((x+1<dst_cols) && (y<dst_rows))
{
dst[start_addr] = sum.x;
dst[start_addr+1] = sum.y;
}
- else if((x<dst_cols) & (y<dst_rows))
- {
+ else if (x<dst_cols && y<dst_rows)
dst[start_addr] = sum.x;
- }
}
+
__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C4_D0
-(__global const uchar4 * restrict src,
- __global float4 * 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 uchar4 * restrict src,
+ __global float4 * 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);
int y = get_global_id(1);
__local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1];
#ifdef BORDER_CONSTANT
int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
- //read pixels from src
- for(i = 0; i<READ_TIMES_ROW; i++)
+
+ // read pixels from src
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
int current_addr = start_addr+i*LSIZE0;
current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
temp[i] = src[current_addr];
}
+
//judge if read out of boundary
- for(i = 0; i<READ_TIMES_ROW; i++)
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
temp[i]= ELEM(start_x+i*LSIZE0,0,src_whole_cols,(uchar4)0,temp[i]);
temp[i]= ELEM(start_y,0,src_whole_rows,(uchar4)0,temp[i]);
#else
int index[READ_TIMES_ROW];
int s_x,s_y;
- //judge if read out of boundary
- for(i = 0; i<READ_TIMES_ROW; i++)
+
+ // judge if read out of boundary
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
- s_x= ADDR_L(start_x+i*LSIZE0,0,src_whole_cols,start_x+i*LSIZE0);
- s_x= ADDR_R(start_x+i*LSIZE0,src_whole_cols,s_x);
- s_y= ADDR_L(start_y,0,src_whole_rows,start_y);
- s_y= ADDR_R(start_y,src_whole_rows,s_y);
+ s_x = start_x+i*LSIZE0;
+ EXTRAPOLATE(s_x, src_whole_cols);
+ s_y = start_y;
+ EXTRAPOLATE(s_y, src_whole_rows);
index[i]=mad24(s_y,src_step_in_pixel,s_x);
}
+
//read pixels from src
- for(i = 0; i<READ_TIMES_ROW; i++)
- {
+ for (i = 0; i<READ_TIMES_ROW; i++)
temp[i] = src[index[i]];
- }
#endif
//save pixels to lds
- for(i = 0; i<READ_TIMES_ROW; i++)
- {
+ for (i = 0; i<READ_TIMES_ROW; i++)
LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
- }
barrier(CLK_LOCAL_MEM_FENCE);
//read pixels from lds and calculate the result
sum =convert_float4(LDS_DAT[l_y][l_x+RADIUSX])*mat_kernel[RADIUSX];
- for(i=1; i<=RADIUSX; i++)
+ for (i=1; i<=RADIUSX; i++)
{
temp[0]=LDS_DAT[l_y][l_x+RADIUSX-i];
temp[1]=LDS_DAT[l_y][l_x+RADIUSX+i];
sum += convert_float4(temp[0])*mat_kernel[RADIUSX-i]+convert_float4(temp[1])*mat_kernel[RADIUSX+i];
}
//write the result to dst
- if((x<dst_cols) & (y<dst_rows))
+ if (x<dst_cols && y<dst_rows)
{
start_addr = mad24(y,dst_step_in_pixel,x);
dst[start_addr] = sum;
}
__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C1_D5
-(__global const float * 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 float * 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);
int y = get_global_id(1);
__local float LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1];
#ifdef BORDER_CONSTANT
int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
- //read pixels from src
- for(i = 0; i<READ_TIMES_ROW; i++)
+
+ // read pixels from src
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
int current_addr = start_addr+i*LSIZE0;
current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
temp[i] = src[current_addr];
}
- //judge if read out of boundary
- for(i = 0; i<READ_TIMES_ROW; i++)
+
+ // judge if read out of boundary
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
temp[i]= ELEM(start_x+i*LSIZE0,0,src_whole_cols,(float)0,temp[i]);
temp[i]= ELEM(start_y,0,src_whole_rows,(float)0,temp[i]);
#else
int index[READ_TIMES_ROW];
int s_x,s_y;
- //judge if read out of boundary
- for(i = 0; i<READ_TIMES_ROW; i++)
+ // judge if read out of boundary
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
- s_x= ADDR_L(start_x+i*LSIZE0,0,src_whole_cols,start_x+i*LSIZE0);
- s_x= ADDR_R(start_x+i*LSIZE0,src_whole_cols,s_x);
- s_y= ADDR_L(start_y,0,src_whole_rows,start_y);
- s_y= ADDR_R(start_y,src_whole_rows,s_y);
- index[i]=mad24(s_y,src_step_in_pixel,s_x);
+ s_x = start_x + i*LSIZE0, s_y = start_y;
+ EXTRAPOLATE(s_x, src_whole_cols);
+ EXTRAPOLATE(s_y, src_whole_rows);
+
+ index[i]=mad24(s_y, src_step_in_pixel, s_x);
}
- //read pixels from src
- for(i = 0; i<READ_TIMES_ROW; i++)
- {
+ // read pixels from src
+ for (i = 0; i<READ_TIMES_ROW; i++)
temp[i] = src[index[i]];
- }
#endif
//save pixels to lds
- for(i = 0; i<READ_TIMES_ROW; i++)
- {
+ for (i = 0; i<READ_TIMES_ROW; i++)
LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
- }
barrier(CLK_LOCAL_MEM_FENCE);
- //read pixels from lds and calculate the result
+ // read pixels from lds and calculate the result
sum =LDS_DAT[l_y][l_x+RADIUSX]*mat_kernel[RADIUSX];
- for(i=1; i<=RADIUSX; i++)
+ for (i=1; i<=RADIUSX; i++)
{
temp[0]=LDS_DAT[l_y][l_x+RADIUSX-i];
temp[1]=LDS_DAT[l_y][l_x+RADIUSX+i];
sum += temp[0]*mat_kernel[RADIUSX-i]+temp[1]*mat_kernel[RADIUSX+i];
}
- //write the result to dst
- if((x<dst_cols) & (y<dst_rows))
+
+ // write the result to dst
+ if (x<dst_cols && y<dst_rows)
{
start_addr = mad24(y,dst_step_in_pixel,x);
dst[start_addr] = sum;
}
__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C4_D5
-(__global const float4 * restrict src,
- __global float4 * 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 float4 * restrict src,
+ __global float4 * 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);
int y = get_global_id(1);
__local float4 LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1];
#ifdef BORDER_CONSTANT
int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
- //read pixels from src
- for(i = 0; i<READ_TIMES_ROW; i++)
+
+ // read pixels from src
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
int current_addr = start_addr+i*LSIZE0;
current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
temp[i] = src[current_addr];
}
- //judge if read out of boundary
- for(i = 0; i<READ_TIMES_ROW; i++)
+
+ // judge if read out of boundary
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
temp[i]= ELEM(start_x+i*LSIZE0,0,src_whole_cols,(float4)0,temp[i]);
temp[i]= ELEM(start_y,0,src_whole_rows,(float4)0,temp[i]);
#else
int index[READ_TIMES_ROW];
int s_x,s_y;
- //judge if read out of boundary
- for(i = 0; i<READ_TIMES_ROW; i++)
+
+ // judge if read out of boundary
+ for (i = 0; i<READ_TIMES_ROW; i++)
{
- s_x= ADDR_L(start_x+i*LSIZE0,0,src_whole_cols,start_x+i*LSIZE0);
- s_x= ADDR_R(start_x+i*LSIZE0,src_whole_cols,s_x);
- s_y= ADDR_L(start_y,0,src_whole_rows,start_y);
- s_y= ADDR_R(start_y,src_whole_rows,s_y);
+ s_x = start_x + i*LSIZE0, s_y = start_y;
+ EXTRAPOLATE(s_x, src_whole_cols);
+ EXTRAPOLATE(s_y, src_whole_rows);
+
index[i]=mad24(s_y,src_step_in_pixel,s_x);
}
- //read pixels from src
- for(i = 0; i<READ_TIMES_ROW; i++)
- {
+ // read pixels from src
+ for (i = 0; i<READ_TIMES_ROW; i++)
temp[i] = src[index[i]];
- }
#endif
- //save pixels to lds
- for(i = 0; i<READ_TIMES_ROW; i++)
- {
+ // save pixels to lds
+ for (i = 0; i<READ_TIMES_ROW; i++)
LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
- }
barrier(CLK_LOCAL_MEM_FENCE);
- //read pixels from lds and calculate the result
+ // read pixels from lds and calculate the result
sum =LDS_DAT[l_y][l_x+RADIUSX]*mat_kernel[RADIUSX];
- for(i=1; i<=RADIUSX; i++)
+ for (i=1; i<=RADIUSX; i++)
{
temp[0]=LDS_DAT[l_y][l_x+RADIUSX-i];
temp[1]=LDS_DAT[l_y][l_x+RADIUSX+i];
sum += temp[0]*mat_kernel[RADIUSX-i]+temp[1]*mat_kernel[RADIUSX+i];
}
- //write the result to dst
- if((x<dst_cols) & (y<dst_rows))
+
+ // write the result to dst
+ if (x<dst_cols && y<dst_rows)
{
start_addr = mad24(y,dst_step_in_pixel,x);
dst[start_addr] = sum;
}
-
}
//
// * 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.
//
//M*/
-
-#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))
-#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (t_edge) :(i))
-#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (b_edge)-1 :(addr))
-#endif
-
+#ifdef BORDER_CONSTANT
+#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
+#elif defined BORDER_REPLICATE
+#define EXTRAPOLATE(x, maxV) \
+ { \
+ x = max(min(x, maxV - 1), 0); \
+ }
+#elif defined BORDER_WRAP
+#define EXTRAPOLATE(x, maxV) \
+ { \
+ if (x < 0) \
+ x -= ((x - maxV + 1) / maxV) * maxV; \
+ if (x >= maxV) \
+ x %= maxV; \
+ }
+#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
+#define EXTRAPOLATE_(x, maxV, delta) \
+ { \
+ if (maxV == 1) \
+ x = 0; \
+ else \
+ do \
+ { \
+ if ( x < 0 ) \
+ x = -x - 1 + delta; \
+ else \
+ x = maxV - 1 - (x - maxV) - delta; \
+ } \
+ while (x >= maxV || x < 0); \
+ }
#ifdef BORDER_REFLECT
-//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
-#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i)-1 : (i))
-#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr))
-#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i)-1 : (i))
-#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr))
-#endif
-
-#ifdef BORDER_REFLECT_101
-//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
-#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i))
-#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
-#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i) : (i))
-#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr))
+#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 0)
+#else
+#define EXTRAPOLATE(x, maxV) EXTRAPOLATE_(x, maxV, 1)
#endif
-
-//blur function does not support BORDER_WRAP
-#ifdef BORDER_WRAP
-//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
-#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (i)+(r_edge) : (i))
-#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (i)-(r_edge) : (addr))
-#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (i)+(b_edge) : (i))
-#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (i)-(b_edge) : (addr))
+#else
+#error No extrapolation method
#endif
__kernel void
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;
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];
}
if(col < (THREADS-(ksX-1)))
{
int4 currVal;
-
int howManyAll = (2*anX+1)*(ksY);
//find variance of all data
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);
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;
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);
#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];
}
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
#if defined (DOUBLE_SUPPORT)
-#ifdef cl_khr_fp64
-#pragma OPENCL EXTENSION cl_khr_fp64:enable
-#elif defined (cl_amd_fp64)
+#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
#endif
#ifdef BORDER_CONSTANT
-//BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii
-#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
-#endif
-
-#ifdef BORDER_REPLICATE
-//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
-#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (l_edge) : (addr)
-#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (r_edge)-1 : (addr)
-#endif
-
+#define EXTRAPOLATE(x, y, v) v = scalar;
+#elif defined BORDER_REPLICATE
+#define EXTRAPOLATE(x, y, v) \
+ { \
+ x = max(min(x, src_cols - 1), 0); \
+ y = max(min(y, src_rows - 1), 0); \
+ v = src[mad24(y, src_step, x + src_offset)]; \
+ }
+#elif defined BORDER_WRAP
+#define EXTRAPOLATE(x, y, v) \
+ { \
+ if (x < 0) \
+ x -= ((x - src_cols + 1) / src_cols) * src_cols; \
+ if (x >= src_cols) \
+ x %= src_cols; \
+ \
+ if (y < 0) \
+ y -= ((y - src_rows + 1) / src_rows) * src_rows; \
+ if( y >= src_rows ) \
+ y %= src_rows; \
+ v = src[mad24(y, src_step, x + src_offset)]; \
+ }
+#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
#ifdef BORDER_REFLECT
-//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
-#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i)-1 : (addr)
-#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)
+#define DELTA int delta = 0
+#else
+#define DELTA int delta = 1
#endif
-
-#ifdef BORDER_REFLECT_101
-//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
-#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i) : (addr)
-#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)
+#define EXTRAPOLATE(x, y, v) \
+ { \
+ DELTA; \
+ if (src_cols == 1) \
+ x = 0; \
+ else \
+ do \
+ { \
+ if( x < 0 ) \
+ x = -x - 1 + delta; \
+ else \
+ x = src_cols - 1 - (x - src_cols) - delta; \
+ } \
+ while (x >= src_cols || x < 0); \
+ \
+ if (src_rows == 1) \
+ y = 0; \
+ else \
+ do \
+ { \
+ if( y < 0 ) \
+ y = -y - 1 + delta; \
+ else \
+ y = src_rows - 1 - (y - src_rows) - delta; \
+ } \
+ while (y >= src_rows || y < 0); \
+ v = src[mad24(y, src_step, x + src_offset)]; \
+ }
+#else
+#error No extrapolation method
#endif
-#ifdef BORDER_WRAP
-//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
-#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (i)+(r_edge) : (addr)
-#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (i)-(r_edge) : (addr)
-#endif
+#define NEED_EXTRAPOLATION(gx, gy) (gx >= src_cols || gy >= src_rows || gx < 0 || gy < 0)
__kernel void copymakeborder
(__global const GENTYPE *src,
__global GENTYPE *dst,
- const int dst_cols,
- const int dst_rows,
- const int src_cols,
- const int src_rows,
- const int src_step_in_pixel,
- const int src_offset_in_pixel,
- const int dst_step_in_pixel,
- const int dst_offset_in_pixel,
- const int top,
- const int left,
- const GENTYPE val
- )
+ int dst_cols, int dst_rows,
+ int src_cols, int src_rows,
+ int src_step, int src_offset,
+ int dst_step, int dst_offset,
+ int top, int left, GENTYPE scalar)
{
int x = get_global_id(0);
int y = get_global_id(1);
- int src_x = x-left;
- int src_y = y-top;
- int src_addr = mad24(src_y,src_step_in_pixel,src_x+src_offset_in_pixel);
- int dst_addr = mad24(y,dst_step_in_pixel,x+dst_offset_in_pixel);
- int con = (src_x >= 0) && (src_x < src_cols) && (src_y >= 0) && (src_y < src_rows);
- if(con)
- {
- dst[dst_addr] = src[src_addr];
- }
- else
- {
- #ifdef BORDER_CONSTANT
- //write the result to dst
- if((x<dst_cols) && (y<dst_rows))
- {
- dst[dst_addr] = val;
- }
- #else
- int s_x,s_y;
- //judge if read out of boundary
- s_x= ADDR_L(src_x,0,src_cols,src_x);
- s_x= ADDR_R(src_x,src_cols,s_x);
- s_y= ADDR_L(src_y,0,src_rows,src_y);
- s_y= ADDR_R(src_y,src_rows,s_y);
- src_addr=mad24(s_y,src_step_in_pixel,s_x+src_offset_in_pixel);
- //write the result to dst
- if((x<dst_cols) && (y<dst_rows))
- {
- dst[dst_addr] = src[src_addr];
- }
- #endif
- }
-}
-__kernel void copymakeborder_C1_D0
- (__global const uchar *src,
- __global uchar *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 uchar val
- )
-{
- int x = get_global_id(0)<<2;
- 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+3 < src_cols) && (src_y >= 0) && (src_y < src_rows);
- if(con)
+ if (x < dst_cols && y < dst_rows)
{
- uchar4 tmp = vload4(0,src+src_addr);
- *(__global uchar4*)(dst+dst_addr) = tmp;
- }
- else
- {
- #ifdef BORDER_CONSTANT
- //write the result to dst
- if((((src_x<0) && (src_x+3>=0))||(src_x < src_cols) && (src_x+3 >= src_cols)) && (src_y >= 0) && (src_y < src_rows))
- {
- int4 addr;
- uchar4 tmp;
- addr.x = ((src_x < 0) || (src_x>= src_cols)) ? 0 : src_addr;
- addr.y = ((src_x+1 < 0) || (src_x+1>= src_cols)) ? 0 : (src_addr+1);
- addr.z = ((src_x+2 < 0) || (src_x+2>= src_cols)) ? 0 : (src_addr+2);
- addr.w = ((src_x+3 < 0) || (src_x+3>= src_cols)) ? 0 : (src_addr+3);
- tmp.x = src[addr.x];
- tmp.y = src[addr.y];
- tmp.z = src[addr.z];
- tmp.w = src[addr.w];
- tmp.x = (src_x >=0)&&(src_x < src_cols) ? tmp.x : val;
- tmp.y = (src_x+1 >=0)&&(src_x +1 < src_cols) ? tmp.y : val;
- tmp.z = (src_x+2 >=0)&&(src_x +2 < src_cols) ? tmp.z : val;
- tmp.w = (src_x+3 >=0)&&(src_x +3 < src_cols) ? tmp.w : val;
- *(__global uchar4*)(dst+dst_addr) = tmp;
- }
- else if((x<dst_cols) && (y<dst_rows))
- {
- *(__global uchar4*)(dst+dst_addr) = (uchar4)val;
- }
- #else
- int4 s_x;
- int s_y;
- //judge if read out of boundary
- s_x.x= ADDR_L(src_x,0,src_cols,src_x);
- s_x.y= ADDR_L(src_x+1,0,src_cols,src_x+1);
- s_x.z= ADDR_L(src_x+2,0,src_cols,src_x+2);
- s_x.w= ADDR_L(src_x+3,0,src_cols,src_x+3);
- s_x.x= ADDR_R(src_x,src_cols,s_x.x);
- s_x.y= ADDR_R(src_x+1,src_cols,s_x.y);
- s_x.z= ADDR_R(src_x+2,src_cols,s_x.z);
- s_x.w= ADDR_R(src_x+3,src_cols,s_x.w);
- s_y= ADDR_L(src_y,0,src_rows,src_y);
- s_y= ADDR_R(src_y,src_rows,s_y);
- int4 src_addr4=mad24((int4)s_y,(int4)src_step_in_pixel,s_x+(int4)src_offset_in_pixel);
- //write the result to dst
- if((x<dst_cols) && (y<dst_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
{
- uchar4 tmp;
- tmp.x = src[src_addr4.x];
- tmp.y = src[src_addr4.y];
- tmp.z = src[src_addr4.z];
- tmp.w = src[src_addr4.w];
- *(__global uchar4*)(dst+dst_addr) = tmp;
+ int src_index = mad24(src_y, src_step, src_x + src_offset);
+ dst[dst_index] = src[src_index];
}
- #endif
}
}
//
// * 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.
globalHist[mad24(gx, hist_step, lid)] = bin1+bin2+bin3+bin4;
}
-__kernel void __attribute__((reqd_work_group_size(1,HISTOGRAM256_BIN_COUNT,1)))calc_sub_hist_border_D0(
- __global const uchar* src,
- int src_step, int src_offset,
- __global int* globalHist,
- int left_col, int cols,
- int rows, int hist_step)
+__kernel void __attribute__((reqd_work_group_size(1,HISTOGRAM256_BIN_COUNT,1)))
+calc_sub_hist_border_D0(__global const uchar* src, int src_step, int src_offset,
+ __global int* globalHist, int left_col, int cols,
+ int rows, int hist_step)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
globalHist[mad24(rowIndex, hist_step, lidy)] += subhist[lidy];
}
+
__kernel __attribute__((reqd_work_group_size(256,1,1)))void merge_hist(__global int* buf,
__global int* hist,
int src_step)
hist[gx] = data[0];
}
-__kernel __attribute__((reqd_work_group_size(256,1,1)))void calLUT(
- __global uchar * dst,
- __constant int * hist,
- int total)
+__kernel __attribute__((reqd_work_group_size(256,1,1)))
+void calLUT(__global uchar * dst, __constant int * hist, int total)
{
int lid = get_local_id(0);
- __local int sumhist[HISTOGRAM256_BIN_COUNT+1];
+ __local int sumhist[HISTOGRAM256_BIN_COUNT];
+ __local float scale;
- sumhist[lid]=hist[lid];
+ sumhist[lid] = hist[lid];
barrier(CLK_LOCAL_MEM_FENCE);
- if(lid==0)
+ if (lid == 0)
{
- int sum = 0;
- int i = 0;
- while (!sumhist[i]) ++i;
- sumhist[HISTOGRAM256_BIN_COUNT] = sumhist[i];
- for(sumhist[i++] = 0; i<HISTOGRAM256_BIN_COUNT; i++)
+ int sum = 0, i = 0;
+ while (!sumhist[i])
+ ++i;
+
+ if (total == sumhist[i])
{
- sum+=sumhist[i];
- sumhist[i]=sum;
+ scale = 1;
+ for (int j = 0; j < HISTOGRAM256_BIN_COUNT; ++j)
+ sumhist[i] = i;
+ }
+ else
+ {
+ scale = 255.f/(total - sumhist[i]);
+
+ for (sumhist[i++] = 0; i < HISTOGRAM256_BIN_COUNT; i++)
+ {
+ sum += sumhist[i];
+ sumhist[i] = sum;
+ }
}
}
+
barrier(CLK_LOCAL_MEM_FENCE);
- float scale = 255.f/(total - sumhist[HISTOGRAM256_BIN_COUNT]);
- dst[lid]= lid == 0 ? 0 : convert_uchar_sat(convert_float(sumhist[lid])*scale);
+ dst[lid]= convert_uchar_sat_rte(convert_float(sumhist[lid])*scale);
}
+
/*
///////////////////////////////equalizeHist//////////////////////////////////////////////////
__kernel __attribute__((reqd_work_group_size(256,1,1)))void equalizeHist(
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
swap(M, bufM);
- finish();
-
optflow_farneback::updateFlowOcl(M, flowx, flowy);
if (updateMatrices)
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
using namespace testing;
using namespace std;
+static bool relativeError(double actual, double expected, double eps)
+{
+ return std::abs(actual - expected) / actual < eps;
+}
+
//////////////////////////////// LUT /////////////////////////////////////////////////
PARAM_TEST_CASE(Lut, MatDepth, MatDepth, bool, bool)
depth = GET_PARAM(0);
cn = GET_PARAM(1);
use_roi = GET_PARAM(2);
-
- val = cv::Scalar(rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0),
- rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0));
}
void random_roi()
generateOclMat(gdst1_whole, gdst1_roi, dst1, roiSize, dst1Border);
generateOclMat(gdst2_whole, gdst2_roi, dst2, roiSize, dst2Border);
generateOclMat(gmask_whole, gmask_roi, mask, roiSize, maskBorder);
+
+ val = cv::Scalar(rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0),
+ rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0));
}
void Near(double threshold = 0.)
{
random_roi();
- cv::multiply(val[0], src1_roi, dst1_roi);
+ cv::multiply(Scalar::all(val[0]), src1_roi, dst1_roi);
cv::ocl::multiply(val[0], gsrc1_roi, gdst1_roi);
Near(gdst1_roi.depth() >= CV_32F ? 1e-3 : 1);
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);
}
}
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);
}
}
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
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,
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),
//
// * 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.
//
// * 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.
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);
}
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
//
// * 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.
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);
{
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);
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
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);
{
float* btvWeights_ = NULL;
size_t btvWeights_size = 0;
+ oclMat c_btvRegWeights;
}
}
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);
}
-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();
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));
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
{
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);
}
btvWeights_ = &btvWeights[0];
btvWeights_size = size;
+ Mat btvWeights_mheader(1, static_cast<int>(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;
oclMat highRes_;
vector<oclMat> diffTerms_;
- vector<oclMat> a_, b_, c_;
+ oclMat a_, b_, c_, d_;
oclMat regTerm_;
};
btvKernelSize_ = 7;
blurKernelSize_ = 5;
blurSigma_ = 0.0;
- opticalFlow_ = createOptFlow_DualTVL1_OCL();
+ opticalFlow_ = createOptFlow_Farneback_OCL();
curBlurKernelSize_ = -1;
curBlurSigma_ = -1.0;
// 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)
highRes_.release();
diffTerms_.clear();
- a_.clear();
- b_.clear();
- c_.clear();
+ a_.release();
+ b_.release();
+ c_.release();
regTerm_.release();
+ c_btvRegWeights.release();
}
////////////////////////////////////////////////////////////
//
// * 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.
//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);
}
__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);
if(channels == 1)
{
dst[y * scale * dst_step + x * scale] = src[y * src_step + x];
- }else if(channels == 3)
- {
- dst[y * channels * scale * dst_step + 3 * x * scale + 0] = src[y * channels * src_step + 3 * x + 0];
- dst[y * channels * scale * dst_step + 3 * x * scale + 1] = src[y * channels * src_step + 3 * x + 1];
- dst[y * channels * scale * dst_step + 3 * x * scale + 2] = src[y * channels * src_step + 3 * x + 2];
- }else
+ }
+ else
{
- dst[y * channels * scale * dst_step + 4 * x * scale + 0] = src[y * channels * src_step + 4 * x + 0];
- dst[y * channels * scale * dst_step + 4 * x * scale + 1] = src[y * channels * src_step + 4 * x + 1];
- dst[y * channels * scale * dst_step + 4 * x * scale + 2] = src[y * channels * src_step + 4 * x + 2];
- dst[y * channels * scale * dst_step + 4 * x * scale + 3] = src[y * channels * src_step + 4 * x + 3];
+ vstore4(vload4(0, src + y * channels * src_step + 4 * x), 0, dst + y * channels * scale * dst_step + 4 * x * scale);
}
}
}
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;
}
__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);
{
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;
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)
{
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);
}
}
}
# 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}\"
" )
if(useOcl)
{
- MEASURE_TIME(superRes->nextFrame(result_));
+ MEASURE_TIME(
+ {
+ superRes->nextFrame(result_);
+ ocl::finish();
+ });
}
else
#endif
{
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<string>("i");
int ks = cmd.get<int>("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;
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;
}
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<bool>("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<bool>("camera");
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
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;
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)
{
break;
}
- for(;;)
+ for (;;)
{
cap >> frame;
if (frame.empty())
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;
}
Ptr<CLAHE> 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)
"{ i input | | specify input image }"
"{ c camera | 0 | specify camera id }"
"{ s use_cpu | false | use cpu algorithm }"
- "{ o output | clahe_output.jpg | specify output save path}";
+ "{ o output | clahe_output.jpg | specify output save path}"
+ "{ h help | false | print help message }";
+
+ cv::CommandLineParser cmd(argc, argv, keys);
+ if (cmd.has("help"))
+ {
+ cout << "Usage : clahe [options]" << endl;
+ cout << "Available options:" << endl;
+ cmd.printMessage();
+ return EXIT_SUCCESS;
+ }
- CommandLineParser cmd(argc, argv, keys);
- string infile = cmd.get<string>("i");
- outfile = cmd.get<string>("o");
+ string infile = cmd.get<string>("i"), outfile = cmd.get<string>("o");
int camid = cmd.get<int>("c");
bool use_cpu = cmd.get<bool>("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);
}
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;
}
{
work_begin = getTickCount();
}
+
static void workEnd()
{
work_end += (getTickCount() - work_begin);
return work_end /((double)cvGetTickFrequency() * 1000.);
}
-void detect( Mat& img, vector<Rect>& faces,
+
+static void detect( Mat& img, vector<Rect>& faces,
ocl::OclCascadeClassifier& cascade,
double scale, bool calTime);
-void detectCPU( Mat& img, vector<Rect>& faces,
+static void detectCPU( Mat& img, vector<Rect>& faces,
CascadeClassifier& cascade,
double scale, bool calTime);
-void Draw(Mat& img, vector<Rect>& faces, double scale);
+static void Draw(Mat& img, vector<Rect>& faces, double scale);
// This function test if gpu_rst matches cpu_rst.
// Else if will return (total diff of each cpu and gpu rects covered pixels)/(total cpu rects covered pixels)
double checkRectSimilarity(Size sz, vector<Rect>& cpu_rst, vector<Rect>& gpu_rst);
-
int main( int argc, const char** argv )
{
const char* keys =
CommandLineParser cmd(argc, argv, keys);
if (cmd.get<bool>("help"))
{
+ cout << "Usage : facedetect [options]" << endl;
cout << "Available options:" << endl;
cmd.printMessage();
- return 0;
+ return EXIT_SUCCESS;
}
+
CvCapture* capture = 0;
Mat frame, frameCopy, image;
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() )
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 )
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
{
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);
"{ 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
{
{
return cout << "unknown exception" << endl, 1;
}
- return 0;
+ return EXIT_SUCCESS;
}
App::App(CommandLineParser& cmd)
d_mat.download(mat);
}
-static void drawArrows(Mat& frame, const vector<Point2f>& prevPts, const vector<Point2f>& nextPts, const vector<uchar>& status, Scalar line_color = Scalar(0, 0, 255))
+static void drawArrows(Mat& frame, const vector<Point2f>& prevPts, const vector<Point2f>& nextPts, const vector<uchar>& status,
+ Scalar line_color = Scalar(0, 0, 255))
{
for (size_t i = 0; i < prevPts.size(); ++i)
{
cout << "Usage: pyrlk_optical_flow [options]" << endl;
cout << "Available options:" << endl;
cmd.printMessage();
- return 0;
+ return EXIT_SUCCESS;
}
bool defaultPicturesFail = false;
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 == "")
+ if(vdofile.empty())
cout << "Capture from CAM " << c << " didn't work" << endl;
else
cout << "Capture from file " << vdofile << " failed" <<endl;
if (defaultPicturesFail)
- {
- return -1;
- }
+ return EXIT_FAILURE;
goto nocamera;
}
}
if( waitKey( 10 ) >= 0 )
- goto _cleanup_;
+ break;
}
- waitKey(0);
-
-_cleanup_:
capture.release();
}
else
waitKey();
- return 0;
+ return EXIT_SUCCESS;
}
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(
{
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<string>("i");
string outfile = cmd.get<string>("o");
- if(inputName.empty())
+
+ if(cmd.get<bool>("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<vector<Point> > 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
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
imwrite(outfile, result);
waitKey(0);
- return 0;
+ return EXIT_SUCCESS;
}
"{ 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<bool>("help"))
{
cmd.printMessage();
return 0;
}
+
try
{
App app(cmd);
{
cout << "error: " << e.what() << endl;
}
- return 0;
+
+ return EXIT_SUCCESS;
}
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<string>("l");
r_img = cmd.get<string>("r");
string mstr = cmd.get<string>("m");
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.);
}
}
};
-Mat drawGoodMatches(
+static Mat drawGoodMatches(
const Mat& cpu_img1,
const Mat& cpu_img2,
const std::vector<KeyPoint>& keypoints1,
return img_matches;
}
-}
////////////////////////////////////////////////////
// This program demonstrates the usage of SURF_OCL.
// use cpu findHomography interface to calculate the transformation matrix
"{ 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<bool>("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;
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
imshow("ocl surf matches", ocl_img_matches);
}
waitKey(0);
- return 0;
+ return EXIT_SUCCESS;
}
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<string>("l");
string fname1 = cmd.get<string>("r");
string vdofile = cmd.get<string>("v");
cv::Ptr<cv::DenseOpticalFlow> 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)
{
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" <<endl;
- if (defaultPicturesFail)
- {
- return -1;
- }
goto nocamera;
}
}
if( waitKey( 10 ) >= 0 )
- goto _cleanup_;
+ break;
}
- waitKey(0);
-
-_cleanup_:
capture.release();
}
else
waitKey();
- return 0;
+ return EXIT_SUCCESS;
}
projects / profile / ivi / opencv.git / commitdiff