From 118709fd9f3d5ad036fdd596dc1e777ee9703589 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 1 Jan 2014 23:46:19 +0400 Subject: [PATCH] added OpenCL accelerated warpers --- cmake/OpenCVCompilerOptions.cmake~ | 308 +++++++++++++++++++++ modules/stitching/doc/warpers.rst~ | 263 ++++++++++++++++++ .../include/opencv2/stitching/detail/warpers.hpp | 43 ++- .../include/opencv2/stitching/warpers.hpp | 18 ++ modules/stitching/src/opencl/warpers.cl | 148 ++++++++++ modules/stitching/src/precomp.hpp | 1 + modules/stitching/src/warpers_ocl.cpp | 184 ++++++++++++ modules/stitching/test/ocl/test_warpers.cpp | 149 ++++++++++ samples/cpp/stitching_detailed.cpp | 54 +++- 9 files changed, 1152 insertions(+), 16 deletions(-) create mode 100644 cmake/OpenCVCompilerOptions.cmake~ create mode 100644 modules/stitching/doc/warpers.rst~ create mode 100644 modules/stitching/src/opencl/warpers.cl create mode 100644 modules/stitching/src/warpers_ocl.cpp create mode 100644 modules/stitching/test/ocl/test_warpers.cpp diff --git a/cmake/OpenCVCompilerOptions.cmake~ b/cmake/OpenCVCompilerOptions.cmake~ new file mode 100644 index 0000000..59b19b6 --- /dev/null +++ b/cmake/OpenCVCompilerOptions.cmake~ @@ -0,0 +1,308 @@ +if(MINGW OR (X86 AND UNIX AND NOT APPLE)) + # mingw compiler is known to produce unstable SSE code with -O3 hence we are trying to use -O2 instead + if(CMAKE_COMPILER_IS_GNUCXX) + foreach(flags CMAKE_CXX_FLAGS CMAKE_CXX_FLAGS_RELEASE CMAKE_CXX_FLAGS_DEBUG) + string(REPLACE "-O3" "-O2" ${flags} "${${flags}}") + endforeach() + endif() + + if(CMAKE_COMPILER_IS_GNUCC) + foreach(flags CMAKE_C_FLAGS CMAKE_C_FLAGS_RELEASE CMAKE_C_FLAGS_DEBUG) + string(REPLACE "-O3" "-O2" ${flags} "${${flags}}") + endforeach() + endif() +endif() + +if(MSVC) + string(REGEX REPLACE "^ *| * $" "" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") + string(REGEX REPLACE "^ *| * $" "" CMAKE_CXX_FLAGS_INIT "${CMAKE_CXX_FLAGS_INIT}") + if(CMAKE_CXX_FLAGS STREQUAL CMAKE_CXX_FLAGS_INIT) + # override cmake default exception handling option + string(REPLACE "/EHsc" "/EHa" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}" CACHE STRING "Flags used by the compiler during all build types." FORCE) + endif() +endif() + +set(OPENCV_EXTRA_FLAGS "") +set(OPENCV_EXTRA_C_FLAGS "") +set(OPENCV_EXTRA_CXX_FLAGS "") +set(OPENCV_EXTRA_FLAGS_RELEASE "") +set(OPENCV_EXTRA_FLAGS_DEBUG "") +set(OPENCV_EXTRA_EXE_LINKER_FLAGS "") +set(OPENCV_EXTRA_EXE_LINKER_FLAGS_RELEASE "") +set(OPENCV_EXTRA_EXE_LINKER_FLAGS_DEBUG "") + +macro(add_extra_compiler_option option) + if(CMAKE_BUILD_TYPE) + set(CMAKE_TRY_COMPILE_CONFIGURATION ${CMAKE_BUILD_TYPE}) + endif() + ocv_check_flag_support(CXX "${option}" _varname "${OPENCV_EXTRA_CXX_FLAGS} ${ARGN}") + if(${_varname}) + set(OPENCV_EXTRA_CXX_FLAGS "${OPENCV_EXTRA_CXX_FLAGS} ${option}") + endif() + + ocv_check_flag_support(C "${option}" _varname "${OPENCV_EXTRA_C_FLAGS} ${ARGN}") + if(${_varname}) + set(OPENCV_EXTRA_C_FLAGS "${OPENCV_EXTRA_C_FLAGS} ${option}") + endif() +endmacro() + +# OpenCV fails some tests when 'char' is 'unsigned' by default +add_extra_compiler_option(-fsigned-char) + +if(MINGW) + # http://gcc.gnu.org/bugzilla/show_bug.cgi?id=40838 + # here we are trying to workaround the problem + add_extra_compiler_option(-mstackrealign) + if(NOT HAVE_CXX_MSTACKREALIGN) + add_extra_compiler_option(-mpreferred-stack-boundary=2) + endif() +endif() + +if(CMAKE_COMPILER_IS_GNUCXX) + # High level of warnings. + add_extra_compiler_option(-W) + add_extra_compiler_option(-Wall) + add_extra_compiler_option(-Werror=return-type) + add_extra_compiler_option(-Werror=non-virtual-dtor) + add_extra_compiler_option(-Werror=address) + add_extra_compiler_option(-Werror=sequence-point) + add_extra_compiler_option(-Wformat) + add_extra_compiler_option(-Werror=format-security -Wformat) + add_extra_compiler_option(-Wmissing-declarations) + add_extra_compiler_option(-Wmissing-prototypes) + add_extra_compiler_option(-Wstrict-prototypes) + add_extra_compiler_option(-Wundef) + add_extra_compiler_option(-Winit-self) + add_extra_compiler_option(-Wpointer-arith) + add_extra_compiler_option(-Wshadow) + add_extra_compiler_option(-Wsign-promo) + + if(ENABLE_NOISY_WARNINGS) + add_extra_compiler_option(-Wcast-align) + add_extra_compiler_option(-Wstrict-aliasing=2) + else() + add_extra_compiler_option(-Wno-narrowing) + add_extra_compiler_option(-Wno-delete-non-virtual-dtor) + add_extra_compiler_option(-Wno-unnamed-type-template-args) + endif() + add_extra_compiler_option(-fdiagnostics-show-option) + + # The -Wno-long-long is required in 64bit systems when including sytem headers. + if(X86_64) + add_extra_compiler_option(-Wno-long-long) + endif() + + # We need pthread's + if(UNIX AND NOT ANDROID AND NOT (APPLE AND CMAKE_COMPILER_IS_CLANGCXX)) + add_extra_compiler_option(-pthread) + endif() + + if(OPENCV_WARNINGS_ARE_ERRORS) + add_extra_compiler_option(-Werror) + endif() + + if(X86 AND NOT MINGW64 AND NOT X86_64 AND NOT APPLE) + add_extra_compiler_option(-march=i686) + endif() + + # Other optimizations + if(ENABLE_OMIT_FRAME_POINTER) + add_extra_compiler_option(-fomit-frame-pointer) + else() + add_extra_compiler_option(-fno-omit-frame-pointer) + endif() + if(ENABLE_FAST_MATH) + add_extra_compiler_option(-ffast-math) + endif() + if(ENABLE_POWERPC) + add_extra_compiler_option("-mcpu=G3 -mtune=G5") + endif() + if(ENABLE_SSE) + add_extra_compiler_option(-msse) + endif() + if(ENABLE_SSE2) + add_extra_compiler_option(-msse2) + endif() + if (ENABLE_NEON) + add_extra_compiler_option("-mfpu=neon") + endif() + if (ENABLE_VFPV3 AND NOT ENABLE_NEON) + add_extra_compiler_option("-mfpu=vfpv3") + endif() + + # SSE3 and further should be disabled under MingW because it generates compiler errors + if(NOT MINGW) + if(ENABLE_AVX) + add_extra_compiler_option(-mavx) + endif() + + # GCC depresses SSEx instructions when -mavx is used. Instead, it generates new AVX instructions or AVX equivalence for all SSEx instructions when needed. + if(NOT OPENCV_EXTRA_CXX_FLAGS MATCHES "-mavx") + if(ENABLE_SSE3) + add_extra_compiler_option(-msse3) + endif() + + if(ENABLE_SSSE3) + add_extra_compiler_option(-mssse3) + endif() + + if(ENABLE_SSE41) + add_extra_compiler_option(-msse4.1) + endif() + + if(ENABLE_SSE42) + add_extra_compiler_option(-msse4.2) + endif() + endif() + endif(NOT MINGW) + + if(X86 OR X86_64) + if(NOT APPLE AND CMAKE_SIZEOF_VOID_P EQUAL 4) + if(OPENCV_EXTRA_CXX_FLAGS MATCHES "-m(sse2|avx)") + add_extra_compiler_option(-mfpmath=sse)# !! important - be on the same wave with x64 compilers + else() + add_extra_compiler_option(-mfpmath=387) + endif() + endif() + endif() + + if(ENABLE_NEON) + add_extra_compiler_option(-mfpu=neon) + endif() + + # Profiling? + if(ENABLE_PROFILING) + add_extra_compiler_option("-pg -g") + # turn off incompatible options + foreach(flags CMAKE_CXX_FLAGS CMAKE_C_FLAGS CMAKE_CXX_FLAGS_RELEASE CMAKE_C_FLAGS_RELEASE CMAKE_CXX_FLAGS_DEBUG CMAKE_C_FLAGS_DEBUG + OPENCV_EXTRA_FLAGS_RELEASE OPENCV_EXTRA_FLAGS_DEBUG OPENCV_EXTRA_C_FLAGS OPENCV_EXTRA_CXX_FLAGS) + string(REPLACE "-fomit-frame-pointer" "" ${flags} "${${flags}}") + string(REPLACE "-ffunction-sections" "" ${flags} "${${flags}}") + endforeach() + elseif(NOT APPLE AND NOT ANDROID) + # Remove unreferenced functions: function level linking + add_extra_compiler_option(-ffunction-sections) + endif() + + set(OPENCV_EXTRA_FLAGS_RELEASE "${OPENCV_EXTRA_FLAGS_RELEASE} -DNDEBUG") + set(OPENCV_EXTRA_FLAGS_DEBUG "${OPENCV_EXTRA_FLAGS_DEBUG} -O0 -DDEBUG -D_DEBUG") +endif() + +if(MSVC) + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /D _CRT_SECURE_NO_DEPRECATE /D _CRT_NONSTDC_NO_DEPRECATE /D _SCL_SECURE_NO_WARNINGS") + # 64-bit portability warnings, in MSVC80 + if(MSVC80) + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /Wp64") + endif() + + if(BUILD_WITH_DEBUG_INFO) + set(OPENCV_EXTRA_EXE_LINKER_FLAGS_RELEASE "${OPENCV_EXTRA_EXE_LINKER_FLAGS_RELEASE} /debug") + endif() + + # Remove unreferenced functions: function level linking + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /Gy") + if(NOT MSVC_VERSION LESS 1400) + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /bigobj") + endif() + if(BUILD_WITH_DEBUG_INFO) + set(OPENCV_EXTRA_FLAGS_RELEASE "${OPENCV_EXTRA_FLAGS_RELEASE} /Zi") + endif() + + if(ENABLE_AVX AND NOT MSVC_VERSION LESS 1600) + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /arch:AVX") + endif() + + if(ENABLE_SSE4_1 AND CV_ICC AND NOT OPENCV_EXTRA_FLAGS MATCHES "/arch:") + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /arch:SSE4.1") + endif() + + if(ENABLE_SSE3 AND CV_ICC AND NOT OPENCV_EXTRA_FLAGS MATCHES "/arch:") + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /arch:SSE3") + endif() + + if(NOT MSVC64) + # 64-bit MSVC compiler uses SSE/SSE2 by default + if(ENABLE_SSE2 AND NOT OPENCV_EXTRA_FLAGS MATCHES "/arch:") + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /arch:SSE2") + endif() + if(ENABLE_SSE AND NOT OPENCV_EXTRA_FLAGS MATCHES "/arch:") + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /arch:SSE") + endif() + endif() + + if(ENABLE_SSE OR ENABLE_SSE2 OR ENABLE_SSE3 OR ENABLE_SSE4_1 OR ENABLE_AVX) + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /Oi") + endif() + + if(X86 OR X86_64) + if(CMAKE_SIZEOF_VOID_P EQUAL 4 AND ENABLE_SSE2) + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /fp:fast") # !! important - be on the same wave with x64 compilers + endif() + endif() + + if(OPENCV_WARNINGS_ARE_ERRORS) + set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS} /WX") + endif() +endif() + +# Extra link libs if the user selects building static libs: +if(NOT BUILD_SHARED_LIBS AND CMAKE_COMPILER_IS_GNUCXX AND NOT ANDROID) + # Android does not need these settings because they are already set by toolchain file + set(OPENCV_LINKER_LIBS ${OPENCV_LINKER_LIBS} stdc++) + set(OPENCV_EXTRA_FLAGS "-fPIC ${OPENCV_EXTRA_FLAGS}") +endif() + +# Add user supplied extra options (optimization, etc...) +# ========================================================== +set(OPENCV_EXTRA_FLAGS "${OPENCV_EXTRA_FLAGS}" CACHE INTERNAL "Extra compiler options") +set(OPENCV_EXTRA_C_FLAGS "${OPENCV_EXTRA_C_FLAGS}" CACHE INTERNAL "Extra compiler options for C sources") +set(OPENCV_EXTRA_CXX_FLAGS "${OPENCV_EXTRA_CXX_FLAGS}" CACHE INTERNAL "Extra compiler options for C++ sources") +set(OPENCV_EXTRA_FLAGS_RELEASE "${OPENCV_EXTRA_FLAGS_RELEASE}" CACHE INTERNAL "Extra compiler options for Release build") +set(OPENCV_EXTRA_FLAGS_DEBUG "${OPENCV_EXTRA_FLAGS_DEBUG}" CACHE INTERNAL "Extra compiler options for Debug build") +set(OPENCV_EXTRA_EXE_LINKER_FLAGS "${OPENCV_EXTRA_EXE_LINKER_FLAGS}" CACHE INTERNAL "Extra linker flags") +set(OPENCV_EXTRA_EXE_LINKER_FLAGS_RELEASE "${OPENCV_EXTRA_EXE_LINKER_FLAGS_RELEASE}" CACHE INTERNAL "Extra linker flags for Release build") +set(OPENCV_EXTRA_EXE_LINKER_FLAGS_DEBUG "${OPENCV_EXTRA_EXE_LINKER_FLAGS_DEBUG}" CACHE INTERNAL "Extra linker flags for Debug build") + +# set default visibility to hidden +if(CMAKE_COMPILER_IS_GNUCXX AND CMAKE_OPENCV_GCC_VERSION_NUM GREATER 399) + add_extra_compiler_option(-fvisibility=hidden) + add_extra_compiler_option(-fvisibility-inlines-hidden) +endif() + +#combine all "extra" options +set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${OPENCV_EXTRA_FLAGS} ${OPENCV_EXTRA_C_FLAGS}") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OPENCV_EXTRA_FLAGS} ${OPENCV_EXTRA_CXX_FLAGS}") +set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} ${OPENCV_EXTRA_FLAGS_RELEASE}") +set(CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE} ${OPENCV_EXTRA_FLAGS_RELEASE}") +set(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} ${OPENCV_EXTRA_FLAGS_DEBUG}") +set(CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS_DEBUG} ${OPENCV_EXTRA_FLAGS_DEBUG}") +set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} ${OPENCV_EXTRA_EXE_LINKER_FLAGS}") +set(CMAKE_EXE_LINKER_FLAGS_RELEASE "${CMAKE_EXE_LINKER_FLAGS_RELEASE} ${OPENCV_EXTRA_EXE_LINKER_FLAGS_RELEASE}") +set(CMAKE_EXE_LINKER_FLAGS_DEBUG "${CMAKE_EXE_LINKER_FLAGS_DEBUG} ${OPENCV_EXTRA_EXE_LINKER_FLAGS_DEBUG}") + +if(MSVC) + # avoid warnings from MSVC about overriding the /W* option + # we replace /W3 with /W4 only for C++ files, + # since all the 3rd-party libraries OpenCV uses are in C, + # and we do not care about their warnings. + string(REPLACE "/W3" "/W4" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") + string(REPLACE "/W3" "/W4" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}") + string(REPLACE "/W3" "/W4" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}") + + if(NOT ENABLE_NOISY_WARNINGS AND MSVC_VERSION EQUAL 1400) + ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4510 /wd4610 /wd4312 /wd4201 /wd4244 /wd4328 /wd4267) + endif() + + # allow extern "C" functions throw exceptions + foreach(flags CMAKE_C_FLAGS CMAKE_C_FLAGS_RELEASE CMAKE_C_FLAGS_RELEASE CMAKE_CXX_FLAGS CMAKE_CXX_FLAGS_RELEASE CMAKE_CXX_FLAGS_DEBUG) + string(REPLACE "/EHsc-" "/EHs" ${flags} "${${flags}}") + string(REPLACE "/EHsc" "/EHs" ${flags} "${${flags}}") + + string(REPLACE "/Zm1000" "" ${flags} "${${flags}}") + endforeach() + + if(NOT ENABLE_NOISY_WARNINGS) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4251") #class 'std::XXX' needs to have dll-interface to be used by clients of YYY + endif() +endif() diff --git a/modules/stitching/doc/warpers.rst~ b/modules/stitching/doc/warpers.rst~ new file mode 100644 index 0000000..1025ffa --- /dev/null +++ b/modules/stitching/doc/warpers.rst~ @@ -0,0 +1,263 @@ +Images Warping +============== + +.. highlight:: cpp + +detail::RotationWarper +---------------------- +.. ocv:class:: detail::RotationWarper + +Rotation-only model image warper interface. :: + + class CV_EXPORTS RotationWarper + { + public: + virtual ~RotationWarper() {} + + virtual Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R) = 0; + + virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) = 0; + + virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, + Mat &dst) = 0; + + virtual void warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, + Size dst_size, Mat &dst) = 0; + + virtual Rect warpRoi(Size src_size, const Mat &K, const Mat &R) = 0; + }; + +detail::RotationWarper::warpPoint +--------------------------------- + +Projects the image point. + +.. ocv:function:: Point2f detail::RotationWarper::warpPoint(const Point2f &pt, const Mat &K, const Mat &R) + + :param pt: Source point + + :param K: Camera intrinsic parameters + + :param R: Camera rotation matrix + + :return: Projected point + +detail::RotationWarper::buildMaps +--------------------------------- + +Builds the projection maps according to the given camera data. + +.. ocv:function:: Rect detail::RotationWarper::buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) + + :param src_size: Source image size + + :param K: Camera intrinsic parameters + + :param R: Camera rotation matrix + + :param xmap: Projection map for the x axis + + :param ymap: Projection map for the y axis + + :return: Projected image minimum bounding box + +detail::RotationWarper::warp +---------------------------- + +Projects the image. + +.. ocv:function:: Point detail::RotationWarper::warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Mat &dst) + + :param src: Source image + + :param K: Camera intrinsic parameters + + :param R: Camera rotation matrix + + :param interp_mode: Interpolation mode + + :param border_mode: Border extrapolation mode + + :param dst: Projected image + + :return: Project image top-left corner + +detail::RotationWarper::warpBackward +------------------------------------ + +Projects the image backward. + +.. ocv:function:: void detail::RotationWarper::warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Size dst_size, Mat &dst) + + :param src: Projected image + + :param K: Camera intrinsic parameters + + :param R: Camera rotation matrix + + :param interp_mode: Interpolation mode + + :param border_mode: Border extrapolation mode + + :param dst_size: Backward-projected image size + + :param dst: Backward-projected image + +detail::RotationWarper::warpRoi +------------------------------- + +.. ocv:function:: Rect detail::RotationWarper::warpRoi(Size src_size, const Mat &K, const Mat &R) + + :param src_size: Source image bounding box + + :param K: Camera intrinsic parameters + + :param R: Camera rotation matrix + + :return: Projected image minimum bounding box + +detail::ProjectorBase +--------------------- +.. ocv:struct:: detail::ProjectorBase + +Base class for warping logic implementation. :: + + struct CV_EXPORTS ProjectorBase + { + void setCameraParams(const Mat &K = Mat::eye(3, 3, CV_32F), + const Mat &R = Mat::eye(3, 3, CV_32F), + const Mat &T = Mat::zeros(3, 1, CV_32F)); + + float scale; + float k[9]; + float rinv[9]; + float r_kinv[9]; + float k_rinv[9]; + float t[3]; + }; + +detail::RotationWarperBase +-------------------------- +.. ocv:class:: detail::RotationWarperBase + +Base class for rotation-based warper using a `detail::ProjectorBase`_ derived class. :: + + template + class CV_EXPORTS RotationWarperBase : public RotationWarper + { + public: + Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R); + + Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap); + + Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, + Mat &dst); + + void warpBackward(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, + Size dst_size, Mat &dst); + + Rect warpRoi(Size src_size, const Mat &K, const Mat &R); + + protected: + + // Detects ROI of the destination image. It's correct for any projection. + virtual void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br); + + // Detects ROI of the destination image by walking over image border. + // Correctness for any projection isn't guaranteed. + void detectResultRoiByBorder(Size src_size, Point &dst_tl, Point &dst_br); + + P projector_; + }; + +detail::PlaneWarper +------------------- +.. ocv:class:: detail::PlaneWarper : public detail::RotationWarperBase + +Warper that maps an image onto the z = 1 plane. :: + + class CV_EXPORTS PlaneWarper : public RotationWarperBase + { + public: + PlaneWarper(float scale = 1.f) { projector_.scale = scale; } + + void setScale(float scale) { projector_.scale = scale; } + + Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R, const Mat &T); + + Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap); + + Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, + Mat &dst); + + Rect warpRoi(Size src_size, const Mat &K, const Mat &R, const Mat &T); + + protected: + void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br); + }; + +.. seealso:: :ocv:class:`detail::RotationWarper` + +detail::PlaneWarper::PlaneWarper +-------------------------------- + +Construct an instance of the plane warper class. + +.. ocv:function:: void detail::PlaneWarper::PlaneWarper(float scale = 1.f) + + :param scale: Projected image scale multiplier + +detail::SphericalWarper +----------------------- +.. ocv:class:: detail::SphericalWarper : public detail::RotationWarperBase + +Warper that maps an image onto the unit sphere located at the origin. :: + + class CV_EXPORTS SphericalWarper : public RotationWarperBase + { + public: + SphericalWarper(float scale) { projector_.scale = scale; } + + protected: + void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br); + }; + +.. seealso:: :ocv:class:`detail::RotationWarper` + +detail::SphericalWarper::SphericalWarper +---------------------------------------- + +Construct an instance of the spherical warper class. + +.. ocv:function:: void detail::SphericalWarper::SphericalWarper(float scale) + + :param scale: Projected image scale multiplier + +detail::CylindricalWarper +------------------------- +.. ocv:class:: detail::CylindricalWarper : public detail::RotationWarperBase + +Warper that maps an image onto the x*x + z*z = 1 cylinder. :: + + class CV_EXPORTS CylindricalWarper : public RotationWarperBase + { + public: + CylindricalWarper(float scale) { projector_.scale = scale; } + + protected: + void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br) + { + RotationWarperBase::detectResultRoiByBorder(src_size, dst_tl, dst_br); + } + }; + +.. seealso:: :ocv:class:`detail::RotationWarper` + +detail::CylindricalWarper::CylindricalWarper +-------------------------------------------- + +Construct an instance of the cylindrical warper class. + +.. ocv:function:: void detail::CylindricalWarper::CylindricalWarper(float scale) + + :param scale: Projected image scale multiplier diff --git a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp index 8906d88..ad4a645 100644 --- a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp @@ -134,9 +134,9 @@ public: Point2f warpPoint(const Point2f &pt, const Mat &K, const Mat &R, const Mat &T); - Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap); + virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap); - Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, + virtual Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, Mat &dst); Rect warpRoi(Size src_size, const Mat &K, const Mat &R, const Mat &T); @@ -503,6 +503,45 @@ protected: } }; +/////////////////////////////////////// OpenCL Accelerated Warpers ///////////////////////////////////// + +class CV_EXPORTS PlaneWarperOcl : public PlaneWarper +{ +public: + PlaneWarperOcl(float scale = 1.f) : PlaneWarper(scale) { } + + virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap) + { + return buildMaps(src_size, K, R, Mat::zeros(3, 1, CV_32FC1), xmap, ymap); + } + + virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Mat &dst) + { + return warp(src, K, R, Mat::zeros(3, 1, CV_32FC1), interp_mode, border_mode, dst); + } + + virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, const Mat &T, Mat &xmap, Mat &ymap); + virtual Point warp(const Mat &src, const Mat &K, const Mat &R, const Mat &T, int interp_mode, int border_mode, Mat &dst); +}; + +class CV_EXPORTS SphericalWarperOcl : public SphericalWarper +{ +public: + SphericalWarperOcl(float scale) : SphericalWarper(scale) { } + + virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap); + virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Mat &dst); +}; + +class CV_EXPORTS CylindricalWarperOcl : public CylindricalWarper +{ +public: + CylindricalWarperOcl(float scale) : CylindricalWarper(scale) { } + + virtual Rect buildMaps(Size src_size, const Mat &K, const Mat &R, Mat &xmap, Mat &ymap); + virtual Point warp(const Mat &src, const Mat &K, const Mat &R, int interp_mode, int border_mode, Mat &dst); +}; + } // namespace detail } // namespace cv diff --git a/modules/stitching/include/opencv2/stitching/warpers.hpp b/modules/stitching/include/opencv2/stitching/warpers.hpp index da5fe26..cdcb35c 100644 --- a/modules/stitching/include/opencv2/stitching/warpers.hpp +++ b/modules/stitching/include/opencv2/stitching/warpers.hpp @@ -167,6 +167,24 @@ public: }; #endif +class PlaneWarperOcl: public WarperCreator +{ +public: + Ptr create(float scale) const { return makePtr(scale); } +}; + +class SphericalWarperOcl: public WarperCreator +{ +public: + Ptr create(float scale) const { return makePtr(scale); } +}; + +class CylindricalWarperOcl: public WarperCreator +{ +public: + Ptr create(float scale) const { return makePtr(scale); } +}; + } // namespace cv #endif // __OPENCV_STITCHING_WARPER_CREATORS_HPP__ diff --git a/modules/stitching/src/opencl/warpers.cl b/modules/stitching/src/opencl/warpers.cl new file mode 100644 index 0000000..032ddf3 --- /dev/null +++ b/modules/stitching/src/opencl/warpers.cl @@ -0,0 +1,148 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Peng Xiao, pengxiao@multicorewareinc.com +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +__kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset, + __global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols, + __constant float * ck_rinv, __constant float * ct, + int tl_u, int tl_v, float scale) +{ + int du = get_global_id(0); + int dv = get_global_id(1); + + if (du < cols && dv < rows) + { + __global float * xmap = (__global float *)(xmapptr + mad24(dv, xmap_step, xmap_offset + du * (int)sizeof(float))); + __global float * ymap = (__global float *)(ymapptr + mad24(dv, ymap_step, ymap_offset + du * (int)sizeof(float))); + + float u = tl_u + du; + float v = tl_v + dv; + float x, y; + + float x_ = u / scale - ct[0]; + float y_ = v / scale - ct[1]; + + float z; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * (1 - ct[2]); + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * (1 - ct[2]); + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * (1 - ct[2]); + + x /= z; + y /= z; + + xmap[0] = x; + ymap[0] = y; + } +} + +__kernel void buildWarpCylindricalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset, + __global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols, + __constant float * ck_rinv, int tl_u, int tl_v, float scale) +{ + int du = get_global_id(0); + int dv = get_global_id(1); + + if (du < cols && dv < rows) + { + __global float * xmap = (__global float *)(xmapptr + mad24(dv, xmap_step, xmap_offset + du * (int)sizeof(float))); + __global float * ymap = (__global float *)(ymapptr + mad24(dv, ymap_step, ymap_offset + du * (int)sizeof(float))); + + float u = tl_u + du; + float v = tl_v + dv; + float x, y; + + u /= scale; + float x_ = sin(u); + float y_ = v / scale; + float z_ = cos(u); + + float z; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; + + if (z > 0) x /= z, y /= z; + else x = y = -1; + + xmap[0] = x; + ymap[0] = y; + } +} + +__kernel void buildWarpSphericalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset, + __global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols, + __constant float * ck_rinv, int tl_u, int tl_v, float scale) +{ + int du = get_global_id(0); + int dv = get_global_id(1); + + if (du < cols && dv < rows) + { + __global float * xmap = (__global float *)(xmapptr + mad24(dv, xmap_step, xmap_offset + du * (int)sizeof(float))); + __global float * ymap = (__global float *)(ymapptr + mad24(dv, ymap_step, ymap_offset + du * (int)sizeof(float))); + + float u = tl_u + du; + float v = tl_v + dv; + float x, y; + + v /= scale; + u /= scale; + + float sinv = sin(v); + float x_ = sinv * sin(u); + float y_ = -cos(v); + float z_ = sinv * cos(u); + + float z; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; + + if (z > 0) x /= z, y /= z; + else x = y = -1; + + xmap[0] = x; + ymap[0] = y; + } +} diff --git a/modules/stitching/src/precomp.hpp b/modules/stitching/src/precomp.hpp index 0444517..499202f 100644 --- a/modules/stitching/src/precomp.hpp +++ b/modules/stitching/src/precomp.hpp @@ -53,6 +53,7 @@ #include #include #include "opencv2/core.hpp" +#include "opencv2/core/ocl.hpp" #include "opencv2/core/utility.hpp" #include "opencv2/stitching.hpp" #include "opencv2/stitching/detail/autocalib.hpp" diff --git a/modules/stitching/src/warpers_ocl.cpp b/modules/stitching/src/warpers_ocl.cpp new file mode 100644 index 0000000..7735cbb --- /dev/null +++ b/modules/stitching/src/warpers_ocl.cpp @@ -0,0 +1,184 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "precomp.hpp" +#include "opencl_kernels.hpp" + +namespace cv { +namespace detail { + +/////////////////////////////////////////// PlaneWarperOcl //////////////////////////////////////////// + +Rect PlaneWarperOcl::buildMaps(Size src_size, const Mat & K, const Mat & R, const Mat & T, Mat & xmap, Mat & ymap) +{ + projector_.setCameraParams(K, R); + + Point dst_tl, dst_br; + detectResultRoi(src_size, dst_tl, dst_br); + + if (ocl::useOpenCL()) + { + ocl::Kernel k("buildWarpPlaneMaps", ocl::stitching::warpers_oclsrc); + if (!k.empty()) + { + xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1); + ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1); + + Mat r_kinv(1, 9, CV_32FC1, projector_.r_kinv), t(1, 3, CV_32FC1, projector_.t); + UMat uxmap = xmap.getUMat(ACCESS_WRITE), uymap = ymap.getUMat(ACCESS_WRITE), + ur_kinv = r_kinv.getUMat(ACCESS_READ), ut = t.getUMat(ACCESS_READ); + + k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), + ocl::KernelArg::PtrReadOnly(ur_kinv), ocl::KernelArg::PtrReadOnly(ut), + dst_tl.x, dst_tl.y, projector_.scale); + + size_t globalsize[2] = { xmap.cols, xmap.rows }; + if (k.run(2, globalsize, NULL, true)) + return Rect(dst_tl, dst_br); + } + } + + return PlaneWarper::buildMaps(src_size, K, R, T, xmap, ymap); +} + +Point PlaneWarperOcl::warp(const Mat & src, const Mat & K, const Mat & R, const Mat & T, int interp_mode, int border_mode, Mat & dst) +{ + Mat uxmap, uymap; + Rect dst_roi = buildMaps(src.size(), K, R, T, uxmap, uymap); + + dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); + UMat udst = dst.getUMat(ACCESS_WRITE); + remap(src, udst, uxmap, uymap, interp_mode, border_mode); + + return dst_roi.tl(); +} + +/////////////////////////////////////////// SphericalWarperOcl //////////////////////////////////////// + +Rect SphericalWarperOcl::buildMaps(Size src_size, const Mat & K, const Mat &R, Mat & xmap, Mat & ymap) +{ + projector_.setCameraParams(K, R); + + Point dst_tl, dst_br; + detectResultRoi(src_size, dst_tl, dst_br); + + if (ocl::useOpenCL()) + { + ocl::Kernel k("buildWarpSphericalMaps", ocl::stitching::warpers_oclsrc); + if (!k.empty()) + { + xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1); + ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1); + + Mat r_kinv(1, 9, CV_32FC1, projector_.r_kinv); + UMat uxmap = xmap.getUMat(ACCESS_WRITE), uymap = ymap.getUMat(ACCESS_WRITE), ur_kinv = r_kinv.getUMat(ACCESS_READ); + + k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), + ocl::KernelArg::PtrReadOnly(ur_kinv), dst_tl.x, dst_tl.y, projector_.scale); + + size_t globalsize[2] = { xmap.cols, xmap.rows }; + if (k.run(2, globalsize, NULL, true)) + return Rect(dst_tl, dst_br); + } + } + + return SphericalWarper::buildMaps(src_size, K, R, xmap, ymap); +} + +Point SphericalWarperOcl::warp(const Mat & src, const Mat & K, const Mat & R, int interp_mode, int border_mode, Mat & dst) +{ + Mat uxmap, uymap; + Rect dst_roi = buildMaps(src.size(), K, R, uxmap, uymap); + + dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); + UMat udst = dst.getUMat(ACCESS_WRITE); + remap(src, udst, uxmap, uymap, interp_mode, border_mode); + + return dst_roi.tl(); +} + +/////////////////////////////////////////// CylindricalWarperOcl //////////////////////////////////////// + +Rect CylindricalWarperOcl::buildMaps(Size src_size, const Mat & K, const Mat & R, Mat & xmap, Mat & ymap) +{ + projector_.setCameraParams(K, R); + + Point dst_tl, dst_br; + detectResultRoi(src_size, dst_tl, dst_br); + + if (ocl::useOpenCL()) + { + ocl::Kernel k("buildWarpCylindricalMaps", ocl::stitching::warpers_oclsrc); + if (!k.empty()) + { + xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1); + ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32FC1); + + Mat r_kinv(1, 9, CV_32FC1, projector_.r_kinv); + UMat uxmap = xmap.getUMat(ACCESS_WRITE), uymap = ymap.getUMat(ACCESS_WRITE), ur_kinv = r_kinv.getUMat(ACCESS_READ); + + k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), + ocl::KernelArg::PtrReadOnly(ur_kinv), dst_tl.x, dst_tl.y, projector_.scale); + + size_t globalsize[2] = { xmap.cols, xmap.rows }; + if (k.run(2, globalsize, NULL, true)) + return Rect(dst_tl, dst_br); + } + } + + return CylindricalWarper::buildMaps(src_size, K, R, xmap, ymap); +} + +Point CylindricalWarperOcl::warp(const Mat & src, const Mat & K, const Mat & R, int interp_mode, int border_mode, Mat & dst) +{ + Mat uxmap, uymap; + Rect dst_roi = buildMaps(src.size(), K, R, uxmap, uymap); + + dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); + UMat udst = dst.getUMat(ACCESS_WRITE); + remap(src, udst, uxmap, uymap, interp_mode, border_mode); + + return dst_roi.tl(); +} + +} // namespace detail +} // namespace cv diff --git a/modules/stitching/test/ocl/test_warpers.cpp b/modules/stitching/test/ocl/test_warpers.cpp new file mode 100644 index 0000000..83f0834 --- /dev/null +++ b/modules/stitching/test/ocl/test_warpers.cpp @@ -0,0 +1,149 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the OpenCV Foundation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "test_precomp.hpp" +#include "opencv2/ts/ocl_test.hpp" +#include "opencv2/stitching/warpers.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +///////////////////////// WarperTestBase /////////////////////////// + +struct WarperTestBase : + public Test, public TestUtils +{ + Mat src, dst, xmap, ymap; + Mat udst, uxmap, uymap; + Mat K, R; + + virtual void generateTestData() + { + Size size = randomSize(1, MAX_VALUE); + + src = randomMat(size, CV_32FC1, -500, 500); + + K = Mat::eye(3, 3, CV_32FC1); + R = Mat::eye(3, 3, CV_32FC1); + } + + void Near(double threshold = 0.) + { + EXPECT_MAT_NEAR(xmap, uxmap, threshold); + EXPECT_MAT_NEAR(ymap, uymap, threshold); + EXPECT_MAT_NEAR(dst, udst, threshold); + } +}; + +//////////////////////////////// SphericalWarperOcl ///////////////////////////////////////////////// + +typedef WarperTestBase SphericalWarperOclTest; + +OCL_TEST_F(SphericalWarperOclTest, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + Ptr creator = makePtr(); + Ptr warper = creator->create(2.0); + + OCL_OFF(warper->buildMaps(src.size(), K, R, xmap, ymap)); + OCL_ON(warper->buildMaps(src.size(), K, R, uxmap, uymap)); + + OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); + OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); + + Near(1e-5); + } +} + +//////////////////////////////// CylindricalWarperOcl ///////////////////////////////////////////////// + +typedef WarperTestBase CylindricalWarperOclTest; + +OCL_TEST_F(CylindricalWarperOclTest, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + Ptr creator = makePtr(); + Ptr warper = creator->create(2.0); + + OCL_OFF(warper->buildMaps(src.size(), K, R, xmap, ymap)); + OCL_ON(warper->buildMaps(src.size(), K, R, uxmap, uymap)); + + OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); + OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); + + Near(1e-5); + } +} + +//////////////////////////////// PlaneWarperOcl ///////////////////////////////////////////////// + +typedef WarperTestBase PlaneWarperOclTest; + +OCL_TEST_F(PlaneWarperOclTest, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + Ptr creator = makePtr(); + Ptr warper = creator->create(2.0); + + OCL_OFF(warper->buildMaps(src.size(), K, R, xmap, ymap)); + OCL_ON(warper->buildMaps(src.size(), K, R, uxmap, uymap)); + + OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); + OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); + + Near(1e-5); + } +} + +} } // namespace cvtest::ocl + +#endif // HAVE_OPENCL diff --git a/samples/cpp/stitching_detailed.cpp b/samples/cpp/stitching_detailed.cpp index b6eefc6..5eb3df4 100644 --- a/samples/cpp/stitching_detailed.cpp +++ b/samples/cpp/stitching_detailed.cpp @@ -71,8 +71,11 @@ static void printUsage() " --preview\n" " Run stitching in the preview mode. Works faster than usual mode,\n" " but output image will have lower resolution.\n" - " --try_gpu (yes|no)\n" - " Try to use GPU. The default value is 'no'. All default values\n" + " --try_cuda (yes|no)\n" + " Try to use CUDA. The default value is 'no'. All default values\n" + " are for CPU mode.\n" + " --try_ocl (yes|no)\n" + " Try to use OpenCL. The default value is 'no'. All default values\n" " are for CPU mode.\n" "\nMotion Estimation Flags:\n" " --work_megapix \n" @@ -123,7 +126,8 @@ static void printUsage() // Default command line args vector img_names; bool preview = false; -bool try_gpu = false; +bool try_cuda = false; +bool try_ocl = false; double work_megapix = 0.6; double seam_megapix = 0.1; double compose_megapix = -1; @@ -161,15 +165,28 @@ static int parseCmdArgs(int argc, char** argv) { preview = true; } - else if (string(argv[i]) == "--try_gpu") + else if (string(argv[i]) == "--try_cuda") { if (string(argv[i + 1]) == "no") - try_gpu = false; + try_cuda = false; else if (string(argv[i + 1]) == "yes") - try_gpu = true; + try_cuda = true; else { - cout << "Bad --try_gpu flag value\n"; + cout << "Bad --try_cuda flag value\n"; + return -1; + } + i++; + } + else if (string(argv[i]) == "--try_ocl") + { + if (string(argv[i + 1]) == "no") + try_ocl = false; + else if (string(argv[i + 1]) == "yes") + try_ocl = true; + else + { + cout << "Bad --try_ocl flag value\n"; return -1; } i++; @@ -357,7 +374,7 @@ int main(int argc, char* argv[]) if (features_type == "surf") { #ifdef HAVE_OPENCV_NONFREE - if (try_gpu && cuda::getCudaEnabledDeviceCount() > 0) + if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) finder = makePtr(); else #endif @@ -430,7 +447,7 @@ int main(int argc, char* argv[]) t = getTickCount(); #endif vector pairwise_matches; - BestOf2NearestMatcher matcher(try_gpu, match_conf); + BestOf2NearestMatcher matcher(try_cuda, match_conf); matcher(features, pairwise_matches); matcher.collectGarbage(); LOGLN("Pairwise matching, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); @@ -552,8 +569,17 @@ int main(int argc, char* argv[]) // Warp images and their masks Ptr warper_creator; + if (try_ocl) + { + if (warp_type == "plane") + warper_creator = makePtr(); + else if (warp_type == "cylindrical") + warper_creator = makePtr(); + else if (warp_type == "spherical") + warper_creator = makePtr(); + } #ifdef HAVE_OPENCV_CUDAWARPING - if (try_gpu && cuda::getCudaEnabledDeviceCount() > 0) + else if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) { if (warp_type == "plane") warper_creator = makePtr(); @@ -636,7 +662,7 @@ int main(int argc, char* argv[]) else if (seam_find_type == "gc_color") { #ifdef HAVE_OPENCV_CUDA - if (try_gpu && cuda::getCudaEnabledDeviceCount() > 0) + if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) seam_finder = makePtr(GraphCutSeamFinderBase::COST_COLOR); else #endif @@ -645,7 +671,7 @@ int main(int argc, char* argv[]) else if (seam_find_type == "gc_colorgrad") { #ifdef HAVE_OPENCV_CUDA - if (try_gpu && cuda::getCudaEnabledDeviceCount() > 0) + if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) seam_finder = makePtr(GraphCutSeamFinderBase::COST_COLOR_GRAD); else #endif @@ -755,11 +781,11 @@ int main(int argc, char* argv[]) if (!blender) { - blender = Blender::createDefault(blend_type, try_gpu); + blender = Blender::createDefault(blend_type, try_cuda); Size dst_sz = resultRoi(corners, sizes).size(); float blend_width = sqrt(static_cast(dst_sz.area())) * blend_strength / 100.f; if (blend_width < 1.f) - blender = Blender::createDefault(Blender::NO, try_gpu); + blender = Blender::createDefault(Blender::NO, try_cuda); else if (blend_type == Blender::MULTI_BAND) { MultiBandBlender* mb = dynamic_cast(blender.get()); -- 2.7.4