From 8b6c5ade0e527db9935fb7485f12be5c2295abf9 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Thu, 13 Feb 2014 18:10:37 +0400 Subject: [PATCH 01/15] TAPI: added perf test for stitching --- modules/stitching/perf/opencl/perf_stitch.cpp | 144 ++++++++++++++++++ modules/ts/include/opencv2/ts/ocl_test.hpp | 25 +++ 2 files changed, 169 insertions(+) create mode 100644 modules/stitching/perf/opencl/perf_stitch.cpp diff --git a/modules/stitching/perf/opencl/perf_stitch.cpp b/modules/stitching/perf/opencl/perf_stitch.cpp new file mode 100644 index 000000000..343472699 --- /dev/null +++ b/modules/stitching/perf/opencl/perf_stitch.cpp @@ -0,0 +1,144 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. +// +// Copyright (C) 2014, Itseez, Inc, all rights reserved. + +#include "perf_precomp.hpp" +#include "opencv2/ts/ocl_perf.hpp" + +using namespace cv; +using namespace perf; +using namespace cvtest::ocl; +using namespace std; +using namespace std::tr1; + +#define SURF_MATCH_CONFIDENCE 0.65f +#define ORB_MATCH_CONFIDENCE 0.3f +#define WORK_MEGAPIX 0.6 + +typedef TestBaseWithParam stitch; + +#ifdef HAVE_OPENCV_NONFREE_TODO_FIND_WHY_SURF_IS_NOT_ABLE_TO_STITCH_PANOS +#define TEST_DETECTORS testing::Values("surf", "orb") +#else +#define TEST_DETECTORS testing::Values("orb") +#endif + +OCL_PERF_TEST_P(stitch, a123, TEST_DETECTORS) +{ + UMat pano; + + vector _imgs; + _imgs.push_back( imread( getDataPath("stitching/a1.png") ) ); + _imgs.push_back( imread( getDataPath("stitching/a2.png") ) ); + _imgs.push_back( imread( getDataPath("stitching/a3.png") ) ); + vector imgs = ToUMat(_imgs); + + Ptr featuresFinder = GetParam() == "orb" + ? Ptr(new detail::OrbFeaturesFinder()) + : Ptr(new detail::SurfFeaturesFinder()); + + Ptr featuresMatcher = GetParam() == "orb" + ? makePtr(false, ORB_MATCH_CONFIDENCE) + : makePtr(false, SURF_MATCH_CONFIDENCE); + + declare.iterations(20); + + while(next()) + { + Stitcher stitcher = Stitcher::createDefault(); + stitcher.setFeaturesFinder(featuresFinder); + stitcher.setFeaturesMatcher(featuresMatcher); + stitcher.setWarper(makePtr()); + stitcher.setRegistrationResol(WORK_MEGAPIX); + + startTimer(); + stitcher.stitch(imgs, pano); + stopTimer(); + } + + EXPECT_NEAR(pano.size().width, 1182, 50); + EXPECT_NEAR(pano.size().height, 682, 30); + + SANITY_CHECK_NOTHING(); +} + +OCL_PERF_TEST_P(stitch, b12, TEST_DETECTORS) +{ + UMat pano; + + vector imgs; + imgs.push_back( imread( getDataPath("stitching/b1.png") ) ); + imgs.push_back( imread( getDataPath("stitching/b2.png") ) ); + + Ptr featuresFinder = GetParam() == "orb" + ? Ptr(new detail::OrbFeaturesFinder()) + : Ptr(new detail::SurfFeaturesFinder()); + + Ptr featuresMatcher = GetParam() == "orb" + ? makePtr(false, ORB_MATCH_CONFIDENCE) + : makePtr(false, SURF_MATCH_CONFIDENCE); + + declare.iterations(20); + + while(next()) + { + Stitcher stitcher = Stitcher::createDefault(); + stitcher.setFeaturesFinder(featuresFinder); + stitcher.setFeaturesMatcher(featuresMatcher); + stitcher.setWarper(makePtr()); + stitcher.setRegistrationResol(WORK_MEGAPIX); + + startTimer(); + stitcher.stitch(imgs, pano); + stopTimer(); + } + + EXPECT_NEAR(pano.size().width, 1124, 50); + EXPECT_NEAR(pano.size().height, 644, 30); + + SANITY_CHECK_NOTHING(); +} + +OCL_PERF_TEST_P(stitch, boat, TEST_DETECTORS) +{ + UMat pano; + + vector _imgs; + _imgs.push_back( imread( getDataPath("stitching/boat1.jpg") ) ); + _imgs.push_back( imread( getDataPath("stitching/boat2.jpg") ) ); + _imgs.push_back( imread( getDataPath("stitching/boat3.jpg") ) ); + _imgs.push_back( imread( getDataPath("stitching/boat4.jpg") ) ); + _imgs.push_back( imread( getDataPath("stitching/boat5.jpg") ) ); + _imgs.push_back( imread( getDataPath("stitching/boat6.jpg") ) ); + vector imgs = ToUMat(_imgs); + + Ptr featuresFinder = GetParam() == "orb" + ? Ptr(new detail::OrbFeaturesFinder()) + : Ptr(new detail::SurfFeaturesFinder()); + + Ptr featuresMatcher = GetParam() == "orb" + ? makePtr(false, ORB_MATCH_CONFIDENCE) + : makePtr(false, SURF_MATCH_CONFIDENCE); + + declare.iterations(20); + + while(next()) + { + Stitcher stitcher = Stitcher::createDefault(); + stitcher.setFeaturesFinder(featuresFinder); + stitcher.setFeaturesMatcher(featuresMatcher); + stitcher.setWarper(makePtr()); + stitcher.setRegistrationResol(WORK_MEGAPIX); + + startTimer(); + stitcher.stitch(imgs, pano); + stopTimer(); + } + + EXPECT_NEAR(pano.size().width, 10789, 200); + EXPECT_NEAR(pano.size().height, 2663, 100); + + SANITY_CHECK_NOTHING(); +} diff --git a/modules/ts/include/opencv2/ts/ocl_test.hpp b/modules/ts/include/opencv2/ts/ocl_test.hpp index 5dd25dba4..43c01b2da 100644 --- a/modules/ts/include/opencv2/ts/ocl_test.hpp +++ b/modules/ts/include/opencv2/ts/ocl_test.hpp @@ -57,6 +57,31 @@ namespace ocl { using namespace cv; using namespace testing; +inline std::vector ToUMat(const std::vector& src) +{ + std::vector dst; + dst.resize(src.size()); + for (size_t i = 0; i < src.size(); ++i) + { + src[i].copyTo(dst[i]); + } + return dst; +} + +inline UMat ToUMat(const Mat& src) +{ + UMat dst; + src.copyTo(dst); + return dst; +} + +inline UMat ToUMat(InputArray src) +{ + UMat dst; + src.getMat().copyTo(dst); + return dst; +} + extern int test_loop_times; #define MAX_VALUE 357 From 1d9808e5d5c6f53fe25077f77279404bab205755 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Fri, 14 Feb 2014 15:36:04 +0400 Subject: [PATCH 02/15] TAPI: stitching, replaced Mat->UMat/_Array --- modules/core/include/opencv2/core/mat.hpp | 3 + modules/core/src/matrix.cpp | 37 +++++ .../stitching/include/opencv2/stitching.hpp | 20 +-- .../opencv2/stitching/detail/blenders.hpp | 38 +++--- .../stitching/detail/exposure_compensate.hpp | 28 ++-- .../opencv2/stitching/detail/matchers.hpp | 16 +-- .../opencv2/stitching/detail/seam_finders.hpp | 30 ++-- .../include/opencv2/stitching/detail/util.hpp | 2 +- modules/stitching/src/blenders.cpp | 129 +++++++++++------- modules/stitching/src/exposure_compensate.cpp | 34 ++--- modules/stitching/src/matchers.cpp | 39 +++--- modules/stitching/src/seam_finders.cpp | 41 +++--- modules/stitching/src/stitcher.cpp | 52 +++---- modules/stitching/src/util.cpp | 2 +- samples/cpp/stitching_detailed.cpp | 10 +- 15 files changed, 280 insertions(+), 201 deletions(-) diff --git a/modules/core/include/opencv2/core/mat.hpp b/modules/core/include/opencv2/core/mat.hpp index 6b8368fd5..d921f7565 100644 --- a/modules/core/include/opencv2/core/mat.hpp +++ b/modules/core/include/opencv2/core/mat.hpp @@ -218,6 +218,9 @@ public: virtual void release() const; virtual void clear() const; virtual void setTo(const _InputArray& value, const _InputArray & mask = _InputArray()) const; + + void assign(const UMat& u) const; + void assign(const Mat& m) const; }; diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index 4efba4654..32ccd0377 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -2592,6 +2592,43 @@ void _OutputArray::setTo(const _InputArray& arr, const _InputArray & mask) const CV_Error(Error::StsNotImplemented, ""); } + +void _OutputArray::assign(const UMat& u) const +{ + int k = kind(); + if (k == UMAT) + { + *(UMat*)obj = u; + } + else if (k == MAT) + { + u.copyTo(*(Mat*)obj); // TODO check u.getMat() + } + else + { + CV_Error(Error::StsNotImplemented, ""); + } +} + + +void _OutputArray::assign(const Mat& m) const +{ + int k = kind(); + if (k == UMAT) + { + m.copyTo(*(UMat*)obj); // TODO check m.getUMat() + } + else if (k == MAT) + { + *(Mat*)obj = m; + } + else + { + CV_Error(Error::StsNotImplemented, ""); + } +} + + static _InputOutputArray _none; InputOutputArray noArray() { return _none; } diff --git a/modules/stitching/include/opencv2/stitching.hpp b/modules/stitching/include/opencv2/stitching.hpp index 2c48f2f2a..b647c0fe2 100644 --- a/modules/stitching/include/opencv2/stitching.hpp +++ b/modules/stitching/include/opencv2/stitching.hpp @@ -98,8 +98,8 @@ public: void setFeaturesMatcher(Ptr features_matcher) { features_matcher_ = features_matcher; } - const cv::Mat& matchingMask() const { return matching_mask_; } - void setMatchingMask(const cv::Mat &mask) + const cv::UMat& matchingMask() const { return matching_mask_; } + void setMatchingMask(const cv::UMat &mask) { CV_Assert(mask.type() == CV_8U && mask.cols == mask.rows); matching_mask_ = mask.clone(); @@ -127,14 +127,14 @@ public: const Ptr blender() const { return blender_; } void setBlender(Ptr b) { blender_ = b; } - Status estimateTransform(InputArray images); - Status estimateTransform(InputArray images, const std::vector > &rois); + Status estimateTransform(InputArrayOfArrays images); + Status estimateTransform(InputArrayOfArrays images, const std::vector > &rois); Status composePanorama(OutputArray pano); - Status composePanorama(InputArray images, OutputArray pano); + Status composePanorama(InputArrayOfArrays images, OutputArray pano); - Status stitch(InputArray images, OutputArray pano); - Status stitch(InputArray images, const std::vector > &rois, OutputArray pano); + Status stitch(InputArrayOfArrays images, OutputArray pano); + Status stitch(InputArrayOfArrays images, const std::vector > &rois, OutputArray pano); std::vector component() const { return indices_; } std::vector cameras() const { return cameras_; } @@ -152,7 +152,7 @@ private: double conf_thresh_; Ptr features_finder_; Ptr features_matcher_; - cv::Mat matching_mask_; + cv::UMat matching_mask_; Ptr bundle_adjuster_; bool do_wave_correct_; detail::WaveCorrectKind wave_correct_kind_; @@ -161,12 +161,12 @@ private: Ptr seam_finder_; Ptr blender_; - std::vector imgs_; + std::vector imgs_; std::vector > rois_; std::vector full_img_sizes_; std::vector features_; std::vector pairwise_matches_; - std::vector seam_est_imgs_; + std::vector seam_est_imgs_; std::vector indices_; std::vector cameras_; double work_scale_; diff --git a/modules/stitching/include/opencv2/stitching/detail/blenders.hpp b/modules/stitching/include/opencv2/stitching/detail/blenders.hpp index 8d47d3225..f91a0eae6 100644 --- a/modules/stitching/include/opencv2/stitching/detail/blenders.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/blenders.hpp @@ -60,11 +60,11 @@ public: void prepare(const std::vector &corners, const std::vector &sizes); virtual void prepare(Rect dst_roi); - virtual void feed(const Mat &img, const Mat &mask, Point tl); - virtual void blend(Mat &dst, Mat &dst_mask); + virtual void feed(InputArray img, InputArray mask, Point tl); + virtual void blend(InputOutputArray dst, InputOutputArray dst_mask); protected: - Mat dst_, dst_mask_; + UMat dst_, dst_mask_; Rect dst_roi_; }; @@ -78,18 +78,18 @@ public: void setSharpness(float val) { sharpness_ = val; } void prepare(Rect dst_roi); - void feed(const Mat &img, const Mat &mask, Point tl); - void blend(Mat &dst, Mat &dst_mask); + void feed(InputArray img, InputArray mask, Point tl); + void blend(InputOutputArray dst, InputOutputArray dst_mask); // Creates weight maps for fixed set of source images by their masks and top-left corners. // Final image can be obtained by simple weighting of the source images. - Rect createWeightMaps(const std::vector &masks, const std::vector &corners, - std::vector &weight_maps); + Rect createWeightMaps(const std::vector &masks, const std::vector &corners, + std::vector &weight_maps); private: float sharpness_; - Mat weight_map_; - Mat dst_weight_map_; + UMat weight_map_; + UMat dst_weight_map_; }; inline FeatherBlender::FeatherBlender(float _sharpness) { setSharpness(_sharpness); } @@ -104,13 +104,13 @@ public: void setNumBands(int val) { actual_num_bands_ = val; } void prepare(Rect dst_roi); - void feed(const Mat &img, const Mat &mask, Point tl); - void blend(Mat &dst, Mat &dst_mask); + void feed(InputArray img, InputArray mask, Point tl); + void blend(InputOutputArray dst, InputOutputArray dst_mask); private: int actual_num_bands_, num_bands_; - std::vector dst_pyr_laplace_; - std::vector dst_band_weights_; + std::vector dst_pyr_laplace_; + std::vector dst_band_weights_; Rect dst_roi_final_; bool can_use_gpu_; int weight_type_; //CV_32F or CV_16S @@ -120,16 +120,16 @@ private: ////////////////////////////////////////////////////////////////////////////// // Auxiliary functions -void CV_EXPORTS normalizeUsingWeightMap(const Mat& weight, Mat& src); +void CV_EXPORTS normalizeUsingWeightMap(InputArray weight, InputOutputArray src); -void CV_EXPORTS createWeightMap(const Mat& mask, float sharpness, Mat& weight); +void CV_EXPORTS createWeightMap(InputArray mask, float sharpness, InputOutputArray weight); -void CV_EXPORTS createLaplacePyr(const Mat &img, int num_levels, std::vector& pyr); -void CV_EXPORTS createLaplacePyrGpu(const Mat &img, int num_levels, std::vector& pyr); +void CV_EXPORTS createLaplacePyr(InputArray img, int num_levels, std::vector& pyr); +void CV_EXPORTS createLaplacePyrGpu(InputArray img, int num_levels, std::vector& pyr); // Restores source image -void CV_EXPORTS restoreImageFromLaplacePyr(std::vector& pyr); -void CV_EXPORTS restoreImageFromLaplacePyrGpu(std::vector& pyr); +void CV_EXPORTS restoreImageFromLaplacePyr(std::vector& pyr); +void CV_EXPORTS restoreImageFromLaplacePyrGpu(std::vector& pyr); } // namespace detail } // namespace cv diff --git a/modules/stitching/include/opencv2/stitching/detail/exposure_compensate.hpp b/modules/stitching/include/opencv2/stitching/detail/exposure_compensate.hpp index 84a8ce4fe..5626b06f2 100644 --- a/modules/stitching/include/opencv2/stitching/detail/exposure_compensate.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/exposure_compensate.hpp @@ -56,29 +56,29 @@ public: enum { NO, GAIN, GAIN_BLOCKS }; static Ptr createDefault(int type); - void feed(const std::vector &corners, const std::vector &images, - const std::vector &masks); - virtual void feed(const std::vector &corners, const std::vector &images, - const std::vector > &masks) = 0; - virtual void apply(int index, Point corner, Mat &image, const Mat &mask) = 0; + void feed(const std::vector &corners, const std::vector &images, + const std::vector &masks); + virtual void feed(const std::vector &corners, const std::vector &images, + const std::vector > &masks) = 0; + virtual void apply(int index, Point corner, InputOutputArray image, InputArray mask) = 0; }; class CV_EXPORTS NoExposureCompensator : public ExposureCompensator { public: - void feed(const std::vector &/*corners*/, const std::vector &/*images*/, - const std::vector > &/*masks*/) { } - void apply(int /*index*/, Point /*corner*/, Mat &/*image*/, const Mat &/*mask*/) { } + void feed(const std::vector &/*corners*/, const std::vector &/*images*/, + const std::vector > &/*masks*/) { } + void apply(int /*index*/, Point /*corner*/, InputOutputArray /*image*/, InputArray /*mask*/) { } }; class CV_EXPORTS GainCompensator : public ExposureCompensator { public: - void feed(const std::vector &corners, const std::vector &images, - const std::vector > &masks); - void apply(int index, Point corner, Mat &image, const Mat &mask); + void feed(const std::vector &corners, const std::vector &images, + const std::vector > &masks); + void apply(int index, Point corner, InputOutputArray image, InputArray mask); std::vector gains() const; private: @@ -91,9 +91,9 @@ class CV_EXPORTS BlocksGainCompensator : public ExposureCompensator public: BlocksGainCompensator(int bl_width = 32, int bl_height = 32) : bl_width_(bl_width), bl_height_(bl_height) {} - void feed(const std::vector &corners, const std::vector &images, - const std::vector > &masks); - void apply(int index, Point corner, Mat &image, const Mat &mask); + void feed(const std::vector &corners, const std::vector &images, + const std::vector > &masks); + void apply(int index, Point corner, InputOutputArray image, InputArray mask); private: int bl_width_, bl_height_; diff --git a/modules/stitching/include/opencv2/stitching/detail/matchers.hpp b/modules/stitching/include/opencv2/stitching/detail/matchers.hpp index 02f86608c..8c3284a46 100644 --- a/modules/stitching/include/opencv2/stitching/detail/matchers.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/matchers.hpp @@ -60,7 +60,7 @@ struct CV_EXPORTS ImageFeatures int img_idx; Size img_size; std::vector keypoints; - Mat descriptors; + UMat descriptors; }; @@ -68,12 +68,12 @@ class CV_EXPORTS FeaturesFinder { public: virtual ~FeaturesFinder() {} - void operator ()(const Mat &image, ImageFeatures &features); - void operator ()(const Mat &image, ImageFeatures &features, const std::vector &rois); + void operator ()(InputArray image, ImageFeatures &features); + void operator ()(InputArray image, ImageFeatures &features, const std::vector &rois); virtual void collectGarbage() {} protected: - virtual void find(const Mat &image, ImageFeatures &features) = 0; + virtual void find(InputArray image, ImageFeatures &features) = 0; }; @@ -84,7 +84,7 @@ public: int num_octaves_descr = /*4*/3, int num_layers_descr = /*2*/4); private: - void find(const Mat &image, ImageFeatures &features); + void find(InputArray image, ImageFeatures &features); Ptr detector_; Ptr extractor_; @@ -97,7 +97,7 @@ public: OrbFeaturesFinder(Size _grid_size = Size(3,1), int nfeatures=1500, float scaleFactor=1.3f, int nlevels=5); private: - void find(const Mat &image, ImageFeatures &features); + void find(InputArray image, ImageFeatures &features); Ptr orb; Size grid_size; @@ -114,7 +114,7 @@ public: void collectGarbage(); private: - void find(const Mat &image, ImageFeatures &features); + void find(InputArray image, ImageFeatures &features); cuda::GpuMat image_; cuda::GpuMat gray_image_; @@ -151,7 +151,7 @@ public: MatchesInfo& matches_info) { match(features1, features2, matches_info); } void operator ()(const std::vector &features, std::vector &pairwise_matches, - const cv::Mat &mask = cv::Mat()); + const cv::UMat &mask = cv::UMat()); bool isThreadSafe() const { return is_thread_safe_; } diff --git a/modules/stitching/include/opencv2/stitching/detail/seam_finders.hpp b/modules/stitching/include/opencv2/stitching/detail/seam_finders.hpp index 24b7db6b5..5f085c1a4 100644 --- a/modules/stitching/include/opencv2/stitching/detail/seam_finders.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/seam_finders.hpp @@ -54,32 +54,32 @@ class CV_EXPORTS SeamFinder { public: virtual ~SeamFinder() {} - virtual void find(const std::vector &src, const std::vector &corners, - std::vector &masks) = 0; + virtual void find(const std::vector &src, const std::vector &corners, + std::vector &masks) = 0; }; class CV_EXPORTS NoSeamFinder : public SeamFinder { public: - void find(const std::vector&, const std::vector&, std::vector&) {} + void find(const std::vector&, const std::vector&, std::vector&) {} }; class CV_EXPORTS PairwiseSeamFinder : public SeamFinder { public: - virtual void find(const std::vector &src, const std::vector &corners, - std::vector &masks); + virtual void find(const std::vector &src, const std::vector &corners, + std::vector &masks); protected: void run(); virtual void findInPair(size_t first, size_t second, Rect roi) = 0; - std::vector images_; + std::vector images_; std::vector sizes_; std::vector corners_; - std::vector masks_; + std::vector masks_; }; @@ -87,7 +87,7 @@ class CV_EXPORTS VoronoiSeamFinder : public PairwiseSeamFinder { public: virtual void find(const std::vector &size, const std::vector &corners, - std::vector &masks); + std::vector &masks); private: void findInPair(size_t first, size_t second, Rect roi); }; @@ -103,8 +103,8 @@ public: CostFunction costFunction() const { return costFunc_; } void setCostFunction(CostFunction val) { costFunc_ = val; } - virtual void find(const std::vector &src, const std::vector &corners, - std::vector &masks); + virtual void find(const std::vector &src, const std::vector &corners, + std::vector &masks); private: enum ComponentState @@ -154,7 +154,7 @@ private: }; void process( - const Mat &image1, const Mat &image2, Point tl1, Point tl2, Mat &mask1, Mat &mask2); + const Mat &image1, const Mat &image2, Point tl1, Point tl2, Mat &mask1, Mat &mask2); void findComponents(); @@ -217,8 +217,8 @@ public: ~GraphCutSeamFinder(); - void find(const std::vector &src, const std::vector &corners, - std::vector &masks); + void find(const std::vector &src, const std::vector &corners, + std::vector &masks); private: // To avoid GCGraph dependency @@ -236,8 +236,8 @@ public: : cost_type_(cost_type), terminal_cost_(terminal_cost), bad_region_penalty_(bad_region_penalty) {} - void find(const std::vector &src, const std::vector &corners, - std::vector &masks); + void find(const std::vector &src, const std::vector &corners, + std::vector &masks); void findInPair(size_t first, size_t second, Rect roi); private: diff --git a/modules/stitching/include/opencv2/stitching/detail/util.hpp b/modules/stitching/include/opencv2/stitching/detail/util.hpp index 6f334b9ad..561880cc2 100644 --- a/modules/stitching/include/opencv2/stitching/detail/util.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/util.hpp @@ -145,7 +145,7 @@ private: // Auxiliary functions CV_EXPORTS bool overlapRoi(Point tl1, Point tl2, Size sz1, Size sz2, Rect &roi); -CV_EXPORTS Rect resultRoi(const std::vector &corners, const std::vector &images); +CV_EXPORTS Rect resultRoi(const std::vector &corners, const std::vector &images); CV_EXPORTS Rect resultRoi(const std::vector &corners, const std::vector &sizes); CV_EXPORTS Point resultTl(const std::vector &corners); diff --git a/modules/stitching/src/blenders.cpp b/modules/stitching/src/blenders.cpp index 446bfc131..5012d9ba4 100644 --- a/modules/stitching/src/blenders.cpp +++ b/modules/stitching/src/blenders.cpp @@ -76,8 +76,13 @@ void Blender::prepare(Rect dst_roi) } -void Blender::feed(const Mat &img, const Mat &mask, Point tl) +void Blender::feed(InputArray _img, InputArray _mask, Point tl) { + Mat img = _img.getMat(); + Mat mask = _mask.getMat(); + Mat dst = dst_.getMat(ACCESS_RW); + Mat dst_mask = dst_mask_.getMat(ACCESS_RW); + CV_Assert(img.type() == CV_16SC3); CV_Assert(mask.type() == CV_8U); int dx = tl.x - dst_roi_.x; @@ -86,9 +91,9 @@ void Blender::feed(const Mat &img, const Mat &mask, Point tl) for (int y = 0; y < img.rows; ++y) { const Point3_ *src_row = img.ptr >(y); - Point3_ *dst_row = dst_.ptr >(dy + y); + Point3_ *dst_row = dst.ptr >(dy + y); const uchar *mask_row = mask.ptr(y); - uchar *dst_mask_row = dst_mask_.ptr(dy + y); + uchar *dst_mask_row = dst_mask.ptr(dy + y); for (int x = 0; x < img.cols; ++x) { @@ -100,11 +105,11 @@ void Blender::feed(const Mat &img, const Mat &mask, Point tl) } -void Blender::blend(Mat &dst, Mat &dst_mask) +void Blender::blend(InputOutputArray dst, InputOutputArray dst_mask) { - dst_.setTo(Scalar::all(0), dst_mask_ == 0); - dst = dst_; - dst_mask = dst_mask_; + dst_.setTo(Scalar::all(0), dst_mask_.getMat(ACCESS_READ) == 0); // TODO + dst.assign(dst_); + dst_mask.assign(dst_mask_); dst_.release(); dst_mask_.release(); } @@ -118,21 +123,27 @@ void FeatherBlender::prepare(Rect dst_roi) } -void FeatherBlender::feed(const Mat &img, const Mat &mask, Point tl) +void FeatherBlender::feed(InputArray _img, InputArray mask, Point tl) { + Mat img = _img.getMat(); + Mat dst = dst_.getMat(ACCESS_RW); + CV_Assert(img.type() == CV_16SC3); CV_Assert(mask.type() == CV_8U); createWeightMap(mask, sharpness_, weight_map_); + Mat weight_map = weight_map_.getMat(ACCESS_READ); + Mat dst_weight_map = dst_weight_map_.getMat(ACCESS_RW); + int dx = tl.x - dst_roi_.x; int dy = tl.y - dst_roi_.y; for (int y = 0; y < img.rows; ++y) { const Point3_* src_row = img.ptr >(y); - Point3_* dst_row = dst_.ptr >(dy + y); - const float* weight_row = weight_map_.ptr(y); - float* dst_weight_row = dst_weight_map_.ptr(dy + y); + Point3_* dst_row = dst.ptr >(dy + y); + const float* weight_row = weight_map.ptr(y); + float* dst_weight_row = dst_weight_map.ptr(dy + y); for (int x = 0; x < img.cols; ++x) { @@ -145,16 +156,16 @@ void FeatherBlender::feed(const Mat &img, const Mat &mask, Point tl) } -void FeatherBlender::blend(Mat &dst, Mat &dst_mask) +void FeatherBlender::blend(InputOutputArray dst, InputOutputArray dst_mask) { normalizeUsingWeightMap(dst_weight_map_, dst_); - dst_mask_ = dst_weight_map_ > WEIGHT_EPS; + dst_mask_ = ((Mat)(dst_weight_map_.getMat(ACCESS_READ) > WEIGHT_EPS)).getUMat(ACCESS_READ); Blender::blend(dst, dst_mask); } -Rect FeatherBlender::createWeightMaps(const std::vector &masks, const std::vector &corners, - std::vector &weight_maps) +Rect FeatherBlender::createWeightMaps(const std::vector &masks, const std::vector &corners, + std::vector &weight_maps) { weight_maps.resize(masks.size()); for (size_t i = 0; i < masks.size(); ++i) @@ -168,7 +179,7 @@ Rect FeatherBlender::createWeightMaps(const std::vector &masks, const std:: { Rect roi(corners[i].x - dst_roi.x, corners[i].y - dst_roi.y, weight_maps[i].cols, weight_maps[i].rows); - weights_sum(roi) += weight_maps[i]; + add(weights_sum(roi), weight_maps[i], weights_sum(roi)); } for (size_t i = 0; i < weight_maps.size(); ++i) @@ -233,8 +244,9 @@ void MultiBandBlender::prepare(Rect dst_roi) } -void MultiBandBlender::feed(const Mat &img, const Mat &mask, Point tl) +void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) { + Mat img = _img.getMat(); CV_Assert(img.type() == CV_16SC3 || img.type() == CV_8UC3); CV_Assert(mask.type() == CV_8U); @@ -269,27 +281,27 @@ void MultiBandBlender::feed(const Mat &img, const Mat &mask, Point tl) int right = br_new.x - tl.x - img.cols; // Create the source image Laplacian pyramid - Mat img_with_border; - copyMakeBorder(img, img_with_border, top, bottom, left, right, + UMat img_with_border; + copyMakeBorder(_img, img_with_border, top, bottom, left, right, BORDER_REFLECT); - std::vector src_pyr_laplace; + std::vector src_pyr_laplace; if (can_use_gpu_ && img_with_border.depth() == CV_16S) createLaplacePyrGpu(img_with_border, num_bands_, src_pyr_laplace); else createLaplacePyr(img_with_border, num_bands_, src_pyr_laplace); // Create the weight map Gaussian pyramid - Mat weight_map; - std::vector weight_pyr_gauss(num_bands_ + 1); + UMat weight_map; + std::vector weight_pyr_gauss(num_bands_ + 1); if(weight_type_ == CV_32F) { - mask.convertTo(weight_map, CV_32F, 1./255.); + mask.getUMat().convertTo(weight_map, CV_32F, 1./255.); } - else// weight_type_ == CV_16S + else // weight_type_ == CV_16S { - mask.convertTo(weight_map, CV_16S); - add(weight_map, 1, weight_map, mask != 0); + mask.getUMat().convertTo(weight_map, CV_16S); + add(weight_map, Scalar::all(1), weight_map, mask.getMat(ACCESS_READ) != 0); // TODO } copyMakeBorder(weight_map, weight_pyr_gauss[0], top, bottom, left, right, BORDER_CONSTANT); @@ -307,13 +319,17 @@ void MultiBandBlender::feed(const Mat &img, const Mat &mask, Point tl) { for (int i = 0; i <= num_bands_; ++i) { + Mat _src_pyr_laplace = src_pyr_laplace[i].getMat(ACCESS_READ); + Mat _dst_pyr_laplace = dst_pyr_laplace_[i].getMat(ACCESS_RW); + Mat _weight_pyr_gauss = weight_pyr_gauss[i].getMat(ACCESS_READ); + Mat _dst_band_weights = dst_band_weights_[i].getMat(ACCESS_RW); for (int y = y_tl; y < y_br; ++y) { int y_ = y - y_tl; - const Point3_* src_row = src_pyr_laplace[i].ptr >(y_); - Point3_* dst_row = dst_pyr_laplace_[i].ptr >(y); - const float* weight_row = weight_pyr_gauss[i].ptr(y_); - float* dst_weight_row = dst_band_weights_[i].ptr(y); + const Point3_* src_row = _src_pyr_laplace.ptr >(y_); + Point3_* dst_row = _dst_pyr_laplace.ptr >(y); + const float* weight_row = _weight_pyr_gauss.ptr(y_); + float* dst_weight_row = _dst_band_weights.ptr(y); for (int x = x_tl; x < x_br; ++x) { @@ -332,13 +348,17 @@ void MultiBandBlender::feed(const Mat &img, const Mat &mask, Point tl) { for (int i = 0; i <= num_bands_; ++i) { + Mat _src_pyr_laplace = src_pyr_laplace[i].getMat(ACCESS_READ); + Mat _dst_pyr_laplace = dst_pyr_laplace_[i].getMat(ACCESS_RW); + Mat _weight_pyr_gauss = weight_pyr_gauss[i].getMat(ACCESS_READ); + Mat _dst_band_weights = dst_band_weights_[i].getMat(ACCESS_RW); for (int y = y_tl; y < y_br; ++y) { int y_ = y - y_tl; - const Point3_* src_row = src_pyr_laplace[i].ptr >(y_); - Point3_* dst_row = dst_pyr_laplace_[i].ptr >(y); - const short* weight_row = weight_pyr_gauss[i].ptr(y_); - short* dst_weight_row = dst_band_weights_[i].ptr(y); + const Point3_* src_row = _src_pyr_laplace.ptr >(y_); + Point3_* dst_row = _dst_pyr_laplace.ptr >(y); + const short* weight_row = _weight_pyr_gauss.ptr(y_); + short* dst_weight_row = _dst_band_weights.ptr(y); for (int x = x_tl; x < x_br; ++x) { @@ -356,7 +376,7 @@ void MultiBandBlender::feed(const Mat &img, const Mat &mask, Point tl) } -void MultiBandBlender::blend(Mat &dst, Mat &dst_mask) +void MultiBandBlender::blend(InputOutputArray dst, InputOutputArray dst_mask) { for (int i = 0; i <= num_bands_; ++i) normalizeUsingWeightMap(dst_band_weights_[i], dst_pyr_laplace_[i]); @@ -368,7 +388,7 @@ void MultiBandBlender::blend(Mat &dst, Mat &dst_mask) dst_ = dst_pyr_laplace_[0]; dst_ = dst_(Range(0, dst_roi_final_.height), Range(0, dst_roi_final_.width)); - dst_mask_ = dst_band_weights_[0] > WEIGHT_EPS; + dst_mask_ = ((Mat)(dst_band_weights_[0].getMat(ACCESS_READ) > WEIGHT_EPS)).getUMat(ACCESS_READ); dst_mask_ = dst_mask_(Range(0, dst_roi_final_.height), Range(0, dst_roi_final_.width)); dst_pyr_laplace_.clear(); dst_band_weights_.clear(); @@ -380,12 +400,15 @@ void MultiBandBlender::blend(Mat &dst, Mat &dst_mask) ////////////////////////////////////////////////////////////////////////////// // Auxiliary functions -void normalizeUsingWeightMap(const Mat& weight, Mat& src) +void normalizeUsingWeightMap(InputArray _weight, InputOutputArray _src) { #ifdef HAVE_TEGRA_OPTIMIZATION if(tegra::normalizeUsingWeightMap(weight, src)) return; #endif + Mat weight = _weight.getMat(); + Mat src = _src.getMat(); + CV_Assert(src.type() == CV_16SC3); if(weight.type() == CV_32FC1) @@ -424,15 +447,17 @@ void normalizeUsingWeightMap(const Mat& weight, Mat& src) } -void createWeightMap(const Mat &mask, float sharpness, Mat &weight) +void createWeightMap(InputArray mask, float sharpness, InputOutputArray weight) { CV_Assert(mask.type() == CV_8U); distanceTransform(mask, weight, DIST_L1, 3); - threshold(weight * sharpness, weight, 1.f, 1.f, THRESH_TRUNC); + UMat tmp; + multiply(weight, sharpness, tmp); + threshold(tmp, weight, 1.f, 1.f, THRESH_TRUNC); } -void createLaplacePyr(const Mat &img, int num_levels, std::vector &pyr) +void createLaplacePyr(InputArray img, int num_levels, std::vector &pyr) { #ifdef HAVE_TEGRA_OPTIMIZATION if(tegra::createLaplacePyr(img, num_levels, pyr)) @@ -445,18 +470,18 @@ void createLaplacePyr(const Mat &img, int num_levels, std::vector &pyr) { if(num_levels == 0) { - img.convertTo(pyr[0], CV_16S); + img.getUMat().convertTo(pyr[0], CV_16S); return; } - Mat downNext; - Mat current = img; + UMat downNext; + UMat current = img.getUMat(); pyrDown(img, downNext); for(int i = 1; i < num_levels; ++i) { - Mat lvl_up; - Mat lvl_down; + UMat lvl_up; + UMat lvl_down; pyrDown(downNext, lvl_down); pyrUp(downNext, lvl_up, current.size()); @@ -467,7 +492,7 @@ void createLaplacePyr(const Mat &img, int num_levels, std::vector &pyr) } { - Mat lvl_up; + UMat lvl_up; pyrUp(downNext, lvl_up, current.size()); subtract(current, lvl_up, pyr[num_levels-1], noArray(), CV_16S); @@ -476,10 +501,10 @@ void createLaplacePyr(const Mat &img, int num_levels, std::vector &pyr) } else { - pyr[0] = img; + pyr[0] = img.getUMat(); for (int i = 0; i < num_levels; ++i) pyrDown(pyr[i], pyr[i + 1]); - Mat tmp; + UMat tmp; for (int i = 0; i < num_levels; ++i) { pyrUp(pyr[i + 1], tmp, pyr[i].size()); @@ -489,7 +514,7 @@ void createLaplacePyr(const Mat &img, int num_levels, std::vector &pyr) } -void createLaplacePyrGpu(const Mat &img, int num_levels, std::vector &pyr) +void createLaplacePyrGpu(InputArray img, int num_levels, std::vector &pyr) { #if defined(HAVE_OPENCV_CUDAARITHM) && defined(HAVE_OPENCV_CUDAWARPING) pyr.resize(num_levels + 1); @@ -517,11 +542,11 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, std::vector &pyr) } -void restoreImageFromLaplacePyr(std::vector &pyr) +void restoreImageFromLaplacePyr(std::vector &pyr) { if (pyr.empty()) return; - Mat tmp; + UMat tmp; for (size_t i = pyr.size() - 1; i > 0; --i) { pyrUp(pyr[i], tmp, pyr[i - 1].size()); @@ -530,7 +555,7 @@ void restoreImageFromLaplacePyr(std::vector &pyr) } -void restoreImageFromLaplacePyrGpu(std::vector &pyr) +void restoreImageFromLaplacePyrGpu(std::vector &pyr) { #if defined(HAVE_OPENCV_CUDAARITHM) && defined(HAVE_OPENCV_CUDAWARPING) if (pyr.empty()) diff --git a/modules/stitching/src/exposure_compensate.cpp b/modules/stitching/src/exposure_compensate.cpp index 78ce6d371..32112d69f 100644 --- a/modules/stitching/src/exposure_compensate.cpp +++ b/modules/stitching/src/exposure_compensate.cpp @@ -58,18 +58,18 @@ Ptr ExposureCompensator::createDefault(int type) } -void ExposureCompensator::feed(const std::vector &corners, const std::vector &images, - const std::vector &masks) +void ExposureCompensator::feed(const std::vector &corners, const std::vector &images, + const std::vector &masks) { - std::vector > level_masks; + std::vector > level_masks; for (size_t i = 0; i < masks.size(); ++i) level_masks.push_back(std::make_pair(masks[i], 255)); feed(corners, images, level_masks); } -void GainCompensator::feed(const std::vector &corners, const std::vector &images, - const std::vector > &masks) +void GainCompensator::feed(const std::vector &corners, const std::vector &images, + const std::vector > &masks) { LOGLN("Exposure compensation..."); #if ENABLE_LOG @@ -93,11 +93,11 @@ void GainCompensator::feed(const std::vector &corners, const std::vector< Rect roi; if (overlapRoi(corners[i], corners[j], images[i].size(), images[j].size(), roi)) { - subimg1 = images[i](Rect(roi.tl() - corners[i], roi.br() - corners[i])); - subimg2 = images[j](Rect(roi.tl() - corners[j], roi.br() - corners[j])); + subimg1 = images[i](Rect(roi.tl() - corners[i], roi.br() - corners[i])).getMat(ACCESS_READ); + subimg2 = images[j](Rect(roi.tl() - corners[j], roi.br() - corners[j])).getMat(ACCESS_READ); - submask1 = masks[i].first(Rect(roi.tl() - corners[i], roi.br() - corners[i])); - submask2 = masks[j].first(Rect(roi.tl() - corners[j], roi.br() - corners[j])); + submask1 = masks[i].first(Rect(roi.tl() - corners[i], roi.br() - corners[i])).getMat(ACCESS_READ); + submask2 = masks[j].first(Rect(roi.tl() - corners[j], roi.br() - corners[j])).getMat(ACCESS_READ); intersect = (submask1 == masks[i].second) & (submask2 == masks[j].second); N(i, j) = N(j, i) = std::max(1, countNonZero(intersect)); @@ -145,9 +145,9 @@ void GainCompensator::feed(const std::vector &corners, const std::vector< } -void GainCompensator::apply(int index, Point /*corner*/, Mat &image, const Mat &/*mask*/) +void GainCompensator::apply(int index, Point /*corner*/, InputOutputArray image, InputArray /*mask*/) { - image *= gains_(index, 0); + multiply(image, gains_(index, 0), image); } @@ -160,8 +160,8 @@ std::vector GainCompensator::gains() const } -void BlocksGainCompensator::feed(const std::vector &corners, const std::vector &images, - const std::vector > &masks) +void BlocksGainCompensator::feed(const std::vector &corners, const std::vector &images, + const std::vector > &masks) { CV_Assert(corners.size() == images.size() && images.size() == masks.size()); @@ -169,8 +169,8 @@ void BlocksGainCompensator::feed(const std::vector &corners, const std::v std::vector bl_per_imgs(num_images); std::vector block_corners; - std::vector block_images; - std::vector > block_masks; + std::vector block_images; + std::vector > block_masks; // Construct blocks for gain compensator for (int img_idx = 0; img_idx < num_images; ++img_idx) @@ -220,8 +220,10 @@ void BlocksGainCompensator::feed(const std::vector &corners, const std::v } -void BlocksGainCompensator::apply(int index, Point /*corner*/, Mat &image, const Mat &/*mask*/) +void BlocksGainCompensator::apply(int index, Point /*corner*/, InputOutputArray _image, InputArray /*mask*/) { + Mat image = _image.getMat(); + CV_Assert(image.type() == CV_8UC3); Mat_ gain_map; diff --git a/modules/stitching/src/matchers.cpp b/modules/stitching/src/matchers.cpp index f463518b6..1ce1fc3b2 100644 --- a/modules/stitching/src/matchers.cpp +++ b/modules/stitching/src/matchers.cpp @@ -264,14 +264,14 @@ void GpuMatcher::collectGarbage() namespace cv { namespace detail { -void FeaturesFinder::operator ()(const Mat &image, ImageFeatures &features) +void FeaturesFinder::operator ()(InputArray image, ImageFeatures &features) { find(image, features); features.img_size = image.size(); } -void FeaturesFinder::operator ()(const Mat &image, ImageFeatures &features, const std::vector &rois) +void FeaturesFinder::operator ()(InputArray image, ImageFeatures &features, const std::vector &rois) { std::vector roi_features(rois.size()); size_t total_kps_count = 0; @@ -279,7 +279,7 @@ void FeaturesFinder::operator ()(const Mat &image, ImageFeatures &features, cons for (size_t i = 0; i < rois.size(); ++i) { - find(image(rois[i]), roi_features[i]); + find(image.getUMat()(rois[i]), roi_features[i]); total_kps_count += roi_features[i].keypoints.size(); total_descriptors_height += roi_features[i].descriptors.rows; } @@ -300,7 +300,7 @@ void FeaturesFinder::operator ()(const Mat &image, ImageFeatures &features, cons features.keypoints[kp_idx].pt.x += (float)rois[i].x; features.keypoints[kp_idx].pt.y += (float)rois[i].y; } - Mat subdescr = features.descriptors.rowRange( + UMat subdescr = features.descriptors.rowRange( descr_offset, descr_offset + roi_features[i].descriptors.rows); roi_features[i].descriptors.copyTo(subdescr); descr_offset += roi_features[i].descriptors.rows; @@ -337,9 +337,9 @@ SurfFeaturesFinder::SurfFeaturesFinder(double hess_thresh, int num_octaves, int } } -void SurfFeaturesFinder::find(const Mat &image, ImageFeatures &features) +void SurfFeaturesFinder::find(InputArray image, ImageFeatures &features) { - Mat gray_image; + UMat gray_image; CV_Assert((image.type() == CV_8UC3) || (image.type() == CV_8UC1)); if(image.type() == CV_8UC3) { @@ -347,7 +347,7 @@ void SurfFeaturesFinder::find(const Mat &image, ImageFeatures &features) } else { - gray_image = image; + gray_image = image.getUMat(); } if (!surf) { @@ -356,7 +356,7 @@ void SurfFeaturesFinder::find(const Mat &image, ImageFeatures &features) } else { - Mat descriptors; + UMat descriptors; (*surf)(gray_image, Mat(), features.keypoints, descriptors); features.descriptors = descriptors.reshape(1, (int)features.keypoints.size()); } @@ -368,9 +368,9 @@ OrbFeaturesFinder::OrbFeaturesFinder(Size _grid_size, int n_features, float scal orb = makePtr(n_features * (99 + grid_size.area())/100/grid_size.area(), scaleFactor, nlevels); } -void OrbFeaturesFinder::find(const Mat &image, ImageFeatures &features) +void OrbFeaturesFinder::find(InputArray image, ImageFeatures &features) { - Mat gray_image; + UMat gray_image; CV_Assert((image.type() == CV_8UC3) || (image.type() == CV_8UC4) || (image.type() == CV_8UC1)); @@ -379,7 +379,7 @@ void OrbFeaturesFinder::find(const Mat &image, ImageFeatures &features) } else if (image.type() == CV_8UC4) { cvtColor(image, gray_image, COLOR_BGRA2GRAY); } else if (image.type() == CV_8UC1) { - gray_image=image; + gray_image = image.getUMat(); } else { CV_Error(Error::StsUnsupportedFormat, ""); } @@ -392,7 +392,8 @@ void OrbFeaturesFinder::find(const Mat &image, ImageFeatures &features) features.descriptors.release(); std::vector points; - Mat descriptors; + Mat _descriptors; + UMat descriptors; for (int r = 0; r < grid_size.height; ++r) for (int c = 0; c < grid_size.width; ++c) @@ -408,13 +409,13 @@ void OrbFeaturesFinder::find(const Mat &image, ImageFeatures &features) // << " xl=" << xl << ", xr=" << xr << ", gray_image.data=" << ((size_t)gray_image.data) << ", " // << "gray_image.dims=" << gray_image.dims << "\n"); - Mat gray_image_part=gray_image(Range(yl, yr), Range(xl, xr)); + UMat gray_image_part=gray_image(Range(yl, yr), Range(xl, xr)); // LOGLN("OrbFeaturesFinder::find: gray_image_part.empty=" << (gray_image_part.empty()?"true":"false") << ", " // << " gray_image_part.size()=(" << gray_image_part.size().width << "x" << gray_image_part.size().height << "), " // << " gray_image_part.dims=" << gray_image_part.dims << ", " // << " gray_image_part.data=" << ((size_t)gray_image_part.data) << "\n"); - (*orb)(gray_image_part, Mat(), points, descriptors); + (*orb)(gray_image_part, UMat(), points, descriptors); features.keypoints.reserve(features.keypoints.size() + points.size()); for (std::vector::iterator kp = points.begin(); kp != points.end(); ++kp) @@ -423,8 +424,10 @@ void OrbFeaturesFinder::find(const Mat &image, ImageFeatures &features) kp->pt.y += yl; features.keypoints.push_back(*kp); } - features.descriptors.push_back(descriptors); + _descriptors.push_back(descriptors.getMat(ACCESS_READ)); } + + features.descriptors = _descriptors.getUMat(ACCESS_READ); } } @@ -442,7 +445,7 @@ SurfFeaturesFinderGpu::SurfFeaturesFinderGpu(double hess_thresh, int num_octaves } -void SurfFeaturesFinderGpu::find(const Mat &image, ImageFeatures &features) +void SurfFeaturesFinderGpu::find(InputArray image, ImageFeatures &features) { CV_Assert(image.depth() == CV_8U); @@ -499,12 +502,12 @@ const MatchesInfo& MatchesInfo::operator =(const MatchesInfo &other) ////////////////////////////////////////////////////////////////////////////// void FeaturesMatcher::operator ()(const std::vector &features, std::vector &pairwise_matches, - const Mat &mask) + const UMat &mask) { const int num_images = static_cast(features.size()); CV_Assert(mask.empty() || (mask.type() == CV_8U && mask.cols == num_images && mask.rows)); - Mat_ mask_(mask); + Mat_ mask_(mask.getMat(ACCESS_READ)); if (mask_.empty()) mask_ = Mat::ones(num_images, num_images, CV_8U); diff --git a/modules/stitching/src/seam_finders.cpp b/modules/stitching/src/seam_finders.cpp index 2e5117e5b..4d5c8d163 100644 --- a/modules/stitching/src/seam_finders.cpp +++ b/modules/stitching/src/seam_finders.cpp @@ -46,8 +46,8 @@ namespace cv { namespace detail { -void PairwiseSeamFinder::find(const std::vector &src, const std::vector &corners, - std::vector &masks) +void PairwiseSeamFinder::find(const std::vector &src, const std::vector &corners, + std::vector &masks) { LOGLN("Finding seams..."); if (src.size() == 0) @@ -84,7 +84,7 @@ void PairwiseSeamFinder::run() void VoronoiSeamFinder::find(const std::vector &sizes, const std::vector &corners, - std::vector &masks) + std::vector &masks) { LOGLN("Finding seams..."); if (sizes.size() == 0) @@ -110,7 +110,7 @@ void VoronoiSeamFinder::findInPair(size_t first, size_t second, Rect roi) Mat submask2(roi.height + 2 * gap, roi.width + 2 * gap, CV_8U); Size img1 = sizes_[first], img2 = sizes_[second]; - Mat mask1 = masks_[first], mask2 = masks_[second]; + Mat mask1 = masks_[first].getMat(ACCESS_READ), mask2 = masks_[second].getMat(ACCESS_READ); Point tl1 = corners_[first], tl2 = corners_[second]; // Cut submasks with some gap @@ -160,7 +160,7 @@ void VoronoiSeamFinder::findInPair(size_t first, size_t second, Rect roi) DpSeamFinder::DpSeamFinder(CostFunction costFunc) : costFunc_(costFunc) {} -void DpSeamFinder::find(const std::vector &src, const std::vector &corners, std::vector &masks) +void DpSeamFinder::find(const std::vector &src, const std::vector &corners, std::vector &masks) { LOGLN("Finding seams..."); #if ENABLE_LOG @@ -176,13 +176,18 @@ void DpSeamFinder::find(const std::vector &src, const std::vector &c for (size_t j = i+1; j < src.size(); ++j) pairs.push_back(std::make_pair(i, j)); - sort(pairs.begin(), pairs.end(), ImagePairLess(src, corners)); + { + std::vector _src(src.size()); + for (size_t i = 0; i < src.size(); ++i) _src[i] = src[i].getMat(ACCESS_READ); + sort(pairs.begin(), pairs.end(), ImagePairLess(_src, corners)); + } std::reverse(pairs.begin(), pairs.end()); for (size_t i = 0; i < pairs.size(); ++i) { size_t i0 = pairs[i].first, i1 = pairs[i].second; - process(src[i0], src[i1], corners[i0], corners[i1], masks[i0], masks[i1]); + Mat mask0 = masks[i0].getMat(ACCESS_RW), mask1 = masks[i1].getMat(ACCESS_RW); + process(src[i0].getMat(ACCESS_READ), src[i1].getMat(ACCESS_READ), corners[i0], corners[i1], mask0, mask1); } LOGLN("Finding seams, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); @@ -1055,7 +1060,7 @@ public: ~Impl() {} - void find(const std::vector &src, const std::vector &corners, std::vector &masks); + void find(const std::vector &src, const std::vector &corners, std::vector &masks); void findInPair(size_t first, size_t second, Rect roi); private: @@ -1072,8 +1077,8 @@ private: }; -void GraphCutSeamFinder::Impl::find(const std::vector &src, const std::vector &corners, - std::vector &masks) +void GraphCutSeamFinder::Impl::find(const std::vector &src, const std::vector &corners, + std::vector &masks) { // Compute gradients dx_.resize(src.size()); @@ -1207,10 +1212,10 @@ void GraphCutSeamFinder::Impl::setGraphWeightsColorGrad( void GraphCutSeamFinder::Impl::findInPair(size_t first, size_t second, Rect roi) { - Mat img1 = images_[first], img2 = images_[second]; + Mat img1 = images_[first].getMat(ACCESS_READ), img2 = images_[second].getMat(ACCESS_READ); Mat dx1 = dx_[first], dx2 = dx_[second]; Mat dy1 = dy_[first], dy2 = dy_[second]; - Mat mask1 = masks_[first], mask2 = masks_[second]; + Mat mask1 = masks_[first].getMat(ACCESS_RW), mask2 = masks_[second].getMat(ACCESS_RW); Point tl1 = corners_[first], tl2 = corners_[second]; const int gap = 10; @@ -1309,16 +1314,16 @@ GraphCutSeamFinder::GraphCutSeamFinder(int cost_type, float terminal_cost, float GraphCutSeamFinder::~GraphCutSeamFinder() {} -void GraphCutSeamFinder::find(const std::vector &src, const std::vector &corners, - std::vector &masks) +void GraphCutSeamFinder::find(const std::vector &src, const std::vector &corners, + std::vector &masks) { impl_->find(src, corners, masks); } #ifdef HAVE_OPENCV_CUDA -void GraphCutSeamFinderGpu::find(const std::vector &src, const std::vector &corners, - std::vector &masks) +void GraphCutSeamFinderGpu::find(const std::vector &src, const std::vector &corners, + std::vector &masks) { // Compute gradients dx_.resize(src.size()); @@ -1350,10 +1355,10 @@ void GraphCutSeamFinderGpu::find(const std::vector &src, const std::vector< void GraphCutSeamFinderGpu::findInPair(size_t first, size_t second, Rect roi) { - Mat img1 = images_[first], img2 = images_[second]; + Mat img1 = images_[first].getMat(ACCESS_READ), img2 = images_[second].getMat(ACCESS_READ); Mat dx1 = dx_[first], dx2 = dx_[second]; Mat dy1 = dy_[first], dy2 = dy_[second]; - Mat mask1 = masks_[first], mask2 = masks_[second]; + Mat mask1 = masks_[first].getMat(ACCESS_READ), mask2 = masks_[second].getMat(ACCESS_READ); Point tl1 = corners_[first], tl2 = corners_[second]; const int gap = 10; diff --git a/modules/stitching/src/stitcher.cpp b/modules/stitching/src/stitcher.cpp index 5683ec308..ae4672639 100644 --- a/modules/stitching/src/stitcher.cpp +++ b/modules/stitching/src/stitcher.cpp @@ -86,15 +86,15 @@ Stitcher Stitcher::createDefault(bool try_use_gpu) } -Stitcher::Status Stitcher::estimateTransform(InputArray images) +Stitcher::Status Stitcher::estimateTransform(InputArrayOfArrays images) { return estimateTransform(images, std::vector >()); } -Stitcher::Status Stitcher::estimateTransform(InputArray images, const std::vector > &rois) +Stitcher::Status Stitcher::estimateTransform(InputArrayOfArrays images, const std::vector > &rois) { - images.getMatVector(imgs_); + images.getUMatVector(imgs_); rois_ = rois; Status status; @@ -112,21 +112,21 @@ Stitcher::Status Stitcher::estimateTransform(InputArray images, const std::vecto Stitcher::Status Stitcher::composePanorama(OutputArray pano) { - return composePanorama(std::vector(), pano); + return composePanorama(std::vector(), pano); } -Stitcher::Status Stitcher::composePanorama(InputArray images, OutputArray pano) +Stitcher::Status Stitcher::composePanorama(InputArrayOfArrays images, OutputArray pano) { LOGLN("Warping images (auxiliary)... "); - std::vector imgs; - images.getMatVector(imgs); + std::vector imgs; + images.getUMatVector(imgs); if (!imgs.empty()) { CV_Assert(imgs.size() == imgs_.size()); - Mat img; + UMat img; seam_est_imgs_.resize(imgs.size()); for (size_t i = 0; i < imgs.size(); ++i) @@ -136,8 +136,8 @@ Stitcher::Status Stitcher::composePanorama(InputArray images, OutputArray pano) seam_est_imgs_[i] = img.clone(); } - std::vector seam_est_imgs_subset; - std::vector imgs_subset; + std::vector seam_est_imgs_subset; + std::vector imgs_subset; for (size_t i = 0; i < indices_.size(); ++i) { @@ -149,17 +149,17 @@ Stitcher::Status Stitcher::composePanorama(InputArray images, OutputArray pano) imgs_ = imgs_subset; } - Mat &pano_ = pano.getMatRef(); + UMat pano_; #if ENABLE_LOG int64 t = getTickCount(); #endif std::vector corners(imgs_.size()); - std::vector masks_warped(imgs_.size()); - std::vector images_warped(imgs_.size()); + std::vector masks_warped(imgs_.size()); + std::vector images_warped(imgs_.size()); std::vector sizes(imgs_.size()); - std::vector masks(imgs_.size()); + std::vector masks(imgs_.size()); // Prepare image masks for (size_t i = 0; i < imgs_.size(); ++i) @@ -185,7 +185,7 @@ Stitcher::Status Stitcher::composePanorama(InputArray images, OutputArray pano) w->warp(masks[i], K, cameras_[i].R, INTER_NEAREST, BORDER_CONSTANT, masks_warped[i]); } - std::vector images_warped_f(imgs_.size()); + std::vector images_warped_f(imgs_.size()); for (size_t i = 0; i < imgs_.size(); ++i) images_warped[i].convertTo(images_warped_f[i], CV_32F); @@ -206,8 +206,8 @@ Stitcher::Status Stitcher::composePanorama(InputArray images, OutputArray pano) t = getTickCount(); #endif - Mat img_warped, img_warped_s; - Mat dilated_mask, seam_mask, mask, mask_warped; + UMat img_warped, img_warped_s; + UMat dilated_mask, seam_mask, mask, mask_warped; //double compose_seam_aspect = 1; double compose_work_aspect = 1; @@ -216,7 +216,7 @@ Stitcher::Status Stitcher::composePanorama(InputArray images, OutputArray pano) double compose_scale = 1; bool is_compose_scale_set = false; - Mat full_img, img; + UMat full_img, img; for (size_t img_idx = 0; img_idx < imgs_.size(); ++img_idx) { LOGLN("Compositing image #" << indices_[img_idx] + 1); @@ -290,7 +290,7 @@ Stitcher::Status Stitcher::composePanorama(InputArray images, OutputArray pano) dilate(masks_warped[img_idx], dilated_mask, Mat()); resize(dilated_mask, seam_mask, mask_warped.size()); - mask_warped = seam_mask & mask_warped; + bitwise_and(seam_mask, mask_warped, mask_warped); if (!is_blender_prepared) { @@ -302,7 +302,7 @@ Stitcher::Status Stitcher::composePanorama(InputArray images, OutputArray pano) blender_->feed(img_warped_s, mask_warped, corners[img_idx]); } - Mat result, result_mask; + UMat result, result_mask; blender_->blend(result, result_mask); LOGLN("Compositing, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); @@ -311,11 +311,13 @@ Stitcher::Status Stitcher::composePanorama(InputArray images, OutputArray pano) // so convert it to avoid user confusing result.convertTo(pano_, CV_8U); + pano.assign(pano_); + return OK; } -Stitcher::Status Stitcher::stitch(InputArray images, OutputArray pano) +Stitcher::Status Stitcher::stitch(InputArrayOfArrays images, OutputArray pano) { Status status = estimateTransform(images); if (status != OK) @@ -324,7 +326,7 @@ Stitcher::Status Stitcher::stitch(InputArray images, OutputArray pano) } -Stitcher::Status Stitcher::stitch(InputArray images, const std::vector > &rois, OutputArray pano) +Stitcher::Status Stitcher::stitch(InputArrayOfArrays images, const std::vector > &rois, OutputArray pano) { Status status = estimateTransform(images, rois); if (status != OK) @@ -346,7 +348,7 @@ Stitcher::Status Stitcher::matchImages() seam_scale_ = 1; bool is_work_scale_set = false; bool is_seam_scale_set = false; - Mat full_img, img; + UMat full_img, img; features_.resize(imgs_.size()); seam_est_imgs_.resize(imgs_.size()); full_img_sizes_.resize(imgs_.size()); @@ -420,8 +422,8 @@ Stitcher::Status Stitcher::matchImages() // Leave only images we are sure are from the same panorama indices_ = detail::leaveBiggestComponent(features_, pairwise_matches_, (float)conf_thresh_); - std::vector seam_est_imgs_subset; - std::vector imgs_subset; + std::vector seam_est_imgs_subset; + std::vector imgs_subset; std::vector full_img_sizes_subset; for (size_t i = 0; i < indices_.size(); ++i) { diff --git a/modules/stitching/src/util.cpp b/modules/stitching/src/util.cpp index f6abf9e48..5e026f098 100644 --- a/modules/stitching/src/util.cpp +++ b/modules/stitching/src/util.cpp @@ -113,7 +113,7 @@ bool overlapRoi(Point tl1, Point tl2, Size sz1, Size sz2, Rect &roi) } -Rect resultRoi(const std::vector &corners, const std::vector &images) +Rect resultRoi(const std::vector &corners, const std::vector &images) { std::vector sizes(images.size()); for (size_t i = 0; i < images.size(); ++i) diff --git a/samples/cpp/stitching_detailed.cpp b/samples/cpp/stitching_detailed.cpp index 5eb3df46c..93389a919 100644 --- a/samples/cpp/stitching_detailed.cpp +++ b/samples/cpp/stitching_detailed.cpp @@ -348,7 +348,9 @@ int main(int argc, char* argv[]) int64 app_start_time = getTickCount(); #endif +#if 0 cv::setBreakOnError(true); +#endif int retval = parseCmdArgs(argc, argv); if (retval) @@ -554,10 +556,10 @@ int main(int argc, char* argv[]) #endif vector corners(num_images); - vector masks_warped(num_images); - vector images_warped(num_images); + vector masks_warped(num_images); + vector images_warped(num_images); vector sizes(num_images); - vector masks(num_images); + vector masks(num_images); // Preapre images masks for (int i = 0; i < num_images; ++i) @@ -645,7 +647,7 @@ int main(int argc, char* argv[]) warper->warp(masks[i], K, cameras[i].R, INTER_NEAREST, BORDER_CONSTANT, masks_warped[i]); } - vector images_warped_f(num_images); + vector images_warped_f(num_images); for (int i = 0; i < num_images; ++i) images_warped[i].convertTo(images_warped_f[i], CV_32F); From 89e3e448f58ebff525888ed97f6b7f7768ab567f Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Wed, 12 Mar 2014 15:54:55 +0400 Subject: [PATCH 03/15] features2d: workaround for 'utrainDescCollection' issue (PR #2142) --- modules/features2d/src/matchers.cpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/modules/features2d/src/matchers.cpp b/modules/features2d/src/matchers.cpp index 6bc925f8a..91ee9d9bf 100644 --- a/modules/features2d/src/matchers.cpp +++ b/modules/features2d/src/matchers.cpp @@ -1232,6 +1232,13 @@ void FlannBasedMatcher::train() { if( !flannIndex || mergedDescriptors.size() < addedDescCount ) { + // FIXIT: Workaround for 'utrainDescCollection' issue (PR #2142) + if (!utrainDescCollection.empty()) + { + CV_Assert(trainDescCollection.size() == 0); + for (size_t i = 0; i < utrainDescCollection.size(); ++i) + trainDescCollection.push_back(utrainDescCollection[i].getMat(ACCESS_READ)); + } mergedDescriptors.set( trainDescCollection ); flannIndex = makePtr( mergedDescriptors.getDescriptors(), *indexParams ); } From c1ea6f3c4253b2a8b4ae897e1d129547babf236a Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Wed, 12 Mar 2014 14:54:22 +0400 Subject: [PATCH 04/15] TAPI: stitching: improve warpers --- .../include/opencv2/stitching/detail/util.hpp | 1 + .../opencv2/stitching/detail/warpers.hpp | 43 +--- .../opencv2/stitching/detail/warpers_inl.hpp | 2 +- .../include/opencv2/stitching/warpers.hpp | 18 -- .../stitching/perf/opencl/perf_warpers.cpp | 24 +-- modules/stitching/src/precomp.hpp | 1 + modules/stitching/src/stitcher.cpp | 8 +- modules/stitching/src/warpers.cpp | 121 +++++++++++- modules/stitching/src/warpers_ocl.cpp | 187 ------------------ modules/stitching/test/ocl/test_warpers.cpp | 41 ++-- samples/cpp/stitching_detailed.cpp | 28 +-- 11 files changed, 149 insertions(+), 325 deletions(-) delete mode 100644 modules/stitching/src/warpers_ocl.cpp diff --git a/modules/stitching/include/opencv2/stitching/detail/util.hpp b/modules/stitching/include/opencv2/stitching/detail/util.hpp index 561880cc2..6b1c5f34f 100644 --- a/modules/stitching/include/opencv2/stitching/detail/util.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/util.hpp @@ -71,6 +71,7 @@ #define LOG_(_level, _msg) \ for(;;) \ { \ + using namespace std; \ if ((_level) >= ::cv::detail::stitchingLogLevel()) \ { \ LOG_STITCHING_MSG(_msg); \ diff --git a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp index 093f07cc1..c8869f116 100644 --- a/modules/stitching/include/opencv2/stitching/detail/warpers.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/warpers.hpp @@ -160,6 +160,8 @@ class CV_EXPORTS SphericalWarper : public RotationWarperBase public: SphericalWarper(float scale) { projector_.scale = scale; } + Rect buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap); + Point warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst); protected: void detectResultRoi(Size src_size, Point &dst_tl, Point &dst_br); }; @@ -178,6 +180,8 @@ class CV_EXPORTS CylindricalWarper : public RotationWarperBase Point RotationWarperBase

::warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst) { - Mat xmap, ymap; + UMat xmap, ymap; Rect dst_roi = buildMaps(src.size(), K, R, xmap, ymap); dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); diff --git a/modules/stitching/include/opencv2/stitching/warpers.hpp b/modules/stitching/include/opencv2/stitching/warpers.hpp index cdcb35c20..da5fe2618 100644 --- a/modules/stitching/include/opencv2/stitching/warpers.hpp +++ b/modules/stitching/include/opencv2/stitching/warpers.hpp @@ -167,24 +167,6 @@ 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/perf/opencl/perf_warpers.cpp b/modules/stitching/perf/opencl/perf_warpers.cpp index 21fe22da7..6a8be4ebe 100644 --- a/modules/stitching/perf/opencl/perf_warpers.cpp +++ b/modules/stitching/perf/opencl/perf_warpers.cpp @@ -63,24 +63,12 @@ public: explicit WarperBase(int type, Size srcSize) { Ptr creator; - if (cv::ocl::useOpenCL()) - { - if (type == SphericalWarperType) - creator = makePtr(); - else if (type == CylindricalWarperType) - creator = makePtr(); - else if (type == PlaneWarperType) - creator = makePtr(); - } - else - { - if (type == SphericalWarperType) - creator = makePtr(); - else if (type == CylindricalWarperType) - creator = makePtr(); - else if (type == PlaneWarperType) - creator = makePtr(); - } + if (type == SphericalWarperType) + creator = makePtr(); + else if (type == CylindricalWarperType) + creator = makePtr(); + else if (type == PlaneWarperType) + creator = makePtr(); CV_Assert(!creator.empty()); K = Mat::eye(3, 3, CV_32FC1); diff --git a/modules/stitching/src/precomp.hpp b/modules/stitching/src/precomp.hpp index 499202fa0..18ce41309 100644 --- a/modules/stitching/src/precomp.hpp +++ b/modules/stitching/src/precomp.hpp @@ -51,6 +51,7 @@ #include #include #include +#include #include #include "opencv2/core.hpp" #include "opencv2/core/ocl.hpp" diff --git a/modules/stitching/src/stitcher.cpp b/modules/stitching/src/stitcher.cpp index ae4672639..2d1288663 100644 --- a/modules/stitching/src/stitcher.cpp +++ b/modules/stitching/src/stitcher.cpp @@ -309,9 +309,7 @@ Stitcher::Status Stitcher::composePanorama(InputArrayOfArrays images, OutputArra // Preliminary result is in CV_16SC3 format, but all values are in [0,255] range, // so convert it to avoid user confusing - result.convertTo(pano_, CV_8U); - - pano.assign(pano_); + result.convertTo(pano, CV_8U); return OK; } @@ -456,7 +454,7 @@ Stitcher::Status Stitcher::estimateCameraParams() Mat R; cameras_[i].R.convertTo(R, CV_32F); cameras_[i].R = R; - LOGLN("Initial intrinsic parameters #" << indices_[i] + 1 << ":\n " << cameras_[i].K()); + //LOGLN("Initial intrinsic parameters #" << indices_[i] + 1 << ":\n " << cameras_[i].K()); } bundle_adjuster_->setConfThresh(conf_thresh_); @@ -467,7 +465,7 @@ Stitcher::Status Stitcher::estimateCameraParams() std::vector focals; for (size_t i = 0; i < cameras_.size(); ++i) { - LOGLN("Camera #" << indices_[i] + 1 << ":\n" << cameras_[i].K()); + //LOGLN("Camera #" << indices_[i] + 1 << ":\n" << cameras_[i].K()); focals.push_back(cameras_[i].focal); } diff --git a/modules/stitching/src/warpers.cpp b/modules/stitching/src/warpers.cpp index eb15d44c0..a05ad1f5d 100644 --- a/modules/stitching/src/warpers.cpp +++ b/modules/stitching/src/warpers.cpp @@ -41,6 +41,7 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" namespace cv { namespace detail { @@ -86,7 +87,6 @@ Point2f PlaneWarper::warpPoint(const Point2f &pt, InputArray K, InputArray R, In return uv; } - Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, InputArray T, OutputArray _xmap, OutputArray _ymap) { projector_.setCameraParams(K, R, T); @@ -94,8 +94,29 @@ Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, InputArra Point dst_tl, dst_br; detectResultRoi(src_size, dst_tl, dst_br); - _xmap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); - _ymap.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, CV_32F); + Size dsize(dst_br.x - dst_tl.x + 1, dst_br.y - dst_tl.y + 1); + _xmap.create(dsize, CV_32FC1); + _ymap.create(dsize, CV_32FC1); + + if (ocl::useOpenCL()) // TODO !!!!! check T + { + ocl::Kernel k("buildWarpPlaneMaps", ocl::stitching::warpers_oclsrc); + if (!k.empty()) + { + + Mat k_rinv(1, 9, CV_32FC1, projector_.k_rinv), t(1, 3, CV_32FC1, projector_.t); + UMat uxmap = _xmap.getUMat(), uymap = _ymap.getUMat(), + uk_rinv = k_rinv.getUMat(ACCESS_READ), ut = t.getUMat(ACCESS_READ); + + k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), + ocl::KernelArg::PtrReadOnly(uk_rinv), ocl::KernelArg::PtrReadOnly(ut), + dst_tl.x, dst_tl.y, projector_.scale); + + size_t globalsize[2] = { dsize.width, dsize.height }; + if (k.run(2, globalsize, NULL, true)) + return Rect(dst_tl, dst_br); + } + } Mat xmap = _xmap.getMat(), ymap = _ymap.getMat(); @@ -117,11 +138,11 @@ Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, InputArra Point PlaneWarper::warp(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, OutputArray dst) { - Mat xmap, ymap; - Rect dst_roi = buildMaps(src.size(), K, R, T, xmap, ymap); + UMat 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()); - remap(src, dst, xmap, ymap, interp_mode, border_mode); + remap(src, dst, uxmap, uymap, interp_mode, border_mode); return dst_roi.tl(); } @@ -341,5 +362,93 @@ void SphericalPortraitWarper::detectResultRoi(Size src_size, Point &dst_tl, Poin dst_br.y = static_cast(br_vf); } +/////////////////////////////////////////// SphericalWarper //////////////////////////////////////// + +Rect SphericalWarper::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) +{ + if (ocl::useOpenCL()) + { + ocl::Kernel k("buildWarpSphericalMaps", ocl::stitching::warpers_oclsrc); + if (!k.empty()) + { + projector_.setCameraParams(K, R); + + Point dst_tl, dst_br; + detectResultRoi(src_size, dst_tl, dst_br); + + Size dsize(dst_br.x - dst_tl.x + 1, dst_br.y - dst_tl.y + 1); + xmap.create(dsize, CV_32FC1); + ymap.create(dsize, CV_32FC1); + + Mat k_rinv(1, 9, CV_32FC1, projector_.k_rinv); + UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), uk_rinv = k_rinv.getUMat(ACCESS_READ); + + k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), + ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, projector_.scale); + + size_t globalsize[2] = { dsize.width, dsize.height }; + if (k.run(2, globalsize, NULL, true)) + return Rect(dst_tl, dst_br); + } + } + + return RotationWarperBase::buildMaps(src_size, K, R, xmap, ymap); +} + +Point SphericalWarper::warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst) +{ + UMat uxmap, uymap; + Rect dst_roi = buildMaps(src.size(), K, R, uxmap, uymap); + + dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); + remap(src, dst, uxmap, uymap, interp_mode, border_mode); + + return dst_roi.tl(); +} + +/////////////////////////////////////////// CylindricalWarper //////////////////////////////////////// + +Rect CylindricalWarper::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray ymap) +{ + if (ocl::useOpenCL()) + { + ocl::Kernel k("buildWarpCylindricalMaps", ocl::stitching::warpers_oclsrc); + if (!k.empty()) + { + projector_.setCameraParams(K, R); + + Point dst_tl, dst_br; + detectResultRoi(src_size, dst_tl, dst_br); + + Size dsize(dst_br.x - dst_tl.x + 1, dst_br.y - dst_tl.y + 1); + xmap.create(dsize, CV_32FC1); + ymap.create(dsize, CV_32FC1); + + Mat k_rinv(1, 9, CV_32FC1, projector_.k_rinv); + UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), uk_rinv = k_rinv.getUMat(ACCESS_READ); + + k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), + ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, projector_.scale); + + size_t globalsize[2] = { dsize.width, dsize.height }; + if (k.run(2, globalsize, NULL, true)) + return Rect(dst_tl, dst_br); + } + } + + return RotationWarperBase::buildMaps(src_size, K, R, xmap, ymap); +} + +Point CylindricalWarper::warp(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst) +{ + UMat uxmap, uymap; + Rect dst_roi = buildMaps(src.size(), K, R, uxmap, uymap); + + dst.create(dst_roi.height + 1, dst_roi.width + 1, src.type()); + remap(src, dst, uxmap, uymap, interp_mode, border_mode); + + return dst_roi.tl(); +} + } // namespace detail } // namespace cv diff --git a/modules/stitching/src/warpers_ocl.cpp b/modules/stitching/src/warpers_ocl.cpp deleted file mode 100644 index ef8f31677..000000000 --- a/modules/stitching/src/warpers_ocl.cpp +++ /dev/null @@ -1,187 +0,0 @@ -/*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, InputArray K, InputArray R, InputArray T, OutputArray xmap, OutputArray ymap) -{ - projector_.setCameraParams(K, R, T); - - 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()) - { - Size dsize(dst_br.x - dst_tl.x + 1, dst_br.y - dst_tl.y + 1); - xmap.create(dsize, CV_32FC1); - ymap.create(dsize, CV_32FC1); - - Mat k_rinv(1, 9, CV_32FC1, projector_.k_rinv), t(1, 3, CV_32FC1, projector_.t); - UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), - uk_rinv = k_rinv.getUMat(ACCESS_READ), ut = t.getUMat(ACCESS_READ); - - k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), - ocl::KernelArg::PtrReadOnly(uk_rinv), ocl::KernelArg::PtrReadOnly(ut), - dst_tl.x, dst_tl.y, projector_.scale); - - size_t globalsize[2] = { dsize.width, dsize.height }; - 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(InputArray src, InputArray K, InputArray R, InputArray T, int interp_mode, int border_mode, OutputArray dst) -{ - UMat 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(); - remap(src, udst, uxmap, uymap, interp_mode, border_mode); - - return dst_roi.tl(); -} - -/////////////////////////////////////////// SphericalWarperOcl //////////////////////////////////////// - -Rect SphericalWarperOcl::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray 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()) - { - Size dsize(dst_br.x - dst_tl.x + 1, dst_br.y - dst_tl.y + 1); - xmap.create(dsize, CV_32FC1); - ymap.create(dsize, CV_32FC1); - - Mat k_rinv(1, 9, CV_32FC1, projector_.k_rinv); - UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), uk_rinv = k_rinv.getUMat(ACCESS_READ); - - k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), - ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, projector_.scale); - - size_t globalsize[2] = { dsize.width, dsize.height }; - 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(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst) -{ - UMat 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(); - remap(src, udst, uxmap, uymap, interp_mode, border_mode); - - return dst_roi.tl(); -} - -/////////////////////////////////////////// CylindricalWarperOcl //////////////////////////////////////// - -Rect CylindricalWarperOcl::buildMaps(Size src_size, InputArray K, InputArray R, OutputArray xmap, OutputArray 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()) - { - Size dsize(dst_br.x - dst_tl.x + 1, dst_br.y - dst_tl.y + 1); - xmap.create(dsize, CV_32FC1); - ymap.create(dsize, CV_32FC1); - - Mat k_rinv(1, 9, CV_32FC1, projector_.k_rinv); - UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), uk_rinv = k_rinv.getUMat(ACCESS_READ); - - k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), - ocl::KernelArg::PtrReadOnly(uk_rinv), dst_tl.x, dst_tl.y, projector_.scale); - - size_t globalsize[2] = { dsize.width, dsize.height }; - 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(InputArray src, InputArray K, InputArray R, int interp_mode, int border_mode, OutputArray dst) -{ - UMat 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(); - 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 index 94050e966..43f0e9741 100644 --- a/modules/stitching/test/ocl/test_warpers.cpp +++ b/modules/stitching/test/ocl/test_warpers.cpp @@ -48,13 +48,11 @@ namespace cvtest { namespace ocl { -///////////////////////// WarperTestBase /////////////////////////// - struct WarperTestBase : public Test, public TestUtils { Mat src, dst, xmap, ymap; - Mat udst, uxmap, uymap; + UMat usrc, udst, uxmap, uymap; Mat K, R; virtual void generateTestData() @@ -62,6 +60,7 @@ struct WarperTestBase : Size size = randomSize(1, MAX_VALUE); src = randomMat(size, CV_32FC1, -500, 500); + src.copyTo(usrc); K = Mat::eye(3, 3, CV_32FC1); float angle = (float)(30.0 * CV_PI / 180.0); @@ -81,70 +80,64 @@ struct WarperTestBase : } }; -//////////////////////////////// SphericalWarperOcl ///////////////////////////////////////////////// +typedef WarperTestBase SphericalWarperTest; -typedef WarperTestBase SphericalWarperOclTest; - -OCL_TEST_F(SphericalWarperOclTest, Mat) +OCL_TEST_F(SphericalWarperTest, Mat) { for (int j = 0; j < test_loop_times; j++) { generateTestData(); - Ptr creator = makePtr(); + 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_ON(warper->buildMaps(usrc.size(), K, R, uxmap, uymap)); OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); - OCL_ON(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); + OCL_ON(warper->warp(usrc, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); Near(1e-4); } } -//////////////////////////////// CylindricalWarperOcl ///////////////////////////////////////////////// +typedef WarperTestBase CylindricalWarperTest; -typedef WarperTestBase CylindricalWarperOclTest; - -OCL_TEST_F(CylindricalWarperOclTest, Mat) +OCL_TEST_F(CylindricalWarperTest, Mat) { for (int j = 0; j < test_loop_times; j++) { generateTestData(); - Ptr creator = makePtr(); + 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_ON(warper->buildMaps(usrc.size(), K, R, uxmap, uymap)); OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); - OCL_ON(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); + OCL_ON(warper->warp(usrc, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); Near(1e-4); } } -//////////////////////////////// PlaneWarperOcl ///////////////////////////////////////////////// +typedef WarperTestBase PlaneWarperTest; -typedef WarperTestBase PlaneWarperOclTest; - -OCL_TEST_F(PlaneWarperOclTest, Mat) +OCL_TEST_F(PlaneWarperTest, Mat) { for (int j = 0; j < test_loop_times; j++) { generateTestData(); - Ptr creator = makePtr(); + 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_ON(warper->buildMaps(usrc.size(), K, R, uxmap, uymap)); OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); - OCL_ON(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); + OCL_ON(warper->warp(usrc, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); Near(1e-4); } diff --git a/samples/cpp/stitching_detailed.cpp b/samples/cpp/stitching_detailed.cpp index 93389a919..df0a9abc8 100644 --- a/samples/cpp/stitching_detailed.cpp +++ b/samples/cpp/stitching_detailed.cpp @@ -74,9 +74,6 @@ static void printUsage() " --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" " Resolution for image registration step. The default is 0.6 Mpx.\n" @@ -127,7 +124,6 @@ static void printUsage() vector img_names; bool preview = false; bool try_cuda = false; -bool try_ocl = false; double work_megapix = 0.6; double seam_megapix = 0.1; double compose_megapix = -1; @@ -178,19 +174,6 @@ static int parseCmdArgs(int argc, char** argv) } 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++; - } else if (string(argv[i]) == "--work_megapix") { work_megapix = atof(argv[i + 1]); @@ -571,17 +554,8 @@ 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 - else if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) + if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0) { if (warp_type == "plane") warper_creator = makePtr(); From 2737e3c99bc1c9b656d322f2d895e7e67d1960c1 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Wed, 26 Feb 2014 17:15:46 +0400 Subject: [PATCH 05/15] TAPI: stitching: workaround for problem getUMat() call --- modules/stitching/src/matchers.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/modules/stitching/src/matchers.cpp b/modules/stitching/src/matchers.cpp index 1ce1fc3b2..ceb3d3da7 100644 --- a/modules/stitching/src/matchers.cpp +++ b/modules/stitching/src/matchers.cpp @@ -427,7 +427,9 @@ void OrbFeaturesFinder::find(InputArray image, ImageFeatures &features) _descriptors.push_back(descriptors.getMat(ACCESS_READ)); } - features.descriptors = _descriptors.getUMat(ACCESS_READ); + // TODO optimize copyTo() + //features.descriptors = _descriptors.getUMat(ACCESS_READ); + _descriptors.copyTo(features.descriptors); } } From 7a5c1babde8b6a4f8b5820e32790c5e88f35a680 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Thu, 27 Feb 2014 18:09:56 +0400 Subject: [PATCH 06/15] TAPI: stitching: optimize exposure --- .../stitching/detail/exposure_compensate.hpp | 2 +- modules/stitching/src/exposure_compensate.cpp | 25 +++++++++++-------- 2 files changed, 15 insertions(+), 12 deletions(-) diff --git a/modules/stitching/include/opencv2/stitching/detail/exposure_compensate.hpp b/modules/stitching/include/opencv2/stitching/detail/exposure_compensate.hpp index 5626b06f2..9cd8b32f4 100644 --- a/modules/stitching/include/opencv2/stitching/detail/exposure_compensate.hpp +++ b/modules/stitching/include/opencv2/stitching/detail/exposure_compensate.hpp @@ -97,7 +97,7 @@ public: private: int bl_width_, bl_height_; - std::vector > gain_maps_; + std::vector gain_maps_; }; } // namespace detail diff --git a/modules/stitching/src/exposure_compensate.cpp b/modules/stitching/src/exposure_compensate.cpp index 32112d69f..1f04fff9f 100644 --- a/modules/stitching/src/exposure_compensate.cpp +++ b/modules/stitching/src/exposure_compensate.cpp @@ -208,11 +208,14 @@ void BlocksGainCompensator::feed(const std::vector &corners, const std::v for (int img_idx = 0; img_idx < num_images; ++img_idx) { Size bl_per_img = bl_per_imgs[img_idx]; - gain_maps_[img_idx].create(bl_per_img); + gain_maps_[img_idx].create(bl_per_img, CV_32F); - for (int by = 0; by < bl_per_img.height; ++by) - for (int bx = 0; bx < bl_per_img.width; ++bx, ++bl_idx) - gain_maps_[img_idx](by, bx) = static_cast(gains[bl_idx]); + { + Mat_ gain_map = gain_maps_[img_idx].getMat(ACCESS_WRITE); + for (int by = 0; by < bl_per_img.height; ++by) + for (int bx = 0; bx < bl_per_img.width; ++bx, ++bl_idx) + gain_map(by, bx) = static_cast(gains[bl_idx]); + } sepFilter2D(gain_maps_[img_idx], gain_maps_[img_idx], CV_32F, ker, ker); sepFilter2D(gain_maps_[img_idx], gain_maps_[img_idx], CV_32F, ker, ker); @@ -222,16 +225,16 @@ void BlocksGainCompensator::feed(const std::vector &corners, const std::v void BlocksGainCompensator::apply(int index, Point /*corner*/, InputOutputArray _image, InputArray /*mask*/) { - Mat image = _image.getMat(); + CV_Assert(_image.type() == CV_8UC3); - CV_Assert(image.type() == CV_8UC3); - - Mat_ gain_map; - if (gain_maps_[index].size() == image.size()) - gain_map = gain_maps_[index]; + UMat u_gain_map; + if (gain_maps_[index].size() == _image.size()) + u_gain_map = gain_maps_[index]; else - resize(gain_maps_[index], gain_map, image.size(), 0, 0, INTER_LINEAR); + resize(gain_maps_[index], u_gain_map, _image.size(), 0, 0, INTER_LINEAR); + Mat_ gain_map = u_gain_map.getMat(ACCESS_READ); + Mat image = _image.getMat(); for (int y = 0; y < image.rows; ++y) { const float* gain_row = gain_map.ptr(y); From c4a31a1ab161c8cb0a3977e52df5abc41666380a Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Wed, 26 Feb 2014 17:01:45 +0400 Subject: [PATCH 07/15] TAPI: stitching: optimize compare operation --- modules/stitching/src/blenders.cpp | 14 +++++++++----- modules/stitching/src/warpers.cpp | 2 +- 2 files changed, 10 insertions(+), 6 deletions(-) diff --git a/modules/stitching/src/blenders.cpp b/modules/stitching/src/blenders.cpp index 5012d9ba4..922c524d0 100644 --- a/modules/stitching/src/blenders.cpp +++ b/modules/stitching/src/blenders.cpp @@ -107,7 +107,9 @@ void Blender::feed(InputArray _img, InputArray _mask, Point tl) void Blender::blend(InputOutputArray dst, InputOutputArray dst_mask) { - dst_.setTo(Scalar::all(0), dst_mask_.getMat(ACCESS_READ) == 0); // TODO + UMat mask; + compare(dst_mask_, 0, mask, CMP_EQ); + dst_.setTo(Scalar::all(0), mask); dst.assign(dst_); dst_mask.assign(dst_mask_); dst_.release(); @@ -159,7 +161,7 @@ void FeatherBlender::feed(InputArray _img, InputArray mask, Point tl) void FeatherBlender::blend(InputOutputArray dst, InputOutputArray dst_mask) { normalizeUsingWeightMap(dst_weight_map_, dst_); - dst_mask_ = ((Mat)(dst_weight_map_.getMat(ACCESS_READ) > WEIGHT_EPS)).getUMat(ACCESS_READ); + compare(dst_weight_map_, WEIGHT_EPS, dst_mask_, CMP_GT); Blender::blend(dst, dst_mask); } @@ -301,7 +303,9 @@ void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) else // weight_type_ == CV_16S { mask.getUMat().convertTo(weight_map, CV_16S); - add(weight_map, Scalar::all(1), weight_map, mask.getMat(ACCESS_READ) != 0); // TODO + UMat add_mask; + compare(mask, 0, add_mask, CMP_NE); + add(weight_map, Scalar::all(1), weight_map, add_mask); } copyMakeBorder(weight_map, weight_pyr_gauss[0], top, bottom, left, right, BORDER_CONSTANT); @@ -388,8 +392,8 @@ void MultiBandBlender::blend(InputOutputArray dst, InputOutputArray dst_mask) dst_ = dst_pyr_laplace_[0]; dst_ = dst_(Range(0, dst_roi_final_.height), Range(0, dst_roi_final_.width)); - dst_mask_ = ((Mat)(dst_band_weights_[0].getMat(ACCESS_READ) > WEIGHT_EPS)).getUMat(ACCESS_READ); - dst_mask_ = dst_mask_(Range(0, dst_roi_final_.height), Range(0, dst_roi_final_.width)); + UMat _dst_mask; + compare(dst_band_weights_[0](Range(0, dst_roi_final_.height), Range(0, dst_roi_final_.width)), WEIGHT_EPS, dst_mask_, CMP_GT); dst_pyr_laplace_.clear(); dst_band_weights_.clear(); diff --git a/modules/stitching/src/warpers.cpp b/modules/stitching/src/warpers.cpp index a05ad1f5d..b6d1f8a8a 100644 --- a/modules/stitching/src/warpers.cpp +++ b/modules/stitching/src/warpers.cpp @@ -98,7 +98,7 @@ Rect PlaneWarper::buildMaps(Size src_size, InputArray K, InputArray R, InputArra _xmap.create(dsize, CV_32FC1); _ymap.create(dsize, CV_32FC1); - if (ocl::useOpenCL()) // TODO !!!!! check T + if (ocl::useOpenCL()) { ocl::Kernel k("buildWarpPlaneMaps", ocl::stitching::warpers_oclsrc); if (!k.empty()) From a7f69a37e345146bb5d23dc72053ae94bb1ea241 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Tue, 25 Feb 2014 13:41:07 +0400 Subject: [PATCH 08/15] stitching: use BORDER_CONSTANT in warp --- modules/stitching/src/stitcher.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/modules/stitching/src/stitcher.cpp b/modules/stitching/src/stitcher.cpp index 2d1288663..db79e0d56 100644 --- a/modules/stitching/src/stitcher.cpp +++ b/modules/stitching/src/stitcher.cpp @@ -179,7 +179,7 @@ Stitcher::Status Stitcher::composePanorama(InputArrayOfArrays images, OutputArra K(1,1) *= (float)seam_work_aspect_; K(1,2) *= (float)seam_work_aspect_; - corners[i] = w->warp(seam_est_imgs_[i], K, cameras_[i].R, INTER_LINEAR, BORDER_REFLECT, images_warped[i]); + corners[i] = w->warp(seam_est_imgs_[i], K, cameras_[i].R, INTER_LINEAR, BORDER_CONSTANT, images_warped[i]); sizes[i] = images_warped[i].size(); w->warp(masks[i], K, cameras_[i].R, INTER_NEAREST, BORDER_CONSTANT, masks_warped[i]); @@ -271,7 +271,7 @@ Stitcher::Status Stitcher::composePanorama(InputArrayOfArrays images, OutputArra cameras_[img_idx].K().convertTo(K, CV_32F); // Warp the current image - w->warp(img, K, cameras_[img_idx].R, INTER_LINEAR, BORDER_REFLECT, img_warped); + w->warp(img, K, cameras_[img_idx].R, INTER_LINEAR, BORDER_CONSTANT, img_warped); // Warp the current image mask mask.create(img_size, CV_8U); From c22d92c1cb5434eabce20a068adf444dee76e0ee Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Wed, 26 Feb 2014 15:15:20 +0400 Subject: [PATCH 09/15] stitching: extend logging --- modules/stitching/src/blenders.cpp | 25 +++++++++++++++-- modules/stitching/src/stitcher.cpp | 43 ++++++++++++++++++++++++++++++ 2 files changed, 66 insertions(+), 2 deletions(-) diff --git a/modules/stitching/src/blenders.cpp b/modules/stitching/src/blenders.cpp index 922c524d0..acb4987d7 100644 --- a/modules/stitching/src/blenders.cpp +++ b/modules/stitching/src/blenders.cpp @@ -248,7 +248,11 @@ void MultiBandBlender::prepare(Rect dst_roi) void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) { - Mat img = _img.getMat(); +#if ENABLE_LOG + int64 t = getTickCount(); +#endif + + UMat img = _img.getUMat(); CV_Assert(img.type() == CV_16SC3 || img.type() == CV_8UC3); CV_Assert(mask.type() == CV_8U); @@ -286,12 +290,22 @@ void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) UMat img_with_border; copyMakeBorder(_img, img_with_border, top, bottom, left, right, BORDER_REFLECT); + LOGLN(" Add border to the source image, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); +#if ENABLE_LOG + t = getTickCount(); +#endif + std::vector src_pyr_laplace; if (can_use_gpu_ && img_with_border.depth() == CV_16S) createLaplacePyrGpu(img_with_border, num_bands_, src_pyr_laplace); else createLaplacePyr(img_with_border, num_bands_, src_pyr_laplace); + LOGLN(" Create the source image Laplacian pyramid, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); +#if ENABLE_LOG + t = getTickCount(); +#endif + // Create the weight map Gaussian pyramid UMat weight_map; std::vector weight_pyr_gauss(num_bands_ + 1); @@ -313,6 +327,11 @@ void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) for (int i = 0; i < num_bands_; ++i) pyrDown(weight_pyr_gauss[i], weight_pyr_gauss[i + 1]); + LOGLN(" Create the weight map Gaussian pyramid, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); +#if ENABLE_LOG + t = getTickCount(); +#endif + int y_tl = tl_new.y - dst_roi_.y; int y_br = br_new.y - dst_roi_.y; int x_tl = tl_new.x - dst_roi_.x; @@ -348,7 +367,7 @@ void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) x_br /= 2; y_br /= 2; } } - else// weight_type_ == CV_16S + else // weight_type_ == CV_16S { for (int i = 0; i <= num_bands_; ++i) { @@ -377,6 +396,8 @@ void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) x_br /= 2; y_br /= 2; } } + + LOGLN(" Add weighted layer of the source image to the final Laplacian pyramid layer, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); } diff --git a/modules/stitching/src/stitcher.cpp b/modules/stitching/src/stitcher.cpp index db79e0d56..ddfdb5084 100644 --- a/modules/stitching/src/stitcher.cpp +++ b/modules/stitching/src/stitcher.cpp @@ -220,6 +220,9 @@ Stitcher::Status Stitcher::composePanorama(InputArrayOfArrays images, OutputArra for (size_t img_idx = 0; img_idx < imgs_.size(); ++img_idx) { LOGLN("Compositing image #" << indices_[img_idx] + 1); +#if ENABLE_LOG + int64 compositing_t = getTickCount(); +#endif // Read image and resize it if necessary full_img = imgs_[img_idx]; @@ -261,25 +264,48 @@ Stitcher::Status Stitcher::composePanorama(InputArrayOfArrays images, OutputArra } } if (std::abs(compose_scale - 1) > 1e-1) + { +#if ENABLE_LOG + int64 resize_t = getTickCount(); +#endif resize(full_img, img, Size(), compose_scale, compose_scale); + LOGLN(" resize time: " << ((getTickCount() - resize_t) / getTickFrequency()) << " sec"); + } else img = full_img; full_img.release(); Size img_size = img.size(); + LOGLN(" after resize time: " << ((getTickCount() - compositing_t) / getTickFrequency()) << " sec"); + Mat K; cameras_[img_idx].K().convertTo(K, CV_32F); +#if ENABLE_LOG + int64 pt = getTickCount(); +#endif // Warp the current image w->warp(img, K, cameras_[img_idx].R, INTER_LINEAR, BORDER_CONSTANT, img_warped); + LOGLN(" warp the current image: " << ((getTickCount() - pt) / getTickFrequency()) << " sec"); +#if ENABLE_LOG + pt = getTickCount(); +#endif // Warp the current image mask mask.create(img_size, CV_8U); mask.setTo(Scalar::all(255)); w->warp(mask, K, cameras_[img_idx].R, INTER_NEAREST, BORDER_CONSTANT, mask_warped); + LOGLN(" warp the current image mask: " << ((getTickCount() - pt) / getTickFrequency()) << " sec"); +#if ENABLE_LOG + pt = getTickCount(); +#endif // Compensate exposure exposure_comp_->apply((int)img_idx, corners[img_idx], img_warped, mask_warped); + LOGLN(" compensate exposure: " << ((getTickCount() - pt) / getTickFrequency()) << " sec"); +#if ENABLE_LOG + pt = getTickCount(); +#endif img_warped.convertTo(img_warped_s, CV_16S); img_warped.release(); @@ -292,18 +318,35 @@ Stitcher::Status Stitcher::composePanorama(InputArrayOfArrays images, OutputArra bitwise_and(seam_mask, mask_warped, mask_warped); + LOGLN(" other: " << ((getTickCount() - pt) / getTickFrequency()) << " sec"); +#if ENABLE_LOG + pt = getTickCount(); +#endif + if (!is_blender_prepared) { blender_->prepare(corners, sizes); is_blender_prepared = true; } + LOGLN(" other2: " << ((getTickCount() - pt) / getTickFrequency()) << " sec"); + + LOGLN(" feed..."); +#if ENABLE_LOG + int64 feed_t = getTickCount(); +#endif // Blend the current image blender_->feed(img_warped_s, mask_warped, corners[img_idx]); + LOGLN(" feed time: " << ((getTickCount() - feed_t) / getTickFrequency()) << " sec"); + LOGLN("Compositing ## time: " << ((getTickCount() - compositing_t) / getTickFrequency()) << " sec"); } +#if ENABLE_LOG + int64 blend_t = getTickCount(); +#endif UMat result, result_mask; blender_->blend(result, result_mask); + LOGLN("blend time: " << ((getTickCount() - blend_t) / getTickFrequency()) << " sec"); LOGLN("Compositing, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); From 06738468af9b706e423a8ccbe98374f288f42c18 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Wed, 26 Feb 2014 19:02:36 +0400 Subject: [PATCH 10/15] TAPI: stiching: add custom OpenCL kernels for MultiBandBlender --- modules/core/include/opencv2/core/ocl.hpp | 2 + modules/core/include/opencv2/core/utility.hpp | 5 + modules/core/src/ocl.cpp | 19 +- modules/stitching/src/blenders.cpp | 197 +++++++----- .../stitching/src/opencl/multibandblend.cl | 282 ++++++++++++++++++ 5 files changed, 431 insertions(+), 74 deletions(-) create mode 100644 modules/stitching/src/opencl/multibandblend.cl diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index fdb6f9a0a..254cb1074 100644 --- a/modules/core/include/opencv2/core/ocl.hpp +++ b/modules/core/include/opencv2/core/ocl.hpp @@ -598,6 +598,8 @@ CV_EXPORTS int predictOptimalVectorWidth(InputArray src1, InputArray src2 = noAr InputArray src4 = noArray(), InputArray src5 = noArray(), InputArray src6 = noArray(), InputArray src7 = noArray(), InputArray src8 = noArray(), InputArray src9 = noArray()); +CV_EXPORTS void buildOptionsAddMatrixDescription(String& buildOptions, const String& name, InputArray _m); + class CV_EXPORTS Image2D { public: diff --git a/modules/core/include/opencv2/core/utility.hpp b/modules/core/include/opencv2/core/utility.hpp index 3e844ccf4..a8957f711 100644 --- a/modules/core/include/opencv2/core/utility.hpp +++ b/modules/core/include/opencv2/core/utility.hpp @@ -495,6 +495,11 @@ template<> inline std::string CommandLineParser::get(const String& } #endif // OPENCV_NOSTL +#if !defined(OPENCV_SKIP_SUPPRESS_WARNING) || !OPENCV_SKIP_SUPPRESS_WARNING +// Use this to bypass "warning C4127: conditional expression is constant" +template T SuppressWarning(T v) { return v; } +#endif + } //namespace cv #endif //__OPENCV_CORE_UTILITY_H__ diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 24190c52c..2017300fa 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -4404,7 +4404,24 @@ int predictOptimalVectorWidth(InputArray src1, InputArray src2, InputArray src3, #undef PROCESS_SRC -/////////////////////////////////////////// Image2D //////////////////////////////////////////////////// + +// TODO Make this as a method of OpenCL "BuildOptions" class +void buildOptionsAddMatrixDescription(String& buildOptions, const String& name, InputArray _m) +{ + if (!buildOptions.empty()) + buildOptions += " "; + int type = _m.type(), depth = CV_MAT_DEPTH(type); + buildOptions += format( + "-D %s_T=%s -D %s_T1=%s -D %s_CN=%d -D %s_TSIZE=%d -D %s_T1SIZE=%d -D %s_DEPTH=%d", + name.c_str(), ocl::typeToStr(type), + name.c_str(), ocl::typeToStr(CV_MAKE_TYPE(depth, 1)), + name.c_str(), (int)CV_MAT_CN(type), + name.c_str(), (int)CV_ELEM_SIZE(type), + name.c_str(), (int)CV_ELEM_SIZE1(type), + name.c_str(), (int)depth + ); +} + struct Image2D::Impl { diff --git a/modules/stitching/src/blenders.cpp b/modules/stitching/src/blenders.cpp index acb4987d7..a82da97ea 100644 --- a/modules/stitching/src/blenders.cpp +++ b/modules/stitching/src/blenders.cpp @@ -41,6 +41,7 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" namespace cv { namespace detail { @@ -245,6 +246,31 @@ void MultiBandBlender::prepare(Rect dst_roi) } } +#ifdef HAVE_OPENCL +static bool ocl_MultiBandBlender_feed(InputArray _src, InputArray _weight, + InputOutputArray _dst, InputOutputArray _dst_weight) +{ + String buildOptions = "-D DEFINE_feed"; + ocl::buildOptionsAddMatrixDescription(buildOptions, "src", _src); + ocl::buildOptionsAddMatrixDescription(buildOptions, "weight", _weight); + ocl::buildOptionsAddMatrixDescription(buildOptions, "dst", _dst); + ocl::buildOptionsAddMatrixDescription(buildOptions, "dstWeight", _dst_weight); + ocl::Kernel k("feed", ocl::stitching::multibandblend_oclsrc, buildOptions); + if (k.empty()) + return false; + + UMat src = _src.getUMat(); + + k.args(ocl::KernelArg::ReadOnly(src), + ocl::KernelArg::ReadOnly(_weight.getUMat()), + ocl::KernelArg::ReadWrite(_dst.getUMat()), + ocl::KernelArg::ReadWrite(_dst_weight.getUMat()) + ); + + size_t globalsize[2] = {src.cols, src.rows }; + return k.run(2, globalsize, NULL, false); +} +#endif void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) { @@ -338,63 +364,61 @@ void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) int x_br = br_new.x - dst_roi_.x; // Add weighted layer of the source image to the final Laplacian pyramid layer - if(weight_type_ == CV_32F) + for (int i = 0; i <= num_bands_; ++i) { - for (int i = 0; i <= num_bands_; ++i) + Rect rc(x_tl, y_tl, x_br - x_tl, y_br - y_tl); + CV_OPENCL_RUN(SuppressWarning(true), + ocl_MultiBandBlender_feed(src_pyr_laplace[i], weight_pyr_gauss[i], + dst_pyr_laplace_[i](rc), + dst_band_weights_[i](rc)), + goto next_band;) { Mat _src_pyr_laplace = src_pyr_laplace[i].getMat(ACCESS_READ); - Mat _dst_pyr_laplace = dst_pyr_laplace_[i].getMat(ACCESS_RW); + Mat _dst_pyr_laplace = dst_pyr_laplace_[i](rc).getMat(ACCESS_RW); Mat _weight_pyr_gauss = weight_pyr_gauss[i].getMat(ACCESS_READ); - Mat _dst_band_weights = dst_band_weights_[i].getMat(ACCESS_RW); - for (int y = y_tl; y < y_br; ++y) + Mat _dst_band_weights = dst_band_weights_[i](rc).getMat(ACCESS_RW); + if(weight_type_ == CV_32F) { - int y_ = y - y_tl; - const Point3_* src_row = _src_pyr_laplace.ptr >(y_); - Point3_* dst_row = _dst_pyr_laplace.ptr >(y); - const float* weight_row = _weight_pyr_gauss.ptr(y_); - float* dst_weight_row = _dst_band_weights.ptr(y); - - for (int x = x_tl; x < x_br; ++x) + for (int y = 0; y < rc.height; ++y) { - int x_ = x - x_tl; - dst_row[x].x += static_cast(src_row[x_].x * weight_row[x_]); - dst_row[x].y += static_cast(src_row[x_].y * weight_row[x_]); - dst_row[x].z += static_cast(src_row[x_].z * weight_row[x_]); - dst_weight_row[x] += weight_row[x_]; + const Point3_* src_row = _src_pyr_laplace.ptr >(y); + Point3_* dst_row = _dst_pyr_laplace.ptr >(y); + const float* weight_row = _weight_pyr_gauss.ptr(y); + float* dst_weight_row = _dst_band_weights.ptr(y); + + for (int x = 0; x < rc.width; ++x) + { + dst_row[x].x += static_cast(src_row[x].x * weight_row[x]); + dst_row[x].y += static_cast(src_row[x].y * weight_row[x]); + dst_row[x].z += static_cast(src_row[x].z * weight_row[x]); + dst_weight_row[x] += weight_row[x]; + } } } - x_tl /= 2; y_tl /= 2; - x_br /= 2; y_br /= 2; - } - } - else // weight_type_ == CV_16S - { - for (int i = 0; i <= num_bands_; ++i) - { - Mat _src_pyr_laplace = src_pyr_laplace[i].getMat(ACCESS_READ); - Mat _dst_pyr_laplace = dst_pyr_laplace_[i].getMat(ACCESS_RW); - Mat _weight_pyr_gauss = weight_pyr_gauss[i].getMat(ACCESS_READ); - Mat _dst_band_weights = dst_band_weights_[i].getMat(ACCESS_RW); - for (int y = y_tl; y < y_br; ++y) + else // weight_type_ == CV_16S { - int y_ = y - y_tl; - const Point3_* src_row = _src_pyr_laplace.ptr >(y_); - Point3_* dst_row = _dst_pyr_laplace.ptr >(y); - const short* weight_row = _weight_pyr_gauss.ptr(y_); - short* dst_weight_row = _dst_band_weights.ptr(y); - - for (int x = x_tl; x < x_br; ++x) + for (int y = 0; y < y_br - y_tl; ++y) { - int x_ = x - x_tl; - dst_row[x].x += short((src_row[x_].x * weight_row[x_]) >> 8); - dst_row[x].y += short((src_row[x_].y * weight_row[x_]) >> 8); - dst_row[x].z += short((src_row[x_].z * weight_row[x_]) >> 8); - dst_weight_row[x] += weight_row[x_]; + const Point3_* src_row = _src_pyr_laplace.ptr >(y); + Point3_* dst_row = _dst_pyr_laplace.ptr >(y); + const short* weight_row = _weight_pyr_gauss.ptr(y); + short* dst_weight_row = _dst_band_weights.ptr(y); + + for (int x = 0; x < x_br - x_tl; ++x) + { + dst_row[x].x += short((src_row[x].x * weight_row[x]) >> 8); + dst_row[x].y += short((src_row[x].y * weight_row[x]) >> 8); + dst_row[x].z += short((src_row[x].z * weight_row[x]) >> 8); + dst_weight_row[x] += weight_row[x]; + } } } - x_tl /= 2; y_tl /= 2; - x_br /= 2; y_br /= 2; } +#ifdef HAVE_OPENCL +next_band: +#endif + x_tl /= 2; y_tl /= 2; + x_br /= 2; y_br /= 2; } LOGLN(" Add weighted layer of the source image to the final Laplacian pyramid layer, time: " << ((getTickCount() - t) / getTickFrequency()) << " sec"); @@ -411,10 +435,10 @@ void MultiBandBlender::blend(InputOutputArray dst, InputOutputArray dst_mask) else restoreImageFromLaplacePyr(dst_pyr_laplace_); - dst_ = dst_pyr_laplace_[0]; - dst_ = dst_(Range(0, dst_roi_final_.height), Range(0, dst_roi_final_.width)); + Rect dst_rc(0, 0, dst_roi_final_.width, dst_roi_final_.height); + dst_ = dst_pyr_laplace_[0](dst_rc); UMat _dst_mask; - compare(dst_band_weights_[0](Range(0, dst_roi_final_.height), Range(0, dst_roi_final_.width)), WEIGHT_EPS, dst_mask_, CMP_GT); + compare(dst_band_weights_[0](dst_rc), WEIGHT_EPS, dst_mask_, CMP_GT); dst_pyr_laplace_.clear(); dst_band_weights_.clear(); @@ -425,47 +449,74 @@ void MultiBandBlender::blend(InputOutputArray dst, InputOutputArray dst_mask) ////////////////////////////////////////////////////////////////////////////// // Auxiliary functions +#ifdef HAVE_OPENCL +static bool ocl_normalizeUsingWeightMap(InputArray _weight, InputOutputArray _mat) +{ + String buildOptions = "-D DEFINE_normalizeUsingWeightMap"; + ocl::buildOptionsAddMatrixDescription(buildOptions, "mat", _mat); + ocl::buildOptionsAddMatrixDescription(buildOptions, "weight", _weight); + ocl::Kernel k("normalizeUsingWeightMap", ocl::stitching::multibandblend_oclsrc, buildOptions); + if (k.empty()) + return false; + + UMat mat = _mat.getUMat(); + + k.args(ocl::KernelArg::ReadWrite(mat), + ocl::KernelArg::ReadOnly(_weight.getUMat()) + ); + + size_t globalsize[2] = {mat.cols, mat.rows }; + return k.run(2, globalsize, NULL, false); +} +#endif + void normalizeUsingWeightMap(InputArray _weight, InputOutputArray _src) { #ifdef HAVE_TEGRA_OPTIMIZATION if(tegra::normalizeUsingWeightMap(weight, src)) return; #endif - Mat weight = _weight.getMat(); - Mat src = _src.getMat(); - CV_Assert(src.type() == CV_16SC3); - - if(weight.type() == CV_32FC1) + CV_OPENCL_RUN(SuppressWarning(true), + ocl_normalizeUsingWeightMap(_weight, _src), + return;) { - for (int y = 0; y < src.rows; ++y) - { - Point3_ *row = src.ptr >(y); - const float *weight_row = weight.ptr(y); + Mat weight = _weight.getMat(); + Mat src = _src.getMat(); - for (int x = 0; x < src.cols; ++x) + CV_Assert(src.type() == CV_16SC3); + + if(weight.type() == CV_32FC1) + { + for (int y = 0; y < src.rows; ++y) { - row[x].x = static_cast(row[x].x / (weight_row[x] + WEIGHT_EPS)); - row[x].y = static_cast(row[x].y / (weight_row[x] + WEIGHT_EPS)); - row[x].z = static_cast(row[x].z / (weight_row[x] + WEIGHT_EPS)); + Point3_ *row = src.ptr >(y); + const float *weight_row = weight.ptr(y); + + for (int x = 0; x < src.cols; ++x) + { + row[x].x = static_cast(row[x].x / (weight_row[x] + WEIGHT_EPS)); + row[x].y = static_cast(row[x].y / (weight_row[x] + WEIGHT_EPS)); + row[x].z = static_cast(row[x].z / (weight_row[x] + WEIGHT_EPS)); + } } } - } - else - { - CV_Assert(weight.type() == CV_16SC1); - - for (int y = 0; y < src.rows; ++y) + else { - const short *weight_row = weight.ptr(y); - Point3_ *row = src.ptr >(y); + CV_Assert(weight.type() == CV_16SC1); - for (int x = 0; x < src.cols; ++x) + for (int y = 0; y < src.rows; ++y) { - int w = weight_row[x] + 1; - row[x].x = static_cast((row[x].x << 8) / w); - row[x].y = static_cast((row[x].y << 8) / w); - row[x].z = static_cast((row[x].z << 8) / w); + const short *weight_row = weight.ptr(y); + Point3_ *row = src.ptr >(y); + + for (int x = 0; x < src.cols; ++x) + { + int w = weight_row[x] + 1; + row[x].x = static_cast((row[x].x << 8) / w); + row[x].y = static_cast((row[x].y << 8) / w); + row[x].z = static_cast((row[x].z << 8) / w); + } } } } diff --git a/modules/stitching/src/opencl/multibandblend.cl b/modules/stitching/src/opencl/multibandblend.cl new file mode 100644 index 000000000..d42ad82fb --- /dev/null +++ b/modules/stitching/src/opencl/multibandblend.cl @@ -0,0 +1,282 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. +// +// Copyright (C) 2014, Itseez, Inc, all rights reserved. + +// +// Common preprocessors macro +// + +// +// TODO: Move this common code into "header" file +// + +#ifndef NL // New Line: for preprocessor debugging +#define NL +#endif + +#define REF(x) x +#define __CAT(x, y) x##y +#define CAT(x, y) __CAT(x, y) + +// +// All matrixes are come with this description ("name" is a name of matrix): +// * name_CN - number of channels (1,2,3,4) +// * name_DEPTH - numeric value of CV_MAT_DEPTH(type). See CV_8U, CV_32S, etc macro below. +// +// Currently we also pass these attributes (to reduce this macro block): +// * name_T - datatype (int, float, uchar4, float4) +// * name_T1 - datatype for one channel (int, float, uchar). +// It is equal to result of "T1(name_T)" macro +// * name_TSIZE - CV_ELEM_SIZE(type). +// We can't use sizeof(name_T) here, because sizeof(float3) is usually equal to 8, not 6. +// * name_T1SIZE - CV_ELEM_SIZE1(type) +// + +// +// Usage sample: +// +// #define workType TYPE(float, src_CN) +// #define convertToWorkType CONVERT_TO(workType) +// #define convertWorkTypeToDstType CONVERT(workType, dst_T) +// +// __kernel void kernelFn(DECLARE_MAT_ARG(src), DECLARE_MAT_ARG(dst)) +// { +// const int x = get_global_id(0); +// const int y = get_global_id(1); +// +// if (x < srcWidth && y < srcHeight) +// { +// int src_byteOffset = MAT_BYTE_OFFSET(src, x, y); +// int dst_byteOffset = MAT_BYTE_OFFSET(dst, x, y); +// workType value = convertToWorkType(LOAD_MAT_AT(src, src_byteOffset)); +// +// ... value processing ... +// +// STORE_MAT_AT(dst, dst_byteOffset, convertWorkTypeToDstType(value)); +// } +// } +// + +#define DECLARE_MAT_ARG(name) \ + __global uchar* restrict name ## Ptr, \ + int name ## StepBytes, \ + int name ## Offset, \ + int name ## Height, \ + int name ## Width NL + +#define MAT_BYTE_OFFSET(name, x, y) mad24((y)/* + name ## OffsetY*/, name ## StepBytes, ((x)/* + name ## OffsetX*/) * (int)(name ## _TSIZE) + name ## Offset) +#define MAT_RELATIVE_BYTE_OFFSET(name, x, y) mad24(y, name ## StepBytes, (x) * (int)(name ## _TSIZE)) + +#define __LOAD_MAT_AT(name, byteOffset) *((const __global name ## _T*)(name ## Ptr + (byteOffset))) +#define __vload_CN__(name_cn) vload ## name_cn +#define __vload_CN_(name_cn) __vload_CN__(name_cn) +#define __vload_CN(name) __vload_CN_(name ## _CN) +#define __LOAD_MAT_AT_vload(name, byteOffset) __vload_CN(name)(0, ((const __global name ## _T1*)(name ## Ptr + (byteOffset)))) +#define __LOAD_MAT_AT_1 __LOAD_MAT_AT +#define __LOAD_MAT_AT_2 __LOAD_MAT_AT +#define __LOAD_MAT_AT_3 __LOAD_MAT_AT_vload +#define __LOAD_MAT_AT_4 __LOAD_MAT_AT +#define __LOAD_MAT_AT_CN__(name_cn) __LOAD_MAT_AT_ ## name_cn +#define __LOAD_MAT_AT_CN_(name_cn) __LOAD_MAT_AT_CN__(name_cn) +#define __LOAD_MAT_AT_CN(name) __LOAD_MAT_AT_CN_(name ## _CN) +#define LOAD_MAT_AT(name, byteOffset) __LOAD_MAT_AT_CN(name)(name, byteOffset) + +#define __STORE_MAT_AT(name, byteOffset, v) *((__global name ## _T*)(name ## Ptr + (byteOffset))) = v +#define __vstore_CN__(name_cn) vstore ## name_cn +#define __vstore_CN_(name_cn) __vstore_CN__(name_cn) +#define __vstore_CN(name) __vstore_CN_(name ## _CN) +#define __STORE_MAT_AT_vstore(name, byteOffset, v) __vstore_CN(name)(v, 0, ((__global name ## _T1*)(name ## Ptr + (byteOffset)))) +#define __STORE_MAT_AT_1 __STORE_MAT_AT +#define __STORE_MAT_AT_2 __STORE_MAT_AT +#define __STORE_MAT_AT_3 __STORE_MAT_AT_vstore +#define __STORE_MAT_AT_4 __STORE_MAT_AT +#define __STORE_MAT_AT_CN__(name_cn) __STORE_MAT_AT_ ## name_cn +#define __STORE_MAT_AT_CN_(name_cn) __STORE_MAT_AT_CN__(name_cn) +#define __STORE_MAT_AT_CN(name) __STORE_MAT_AT_CN_(name ## _CN) +#define STORE_MAT_AT(name, byteOffset, v) __STORE_MAT_AT_CN(name)(name, byteOffset, v) + +#define T1_uchar uchar +#define T1_uchar2 uchar +#define T1_uchar3 uchar +#define T1_uchar4 uchar +#define T1_char char +#define T1_char2 char +#define T1_char3 char +#define T1_char4 char +#define T1_ushort ushort +#define T1_ushort2 ushort +#define T1_ushort3 ushort +#define T1_ushort4 ushort +#define T1_short short +#define T1_short2 short +#define T1_short3 short +#define T1_short4 short +#define T1_int int +#define T1_int2 int +#define T1_int3 int +#define T1_int4 int +#define T1_float float +#define T1_float2 float +#define T1_float3 float +#define T1_float4 float +#define T1_double double +#define T1_double2 double +#define T1_double3 double +#define T1_double4 double +#define T1(type) REF(CAT(T1_, REF(type))) + +#define uchar1 uchar +#define char1 char +#define short1 short +#define ushort1 ushort +#define int1 int +#define float1 float +#define double1 double +#define TYPE(type, cn) REF(CAT(REF(type), REF(cn))) + +#define __CONVERT_MODE_uchar_uchar __NO_CONVERT +#define __CONVERT_MODE_uchar_char __CONVERT_sat +#define __CONVERT_MODE_uchar_ushort __CONVERT +#define __CONVERT_MODE_uchar_short __CONVERT +#define __CONVERT_MODE_uchar_int __CONVERT +#define __CONVERT_MODE_uchar_float __CONVERT +#define __CONVERT_MODE_uchar_double __CONVERT +#define __CONVERT_MODE_char_uchar __CONVERT_sat +#define __CONVERT_MODE_char_char __NO_CONVERT +#define __CONVERT_MODE_char_ushort __CONVERT_sat +#define __CONVERT_MODE_char_short __CONVERT +#define __CONVERT_MODE_char_int __CONVERT +#define __CONVERT_MODE_char_float __CONVERT +#define __CONVERT_MODE_char_double __CONVERT +#define __CONVERT_MODE_ushort_uchar __CONVERT_sat +#define __CONVERT_MODE_ushort_char __CONVERT_sat +#define __CONVERT_MODE_ushort_ushort __NO_CONVERT +#define __CONVERT_MODE_ushort_short __CONVERT_sat +#define __CONVERT_MODE_ushort_int __CONVERT +#define __CONVERT_MODE_ushort_float __CONVERT +#define __CONVERT_MODE_ushort_double __CONVERT +#define __CONVERT_MODE_short_uchar __CONVERT_sat +#define __CONVERT_MODE_short_char __CONVERT_sat +#define __CONVERT_MODE_short_ushort __CONVERT_sat +#define __CONVERT_MODE_short_short __NO_CONVERT +#define __CONVERT_MODE_short_int __CONVERT +#define __CONVERT_MODE_short_float __CONVERT +#define __CONVERT_MODE_short_double __CONVERT +#define __CONVERT_MODE_int_uchar __CONVERT_sat +#define __CONVERT_MODE_int_char __CONVERT_sat +#define __CONVERT_MODE_int_ushort __CONVERT_sat +#define __CONVERT_MODE_int_short __CONVERT_sat +#define __CONVERT_MODE_int_int __NO_CONVERT +#define __CONVERT_MODE_int_float __CONVERT +#define __CONVERT_MODE_int_double __CONVERT +#define __CONVERT_MODE_float_uchar __CONVERT_sat_rte +#define __CONVERT_MODE_float_char __CONVERT_sat_rte +#define __CONVERT_MODE_float_ushort __CONVERT_sat_rte +#define __CONVERT_MODE_float_short __CONVERT_sat_rte +#define __CONVERT_MODE_float_int __CONVERT_rte +#define __CONVERT_MODE_float_float __NO_CONVERT +#define __CONVERT_MODE_float_double __CONVERT +#define __CONVERT_MODE_double_uchar __CONVERT_sat_rte +#define __CONVERT_MODE_double_char __CONVERT_sat_rte +#define __CONVERT_MODE_double_ushort __CONVERT_sat_rte +#define __CONVERT_MODE_double_short __CONVERT_sat_rte +#define __CONVERT_MODE_double_int __CONVERT_rte +#define __CONVERT_MODE_double_float __CONVERT +#define __CONVERT_MODE_double_double __NO_CONVERT +#define __CONVERT_MODE(srcType, dstType) CAT(__CONVERT_MODE_, CAT(REF(T1(srcType)), CAT(_, REF(T1(dstType))))) + +#define __ROUND_MODE__NO_CONVERT +#define __ROUND_MODE__CONVERT // nothing +#define __ROUND_MODE__CONVERT_rte _rte +#define __ROUND_MODE__CONVERT_sat _sat +#define __ROUND_MODE__CONVERT_sat_rte _sat_rte +#define ROUND_MODE(srcType, dstType) CAT(__ROUND_MODE_, __CONVERT_MODE(srcType, dstType)) + +#define __CONVERT_ROUND(dstType, roundMode) CAT(CAT(convert_, REF(dstType)), roundMode) +#define __NO_CONVERT(dstType) // nothing +#define __CONVERT(dstType) __CONVERT_ROUND(dstType,) +#define __CONVERT_rte(dstType) __CONVERT_ROUND(dstType,_rte) +#define __CONVERT_sat(dstType) __CONVERT_ROUND(dstType,_sat) +#define __CONVERT_sat_rte(dstType) __CONVERT_ROUND(dstType,_sat_rte) +#define CONVERT(srcType, dstType) REF(__CONVERT_MODE(srcType,dstType))(dstType) +#define CONVERT_TO(dstType) __CONVERT_ROUND(dstType,) + +// OpenCV depths +#define CV_8U 0 +#define CV_8S 1 +#define CV_16U 2 +#define CV_16S 3 +#define CV_32S 4 +#define CV_32F 5 +#define CV_64F 6 + +// +// End of common preprocessors macro +// + + + +#if defined(DEFINE_feed) + +#define workType TYPE(weight_T1, src_CN) +#define convertSrcToWorkType CONVERT_TO(workType) +#define convertWorkTypeToDstType CONVERT(workType, dst_T) + +__kernel void feed( + DECLARE_MAT_ARG(src), DECLARE_MAT_ARG(weight), + DECLARE_MAT_ARG(dst), DECLARE_MAT_ARG(dstWeight) +) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + if (x < srcWidth && y < srcHeight) + { + int src_byteOffset = MAT_BYTE_OFFSET(src, x, y); + int weight_byteOffset = MAT_BYTE_OFFSET(weight, x, y); + int dst_byteOffset = MAT_BYTE_OFFSET(dst, x, y); + int dstWeight_byteOffset = MAT_BYTE_OFFSET(dstWeight, x, y); + + weight_T w = LOAD_MAT_AT(weight, weight_byteOffset); + workType src_value = convertSrcToWorkType(LOAD_MAT_AT(src, src_byteOffset)); + STORE_MAT_AT(dst, dst_byteOffset, LOAD_MAT_AT(dst, dst_byteOffset) + convertWorkTypeToDstType(src_value * w)); + STORE_MAT_AT(dstWeight, dstWeight_byteOffset, LOAD_MAT_AT(dstWeight, dstWeight_byteOffset) + w); + } +} + +#endif + +#if defined(DEFINE_normalizeUsingWeightMap) + +#define workType TYPE(weight_T1, mat_CN) +#define convertSrcToWorkType CONVERT_TO(workType) +#define convertWorkTypeToDstType CONVERT(workType, mat_T) + +#if weight_DEPTH >= CV_32F +#define WEIGHT_EPS 1e-5f +#else +#define WEIGHT_EPS 0 +#endif + +__kernel void normalizeUsingWeightMap( + DECLARE_MAT_ARG(mat), DECLARE_MAT_ARG(weight) +) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + if (x < matWidth && y < matHeight) + { + int mat_byteOffset = MAT_BYTE_OFFSET(mat, x, y); + int weight_byteOffset = MAT_BYTE_OFFSET(weight, x, y); + + weight_T w = LOAD_MAT_AT(weight, weight_byteOffset); + workType value = convertSrcToWorkType(LOAD_MAT_AT(mat, mat_byteOffset)); + value = value / (w + WEIGHT_EPS); + STORE_MAT_AT(mat, mat_byteOffset, convertWorkTypeToDstType(value)); + } +} + +#endif From e6cc1be7e8d6e9628b07f3d347dd32a36f0f7478 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Fri, 21 Feb 2014 17:58:33 +0400 Subject: [PATCH 11/15] stitching: allow to use dynamic DescriptorMatcher --- modules/core/src/matrix.cpp | 15 +++++++++++++++ modules/features2d/src/matchers.cpp | 6 +++--- modules/stitching/src/matchers.cpp | 28 +++++++++++++++++++--------- 3 files changed, 37 insertions(+), 12 deletions(-) diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index 32ccd0377..f7aded731 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -1369,6 +1369,21 @@ void _InputArray::getUMatVector(std::vector& umv) const return; } + if( k == UMAT ) + { + UMat& v = *(UMat*)obj; + umv.resize(1); + umv[0] = v; + return; + } + if( k == MAT ) + { + Mat& v = *(Mat*)obj; + umv.resize(1); + umv[0] = v.getUMat(accessFlags); + return; + } + CV_Error(Error::StsNotImplemented, "Unknown/unsupported array type"); } diff --git a/modules/features2d/src/matchers.cpp b/modules/features2d/src/matchers.cpp index 91ee9d9bf..2b5605031 100644 --- a/modules/features2d/src/matchers.cpp +++ b/modules/features2d/src/matchers.cpp @@ -330,7 +330,7 @@ static bool ocl_match2Dispatcher(InputArray query, InputArray train, const UMat static bool ocl_kmatchDispatcher(InputArray query, InputArray train, const UMat &trainIdx, const UMat &distance, int distType) { - return ocl_match2Dispatcher(query, train, trainIdx, distance, distType); + return ocl_match2Dispatcher(query, train, trainIdx, distance, distType); } static bool ocl_knnMatchSingle(InputArray query, InputArray train, UMat &trainIdx, @@ -1209,8 +1209,8 @@ FlannBasedMatcher::FlannBasedMatcher( const Ptr& _indexParam void FlannBasedMatcher::add( InputArrayOfArrays _descriptors ) { DescriptorMatcher::add( _descriptors ); - std::vector descriptors; - _descriptors.getMatVector(descriptors); + std::vector descriptors; + _descriptors.getUMatVector(descriptors); for( size_t i = 0; i < descriptors.size(); i++ ) { diff --git a/modules/stitching/src/matchers.cpp b/modules/stitching/src/matchers.cpp index ceb3d3da7..c303c4aab 100644 --- a/modules/stitching/src/matchers.cpp +++ b/modules/stitching/src/matchers.cpp @@ -155,21 +155,31 @@ void CpuMatcher::match(const ImageFeatures &features1, const ImageFeatures &feat matches_info.matches.clear(); - Ptr indexParams = makePtr(); - Ptr searchParams = makePtr(); - - if (features2.descriptors.depth() == CV_8U) + Ptr matcher; +#if 0 // TODO check this + if (ocl::useOpenCL()) { - indexParams->setAlgorithm(cvflann::FLANN_INDEX_LSH); - searchParams->setAlgorithm(cvflann::FLANN_INDEX_LSH); + matcher = makePtr((int)NORM_L2); } + else +#endif + { + Ptr indexParams = makePtr(); + Ptr searchParams = makePtr(); - FlannBasedMatcher matcher(indexParams, searchParams); + if (features2.descriptors.depth() == CV_8U) + { + indexParams->setAlgorithm(cvflann::FLANN_INDEX_LSH); + searchParams->setAlgorithm(cvflann::FLANN_INDEX_LSH); + } + + matcher = makePtr(indexParams, searchParams); + } std::vector< std::vector > pair_matches; MatchesSet matches; // Find 1->2 matches - matcher.knnMatch(features1.descriptors, features2.descriptors, pair_matches, 2); + matcher->knnMatch(features1.descriptors, features2.descriptors, pair_matches, 2); for (size_t i = 0; i < pair_matches.size(); ++i) { if (pair_matches[i].size() < 2) @@ -186,7 +196,7 @@ void CpuMatcher::match(const ImageFeatures &features1, const ImageFeatures &feat // Find 2->1 matches pair_matches.clear(); - matcher.knnMatch(features2.descriptors, features1.descriptors, pair_matches, 2); + matcher->knnMatch(features2.descriptors, features1.descriptors, pair_matches, 2); for (size_t i = 0; i < pair_matches.size(); ++i) { if (pair_matches[i].size() < 2) From 026b13b3dbe3d01778d74fd254dfa9f72adec9c3 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Wed, 12 Mar 2014 19:33:30 +0400 Subject: [PATCH 12/15] TAPI: stitching: blender: fix OpenCL path & adjust test condition --- modules/stitching/src/opencl/multibandblend.cl | 8 ++++---- modules/stitching/test/test_blenders.cpp | 4 ++-- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/modules/stitching/src/opencl/multibandblend.cl b/modules/stitching/src/opencl/multibandblend.cl index d42ad82fb..72d3de0fb 100644 --- a/modules/stitching/src/opencl/multibandblend.cl +++ b/modules/stitching/src/opencl/multibandblend.cl @@ -222,7 +222,7 @@ #define workType TYPE(weight_T1, src_CN) #define convertSrcToWorkType CONVERT_TO(workType) -#define convertWorkTypeToDstType CONVERT(workType, dst_T) +#define convertToDstType CONVERT_TO(dst_T) // sat_rte provides incompatible results with CPU path __kernel void feed( DECLARE_MAT_ARG(src), DECLARE_MAT_ARG(weight), @@ -241,7 +241,7 @@ __kernel void feed( weight_T w = LOAD_MAT_AT(weight, weight_byteOffset); workType src_value = convertSrcToWorkType(LOAD_MAT_AT(src, src_byteOffset)); - STORE_MAT_AT(dst, dst_byteOffset, LOAD_MAT_AT(dst, dst_byteOffset) + convertWorkTypeToDstType(src_value * w)); + STORE_MAT_AT(dst, dst_byteOffset, LOAD_MAT_AT(dst, dst_byteOffset) + convertToDstType(src_value * w)); STORE_MAT_AT(dstWeight, dstWeight_byteOffset, LOAD_MAT_AT(dstWeight, dstWeight_byteOffset) + w); } } @@ -252,7 +252,7 @@ __kernel void feed( #define workType TYPE(weight_T1, mat_CN) #define convertSrcToWorkType CONVERT_TO(workType) -#define convertWorkTypeToDstType CONVERT(workType, mat_T) +#define convertToDstType CONVERT_TO(mat_T) // sat_rte provides incompatible results with CPU path #if weight_DEPTH >= CV_32F #define WEIGHT_EPS 1e-5f @@ -275,7 +275,7 @@ __kernel void normalizeUsingWeightMap( weight_T w = LOAD_MAT_AT(weight, weight_byteOffset); workType value = convertSrcToWorkType(LOAD_MAT_AT(mat, mat_byteOffset)); value = value / (w + WEIGHT_EPS); - STORE_MAT_AT(mat, mat_byteOffset, convertWorkTypeToDstType(value)); + STORE_MAT_AT(mat, mat_byteOffset, convertToDstType(value)); } } diff --git a/modules/stitching/test/test_blenders.cpp b/modules/stitching/test/test_blenders.cpp index cb84482f2..f7c7835de 100644 --- a/modules/stitching/test/test_blenders.cpp +++ b/modules/stitching/test/test_blenders.cpp @@ -73,6 +73,6 @@ TEST(MultiBandBlender, CanBlendTwoImages) Mat result; result_s.convertTo(result, CV_8U); Mat expected = imread(string(cvtest::TS::ptr()->get_data_path()) + "stitching/baboon_lena.png"); - double rmsErr = cvtest::norm(expected, result, NORM_L2) / sqrt(double(expected.size().area())); - ASSERT_LT(rmsErr, 1e-3); + double psnr = cvtest::PSNR(expected, result); + EXPECT_GE(psnr, 50); } From 588658d1a3f16ef9c61fa71b5d2ef0630030dbf4 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Mon, 17 Mar 2014 18:54:53 +0400 Subject: [PATCH 13/15] tapi: stitching: fix documentation --- modules/stitching/doc/blenders.rst | 4 ++-- modules/stitching/doc/exposure_compensation.rst | 6 +++--- modules/stitching/doc/high_level.rst | 10 +++++----- modules/stitching/doc/matching.rst | 8 ++++---- modules/stitching/doc/seam_estimation.rst | 2 +- 5 files changed, 15 insertions(+), 15 deletions(-) diff --git a/modules/stitching/doc/blenders.rst b/modules/stitching/doc/blenders.rst index 031a32664..39d9571c3 100644 --- a/modules/stitching/doc/blenders.rst +++ b/modules/stitching/doc/blenders.rst @@ -43,7 +43,7 @@ detail::Blender::feed Processes the image. -.. ocv:function:: void detail::Blender::feed(const Mat &img, const Mat &mask, Point tl) +.. ocv:function:: void detail::Blender::feed(InputArray img, InputArray mask, Point tl) :param img: Source image @@ -56,7 +56,7 @@ detail::Blender::blend Blends and returns the final pano. -.. ocv:function:: void detail::Blender::blend(Mat &dst, Mat &dst_mask) +.. ocv:function:: void detail::Blender::blend(InputOutputArray dst, InputOutputArray dst_mask) :param dst: Final pano diff --git a/modules/stitching/doc/exposure_compensation.rst b/modules/stitching/doc/exposure_compensation.rst index ec0d5db32..b2778f124 100644 --- a/modules/stitching/doc/exposure_compensation.rst +++ b/modules/stitching/doc/exposure_compensation.rst @@ -27,9 +27,9 @@ Base class for all exposure compensators. :: detail::ExposureCompensator::feed ---------------------------------- -.. ocv:function:: void detail::ExposureCompensator::feed(const std::vector &corners, const std::vector &images, const std::vector &masks) +.. ocv:function:: void detail::ExposureCompensator::feed(const std::vector &corners, const std::vector &images, const std::vector &masks) -.. ocv:function:: void detail::ExposureCompensator::feed(const std::vector &corners, const std::vector &images, const std::vector > &masks) +.. ocv:function:: void detail::ExposureCompensator::feed(const std::vector &corners, const std::vector &images, const std::vector > &masks) :param corners: Source image top-left corners @@ -42,7 +42,7 @@ detil::ExposureCompensator::apply Compensate exposure in the specified image. -.. ocv:function:: void detail::ExposureCompensator::apply(int index, Point corner, Mat &image, const Mat &mask) +.. ocv:function:: void detail::ExposureCompensator::apply(int index, Point corner, InputOutputArray image, InputArray mask) :param index: Image index diff --git a/modules/stitching/doc/high_level.rst b/modules/stitching/doc/high_level.rst index 955f8f248..4cc5f9914 100644 --- a/modules/stitching/doc/high_level.rst +++ b/modules/stitching/doc/high_level.rst @@ -110,9 +110,9 @@ These functions try to match the given images and to estimate rotations of each .. note:: Use the functions only if you're aware of the stitching pipeline, otherwise use :ocv:func:`Stitcher::stitch`. -.. ocv:function:: Status Stitcher::estimateTransform(InputArray images) +.. ocv:function:: Status Stitcher::estimateTransform(InputArrayOfArrays images) -.. ocv:function:: Status Stitcher::estimateTransform(InputArray images, const std::vector > &rois) +.. ocv:function:: Status Stitcher::estimateTransform(InputArrayOfArrays images, const std::vector > &rois) :param images: Input images. @@ -129,7 +129,7 @@ These functions try to compose the given images (or images stored internally fro .. ocv:function:: Status Stitcher::composePanorama(OutputArray pano) -.. ocv:function:: Status Stitcher::composePanorama(InputArray images, OutputArray pano) +.. ocv:function:: Status Stitcher::composePanorama(InputArrayOfArrays images, OutputArray pano) :param images: Input images. @@ -142,9 +142,9 @@ Stitcher::stitch These functions try to stitch the given images. -.. ocv:function:: Status Stitcher::stitch(InputArray images, OutputArray pano) +.. ocv:function:: Status Stitcher::stitch(InputArrayOfArrays images, OutputArray pano) -.. ocv:function:: Status Stitcher::stitch(InputArray images, const std::vector > &rois, OutputArray pano) +.. ocv:function:: Status Stitcher::stitch(InputArrayOfArrays images, const std::vector > &rois, OutputArray pano) :param images: Input images. diff --git a/modules/stitching/doc/matching.rst b/modules/stitching/doc/matching.rst index 2f9dabf72..9f112d0a0 100644 --- a/modules/stitching/doc/matching.rst +++ b/modules/stitching/doc/matching.rst @@ -40,9 +40,9 @@ detail::FeaturesFinder::operator() Finds features in the given image. -.. ocv:function:: void detail::FeaturesFinder::operator ()(const Mat &image, ImageFeatures &features) +.. ocv:function:: void detail::FeaturesFinder::operator ()(InputArray image, ImageFeatures &features) -.. ocv:function:: void detail::FeaturesFinder::operator ()(const Mat &image, ImageFeatures &features, const std::vector &rois) +.. ocv:function:: void detail::FeaturesFinder::operator ()(InputArray image, ImageFeatures &features, const std::vector &rois) :param image: Source image @@ -64,7 +64,7 @@ detail::FeaturesFinder::find This method must implement features finding logic in order to make the wrappers `detail::FeaturesFinder::operator()`_ work. -.. ocv:function:: void detail::FeaturesFinder::find(const Mat &image, ImageFeatures &features) +.. ocv:function:: void detail::FeaturesFinder::find(InputArray image, ImageFeatures &features) :param image: Source image @@ -171,7 +171,7 @@ Performs images matching. :param matches_info: Found matches -.. ocv:function:: void detail::FeaturesMatcher::operator ()( const std::vector & features, std::vector & pairwise_matches, const Mat & mask=Mat() ) +.. ocv:function:: void detail::FeaturesMatcher::operator ()( const std::vector & features, std::vector & pairwise_matches, const UMat & mask=UMat() ) :param features: Features of the source images diff --git a/modules/stitching/doc/seam_estimation.rst b/modules/stitching/doc/seam_estimation.rst index 119fb1e15..e7a3e47b2 100644 --- a/modules/stitching/doc/seam_estimation.rst +++ b/modules/stitching/doc/seam_estimation.rst @@ -22,7 +22,7 @@ detail::SeamFinder::find Estimates seams. -.. ocv:function:: void detail::SeamFinder::find(const std::vector &src, const std::vector &corners, std::vector &masks) +.. ocv:function:: void detail::SeamFinder::find(const std::vector &src, const std::vector &corners, std::vector &masks) :param src: Source images From 3578f0afe95b9e182f86476629f333a9beaf6517 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Thu, 13 Mar 2014 16:39:54 +0400 Subject: [PATCH 14/15] ocl: workaround for synchronization issue in ::map() --- modules/core/src/ocl.cpp | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 2017300fa..9b4cf7260 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -3798,11 +3798,16 @@ public: cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr(); - if( u->refcount == 0 ) + // FIXIT Workaround for UMat synchronization issue + // if( u->refcount == 0 ) { if( !u->copyOnMap() ) { - CV_Assert(u->data == 0); + if (u->data) // FIXIT Workaround for UMat synchronization issue + { + //CV_Assert(u->hostCopyObsolete() == false); + return; + } // because there can be other map requests for the same UMat with different access flags, // we use the universal (read-write) access mode. cl_int retval = 0; @@ -3844,6 +3849,10 @@ public: UMatDataAutoLock autolock(u); + // FIXIT Workaround for UMat synchronization issue + if(u->refcount > 0) + return; + cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr(); cl_int retval = 0; if( !u->copyOnMap() && u->data ) From 925178749acef34d40b49aa5178c0015b3f25215 Mon Sep 17 00:00:00 2001 From: Andrey Pavlenko Date: Thu, 3 Apr 2014 11:26:25 +0400 Subject: [PATCH 15/15] removing `SuppressWarning` and `CV_OPENCL_RUN` --- modules/core/include/opencv2/core/utility.hpp | 5 ----- modules/stitching/src/blenders.cpp | 21 +++++++++---------- 2 files changed, 10 insertions(+), 16 deletions(-) diff --git a/modules/core/include/opencv2/core/utility.hpp b/modules/core/include/opencv2/core/utility.hpp index a8957f711..3e844ccf4 100644 --- a/modules/core/include/opencv2/core/utility.hpp +++ b/modules/core/include/opencv2/core/utility.hpp @@ -495,11 +495,6 @@ template<> inline std::string CommandLineParser::get(const String& } #endif // OPENCV_NOSTL -#if !defined(OPENCV_SKIP_SUPPRESS_WARNING) || !OPENCV_SKIP_SUPPRESS_WARNING -// Use this to bypass "warning C4127: conditional expression is constant" -template T SuppressWarning(T v) { return v; } -#endif - } //namespace cv #endif //__OPENCV_CORE_UTILITY_H__ diff --git a/modules/stitching/src/blenders.cpp b/modules/stitching/src/blenders.cpp index a82da97ea..03aad752f 100644 --- a/modules/stitching/src/blenders.cpp +++ b/modules/stitching/src/blenders.cpp @@ -367,11 +367,11 @@ void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) for (int i = 0; i <= num_bands_; ++i) { Rect rc(x_tl, y_tl, x_br - x_tl, y_br - y_tl); - CV_OPENCL_RUN(SuppressWarning(true), - ocl_MultiBandBlender_feed(src_pyr_laplace[i], weight_pyr_gauss[i], - dst_pyr_laplace_[i](rc), - dst_band_weights_[i](rc)), - goto next_band;) +#ifdef HAVE_OPENCL + if ( !cv::ocl::useOpenCL() || + !ocl_MultiBandBlender_feed(src_pyr_laplace[i], weight_pyr_gauss[i], + dst_pyr_laplace_[i](rc), dst_band_weights_[i](rc)) ) +#endif { Mat _src_pyr_laplace = src_pyr_laplace[i].getMat(ACCESS_READ); Mat _dst_pyr_laplace = dst_pyr_laplace_[i](rc).getMat(ACCESS_RW); @@ -414,9 +414,7 @@ void MultiBandBlender::feed(InputArray _img, InputArray mask, Point tl) } } } -#ifdef HAVE_OPENCL -next_band: -#endif + x_tl /= 2; y_tl /= 2; x_br /= 2; y_br /= 2; } @@ -477,9 +475,10 @@ void normalizeUsingWeightMap(InputArray _weight, InputOutputArray _src) return; #endif - CV_OPENCL_RUN(SuppressWarning(true), - ocl_normalizeUsingWeightMap(_weight, _src), - return;) +#ifdef HAVE_OPENCL + if ( !cv::ocl::useOpenCL() || + !ocl_normalizeUsingWeightMap(_weight, _src) ) +#endif { Mat weight = _weight.getMat(); Mat src = _src.getMat();