#else
// CUBLAS works with column-major matrices
- GpuMat src1 = _src1.getGpuMat();
- GpuMat src2 = _src2.getGpuMat();
- GpuMat src3 = _src3.getGpuMat();
+ GpuMat src1 = getInputMat(_src1, stream);
+ GpuMat src2 = getInputMat(_src2, stream);
+ GpuMat src3 = getInputMat(_src3, stream);
CV_Assert( src1.type() == CV_32FC1 || src1.type() == CV_32FC2 || src1.type() == CV_64FC1 || src1.type() == CV_64FC2 );
CV_Assert( src2.type() == src1.type() && (src3.empty() || src3.type() == src1.type()) );
CV_Assert( src1Size.width == src2Size.height );
CV_Assert( src3.empty() || src3Size == dstSize );
- _dst.create(dstSize, src1.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, dstSize, src1.type(), stream);
if (beta != 0)
{
}
cublasSafeCall( cublasDestroy_v2(handle) );
+
+ syncOutput(dst, _dst, stream);
#endif
}
(void) stream;
throw_no_cuda();
#else
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
CV_Assert( src.type() == CV_32FC1 || src.type() == CV_32FC2 );
void ConvolutionImpl::convolve(InputArray _image, InputArray _templ, OutputArray _result, bool ccorr, Stream& _stream)
{
- GpuMat image = _image.getGpuMat();
- GpuMat templ = _templ.getGpuMat();
+ GpuMat image = getInputMat(_image, _stream);
+ GpuMat templ = getInputMat(_templ, _stream);
CV_Assert( image.type() == CV_32FC1 );
CV_Assert( templ.type() == CV_32FC1 );
create(image.size(), templ.size());
- _result.create(result_size, CV_32FC1);
- GpuMat result = _result.getGpuMat();
+ GpuMat result = getOutputMat(_result, result_size, CV_32FC1, _stream);
cudaStream_t stream = StreamAccessor::getStream(_stream);
cufftSafeCall( cufftDestroy(planR2C) );
cufftSafeCall( cufftDestroy(planC2R) );
+
+ syncOutput(result, _result, _stream);
}
}
{NppMirror<CV_32F, nppiMirror_32f_C1R>::call, 0, NppMirror<CV_32F, nppiMirror_32f_C3R>::call, NppMirror<CV_32F, nppiMirror_32f_C4R>::call}
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32S || src.depth() == CV_32F);
CV_Assert(src.channels() == 1 || src.channels() == 3 || src.channels() == 4);
_dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
funcs[src.depth()][src.channels() - 1](src, dst, flipCode, StreamAccessor::getStream(stream));
+
+ syncOutput(dst, _dst, stream);
}
#endif /* !defined (HAVE_CUDA) */
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudev.hpp"
+#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
namespace
__device__ __forceinline__ D operator ()(T1 a, T2 b) const
{
- return saturate_cast<D>(a * alpha + b * beta + gamma);
+ return cudev::saturate_cast<D>(a * alpha + b * beta + gamma);
}
};
}
};
- GpuMat src1 = _src1.getGpuMat();
- GpuMat src2 = _src2.getGpuMat();
+ GpuMat src1 = getInputMat(_src1, stream);
+ GpuMat src2 = getInputMat(_src2, stream);
int sdepth1 = src1.depth();
int sdepth2 = src2.depth();
ddepth = ddepth >= 0 ? CV_MAT_DEPTH(ddepth) : std::max(sdepth1, sdepth2);
const int cn = src1.channels();
- CV_DbgAssert( src2.size() == src1.size() && src2.channels() == cn );
- CV_DbgAssert( sdepth1 <= CV_64F && sdepth2 <= CV_64F && ddepth <= CV_64F );
+ CV_Assert( src2.size() == src1.size() && src2.channels() == cn );
+ CV_Assert( sdepth1 <= CV_64F && sdepth2 <= CV_64F && ddepth <= CV_64F );
- _dst.create(src1.size(), CV_MAKE_TYPE(ddepth, cn));
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src1.size(), CV_MAKE_TYPE(ddepth, cn), stream);
- GpuMat src1_ = src1.reshape(1);
- GpuMat src2_ = src2.reshape(1);
- GpuMat dst_ = dst.reshape(1);
+ GpuMat src1_single = src1.reshape(1);
+ GpuMat src2_single = src2.reshape(1);
+ GpuMat dst_single = dst.reshape(1);
if (sdepth1 > sdepth2)
{
- src1_.swap(src2_);
+ src1_single.swap(src2_single);
std::swap(alpha, beta);
std::swap(sdepth1, sdepth2);
}
if (!func)
CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types");
- func(src1_, alpha, src2_, beta, gamma, dst_, stream);
+ func(src1_single, alpha, src2_single, beta, gamma, dst_single, stream);
+
+ syncOutput(dst, _dst, stream);
}
#endif
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudev.hpp"
+#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
void bitMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int op);
void cv::cuda::bitwise_not(InputArray _src, OutputArray _dst, InputArray _mask, Stream& stream)
{
- GpuMat src = _src.getGpuMat();
- GpuMat mask = _mask.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
+ GpuMat mask = getInputMat(_mask, stream);
const int depth = src.depth();
CV_DbgAssert( depth <= CV_32F );
CV_DbgAssert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src.size()) );
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
if (mask.empty())
{
gridTransformUnary(vsrc, vdst, bit_not<uchar>(), singleMaskChannels(globPtr<uchar>(mask), src.channels()), stream);
}
}
+
+ syncOutput(dst, _dst, stream);
}
//////////////////////////////////////////////////////////////////////////////
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudev.hpp"
+#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
namespace
{ copyMakeBorderImpl<float , 1> , 0 /*copyMakeBorderImpl<float , 2>*/, copyMakeBorderImpl<float , 3> , copyMakeBorderImpl<float ,4> }
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
const int depth = src.depth();
const int cn = src.channels();
CV_Assert( depth <= CV_32F && cn <= 4 );
CV_Assert( borderType == BORDER_REFLECT_101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP );
- _dst.create(src.rows + top + bottom, src.cols + left + right, src.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.rows + top + bottom, src.cols + left + right, src.type(), stream);
const func_t func = funcs[depth][cn - 1];
CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types");
func(src, dst, top, left, borderType, value, stream);
+
+ syncOutput(dst, _dst, stream);
}
#endif
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudev.hpp"
+#include "opencv2/core/private.cuda.hpp"
using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
namespace
void LookUpTableImpl::transform(InputArray _src, OutputArray _dst, Stream& stream)
{
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
const int cn = src.channels();
const int lut_cn = d_lut.channels();
CV_Assert( src.type() == CV_8UC1 || src.type() == CV_8UC3 );
CV_Assert( lut_cn == 1 || lut_cn == cn );
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
if (lut_cn == 1)
{
dst3.assign(lut_(src3, tbl), stream);
}
+
+ syncOutput(dst, _dst, stream);
}
}
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudev.hpp"
+#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
namespace
absMat<double>
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
- const int depth = src.depth();
+ CV_Assert( src.depth() <= CV_64F );
- CV_DbgAssert( depth <= CV_64F );
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
- funcs[depth](src.reshape(1), dst.reshape(1), stream);
+ syncOutput(dst, _dst, stream);
}
//////////////////////////////////////////////////////////////////////////////
{
__device__ __forceinline__ T operator ()(T x) const
{
- return saturate_cast<T>(x * x);
+ return cudev::saturate_cast<T>(x * x);
}
};
sqrMat<double>
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
- const int depth = src.depth();
+ CV_Assert( src.depth() <= CV_64F );
- CV_DbgAssert( depth <= CV_64F );
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
- funcs[depth](src.reshape(1), dst.reshape(1), stream);
+ syncOutput(dst, _dst, stream);
}
//////////////////////////////////////////////////////////////////////////////
sqrtMat<double>
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
- const int depth = src.depth();
+ CV_Assert( src.depth() <= CV_64F );
- CV_DbgAssert( depth <= CV_64F );
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
- funcs[depth](src.reshape(1), dst.reshape(1), stream);
+ syncOutput(dst, _dst, stream);
}
////////////////////////////////////////////////////////////////////////
__device__ __forceinline__ T operator ()(T x) const
{
exp_func<T> f;
- return saturate_cast<T>(f(x));
+ return cudev::saturate_cast<T>(f(x));
}
};
expMat<double>
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
- const int depth = src.depth();
+ CV_Assert( src.depth() <= CV_64F );
- CV_DbgAssert( depth <= CV_64F );
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
- funcs[depth](src.reshape(1), dst.reshape(1), stream);
+ syncOutput(dst, _dst, stream);
}
////////////////////////////////////////////////////////////////////////
logMat<double>
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
- const int depth = src.depth();
+ CV_Assert( src.depth() <= CV_64F );
- CV_DbgAssert( depth <= CV_64F );
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ funcs[src.depth()](src.reshape(1), dst.reshape(1), stream);
- funcs[depth](src.reshape(1), dst.reshape(1), stream);
+ syncOutput(dst, _dst, stream);
}
////////////////////////////////////////////////////////////////////////
__device__ __forceinline__ T operator()(T e) const
{
- return saturate_cast<T>(__powf((float)e, power));
+ return cudev::saturate_cast<T>(__powf((float)e, power));
}
};
template<typename T> struct PowOp<T, true> : unary_function<T, T>
__device__ __forceinline__ T operator()(T e) const
{
- T res = saturate_cast<T>(__powf((float)e, power));
+ T res = cudev::saturate_cast<T>(__powf((float)e, power));
if ((e < 0) && (1 & static_cast<int>(power)))
res *= -1;
powMat<double>
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
- const int depth = src.depth();
+ CV_Assert( src.depth() <= CV_64F );
- CV_DbgAssert(depth <= CV_64F);
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ funcs[src.depth()](src.reshape(1), power, dst.reshape(1), stream);
- funcs[depth](src.reshape(1), power, dst.reshape(1), stream);
+ syncOutput(dst, _dst, stream);
}
#endif
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudev.hpp"
+#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
//////////////////////////////////////////////////////////////////////////////
{
(void) flags;
- GpuMat src1 = _src1.getGpuMat();
- GpuMat src2 = _src2.getGpuMat();
+ GpuMat src1 = getInputMat(_src1, stream);
+ GpuMat src2 = getInputMat(_src2, stream);
CV_Assert( src1.type() == src2.type() && src1.type() == CV_32FC2 );
CV_Assert( src1.size() == src2.size() );
- _dst.create(src1.size(), CV_32FC2);
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src1.size(), CV_32FC2, stream);
if (conjB)
gridTransformBinary(globPtr<float2>(src1), globPtr<float2>(src2), globPtr<float2>(dst), comlex_mul_conj(), stream);
else
gridTransformBinary(globPtr<float2>(src1), globPtr<float2>(src2), globPtr<float2>(dst), comlex_mul(), stream);
+
+ syncOutput(dst, _dst, stream);
}
void cv::cuda::mulAndScaleSpectrums(InputArray _src1, InputArray _src2, OutputArray _dst, int flags, float scale, bool conjB, Stream& stream)
{
(void) flags;
- GpuMat src1 = _src1.getGpuMat();
- GpuMat src2 = _src2.getGpuMat();
+ GpuMat src1 = getInputMat(_src1, stream);
+ GpuMat src2 = getInputMat(_src2, stream);
CV_Assert( src1.type() == src2.type() && src1.type() == CV_32FC2);
CV_Assert( src1.size() == src2.size() );
- _dst.create(src1.size(), CV_32FC2);
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src1.size(), CV_32FC2, stream);
if (conjB)
{
op.scale = scale;
gridTransformBinary(globPtr<float2>(src1), globPtr<float2>(src2), globPtr<float2>(dst), op, stream);
}
+
+ syncOutput(dst, _dst, stream);
}
#endif
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudev.hpp"
+#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
void cv::cuda::magnitude(InputArray _x, InputArray _y, OutputArray _dst, Stream& stream)
{
- GpuMat x = _x.getGpuMat();
- GpuMat y = _y.getGpuMat();
+ GpuMat x = getInputMat(_x, stream);
+ GpuMat y = getInputMat(_y, stream);
- CV_DbgAssert( x.depth() == CV_32F );
- CV_DbgAssert( y.type() == x.type() && y.size() == x.size() );
+ CV_Assert( x.depth() == CV_32F );
+ CV_Assert( y.type() == x.type() && y.size() == x.size() );
- _dst.create(x.size(), CV_32FC1);
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, x.size(), CV_32FC1, stream);
GpuMat_<float> xc(x.reshape(1));
GpuMat_<float> yc(y.reshape(1));
GpuMat_<float> magc(dst.reshape(1));
gridTransformBinary(xc, yc, magc, magnitude_func<float>(), stream);
+
+ syncOutput(dst, _dst, stream);
}
void cv::cuda::magnitudeSqr(InputArray _x, InputArray _y, OutputArray _dst, Stream& stream)
{
- GpuMat x = _x.getGpuMat();
- GpuMat y = _y.getGpuMat();
+ GpuMat x = getInputMat(_x, stream);
+ GpuMat y = getInputMat(_y, stream);
- CV_DbgAssert( x.depth() == CV_32F );
- CV_DbgAssert( y.type() == x.type() && y.size() == x.size() );
+ CV_Assert( x.depth() == CV_32F );
+ CV_Assert( y.type() == x.type() && y.size() == x.size() );
- _dst.create(x.size(), CV_32FC1);
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, x.size(), CV_32FC1, stream);
GpuMat_<float> xc(x.reshape(1));
GpuMat_<float> yc(y.reshape(1));
GpuMat_<float> magc(dst.reshape(1));
gridTransformBinary(xc, yc, magc, magnitude_sqr_func<float>(), stream);
+
+ syncOutput(dst, _dst, stream);
}
void cv::cuda::phase(InputArray _x, InputArray _y, OutputArray _dst, bool angleInDegrees, Stream& stream)
{
- GpuMat x = _x.getGpuMat();
- GpuMat y = _y.getGpuMat();
+ GpuMat x = getInputMat(_x, stream);
+ GpuMat y = getInputMat(_y, stream);
- CV_DbgAssert( x.depth() == CV_32F );
- CV_DbgAssert( y.type() == x.type() && y.size() == x.size() );
+ CV_Assert( x.depth() == CV_32F );
+ CV_Assert( y.type() == x.type() && y.size() == x.size() );
- _dst.create(x.size(), CV_32FC1);
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, x.size(), CV_32FC1, stream);
GpuMat_<float> xc(x.reshape(1));
GpuMat_<float> yc(y.reshape(1));
gridTransformBinary(xc, yc, anglec, direction_func<float, true>(), stream);
else
gridTransformBinary(xc, yc, anglec, direction_func<float, false>(), stream);
+
+ syncOutput(dst, _dst, stream);
}
void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, OutputArray _angle, bool angleInDegrees, Stream& stream)
{
- GpuMat x = _x.getGpuMat();
- GpuMat y = _y.getGpuMat();
-
- CV_DbgAssert( x.depth() == CV_32F );
- CV_DbgAssert( y.type() == x.type() && y.size() == x.size() );
+ GpuMat x = getInputMat(_x, stream);
+ GpuMat y = getInputMat(_y, stream);
- _mag.create(x.size(), CV_32FC1);
- GpuMat mag = _mag.getGpuMat();
+ CV_Assert( x.depth() == CV_32F );
+ CV_Assert( y.type() == x.type() && y.size() == x.size() );
- _angle.create(x.size(), CV_32FC1);
- GpuMat angle = _angle.getGpuMat();
+ GpuMat mag = getOutputMat(_mag, x.size(), CV_32FC1, stream);
+ GpuMat angle = getOutputMat(_angle, x.size(), CV_32FC1, stream);
GpuMat_<float> xc(x.reshape(1));
GpuMat_<float> yc(y.reshape(1));
binaryTupleAdapter<0, 1>(direction_func<float, false>())),
stream);
}
+
+ syncOutput(mag, _mag, stream);
+ syncOutput(angle, _angle, stream);
}
namespace
void cv::cuda::polarToCart(InputArray _mag, InputArray _angle, OutputArray _x, OutputArray _y, bool angleInDegrees, Stream& _stream)
{
- GpuMat mag = _mag.getGpuMat();
- GpuMat angle = _angle.getGpuMat();
-
- CV_DbgAssert( angle.depth() == CV_32F );
- CV_DbgAssert( mag.empty() || (mag.type() == angle.type() && mag.size() == angle.size()) );
+ GpuMat mag = getInputMat(_mag, _stream);
+ GpuMat angle = getInputMat(_angle, _stream);
- _x.create(angle.size(), CV_32FC1);
- GpuMat x = _x.getGpuMat();
+ CV_Assert( angle.depth() == CV_32F );
+ CV_Assert( mag.empty() || (mag.type() == angle.type() && mag.size() == angle.size()) );
- _y.create(angle.size(), CV_32FC1);
- GpuMat y = _y.getGpuMat();
+ GpuMat x = getOutputMat(_x, angle.size(), CV_32FC1, _stream);
+ GpuMat y = getOutputMat(_y, angle.size(), CV_32FC1, _stream);
GpuMat_<float> xc(x.reshape(1));
GpuMat_<float> yc(y.reshape(1));
CV_CUDEV_SAFE_CALL( cudaGetLastError() );
+ syncOutput(x, _x, _stream);
+ syncOutput(y, _y, _stream);
+
if (stream == 0)
CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
}
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudev.hpp"
+#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
namespace
void cv::cuda::reduce(InputArray _src, OutputArray _dst, int dim, int reduceOp, int dtype, Stream& stream)
{
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
CV_Assert( src.channels() <= 4 );
CV_Assert( dim == 0 || dim == 1 );
if (dtype < 0)
dtype = src.depth();
- _dst.create(1, dim == 0 ? src.cols : src.rows, CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src.channels()));
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, 1, dim == 0 ? src.cols : src.rows, CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src.channels()), stream);
if (dim == 0)
{
func(src, dst, reduceOp, stream);
}
+
+ syncOutput(dst, _dst, stream);
}
#endif
#include "opencv2/cudaarithm.hpp"
#include "opencv2/cudev.hpp"
+#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
////////////////////////////////////////////////////////////////////////
void mergeImpl(const GpuMat* src, size_t n, cv::OutputArray _dst, Stream& stream)
{
- CV_DbgAssert( src != 0 );
- CV_DbgAssert( n > 0 && n <= 4 );
+ CV_Assert( src != 0 );
+ CV_Assert( n > 0 && n <= 4 );
const int depth = src[0].depth();
const cv::Size size = src[0].size();
-#ifdef _DEBUG
for (size_t i = 0; i < n; ++i)
{
CV_Assert( src[i].size() == size );
CV_Assert( src[i].depth() == depth );
CV_Assert( src[i].channels() == 1 );
}
-#endif
if (n == 1)
{
const int channels = static_cast<int>(n);
- _dst.create(size, CV_MAKE_TYPE(depth, channels));
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, size, CV_MAKE_TYPE(depth, channels), stream);
const func_t func = funcs[channels - 2][CV_ELEM_SIZE(depth) / 2];
CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported channel count or data type");
func(src, dst, stream);
+
+ syncOutput(dst, _dst, stream);
}
}
}
{SplitFunc<4, uchar>::call, SplitFunc<4, ushort>::call, SplitFunc<4, int>::call, 0, SplitFunc<4, double>::call}
};
- CV_DbgAssert( dst != 0 );
+ CV_Assert( dst != 0 );
const int depth = src.depth();
const int channels = src.channels();
- CV_DbgAssert( channels <= 4 );
+ CV_Assert( channels <= 4 );
if (channels == 0)
return;
void cv::cuda::split(InputArray _src, GpuMat* dst, Stream& stream)
{
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
splitImpl(src, dst, stream);
}
void cv::cuda::split(InputArray _src, std::vector<GpuMat>& dst, Stream& stream)
{
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
dst.resize(src.channels());
if (src.channels() > 0)
splitImpl(src, &dst[0], stream);
#include "opencv2/cudev.hpp"
#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
namespace
double cv::cuda::threshold(InputArray _src, OutputArray _dst, double thresh, double maxVal, int type, Stream& stream)
{
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
const int depth = src.depth();
- CV_DbgAssert( src.channels() == 1 && depth <= CV_64F );
- CV_DbgAssert( type <= 4 /*THRESH_TOZERO_INV*/ );
+ CV_Assert( src.channels() == 1 && depth <= CV_64F );
+ CV_Assert( type <= 4 /*THRESH_TOZERO_INV*/ );
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
if (depth == CV_32F && type == 2 /*THRESH_TRUNC*/)
{
funcs[depth](src, dst, thresh, maxVal, type, stream);
}
+ syncOutput(dst, _dst, stream);
+
return thresh;
}
#include "opencv2/cudev.hpp"
#include "opencv2/core/private.cuda.hpp"
+using namespace cv;
+using namespace cv::cuda;
using namespace cv::cudev;
void cv::cuda::transpose(InputArray _src, OutputArray _dst, Stream& stream)
{
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
const size_t elemSize = src.elemSize();
CV_Assert( elemSize == 1 || elemSize == 4 || elemSize == 8 );
- _dst.create( src.cols, src.rows, src.type() );
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.cols, src.rows, src.type(), stream);
if (elemSize == 1)
{
{
gridTranspose(globPtr<double>(src), globPtr<double>(dst), stream);
}
+
+ syncOutput(dst, _dst, stream);
}
#endif
GpuMat src1;
if (!isScalar1)
- src1 = _src1.getGpuMat();
+ src1 = getInputMat(_src1, stream);
GpuMat src2;
if (!isScalar2)
- src2 = _src2.getGpuMat();
+ src2 = getInputMat(_src2, stream);
Mat scalar;
if (isScalar1)
scalar.convertTo(Mat_<double>(scalar.rows, scalar.cols, &val[0]), CV_64F);
}
- GpuMat mask = _mask.getGpuMat();
+ GpuMat mask = getInputMat(_mask, stream);
const int sdepth = src1.empty() ? src2.depth() : src1.depth();
const int cn = src1.empty() ? src2.channels() : src1.channels();
CV_Error(Error::StsUnsupportedFormat, "The device doesn't support double");
}
- _dst.create(size, CV_MAKE_TYPE(ddepth, cn));
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, size, CV_MAKE_TYPE(ddepth, cn), stream);
if (isScalar1)
mat_scalar_func(src2, val, true, dst, mask, scale, stream, op);
mat_scalar_func(src1, val, false, dst, mask, scale, stream, op);
else
mat_mat_func(src1, src2, dst, mask, scale, stream, op);
+
+ syncOutput(dst, _dst, stream);
}
}
{
if (_src1.type() == CV_8UC4 && _src2.type() == CV_32FC1)
{
- GpuMat src1 = _src1.getGpuMat();
- GpuMat src2 = _src2.getGpuMat();
+ GpuMat src1 = getInputMat(_src1, stream);
+ GpuMat src2 = getInputMat(_src2, stream);
CV_Assert( src1.size() == src2.size() );
- _dst.create(src1.size(), src1.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src1.size(), src1.type(), stream);
mulMat_8uc4_32f(src1, src2, dst, stream);
+
+ syncOutput(dst, _dst, stream);
}
else if (_src1.type() == CV_16SC4 && _src2.type() == CV_32FC1)
{
- GpuMat src1 = _src1.getGpuMat();
- GpuMat src2 = _src2.getGpuMat();
+ GpuMat src1 = getInputMat(_src1, stream);
+ GpuMat src2 = getInputMat(_src2, stream);
CV_Assert( src1.size() == src2.size() );
- _dst.create(src1.size(), src1.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src1.size(), src1.type(), stream);
mulMat_16sc4_32f(src1, src2, dst, stream);
+
+ syncOutput(dst, _dst, stream);
}
else
{
{
if (_src1.type() == CV_8UC4 && _src2.type() == CV_32FC1)
{
- GpuMat src1 = _src1.getGpuMat();
- GpuMat src2 = _src2.getGpuMat();
+ GpuMat src1 = getInputMat(_src1, stream);
+ GpuMat src2 = getInputMat(_src2, stream);
CV_Assert( src1.size() == src2.size() );
- _dst.create(src1.size(), src1.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src1.size(), src1.type(), stream);
divMat_8uc4_32f(src1, src2, dst, stream);
+
+ syncOutput(dst, _dst, stream);
}
else if (_src1.type() == CV_16SC4 && _src2.type() == CV_32FC1)
{
- GpuMat src1 = _src1.getGpuMat();
- GpuMat src2 = _src2.getGpuMat();
+ GpuMat src1 = getInputMat(_src1, stream);
+ GpuMat src2 = getInputMat(_src2, stream);
CV_Assert( src1.size() == src2.size() );
- _dst.create(src1.size(), src1.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src1.size(), src1.type(), stream);
divMat_16sc4_32f(src1, src2, dst, stream);
+
+ syncOutput(dst, _dst, stream);
}
else
{
{NppShift<CV_32S, 1, nppiRShiftC_32s_C1R>::call, 0, NppShift<CV_32S, 3, nppiRShiftC_32s_C3R>::call, NppShift<CV_32S, 4, nppiRShiftC_32s_C4R>::call},
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
CV_Assert( src.depth() < CV_32F );
CV_Assert( src.channels() == 1 || src.channels() == 3 || src.channels() == 4 );
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
funcs[src.depth()][src.channels() - 1](src, val, dst, StreamAccessor::getStream(stream));
+
+ syncOutput(dst, _dst, stream);
}
void cv::cuda::lshift(InputArray _src, Scalar_<int> val, OutputArray _dst, Stream& stream)
{NppShift<CV_32S, 1, nppiLShiftC_32s_C1R>::call, 0, NppShift<CV_32S, 3, nppiLShiftC_32s_C3R>::call, NppShift<CV_32S, 4, nppiLShiftC_32s_C4R>::call},
};
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32S );
CV_Assert( src.channels() == 1 || src.channels() == 3 || src.channels() == 4 );
- _dst.create(src.size(), src.type());
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
funcs[src.depth()][src.channels() - 1](src, val, dst, StreamAccessor::getStream(stream));
+
+ syncOutput(dst, _dst, stream);
}
//////////////////////////////////////////////////////////////////////////////
void cv::cuda::magnitude(InputArray _src, OutputArray _dst, Stream& stream)
{
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
- _dst.create(src.size(), CV_32FC1);
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.size(), CV_32FC1, stream);
npp_magnitude(src, dst, nppiMagnitude_32fc32f_C1R, StreamAccessor::getStream(stream));
+
+ syncOutput(dst, _dst, stream);
}
void cv::cuda::magnitudeSqr(InputArray _src, OutputArray _dst, Stream& stream)
{
- GpuMat src = _src.getGpuMat();
+ GpuMat src = getInputMat(_src, stream);
- _dst.create(src.size(), CV_32FC1);
- GpuMat dst = _dst.getGpuMat();
+ GpuMat dst = getOutputMat(_dst, src.size(), CV_32FC1, stream);
npp_magnitude(src, dst, nppiMagnitudeSqr_32fc32f_C1R, StreamAccessor::getStream(stream));
+
+ syncOutput(dst, _dst, stream);
}
#endif