Commit 6960e154 authored by Kirill Kornyakov's avatar Kirill Kornyakov

GPU module update: _GPU suffix removed, some namespaces renamed, minor refactorings.

parent a38e5111
...@@ -344,10 +344,10 @@ namespace cv ...@@ -344,10 +344,10 @@ namespace cv
////////////////////////////// Image processing ////////////////////////////// ////////////////////////////// Image processing //////////////////////////////
// DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation. // DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation.
// xymap.type() == xymap.type() == CV_32FC1 // xymap.type() == xymap.type() == CV_32FC1
CV_EXPORTS void remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, GpuMat& dst); CV_EXPORTS void remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap);
// Does mean shift filtering on GPU. // Does mean shift filtering on GPU.
CV_EXPORTS void meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, CV_EXPORTS void meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr,
TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1)); TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1));
// Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV. // Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV.
...@@ -362,18 +362,18 @@ namespace cv ...@@ -362,18 +362,18 @@ namespace cv
// The output is a 4-channel floating-point (CV_32FC4) matrix. // The output is a 4-channel floating-point (CV_32FC4) matrix.
// Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map. // Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map.
// Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify. // Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify.
CV_EXPORTS void reprojectImageTo3D_GPU(const GpuMat& disp, GpuMat& xyzw, const Mat& Q); CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q);
// Acync version // Acync version
CV_EXPORTS void reprojectImageTo3D_GPU(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream); CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream);
CV_EXPORTS void cvtColor_GPU(const GpuMat& src, GpuMat& dst, int code, int dcn = 0); CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn = 0);
CV_EXPORTS void cvtColor_GPU(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream); CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream);
//////////////////////////////// StereoBM_GPU //////////////////////////////// //////////////////////////////// StereoBM_GPU ////////////////////////////////
class CV_EXPORTS StereoBM_GPU class CV_EXPORTS StereoBM_GPU
{ {
public: public:
enum { BASIC_PRESET = 0, PREFILTER_XSOBEL = 1 }; enum { BASIC_PRESET = 0, PREFILTER_XSOBEL = 1 };
enum { DEFAULT_NDISP = 64, DEFAULT_WINSZ = 19 }; enum { DEFAULT_NDISP = 64, DEFAULT_WINSZ = 19 };
......
...@@ -45,7 +45,6 @@ ...@@ -45,7 +45,6 @@
#include "safe_call.hpp" #include "safe_call.hpp"
using namespace cv::gpu; using namespace cv::gpu;
using namespace cv::gpu::impl;
#ifndef FLT_MAX #ifndef FLT_MAX
#define FLT_MAX 3.402823466e+38F #define FLT_MAX 3.402823466e+38F
......
...@@ -45,7 +45,6 @@ ...@@ -45,7 +45,6 @@
#include "safe_call.hpp" #include "safe_call.hpp"
using namespace cv::gpu; using namespace cv::gpu;
using namespace cv::gpu::impl;
#ifndef FLT_MAX #ifndef FLT_MAX
#define FLT_MAX 3.402823466e+30F #define FLT_MAX 3.402823466e+30F
......
...@@ -44,7 +44,6 @@ ...@@ -44,7 +44,6 @@
#include "saturate_cast.hpp" #include "saturate_cast.hpp"
using namespace cv::gpu; using namespace cv::gpu;
using namespace cv::gpu::impl;
#ifndef CV_DESCALE #ifndef CV_DESCALE
#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n)) #define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))
...@@ -167,7 +166,7 @@ namespace imgproc ...@@ -167,7 +166,7 @@ namespace imgproc
} }
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace improc
{ {
template <typename T> template <typename T>
void RGB2RGB_caller(const DevMem2D_<T>& src, int srccn, const DevMem2D_<T>& dst, int dstcn, int bidx, cudaStream_t stream) void RGB2RGB_caller(const DevMem2D_<T>& src, int srccn, const DevMem2D_<T>& dst, int dstcn, int bidx, cudaStream_t stream)
...@@ -377,7 +376,7 @@ namespace imgproc ...@@ -377,7 +376,7 @@ namespace imgproc
//}; //};
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace improc
{ {
template <typename T> template <typename T>
void Gray2RGB_caller(const DevMem2D_<T>& src, const DevMem2D_<T>& dst, int dstcn, cudaStream_t stream) void Gray2RGB_caller(const DevMem2D_<T>& src, const DevMem2D_<T>& dst, int dstcn, cudaStream_t stream)
...@@ -627,7 +626,7 @@ namespace imgproc ...@@ -627,7 +626,7 @@ namespace imgproc
} }
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace improc
{ {
void RGB2Gray_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream) void RGB2Gray_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream)
{ {
......
...@@ -45,7 +45,6 @@ ...@@ -45,7 +45,6 @@
#include "safe_call.hpp" #include "safe_call.hpp"
using namespace cv::gpu; using namespace cv::gpu;
using namespace cv::gpu::impl;
#ifndef FLT_MAX #ifndef FLT_MAX
#define FLT_MAX 3.402823466e+30F #define FLT_MAX 3.402823466e+30F
......
...@@ -56,10 +56,10 @@ namespace cv ...@@ -56,10 +56,10 @@ namespace cv
typedef unsigned short ushort; typedef unsigned short ushort;
typedef unsigned int uint; typedef unsigned int uint;
namespace impl static inline int divUp(int a, int b) { return (a % b == 0) ? a/b : a/b + 1; }
{
static inline int divUp(int a, int b) { return (a % b == 0) ? a/b : a/b + 1; }
namespace matrix_operations
{
extern "C" void copy_to_with_mask(const DevMem2D& src, DevMem2D dst, int depth, const DevMem2D& mask, int channels, const cudaStream_t & stream = 0); extern "C" void copy_to_with_mask(const DevMem2D& src, DevMem2D dst, int depth, const DevMem2D& mask, int channels, const cudaStream_t & stream = 0);
extern "C" void set_to_without_mask (DevMem2D dst, int depth, const double *scalar, int channels, const cudaStream_t & stream = 0); extern "C" void set_to_without_mask (DevMem2D dst, int depth, const double *scalar, int channels, const cudaStream_t & stream = 0);
......
...@@ -123,7 +123,7 @@ namespace imgproc ...@@ -123,7 +123,7 @@ namespace imgproc
} }
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace improc
{ {
void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst) void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst)
{ {
...@@ -231,7 +231,7 @@ namespace imgproc ...@@ -231,7 +231,7 @@ namespace imgproc
} }
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace improc
{ {
extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps) extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps)
{ {
...@@ -354,7 +354,7 @@ namespace imgproc ...@@ -354,7 +354,7 @@ namespace imgproc
} }
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace improc
{ {
void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream) void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream)
{ {
...@@ -420,7 +420,7 @@ namespace imgproc ...@@ -420,7 +420,7 @@ namespace imgproc
} }
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace improc
{ {
template <typename T> template <typename T>
inline void reprojectImageTo3D_caller(const DevMem2D_<T>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream) inline void reprojectImageTo3D_caller(const DevMem2D_<T>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream)
......
...@@ -47,7 +47,7 @@ ...@@ -47,7 +47,7 @@
#include "saturate_cast.hpp" #include "saturate_cast.hpp"
using namespace cv::gpu; using namespace cv::gpu;
using namespace cv::gpu::impl; using namespace cv::gpu::matrix_operations;
namespace mat_operators namespace mat_operators
...@@ -261,7 +261,7 @@ namespace cv ...@@ -261,7 +261,7 @@ namespace cv
{ {
namespace gpu namespace gpu
{ {
namespace impl namespace matrix_operations
{ {
/////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////////////////
......
...@@ -316,7 +316,7 @@ __global__ void stereoKernel(unsigned char *left, unsigned char *right, size_t i ...@@ -316,7 +316,7 @@ __global__ void stereoKernel(unsigned char *left, unsigned char *right, size_t i
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace bm
{ {
template<int RADIUS> void kernel_caller(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int maxdisp, const cudaStream_t & stream) template<int RADIUS> void kernel_caller(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int maxdisp, const cudaStream_t & stream)
{ {
...@@ -408,7 +408,7 @@ extern "C" __global__ void prefilter_kernel(unsigned char *output, size_t step, ...@@ -408,7 +408,7 @@ extern "C" __global__ void prefilter_kernel(unsigned char *output, size_t step,
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace bm
{ {
extern "C" void prefilter_xsobel(const DevMem2D& input, const DevMem2D& output, int prefilterCap) extern "C" void prefilter_xsobel(const DevMem2D& input, const DevMem2D& output, int prefilterCap)
{ {
...@@ -530,7 +530,7 @@ extern "C" __global__ void textureness_kernel(unsigned char *disp, size_t disp_s ...@@ -530,7 +530,7 @@ extern "C" __global__ void textureness_kernel(unsigned char *disp, size_t disp_s
} }
} }
namespace cv { namespace gpu { namespace impl namespace cv { namespace gpu { namespace bm
{ {
extern "C" void postfilter_textureness(const DevMem2D& input, int winsz, float avgTexturenessThreshold, const DevMem2D& disp) extern "C" void postfilter_textureness(const DevMem2D& input, int winsz, float avgTexturenessThreshold, const DevMem2D& disp)
{ {
......
...@@ -158,12 +158,12 @@ void cv::gpu::Stream::enqueueCopy(const GpuMat& src, GpuMat& dst) { devcopy(src, ...@@ -158,12 +158,12 @@ void cv::gpu::Stream::enqueueCopy(const GpuMat& src, GpuMat& dst) { devcopy(src,
void cv::gpu::Stream::enqueueMemSet(const GpuMat& src, Scalar val) void cv::gpu::Stream::enqueueMemSet(const GpuMat& src, Scalar val)
{ {
impl::set_to_without_mask(src, src.depth(), val.val, src.channels(), impl->stream); matrix_operations::set_to_without_mask(src, src.depth(), val.val, src.channels(), impl->stream);
} }
void cv::gpu::Stream::enqueueMemSet(const GpuMat& src, Scalar val, const GpuMat& mask) void cv::gpu::Stream::enqueueMemSet(const GpuMat& src, Scalar val, const GpuMat& mask)
{ {
impl::set_to_with_mask(src, src.depth(), val.val, mask, src.channels(), impl->stream); matrix_operations::set_to_with_mask(src, src.depth(), val.val, mask, src.channels(), impl->stream);
} }
void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int rtype, double alpha, double beta) void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int rtype, double alpha, double beta)
...@@ -188,7 +188,7 @@ void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int rtype, ...@@ -188,7 +188,7 @@ void cv::gpu::Stream::enqueueConvert(const GpuMat& src, GpuMat& dst, int rtype,
psrc = &(temp = src); psrc = &(temp = src);
dst.create( src.size(), rtype ); dst.create( src.size(), rtype );
impl::convert_to(*psrc, sdepth, dst, ddepth, psrc->channels(), alpha, beta, impl->stream); matrix_operations::convert_to(*psrc, sdepth, dst, ddepth, psrc->channels(), alpha, beta, impl->stream);
} }
......
...@@ -47,20 +47,20 @@ using namespace cv::gpu; ...@@ -47,20 +47,20 @@ using namespace cv::gpu;
#if !defined (HAVE_CUDA) #if !defined (HAVE_CUDA)
void cv::gpu::remap(const GpuMat&, const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::remap( const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap ){ throw_nogpu(); }
void cv::gpu::meanShiftFiltering_GPU(const GpuMat&, GpuMat&, int, int, TermCriteria ) { throw_nogpu(); } void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria ) { throw_nogpu(); }
void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int) { throw_nogpu(); } void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int) { throw_nogpu(); }
void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, const Stream&) { throw_nogpu(); } void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, const Stream&) { throw_nogpu(); }
void cv::gpu::reprojectImageTo3D_GPU(const GpuMat&, GpuMat&, const Mat&) { throw_nogpu(); } void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&) { throw_nogpu(); }
void cv::gpu::reprojectImageTo3D_GPU(const GpuMat&, GpuMat&, const Mat&, const Stream&) { throw_nogpu(); } void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, const Stream&) { throw_nogpu(); }
void cv::gpu::cvtColor_GPU(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); } void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); }
void cv::gpu::cvtColor_GPU(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); } void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); }
#else /* !defined (HAVE_CUDA) */ #else /* !defined (HAVE_CUDA) */
namespace cv { namespace gpu namespace cv { namespace gpu
{ {
namespace impl namespace improc
{ {
void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst); void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst); void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
...@@ -90,10 +90,10 @@ namespace cv { namespace gpu ...@@ -90,10 +90,10 @@ namespace cv { namespace gpu
//////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////
// remap // remap
void cv::gpu::remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, GpuMat& dst) void cv::gpu::remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap)
{ {
typedef void (*remap_gpu_t)(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst); typedef void (*remap_gpu_t)(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
static const remap_gpu_t callers[] = {impl::remap_gpu_1c, 0, impl::remap_gpu_3c}; static const remap_gpu_t callers[] = {improc::remap_gpu_1c, 0, improc::remap_gpu_3c};
CV_Assert((src.type() == CV_8U || src.type() == CV_8UC3) && xmap.type() == CV_32F && ymap.type() == CV_32F); CV_Assert((src.type() == CV_8U || src.type() == CV_8UC3) && xmap.type() == CV_32F && ymap.type() == CV_32F);
...@@ -111,7 +111,7 @@ void cv::gpu::remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, G ...@@ -111,7 +111,7 @@ void cv::gpu::remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, G
//////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////
// meanShiftFiltering_GPU // meanShiftFiltering_GPU
void cv::gpu::meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria) void cv::gpu::meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria)
{ {
if( src.empty() ) if( src.empty() )
CV_Error( CV_StsBadArg, "The input image is empty" ); CV_Error( CV_StsBadArg, "The input image is empty" );
...@@ -131,7 +131,7 @@ void cv::gpu::meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int ...@@ -131,7 +131,7 @@ void cv::gpu::meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int
eps = 1.f; eps = 1.f;
eps = (float)std::max(criteria.epsilon, 0.0); eps = (float)std::max(criteria.epsilon, 0.0);
impl::meanShiftFiltering_gpu(src, dst, sp, sr, maxIter, eps); improc::meanShiftFiltering_gpu(src, dst, sp, sr, maxIter, eps);
} }
//////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////
...@@ -147,7 +147,7 @@ namespace ...@@ -147,7 +147,7 @@ namespace
out = dst; out = dst;
out.create(src.size(), CV_8UC4); out.create(src.size(), CV_8UC4);
impl::drawColorDisp_gpu((DevMem2D_<T>)src, out, ndisp, stream); improc::drawColorDisp_gpu((DevMem2D_<T>)src, out, ndisp, stream);
dst = out; dst = out;
} }
...@@ -180,7 +180,7 @@ namespace ...@@ -180,7 +180,7 @@ namespace
void reprojectImageTo3D_caller(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream) void reprojectImageTo3D_caller(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream)
{ {
xyzw.create(disp.rows, disp.cols, CV_32FC4); xyzw.create(disp.rows, disp.cols, CV_32FC4);
impl::reprojectImageTo3D_gpu((DevMem2D_<T>)disp, xyzw, Q.ptr<float>(), stream); improc::reprojectImageTo3D_gpu((DevMem2D_<T>)disp, xyzw, Q.ptr<float>(), stream);
} }
typedef void (*reprojectImageTo3D_caller_t)(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream); typedef void (*reprojectImageTo3D_caller_t)(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream);
...@@ -188,14 +188,14 @@ namespace ...@@ -188,14 +188,14 @@ namespace
const reprojectImageTo3D_caller_t reprojectImageTo3D_callers[] = {reprojectImageTo3D_caller<unsigned char>, 0, 0, reprojectImageTo3D_caller<short>, 0, 0, 0, 0}; const reprojectImageTo3D_caller_t reprojectImageTo3D_callers[] = {reprojectImageTo3D_caller<unsigned char>, 0, 0, reprojectImageTo3D_caller<short>, 0, 0, 0, 0};
} }
void cv::gpu::reprojectImageTo3D_GPU(const GpuMat& disp, GpuMat& xyzw, const Mat& Q) void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q)
{ {
CV_Assert((disp.type() == CV_8U || disp.type() == CV_16S) && Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4); CV_Assert((disp.type() == CV_8U || disp.type() == CV_16S) && Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4);
reprojectImageTo3D_callers[disp.type()](disp, xyzw, Q, 0); reprojectImageTo3D_callers[disp.type()](disp, xyzw, Q, 0);
} }
void cv::gpu::reprojectImageTo3D_GPU(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream) void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream)
{ {
CV_Assert((disp.type() == CV_8U || disp.type() == CV_16S) && Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4); CV_Assert((disp.type() == CV_8U || disp.type() == CV_16S) && Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4);
...@@ -229,11 +229,11 @@ namespace ...@@ -229,11 +229,11 @@ namespace
out.create(sz, CV_MAKETYPE(depth, dcn)); out.create(sz, CV_MAKETYPE(depth, dcn));
if( depth == CV_8U ) if( depth == CV_8U )
impl::RGB2RGB_gpu((DevMem2D)src, scn, (DevMem2D)out, dcn, bidx, stream); improc::RGB2RGB_gpu((DevMem2D)src, scn, (DevMem2D)out, dcn, bidx, stream);
else if( depth == CV_16U ) else if( depth == CV_16U )
impl::RGB2RGB_gpu((DevMem2D_<unsigned short>)src, scn, (DevMem2D_<unsigned short>)out, dcn, bidx, stream); improc::RGB2RGB_gpu((DevMem2D_<unsigned short>)src, scn, (DevMem2D_<unsigned short>)out, dcn, bidx, stream);
else else
impl::RGB2RGB_gpu((DevMem2Df)src, scn, (DevMem2Df)out, dcn, bidx, stream); improc::RGB2RGB_gpu((DevMem2Df)src, scn, (DevMem2Df)out, dcn, bidx, stream);
break; break;
//case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555: //case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:
...@@ -270,11 +270,11 @@ namespace ...@@ -270,11 +270,11 @@ namespace
bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2; bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;
if( depth == CV_8U ) if( depth == CV_8U )
impl::RGB2Gray_gpu((DevMem2D)src, scn, (DevMem2D)out, bidx, stream); improc::RGB2Gray_gpu((DevMem2D)src, scn, (DevMem2D)out, bidx, stream);
else if( depth == CV_16U ) else if( depth == CV_16U )
impl::RGB2Gray_gpu((DevMem2D_<unsigned short>)src, scn, (DevMem2D_<unsigned short>)out, bidx, stream); improc::RGB2Gray_gpu((DevMem2D_<unsigned short>)src, scn, (DevMem2D_<unsigned short>)out, bidx, stream);
else else
impl::RGB2Gray_gpu((DevMem2Df)src, scn, (DevMem2Df)out, bidx, stream); improc::RGB2Gray_gpu((DevMem2Df)src, scn, (DevMem2Df)out, bidx, stream);
break; break;
//case CV_BGR5652GRAY: case CV_BGR5552GRAY: //case CV_BGR5652GRAY: case CV_BGR5552GRAY:
...@@ -291,11 +291,11 @@ namespace ...@@ -291,11 +291,11 @@ namespace
out.create(sz, CV_MAKETYPE(depth, dcn)); out.create(sz, CV_MAKETYPE(depth, dcn));
if( depth == CV_8U ) if( depth == CV_8U )
impl::Gray2RGB_gpu((DevMem2D)src, (DevMem2D)out, dcn, stream); improc::Gray2RGB_gpu((DevMem2D)src, (DevMem2D)out, dcn, stream);
else if( depth == CV_16U ) else if( depth == CV_16U )
impl::Gray2RGB_gpu((DevMem2D_<unsigned short>)src, (DevMem2D_<unsigned short>)out, dcn, stream); improc::Gray2RGB_gpu((DevMem2D_<unsigned short>)src, (DevMem2D_<unsigned short>)out, dcn, stream);
else else
impl::Gray2RGB_gpu((DevMem2Df)src, (DevMem2Df)out, dcn, stream); improc::Gray2RGB_gpu((DevMem2Df)src, (DevMem2Df)out, dcn, stream);
break; break;
//case CV_GRAY2BGR565: case CV_GRAY2BGR555: //case CV_GRAY2BGR565: case CV_GRAY2BGR555:
...@@ -516,12 +516,12 @@ namespace ...@@ -516,12 +516,12 @@ namespace
} }
} }
void cv::gpu::cvtColor_GPU(const GpuMat& src, GpuMat& dst, int code, int dcn) void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn)
{ {
cvtColor_caller(src, dst, code, dcn, 0); cvtColor_caller(src, dst, code, dcn, 0);
} }
void cv::gpu::cvtColor_GPU(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream) void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream)
{ {
cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream)); cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream));
} }
......
...@@ -120,7 +120,7 @@ void cv::gpu::GpuMat::copyTo( GpuMat& mat, const GpuMat& mask ) const ...@@ -120,7 +120,7 @@ void cv::gpu::GpuMat::copyTo( GpuMat& mat, const GpuMat& mask ) const
else else
{ {
mat.create(size(), type()); mat.create(size(), type());
cv::gpu::impl::copy_to_with_mask(*this, mat, depth(), mask, channels()); cv::gpu::matrix_operations::copy_to_with_mask(*this, mat, depth(), mask, channels());
} }
} }
...@@ -146,12 +146,12 @@ void cv::gpu::GpuMat::convertTo( GpuMat& dst, int rtype, double alpha, double be ...@@ -146,12 +146,12 @@ void cv::gpu::GpuMat::convertTo( GpuMat& dst, int rtype, double alpha, double be
psrc = &(temp = *this); psrc = &(temp = *this);
dst.create( size(), rtype ); dst.create( size(), rtype );
impl::convert_to(*psrc, sdepth, dst, ddepth, psrc->channels(), alpha, beta); matrix_operations::convert_to(*psrc, sdepth, dst, ddepth, psrc->channels(), alpha, beta);
} }
GpuMat& GpuMat::operator = (const Scalar& s) GpuMat& GpuMat::operator = (const Scalar& s)
{ {
impl::set_to_without_mask( *this, depth(), s.val, channels()); matrix_operations::set_to_without_mask( *this, depth(), s.val, channels());
return *this; return *this;
} }
...@@ -162,9 +162,9 @@ GpuMat& GpuMat::setTo(const Scalar& s, const GpuMat& mask) ...@@ -162,9 +162,9 @@ GpuMat& GpuMat::setTo(const Scalar& s, const GpuMat& mask)
CV_DbgAssert(!this->empty()); CV_DbgAssert(!this->empty());
if (mask.empty()) if (mask.empty())
impl::set_to_without_mask( *this, depth(), s.val, channels()); matrix_operations::set_to_without_mask( *this, depth(), s.val, channels());
else else
impl::set_to_with_mask( *this, depth(), s.val, mask, channels()); matrix_operations::set_to_with_mask( *this, depth(), s.val, mask, channels());
return *this; return *this;
} }
......
...@@ -58,7 +58,7 @@ void cv::gpu::StereoBM_GPU::operator() ( const GpuMat&, const GpuMat&, GpuMat&, ...@@ -58,7 +58,7 @@ void cv::gpu::StereoBM_GPU::operator() ( const GpuMat&, const GpuMat&, GpuMat&,
namespace cv { namespace gpu namespace cv { namespace gpu
{ {
namespace impl namespace bm
{ {
//extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int ndisp, int winsz, const DevMem2D_<uint>& minSSD_buf); //extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int ndisp, int winsz, const DevMem2D_<uint>& minSSD_buf);
extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int ndisp, int winsz, const DevMem2D_<uint>& minSSD_buf, const cudaStream_t & stream); extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int ndisp, int winsz, const DevMem2D_<uint>& minSSD_buf, const cudaStream_t & stream);
...@@ -115,17 +115,17 @@ static void stereo_bm_gpu_operator ( GpuMat& minSSD, GpuMat& leBuf, GpuMat& ri ...@@ -115,17 +115,17 @@ static void stereo_bm_gpu_operator ( GpuMat& minSSD, GpuMat& leBuf, GpuMat& ri
leBuf.create( left.size(), left.type()); leBuf.create( left.size(), left.type());
riBuf.create(right.size(), right.type()); riBuf.create(right.size(), right.type());
impl::prefilter_xsobel( left, leBuf); bm::prefilter_xsobel( left, leBuf);
impl::prefilter_xsobel(right, riBuf); bm::prefilter_xsobel(right, riBuf);
le_for_bm = leBuf; le_for_bm = leBuf;
ri_for_bm = riBuf; ri_for_bm = riBuf;
} }
impl::stereoBM_GPU(le_for_bm, ri_for_bm, disparity, ndisp, winSize, minSSD, stream); bm::stereoBM_GPU(le_for_bm, ri_for_bm, disparity, ndisp, winSize, minSSD, stream);
if (avergeTexThreshold) if (avergeTexThreshold)
impl::postfilter_textureness(le_for_bm, winSize, avergeTexThreshold, disparity); bm::postfilter_textureness(le_for_bm, winSize, avergeTexThreshold, disparity);
} }
......
...@@ -76,7 +76,7 @@ void CV_GpuMeanShiftTest::run(int) ...@@ -76,7 +76,7 @@ void CV_GpuMeanShiftTest::run(int)
cvtColor(img, rgba, CV_BGR2BGRA); cvtColor(img, rgba, CV_BGR2BGRA);
cv::gpu::GpuMat res; cv::gpu::GpuMat res;
cv::gpu::meanShiftFiltering_GPU( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad ); cv::gpu::meanShiftFiltering( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad );
if (res.type() != CV_8UC4) if (res.type() != CV_8UC4)
{ {
ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT); ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);
......
...@@ -51,6 +51,7 @@ class CV_GpuStereoBMTest : public CvTest ...@@ -51,6 +51,7 @@ class CV_GpuStereoBMTest : public CvTest
{ {
public: public:
CV_GpuStereoBMTest(); CV_GpuStereoBMTest();
protected: protected:
void run(int); void run(int);
}; };
...@@ -61,9 +62,9 @@ void CV_GpuStereoBMTest::run(int ) ...@@ -61,9 +62,9 @@ void CV_GpuStereoBMTest::run(int )
{ {
cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-L.png", 0); cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-L.png", 0);
cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-R.png", 0); cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-R.png", 0);
cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", 0); cv::Mat img_reference = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", 0);
if (img_l.empty() || img_r.empty() || img_template.empty()) if (img_l.empty() || img_r.empty() || img_reference.empty())
{ {
ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA); ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
return; return;
...@@ -71,14 +72,11 @@ void CV_GpuStereoBMTest::run(int ) ...@@ -71,14 +72,11 @@ void CV_GpuStereoBMTest::run(int )
cv::gpu::GpuMat disp; cv::gpu::GpuMat disp;
cv::gpu::StereoBM_GPU bm(0, 128, 19); cv::gpu::StereoBM_GPU bm(0, 128, 19);
bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp); bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
//cv::imwrite(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", disp); disp.convertTo(disp, img_reference.type());
double norm = cv::norm(disp, img_reference, cv::NORM_INF);
disp.convertTo(disp, img_template.type());
double norm = cv::norm(disp, img_template, cv::NORM_INF);
if (norm >= 100) if (norm >= 100)
ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm); ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm);
ts->set_failed_test_info((norm < 100) ? CvTS::OK : CvTS::FAIL_GENERIC); ts->set_failed_test_info((norm < 100) ? CvTS::OK : CvTS::FAIL_GENERIC);
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment