Merge pull request #2112 from ilya-lavrenov:tapi_stitching
commit
d64f05cd59
10 changed files with 693 additions and 117 deletions
@ -0,0 +1,148 @@ |
||||
/*M/////////////////////////////////////////////////////////////////////////////////////// |
||||
// |
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. |
||||
// |
||||
// By downloading, copying, installing or using the software you agree to this license. |
||||
// If you do not agree to this license, do not download, install, |
||||
// copy or use the software. |
||||
// |
||||
// |
||||
// License Agreement |
||||
// For Open Source Computer Vision Library |
||||
// |
||||
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. |
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. |
||||
// Third party copyrights are property of their respective owners. |
||||
// |
||||
// @Authors |
||||
// Peng Xiao, pengxiao@multicorewareinc.com |
||||
// |
||||
// Redistribution and use in source and binary forms, with or without modification, |
||||
// are permitted provided that the following conditions are met: |
||||
// |
||||
// * Redistribution's of source code must retain the above copyright notice, |
||||
// this list of conditions and the following disclaimer. |
||||
// |
||||
// * Redistribution's in binary form must reproduce the above copyright notice, |
||||
// this list of conditions and the following disclaimer in the documentation |
||||
// and/or other materials provided with the distribution. |
||||
// |
||||
// * The name of the copyright holders may not be used to endorse or promote products |
||||
// derived from this software without specific prior written permission. |
||||
// |
||||
// This software is provided by the copyright holders and contributors as is and |
||||
// any express or implied warranties, including, but not limited to, the implied |
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed. |
||||
// In no event shall the Intel Corporation or contributors be liable for any direct, |
||||
// indirect, incidental, special, exemplary, or consequential damages |
||||
// (including, but not limited to, procurement of substitute goods or services; |
||||
// loss of use, data, or profits; or business interruption) however caused |
||||
// and on any theory of liability, whether in contract, strict liability, |
||||
// or tort (including negligence or otherwise) arising in any way out of |
||||
// the use of this software, even if advised of the possibility of such damage. |
||||
// |
||||
//M*/ |
||||
|
||||
__kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset, |
||||
__global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols, |
||||
__constant float * ck_rinv, __constant float * ct, |
||||
int tl_u, int tl_v, float scale) |
||||
{ |
||||
int du = get_global_id(0); |
||||
int dv = get_global_id(1); |
||||
|
||||
if (du < cols && dv < rows) |
||||
{ |
||||
__global float * xmap = (__global float *)(xmapptr + mad24(dv, xmap_step, xmap_offset + du * (int)sizeof(float))); |
||||
__global float * ymap = (__global float *)(ymapptr + mad24(dv, ymap_step, ymap_offset + du * (int)sizeof(float))); |
||||
|
||||
float u = tl_u + du; |
||||
float v = tl_v + dv; |
||||
float x, y; |
||||
|
||||
float x_ = u / scale - ct[0]; |
||||
float y_ = v / scale - ct[1]; |
||||
|
||||
float z; |
||||
x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * (1 - ct[2]); |
||||
y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * (1 - ct[2]); |
||||
z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * (1 - ct[2]); |
||||
|
||||
x /= z; |
||||
y /= z; |
||||
|
||||
xmap[0] = x; |
||||
ymap[0] = y; |
||||
} |
||||
} |
||||
|
||||
__kernel void buildWarpCylindricalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset, |
||||
__global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols, |
||||
__constant float * ck_rinv, int tl_u, int tl_v, float scale) |
||||
{ |
||||
int du = get_global_id(0); |
||||
int dv = get_global_id(1); |
||||
|
||||
if (du < cols && dv < rows) |
||||
{ |
||||
__global float * xmap = (__global float *)(xmapptr + mad24(dv, xmap_step, xmap_offset + du * (int)sizeof(float))); |
||||
__global float * ymap = (__global float *)(ymapptr + mad24(dv, ymap_step, ymap_offset + du * (int)sizeof(float))); |
||||
|
||||
float u = tl_u + du; |
||||
float v = tl_v + dv; |
||||
float x, y; |
||||
|
||||
u /= scale; |
||||
float x_ = sin(u); |
||||
float y_ = v / scale; |
||||
float z_ = cos(u); |
||||
|
||||
float z; |
||||
x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; |
||||
y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; |
||||
z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; |
||||
|
||||
if (z > 0) x /= z, y /= z; |
||||
else x = y = -1; |
||||
|
||||
xmap[0] = x; |
||||
ymap[0] = y; |
||||
} |
||||
} |
||||
|
||||
__kernel void buildWarpSphericalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset, |
||||
__global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols, |
||||
__constant float * ck_rinv, int tl_u, int tl_v, float scale) |
||||
{ |
||||
int du = get_global_id(0); |
||||
int dv = get_global_id(1); |
||||
|
||||
if (du < cols && dv < rows) |
||||
{ |
||||
__global float * xmap = (__global float *)(xmapptr + mad24(dv, xmap_step, xmap_offset + du * (int)sizeof(float))); |
||||
__global float * ymap = (__global float *)(ymapptr + mad24(dv, ymap_step, ymap_offset + du * (int)sizeof(float))); |
||||
|
||||
float u = tl_u + du; |
||||
float v = tl_v + dv; |
||||
float x, y; |
||||
|
||||
v /= scale; |
||||
u /= scale; |
||||
|
||||
float sinv = sin(v); |
||||
float x_ = sinv * sin(u); |
||||
float y_ = -cos(v); |
||||
float z_ = sinv * cos(u); |
||||
|
||||
float z; |
||||
x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; |
||||
y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; |
||||
z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; |
||||
|
||||
if (z > 0) x /= z, y /= z; |
||||
else x = y = -1; |
||||
|
||||
xmap[0] = x; |
||||
ymap[0] = y; |
||||
} |
||||
} |
@ -0,0 +1,187 @@ |
||||
/*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); |
||||
|
||||
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 r_kinv(1, 9, CV_32FC1, projector_.r_kinv), t(1, 3, CV_32FC1, projector_.t); |
||||
UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), |
||||
ur_kinv = r_kinv.getUMat(ACCESS_READ), ut = t.getUMat(ACCESS_READ); |
||||
|
||||
k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), |
||||
ocl::KernelArg::PtrReadOnly(ur_kinv), ocl::KernelArg::PtrReadOnly(ut), |
||||
dst_tl.x, dst_tl.y, projector_.scale); |
||||
|
||||
size_t globalsize[2] = { 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 r_kinv(1, 9, CV_32FC1, projector_.r_kinv); |
||||
UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), ur_kinv = r_kinv.getUMat(ACCESS_READ); |
||||
|
||||
k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), |
||||
ocl::KernelArg::PtrReadOnly(ur_kinv), dst_tl.x, dst_tl.y, projector_.scale); |
||||
|
||||
size_t globalsize[2] = { 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 r_kinv(1, 9, CV_32FC1, projector_.r_kinv); |
||||
UMat uxmap = xmap.getUMat(), uymap = ymap.getUMat(), ur_kinv = r_kinv.getUMat(ACCESS_READ); |
||||
|
||||
k.args(ocl::KernelArg::WriteOnlyNoSize(uxmap), ocl::KernelArg::WriteOnly(uymap), |
||||
ocl::KernelArg::PtrReadOnly(ur_kinv), dst_tl.x, dst_tl.y, projector_.scale); |
||||
|
||||
size_t globalsize[2] = { 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
|
@ -0,0 +1,149 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other materials provided with the distribution.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors "as is" and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the OpenCV Foundation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#include "test_precomp.hpp" |
||||
#include "opencv2/ts/ocl_test.hpp" |
||||
#include "opencv2/stitching/warpers.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
///////////////////////// WarperTestBase ///////////////////////////
|
||||
|
||||
struct WarperTestBase : |
||||
public Test, public TestUtils |
||||
{ |
||||
Mat src, dst, xmap, ymap; |
||||
Mat udst, uxmap, uymap; |
||||
Mat K, R; |
||||
|
||||
virtual void generateTestData() |
||||
{ |
||||
Size size = randomSize(1, MAX_VALUE); |
||||
|
||||
src = randomMat(size, CV_32FC1, -500, 500); |
||||
|
||||
K = Mat::eye(3, 3, CV_32FC1); |
||||
R = Mat::eye(3, 3, CV_32FC1); |
||||
} |
||||
|
||||
void Near(double threshold = 0.) |
||||
{ |
||||
EXPECT_MAT_NEAR(xmap, uxmap, threshold); |
||||
EXPECT_MAT_NEAR(ymap, uymap, threshold); |
||||
EXPECT_MAT_NEAR(dst, udst, threshold); |
||||
} |
||||
}; |
||||
|
||||
//////////////////////////////// SphericalWarperOcl /////////////////////////////////////////////////
|
||||
|
||||
typedef WarperTestBase SphericalWarperOclTest; |
||||
|
||||
OCL_TEST_F(SphericalWarperOclTest, Mat) |
||||
{ |
||||
for (int j = 0; j < test_loop_times; j++) |
||||
{ |
||||
generateTestData(); |
||||
|
||||
Ptr<WarperCreator> creator = makePtr<SphericalWarperOcl>(); |
||||
Ptr<detail::RotationWarper> warper = creator->create(2.0); |
||||
|
||||
OCL_OFF(warper->buildMaps(src.size(), K, R, xmap, ymap)); |
||||
OCL_ON(warper->buildMaps(src.size(), K, R, uxmap, uymap)); |
||||
|
||||
OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); |
||||
OCL_ON(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); |
||||
|
||||
Near(1e-4); |
||||
} |
||||
} |
||||
|
||||
//////////////////////////////// CylindricalWarperOcl /////////////////////////////////////////////////
|
||||
|
||||
typedef WarperTestBase CylindricalWarperOclTest; |
||||
|
||||
OCL_TEST_F(CylindricalWarperOclTest, Mat) |
||||
{ |
||||
for (int j = 0; j < test_loop_times; j++) |
||||
{ |
||||
generateTestData(); |
||||
|
||||
Ptr<WarperCreator> creator = makePtr<CylindricalWarperOcl>(); |
||||
Ptr<detail::RotationWarper> warper = creator->create(2.0); |
||||
|
||||
OCL_OFF(warper->buildMaps(src.size(), K, R, xmap, ymap)); |
||||
OCL_ON(warper->buildMaps(src.size(), K, R, uxmap, uymap)); |
||||
|
||||
OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); |
||||
OCL_ON(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); |
||||
|
||||
Near(1e-4); |
||||
} |
||||
} |
||||
|
||||
//////////////////////////////// PlaneWarperOcl /////////////////////////////////////////////////
|
||||
|
||||
typedef WarperTestBase PlaneWarperOclTest; |
||||
|
||||
OCL_TEST_F(PlaneWarperOclTest, Mat) |
||||
{ |
||||
for (int j = 0; j < test_loop_times; j++) |
||||
{ |
||||
generateTestData(); |
||||
|
||||
Ptr<WarperCreator> creator = makePtr<PlaneWarperOcl>(); |
||||
Ptr<detail::RotationWarper> warper = creator->create(2.0); |
||||
|
||||
OCL_OFF(warper->buildMaps(src.size(), K, R, xmap, ymap)); |
||||
OCL_ON(warper->buildMaps(src.size(), K, R, uxmap, uymap)); |
||||
|
||||
OCL_OFF(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, dst)); |
||||
OCL_ON(warper->warp(src, K, R, INTER_LINEAR, BORDER_REPLICATE, udst)); |
||||
|
||||
Near(1e-5); |
||||
} |
||||
} |
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
Loading…
Reference in new issue