mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-29 05:29:54 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Merge https://github.com/Itseez/opencv

Browse Source
This commit is contained in:
Nghia Ho 2013-08-14 22:38:39 +10:00
commit 2c9c4188ae
24 changed files with 2398 additions and 3160 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

233
doc/tutorials/ml/non_linear_svms/non_linear_svms.rst
View File

File diff suppressed because one or more lines are too long

4
modules/core/CMakeLists.txt
View File

@ -1,5 +1,5 @@
set(the_description "The Core Functionality")
ocv_add_module(core ${ZLIB_LIBRARIES})
ocv_add_module(core ${ZLIB_LIBRARIES} OPTIONAL opencv_cudev)
ocv_module_include_directories(${ZLIB_INCLUDE_DIR})
if (HAVE_WINRT)
@ -7,7 +7,7 @@ if (HAVE_WINRT)
endif()
if(HAVE_CUDA)
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef)
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wenum-compare -Wunused-function)
endif()
file(GLOB lib_cuda_hdrs "include/opencv2/${name}/cuda/*.hpp" "include/opencv2/${name}/cuda/*.h")

5
modules/core/include/opencv2/core/base.hpp
View File

@ -498,6 +498,11 @@ namespace gpu
class CV_EXPORTS Event;
}
namespace cudev
{
template <typename _Tp> class GpuMat_;
}
} // cv
#endif //__OPENCV_CORE_BASE_HPP__

3
modules/core/include/opencv2/core/mat.hpp
View File

@ -96,6 +96,7 @@ public:
_InputArray(const gpu::GpuMat& d_mat);
_InputArray(const ogl::Buffer& buf);
_InputArray(const gpu::CudaMem& cuda_mem);
template<typename _Tp> _InputArray(const cudev::GpuMat_<_Tp>& m);
virtual Mat getMat(int i=-1) const;
virtual void getMatVector(std::vector<Mat>& mv) const;
@ -144,6 +145,7 @@ public:
_OutputArray(gpu::GpuMat& d_mat);
_OutputArray(ogl::Buffer& buf);
_OutputArray(gpu::CudaMem& cuda_mem);
template<typename _Tp> _OutputArray(cudev::GpuMat_<_Tp>& m);
template<typename _Tp> _OutputArray(std::vector<_Tp>& vec);
template<typename _Tp> _OutputArray(std::vector<std::vector<_Tp> >& vec);
template<typename _Tp> _OutputArray(std::vector<Mat_<_Tp> >& vec);
@ -156,6 +158,7 @@ public:
_OutputArray(const gpu::GpuMat& d_mat);
_OutputArray(const ogl::Buffer& buf);
_OutputArray(const gpu::CudaMem& cuda_mem);
template<typename _Tp> _OutputArray(const cudev::GpuMat_<_Tp>& m);
template<typename _Tp> _OutputArray(const std::vector<_Tp>& vec);
template<typename _Tp> _OutputArray(const std::vector<std::vector<_Tp> >& vec);
template<typename _Tp> _OutputArray(const std::vector<Mat_<_Tp> >& vec);

486
modules/core/src/cuda/gpu_mat.cu Normal file
View File

@ -0,0 +1,486 @@
/*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 "opencv2/opencv_modules.hpp"
#ifndef HAVE_OPENCV_CUDEV
#error "opencv_cudev is required"
#else
#include "opencv2/core/gpu.hpp"
#include "opencv2/cudev.hpp"
using namespace cv;
using namespace cv::gpu;
using namespace cv::cudev;
/////////////////////////////////////////////////////
/// create
void cv::gpu::GpuMat::create(int _rows, int _cols, int _type)
{
CV_DbgAssert( _rows >= 0 && _cols >= 0 );
_type &= Mat::TYPE_MASK;
if (rows == _rows && cols == _cols && type() == _type && data)
return;
if (data)
release();
if (_rows > 0 && _cols > 0)
{
flags = Mat::MAGIC_VAL + _type;
rows = _rows;
cols = _cols;
size_t esz = elemSize();
void* devPtr;
if (rows > 1 && cols > 1)
{
CV_CUDEV_SAFE_CALL( cudaMallocPitch(&devPtr, &step, esz * cols, rows) );
}
else
{
// Single row or single column must be continuous
CV_CUDEV_SAFE_CALL( cudaMalloc(&devPtr, esz * cols * rows) );
step = esz * cols;
}
if (esz * cols == step)
flags |= Mat::CONTINUOUS_FLAG;
int64 _nettosize = static_cast<int64>(step) * rows;
size_t nettosize = static_cast<size_t>(_nettosize);
datastart = data = static_cast<uchar*>(devPtr);
dataend = data + nettosize;
refcount = static_cast<int*>(fastMalloc(sizeof(*refcount)));
*refcount = 1;
}
}
/////////////////////////////////////////////////////
/// release
void cv::gpu::GpuMat::release()
{
if (refcount && CV_XADD(refcount, -1) == 1)
{
cudaFree(datastart);
fastFree(refcount);
}
data = datastart = dataend = 0;
step = rows = cols = 0;
refcount = 0;
}
/////////////////////////////////////////////////////
/// upload
void cv::gpu::GpuMat::upload(InputArray arr)
{
Mat mat = arr.getMat();
CV_DbgAssert( !mat.empty() );
create(mat.size(), mat.type());
CV_CUDEV_SAFE_CALL( cudaMemcpy2D(data, step, mat.data, mat.step, cols * elemSize(), rows, cudaMemcpyHostToDevice) );
}
void cv::gpu::GpuMat::upload(InputArray arr, Stream& _stream)
{
Mat mat = arr.getMat();
CV_DbgAssert( !mat.empty() );
create(mat.size(), mat.type());
cudaStream_t stream = StreamAccessor::getStream(_stream);
CV_CUDEV_SAFE_CALL( cudaMemcpy2DAsync(data, step, mat.data, mat.step, cols * elemSize(), rows, cudaMemcpyHostToDevice, stream) );
}
/////////////////////////////////////////////////////
/// download
void cv::gpu::GpuMat::download(OutputArray _dst) const
{
CV_DbgAssert( !empty() );
_dst.create(size(), type());
Mat dst = _dst.getMat();
CV_CUDEV_SAFE_CALL( cudaMemcpy2D(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToHost) );
}
void cv::gpu::GpuMat::download(OutputArray _dst, Stream& _stream) const
{
CV_DbgAssert( !empty() );
_dst.create(size(), type());
Mat dst = _dst.getMat();
cudaStream_t stream = StreamAccessor::getStream(_stream);
CV_CUDEV_SAFE_CALL( cudaMemcpy2DAsync(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToHost, stream) );
}
/////////////////////////////////////////////////////
/// copyTo
void cv::gpu::GpuMat::copyTo(OutputArray _dst) const
{
CV_DbgAssert( !empty() );
_dst.create(size(), type());
GpuMat dst = _dst.getGpuMat();
CV_CUDEV_SAFE_CALL( cudaMemcpy2D(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToDevice) );
}
void cv::gpu::GpuMat::copyTo(OutputArray _dst, Stream& _stream) const
{
CV_DbgAssert( !empty() );
_dst.create(size(), type());
GpuMat dst = _dst.getGpuMat();
cudaStream_t stream = StreamAccessor::getStream(_stream);
CV_CUDEV_SAFE_CALL( cudaMemcpy2DAsync(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToDevice, stream) );
}
namespace
{
template <size_t size> struct CopyToPolicy : DefaultTransformPolicy
{
};
template <> struct CopyToPolicy<4> : DefaultTransformPolicy
{
enum {
shift = 2
};
};
template <> struct CopyToPolicy<8> : DefaultTransformPolicy
{
enum {
shift = 1
};
};
template <typename T>
void copyWithMask(const GpuMat& src, const GpuMat& dst, const GpuMat& mask, Stream& stream)
{
gridTransform_< CopyToPolicy<sizeof(typename VecTraits<T>::elem_type)> >(globPtr<T>(src), globPtr<T>(dst), identity<T>(), globPtr<uchar>(mask), stream);
}
}
void cv::gpu::GpuMat::copyTo(OutputArray _dst, InputArray _mask, Stream& stream) const
{
CV_DbgAssert( !empty() );
CV_DbgAssert( depth() <= CV_64F && channels() <= 4 );
GpuMat mask = _mask.getGpuMat();
CV_DbgAssert( size() == mask.size() && mask.depth() == CV_8U && (mask.channels() == 1 || mask.channels() == channels()) );
_dst.create(size(), type());
GpuMat dst = _dst.getGpuMat();
typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, const GpuMat& mask, Stream& stream);
static const func_t funcs[9][4] =
{
{0,0,0,0},
{copyWithMask<uchar>, copyWithMask<uchar2>, copyWithMask<uchar3>, copyWithMask<uchar4>},
{copyWithMask<ushort>, copyWithMask<ushort2>, copyWithMask<ushort3>, copyWithMask<ushort4>},
{0,0,0,0},
{copyWithMask<int>, copyWithMask<int2>, copyWithMask<int3>, copyWithMask<int4>},
{0,0,0,0},
{0,0,0,0},
{0,0,0,0},
{copyWithMask<double>, copyWithMask<double2>, copyWithMask<double3>, copyWithMask<double4>}
};
if (mask.channels() == channels())
{
const func_t func = funcs[elemSize1()][0];
CV_DbgAssert( func != 0 );
func(reshape(1), dst.reshape(1), mask.reshape(1), stream);
}
else
{
const func_t func = funcs[elemSize1()][channels() - 1];
CV_DbgAssert( func != 0 );
func(*this, dst, mask, stream);
}
}
/////////////////////////////////////////////////////
/// setTo
namespace
{
template <typename T>
void setToWithOutMask(const GpuMat& mat, Scalar _scalar, Stream& stream)
{
Scalar_<typename VecTraits<T>::elem_type> scalar = _scalar;
gridTransform(constantPtr(VecTraits<T>::make(scalar.val), mat.rows, mat.cols), globPtr<T>(mat), identity<T>(), stream);
}
template <typename T>
void setToWithMask(const GpuMat& mat, const GpuMat& mask, Scalar _scalar, Stream& stream)
{
Scalar_<typename VecTraits<T>::elem_type> scalar = _scalar;
gridTransform(constantPtr(VecTraits<T>::make(scalar.val), mat.rows, mat.cols), globPtr<T>(mat), identity<T>(), globPtr<uchar>(mask), stream);
}
}
GpuMat& cv::gpu::GpuMat::setTo(Scalar value, Stream& stream)
{
CV_DbgAssert( !empty() );
CV_DbgAssert( depth() <= CV_64F && channels() <= 4 );
if (value[0] == 0.0 && value[1] == 0.0 && value[2] == 0.0 && value[3] == 0.0)
{
// Zero fill
if (stream)
CV_CUDEV_SAFE_CALL( cudaMemset2DAsync(data, step, 0, cols * elemSize(), rows, StreamAccessor::getStream(stream)) );
else
CV_CUDEV_SAFE_CALL( cudaMemset2D(data, step, 0, cols * elemSize(), rows) );
return *this;
}
if (depth() == CV_8U)
{
const int cn = channels();
if (cn == 1
|| (cn == 2 && value[0] == value[1])
|| (cn == 3 && value[0] == value[1] && value[0] == value[2])
|| (cn == 4 && value[0] == value[1] && value[0] == value[2] && value[0] == value[3]))
{
const int val = cv::saturate_cast<uchar>(value[0]);
if (stream)
CV_CUDEV_SAFE_CALL( cudaMemset2DAsync(data, step, val, cols * elemSize(), rows, StreamAccessor::getStream(stream)) );
else
CV_CUDEV_SAFE_CALL( cudaMemset2D(data, step, val, cols * elemSize(), rows) );
return *this;
}
}
typedef void (*func_t)(const GpuMat& mat, Scalar scalar, Stream& stream);
static const func_t funcs[7][4] =
{
{setToWithOutMask<uchar>,setToWithOutMask<uchar2>,setToWithOutMask<uchar3>,setToWithOutMask<uchar4>},
{setToWithOutMask<schar>,setToWithOutMask<char2>,setToWithOutMask<char3>,setToWithOutMask<char4>},
{setToWithOutMask<ushort>,setToWithOutMask<ushort2>,setToWithOutMask<ushort3>,setToWithOutMask<ushort4>},
{setToWithOutMask<short>,setToWithOutMask<short2>,setToWithOutMask<short3>,setToWithOutMask<short4>},
{setToWithOutMask<int>,setToWithOutMask<int2>,setToWithOutMask<int3>,setToWithOutMask<int4>},
{setToWithOutMask<float>,setToWithOutMask<float2>,setToWithOutMask<float3>,setToWithOutMask<float4>},
{setToWithOutMask<double>,setToWithOutMask<double2>,setToWithOutMask<double3>,setToWithOutMask<double4>}
};
funcs[depth()][channels() - 1](*this, value, stream);
return *this;
}
GpuMat& cv::gpu::GpuMat::setTo(Scalar value, InputArray _mask, Stream& stream)
{
CV_DbgAssert( !empty() );
CV_DbgAssert( depth() <= CV_64F && channels() <= 4 );
GpuMat mask = _mask.getGpuMat();
CV_DbgAssert( size() == mask.size() && mask.type() == CV_8UC1 );
typedef void (*func_t)(const GpuMat& mat, const GpuMat& mask, Scalar scalar, Stream& stream);
static const func_t funcs[7][4] =
{
{setToWithMask<uchar>,setToWithMask<uchar2>,setToWithMask<uchar3>,setToWithMask<uchar4>},
{setToWithMask<schar>,setToWithMask<char2>,setToWithMask<char3>,setToWithMask<char4>},
{setToWithMask<ushort>,setToWithMask<ushort2>,setToWithMask<ushort3>,setToWithMask<ushort4>},
{setToWithMask<short>,setToWithMask<short2>,setToWithMask<short3>,setToWithMask<short4>},
{setToWithMask<int>,setToWithMask<int2>,setToWithMask<int3>,setToWithMask<int4>},
{setToWithMask<float>,setToWithMask<float2>,setToWithMask<float3>,setToWithMask<float4>},
{setToWithMask<double>,setToWithMask<double2>,setToWithMask<double3>,setToWithMask<double4>}
};
funcs[depth()][channels() - 1](*this, mask, value, stream);
return *this;
}
/////////////////////////////////////////////////////
/// convertTo
namespace
{
template <typename T> struct ConvertToPolicy : DefaultTransformPolicy
{
};
template <> struct ConvertToPolicy<double> : DefaultTransformPolicy
{
enum {
shift = 1
};
};
template <typename T, typename D>
void convertToNoScale(const GpuMat& src, const GpuMat& dst, Stream& stream)
{
typedef typename VecTraits<T>::elem_type src_elem_type;
typedef typename VecTraits<D>::elem_type dst_elem_type;
typedef typename LargerType<src_elem_type, float>::type larger_elem_type;
typedef typename LargerType<float, dst_elem_type>::type scalar_type;
gridTransform_< ConvertToPolicy<scalar_type> >(globPtr<T>(src), globPtr<D>(dst), saturate_cast_func<T, D>(), stream);
}
template <typename T, typename D, typename S> struct Convertor : unary_function<T, D>
{
S alpha;
S beta;
__device__ __forceinline__ D operator ()(typename TypeTraits<T>::parameter_type src) const
{
return cudev::saturate_cast<D>(alpha * src + beta);
}
};
template <typename T, typename D>
void convertToScale(const GpuMat& src, const GpuMat& dst, double alpha, double beta, Stream& stream)
{
typedef typename VecTraits<T>::elem_type src_elem_type;
typedef typename VecTraits<D>::elem_type dst_elem_type;
typedef typename LargerType<src_elem_type, float>::type larger_elem_type;
typedef typename LargerType<float, dst_elem_type>::type scalar_type;
Convertor<T, D, scalar_type> op;
op.alpha = cv::saturate_cast<scalar_type>(alpha);
op.beta = cv::saturate_cast<scalar_type>(beta);
gridTransform_< ConvertToPolicy<scalar_type> >(globPtr<T>(src), globPtr<D>(dst), op, stream);
}
}
void cv::gpu::GpuMat::convertTo(OutputArray _dst, int rtype, Stream& stream) const
{
if (rtype < 0)
rtype = type();
else
rtype = CV_MAKE_TYPE(CV_MAT_DEPTH(rtype), channels());
const int sdepth = depth();
const int ddepth = CV_MAT_DEPTH(rtype);
if (sdepth == ddepth)
{
if (stream)
copyTo(_dst, stream);
else
copyTo(_dst);
return;
}
CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F );
GpuMat src = *this;
_dst.create(size(), rtype);
GpuMat dst = _dst.getGpuMat();
typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream);
static const func_t funcs[7][7] =
{
{0, convertToNoScale<uchar, schar>, convertToNoScale<uchar, ushort>, convertToNoScale<uchar, short>, convertToNoScale<uchar, int>, convertToNoScale<uchar, float>, convertToNoScale<uchar, double>},
{convertToNoScale<schar, uchar>, 0, convertToNoScale<schar, ushort>, convertToNoScale<schar, short>, convertToNoScale<schar, int>, convertToNoScale<schar, float>, convertToNoScale<schar, double>},
{convertToNoScale<ushort, uchar>, convertToNoScale<ushort, schar>, 0, convertToNoScale<ushort, short>, convertToNoScale<ushort, int>, convertToNoScale<ushort, float>, convertToNoScale<ushort, double>},
{convertToNoScale<short, uchar>, convertToNoScale<short, schar>, convertToNoScale<short, ushort>, 0, convertToNoScale<short, int>, convertToNoScale<short, float>, convertToNoScale<short, double>},
{convertToNoScale<int, uchar>, convertToNoScale<int, schar>, convertToNoScale<int, ushort>, convertToNoScale<int, short>, 0, convertToNoScale<int, float>, convertToNoScale<int, double>},
{convertToNoScale<float, uchar>, convertToNoScale<float, schar>, convertToNoScale<float, ushort>, convertToNoScale<float, short>, convertToNoScale<float, int>, 0, convertToNoScale<float, double>},
{convertToNoScale<double, uchar>, convertToNoScale<double, schar>, convertToNoScale<double, ushort>, convertToNoScale<double, short>, convertToNoScale<double, int>, convertToNoScale<double, float>, 0}
};
funcs[sdepth][ddepth](reshape(1), dst.reshape(1), stream);
}
void cv::gpu::GpuMat::convertTo(OutputArray _dst, int rtype, double alpha, double beta, Stream& stream) const
{
if (rtype < 0)
rtype = type();
else
rtype = CV_MAKETYPE(CV_MAT_DEPTH(rtype), channels());
const int sdepth = depth();
const int ddepth = CV_MAT_DEPTH(rtype);
GpuMat src = *this;
_dst.create(size(), rtype);
GpuMat dst = _dst.getGpuMat();
typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, double alpha, double beta, Stream& stream);
static const func_t funcs[7][7] =
{
{convertToScale<uchar, uchar>, convertToScale<uchar, schar>, convertToScale<uchar, ushort>, convertToScale<uchar, short>, convertToScale<uchar, int>, convertToScale<uchar, float>, convertToScale<uchar, double>},
{convertToScale<schar, uchar>, convertToScale<schar, schar>, convertToScale<schar, ushort>, convertToScale<schar, short>, convertToScale<schar, int>, convertToScale<schar, float>, convertToScale<schar, double>},
{convertToScale<ushort, uchar>, convertToScale<ushort, schar>, convertToScale<ushort, ushort>, convertToScale<ushort, short>, convertToScale<ushort, int>, convertToScale<ushort, float>, convertToScale<ushort, double>},
{convertToScale<short, uchar>, convertToScale<short, schar>, convertToScale<short, ushort>, convertToScale<short, short>, convertToScale<short, int>, convertToScale<short, float>, convertToScale<short, double>},
{convertToScale<int, uchar>, convertToScale<int, schar>, convertToScale<int, ushort>, convertToScale<int, short>, convertToScale<int, int>, convertToScale<int, float>, convertToScale<int, double>},
{convertToScale<float, uchar>, convertToScale<float, schar>, convertToScale<float, ushort>, convertToScale<float, short>, convertToScale<float, int>, convertToScale<float, float>, convertToScale<float, double>},
{convertToScale<double, uchar>, convertToScale<double, schar>, convertToScale<double, ushort>, convertToScale<double, short>, convertToScale<double, int>, convertToScale<double, float>, convertToScale<double, double>}
};
funcs[sdepth][ddepth](reshape(1), dst.reshape(1), alpha, beta, stream);
}
#endif

296
modules/core/src/cuda/matrix_operations.cu
View File

@ -1,296 +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 "opencv2/core/cuda/saturate_cast.hpp"
#include "opencv2/core/cuda/transform.hpp"
#include "opencv2/core/cuda/functional.hpp"
#include "opencv2/core/cuda/type_traits.hpp"
#include "opencv2/core/cuda/vec_traits.hpp"
#include "matrix_operations.hpp"
namespace cv { namespace gpu { namespace cudev
{
///////////////////////////////////////////////////////////////////////////
// copyWithMask
template <typename T>
void copyWithMask(PtrStepSzb src, PtrStepSzb dst, int cn, PtrStepSzb mask, bool multiChannelMask, cudaStream_t stream)
{
if (multiChannelMask)
cv::gpu::cudev::transform((PtrStepSz<T>) src, (PtrStepSz<T>) dst, identity<T>(), SingleMask(mask), stream);
else
cv::gpu::cudev::transform((PtrStepSz<T>) src, (PtrStepSz<T>) dst, identity<T>(), SingleMaskChannels(mask, cn), stream);
}
void copyWithMask(PtrStepSzb src, PtrStepSzb dst, size_t elemSize1, int cn, PtrStepSzb mask, bool multiChannelMask, cudaStream_t stream)
{
typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, int cn, PtrStepSzb mask, bool multiChannelMask, cudaStream_t stream);
static const func_t tab[] =
{
0,
copyWithMask<uchar>,
copyWithMask<ushort>,
0,
copyWithMask<int>,
0,
0,
0,
copyWithMask<double>
};
const func_t func = tab[elemSize1];
CV_DbgAssert( func != 0 );
func(src, dst, cn, mask, multiChannelMask, stream);
}
///////////////////////////////////////////////////////////////////////////
// set
template<typename T, class Mask>
__global__ void set(PtrStepSz<T> mat, const Mask mask, const int channels, const typename TypeVec<T, 4>::vec_type value)
{
const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= mat.cols * channels || y >= mat.rows)
return;
const T scalar[4] = {value.x, value.y, value.z, value.w};
if (mask(y, x / channels))
mat(y, x) = scalar[x % channels];
}
template <typename T>
void set(PtrStepSz<T> mat, const T* scalar, int channels, cudaStream_t stream)
{
typedef typename TypeVec<T, 4>::vec_type scalar_t;
dim3 block(32, 8);
dim3 grid(divUp(mat.cols * channels, block.x), divUp(mat.rows, block.y));
set<T><<<grid, block, 0, stream>>>(mat, WithOutMask(), channels, VecTraits<scalar_t>::make(scalar));
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall ( cudaDeviceSynchronize() );
}
template void set<uchar >(PtrStepSz<uchar > mat, const uchar* scalar, int channels, cudaStream_t stream);
template void set<schar >(PtrStepSz<schar > mat, const schar* scalar, int channels, cudaStream_t stream);
template void set<ushort>(PtrStepSz<ushort> mat, const ushort* scalar, int channels, cudaStream_t stream);
template void set<short >(PtrStepSz<short > mat, const short* scalar, int channels, cudaStream_t stream);
template void set<int >(PtrStepSz<int > mat, const int* scalar, int channels, cudaStream_t stream);
template void set<float >(PtrStepSz<float > mat, const float* scalar, int channels, cudaStream_t stream);
template void set<double>(PtrStepSz<double> mat, const double* scalar, int channels, cudaStream_t stream);
template <typename T>
void set(PtrStepSz<T> mat, const T* scalar, PtrStepSzb mask, int channels, cudaStream_t stream)
{
typedef typename TypeVec<T, 4>::vec_type scalar_t;
dim3 block(32, 8);
dim3 grid(divUp(mat.cols * channels, block.x), divUp(mat.rows, block.y));
set<T><<<grid, block, 0, stream>>>(mat, SingleMask(mask), channels, VecTraits<scalar_t>::make(scalar));
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall ( cudaDeviceSynchronize() );
}
template void set<uchar >(PtrStepSz<uchar > mat, const uchar* scalar, PtrStepSzb mask, int channels, cudaStream_t stream);
template void set<schar >(PtrStepSz<schar > mat, const schar* scalar, PtrStepSzb mask, int channels, cudaStream_t stream);
template void set<ushort>(PtrStepSz<ushort> mat, const ushort* scalar, PtrStepSzb mask, int channels, cudaStream_t stream);
template void set<short >(PtrStepSz<short > mat, const short* scalar, PtrStepSzb mask, int channels, cudaStream_t stream);
template void set<int >(PtrStepSz<int > mat, const int* scalar, PtrStepSzb mask, int channels, cudaStream_t stream);
template void set<float >(PtrStepSz<float > mat, const float* scalar, PtrStepSzb mask, int channels, cudaStream_t stream);
template void set<double>(PtrStepSz<double> mat, const double* scalar, PtrStepSzb mask, int channels, cudaStream_t stream);
///////////////////////////////////////////////////////////////////////////
// convert
template <typename T, typename D, typename S> struct Convertor : unary_function<T, D>
{
Convertor(S alpha_, S beta_) : alpha(alpha_), beta(beta_) {}
__device__ __forceinline__ D operator()(typename TypeTraits<T>::ParameterType src) const
{
return saturate_cast<D>(alpha * src + beta);
}
S alpha, beta;
};
namespace detail
{
template <size_t src_size, size_t dst_size, typename F> struct ConvertTraitsDispatcher : DefaultTransformFunctorTraits<F>
{
};
template <typename F> struct ConvertTraitsDispatcher<1, 1, F> : DefaultTransformFunctorTraits<F>
{
enum { smart_shift = 8 };
};
template <typename F> struct ConvertTraitsDispatcher<1, 2, F> : DefaultTransformFunctorTraits<F>
{
enum { smart_shift = 4 };
};
template <typename F> struct ConvertTraitsDispatcher<1, 4, F> : DefaultTransformFunctorTraits<F>
{
enum { smart_block_dim_y = 8 };
enum { smart_shift = 4 };
};
template <typename F> struct ConvertTraitsDispatcher<2, 2, F> : DefaultTransformFunctorTraits<F>
{
enum { smart_shift = 4 };
};
template <typename F> struct ConvertTraitsDispatcher<2, 4, F> : DefaultTransformFunctorTraits<F>
{
enum { smart_shift = 2 };
};
template <typename F> struct ConvertTraitsDispatcher<4, 2, F> : DefaultTransformFunctorTraits<F>
{
enum { smart_block_dim_y = 8 };
enum { smart_shift = 4 };
};
template <typename F> struct ConvertTraitsDispatcher<4, 4, F> : DefaultTransformFunctorTraits<F>
{
enum { smart_block_dim_y = 8 };
enum { smart_shift = 2 };
};
template <typename F> struct ConvertTraits : ConvertTraitsDispatcher<sizeof(typename F::argument_type), sizeof(typename F::result_type), F>
{
};
}
template <typename T, typename D, typename S> struct TransformFunctorTraits< Convertor<T, D, S> > : detail::ConvertTraits< Convertor<T, D, S> >
{
};
template<typename T, typename D, typename S>
void cvt_(PtrStepSzb src, PtrStepSzb dst, double alpha, double beta, cudaStream_t stream)
{
Convertor<T, D, S> op(static_cast<S>(alpha), static_cast<S>(beta));
cv::gpu::cudev::transform((PtrStepSz<T>)src, (PtrStepSz<D>)dst, op, WithOutMask(), stream);
}
void convert(PtrStepSzb src, int sdepth, PtrStepSzb dst, int ddepth, double alpha, double beta, cudaStream_t stream)
{
typedef void (*caller_t)(PtrStepSzb src, PtrStepSzb dst, double alpha, double beta, cudaStream_t stream);
static const caller_t tab[7][7] =
{
{
cvt_<uchar, uchar, float>,
cvt_<uchar, schar, float>,
cvt_<uchar, ushort, float>,
cvt_<uchar, short, float>,
cvt_<uchar, int, float>,
cvt_<uchar, float, float>,
cvt_<uchar, double, double>
},
{
cvt_<schar, uchar, float>,
cvt_<schar, schar, float>,
cvt_<schar, ushort, float>,
cvt_<schar, short, float>,
cvt_<schar, int, float>,
cvt_<schar, float, float>,
cvt_<schar, double, double>
},
{
cvt_<ushort, uchar, float>,
cvt_<ushort, schar, float>,
cvt_<ushort, ushort, float>,
cvt_<ushort, short, float>,
cvt_<ushort, int, float>,
cvt_<ushort, float, float>,
cvt_<ushort, double, double>
},
{
cvt_<short, uchar, float>,
cvt_<short, schar, float>,
cvt_<short, ushort, float>,
cvt_<short, short, float>,
cvt_<short, int, float>,
cvt_<short, float, float>,
cvt_<short, double, double>
},
{
cvt_<int, uchar, float>,
cvt_<int, schar, float>,
cvt_<int, ushort, float>,
cvt_<int, short, float>,
cvt_<int, int, double>,
cvt_<int, float, double>,
cvt_<int, double, double>
},
{
cvt_<float, uchar, float>,
cvt_<float, schar, float>,
cvt_<float, ushort, float>,
cvt_<float, short, float>,
cvt_<float, int, float>,
cvt_<float, float, float>,
cvt_<float, double, double>
},
{
cvt_<double, uchar, double>,
cvt_<double, schar, double>,
cvt_<double, ushort, double>,
cvt_<double, short, double>,
cvt_<double, int, double>,
cvt_<double, float, double>,
cvt_<double, double, double>
}
};
const caller_t func = tab[sdepth][ddepth];
func(src, dst, alpha, beta, stream);
}
}}} // namespace cv { namespace gpu { namespace cudev

57
modules/core/src/cuda/matrix_operations.hpp
View File

@ -1,57 +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.
// Copyright (C) 2013, OpenCV Foundation, 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 "opencv2/core/cuda/common.hpp"
namespace cv { namespace gpu { namespace cudev
{
void copyWithMask(PtrStepSzb src, PtrStepSzb dst, size_t elemSize1, int cn, PtrStepSzb mask, bool multiChannelMask, cudaStream_t stream);
template <typename T>
void set(PtrStepSz<T> mat, const T* scalar, int channels, cudaStream_t stream);
template <typename T>
void set(PtrStepSz<T> mat, const T* scalar, PtrStepSzb mask, int channels, cudaStream_t stream);
void convert(PtrStepSzb src, int sdepth, PtrStepSzb dst, int ddepth, double alpha, double beta, cudaStream_t stream);
}}}

878
modules/core/src/gpu_mat.cpp
View File

@ -46,504 +46,6 @@
using namespace cv;
using namespace cv::gpu;
/////////////////////////// matrix operations /////////////////////////
#ifdef HAVE_CUDA
// CUDA implementation
#include "cuda/matrix_operations.hpp"
namespace
{
template <typename T> void cudaSet_(GpuMat& src, Scalar s, cudaStream_t stream)
{
Scalar_<T> sf = s;
cudev::set<T>(PtrStepSz<T>(src), sf.val, src.channels(), stream);
}
void cudaSet(GpuMat& src, Scalar s, cudaStream_t stream)
{
typedef void (*func_t)(GpuMat& src, Scalar s, cudaStream_t stream);
static const func_t funcs[] =
{
cudaSet_<uchar>,
cudaSet_<schar>,
cudaSet_<ushort>,
cudaSet_<short>,
cudaSet_<int>,
cudaSet_<float>,
cudaSet_<double>
};
funcs[src.depth()](src, s, stream);
}
template <typename T> void cudaSet_(GpuMat& src, Scalar s, PtrStepSzb mask, cudaStream_t stream)
{
Scalar_<T> sf = s;
cudev::set<T>(PtrStepSz<T>(src), sf.val, mask, src.channels(), stream);
}
void cudaSet(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream)
{
typedef void (*func_t)(GpuMat& src, Scalar s, PtrStepSzb mask, cudaStream_t stream);
static const func_t funcs[] =
{
cudaSet_<uchar>,
cudaSet_<schar>,
cudaSet_<ushort>,
cudaSet_<short>,
cudaSet_<int>,
cudaSet_<float>,
cudaSet_<double>
};
funcs[src.depth()](src, s, mask, stream);
}
void cudaCopyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)
{
cudev::copyWithMask(src.reshape(1), dst.reshape(1), src.elemSize1(), src.channels(), mask.reshape(1), mask.channels() != 1, stream);
}
void cudaConvert(const GpuMat& src, GpuMat& dst, cudaStream_t stream)
{
cudev::convert(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), 1.0, 0.0, stream);
}
void cudaConvert(const GpuMat& src, GpuMat& dst, double alpha, double beta, cudaStream_t stream)
{
cudev::convert(src.reshape(1), src.depth(), dst.reshape(1), dst.depth(), alpha, beta, stream);
}
}
// NPP implementation
namespace
{
//////////////////////////////////////////////////////////////////////////
// Convert
template<int SDEPTH, int DDEPTH> struct NppConvertFunc
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef typename NPPTypeTraits<DDEPTH>::npp_type dst_t;
typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, dst_t* pDst, int nDstStep, NppiSize oSizeROI);
};
template<int DDEPTH> struct NppConvertFunc<CV_32F, DDEPTH>
{
typedef typename NPPTypeTraits<DDEPTH>::npp_type dst_t;
typedef NppStatus (*func_ptr)(const Npp32f* pSrc, int nSrcStep, dst_t* pDst, int nDstStep, NppiSize oSizeROI, NppRoundMode eRoundMode);
};
template<int SDEPTH, int DDEPTH, typename NppConvertFunc<SDEPTH, DDEPTH>::func_ptr func> struct NppCvt
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef typename NPPTypeTraits<DDEPTH>::npp_type dst_t;
static void call(const GpuMat& src, GpuMat& dst, cudaStream_t stream)
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
NppStreamHandler h(stream);
nppSafeCall( func(src.ptr<src_t>(), static_cast<int>(src.step), dst.ptr<dst_t>(), static_cast<int>(dst.step), sz) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
template<int DDEPTH, typename NppConvertFunc<CV_32F, DDEPTH>::func_ptr func> struct NppCvt<CV_32F, DDEPTH, func>
{
typedef typename NPPTypeTraits<DDEPTH>::npp_type dst_t;
static void call(const GpuMat& src, GpuMat& dst, cudaStream_t stream)
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
NppStreamHandler h(stream);
nppSafeCall( func(src.ptr<Npp32f>(), static_cast<int>(src.step), dst.ptr<dst_t>(), static_cast<int>(dst.step), sz, NPP_RND_NEAR) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
//////////////////////////////////////////////////////////////////////////
// Set
template<int SDEPTH, int SCN> struct NppSetFunc
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(const src_t values[], src_t* pSrc, int nSrcStep, NppiSize oSizeROI);
};
template<int SDEPTH> struct NppSetFunc<SDEPTH, 1>
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(src_t val, src_t* pSrc, int nSrcStep, NppiSize oSizeROI);
};
template<int SCN> struct NppSetFunc<CV_8S, SCN>
{
typedef NppStatus (*func_ptr)(Npp8s values[], Npp8s* pSrc, int nSrcStep, NppiSize oSizeROI);
};
template<> struct NppSetFunc<CV_8S, 1>
{
typedef NppStatus (*func_ptr)(Npp8s val, Npp8s* pSrc, int nSrcStep, NppiSize oSizeROI);
};
template<int SDEPTH, int SCN, typename NppSetFunc<SDEPTH, SCN>::func_ptr func> struct NppSet
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(GpuMat& src, Scalar s, cudaStream_t stream)
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
Scalar_<src_t> nppS = s;
NppStreamHandler h(stream);
nppSafeCall( func(nppS.val, src.ptr<src_t>(), static_cast<int>(src.step), sz) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
template<int SDEPTH, typename NppSetFunc<SDEPTH, 1>::func_ptr func> struct NppSet<SDEPTH, 1, func>
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(GpuMat& src, Scalar s, cudaStream_t stream)
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
Scalar_<src_t> nppS = s;
NppStreamHandler h(stream);
nppSafeCall( func(nppS[0], src.ptr<src_t>(), static_cast<int>(src.step), sz) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
template<int SDEPTH, int SCN> struct NppSetMaskFunc
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(const src_t values[], src_t* pSrc, int nSrcStep, NppiSize oSizeROI, const Npp8u* pMask, int nMaskStep);
};
template<int SDEPTH> struct NppSetMaskFunc<SDEPTH, 1>
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(src_t val, src_t* pSrc, int nSrcStep, NppiSize oSizeROI, const Npp8u* pMask, int nMaskStep);
};
template<int SDEPTH, int SCN, typename NppSetMaskFunc<SDEPTH, SCN>::func_ptr func> struct NppSetMask
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream)
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
Scalar_<src_t> nppS = s;
NppStreamHandler h(stream);
nppSafeCall( func(nppS.val, src.ptr<src_t>(), static_cast<int>(src.step), sz, mask.ptr<Npp8u>(), static_cast<int>(mask.step)) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
template<int SDEPTH, typename NppSetMaskFunc<SDEPTH, 1>::func_ptr func> struct NppSetMask<SDEPTH, 1, func>
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream)
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
Scalar_<src_t> nppS = s;
NppStreamHandler h(stream);
nppSafeCall( func(nppS[0], src.ptr<src_t>(), static_cast<int>(src.step), sz, mask.ptr<Npp8u>(), static_cast<int>(mask.step)) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
//////////////////////////////////////////////////////////////////////////
// CopyMasked
template<int SDEPTH> struct NppCopyWithMaskFunc
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, src_t* pDst, int nDstStep, NppiSize oSizeROI, const Npp8u* pMask, int nMaskStep);
};
template<int SDEPTH, typename NppCopyWithMaskFunc<SDEPTH>::func_ptr func> struct NppCopyWithMask
{
typedef typename NPPTypeTraits<SDEPTH>::npp_type src_t;
static void call(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
NppStreamHandler h(stream);
nppSafeCall( func(src.ptr<src_t>(), static_cast<int>(src.step), dst.ptr<src_t>(), static_cast<int>(dst.step), sz, mask.ptr<Npp8u>(), static_cast<int>(mask.step)) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
}
// Dispatcher
namespace
{
void copyWithMask(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream = 0)
{
CV_DbgAssert( src.size() == dst.size() && src.type() == dst.type() );
CV_Assert( src.depth() <= CV_64F && src.channels() <= 4 );
CV_Assert( src.size() == mask.size() && mask.depth() == CV_8U && (mask.channels() == 1 || mask.channels() == src.channels()) );
if (src.depth() == CV_64F)
{
CV_Assert( deviceSupports(NATIVE_DOUBLE) );
}
typedef void (*func_t)(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream);
static const func_t funcs[7][4] =
{
/* 8U */ {NppCopyWithMask<CV_8U , nppiCopy_8u_C1MR >::call, cudaCopyWithMask, NppCopyWithMask<CV_8U , nppiCopy_8u_C3MR >::call, NppCopyWithMask<CV_8U , nppiCopy_8u_C4MR >::call},
/* 8S */ {cudaCopyWithMask , cudaCopyWithMask, cudaCopyWithMask , cudaCopyWithMask },
/* 16U */ {NppCopyWithMask<CV_16U, nppiCopy_16u_C1MR>::call, cudaCopyWithMask, NppCopyWithMask<CV_16U, nppiCopy_16u_C3MR>::call, NppCopyWithMask<CV_16U, nppiCopy_16u_C4MR>::call},
/* 16S */ {NppCopyWithMask<CV_16S, nppiCopy_16s_C1MR>::call, cudaCopyWithMask, NppCopyWithMask<CV_16S, nppiCopy_16s_C3MR>::call, NppCopyWithMask<CV_16S, nppiCopy_16s_C4MR>::call},
/* 32S */ {NppCopyWithMask<CV_32S, nppiCopy_32s_C1MR>::call, cudaCopyWithMask, NppCopyWithMask<CV_32S, nppiCopy_32s_C3MR>::call, NppCopyWithMask<CV_32S, nppiCopy_32s_C4MR>::call},
/* 32F */ {NppCopyWithMask<CV_32F, nppiCopy_32f_C1MR>::call, cudaCopyWithMask, NppCopyWithMask<CV_32F, nppiCopy_32f_C3MR>::call, NppCopyWithMask<CV_32F, nppiCopy_32f_C4MR>::call},
/* 64F */ {cudaCopyWithMask , cudaCopyWithMask, cudaCopyWithMask , cudaCopyWithMask }
};
const func_t func = mask.channels() == src.channels() ? funcs[src.depth()][src.channels() - 1] : cudaCopyWithMask;
func(src, dst, mask, stream);
}
void convert(const GpuMat& src, GpuMat& dst, cudaStream_t stream = 0)
{
CV_DbgAssert( src.size() == dst.size() && src.channels() == dst.channels() );
CV_Assert( src.depth() <= CV_64F && src.channels() <= 4 );
CV_Assert( dst.depth() <= CV_64F );
if (src.depth() == CV_64F || dst.depth() == CV_64F)
{
CV_Assert( deviceSupports(NATIVE_DOUBLE) );
}
typedef void (*func_t)(const GpuMat& src, GpuMat& dst, cudaStream_t stream);
static const func_t funcs[7][7][4] =
{
{
/* 8U -> 8U */ {0, 0, 0, 0},
/* 8U -> 8S */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert },
/* 8U -> 16U */ {NppCvt<CV_8U, CV_16U, nppiConvert_8u16u_C1R>::call, cudaConvert, cudaConvert, NppCvt<CV_8U, CV_16U, nppiConvert_8u16u_C4R>::call},
/* 8U -> 16S */ {NppCvt<CV_8U, CV_16S, nppiConvert_8u16s_C1R>::call, cudaConvert, cudaConvert, NppCvt<CV_8U, CV_16S, nppiConvert_8u16s_C4R>::call},
/* 8U -> 32S */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert },
/* 8U -> 32F */ {NppCvt<CV_8U, CV_32F, nppiConvert_8u32f_C1R>::call, cudaConvert, cudaConvert, cudaConvert },
/* 8U -> 64F */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert }
},
{
/* 8S -> 8U */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 8S -> 8S */ {0,0,0,0},
/* 8S -> 16U */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 8S -> 16S */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 8S -> 32S */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 8S -> 32F */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 8S -> 64F */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert}
},
{
/* 16U -> 8U */ {NppCvt<CV_16U, CV_8U , nppiConvert_16u8u_C1R >::call, cudaConvert, cudaConvert, NppCvt<CV_16U, CV_8U, nppiConvert_16u8u_C4R>::call},
/* 16U -> 8S */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert },
/* 16U -> 16U */ {0,0,0,0},
/* 16U -> 16S */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert },
/* 16U -> 32S */ {NppCvt<CV_16U, CV_32S, nppiConvert_16u32s_C1R>::call, cudaConvert, cudaConvert, cudaConvert },
/* 16U -> 32F */ {NppCvt<CV_16U, CV_32F, nppiConvert_16u32f_C1R>::call, cudaConvert, cudaConvert, cudaConvert },
/* 16U -> 64F */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert }
},
{
/* 16S -> 8U */ {NppCvt<CV_16S, CV_8U , nppiConvert_16s8u_C1R >::call, cudaConvert, cudaConvert, NppCvt<CV_16S, CV_8U, nppiConvert_16s8u_C4R>::call},
/* 16S -> 8S */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert },
/* 16S -> 16U */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert },
/* 16S -> 16S */ {0,0,0,0},
/* 16S -> 32S */ {NppCvt<CV_16S, CV_32S, nppiConvert_16s32s_C1R>::call, cudaConvert, cudaConvert, cudaConvert },
/* 16S -> 32F */ {NppCvt<CV_16S, CV_32F, nppiConvert_16s32f_C1R>::call, cudaConvert, cudaConvert, cudaConvert },
/* 16S -> 64F */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert }
},
{
/* 32S -> 8U */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 32S -> 8S */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 32S -> 16U */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 32S -> 16S */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 32S -> 32S */ {0,0,0,0},
/* 32S -> 32F */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 32S -> 64F */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert}
},
{
/* 32F -> 8U */ {NppCvt<CV_32F, CV_8U , nppiConvert_32f8u_C1R >::call, cudaConvert, cudaConvert, cudaConvert},
/* 32F -> 8S */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert},
/* 32F -> 16U */ {NppCvt<CV_32F, CV_16U, nppiConvert_32f16u_C1R>::call, cudaConvert, cudaConvert, cudaConvert},
/* 32F -> 16S */ {NppCvt<CV_32F, CV_16S, nppiConvert_32f16s_C1R>::call, cudaConvert, cudaConvert, cudaConvert},
/* 32F -> 32S */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert},
/* 32F -> 32F */ {0,0,0,0},
/* 32F -> 64F */ {cudaConvert , cudaConvert, cudaConvert, cudaConvert}
},
{
/* 64F -> 8U */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 64F -> 8S */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 64F -> 16U */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 64F -> 16S */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 64F -> 32S */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 64F -> 32F */ {cudaConvert, cudaConvert, cudaConvert, cudaConvert},
/* 64F -> 64F */ {0,0,0,0}
}
};
const bool aligned = isAligned(src.data, 16) && isAligned(dst.data, 16);
if (!aligned)
{
cudaConvert(src, dst, stream);
return;
}
const func_t func = funcs[src.depth()][dst.depth()][src.channels() - 1];
CV_DbgAssert( func != 0 );
func(src, dst, stream);
}
void convert(const GpuMat& src, GpuMat& dst, double alpha, double beta, cudaStream_t stream = 0)
{
CV_DbgAssert( src.size() == dst.size() && src.channels() == dst.channels() );
CV_Assert( src.depth() <= CV_64F && src.channels() <= 4 );
CV_Assert( dst.depth() <= CV_64F );
if (src.depth() == CV_64F || dst.depth() == CV_64F)
{
CV_Assert( deviceSupports(NATIVE_DOUBLE) );
}
cudaConvert(src, dst, alpha, beta, stream);
}
void set(GpuMat& m, Scalar s, cudaStream_t stream = 0)
{
if (s[0] == 0.0 && s[1] == 0.0 && s[2] == 0.0 && s[3] == 0.0)
{
if (stream)
cudaSafeCall( cudaMemset2DAsync(m.data, m.step, 0, m.cols * m.elemSize(), m.rows, stream) );
else
cudaSafeCall( cudaMemset2D(m.data, m.step, 0, m.cols * m.elemSize(), m.rows) );
return;
}
if (m.depth() == CV_8U)
{
int cn = m.channels();
if (cn == 1 || (cn == 2 && s[0] == s[1]) || (cn == 3 && s[0] == s[1] && s[0] == s[2]) || (cn == 4 && s[0] == s[1] && s[0] == s[2] && s[0] == s[3]))
{
int val = saturate_cast<uchar>(s[0]);
if (stream)
cudaSafeCall( cudaMemset2DAsync(m.data, m.step, val, m.cols * m.elemSize(), m.rows, stream) );
else
cudaSafeCall( cudaMemset2D(m.data, m.step, val, m.cols * m.elemSize(), m.rows) );
return;
}
}
typedef void (*func_t)(GpuMat& src, Scalar s, cudaStream_t stream);
static const func_t funcs[7][4] =
{
{NppSet<CV_8U , 1, nppiSet_8u_C1R >::call, cudaSet , cudaSet , NppSet<CV_8U , 4, nppiSet_8u_C4R >::call},
{NppSet<CV_8S , 1, nppiSet_8s_C1R >::call, NppSet<CV_8S , 2, nppiSet_8s_C2R >::call, NppSet<CV_8S, 3, nppiSet_8s_C3R>::call, NppSet<CV_8S , 4, nppiSet_8s_C4R >::call},
{NppSet<CV_16U, 1, nppiSet_16u_C1R>::call, NppSet<CV_16U, 2, nppiSet_16u_C2R>::call, cudaSet , NppSet<CV_16U, 4, nppiSet_16u_C4R>::call},
{NppSet<CV_16S, 1, nppiSet_16s_C1R>::call, NppSet<CV_16S, 2, nppiSet_16s_C2R>::call, cudaSet , NppSet<CV_16S, 4, nppiSet_16s_C4R>::call},
{NppSet<CV_32S, 1, nppiSet_32s_C1R>::call, cudaSet , cudaSet , NppSet<CV_32S, 4, nppiSet_32s_C4R>::call},
{NppSet<CV_32F, 1, nppiSet_32f_C1R>::call, cudaSet , cudaSet , NppSet<CV_32F, 4, nppiSet_32f_C4R>::call},
{cudaSet , cudaSet , cudaSet , cudaSet }
};
CV_Assert( m.depth() <= CV_64F && m.channels() <= 4 );
if (m.depth() == CV_64F)
{
CV_Assert( deviceSupports(NATIVE_DOUBLE) );
}
funcs[m.depth()][m.channels() - 1](m, s, stream);
}
void set(GpuMat& m, Scalar s, const GpuMat& mask, cudaStream_t stream = 0)
{
CV_DbgAssert( !mask.empty() );
CV_Assert( m.depth() <= CV_64F && m.channels() <= 4 );
if (m.depth() == CV_64F)
{
CV_Assert( deviceSupports(NATIVE_DOUBLE) );
}
typedef void (*func_t)(GpuMat& src, Scalar s, const GpuMat& mask, cudaStream_t stream);
static const func_t funcs[7][4] =
{
{NppSetMask<CV_8U , 1, nppiSet_8u_C1MR >::call, cudaSet, cudaSet, NppSetMask<CV_8U , 4, nppiSet_8u_C4MR >::call},
{cudaSet , cudaSet, cudaSet, cudaSet },
{NppSetMask<CV_16U, 1, nppiSet_16u_C1MR>::call, cudaSet, cudaSet, NppSetMask<CV_16U, 4, nppiSet_16u_C4MR>::call},
{NppSetMask<CV_16S, 1, nppiSet_16s_C1MR>::call, cudaSet, cudaSet, NppSetMask<CV_16S, 4, nppiSet_16s_C4MR>::call},
{NppSetMask<CV_32S, 1, nppiSet_32s_C1MR>::call, cudaSet, cudaSet, NppSetMask<CV_32S, 4, nppiSet_32s_C4MR>::call},
{NppSetMask<CV_32F, 1, nppiSet_32f_C1MR>::call, cudaSet, cudaSet, NppSetMask<CV_32F, 4, nppiSet_32f_C4MR>::call},
{cudaSet , cudaSet, cudaSet, cudaSet }
};
funcs[m.depth()][m.channels() - 1](m, s, mask, stream);
}
}
#endif // HAVE_CUDA
cv::gpu::GpuMat::GpuMat(int rows_, int cols_, int type_, void* data_, size_t step_) :
flags(Mat::MAGIC_VAL + (type_ & Mat::TYPE_MASK)), rows(rows_), cols(cols_),
step(step_), data((uchar*)data_), refcount(0),
@ -651,288 +153,6 @@ cv::gpu::GpuMat::GpuMat(const GpuMat& m, Rect roi) :
rows = cols = 0;
}
void cv::gpu::GpuMat::create(int _rows, int _cols, int _type)
{
#ifndef HAVE_CUDA
(void) _rows;
(void) _cols;
(void) _type;
throw_no_cuda();
#else
_type &= Mat::TYPE_MASK;
if (rows == _rows && cols == _cols && type() == _type && data)
return;
if (data)
release();
CV_DbgAssert( _rows >= 0 && _cols >= 0 );
if (_rows > 0 && _cols > 0)
{
flags = Mat::MAGIC_VAL + _type;
rows = _rows;
cols = _cols;
size_t esz = elemSize();
void* devPtr;
if (rows > 1 && cols > 1)
{
cudaSafeCall( cudaMallocPitch(&devPtr, &step, esz * cols, rows) );
}
else
{
// Single row or single column must be continuous
cudaSafeCall( cudaMalloc(&devPtr, esz * cols * rows) );
step = esz * cols;
}
if (esz * cols == step)
flags |= Mat::CONTINUOUS_FLAG;
int64 _nettosize = static_cast<int64>(step) * rows;
size_t nettosize = static_cast<size_t>(_nettosize);
datastart = data = static_cast<uchar*>(devPtr);
dataend = data + nettosize;
refcount = static_cast<int*>(fastMalloc(sizeof(*refcount)));
*refcount = 1;
}
#endif
}
void cv::gpu::GpuMat::release()
{
#ifdef HAVE_CUDA
if (refcount && CV_XADD(refcount, -1) == 1)
{
cudaFree(datastart);
fastFree(refcount);
}
data = datastart = dataend = 0;
step = rows = cols = 0;
refcount = 0;
#endif
}
void cv::gpu::GpuMat::upload(InputArray arr)
{
#ifndef HAVE_CUDA
(void) arr;
throw_no_cuda();
#else
Mat mat = arr.getMat();
CV_DbgAssert( !mat.empty() );
create(mat.size(), mat.type());
cudaSafeCall( cudaMemcpy2D(data, step, mat.data, mat.step, cols * elemSize(), rows, cudaMemcpyHostToDevice) );
#endif
}
void cv::gpu::GpuMat::upload(InputArray arr, Stream& _stream)
{
#ifndef HAVE_CUDA
(void) arr;
(void) _stream;
throw_no_cuda();
#else
Mat mat = arr.getMat();
CV_DbgAssert( !mat.empty() );
create(mat.size(), mat.type());
cudaStream_t stream = StreamAccessor::getStream(_stream);
cudaSafeCall( cudaMemcpy2DAsync(data, step, mat.data, mat.step, cols * elemSize(), rows, cudaMemcpyHostToDevice, stream) );
#endif
}
void cv::gpu::GpuMat::download(OutputArray _dst) const
{
#ifndef HAVE_CUDA
(void) _dst;
throw_no_cuda();
#else
CV_DbgAssert( !empty() );
_dst.create(size(), type());
Mat dst = _dst.getMat();
cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToHost) );
#endif
}
void cv::gpu::GpuMat::download(OutputArray _dst, Stream& _stream) const
{
#ifndef HAVE_CUDA
(void) _dst;
(void) _stream;
throw_no_cuda();
#else
CV_DbgAssert( !empty() );
_dst.create(size(), type());
Mat dst = _dst.getMat();
cudaStream_t stream = StreamAccessor::getStream(_stream);
cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToHost, stream) );
#endif
}
void cv::gpu::GpuMat::copyTo(OutputArray _dst) const
{
#ifndef HAVE_CUDA
(void) _dst;
throw_no_cuda();
#else
CV_DbgAssert( !empty() );
_dst.create(size(), type());
GpuMat dst = _dst.getGpuMat();
cudaSafeCall( cudaMemcpy2D(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToDevice) );
#endif
}
void cv::gpu::GpuMat::copyTo(OutputArray _dst, Stream& _stream) const
{
#ifndef HAVE_CUDA
(void) _dst;
(void) _stream;
throw_no_cuda();
#else
CV_DbgAssert( !empty() );
_dst.create(size(), type());
GpuMat dst = _dst.getGpuMat();
cudaStream_t stream = StreamAccessor::getStream(_stream);
cudaSafeCall( cudaMemcpy2DAsync(dst.data, dst.step, data, step, cols * elemSize(), rows, cudaMemcpyDeviceToDevice, stream) );
#endif
}
void cv::gpu::GpuMat::copyTo(OutputArray _dst, InputArray _mask, Stream& _stream) const
{
#ifndef HAVE_CUDA
(void) _dst;
(void) _mask;
(void) _stream;
throw_no_cuda();
#else
CV_DbgAssert( !empty() );
_dst.create(size(), type());
GpuMat dst = _dst.getGpuMat();
GpuMat mask = _mask.getGpuMat();
cudaStream_t stream = StreamAccessor::getStream(_stream);
::copyWithMask(*this, dst, mask, stream);
#endif
}
GpuMat& cv::gpu::GpuMat::setTo(Scalar s, Stream& _stream)
{
#ifndef HAVE_CUDA
(void) s;
(void) _stream;
throw_no_cuda();
#else
CV_DbgAssert( !empty() );
cudaStream_t stream = StreamAccessor::getStream(_stream);
::set(*this, s, stream);
#endif
return *this;
}
GpuMat& cv::gpu::GpuMat::setTo(Scalar s, InputArray _mask, Stream& _stream)
{
#ifndef HAVE_CUDA
(void) s;
(void) _mask;
(void) _stream;
throw_no_cuda();
#else
CV_DbgAssert( !empty() );
GpuMat mask = _mask.getGpuMat();
cudaStream_t stream = StreamAccessor::getStream(_stream);
::set(*this, s, mask, stream);
#endif
return *this;
}
void cv::gpu::GpuMat::convertTo(OutputArray _dst, int rtype, Stream& _stream) const
{
#ifndef HAVE_CUDA
(void) _dst;
(void) rtype;
(void) _stream;
throw_no_cuda();
#else
if (rtype < 0)
rtype = type();
else
rtype = CV_MAKETYPE(CV_MAT_DEPTH(rtype), channels());
const int sdepth = depth();
const int ddepth = CV_MAT_DEPTH(rtype);
if (sdepth == ddepth)
{
if (_stream)
copyTo(_dst, _stream);
else
copyTo(_dst);
return;
}
GpuMat src = *this;
_dst.create(size(), rtype);
GpuMat dst = _dst.getGpuMat();
cudaStream_t stream = StreamAccessor::getStream(_stream);
::convert(src, dst, stream);
#endif
}
void cv::gpu::GpuMat::convertTo(OutputArray _dst, int rtype, double alpha, double beta, Stream& _stream) const
{
#ifndef HAVE_CUDA
(void) _dst;
(void) rtype;
(void) alpha;
(void) beta;
(void) _stream;
throw_no_cuda();
#else
if (rtype < 0)
rtype = type();
else
rtype = CV_MAKETYPE(CV_MAT_DEPTH(rtype), channels());
GpuMat src = *this;
_dst.create(size(), rtype);
GpuMat dst = _dst.getGpuMat();
cudaStream_t stream = StreamAccessor::getStream(_stream);
::convert(src, dst, alpha, beta, stream);
#endif
}
GpuMat cv::gpu::GpuMat::reshape(int new_cn, int new_rows) const
{
GpuMat hdr = *this;
@ -1124,3 +344,101 @@ GpuMat cv::gpu::allocMatFromBuf(int rows, int cols, int type, GpuMat& mat)
return mat = GpuMat(rows, cols, type);
}
#ifndef HAVE_CUDA
void cv::gpu::GpuMat::create(int _rows, int _cols, int _type)
{
(void) _rows;
(void) _cols;
(void) _type;
throw_no_cuda();
}
void cv::gpu::GpuMat::release()
{
}
void cv::gpu::GpuMat::upload(InputArray arr)
{
(void) arr;
throw_no_cuda();
}
void cv::gpu::GpuMat::upload(InputArray arr, Stream& _stream)
{
(void) arr;
(void) _stream;
throw_no_cuda();
}
void cv::gpu::GpuMat::download(OutputArray _dst) const
{
(void) _dst;
throw_no_cuda();
}
void cv::gpu::GpuMat::download(OutputArray _dst, Stream& _stream) const
{
(void) _dst;
(void) _stream;
throw_no_cuda();
}
void cv::gpu::GpuMat::copyTo(OutputArray _dst) const
{
(void) _dst;
throw_no_cuda();
}
void cv::gpu::GpuMat::copyTo(OutputArray _dst, Stream& _stream) const
{
(void) _dst;
(void) _stream;
throw_no_cuda();
}
void cv::gpu::GpuMat::copyTo(OutputArray _dst, InputArray _mask, Stream& _stream) const
{
(void) _dst;
(void) _mask;
(void) _stream;
throw_no_cuda();
}
GpuMat& cv::gpu::GpuMat::setTo(Scalar s, Stream& _stream)
{
(void) s;
(void) _stream;
throw_no_cuda();
return *this;
}
GpuMat& cv::gpu::GpuMat::setTo(Scalar s, InputArray _mask, Stream& _stream)
{
(void) s;
(void) _mask;
(void) _stream;
throw_no_cuda();
return *this;
}
void cv::gpu::GpuMat::convertTo(OutputArray _dst, int rtype, Stream& _stream) const
{
(void) _dst;
(void) rtype;
(void) _stream;
throw_no_cuda();
}
void cv::gpu::GpuMat::convertTo(OutputArray _dst, int rtype, double alpha, double beta, Stream& _stream) const
{
(void) _dst;
(void) rtype;
(void) alpha;
(void) beta;
(void) _stream;
throw_no_cuda();
}
#endif

189
modules/cudev/include/opencv2/cudev/grid/copy.hpp
View File

@ -50,6 +50,7 @@
#include "../util/tuple.hpp"
#include "../ptr2d/traits.hpp"
#include "../ptr2d/gpumat.hpp"
#include "../ptr2d/glob.hpp"
#include "../ptr2d/mask.hpp"
#include "../ptr2d/zip.hpp"
#include "detail/copy.hpp"
@ -69,6 +70,18 @@ __host__ void gridCopy_(const SrcPtr& src, GpuMat_<DstType>& dst, const MaskPtr&
grid_copy_detail::copy<Policy>(shrinkPtr(src), shrinkPtr(dst), shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename DstType, class MaskPtr>
__host__ void gridCopy_(const SrcPtr& src, const GlobPtrSz<DstType>& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(dst) == rows && getCols(dst) == cols );
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
grid_copy_detail::copy<Policy>(shrinkPtr(src), shrinkPtr(dst), shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename DstType>
__host__ void gridCopy_(const SrcPtr& src, GpuMat_<DstType>& dst, Stream& stream = Stream::Null())
{
@ -80,6 +93,17 @@ __host__ void gridCopy_(const SrcPtr& src, GpuMat_<DstType>& dst, Stream& stream
grid_copy_detail::copy<Policy>(shrinkPtr(src), shrinkPtr(dst), WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename DstType>
__host__ void gridCopy_(const SrcPtr& src, const GlobPtrSz<DstType>& dst, Stream& stream = Stream::Null())
{
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(dst) == rows && getCols(dst) == cols );
grid_copy_detail::copy<Policy>(shrinkPtr(src), shrinkPtr(dst), WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, class MaskPtr>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
@ -100,6 +124,25 @@ __host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, class MaskPtr>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1> >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<SrcPtrTuple>::value == 2, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
grid_copy_detail::copy_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst))),
shrinkPtr(mask),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, Stream& stream = Stream::Null())
{
@ -118,6 +161,24 @@ __host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1> >& dst, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<SrcPtrTuple>::value == 2, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
grid_copy_detail::copy_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst))),
WithOutMask(),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, typename D2, class MaskPtr>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>& >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
@ -139,6 +200,26 @@ __host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, typename D2, class MaskPtr>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2> >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<SrcPtrTuple>::value == 3, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(get<2>(dst)) == rows && getCols(get<2>(dst)) == cols );
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
grid_copy_detail::copy_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst), get<2>(dst))),
shrinkPtr(mask),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, typename D2>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>& >& dst, Stream& stream = Stream::Null())
{
@ -158,6 +239,25 @@ __host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, typename D2>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2> >& dst, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<SrcPtrTuple>::value == 3, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(get<2>(dst)) == rows && getCols(get<2>(dst)) == cols );
grid_copy_detail::copy_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst), get<2>(dst))),
WithOutMask(),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, typename D2, typename D3, class MaskPtr>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>&, GpuMat_<D3>& >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
@ -180,6 +280,27 @@ __host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, typename D2, typename D3, class MaskPtr>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2>, GlobPtrSz<D3> >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<SrcPtrTuple>::value == 4, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(get<2>(dst)) == rows && getCols(get<2>(dst)) == cols );
CV_Assert( getRows(get<3>(dst)) == rows && getCols(get<3>(dst)) == cols );
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
grid_copy_detail::copy_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst), get<2>(dst), get<3>(dst))),
shrinkPtr(mask),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, typename D2, typename D3>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>&, GpuMat_<D3>& >& dst, Stream& stream = Stream::Null())
{
@ -200,6 +321,26 @@ __host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtrTuple, typename D0, typename D1, typename D2, typename D3>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2>, GlobPtrSz<D3> >& dst, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<SrcPtrTuple>::value == 4, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(get<2>(dst)) == rows && getCols(get<2>(dst)) == cols );
CV_Assert( getRows(get<3>(dst)) == rows && getCols(get<3>(dst)) == cols );
grid_copy_detail::copy_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst), get<2>(dst), get<3>(dst))),
WithOutMask(),
rows, cols,
StreamAccessor::getStream(stream));
}
// Default Policy
struct DefaultCopyPolicy
@ -216,48 +357,96 @@ __host__ void gridCopy(const SrcPtr& src, GpuMat_<DstType>& dst, const MaskPtr&
gridCopy_<DefaultCopyPolicy>(src, dst, mask, stream);
}
template <class SrcPtr, typename DstType, class MaskPtr>
__host__ void gridCopy(const SrcPtr& src, const GlobPtrSz<DstType>& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, mask, stream);
}
template <class SrcPtr, typename DstType>
__host__ void gridCopy(const SrcPtr& src, GpuMat_<DstType>& dst, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, stream);
}
template <class SrcPtr, typename DstType>
__host__ void gridCopy(const SrcPtr& src, const GlobPtrSz<DstType>& dst, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, class MaskPtr>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, mask, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, class MaskPtr>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1> >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, mask, stream);
}
template <class SrcPtrTuple, typename D0, typename D1>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, stream);
}
template <class SrcPtrTuple, typename D0, typename D1>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1> >& dst, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, typename D2, class MaskPtr>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>& >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, mask, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, typename D2, class MaskPtr>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2> >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, mask, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, typename D2>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>& >& dst, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, typename D2>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2> >& dst, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, typename D2, typename D3, class MaskPtr>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>&, GpuMat_<D3>& >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, mask, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, typename D2, typename D3, class MaskPtr>
__host__ void gridCopy(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2>, GlobPtrSz<D3> >& dst, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, mask, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, typename D2, typename D3>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>&, GpuMat_<D3>& >& dst, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, stream);
}
template <class SrcPtrTuple, typename D0, typename D1, typename D2, typename D3>
__host__ void gridCopy_(const SrcPtrTuple& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2>, GlobPtrSz<D3> >& dst, Stream& stream = Stream::Null())
{
gridCopy_<DefaultCopyPolicy>(src, dst, stream);
}
}}
#endif

232
modules/cudev/include/opencv2/cudev/grid/transform.hpp
View File

@ -50,6 +50,7 @@
#include "../util/tuple.hpp"
#include "../ptr2d/traits.hpp"
#include "../ptr2d/gpumat.hpp"
#include "../ptr2d/glob.hpp"
#include "../ptr2d/mask.hpp"
#include "../ptr2d/zip.hpp"
#include "detail/transform.hpp"
@ -69,6 +70,18 @@ __host__ void gridTransform_(const SrcPtr& src, GpuMat_<DstType>& dst, const UnO
grid_transform_detail::transform<Policy>(shrinkPtr(src), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename DstType, class UnOp, class MaskPtr>
__host__ void gridTransform_(const SrcPtr& src, const GlobPtrSz<DstType>& dst, const UnOp& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(dst) == rows && getCols(dst) == cols );
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
grid_transform_detail::transform<Policy>(shrinkPtr(src), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename DstType, class UnOp>
__host__ void gridTransform_(const SrcPtr& src, GpuMat_<DstType>& dst, const UnOp& op, Stream& stream = Stream::Null())
{
@ -80,6 +93,17 @@ __host__ void gridTransform_(const SrcPtr& src, GpuMat_<DstType>& dst, const UnO
grid_transform_detail::transform<Policy>(shrinkPtr(src), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename DstType, class UnOp>
__host__ void gridTransform_(const SrcPtr& src, const GlobPtrSz<DstType>& dst, const UnOp& op, Stream& stream = Stream::Null())
{
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(dst) == rows && getCols(dst) == cols );
grid_transform_detail::transform<Policy>(shrinkPtr(src), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType, class BinOp, class MaskPtr>
__host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_<DstType>& dst, const BinOp& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
@ -94,6 +118,19 @@ __host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_<D
grid_transform_detail::transform<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType, class BinOp, class MaskPtr>
__host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz<DstType>& dst, const BinOp& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
const int rows = getRows(src1);
const int cols = getCols(src1);
CV_Assert( getRows(dst) == rows && getCols(dst) == cols );
CV_Assert( getRows(src2) == rows && getCols(src2) == cols );
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
grid_transform_detail::transform<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType, class BinOp>
__host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_<DstType>& dst, const BinOp& op, Stream& stream = Stream::Null())
{
@ -107,6 +144,18 @@ __host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_<D
grid_transform_detail::transform<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr1, class SrcPtr2, typename DstType, class BinOp>
__host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GlobPtrSz<DstType>& dst, const BinOp& op, Stream& stream = Stream::Null())
{
const int rows = getRows(src1);
const int cols = getCols(src1);
CV_Assert( getRows(dst) == rows && getCols(dst) == cols );
CV_Assert( getRows(src2) == rows && getCols(src2) == cols );
grid_transform_detail::transform<Policy>(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, class OpTuple, class MaskPtr>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
@ -128,6 +177,26 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, class OpTuple, class MaskPtr>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1> >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<OpTuple>::value == 2, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
grid_transform_detail::transform_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst))),
op,
shrinkPtr(mask),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, class OpTuple>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
@ -147,6 +216,25 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, class OpTuple>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1> >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<OpTuple>::value == 2, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
grid_transform_detail::transform_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst))),
op,
WithOutMask(),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, typename D2, class OpTuple, class MaskPtr>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
@ -169,6 +257,27 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, typename D2, class OpTuple, class MaskPtr>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2> >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<OpTuple>::value == 3, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(get<2>(dst)) == rows && getCols(get<2>(dst)) == cols );
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
grid_transform_detail::transform_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst), get<2>(dst))),
op,
shrinkPtr(mask),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, typename D2, class OpTuple>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>& >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
@ -189,6 +298,26 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, typename D2, class OpTuple>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2> >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<OpTuple>::value == 3, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(get<2>(dst)) == rows && getCols(get<2>(dst)) == cols );
grid_transform_detail::transform_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst), get<2>(dst))),
op,
WithOutMask(),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, typename D2, typename D3, class OpTuple, class MaskPtr>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>&, GpuMat_<D3>& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
@ -212,6 +341,28 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, typename D2, typename D3, class OpTuple, class MaskPtr>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2>, GlobPtrSz<D3> >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<OpTuple>::value == 4, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(get<2>(dst)) == rows && getCols(get<2>(dst)) == cols );
CV_Assert( getRows(get<3>(dst)) == rows && getCols(get<3>(dst)) == cols );
CV_Assert( getRows(mask) == rows && getCols(mask) == cols );
grid_transform_detail::transform_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst), get<2>(dst), get<3>(dst))),
op,
shrinkPtr(mask),
rows, cols,
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, typename D2, typename D3, class OpTuple>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>&, GpuMat_<D3>& >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
@ -233,6 +384,27 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMa
StreamAccessor::getStream(stream));
}
template <class Policy, class SrcPtr, typename D0, typename D1, typename D2, typename D3, class OpTuple>
__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2>, GlobPtrSz<D3> >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
CV_StaticAssert( tuple_size<OpTuple>::value == 4, "" );
const int rows = getRows(src);
const int cols = getCols(src);
CV_Assert( getRows(get<0>(dst)) == rows && getCols(get<0>(dst)) == cols );
CV_Assert( getRows(get<1>(dst)) == rows && getCols(get<1>(dst)) == cols );
CV_Assert( getRows(get<2>(dst)) == rows && getCols(get<2>(dst)) == cols );
CV_Assert( getRows(get<3>(dst)) == rows && getCols(get<3>(dst)) == cols );
grid_transform_detail::transform_tuple<Policy>(shrinkPtr(src),
shrinkPtr(zipPtr(get<0>(dst), get<1>(dst), get<2>(dst), get<3>(dst))),
op,
WithOutMask(),
rows, cols,
StreamAccessor::getStream(stream));
}
// Default Policy
struct DefaultTransformPolicy
@ -250,60 +422,120 @@ __host__ void gridTransform(const SrcPtr& src, GpuMat_<DstType>& dst, const Op&
gridTransform_<DefaultTransformPolicy>(src, dst, op, mask, stream);
}
template <class SrcPtr, typename DstType, class Op, class MaskPtr>
__host__ void gridTransform(const SrcPtr& src, const GlobPtrSz<DstType>& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, mask, stream);
}
template <class SrcPtr, typename DstType, class Op>
__host__ void gridTransform(const SrcPtr& src, GpuMat_<DstType>& dst, const Op& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, stream);
}
template <class SrcPtr, typename DstType, class Op>
__host__ void gridTransform(const SrcPtr& src, const GlobPtrSz<DstType>& dst, const Op& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, stream);
}
template <class SrcPtr1, class SrcPtr2, typename DstType, class Op, class MaskPtr>
__host__ void gridTransform(const SrcPtr1& src1, const SrcPtr1& src2, GpuMat_<DstType>& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src1, src2, dst, op, mask, stream);
}
template <class SrcPtr1, class SrcPtr2, typename DstType, class Op, class MaskPtr>
__host__ void gridTransform(const SrcPtr1& src1, const SrcPtr1& src2, const GlobPtrSz<DstType>& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src1, src2, dst, op, mask, stream);
}
template <class SrcPtr1, class SrcPtr2, typename DstType, class Op>
__host__ void gridTransform(const SrcPtr1& src1, const SrcPtr1& src2, GpuMat_<DstType>& dst, const Op& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src1, src2, dst, op, stream);
}
template <class SrcPtr1, class SrcPtr2, typename DstType, class Op>
__host__ void gridTransform(const SrcPtr1& src1, const SrcPtr1& src2, const GlobPtrSz<DstType>& dst, const Op& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src1, src2, dst, op, stream);
}
template <class SrcPtr, typename D0, typename D1, class OpTuple, class MaskPtr>
__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, mask, stream);
}
template <class SrcPtr, typename D0, typename D1, class OpTuple, class MaskPtr>
__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1> >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, mask, stream);
}
template <class SrcPtr, typename D0, typename D1, class OpTuple>
__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>& >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, stream);
}
template <class SrcPtr, typename D0, typename D1, class OpTuple>
__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1> >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, stream);
}
template <class SrcPtr, typename D0, typename D1, typename D2, class OpTuple, class MaskPtr>
__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, mask, stream);
}
template <class SrcPtr, typename D0, typename D1, typename D2, class OpTuple, class MaskPtr>
__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2> >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, mask, stream);
}
template <class SrcPtr, typename D0, typename D1, typename D2, class OpTuple>
__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>& >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, stream);
}
template <class SrcPtr, typename D0, typename D1, typename D2, class OpTuple>
__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2> >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, stream);
}
template <class SrcPtr, typename D0, typename D1, typename D2, typename D3, class OpTuple, class MaskPtr>
__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>&, GpuMat_<D3>& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, mask, stream);
}
template <class SrcPtr, typename D0, typename D1, typename D2, typename D3, class OpTuple, class MaskPtr>
__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2>, GlobPtrSz<D3> >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, mask, stream);
}
template <class SrcPtr, typename D0, typename D1, typename D2, typename D3, class OpTuple>
__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_<D0>&, GpuMat_<D1>&, GpuMat_<D2>&, GpuMat_<D3>& >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, stream);
}
template <class SrcPtr, typename D0, typename D1, typename D2, typename D3, class OpTuple>
__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz<D0>, GlobPtrSz<D1>, GlobPtrSz<D2>, GlobPtrSz<D3> >& dst, const OpTuple& op, Stream& stream = Stream::Null())
{
gridTransform_<DefaultTransformPolicy>(src, dst, op, stream);
}
}}
#endif

23
modules/cudev/include/opencv2/cudev/ptr2d/detail/gpumat.hpp
View File

@ -335,4 +335,27 @@ __host__ GpuMat_<T>& GpuMat_<T>::assign(const Expr<Body>& expr, Stream& stream)
}}
// Input / Output Arrays
namespace cv {
template<typename _Tp>
__host__ _InputArray::_InputArray(const cudev::GpuMat_<_Tp>& m)
: flags(FIXED_TYPE + GPU_MAT + DataType<_Tp>::type), obj((void*)&m)
{}
template<typename _Tp>
__host__ _OutputArray::_OutputArray(cudev::GpuMat_<_Tp>& m)
: _InputArray(m)
{}
template<typename _Tp>
__host__ _OutputArray::_OutputArray(const cudev::GpuMat_<_Tp>& m)
: _InputArray(m)
{
flags |= FIXED_SIZE;
}
}
#endif

11
modules/cudev/include/opencv2/cudev/ptr2d/glob.hpp
View File

@ -91,6 +91,17 @@ __host__ GlobPtrSz<T> globPtr(T* data, size_t step, int rows, int cols)
return p;
}
template <typename T>
__host__ GlobPtrSz<T> globPtr(const GpuMat& mat)
{
GlobPtrSz<T> p;
p.data = (T*) mat.data;
p.step = mat.step;
p.rows = mat.rows;
p.cols = mat.cols;
return p;
}
template <typename T> struct PtrTraits< GlobPtrSz<T> > : PtrTraitsBase<GlobPtrSz<T>, GlobPtr<T> >
{
};

12
modules/gpufilters/src/filtering.cpp
View File

@ -230,22 +230,22 @@ namespace
switch (srcType)
{
case CV_8UC1:
func_ = cudev::filter2D<uchar, uchar>;
func_ = cv::gpu::cudev::filter2D<uchar, uchar>;
break;
case CV_8UC4:
func_ = cudev::filter2D<uchar4, uchar4>;
func_ = cv::gpu::cudev::filter2D<uchar4, uchar4>;
break;
case CV_16UC1:
func_ = cudev::filter2D<ushort, ushort>;
func_ = cv::gpu::cudev::filter2D<ushort, ushort>;
break;
case CV_16UC4:
func_ = cudev::filter2D<ushort4, ushort4>;
func_ = cv::gpu::cudev::filter2D<ushort4, ushort4>;
break;
case CV_32FC1:
func_ = cudev::filter2D<float, float>;
func_ = cv::gpu::cudev::filter2D<float, float>;
break;
case CV_32FC4:
func_ = cudev::filter2D<float4, float4>;
func_ = cv::gpu::cudev::filter2D<float4, float4>;
break;
}
}

93
modules/gpuimgproc/doc/hough.rst
View File

@ -216,98 +216,19 @@ Creates implementation for :ocv:class:`gpu::HoughCirclesDetector` .
gpu::GeneralizedHough
---------------------
.. ocv:class:: gpu::GeneralizedHough : public Algorithm
Base class for generalized hough transform. ::
class CV_EXPORTS GeneralizedHough : public Algorithm
{
public:
static Ptr<GeneralizedHough> create(int method);
virtual void setTemplate(InputArray templ, int cannyThreshold = 100, Point templCenter = Point(-1, -1)) = 0;
virtual void setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter = Point(-1, -1)) = 0;
virtual void detect(InputArray image, OutputArray positions, int cannyThreshold = 100) = 0;
virtual void detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions) = 0;
virtual void downloadResults(InputArray d_positions, OutputArray h_positions, OutputArray h_votes = noArray()) = 0;
};
Finds arbitrary template in the grayscale image using Generalized Hough Transform.
gpu::GeneralizedHough::create
-----------------------------
Creates implementation for :ocv:class:`gpu::GeneralizedHough` .
.. ocv:function:: Ptr<GeneralizedHough> gpu::GeneralizedHough::create(int method)
:param method: Combination of flags ( ``cv::GeneralizedHough::GHT_POSITION`` , ``cv::GeneralizedHough::GHT_SCALE`` , ``cv::GeneralizedHough::GHT_ROTATION`` ) specifying transformation to find.
For full affine transformations (move + scale + rotation) [Guil1999]_ algorithm is used, otherwise [Ballard1981]_ algorithm is used.
gpu::GeneralizedHough::setTemplate
gpu::createGeneralizedHoughBallard
----------------------------------
Set template to search.
Creates implementation for generalized hough transform from [Ballard1981]_ .
.. ocv:function:: void gpu::GeneralizedHough::setTemplate(InputArray templ, int cannyThreshold = 100, Point templCenter = Point(-1, -1))
.. ocv:function:: void gpu::GeneralizedHough::setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter = Point(-1, -1))
:param templ: Template image. Canny edge detector will be applied to extract template edges.
:param cannyThreshold: Threshold value for Canny edge detector.
:param templCenter: Center for rotation. By default image center will be used.
:param edges: Edge map for template image.
:param dx: First derivative of template image in the vertical direction. Support only ``CV_32S`` type.
:param dy: First derivative of template image in the horizontal direction. Support only ``CV_32S`` type.
.. ocv:function:: Ptr<GeneralizedHoughBallard> gpu::createGeneralizedHoughBallard()
gpu::GeneralizedHough::detect
-----------------------------
Finds template (set by :ocv:func:`gpu::GeneralizedHough::setTemplate` ) in the grayscale image.
gpu::createGeneralizedHoughGuil
-------------------------------
Creates implementation for generalized hough transform from [Guil1999]_ .
.. ocv:function:: void gpu::GeneralizedHough::detect(InputArray image, OutputArray positions, int cannyThreshold = 100)
.. ocv:function:: void gpu::GeneralizedHough::detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions)
:param templ: Input image. Canny edge detector will be applied to extract template edges.
:param positions: Output vector of found objects. Each vector is encoded as a 4-element floating-point vector :math:`(x, y, scale, angle)` .
:param cannyThreshold: Threshold value for Canny edge detector.
:param edges: Edge map for input image.
:param dx: First derivative of input image in the vertical direction. Support only ``CV_32S`` type.
:param dy: First derivative of input image in the horizontal direction. Support only ``CV_32S`` type.
gpu::GeneralizedHough::downloadResults
--------------------------------------
Downloads results from :ocv:func:`gpu::GeneralizedHough::detect` to host memory.
.. ocv:function:: void gpu::GeneralizedHough::downloadResult(InputArray d_positions, OutputArray h_positions, OutputArray h_votes = noArray())
:param d_lines: Result of :ocv:func:`gpu::GeneralizedHough::detect` .
:param h_lines: Output host array.
:param h_votes: Optional output array for votes. Each vector is encoded as a 3-element integer-point vector :math:`(position_votes, scale_votes, angle_votes)` .
.. ocv:function:: Ptr<GeneralizedHoughGuil> gpu::createGeneralizedHoughGuil()

21
modules/gpuimgproc/include/opencv2/gpuimgproc.hpp
View File

@ -283,24 +283,13 @@ CV_EXPORTS Ptr<HoughCirclesDetector> createHoughCirclesDetector(float dp, float
//////////////////////////////////////
// GeneralizedHough
//! finds arbitrary template in the grayscale image using Generalized Hough Transform
//! Ballard, D.H. (1981). Generalizing the Hough transform to detect arbitrary shapes. Pattern Recognition 13 (2): 111-122.
//! Detects position only without traslation and rotation
CV_EXPORTS Ptr<GeneralizedHoughBallard> createGeneralizedHoughBallard();
//! Guil, N., González-Linares, J.M. and Zapata, E.L. (1999). Bidimensional shape detection using an invariant approach. Pattern Recognition 32 (6): 1025-1038.
class CV_EXPORTS GeneralizedHough : public Algorithm
{
public:
static Ptr<GeneralizedHough> create(int method);
//! set template to search
virtual void setTemplate(InputArray templ, int cannyThreshold = 100, Point templCenter = Point(-1, -1)) = 0;
virtual void setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter = Point(-1, -1)) = 0;
//! find template on image
virtual void detect(InputArray image, OutputArray positions, int cannyThreshold = 100) = 0;
virtual void detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions) = 0;
virtual void downloadResults(InputArray d_positions, OutputArray h_positions, OutputArray h_votes = noArray()) = 0;
};
//! Detects position, traslation and rotation
CV_EXPORTS Ptr<GeneralizedHoughGuil> createGeneralizedHoughGuil();
////////////////////////// Corners Detection ///////////////////////////

97
modules/gpuimgproc/perf/perf_hough.cpp
View File

@ -227,23 +227,59 @@ PERF_TEST_P(Sz_Dp_MinDist, HoughCircles,
//////////////////////////////////////////////////////////////////////
// GeneralizedHough
enum { GHT_POSITION = cv::GeneralizedHough::GHT_POSITION,
GHT_SCALE = cv::GeneralizedHough::GHT_SCALE,
GHT_ROTATION = cv::GeneralizedHough::GHT_ROTATION
};
CV_FLAGS(GHMethod, GHT_POSITION, GHT_SCALE, GHT_ROTATION);
DEF_PARAM_TEST(Method_Sz, GHMethod, cv::Size);
PERF_TEST_P(Method_Sz, GeneralizedHough,
Combine(Values(GHMethod(GHT_POSITION), GHMethod(GHT_POSITION | GHT_SCALE), GHMethod(GHT_POSITION | GHT_ROTATION), GHMethod(GHT_POSITION | GHT_SCALE | GHT_ROTATION)),
GPU_TYPICAL_MAT_SIZES))
PERF_TEST_P(Sz, GeneralizedHoughBallard, GPU_TYPICAL_MAT_SIZES)
{
declare.time(10);
const int method = GET_PARAM(0);
const cv::Size imageSize = GET_PARAM(1);
const cv::Size imageSize = GetParam();
const cv::Mat templ = readImage("cv/shared/templ.png", cv::IMREAD_GRAYSCALE);
ASSERT_FALSE(templ.empty());
cv::Mat image(imageSize, CV_8UC1, cv::Scalar::all(0));
templ.copyTo(image(cv::Rect(50, 50, templ.cols, templ.rows)));
cv::Mat edges;
cv::Canny(image, edges, 50, 100);
cv::Mat dx, dy;
cv::Sobel(image, dx, CV_32F, 1, 0);
cv::Sobel(image, dy, CV_32F, 0, 1);
if (PERF_RUN_GPU())
{
cv::Ptr<cv::GeneralizedHoughBallard> alg = cv::gpu::createGeneralizedHoughBallard();
const cv::gpu::GpuMat d_edges(edges);
const cv::gpu::GpuMat d_dx(dx);
const cv::gpu::GpuMat d_dy(dy);
cv::gpu::GpuMat positions;
alg->setTemplate(cv::gpu::GpuMat(templ));
TEST_CYCLE() alg->detect(d_edges, d_dx, d_dy, positions);
GPU_SANITY_CHECK(positions);
}
else
{
cv::Ptr<cv::GeneralizedHoughBallard> alg = cv::createGeneralizedHoughBallard();
cv::Mat positions;
alg->setTemplate(templ);
TEST_CYCLE() alg->detect(edges, dx, dy, positions);
CPU_SANITY_CHECK(positions);
}
}
PERF_TEST_P(Sz, GeneralizedHoughGuil, GPU_TYPICAL_MAT_SIZES)
{
declare.time(10);
const cv::Size imageSize = GetParam();
const cv::Mat templ = readImage("cv/shared/templ.png", cv::IMREAD_GRAYSCALE);
ASSERT_FALSE(templ.empty());
@ -281,39 +317,32 @@ PERF_TEST_P(Method_Sz, GeneralizedHough,
if (PERF_RUN_GPU())
{
cv::Ptr<cv::GeneralizedHoughGuil> alg = cv::gpu::createGeneralizedHoughGuil();
alg->setMaxAngle(90.0);
alg->setAngleStep(2.0);
const cv::gpu::GpuMat d_edges(edges);
const cv::gpu::GpuMat d_dx(dx);
const cv::gpu::GpuMat d_dy(dy);
cv::gpu::GpuMat posAndVotes;
cv::gpu::GpuMat positions;
cv::Ptr<cv::gpu::GeneralizedHough> d_hough = cv::gpu::GeneralizedHough::create(method);
if (method & GHT_ROTATION)
{
d_hough->set("maxAngle", 90.0);
d_hough->set("angleStep", 2.0);
}
alg->setTemplate(cv::gpu::GpuMat(templ));
d_hough->setTemplate(cv::gpu::GpuMat(templ));
TEST_CYCLE() alg->detect(d_edges, d_dx, d_dy, positions);
TEST_CYCLE() d_hough->detect(d_edges, d_dx, d_dy, posAndVotes);
const cv::gpu::GpuMat positions(1, posAndVotes.cols, CV_32FC4, posAndVotes.data);
GPU_SANITY_CHECK(positions);
}
else
{
cv::Ptr<cv::GeneralizedHoughGuil> alg = cv::createGeneralizedHoughGuil();
alg->setMaxAngle(90.0);
alg->setAngleStep(2.0);
cv::Mat positions;
cv::Ptr<cv::GeneralizedHough> hough = cv::GeneralizedHough::create(method);
if (method & GHT_ROTATION)
{
hough->set("maxAngle", 90.0);
hough->set("angleStep", 2.0);
}
alg->setTemplate(templ);
hough->setTemplate(templ);
TEST_CYCLE() hough->detect(edges, dx, dy, positions);
TEST_CYCLE() alg->detect(edges, dx, dy, positions);
CPU_SANITY_CHECK(positions);
}

46
modules/gpuimgproc/src/color.cpp
View File

@ -187,7 +187,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream));
}
void bgr_to_bgr565(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -200,7 +200,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream));
}
void rgb_to_bgr555(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -213,7 +213,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream));
}
void rgb_to_bgr565(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -226,7 +226,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream));
}
void bgra_to_bgr555(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -239,7 +239,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream));
}
void bgra_to_bgr565(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -252,7 +252,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream));
}
void rgba_to_bgr555(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -265,7 +265,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream));
}
void rgba_to_bgr565(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -278,7 +278,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream));
}
void bgr555_to_rgb(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -291,7 +291,7 @@ namespace
_dst.create(src.size(), CV_8UC3);
GpuMat dst = _dst.getGpuMat();
cudev::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream));
}
void bgr565_to_rgb(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -304,7 +304,7 @@ namespace
_dst.create(src.size(), CV_8UC3);
GpuMat dst = _dst.getGpuMat();
cudev::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream));
}
void bgr555_to_bgr(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -317,7 +317,7 @@ namespace
_dst.create(src.size(), CV_8UC3);
GpuMat dst = _dst.getGpuMat();
cudev::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream));
}
void bgr565_to_bgr(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -330,7 +330,7 @@ namespace
_dst.create(src.size(), CV_8UC3);
GpuMat dst = _dst.getGpuMat();
cudev::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream));
}
void bgr555_to_rgba(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -343,7 +343,7 @@ namespace
_dst.create(src.size(), CV_8UC4);
GpuMat dst = _dst.getGpuMat();
cudev::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream));
}
void bgr565_to_rgba(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -356,7 +356,7 @@ namespace
_dst.create(src.size(), CV_8UC4);
GpuMat dst = _dst.getGpuMat();
cudev::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream));
}
void bgr555_to_bgra(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -369,7 +369,7 @@ namespace
_dst.create(src.size(), CV_8UC4);
GpuMat dst = _dst.getGpuMat();
cudev::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream));
}
void bgr565_to_bgra(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -382,7 +382,7 @@ namespace
_dst.create(src.size(), CV_8UC4);
GpuMat dst = _dst.getGpuMat();
cudev::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream));
}
void gray_to_bgr(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -427,7 +427,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream));
}
void gray_to_bgr565(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -440,7 +440,7 @@ namespace
_dst.create(src.size(), CV_8UC2);
GpuMat dst = _dst.getGpuMat();
cudev::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream));
}
void bgr555_to_gray(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -453,7 +453,7 @@ namespace
_dst.create(src.size(), CV_8UC1);
GpuMat dst = _dst.getGpuMat();
cudev::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream));
}
void bgr565_to_gray(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -466,7 +466,7 @@ namespace
_dst.create(src.size(), CV_8UC1);
GpuMat dst = _dst.getGpuMat();
cudev::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream));
cv::gpu::cudev::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream));
}
void rgb_to_gray(InputArray _src, OutputArray _dst, int, Stream& stream)
@ -2145,9 +2145,9 @@ void cv::gpu::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn,
code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1);
if (dcn == 3)
cudev::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
cv::gpu::cudev::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
else
cudev::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
cv::gpu::cudev::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
break;
}
@ -2172,7 +2172,7 @@ void cv::gpu::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn,
const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1,
code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1);
cudev::MHCdemosaic<1>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
cv::gpu::cudev::MHCdemosaic<1>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
break;
}

262
modules/gpuimgproc/src/cuda/generalized_hough.cu
View File

@ -307,268 +307,6 @@ namespace cv { namespace gpu { namespace cudev
return totalCount;
}
////////////////////////////////////////////////////////////////////////
// Ballard_PosScale
__global__ void Ballard_PosScale_calcHist(const unsigned int* coordList, const float* thetaList,
PtrStep<short2> r_table, const int* r_sizes,
PtrStepi hist, const int rows, const int cols,
const float minScale, const float scaleStep, const int scaleRange,
const float idp, const float thetaScale)
{
const unsigned int coord = coordList[blockIdx.x];
float2 p;
p.x = (coord & 0xFFFF);
p.y = (coord >> 16) & 0xFFFF;
const float theta = thetaList[blockIdx.x];
const int n = __float2int_rn(theta * thetaScale);
const short2* r_row = r_table.ptr(n);
const int r_row_size = r_sizes[n];
for (int j = 0; j < r_row_size; ++j)
{
const float2 d = saturate_cast<float2>(r_row[j]);
for (int s = threadIdx.x; s < scaleRange; s += blockDim.x)
{
const float scale = minScale + s * scaleStep;
float2 c = p - scale * d;
c.x *= idp;
c.y *= idp;
if (c.x >= 0 && c.x < cols && c.y >= 0 && c.y < rows)
::atomicAdd(hist.ptr((s + 1) * (rows + 2) + __float2int_rn(c.y + 1)) + __float2int_rn(c.x + 1), 1);
}
}
}
void Ballard_PosScale_calcHist_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
PtrStepSz<short2> r_table, const int* r_sizes,
PtrStepi hist, int rows, int cols,
float minScale, float scaleStep, int scaleRange,
float dp, int levels)
{
const dim3 block(256);
const dim3 grid(pointsCount);
const float idp = 1.0f / dp;
const float thetaScale = levels / (2.0f * CV_PI_F);
Ballard_PosScale_calcHist<<<grid, block>>>(coordList, thetaList,
r_table, r_sizes,
hist, rows, cols,
minScale, scaleStep, scaleRange,
idp, thetaScale);
cudaSafeCall( cudaGetLastError() );
cudaSafeCall( cudaDeviceSynchronize() );
}
__global__ void Ballard_PosScale_findPosInHist(const PtrStepi hist, const int rows, const int cols, const int scaleRange,
float4* out, int3* votes, const int maxSize,
const float minScale, const float scaleStep, const float dp, const int threshold)
{
const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= cols || y >= rows)
return;
for (int s = 0; s < scaleRange; ++s)
{
const float scale = minScale + s * scaleStep;
const int prevScaleIdx = (s) * (rows + 2);
const int curScaleIdx = (s + 1) * (rows + 2);
const int nextScaleIdx = (s + 2) * (rows + 2);
const int curVotes = hist(curScaleIdx + y + 1, x + 1);
if (curVotes > threshold &&
curVotes > hist(curScaleIdx + y + 1, x) &&
curVotes >= hist(curScaleIdx + y + 1, x + 2) &&
curVotes > hist(curScaleIdx + y, x + 1) &&
curVotes >= hist(curScaleIdx + y + 2, x + 1) &&
curVotes > hist(prevScaleIdx + y + 1, x + 1) &&
curVotes >= hist(nextScaleIdx + y + 1, x + 1))
{
const int ind = ::atomicAdd(&g_counter, 1);
if (ind < maxSize)
{
out[ind] = make_float4(x * dp, y * dp, scale, 0.0f);
votes[ind] = make_int3(curVotes, curVotes, 0);
}
}
}
}
int Ballard_PosScale_findPosInHist_gpu(PtrStepi hist, int rows, int cols, int scaleRange, float4* out, int3* votes, int maxSize,
float minScale, float scaleStep, float dp, int threshold)
{
void* counterPtr;
cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
const dim3 block(32, 8);
const dim3 grid(divUp(cols, block.x), divUp(rows, block.y));
cudaSafeCall( cudaFuncSetCacheConfig(Ballard_PosScale_findPosInHist, cudaFuncCachePreferL1) );
Ballard_PosScale_findPosInHist<<<grid, block>>>(hist, rows, cols, scaleRange, out, votes,
maxSize, minScale, scaleStep, dp, threshold);
cudaSafeCall( cudaGetLastError() );
cudaSafeCall( cudaDeviceSynchronize() );
int totalCount;
cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
totalCount = ::min(totalCount, maxSize);
return totalCount;
}
////////////////////////////////////////////////////////////////////////
// Ballard_PosRotation
__global__ void Ballard_PosRotation_calcHist(const unsigned int* coordList, const float* thetaList,
PtrStep<short2> r_table, const int* r_sizes,
PtrStepi hist, const int rows, const int cols,
const float minAngle, const float angleStep, const int angleRange,
const float idp, const float thetaScale)
{
const unsigned int coord = coordList[blockIdx.x];
float2 p;
p.x = (coord & 0xFFFF);
p.y = (coord >> 16) & 0xFFFF;
const float thetaVal = thetaList[blockIdx.x];
for (int a = threadIdx.x; a < angleRange; a += blockDim.x)
{
const float angle = (minAngle + a * angleStep) * (CV_PI_F / 180.0f);
float sinA, cosA;
sincosf(angle, &sinA, &cosA);
float theta = thetaVal - angle;
if (theta < 0)
theta += 2.0f * CV_PI_F;
const int n = __float2int_rn(theta * thetaScale);
const short2* r_row = r_table.ptr(n);
const int r_row_size = r_sizes[n];
for (int j = 0; j < r_row_size; ++j)
{
const float2 d = saturate_cast<float2>(r_row[j]);
const float2 dr = make_float2(d.x * cosA - d.y * sinA, d.x * sinA + d.y * cosA);
float2 c = make_float2(p.x - dr.x, p.y - dr.y);
c.x *= idp;
c.y *= idp;
if (c.x >= 0 && c.x < cols && c.y >= 0 && c.y < rows)
::atomicAdd(hist.ptr((a + 1) * (rows + 2) + __float2int_rn(c.y + 1)) + __float2int_rn(c.x + 1), 1);
}
}
}
void Ballard_PosRotation_calcHist_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
PtrStepSz<short2> r_table, const int* r_sizes,
PtrStepi hist, int rows, int cols,
float minAngle, float angleStep, int angleRange,
float dp, int levels)
{
const dim3 block(256);
const dim3 grid(pointsCount);
const float idp = 1.0f / dp;
const float thetaScale = levels / (2.0f * CV_PI_F);
Ballard_PosRotation_calcHist<<<grid, block>>>(coordList, thetaList,
r_table, r_sizes,
hist, rows, cols,
minAngle, angleStep, angleRange,
idp, thetaScale);
cudaSafeCall( cudaGetLastError() );
cudaSafeCall( cudaDeviceSynchronize() );
}
__global__ void Ballard_PosRotation_findPosInHist(const PtrStepi hist, const int rows, const int cols, const int angleRange,
float4* out, int3* votes, const int maxSize,
const float minAngle, const float angleStep, const float dp, const int threshold)
{
const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= cols || y >= rows)
return;
for (int a = 0; a < angleRange; ++a)
{
const float angle = minAngle + a * angleStep;
const int prevAngleIdx = (a) * (rows + 2);
const int curAngleIdx = (a + 1) * (rows + 2);
const int nextAngleIdx = (a + 2) * (rows + 2);
const int curVotes = hist(curAngleIdx + y + 1, x + 1);
if (curVotes > threshold &&
curVotes > hist(curAngleIdx + y + 1, x) &&
curVotes >= hist(curAngleIdx + y + 1, x + 2) &&
curVotes > hist(curAngleIdx + y, x + 1) &&
curVotes >= hist(curAngleIdx + y + 2, x + 1) &&
curVotes > hist(prevAngleIdx + y + 1, x + 1) &&
curVotes >= hist(nextAngleIdx + y + 1, x + 1))
{
const int ind = ::atomicAdd(&g_counter, 1);
if (ind < maxSize)
{
out[ind] = make_float4(x * dp, y * dp, 1.0f, angle);
votes[ind] = make_int3(curVotes, 0, curVotes);
}
}
}
}
int Ballard_PosRotation_findPosInHist_gpu(PtrStepi hist, int rows, int cols, int angleRange, float4* out, int3* votes, int maxSize,
float minAngle, float angleStep, float dp, int threshold)
{
void* counterPtr;
cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
const dim3 block(32, 8);
const dim3 grid(divUp(cols, block.x), divUp(rows, block.y));
cudaSafeCall( cudaFuncSetCacheConfig(Ballard_PosRotation_findPosInHist, cudaFuncCachePreferL1) );
Ballard_PosRotation_findPosInHist<<<grid, block>>>(hist, rows, cols, angleRange, out, votes,
maxSize, minAngle, angleStep, dp, threshold);
cudaSafeCall( cudaGetLastError() );
cudaSafeCall( cudaDeviceSynchronize() );
int totalCount;
cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
totalCount = ::min(totalCount, maxSize);
return totalCount;
}
////////////////////////////////////////////////////////////////////////
// Guil_Full

1189
modules/gpuimgproc/src/generalized_hough.cpp
View File

File diff suppressed because it is too large Load Diff

12
modules/gpuimgproc/test/test_hough.cpp
View File

@ -193,7 +193,7 @@ PARAM_TEST_CASE(GeneralizedHough, cv::gpu::DeviceInfo, UseRoi)
{
};
GPU_TEST_P(GeneralizedHough, POSITION)
GPU_TEST_P(GeneralizedHough, Ballard)
{
const cv::gpu::DeviceInfo devInfo = GET_PARAM(0);
cv::gpu::setDevice(devInfo.deviceID());
@ -218,16 +218,16 @@ GPU_TEST_P(GeneralizedHough, POSITION)
templ.copyTo(imageROI);
}
cv::Ptr<cv::gpu::GeneralizedHough> hough = cv::gpu::GeneralizedHough::create(cv::GeneralizedHough::GHT_POSITION);
hough->set("votesThreshold", 200);
cv::Ptr<cv::GeneralizedHoughBallard> alg = cv::gpu::createGeneralizedHoughBallard();
alg->setVotesThreshold(200);
hough->setTemplate(loadMat(templ, useRoi));
alg->setTemplate(loadMat(templ, useRoi));
cv::gpu::GpuMat d_pos;
hough->detect(loadMat(image, useRoi), d_pos);
alg->detect(loadMat(image, useRoi), d_pos);
std::vector<cv::Vec4f> pos;
hough->downloadResults(d_pos, pos);
d_pos.download(pos);
ASSERT_EQ(gold_count, pos.size());

4
modules/gpuwarping/src/pyramids.cpp
View File

@ -181,7 +181,7 @@ namespace
const GpuMat& prevLayer = i == 0 ? layer0_ : pyramid_[i - 1];
cudev::pyramid::downsampleX2(prevLayer, pyramid_[i], img.depth(), img.channels(), StreamAccessor::getStream(stream));
cv::gpu::cudev::pyramid::downsampleX2(prevLayer, pyramid_[i], img.depth(), img.channels(), StreamAccessor::getStream(stream));
szLastLayer = szCurLayer;
}
@ -222,7 +222,7 @@ namespace
lastLayer = curLayer;
}
cudev::pyramid::interpolateFrom1(lastLayer, outImg, outImg.depth(), outImg.channels(), StreamAccessor::getStream(stream));
cv::gpu::cudev::pyramid::interpolateFrom1(lastLayer, outImg, outImg.depth(), outImg.channels(), StreamAccessor::getStream(stream));
}
}

121
modules/imgproc/include/opencv2/imgproc.hpp
View File

@ -698,39 +698,104 @@ public:
//! finds arbitrary template in the grayscale image using Generalized Hough Transform
//! Ballard, D.H. (1981). Generalizing the Hough transform to detect arbitrary shapes. Pattern Recognition 13 (2): 111-122.
//! Guil, N., González-Linares, J.M. and Zapata, E.L. (1999). Bidimensional shape detection using an invariant approach. Pattern Recognition 32 (6): 1025-1038.
class CV_EXPORTS GeneralizedHough : public Algorithm
{
public:
enum { GHT_POSITION = 0,
GHT_SCALE = 1,
GHT_ROTATION = 2
};
static Ptr<GeneralizedHough> create(int method);
virtual ~GeneralizedHough();
//! set template to search
void setTemplate(InputArray templ, int cannyThreshold = 100, Point templCenter = Point(-1, -1));
void setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter = Point(-1, -1));
virtual void setTemplate(InputArray templ, Point templCenter = Point(-1, -1)) = 0;
virtual void setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter = Point(-1, -1)) = 0;
//! find template on image
void detect(InputArray image, OutputArray positions, OutputArray votes = cv::noArray(), int cannyThreshold = 100);
void detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions, OutputArray votes = cv::noArray());
virtual void detect(InputArray image, OutputArray positions, OutputArray votes = noArray()) = 0;
virtual void detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions, OutputArray votes = noArray()) = 0;
void release();
//! Canny low threshold.
virtual void setCannyLowThresh(int cannyLowThresh) = 0;
virtual int getCannyLowThresh() const = 0;
protected:
virtual void setTemplateImpl(const Mat& edges, const Mat& dx, const Mat& dy, Point templCenter) = 0;
virtual void detectImpl(const Mat& edges, const Mat& dx, const Mat& dy, OutputArray positions, OutputArray votes) = 0;
virtual void releaseImpl() = 0;
//! Canny high threshold.
virtual void setCannyHighThresh(int cannyHighThresh) = 0;
virtual int getCannyHighThresh() const = 0;
private:
Mat edges_;
Mat dx_;
Mat dy_;
//! Minimum distance between the centers of the detected objects.
virtual void setMinDist(double minDist) = 0;
virtual double getMinDist() const = 0;
//! Inverse ratio of the accumulator resolution to the image resolution.
virtual void setDp(double dp) = 0;
virtual double getDp() const = 0;
//! Maximal size of inner buffers.
virtual void setMaxBufferSize(int maxBufferSize) = 0;
virtual int getMaxBufferSize() const = 0;
};
//! Ballard, D.H. (1981). Generalizing the Hough transform to detect arbitrary shapes. Pattern Recognition 13 (2): 111-122.
//! Detects position only without traslation and rotation
class CV_EXPORTS GeneralizedHoughBallard : public GeneralizedHough
{
public:
//! R-Table levels.
virtual void setLevels(int levels) = 0;
virtual int getLevels() const = 0;
//! The accumulator threshold for the template centers at the detection stage. The smaller it is, the more false positions may be detected.
virtual void setVotesThreshold(int votesThreshold) = 0;
virtual int getVotesThreshold() const = 0;
};
//! Guil, N., González-Linares, J.M. and Zapata, E.L. (1999). Bidimensional shape detection using an invariant approach. Pattern Recognition 32 (6): 1025-1038.
//! Detects position, traslation and rotation
class CV_EXPORTS GeneralizedHoughGuil : public GeneralizedHough
{
public:
//! Angle difference in degrees between two points in feature.
virtual void setXi(double xi) = 0;
virtual double getXi() const = 0;
//! Feature table levels.
virtual void setLevels(int levels) = 0;
virtual int getLevels() const = 0;
//! Maximal difference between angles that treated as equal.
virtual void setAngleEpsilon(double angleEpsilon) = 0;
virtual double getAngleEpsilon() const = 0;
//! Minimal rotation angle to detect in degrees.
virtual void setMinAngle(double minAngle) = 0;
virtual double getMinAngle() const = 0;
//! Maximal rotation angle to detect in degrees.
virtual void setMaxAngle(double maxAngle) = 0;
virtual double getMaxAngle() const = 0;
//! Angle step in degrees.
virtual void setAngleStep(double angleStep) = 0;
virtual double getAngleStep() const = 0;
//! Angle votes threshold.
virtual void setAngleThresh(int angleThresh) = 0;
virtual int getAngleThresh() const = 0;
//! Minimal scale to detect.
virtual void setMinScale(double minScale) = 0;
virtual double getMinScale() const = 0;
//! Maximal scale to detect.
virtual void setMaxScale(double maxScale) = 0;
virtual double getMaxScale() const = 0;
//! Scale step.
virtual void setScaleStep(double scaleStep) = 0;
virtual double getScaleStep() const = 0;
//! Scale votes threshold.
virtual void setScaleThresh(int scaleThresh) = 0;
virtual int getScaleThresh() const = 0;
//! Position votes threshold.
virtual void setPosThresh(int posThresh) = 0;
virtual int getPosThresh() const = 0;
};
@ -1423,6 +1488,14 @@ CV_EXPORTS_W int rotatedRectangleIntersection( const RotatedRect& rect1, const R
CV_EXPORTS Ptr<CLAHE> createCLAHE(double clipLimit = 40.0, Size tileGridSize = Size(8, 8));
//! Ballard, D.H. (1981). Generalizing the Hough transform to detect arbitrary shapes. Pattern Recognition 13 (2): 111-122.
//! Detects position only without traslation and rotation
CV_EXPORTS Ptr<GeneralizedHoughBallard> createGeneralizedHoughBallard();
//! Guil, N., González-Linares, J.M. and Zapata, E.L. (1999). Bidimensional shape detection using an invariant approach. Pattern Recognition 32 (6): 1025-1038.
//! Detects position, traslation and rotation
CV_EXPORTS Ptr<GeneralizedHoughGuil> createGeneralizedHoughGuil();
} // cv
#endif

1155
modules/imgproc/src/generalized_hough.cpp
View File

File diff suppressed because it is too large Load Diff

129
samples/gpu/generalized_hough.cpp
View File

@ -5,13 +5,12 @@
#include "opencv2/core.hpp"
#include "opencv2/core/utility.hpp"
#include "opencv2/imgproc.hpp"
#include "opencv2/gpu.hpp"
#include "opencv2/gpuimgproc.hpp"
#include "opencv2/highgui.hpp"
#include "opencv2/contrib.hpp"
using namespace std;
using namespace cv;
using cv::gpu::GpuMat;
static Mat loadImage(const string& name)
{
@ -29,8 +28,7 @@ int main(int argc, const char* argv[])
CommandLineParser cmd(argc, argv,
"{ image i | pic1.png | input image }"
"{ template t | templ.png | template image }"
"{ scale s | | estimate scale }"
"{ rotation r | | estimate rotation }"
"{ full | | estimate scale and rotation }"
"{ gpu | | use gpu version }"
"{ minDist | 100 | minimum distance between the centers of the detected objects }"
"{ levels | 360 | R-Table levels }"
@ -45,7 +43,7 @@ int main(int argc, const char* argv[])
"{ minAngle | 0 | minimal rotation angle to detect in degrees }"
"{ maxAngle | 360 | maximal rotation angle to detect in degrees }"
"{ angleStep | 1 | angle step in degrees }"
"{ maxSize | 1000 | maximal size of inner buffers }"
"{ maxBufSize | 1000 | maximal size of inner buffers }"
"{ help h ? | | print help message }"
);
@ -59,8 +57,7 @@ int main(int argc, const char* argv[])
const string templName = cmd.get<string>("template");
const string imageName = cmd.get<string>("image");
const bool estimateScale = cmd.has("scale");
const bool estimateRotation = cmd.has("rotation");
const bool full = cmd.has("full");
const bool useGpu = cmd.has("gpu");
const double minDist = cmd.get<double>("minDist");
const int levels = cmd.get<int>("levels");
@ -75,7 +72,7 @@ int main(int argc, const char* argv[])
const double minAngle = cmd.get<double>("minAngle");
const double maxAngle = cmd.get<double>("maxAngle");
const double angleStep = cmd.get<double>("angleStep");
const int maxSize = cmd.get<int>("maxSize");
const int maxBufSize = cmd.get<int>("maxBufSize");
if (!cmd.check())
{
@ -86,93 +83,69 @@ int main(int argc, const char* argv[])
Mat templ = loadImage(templName);
Mat image = loadImage(imageName);
int method = cv::GeneralizedHough::GHT_POSITION;
if (estimateScale)
method += cv::GeneralizedHough::GHT_SCALE;
if (estimateRotation)
method += cv::GeneralizedHough::GHT_ROTATION;
Ptr<GeneralizedHough> alg;
if (!full)
{
Ptr<GeneralizedHoughBallard> ballard = useGpu ? gpu::createGeneralizedHoughBallard() : createGeneralizedHoughBallard();
ballard->setMinDist(minDist);
ballard->setLevels(levels);
ballard->setDp(dp);
ballard->setMaxBufferSize(maxBufSize);
ballard->setVotesThreshold(votesThreshold);
alg = ballard;
}
else
{
Ptr<GeneralizedHoughGuil> guil = useGpu ? gpu::createGeneralizedHoughGuil() : createGeneralizedHoughGuil();
guil->setMinDist(minDist);
guil->setLevels(levels);
guil->setDp(dp);
guil->setMaxBufferSize(maxBufSize);
guil->setMinAngle(minAngle);
guil->setMaxAngle(maxAngle);
guil->setAngleStep(angleStep);
guil->setAngleThresh(angleThresh);
guil->setMinScale(minScale);
guil->setMaxScale(maxScale);
guil->setScaleStep(scaleStep);
guil->setScaleThresh(scaleThresh);
guil->setPosThresh(posThresh);
alg = guil;
}
vector<Vec4f> position;
cv::TickMeter tm;
TickMeter tm;
if (useGpu)
{
GpuMat d_templ(templ);
GpuMat d_image(image);
GpuMat d_position;
gpu::GpuMat d_templ(templ);
gpu::GpuMat d_image(image);
gpu::GpuMat d_position;
Ptr<gpu::GeneralizedHough> d_hough = gpu::GeneralizedHough::create(method);
d_hough->set("minDist", minDist);
d_hough->set("levels", levels);
d_hough->set("dp", dp);
d_hough->set("maxSize", maxSize);
if (estimateScale && estimateRotation)
{
d_hough->set("angleThresh", angleThresh);
d_hough->set("scaleThresh", scaleThresh);
d_hough->set("posThresh", posThresh);
}
else
{
d_hough->set("votesThreshold", votesThreshold);
}
if (estimateScale)
{
d_hough->set("minScale", minScale);
d_hough->set("maxScale", maxScale);
d_hough->set("scaleStep", scaleStep);
}
if (estimateRotation)
{
d_hough->set("minAngle", minAngle);
d_hough->set("maxAngle", maxAngle);
d_hough->set("angleStep", angleStep);
}
d_hough->setTemplate(d_templ);
alg->setTemplate(d_templ);
tm.start();
d_hough->detect(d_image, d_position);
d_hough->downloadResults(d_position, position);
alg->detect(d_image, d_position);
d_position.download(position);
tm.stop();
}
else
{
Ptr<GeneralizedHough> hough = GeneralizedHough::create(method);
hough->set("minDist", minDist);
hough->set("levels", levels);
hough->set("dp", dp);
if (estimateScale && estimateRotation)
{
hough->set("angleThresh", angleThresh);
hough->set("scaleThresh", scaleThresh);
hough->set("posThresh", posThresh);
hough->set("maxSize", maxSize);
}
else
{
hough->set("votesThreshold", votesThreshold);
}
if (estimateScale)
{
hough->set("minScale", minScale);
hough->set("maxScale", maxScale);
hough->set("scaleStep", scaleStep);
}
if (estimateRotation)
{
hough->set("minAngle", minAngle);
hough->set("maxAngle", maxAngle);
hough->set("angleStep", angleStep);
}
hough->setTemplate(templ);
alg->setTemplate(templ);
tm.start();
hough->detect(image, position);
alg->detect(image, position);
tm.stop();
}