Merge pull request #2565 from apavlenko:tapi_stitching
Andrey Pavlenko
11 years ago
committed by
OpenCV Buildbot
commit
dab8e920b0
35 changed files with 1161 additions and 624 deletions
@ -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<string> 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<string>("orb") |
||||
#endif |
||||
|
||||
OCL_PERF_TEST_P(stitch, a123, TEST_DETECTORS) |
||||
{ |
||||
UMat pano; |
||||
|
||||
vector<Mat> _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<UMat> imgs = ToUMat(_imgs); |
||||
|
||||
Ptr<detail::FeaturesFinder> featuresFinder = GetParam() == "orb" |
||||
? Ptr<detail::FeaturesFinder>(new detail::OrbFeaturesFinder()) |
||||
: Ptr<detail::FeaturesFinder>(new detail::SurfFeaturesFinder()); |
||||
|
||||
Ptr<detail::FeaturesMatcher> featuresMatcher = GetParam() == "orb" |
||||
? makePtr<detail::BestOf2NearestMatcher>(false, ORB_MATCH_CONFIDENCE) |
||||
: makePtr<detail::BestOf2NearestMatcher>(false, SURF_MATCH_CONFIDENCE); |
||||
|
||||
declare.iterations(20); |
||||
|
||||
while(next()) |
||||
{ |
||||
Stitcher stitcher = Stitcher::createDefault(); |
||||
stitcher.setFeaturesFinder(featuresFinder); |
||||
stitcher.setFeaturesMatcher(featuresMatcher); |
||||
stitcher.setWarper(makePtr<SphericalWarper>()); |
||||
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<Mat> imgs; |
||||
imgs.push_back( imread( getDataPath("stitching/b1.png") ) ); |
||||
imgs.push_back( imread( getDataPath("stitching/b2.png") ) ); |
||||
|
||||
Ptr<detail::FeaturesFinder> featuresFinder = GetParam() == "orb" |
||||
? Ptr<detail::FeaturesFinder>(new detail::OrbFeaturesFinder()) |
||||
: Ptr<detail::FeaturesFinder>(new detail::SurfFeaturesFinder()); |
||||
|
||||
Ptr<detail::FeaturesMatcher> featuresMatcher = GetParam() == "orb" |
||||
? makePtr<detail::BestOf2NearestMatcher>(false, ORB_MATCH_CONFIDENCE) |
||||
: makePtr<detail::BestOf2NearestMatcher>(false, SURF_MATCH_CONFIDENCE); |
||||
|
||||
declare.iterations(20); |
||||
|
||||
while(next()) |
||||
{ |
||||
Stitcher stitcher = Stitcher::createDefault(); |
||||
stitcher.setFeaturesFinder(featuresFinder); |
||||
stitcher.setFeaturesMatcher(featuresMatcher); |
||||
stitcher.setWarper(makePtr<SphericalWarper>()); |
||||
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<Mat> _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<UMat> imgs = ToUMat(_imgs); |
||||
|
||||
Ptr<detail::FeaturesFinder> featuresFinder = GetParam() == "orb" |
||||
? Ptr<detail::FeaturesFinder>(new detail::OrbFeaturesFinder()) |
||||
: Ptr<detail::FeaturesFinder>(new detail::SurfFeaturesFinder()); |
||||
|
||||
Ptr<detail::FeaturesMatcher> featuresMatcher = GetParam() == "orb" |
||||
? makePtr<detail::BestOf2NearestMatcher>(false, ORB_MATCH_CONFIDENCE) |
||||
: makePtr<detail::BestOf2NearestMatcher>(false, SURF_MATCH_CONFIDENCE); |
||||
|
||||
declare.iterations(20); |
||||
|
||||
while(next()) |
||||
{ |
||||
Stitcher stitcher = Stitcher::createDefault(); |
||||
stitcher.setFeaturesFinder(featuresFinder); |
||||
stitcher.setFeaturesMatcher(featuresMatcher); |
||||
stitcher.setWarper(makePtr<SphericalWarper>()); |
||||
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(); |
||||
} |
||||
@ -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 convertToDstType CONVERT_TO(dst_T) // sat_rte provides incompatible results with CPU path |
||||
|
||||
__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) + convertToDstType(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 convertToDstType CONVERT_TO(mat_T) // sat_rte provides incompatible results with CPU path |
||||
|
||||
#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, convertToDstType(value)); |
||||
} |
||||
} |
||||
|
||||
#endif |
||||
@ -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
|
||||
Loading…
Reference in new issue