{\r
#ifndef HAVE_CUBLAS\r
\r
- OPENCV_GPU_UNUSED(src1);\r
- OPENCV_GPU_UNUSED(src2);\r
- OPENCV_GPU_UNUSED(alpha);\r
- OPENCV_GPU_UNUSED(src3);\r
- OPENCV_GPU_UNUSED(beta);\r
- OPENCV_GPU_UNUSED(dst);\r
- OPENCV_GPU_UNUSED(flags);\r
- OPENCV_GPU_UNUSED(stream);\r
-\r
- throw_nogpu();\r
+ CV_Error(CV_StsNotImplemented, "The library was build without CUBLAS");\r
\r
#else\r
\r
CV_Assert(src1.type() == CV_32FC1 || src1.type() == CV_32FC2 || src1.type() == CV_64FC1 || src1.type() == CV_64FC2);\r
CV_Assert(src2.type() == src1.type() && (src3.empty() || src3.type() == src1.type()));\r
\r
+ if (src1.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
bool tr1 = (flags & GEMM_1_T) != 0;\r
bool tr2 = (flags & GEMM_2_T) != 0;\r
bool tr3 = (flags & GEMM_3_T) != 0;\r
}\r
else // if (src.elemSize() == 8)\r
{\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+\r
NppStStreamHandler h(stream);\r
\r
NcvSize32u sz;\r
void cv::gpu::flip(const GpuMat& src, GpuMat& dst, int flipCode, Stream& stream)\r
{\r
typedef void (*func_t)(const GpuMat& src, GpuMat& dst, int flipCode, cudaStream_t stream);\r
-\r
static const func_t funcs[6][4] =\r
{\r
{NppMirror<CV_8U, nppiMirror_8u_C1R>::call, 0, NppMirror<CV_8U, nppiMirror_8u_C3R>::call, NppMirror<CV_8U, nppiMirror_8u_C4R>::call},\r
\r
void cv::gpu::magnitude(const GpuMat& src, GpuMat& dst, Stream& stream)\r
{\r
- ::npp_magnitude(src, dst, nppiMagnitude_32fc32f_C1R, StreamAccessor::getStream(stream));\r
+ npp_magnitude(src, dst, nppiMagnitude_32fc32f_C1R, StreamAccessor::getStream(stream));\r
}\r
\r
void cv::gpu::magnitudeSqr(const GpuMat& src, GpuMat& dst, Stream& stream)\r
{\r
- ::npp_magnitude(src, dst, nppiMagnitudeSqr_32fc32f_C1R, StreamAccessor::getStream(stream));\r
+ npp_magnitude(src, dst, nppiMagnitudeSqr_32fc32f_C1R, StreamAccessor::getStream(stream));\r
}\r
\r
////////////////////////////////////////////////////////////////////////\r
{\r
using namespace ::cv::gpu::device::mathfunc;\r
\r
- CV_DbgAssert(x.size() == y.size() && x.type() == y.type());\r
+ CV_Assert(x.size() == y.size() && x.type() == y.type());\r
CV_Assert(x.depth() == CV_32F);\r
\r
if (mag)\r
{\r
using namespace ::cv::gpu::device::mathfunc;\r
\r
- CV_DbgAssert((mag.empty() || mag.size() == angle.size()) && mag.type() == angle.type());\r
+ CV_Assert((mag.empty() || mag.size() == angle.size()) && mag.type() == angle.type());\r
CV_Assert(mag.depth() == CV_32F);\r
\r
x.create(mag.size(), mag.type());\r
enum { smart_shift = 4 };\r
};\r
\r
- template <typename T> void absdiff_gpu(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream)\r
+ template <typename T> void absdiff_gpu(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream)\r
{\r
cv::gpu::device::transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<T>)dst, Absdiff<T>(), WithOutMask(), stream);\r
}\r
\r
- template void absdiff_gpu<uchar >(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<schar >(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<ushort>(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<short >(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<int >(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<float >(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<double>(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
+ //template void absdiff_gpu<uchar >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void absdiff_gpu<schar >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ //template void absdiff_gpu<ushort>(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void absdiff_gpu<short >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void absdiff_gpu<int >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ //template void absdiff_gpu<float >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void absdiff_gpu<double>(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
\r
template <typename T> struct AbsdiffScalar : unary_function<T, T>\r
{\r
enum { smart_shift = 4 };\r
};\r
\r
- template <typename T> void absdiff_gpu(const DevMem2Db& src1, double val, const DevMem2Db& dst, cudaStream_t stream)\r
+ template <typename T> void absdiff_gpu(const DevMem2Db src1, double val, DevMem2Db dst, cudaStream_t stream)\r
{\r
cudaSafeCall( cudaSetDoubleForDevice(&val) );\r
AbsdiffScalar<T> op(val);\r
cv::gpu::device::transform((DevMem2D_<T>)src1, (DevMem2D_<T>)dst, op, WithOutMask(), stream);\r
}\r
\r
- //template void absdiff_gpu<uchar >(const DevMem2Db& src1, double src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<schar >(const DevMem2Db& src1, double src2, const DevMem2Db& dst, cudaStream_t stream);\r
- //template void absdiff_gpu<ushort>(const DevMem2Db& src1, double src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<short >(const DevMem2Db& src1, double src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<int >(const DevMem2Db& src1, double src2, const DevMem2Db& dst, cudaStream_t stream);\r
- //template void absdiff_gpu<float >(const DevMem2Db& src1, double src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void absdiff_gpu<double>(const DevMem2Db& src1, double src2, const DevMem2Db& dst, cudaStream_t stream);\r
+ //template void absdiff_gpu<uchar >(const DevMem2Db src1, double src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void absdiff_gpu<schar >(const DevMem2Db src1, double src2, DevMem2Db dst, cudaStream_t stream);\r
+ //template void absdiff_gpu<ushort>(const DevMem2Db src1, double src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void absdiff_gpu<short >(const DevMem2Db src1, double src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void absdiff_gpu<int >(const DevMem2Db src1, double src2, DevMem2Db dst, cudaStream_t stream);\r
+ //template void absdiff_gpu<float >(const DevMem2Db src1, double src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void absdiff_gpu<double>(const DevMem2Db src1, double src2, DevMem2Db dst, cudaStream_t stream);\r
\r
//////////////////////////////////////////////////////////////////////////////////////\r
// Compare\r
};\r
\r
template <typename T>\r
- void min_gpu(const DevMem2D_<T>& src1, const DevMem2D_<T>& src2, const DevMem2D_<T>& dst, cudaStream_t stream)\r
+ void min_gpu(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream)\r
{\r
- cv::gpu::device::transform(src1, src2, dst, minimum<T>(), WithOutMask(), stream);\r
+ cv::gpu::device::transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<T>)dst, minimum<T>(), WithOutMask(), stream);\r
}\r
\r
- template void min_gpu<uchar >(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void min_gpu<schar >(const DevMem2D_<schar>& src1, const DevMem2D_<schar>& src2, const DevMem2D_<schar>& dst, cudaStream_t stream);\r
- template void min_gpu<ushort>(const DevMem2D_<ushort>& src1, const DevMem2D_<ushort>& src2, const DevMem2D_<ushort>& dst, cudaStream_t stream);\r
- template void min_gpu<short >(const DevMem2D_<short>& src1, const DevMem2D_<short>& src2, const DevMem2D_<short>& dst, cudaStream_t stream);\r
- template void min_gpu<int >(const DevMem2D_<int>& src1, const DevMem2D_<int>& src2, const DevMem2D_<int>& dst, cudaStream_t stream);\r
- template void min_gpu<float >(const DevMem2D_<float>& src1, const DevMem2D_<float>& src2, const DevMem2D_<float>& dst, cudaStream_t stream);\r
- template void min_gpu<double>(const DevMem2D_<double>& src1, const DevMem2D_<double>& src2, const DevMem2D_<double>& dst, cudaStream_t stream);\r
+ template void min_gpu<uchar >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<schar >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<ushort>(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<short >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<int >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<float >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<double>(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
\r
template <typename T>\r
- void max_gpu(const DevMem2D_<T>& src1, const DevMem2D_<T>& src2, const DevMem2D_<T>& dst, cudaStream_t stream)\r
+ void max_gpu(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream)\r
{\r
- cv::gpu::device::transform(src1, src2, dst, maximum<T>(), WithOutMask(), stream);\r
+ cv::gpu::device::transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<T>)dst, maximum<T>(), WithOutMask(), stream);\r
}\r
\r
- template void max_gpu<uchar >(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void max_gpu<schar >(const DevMem2D_<schar>& src1, const DevMem2D_<schar>& src2, const DevMem2D_<schar>& dst, cudaStream_t stream);\r
- template void max_gpu<ushort>(const DevMem2D_<ushort>& src1, const DevMem2D_<ushort>& src2, const DevMem2D_<ushort>& dst, cudaStream_t stream);\r
- template void max_gpu<short >(const DevMem2D_<short>& src1, const DevMem2D_<short>& src2, const DevMem2D_<short>& dst, cudaStream_t stream);\r
- template void max_gpu<int >(const DevMem2D_<int>& src1, const DevMem2D_<int>& src2, const DevMem2D_<int>& dst, cudaStream_t stream);\r
- template void max_gpu<float >(const DevMem2D_<float>& src1, const DevMem2D_<float>& src2, const DevMem2D_<float>& dst, cudaStream_t stream);\r
- template void max_gpu<double>(const DevMem2D_<double>& src1, const DevMem2D_<double>& src2, const DevMem2D_<double>& dst, cudaStream_t stream);\r
+ template void max_gpu<uchar >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<schar >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<ushort>(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<short >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<int >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<float >(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<double>(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
\r
template <typename T>\r
- void min_gpu(const DevMem2D_<T>& src1, T src2, const DevMem2D_<T>& dst, cudaStream_t stream)\r
+ void min_gpu(const DevMem2Db src, T val, DevMem2Db dst, cudaStream_t stream)\r
{\r
- cv::gpu::device::transform(src1, dst, device::bind2nd(minimum<T>(), src2), WithOutMask(), stream);\r
+ cv::gpu::device::transform((DevMem2D_<T>)src, (DevMem2D_<T>)dst, device::bind2nd(minimum<T>(), val), WithOutMask(), stream);\r
}\r
\r
- template void min_gpu<uchar >(const DevMem2Db& src1, uchar src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void min_gpu<schar >(const DevMem2D_<schar>& src1, schar src2, const DevMem2D_<schar>& dst, cudaStream_t stream);\r
- template void min_gpu<ushort>(const DevMem2D_<ushort>& src1, ushort src2, const DevMem2D_<ushort>& dst, cudaStream_t stream);\r
- template void min_gpu<short >(const DevMem2D_<short>& src1, short src2, const DevMem2D_<short>& dst, cudaStream_t stream);\r
- template void min_gpu<int >(const DevMem2D_<int>& src1, int src2, const DevMem2D_<int>& dst, cudaStream_t stream);\r
- template void min_gpu<float >(const DevMem2D_<float>& src1, float src2, const DevMem2D_<float>& dst, cudaStream_t stream);\r
- template void min_gpu<double>(const DevMem2D_<double>& src1, double src2, const DevMem2D_<double>& dst, cudaStream_t stream);\r
+ template void min_gpu<uchar >(const DevMem2Db src, uchar val, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<schar >(const DevMem2Db src, schar val, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<ushort>(const DevMem2Db src, ushort val, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<short >(const DevMem2Db src, short val, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<int >(const DevMem2Db src, int val, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<float >(const DevMem2Db src, float val, DevMem2Db dst, cudaStream_t stream);\r
+ template void min_gpu<double>(const DevMem2Db src, double val, DevMem2Db dst, cudaStream_t stream);\r
\r
template <typename T>\r
- void max_gpu(const DevMem2D_<T>& src1, T src2, const DevMem2D_<T>& dst, cudaStream_t stream)\r
+ void max_gpu(const DevMem2Db src, T val, DevMem2Db dst, cudaStream_t stream)\r
{\r
- cv::gpu::device::transform(src1, dst, device::bind2nd(maximum<T>(), src2), WithOutMask(), stream);\r
+ cv::gpu::device::transform((DevMem2D_<T>)src, (DevMem2D_<T>)dst, device::bind2nd(maximum<T>(), val), WithOutMask(), stream);\r
}\r
\r
- template void max_gpu<uchar >(const DevMem2Db& src1, uchar src2, const DevMem2Db& dst, cudaStream_t stream);\r
- template void max_gpu<schar >(const DevMem2D_<schar>& src1, schar src2, const DevMem2D_<schar>& dst, cudaStream_t stream);\r
- template void max_gpu<ushort>(const DevMem2D_<ushort>& src1, ushort src2, const DevMem2D_<ushort>& dst, cudaStream_t stream);\r
- template void max_gpu<short >(const DevMem2D_<short>& src1, short src2, const DevMem2D_<short>& dst, cudaStream_t stream);\r
- template void max_gpu<int >(const DevMem2D_<int>& src1, int src2, const DevMem2D_<int>& dst, cudaStream_t stream);\r
- template void max_gpu<float >(const DevMem2D_<float>& src1, float src2, const DevMem2D_<float>& dst, cudaStream_t stream);\r
- template void max_gpu<double>(const DevMem2D_<double>& src1, double src2, const DevMem2D_<double>& dst, cudaStream_t stream);\r
+ template void max_gpu<uchar >(const DevMem2Db src, uchar val, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<schar >(const DevMem2Db src, schar val, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<ushort>(const DevMem2Db src, ushort val, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<short >(const DevMem2Db src, short val, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<int >(const DevMem2Db src, int val, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<float >(const DevMem2Db src, float val, DevMem2Db dst, cudaStream_t stream);\r
+ template void max_gpu<double>(const DevMem2Db src, double val, DevMem2Db dst, cudaStream_t stream);\r
\r
//////////////////////////////////////////////////////////////////////////\r
// threshold\r
//////////////////////////////////////////////////////////////////////////\r
// addWeighted\r
\r
- template <typename T1, typename T2, typename D> struct AddWeighted : binary_function<T1, T2, D>\r
+ namespace detail\r
{\r
- __host__ __device__ __forceinline__ AddWeighted(double alpha_, double beta_, double gamma_) : alpha(alpha_), beta(beta_), gamma(gamma_) {}\r
+ template <typename T> struct UseDouble\r
+ {\r
+ enum {value = 0};\r
+ };\r
+ template <> struct UseDouble<int>\r
+ {\r
+ enum {value = 1};\r
+ };\r
+ template <> struct UseDouble<float>\r
+ {\r
+ enum {value = 1};\r
+ };\r
+ template <> struct UseDouble<double>\r
+ {\r
+ enum {value = 1};\r
+ };\r
+ }\r
+ template <typename T1, typename T2, typename D> struct UseDouble\r
+ {\r
+ enum {value = (detail::UseDouble<T1>::value || detail::UseDouble<T2>::value || detail::UseDouble<D>::value)};\r
+ };\r
\r
- __device__ __forceinline__ D operator ()(typename TypeTraits<T1>::ParameterType a, typename TypeTraits<T2>::ParameterType b) const\r
+ namespace detail\r
+ {\r
+ template <typename T1, typename T2, typename D, bool useDouble> struct AddWeighted;\r
+ template <typename T1, typename T2, typename D> struct AddWeighted<T1, T2, D, false> : binary_function<T1, T2, D>\r
{\r
- return saturate_cast<D>(alpha * a + beta * b + gamma);\r
- }\r
+ AddWeighted(double alpha_, double beta_, double gamma_) : alpha(static_cast<float>(alpha_)), beta(static_cast<float>(beta_)), gamma(static_cast<float>(gamma_)) {}\r
\r
- const double alpha;\r
- const double beta;\r
- const double gamma;\r
+ __device__ __forceinline__ D operator ()(T1 a, T2 b) const\r
+ {\r
+ return saturate_cast<D>(a * alpha + b * beta + gamma);\r
+ }\r
+\r
+ const float alpha;\r
+ const float beta;\r
+ const float gamma;\r
+ };\r
+ template <typename T1, typename T2, typename D> struct AddWeighted<T1, T2, D, true> : binary_function<T1, T2, D>\r
+ {\r
+ AddWeighted(double alpha_, double beta_, double gamma_) : alpha(alpha_), beta(beta_), gamma(gamma_) {}\r
+\r
+ __device__ __forceinline__ D operator ()(T1 a, T2 b) const\r
+ {\r
+ return saturate_cast<D>(a * alpha + b * beta + gamma);\r
+ }\r
+\r
+ const double alpha;\r
+ const double beta;\r
+ const double gamma;\r
+ };\r
+ }\r
+ template <typename T1, typename T2, typename D> struct AddWeighted : detail::AddWeighted<T1, T2, D, UseDouble<T1, T2, D>::value>\r
+ {\r
+ AddWeighted(double alpha_, double beta_, double gamma_) : detail::AddWeighted<T1, T2, D, UseDouble<T1, T2, D>::value>(alpha_, beta_, gamma_) {}\r
};\r
\r
template <> struct TransformFunctorTraits< AddWeighted<ushort, ushort, ushort> > : DefaultTransformFunctorTraits< AddWeighted<ushort, ushort, ushort> >\r
template <typename T1, typename T2, typename D>\r
void addWeighted_gpu(const DevMem2Db& src1, double alpha, const DevMem2Db& src2, double beta, double gamma, const DevMem2Db& dst, cudaStream_t stream)\r
{\r
- cudaSafeCall( cudaSetDoubleForDevice(&alpha) );\r
- cudaSafeCall( cudaSetDoubleForDevice(&beta) );\r
- cudaSafeCall( cudaSetDoubleForDevice(&gamma) );\r
+ if (UseDouble<T1, T2, D>::value)\r
+ {\r
+ cudaSafeCall( cudaSetDoubleForDevice(&alpha) );\r
+ cudaSafeCall( cudaSetDoubleForDevice(&beta) );\r
+ cudaSafeCall( cudaSetDoubleForDevice(&gamma) );\r
+ }\r
\r
AddWeighted<T1, T2, D> op(alpha, beta, gamma);\r
\r
namespace cv { namespace gpu { namespace device\r
{\r
template <typename T>\r
- void absdiff_gpu(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
+ void absdiff_gpu(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
\r
template <typename T>\r
- void absdiff_gpu(const DevMem2Db& src1, double val, const DevMem2Db& dst, cudaStream_t stream);\r
+ void absdiff_gpu(const DevMem2Db src1, double val, DevMem2Db dst, cudaStream_t stream);\r
}}}\r
\r
-void cv::gpu::absdiff(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& s)\r
+namespace\r
{\r
- using namespace ::cv::gpu::device;\r
-\r
- typedef void (*func_t)(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
-\r
- static const func_t funcs[] =\r
- {\r
- absdiff_gpu<unsigned char>, absdiff_gpu<signed char>, absdiff_gpu<unsigned short>, absdiff_gpu<short>, absdiff_gpu<int>, absdiff_gpu<float>, absdiff_gpu<double>\r
- };\r
-\r
- CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
-\r
- dst.create( src1.size(), src1.type() );\r
-\r
- cudaStream_t stream = StreamAccessor::getStream(s);\r
-\r
- NppiSize sz;\r
- sz.width = src1.cols * src1.channels();\r
- sz.height = src1.rows;\r
-\r
- if (src1.depth() == CV_8U)\r
+ template <int DEPTH> struct NppAbsDiffFunc\r
{\r
- NppStreamHandler h(stream);\r
+ typedef typename NppTypeTraits<DEPTH>::npp_t npp_t;\r
\r
- nppSafeCall( nppiAbsDiff_8u_C1R(src1.ptr<Npp8u>(), static_cast<int>(src1.step), src2.ptr<Npp8u>(), static_cast<int>(src2.step),\r
- dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz) );\r
+ typedef NppStatus (*func_t)(const npp_t* src1, int src1_step, const npp_t* src2, int src2_step, npp_t* dst, int dst_step, NppiSize sz);\r
+ };\r
\r
- if (stream == 0)\r
- cudaSafeCall( cudaDeviceSynchronize() );\r
- }\r
- else if (src1.depth() == CV_16U)\r
+ template <int DEPTH, typename NppAbsDiffFunc<DEPTH>::func_t func> struct NppAbsDiff\r
{\r
- NppStreamHandler h(stream);\r
+ typedef typename NppAbsDiffFunc<DEPTH>::npp_t npp_t;\r
\r
- nppSafeCall( nppiAbsDiff_16u_C1R(src1.ptr<Npp16u>(), static_cast<int>(src1.step), src2.ptr<Npp16u>(), static_cast<int>(src2.step),\r
- dst.ptr<Npp16u>(), static_cast<int>(dst.step), sz) );\r
-\r
- if (stream == 0)\r
- cudaSafeCall( cudaDeviceSynchronize() );\r
- }\r
- else if (src1.depth() == CV_32F)\r
- {\r
- NppStreamHandler h(stream);\r
+ static void call(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream)\r
+ {\r
+ NppStreamHandler h(stream);\r
\r
- nppSafeCall( nppiAbsDiff_32f_C1R(src1.ptr<Npp32f>(), static_cast<int>(src1.step), src2.ptr<Npp32f>(), static_cast<int>(src2.step),\r
- dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
+ NppiSize sz;\r
+ sz.width = src1.cols;\r
+ sz.height = src1.rows;\r
\r
- if (stream == 0)\r
- cudaSafeCall( cudaDeviceSynchronize() );\r
- }\r
- else\r
- {\r
- const func_t func = funcs[src1.depth()];\r
- CV_Assert(func != 0);\r
+ nppSafeCall( func((const npp_t*)src1.data, static_cast<int>(src1.step), (const npp_t*)src2.data, static_cast<int>(src2.step),\r
+ (npp_t*)dst.data, static_cast<int>(dst.step), sz) );\r
\r
- func(src1.reshape(1), src2.reshape(1), dst.reshape(1), stream);\r
- }\r
-}\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaDeviceSynchronize() );\r
+ }\r
+ };\r
\r
-namespace\r
-{\r
template <int DEPTH> struct NppAbsDiffCFunc\r
{\r
typedef typename NppTypeTraits<DEPTH>::npp_t npp_t;\r
+ typedef npp_t scalar_t;\r
\r
typedef NppStatus (*func_t)(const npp_t* pSrc1, int nSrc1Step, npp_t* pDst, int nDstStep, NppiSize oSizeROI, npp_t nConstant);\r
};\r
template <> struct NppAbsDiffCFunc<CV_16U>\r
{\r
+ typedef NppTypeTraits<CV_16U>::npp_t npp_t;\r
+ typedef Npp32u scalar_t;\r
+\r
typedef NppStatus (*func_t)(const Npp16u* pSrc1, int nSrc1Step, Npp16u* pDst, int nDstStep, NppiSize oSizeROI, Npp32u nConstant);\r
};\r
\r
template <int DEPTH, typename NppAbsDiffCFunc<DEPTH>::func_t func> struct NppAbsDiffC\r
{\r
- typedef typename NppTypeTraits<DEPTH>::npp_t npp_t;\r
+ typedef typename NppAbsDiffCFunc<DEPTH>::npp_t npp_t;\r
+ typedef typename NppAbsDiffCFunc<DEPTH>::scalar_t scalar_t;\r
\r
- static void call(const DevMem2Db& src1, double val, const DevMem2Db& dst, cudaStream_t stream)\r
+ static void call(const DevMem2Db src1, double val, DevMem2Db dst, cudaStream_t stream)\r
{\r
NppStreamHandler h(stream);\r
\r
sz.width = src1.cols;\r
sz.height = src1.rows;\r
\r
- nppSafeCall( func((const npp_t*)src1.data, static_cast<int>(src1.step), (npp_t*)dst.data, static_cast<int>(dst.step),\r
- sz, static_cast<npp_t>(val)) );\r
+ nppSafeCall( func((const npp_t*)src1.data, static_cast<int>(src1.step),\r
+ (npp_t*)dst.data, static_cast<int>(dst.step), sz, static_cast<scalar_t>(val)) );\r
\r
if (stream == 0)\r
cudaSafeCall( cudaDeviceSynchronize() );\r
};\r
}\r
\r
-void cv::gpu::absdiff(const GpuMat& src1, const Scalar& src2, GpuMat& dst, Stream& s)\r
+void cv::gpu::absdiff(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)\r
{\r
using namespace cv::gpu::device;\r
\r
- typedef void (*func_t)(const DevMem2Db& src1, double val, const DevMem2Db& dst, cudaStream_t stream);\r
+ typedef void (*func_t)(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ static const func_t funcs[] =\r
+ {\r
+ NppAbsDiff<CV_8U, nppiAbsDiff_8u_C1R>::call,\r
+ absdiff_gpu<signed char>,\r
+ NppAbsDiff<CV_16U, nppiAbsDiff_16u_C1R>::call,\r
+ absdiff_gpu<short>,\r
+ absdiff_gpu<int>,\r
+ NppAbsDiff<CV_32F, nppiAbsDiff_32f_C1R>::call,\r
+ absdiff_gpu<double>\r
+ };\r
+\r
+ CV_Assert(src1.depth() <= CV_64F);\r
+ CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
+\r
+ if (src1.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
+ dst.create(src1.size(), src1.type());\r
+\r
+ funcs[src1.depth()](src1.reshape(1), src2.reshape(1), dst.reshape(1), StreamAccessor::getStream(stream));\r
+}\r
\r
+void cv::gpu::absdiff(const GpuMat& src1, const Scalar& src2, GpuMat& dst, Stream& stream)\r
+{\r
+ using namespace cv::gpu::device;\r
+\r
+ typedef void (*func_t)(const DevMem2Db src1, double val, DevMem2Db dst, cudaStream_t stream);\r
static const func_t funcs[] =\r
{\r
NppAbsDiffC<CV_8U, nppiAbsDiffC_8u_C1R>::call,\r
absdiff_gpu<double>\r
};\r
\r
+ CV_Assert(src1.depth() <= CV_64F);\r
CV_Assert(src1.channels() == 1);\r
\r
- dst.create(src1.size(), src1.type());\r
+ if (src1.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
\r
- cudaStream_t stream = StreamAccessor::getStream(s);\r
+ dst.create(src1.size(), src1.type());\r
\r
- funcs[src1.depth()](src1, src2.val[0], dst, stream);\r
+ funcs[src1.depth()](src1, src2.val[0], dst, StreamAccessor::getStream(stream));\r
}\r
\r
//////////////////////////////////////////////////////////////////////////////\r
\r
void cv::gpu::compare(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, int cmpop, Stream& stream)\r
{\r
- using namespace ::cv::gpu::device;\r
+ using namespace cv::gpu::device;\r
\r
typedef void (*func_t)(const DevMem2Db& src1, const DevMem2Db& src2, const DevMem2Db& dst, cudaStream_t stream);\r
-\r
static const func_t funcs[7][4] =\r
{\r
- {compare_eq<unsigned char>, compare_ne<unsigned char>, compare_lt<unsigned char>, compare_le<unsigned char>},\r
- {compare_eq<signed char>, compare_ne<signed char>, compare_lt<signed char>, compare_le<signed char>},\r
+ {compare_eq<unsigned char> , compare_ne<unsigned char> , compare_lt<unsigned char> , compare_le<unsigned char> },\r
+ {compare_eq<signed char> , compare_ne<signed char> , compare_lt<signed char> , compare_le<signed char> },\r
{compare_eq<unsigned short>, compare_ne<unsigned short>, compare_lt<unsigned short>, compare_le<unsigned short>},\r
- {compare_eq<short>, compare_ne<short>, compare_lt<short>, compare_le<short>},\r
- {compare_eq<int>, compare_ne<int>, compare_lt<int>, compare_le<int>},\r
- {compare_eq<float>, compare_ne<float>, compare_lt<float>, compare_le<float>},\r
- {compare_eq<double>, compare_ne<double>, compare_lt<double>, compare_le<double>}\r
+ {compare_eq<short> , compare_ne<short> , compare_lt<short> , compare_le<short> },\r
+ {compare_eq<int> , compare_ne<int> , compare_lt<int> , compare_le<int> },\r
+ {compare_eq<float> , compare_ne<float> , compare_lt<float> , compare_le<float> },\r
+ {compare_eq<double> , compare_ne<double> , compare_lt<double> , compare_le<double> }\r
};\r
\r
+ CV_Assert(src1.depth() <= CV_64F);\r
CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
CV_Assert(cmpop >= CMP_EQ && cmpop <= CMP_NE);\r
\r
+ if (src1.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
static const int codes[] =\r
{\r
0, 2, 3, 2, 3, 1\r
};\r
-\r
const GpuMat* psrc1[] =\r
{\r
&src1, &src2, &src2, &src1, &src1, &src1\r
};\r
-\r
const GpuMat* psrc2[] =\r
{\r
&src2, &src1, &src1, &src2, &src2, &src2\r
{\r
dst.create(src.size(), src.type());\r
\r
- ::cv::gpu::device::bitwiseNotCaller(src.rows, src.cols, src.elemSize1(), dst.channels(), src, dst, stream);\r
+ cv::gpu::device::bitwiseNotCaller(src.rows, src.cols, src.elemSize1(), dst.channels(), src, dst, stream);\r
}\r
\r
-\r
void bitwiseNotCaller(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)\r
{\r
- using namespace ::cv::gpu::device;\r
+ using namespace cv::gpu::device;\r
\r
- typedef void (*Caller)(int, int, int, const PtrStepb, const PtrStepb, PtrStepb, cudaStream_t);\r
-\r
- static Caller callers[] =\r
+ typedef void (*func_t)(int, int, int, const PtrStepb, const PtrStepb, PtrStepb, cudaStream_t);\r
+ static func_t funcs[] =\r
{\r
bitwiseMaskNotCaller<unsigned char>, bitwiseMaskNotCaller<unsigned char>,\r
bitwiseMaskNotCaller<unsigned short>, bitwiseMaskNotCaller<unsigned short>,\r
bitwiseMaskNotCaller<unsigned int>\r
};\r
\r
+ CV_Assert(src.depth() <= CV_64F);\r
CV_Assert(mask.type() == CV_8U && mask.size() == src.size());\r
+\r
dst.create(src.size(), src.type());\r
\r
- Caller caller = callers[src.depth()];\r
- CV_Assert(caller);\r
+ const func_t func = funcs[src.depth()];\r
\r
int cn = src.depth() != CV_64F ? src.channels() : src.channels() * (sizeof(double) / sizeof(unsigned int));\r
- caller(src.rows, src.cols, cn, src, mask, dst, stream);\r
- }\r
\r
+ func(src.rows, src.cols, cn, src, mask, dst, stream);\r
+ }\r
}\r
\r
-\r
void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask, Stream& stream)\r
{\r
if (mask.empty())\r
bitwiseNotCaller(src, dst, mask, StreamAccessor::getStream(stream));\r
}\r
\r
-\r
//////////////////////////////////////////////////////////////////////////////\r
// Binary bitwise logical operations\r
\r
void bitwiseOrCaller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream)\r
{\r
CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
+\r
dst.create(src1.size(), src1.type());\r
\r
- ::cv::gpu::device::bitwiseOrCaller(dst.rows, dst.cols, dst.elemSize1(), dst.channels(), src1, src2, dst, stream);\r
+ cv::gpu::device::bitwiseOrCaller(dst.rows, dst.cols, dst.elemSize1(), dst.channels(), src1, src2, dst, stream);\r
}\r
\r
void bitwiseOrCaller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)\r
{\r
- using namespace ::cv::gpu::device;\r
-\r
- typedef void (*Caller)(int, int, int, const PtrStepb, const PtrStepb, const PtrStepb, PtrStepb, cudaStream_t);\r
+ using namespace cv::gpu::device;\r
\r
- static Caller callers[] =\r
+ typedef void (*func_t)(int, int, int, const PtrStepb, const PtrStepb, const PtrStepb, PtrStepb, cudaStream_t);\r
+ static func_t funcs[] =\r
{\r
bitwiseMaskOrCaller<unsigned char>, bitwiseMaskOrCaller<unsigned char>,\r
bitwiseMaskOrCaller<unsigned short>, bitwiseMaskOrCaller<unsigned short>,\r
bitwiseMaskOrCaller<unsigned int>\r
};\r
\r
+ CV_Assert(src1.depth() <= CV_64F);\r
CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
+ CV_Assert(mask.type() == CV_8U && mask.size() == src1.size());\r
+\r
dst.create(src1.size(), src1.type());\r
\r
- Caller caller = callers[src1.depth()];\r
- CV_Assert(caller);\r
+ const func_t func = funcs[src1.depth()];\r
\r
int cn = dst.depth() != CV_64F ? dst.channels() : dst.channels() * (sizeof(double) / sizeof(unsigned int));\r
- caller(dst.rows, dst.cols, cn, src1, src2, mask, dst, stream);\r
+\r
+ func(dst.rows, dst.cols, cn, src1, src2, mask, dst, stream);\r
}\r
\r
\r
void bitwiseAndCaller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream)\r
{\r
CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
+\r
dst.create(src1.size(), src1.type());\r
\r
- ::cv::gpu::device::bitwiseAndCaller(dst.rows, dst.cols, dst.elemSize1(), dst.channels(), src1, src2, dst, stream);\r
+ cv::gpu::device::bitwiseAndCaller(dst.rows, dst.cols, dst.elemSize1(), dst.channels(), src1, src2, dst, stream);\r
}\r
\r
-\r
void bitwiseAndCaller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)\r
{\r
- using namespace ::cv::gpu::device;\r
+ using namespace cv::gpu::device;\r
\r
- typedef void (*Caller)(int, int, int, const PtrStepb, const PtrStepb, const PtrStepb, PtrStepb, cudaStream_t);\r
-\r
- static Caller callers[] =\r
+ typedef void (*func_t)(int, int, int, const PtrStepb, const PtrStepb, const PtrStepb, PtrStepb, cudaStream_t);\r
+ static func_t funcs[] =\r
{\r
bitwiseMaskAndCaller<unsigned char>, bitwiseMaskAndCaller<unsigned char>,\r
bitwiseMaskAndCaller<unsigned short>, bitwiseMaskAndCaller<unsigned short>,\r
bitwiseMaskAndCaller<unsigned int>\r
};\r
\r
+ CV_Assert(src1.depth() <= CV_64F);\r
CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
+ CV_Assert(mask.type() == CV_8U && mask.size() == src1.size());\r
+\r
dst.create(src1.size(), src1.type());\r
\r
- Caller caller = callers[src1.depth()];\r
- CV_Assert(caller);\r
+ const func_t func = funcs[src1.depth()];\r
\r
int cn = dst.depth() != CV_64F ? dst.channels() : dst.channels() * (sizeof(double) / sizeof(unsigned int));\r
- caller(dst.rows, dst.cols, cn, src1, src2, mask, dst, stream);\r
+\r
+ func(dst.rows, dst.cols, cn, src1, src2, mask, dst, stream);\r
}\r
\r
\r
void bitwiseXorCaller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream)\r
{\r
CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
+\r
dst.create(src1.size(), src1.type());\r
\r
- ::cv::gpu::device::bitwiseXorCaller(dst.rows, dst.cols, dst.elemSize1(), dst.channels(), src1, src2, dst, stream);\r
+ cv::gpu::device::bitwiseXorCaller(dst.rows, dst.cols, dst.elemSize1(), dst.channels(), src1, src2, dst, stream);\r
}\r
\r
-\r
void bitwiseXorCaller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)\r
{\r
- using namespace ::cv::gpu::device;\r
+ using namespace cv::gpu::device;\r
\r
- typedef void (*Caller)(int, int, int, const PtrStepb, const PtrStepb, const PtrStepb, PtrStepb, cudaStream_t);\r
-\r
- static Caller callers[] =\r
+ typedef void (*func_t)(int, int, int, const PtrStepb, const PtrStepb, const PtrStepb, PtrStepb, cudaStream_t);\r
+ static func_t funcs[] =\r
{\r
bitwiseMaskXorCaller<unsigned char>, bitwiseMaskXorCaller<unsigned char>,\r
bitwiseMaskXorCaller<unsigned short>, bitwiseMaskXorCaller<unsigned short>,\r
bitwiseMaskXorCaller<unsigned int>\r
};\r
\r
+ CV_Assert(src1.depth() <= CV_64F);\r
CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
+ CV_Assert(mask.type() == CV_8U && mask.size() == src1.size());\r
+\r
dst.create(src1.size(), src1.type());\r
\r
- Caller caller = callers[src1.depth()];\r
- CV_Assert(caller);\r
+ const func_t func = funcs[src1.depth()];\r
\r
int cn = dst.depth() != CV_64F ? dst.channels() : dst.channels() * (sizeof(double) / sizeof(unsigned int));\r
- caller(dst.rows, dst.cols, cn, src1, src2, mask, dst, stream);\r
+\r
+ func(dst.rows, dst.cols, cn, src1, src2, mask, dst, stream);\r
}\r
}\r
\r
void cv::gpu::bitwise_or(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream& stream)\r
{\r
typedef void (*func_t)(const GpuMat& src, Scalar sc, GpuMat& dst, cudaStream_t stream);\r
-\r
static const func_t funcs[5][4] =\r
{\r
- {NppBitwiseC<CV_8U, 1, nppiOrC_8u_C1R>::call, 0, NppBitwiseC<CV_8U, 3, nppiOrC_8u_C3R>::call, NppBitwiseC<CV_8U, 4, nppiOrC_8u_C4R>::call},\r
+ {NppBitwiseC<CV_8U , 1, nppiOrC_8u_C1R >::call, 0, NppBitwiseC<CV_8U , 3, nppiOrC_8u_C3R >::call, NppBitwiseC<CV_8U , 4, nppiOrC_8u_C4R >::call},\r
{0,0,0,0},\r
{NppBitwiseC<CV_16U, 1, nppiOrC_16u_C1R>::call, 0, NppBitwiseC<CV_16U, 3, nppiOrC_16u_C3R>::call, NppBitwiseC<CV_16U, 4, nppiOrC_16u_C4R>::call},\r
{0,0,0,0},\r
void cv::gpu::bitwise_and(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream& stream)\r
{\r
typedef void (*func_t)(const GpuMat& src, Scalar sc, GpuMat& dst, cudaStream_t stream);\r
-\r
static const func_t funcs[5][4] =\r
{\r
- {NppBitwiseC<CV_8U, 1, nppiAndC_8u_C1R>::call, 0, NppBitwiseC<CV_8U, 3, nppiAndC_8u_C3R>::call, NppBitwiseC<CV_8U, 4, nppiAndC_8u_C4R>::call},\r
+ {NppBitwiseC<CV_8U , 1, nppiAndC_8u_C1R >::call, 0, NppBitwiseC<CV_8U , 3, nppiAndC_8u_C3R >::call, NppBitwiseC<CV_8U , 4, nppiAndC_8u_C4R >::call},\r
{0,0,0,0},\r
{NppBitwiseC<CV_16U, 1, nppiAndC_16u_C1R>::call, 0, NppBitwiseC<CV_16U, 3, nppiAndC_16u_C3R>::call, NppBitwiseC<CV_16U, 4, nppiAndC_16u_C4R>::call},\r
{0,0,0,0},\r
void cv::gpu::bitwise_xor(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream& stream)\r
{\r
typedef void (*func_t)(const GpuMat& src, Scalar sc, GpuMat& dst, cudaStream_t stream);\r
-\r
static const func_t funcs[5][4] =\r
{\r
- {NppBitwiseC<CV_8U, 1, nppiXorC_8u_C1R>::call, 0, NppBitwiseC<CV_8U, 3, nppiXorC_8u_C3R>::call, NppBitwiseC<CV_8U, 4, nppiXorC_8u_C4R>::call},\r
+ {NppBitwiseC<CV_8U , 1, nppiXorC_8u_C1R >::call, 0, NppBitwiseC<CV_8U , 3, nppiXorC_8u_C3R >::call, NppBitwiseC<CV_8U , 4, nppiXorC_8u_C4R >::call},\r
{0,0,0,0},\r
{NppBitwiseC<CV_16U, 1, nppiXorC_16u_C1R>::call, 0, NppBitwiseC<CV_16U, 3, nppiXorC_16u_C3R>::call, NppBitwiseC<CV_16U, 4, nppiXorC_16u_C4R>::call},\r
{0,0,0,0},\r
\r
namespace cv { namespace gpu { namespace device\r
{\r
- template <typename T>\r
- void min_gpu(const DevMem2D_<T>& src1, const DevMem2D_<T>& src2, const DevMem2D_<T>& dst, cudaStream_t stream);\r
-\r
- template <typename T>\r
- void max_gpu(const DevMem2D_<T>& src1, const DevMem2D_<T>& src2, const DevMem2D_<T>& dst, cudaStream_t stream);\r
+ template <typename T> void min_gpu(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ template <typename T> void max_gpu(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
\r
- template <typename T>\r
- void min_gpu(const DevMem2D_<T>& src1, T src2, const DevMem2D_<T>& dst, cudaStream_t stream);\r
-\r
- template <typename T>\r
- void max_gpu(const DevMem2D_<T>& src1, T src2, const DevMem2D_<T>& dst, cudaStream_t stream);\r
+ template <typename T> void min_gpu(const DevMem2Db src, T val, DevMem2Db dst, cudaStream_t stream);\r
+ template <typename T> void max_gpu(const DevMem2Db src, T val, DevMem2Db dst, cudaStream_t stream);\r
}}}\r
\r
-namespace\r
+void cv::gpu::min(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)\r
{\r
- template <typename T>\r
- void min_caller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream)\r
- {\r
- CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
- dst.create(src1.size(), src1.type());\r
- ::cv::gpu::device::min_gpu<T>(src1.reshape(1), src2.reshape(1), dst.reshape(1), stream);\r
- }\r
+ using namespace cv::gpu::device;\r
\r
- template <typename T>\r
- void min_caller(const GpuMat& src1, double src2, GpuMat& dst, cudaStream_t stream)\r
+ typedef void (*func_t)(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
+ static const func_t funcs[] =\r
{\r
- dst.create(src1.size(), src1.type());\r
- ::cv::gpu::device::min_gpu<T>(src1.reshape(1), saturate_cast<T>(src2), dst.reshape(1), stream);\r
- }\r
+ min_gpu<unsigned char>,\r
+ min_gpu<signed char>,\r
+ min_gpu<unsigned short>,\r
+ min_gpu<short>,\r
+ min_gpu<int>,\r
+ min_gpu<float>,\r
+ min_gpu<double>\r
+ };\r
\r
- template <typename T>\r
- void max_caller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream)\r
- {\r
- CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
- dst.create(src1.size(), src1.type());\r
- ::cv::gpu::device::max_gpu<T>(src1.reshape(1), src2.reshape(1), dst.reshape(1), stream);\r
- }\r
+ CV_Assert(src1.depth() <= CV_64F);\r
+ CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
\r
- template <typename T>\r
- void max_caller(const GpuMat& src1, double src2, GpuMat& dst, cudaStream_t stream)\r
+ if (src1.depth() == CV_64F)\r
{\r
- dst.create(src1.size(), src1.type());\r
- ::cv::gpu::device::max_gpu<T>(src1.reshape(1), saturate_cast<T>(src2), dst.reshape(1), stream);\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
}\r
+\r
+ dst.create(src1.size(), src1.type());\r
+\r
+ funcs[src1.depth()](src1.reshape(1), src2.reshape(1), dst.reshape(1), StreamAccessor::getStream(stream));\r
}\r
\r
-void cv::gpu::min(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)\r
+void cv::gpu::max(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)\r
{\r
- CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
- CV_Assert((src1.depth() != CV_64F) ||\r
- (TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE)));\r
+ using namespace cv::gpu::device;\r
\r
- typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream);\r
+ typedef void (*func_t)(const DevMem2Db src1, const DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);\r
static const func_t funcs[] =\r
{\r
- min_caller<unsigned char>, min_caller<signed char>, min_caller<unsigned short>, min_caller<short>, min_caller<int>,\r
- min_caller<float>, min_caller<double>\r
+ max_gpu<unsigned char>,\r
+ max_gpu<signed char>,\r
+ max_gpu<unsigned short>,\r
+ max_gpu<short>,\r
+ max_gpu<int>,\r
+ max_gpu<float>,\r
+ max_gpu<double>\r
};\r
- funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));\r
+\r
+ CV_Assert(src1.depth() <= CV_64F);\r
+ CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
+\r
+ if (src1.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
+ dst.create(src1.size(), src1.type());\r
+\r
+ funcs[src1.depth()](src1.reshape(1), src2.reshape(1), dst.reshape(1), StreamAccessor::getStream(stream));\r
}\r
-void cv::gpu::min(const GpuMat& src1, double src2, GpuMat& dst, Stream& stream)\r
+\r
+namespace\r
{\r
- CV_Assert((src1.depth() != CV_64F) ||\r
- (TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE)));\r
+ template <typename T> void minScalar(const DevMem2Db src, double val, DevMem2Db dst, cudaStream_t stream)\r
+ {\r
+ cv::gpu::device::min_gpu(src, saturate_cast<T>(val), dst, stream);\r
+ }\r
\r
- typedef void (*func_t)(const GpuMat& src1, double src2, GpuMat& dst, cudaStream_t stream);\r
- static const func_t funcs[] =\r
+ template <typename T> void maxScalar(const DevMem2Db src, double val, DevMem2Db dst, cudaStream_t stream)\r
{\r
- min_caller<unsigned char>, min_caller<signed char>, min_caller<unsigned short>, min_caller<short>, min_caller<int>,\r
- min_caller<float>, min_caller<double>\r
- };\r
- funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));\r
+ cv::gpu::device::max_gpu(src, saturate_cast<T>(val), dst, stream);\r
+ }\r
}\r
\r
-void cv::gpu::max(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)\r
+void cv::gpu::min(const GpuMat& src, double val, GpuMat& dst, Stream& stream)\r
{\r
- CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
- CV_Assert((src1.depth() != CV_64F) ||\r
- (TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE)));\r
-\r
- typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream);\r
+ typedef void (*func_t)(const DevMem2Db src1, double src2, DevMem2Db dst, cudaStream_t stream);\r
static const func_t funcs[] =\r
{\r
- max_caller<unsigned char>, max_caller<signed char>, max_caller<unsigned short>, max_caller<short>, max_caller<int>,\r
- max_caller<float>, max_caller<double>\r
+ minScalar<unsigned char>,\r
+ minScalar<signed char>,\r
+ minScalar<unsigned short>,\r
+ minScalar<short>,\r
+ minScalar<int>,\r
+ minScalar<float>,\r
+ minScalar<double>\r
};\r
- funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));\r
+\r
+ CV_Assert(src.depth() <= CV_64F);\r
+ CV_Assert(src.channels() == 1);\r
+\r
+ if (src.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
+ dst.create(src.size(), src.type());\r
+\r
+ funcs[src.depth()](src, val, dst, StreamAccessor::getStream(stream));\r
}\r
\r
-void cv::gpu::max(const GpuMat& src1, double src2, GpuMat& dst, Stream& stream)\r
+void cv::gpu::max(const GpuMat& src, double val, GpuMat& dst, Stream& stream)\r
{\r
- CV_Assert((src1.depth() != CV_64F) ||\r
- (TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE)));\r
-\r
- typedef void (*func_t)(const GpuMat& src1, double src2, GpuMat& dst, cudaStream_t stream);\r
+ typedef void (*func_t)(const DevMem2Db src1, double src2, DevMem2Db dst, cudaStream_t stream);\r
static const func_t funcs[] =\r
{\r
- max_caller<unsigned char>, max_caller<signed char>, max_caller<unsigned short>, max_caller<short>, max_caller<int>,\r
- max_caller<float>, max_caller<double>\r
+ maxScalar<unsigned char>,\r
+ maxScalar<signed char>,\r
+ maxScalar<unsigned short>,\r
+ maxScalar<short>,\r
+ maxScalar<int>,\r
+ maxScalar<float>,\r
+ maxScalar<double>\r
};\r
- funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));\r
+\r
+ CV_Assert(src.depth() <= CV_64F);\r
+ CV_Assert(src.channels() == 1);\r
+\r
+ if (src.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
+ dst.create(src.size(), src.type());\r
+\r
+ funcs[src.depth()](src, val, dst, StreamAccessor::getStream(stream));\r
}\r
\r
////////////////////////////////////////////////////////////////////////\r
CV_Assert(src.channels() == 1 && src.depth() <= CV_64F);\r
CV_Assert(type <= THRESH_TOZERO_INV);\r
\r
+ if (src.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
dst.create(src.size(), src.type());\r
\r
cudaStream_t stream = StreamAccessor::getStream(s);\r
}\r
else\r
{\r
- typedef void (*caller_t)(const GpuMat& src, GpuMat& dst, double thresh, double maxVal, int type, cudaStream_t stream);\r
-\r
- static const caller_t callers[] =\r
+ typedef void (*func_t)(const GpuMat& src, GpuMat& dst, double thresh, double maxVal, int type, cudaStream_t stream);\r
+ static const func_t funcs[] =\r
{\r
threshold_caller<unsigned char>, threshold_caller<signed char>,\r
threshold_caller<unsigned short>, threshold_caller<short>,\r
maxVal = cvRound(maxVal);\r
}\r
\r
- callers[src.depth()](src, dst, thresh, maxVal, type, stream);\r
+ funcs[src.depth()](src, dst, thresh, maxVal, type, stream);\r
}\r
\r
return thresh;\r
\r
namespace cv { namespace gpu { namespace device\r
{\r
- template<typename T>\r
- void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream);\r
+ template<typename T> void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream);\r
}}}\r
\r
void cv::gpu::pow(const GpuMat& src, double power, GpuMat& dst, Stream& stream)\r
using namespace cv::gpu::device;\r
\r
typedef void (*func_t)(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream);\r
-\r
static const func_t funcs[] =\r
{\r
pow_caller<unsigned char>, pow_caller<signed char>,\r
pow_caller<int>, pow_caller<float>, pow_caller<double>\r
};\r
\r
+ CV_Assert(src.depth() <= CV_64F);\r
+\r
+ if (src.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
dst.create(src.size(), src.type());\r
\r
funcs[src.depth()](src.reshape(1), power, dst.reshape(1), StreamAccessor::getStream(stream));\r
NppAlphaComp<CV_16U, nppiAlphaComp_16u_AC4R>::call,\r
0,\r
NppAlphaComp<CV_32S, nppiAlphaComp_32s_AC4R>::call,\r
- NppAlphaComp<CV_32F, nppiAlphaComp_32f_AC4R>::call,\r
- 0\r
+ NppAlphaComp<CV_32F, nppiAlphaComp_32f_AC4R>::call\r
};\r
\r
CV_Assert(img1.type() == CV_8UC4 || img1.type() == CV_16UC4 || img1.type() == CV_32SC4 || img1.type() == CV_32FC4);\r
dst.create(img1.size(), img1.type());\r
\r
const func_t func = funcs[img1.depth()];\r
- CV_Assert(func != 0);\r
\r
func(img1, img2, dst, npp_alpha_ops[alpha_op], StreamAccessor::getStream(stream));\r
}\r
\r
dtype = dtype >= 0 ? CV_MAKETYPE(dtype, src1.channels()) : src1.type();\r
\r
+ CV_Assert(src1.depth() <= CV_64F && src2.depth() <= CV_64F && CV_MAT_DEPTH(dtype) <= CV_64F);\r
+\r
+ if (src1.depth() == CV_64F || src2.depth() == CV_64F || CV_MAT_DEPTH(dtype) == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
dst.create(src1.size(), dtype);\r
\r
const GpuMat* psrc1 = &src1;\r
}\r
\r
const func_t func = funcs[psrc1->depth()][psrc2->depth()][dst.depth()];\r
- CV_Assert(func != 0);\r
+\r
+ if (!func)\r
+ CV_Error(CV_StsUnsupportedFormat, "Unsupported combination of source and destination types");\r
\r
func(psrc1->reshape(1), alpha, psrc2->reshape(1), beta, gamma, dst.reshape(1), StreamAccessor::getStream(stream));\r
}\r
nppSafeCall( nppiMean_StdDev_8u_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step), sz, buf.ptr<Npp8u>(), dbuf, (double*)dbuf + 1) );\r
\r
cudaSafeCall( cudaDeviceSynchronize() );\r
- \r
+\r
double* ptrs[2] = {mean.val, stddev.val};\r
dbuf.download(ptrs);\r
}\r
\r
double cv::gpu::norm(const GpuMat& src, int normType, GpuMat& buf)\r
{\r
+ CV_Assert(normType == NORM_INF || normType == NORM_L1 || normType == NORM_L2);\r
+\r
GpuMat src_single_channel = src.reshape(1);\r
\r
if (normType == NORM_L1)\r
if (normType == NORM_L2)\r
return std::sqrt(sqrSum(src_single_channel, buf)[0]);\r
\r
- if (normType == NORM_INF)\r
- {\r
- double min_val, max_val;\r
- minMax(src_single_channel, &min_val, &max_val, GpuMat(), buf);\r
- return std::max(std::abs(min_val), std::abs(max_val));\r
- }\r
-\r
- CV_Error(CV_StsBadArg, "norm: unsupported norm type");\r
- return 0;\r
+ // NORM_INF\r
+ double min_val, max_val;\r
+ minMax(src_single_channel, &min_val, &max_val, GpuMat(), buf);\r
+ return std::max(std::abs(min_val), std::abs(max_val));\r
}\r
\r
double cv::gpu::norm(const GpuMat& src1, const GpuMat& src2, int normType)\r
{\r
- CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type());\r
-\r
CV_Assert(src1.type() == CV_8UC1);\r
+ CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
CV_Assert(normType == NORM_INF || normType == NORM_L1 || normType == NORM_L2);\r
\r
typedef NppStatus (*npp_norm_diff_func_t)(const Npp8u* pSrc1, int nSrcStep1, const Npp8u* pSrc2, int nSrcStep2,\r
sz.height = src1.rows;\r
\r
int funcIdx = normType >> 1;\r
- \r
+\r
double retVal;\r
\r
DeviceBuffer dbuf;\r
nppSafeCall( npp_norm_diff_func[funcIdx](src1.ptr<Npp8u>(), static_cast<int>(src1.step), src2.ptr<Npp8u>(), static_cast<int>(src2.step), sz, dbuf) );\r
\r
cudaSafeCall( cudaDeviceSynchronize() );\r
- \r
+\r
dbuf.download(&retVal);\r
\r
return retVal;\r
////////////////////////////////////////////////////////////////////////\r
// Sum\r
\r
-namespace cv { namespace gpu { namespace device \r
+namespace cv { namespace gpu { namespace device\r
{\r
- namespace matrix_reductions \r
+ namespace matrix_reductions\r
{\r
namespace sum\r
{\r
}\r
}}}\r
\r
-Scalar cv::gpu::sum(const GpuMat& src) \r
+Scalar cv::gpu::sum(const GpuMat& src)\r
{\r
GpuMat buf;\r
return sum(src, buf);\r
}\r
\r
\r
-Scalar cv::gpu::sum(const GpuMat& src, GpuMat& buf) \r
+Scalar cv::gpu::sum(const GpuMat& src, GpuMat& buf)\r
{\r
- using namespace ::cv::gpu::device::matrix_reductions::sum;\r
+ using namespace cv::gpu::device::matrix_reductions::sum;\r
\r
typedef void (*Caller)(const DevMem2Db, PtrStepb, double*, int);\r
\r
- static Caller multipass_callers[7] = \r
- { \r
- sumMultipassCaller<unsigned char>, sumMultipassCaller<char>, \r
- sumMultipassCaller<unsigned short>, sumMultipassCaller<short>, \r
- sumMultipassCaller<int>, sumMultipassCaller<float>, 0 \r
+ static Caller multipass_callers[] =\r
+ {\r
+ sumMultipassCaller<unsigned char>, sumMultipassCaller<char>,\r
+ sumMultipassCaller<unsigned short>, sumMultipassCaller<short>,\r
+ sumMultipassCaller<int>, sumMultipassCaller<float>\r
};\r
\r
- static Caller singlepass_callers[7] = { \r
- sumCaller<unsigned char>, sumCaller<char>, \r
- sumCaller<unsigned short>, sumCaller<short>, \r
- sumCaller<int>, sumCaller<float>, 0 \r
+ static Caller singlepass_callers[] = {\r
+ sumCaller<unsigned char>, sumCaller<char>,\r
+ sumCaller<unsigned short>, sumCaller<short>,\r
+ sumCaller<int>, sumCaller<float>\r
};\r
\r
+ CV_Assert(src.depth() <= CV_32F);\r
+\r
Size buf_size;\r
- getBufSizeRequired(src.cols, src.rows, src.channels(), buf_size.width, buf_size.height); \r
+ getBufSizeRequired(src.cols, src.rows, src.channels(), buf_size.width, buf_size.height);\r
ensureSizeIsEnough(buf_size, CV_8U, buf);\r
\r
Caller* callers = multipass_callers;\r
callers = singlepass_callers;\r
\r
Caller caller = callers[src.depth()];\r
- if (!caller) CV_Error(CV_StsBadArg, "sum: unsupported type");\r
\r
double result[4];\r
caller(src, buf, result, src.channels());\r
}\r
\r
\r
-Scalar cv::gpu::absSum(const GpuMat& src) \r
+Scalar cv::gpu::absSum(const GpuMat& src)\r
{\r
GpuMat buf;\r
return absSum(src, buf);\r
}\r
\r
\r
-Scalar cv::gpu::absSum(const GpuMat& src, GpuMat& buf) \r
+Scalar cv::gpu::absSum(const GpuMat& src, GpuMat& buf)\r
{\r
- using namespace ::cv::gpu::device::matrix_reductions::sum;\r
+ using namespace cv::gpu::device::matrix_reductions::sum;\r
\r
typedef void (*Caller)(const DevMem2Db, PtrStepb, double*, int);\r
\r
- static Caller multipass_callers[7] = \r
- { \r
- absSumMultipassCaller<unsigned char>, absSumMultipassCaller<char>, \r
- absSumMultipassCaller<unsigned short>, absSumMultipassCaller<short>, \r
- absSumMultipassCaller<int>, absSumMultipassCaller<float>, 0 \r
+ static Caller multipass_callers[] =\r
+ {\r
+ absSumMultipassCaller<unsigned char>, absSumMultipassCaller<char>,\r
+ absSumMultipassCaller<unsigned short>, absSumMultipassCaller<short>,\r
+ absSumMultipassCaller<int>, absSumMultipassCaller<float>\r
};\r
\r
- static Caller singlepass_callers[7] = \r
- { \r
- absSumCaller<unsigned char>, absSumCaller<char>, \r
- absSumCaller<unsigned short>, absSumCaller<short>, \r
- absSumCaller<int>, absSumCaller<float>, 0 \r
+ static Caller singlepass_callers[] =\r
+ {\r
+ absSumCaller<unsigned char>, absSumCaller<char>,\r
+ absSumCaller<unsigned short>, absSumCaller<short>,\r
+ absSumCaller<int>, absSumCaller<float>\r
};\r
\r
+ CV_Assert(src.depth() <= CV_32F);\r
+\r
Size buf_size;\r
- getBufSizeRequired(src.cols, src.rows, src.channels(), buf_size.width, buf_size.height); \r
+ getBufSizeRequired(src.cols, src.rows, src.channels(), buf_size.width, buf_size.height);\r
ensureSizeIsEnough(buf_size, CV_8U, buf);\r
\r
Caller* callers = multipass_callers;\r
callers = singlepass_callers;\r
\r
Caller caller = callers[src.depth()];\r
- if (!caller) CV_Error(CV_StsBadArg, "absSum: unsupported type");\r
\r
double result[4];\r
caller(src, buf, result, src.channels());\r
}\r
\r
\r
-Scalar cv::gpu::sqrSum(const GpuMat& src) \r
+Scalar cv::gpu::sqrSum(const GpuMat& src)\r
{\r
GpuMat buf;\r
return sqrSum(src, buf);\r
}\r
\r
\r
-Scalar cv::gpu::sqrSum(const GpuMat& src, GpuMat& buf) \r
+Scalar cv::gpu::sqrSum(const GpuMat& src, GpuMat& buf)\r
{\r
- using namespace ::cv::gpu::device::matrix_reductions::sum;\r
+ using namespace cv::gpu::device::matrix_reductions::sum;\r
\r
typedef void (*Caller)(const DevMem2Db, PtrStepb, double*, int);\r
\r
- static Caller multipass_callers[7] = \r
- { \r
- sqrSumMultipassCaller<unsigned char>, sqrSumMultipassCaller<char>, \r
- sqrSumMultipassCaller<unsigned short>, sqrSumMultipassCaller<short>, \r
- sqrSumMultipassCaller<int>, sqrSumMultipassCaller<float>, 0 \r
+ static Caller multipass_callers[] =\r
+ {\r
+ sqrSumMultipassCaller<unsigned char>, sqrSumMultipassCaller<char>,\r
+ sqrSumMultipassCaller<unsigned short>, sqrSumMultipassCaller<short>,\r
+ sqrSumMultipassCaller<int>, sqrSumMultipassCaller<float>\r
};\r
\r
- static Caller singlepass_callers[7] = \r
- { \r
- sqrSumCaller<unsigned char>, sqrSumCaller<char>, \r
- sqrSumCaller<unsigned short>, sqrSumCaller<short>, \r
- sqrSumCaller<int>, sqrSumCaller<float>, 0 \r
+ static Caller singlepass_callers[7] =\r
+ {\r
+ sqrSumCaller<unsigned char>, sqrSumCaller<char>,\r
+ sqrSumCaller<unsigned short>, sqrSumCaller<short>,\r
+ sqrSumCaller<int>, sqrSumCaller<float>\r
};\r
\r
+ CV_Assert(src.depth() <= CV_32F);\r
+\r
Caller* callers = multipass_callers;\r
if (TargetArchs::builtWith(GLOBAL_ATOMICS) && DeviceInfo().supports(GLOBAL_ATOMICS))\r
callers = singlepass_callers;\r
\r
Size buf_size;\r
- getBufSizeRequired(src.cols, src.rows, src.channels(), buf_size.width, buf_size.height); \r
+ getBufSizeRequired(src.cols, src.rows, src.channels(), buf_size.width, buf_size.height);\r
ensureSizeIsEnough(buf_size, CV_8U, buf);\r
\r
Caller caller = callers[src.depth()];\r
- if (!caller) CV_Error(CV_StsBadArg, "sqrSum: unsupported type");\r
\r
double result[4];\r
caller(src, buf, result, src.channels());\r
////////////////////////////////////////////////////////////////////////\r
// Find min or max\r
\r
-namespace cv { namespace gpu { namespace device \r
+namespace cv { namespace gpu { namespace device\r
{\r
- namespace matrix_reductions \r
+ namespace matrix_reductions\r
{\r
- namespace minmax \r
+ namespace minmax\r
{\r
void getBufSizeRequired(int cols, int rows, int elem_size, int& bufcols, int& bufrows);\r
- \r
- template <typename T> \r
+\r
+ template <typename T>\r
void minMaxCaller(const DevMem2Db src, double* minval, double* maxval, PtrStepb buf);\r
\r
- template <typename T> \r
+ template <typename T>\r
void minMaxMaskCaller(const DevMem2Db src, const PtrStepb mask, double* minval, double* maxval, PtrStepb buf);\r
\r
- template <typename T> \r
+ template <typename T>\r
void minMaxMultipassCaller(const DevMem2Db src, double* minval, double* maxval, PtrStepb buf);\r
\r
- template <typename T> \r
+ template <typename T>\r
void minMaxMaskMultipassCaller(const DevMem2Db src, const PtrStepb mask, double* minval, double* maxval, PtrStepb buf);\r
}\r
}\r
typedef void (*Caller)(const DevMem2Db, double*, double*, PtrStepb);\r
typedef void (*MaskedCaller)(const DevMem2Db, const PtrStepb, double*, double*, PtrStepb);\r
\r
- static Caller multipass_callers[7] = \r
- { \r
- minMaxMultipassCaller<unsigned char>, minMaxMultipassCaller<char>, \r
- minMaxMultipassCaller<unsigned short>, minMaxMultipassCaller<short>, \r
- minMaxMultipassCaller<int>, minMaxMultipassCaller<float>, 0 \r
+ static Caller multipass_callers[] =\r
+ {\r
+ minMaxMultipassCaller<unsigned char>, minMaxMultipassCaller<char>,\r
+ minMaxMultipassCaller<unsigned short>, minMaxMultipassCaller<short>,\r
+ minMaxMultipassCaller<int>, minMaxMultipassCaller<float>, 0\r
};\r
\r
- static Caller singlepass_callers[7] = \r
- { \r
- minMaxCaller<unsigned char>, minMaxCaller<char>, \r
- minMaxCaller<unsigned short>, minMaxCaller<short>, \r
- minMaxCaller<int>, minMaxCaller<float>, minMaxCaller<double> \r
+ static Caller singlepass_callers[] =\r
+ {\r
+ minMaxCaller<unsigned char>, minMaxCaller<char>,\r
+ minMaxCaller<unsigned short>, minMaxCaller<short>,\r
+ minMaxCaller<int>, minMaxCaller<float>, minMaxCaller<double>\r
};\r
\r
- static MaskedCaller masked_multipass_callers[7] = \r
- { \r
- minMaxMaskMultipassCaller<unsigned char>, minMaxMaskMultipassCaller<char>, \r
+ static MaskedCaller masked_multipass_callers[] =\r
+ {\r
+ minMaxMaskMultipassCaller<unsigned char>, minMaxMaskMultipassCaller<char>,\r
minMaxMaskMultipassCaller<unsigned short>, minMaxMaskMultipassCaller<short>,\r
minMaxMaskMultipassCaller<int>, minMaxMaskMultipassCaller<float>, 0\r
};\r
\r
- static MaskedCaller masked_singlepass_callers[7] =\r
- { \r
- minMaxMaskCaller<unsigned char>, minMaxMaskCaller<char>, \r
- minMaxMaskCaller<unsigned short>, minMaxMaskCaller<short>, \r
- minMaxMaskCaller<int>, minMaxMaskCaller<float>, minMaxMaskCaller<double> \r
+ static MaskedCaller masked_singlepass_callers[] =\r
+ {\r
+ minMaxMaskCaller<unsigned char>, minMaxMaskCaller<char>,\r
+ minMaxMaskCaller<unsigned short>, minMaxMaskCaller<short>,\r
+ minMaxMaskCaller<int>, minMaxMaskCaller<float>, minMaxMaskCaller<double>\r
};\r
\r
+ CV_Assert(src.depth() <= CV_64F);\r
CV_Assert(src.channels() == 1);\r
-\r
CV_Assert(mask.empty() || (mask.type() == CV_8U && src.size() == mask.size()));\r
\r
+ if (src.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
double minVal_; if (!minVal) minVal = &minVal_;\r
double maxVal_; if (!maxVal) maxVal = &maxVal_;\r
- \r
+\r
Size buf_size;\r
getBufSizeRequired(src.cols, src.rows, static_cast<int>(src.elemSize()), buf_size.width, buf_size.height);\r
ensureSizeIsEnough(buf_size, CV_8U, buf);\r
callers = singlepass_callers;\r
\r
Caller caller = callers[src.type()];\r
- if (!caller) CV_Error(CV_StsBadArg, "minMax: unsupported type");\r
+ CV_Assert(caller != 0);\r
caller(src, minVal, maxVal, buf);\r
}\r
else\r
callers = masked_singlepass_callers;\r
\r
MaskedCaller caller = callers[src.type()];\r
- if (!caller) CV_Error(CV_StsBadArg, "minMax: unsupported type");\r
+ CV_Assert(caller != 0);\r
caller(src, mask, minVal, maxVal, buf);\r
}\r
}\r
////////////////////////////////////////////////////////////////////////\r
// Locate min and max\r
\r
-namespace cv { namespace gpu { namespace device \r
+namespace cv { namespace gpu { namespace device\r
{\r
- namespace matrix_reductions \r
+ namespace matrix_reductions\r
{\r
- namespace minmaxloc \r
+ namespace minmaxloc\r
{\r
- void getBufSizeRequired(int cols, int rows, int elem_size, int& b1cols, \r
+ void getBufSizeRequired(int cols, int rows, int elem_size, int& b1cols,\r
int& b1rows, int& b2cols, int& b2rows);\r
\r
- template <typename T> \r
- void minMaxLocCaller(const DevMem2Db src, double* minval, double* maxval, \r
+ template <typename T>\r
+ void minMaxLocCaller(const DevMem2Db src, double* minval, double* maxval,\r
int minloc[2], int maxloc[2], PtrStepb valBuf, PtrStepb locBuf);\r
\r
- template <typename T> \r
- void minMaxLocMaskCaller(const DevMem2Db src, const PtrStepb mask, double* minval, double* maxval, \r
+ template <typename T>\r
+ void minMaxLocMaskCaller(const DevMem2Db src, const PtrStepb mask, double* minval, double* maxval,\r
int minloc[2], int maxloc[2], PtrStepb valBuf, PtrStepb locBuf);\r
\r
- template <typename T> \r
- void minMaxLocMultipassCaller(const DevMem2Db src, double* minval, double* maxval, \r
+ template <typename T>\r
+ void minMaxLocMultipassCaller(const DevMem2Db src, double* minval, double* maxval,\r
int minloc[2], int maxloc[2], PtrStepb valBuf, PtrStepb locBuf);\r
\r
- template <typename T> \r
- void minMaxLocMaskMultipassCaller(const DevMem2Db src, const PtrStepb mask, double* minval, double* maxval, \r
+ template <typename T>\r
+ void minMaxLocMaskMultipassCaller(const DevMem2Db src, const PtrStepb mask, double* minval, double* maxval,\r
int minloc[2], int maxloc[2], PtrStepb valBuf, PtrStepb locBuf);\r
}\r
}\r
}}}\r
\r
void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point* minLoc, Point* maxLoc, const GpuMat& mask)\r
-{ \r
+{\r
GpuMat valBuf, locBuf;\r
minMaxLoc(src, minVal, maxVal, minLoc, maxLoc, mask, valBuf, locBuf);\r
}\r
typedef void (*Caller)(const DevMem2Db, double*, double*, int[2], int[2], PtrStepb, PtrStepb);\r
typedef void (*MaskedCaller)(const DevMem2Db, const PtrStepb, double*, double*, int[2], int[2], PtrStepb, PtrStepb);\r
\r
- static Caller multipass_callers[7] = \r
+ static Caller multipass_callers[] =\r
{\r
- minMaxLocMultipassCaller<unsigned char>, minMaxLocMultipassCaller<char>, \r
- minMaxLocMultipassCaller<unsigned short>, minMaxLocMultipassCaller<short>, \r
- minMaxLocMultipassCaller<int>, minMaxLocMultipassCaller<float>, 0 \r
+ minMaxLocMultipassCaller<unsigned char>, minMaxLocMultipassCaller<char>,\r
+ minMaxLocMultipassCaller<unsigned short>, minMaxLocMultipassCaller<short>,\r
+ minMaxLocMultipassCaller<int>, minMaxLocMultipassCaller<float>, 0\r
};\r
\r
- static Caller singlepass_callers[7] = \r
+ static Caller singlepass_callers[] =\r
{\r
- minMaxLocCaller<unsigned char>, minMaxLocCaller<char>, \r
- minMaxLocCaller<unsigned short>, minMaxLocCaller<short>, \r
- minMaxLocCaller<int>, minMaxLocCaller<float>, minMaxLocCaller<double> \r
+ minMaxLocCaller<unsigned char>, minMaxLocCaller<char>,\r
+ minMaxLocCaller<unsigned short>, minMaxLocCaller<short>,\r
+ minMaxLocCaller<int>, minMaxLocCaller<float>, minMaxLocCaller<double>\r
};\r
\r
- static MaskedCaller masked_multipass_callers[7] = \r
+ static MaskedCaller masked_multipass_callers[] =\r
{\r
minMaxLocMaskMultipassCaller<unsigned char>, minMaxLocMaskMultipassCaller<char>,\r
- minMaxLocMaskMultipassCaller<unsigned short>, minMaxLocMaskMultipassCaller<short>, \r
- minMaxLocMaskMultipassCaller<int>, minMaxLocMaskMultipassCaller<float>, 0 \r
+ minMaxLocMaskMultipassCaller<unsigned short>, minMaxLocMaskMultipassCaller<short>,\r
+ minMaxLocMaskMultipassCaller<int>, minMaxLocMaskMultipassCaller<float>, 0\r
};\r
\r
- static MaskedCaller masked_singlepass_callers[7] = \r
- { \r
- minMaxLocMaskCaller<unsigned char>, minMaxLocMaskCaller<char>, \r
- minMaxLocMaskCaller<unsigned short>, minMaxLocMaskCaller<short>, \r
- minMaxLocMaskCaller<int>, minMaxLocMaskCaller<float>, minMaxLocMaskCaller<double> \r
+ static MaskedCaller masked_singlepass_callers[] =\r
+ {\r
+ minMaxLocMaskCaller<unsigned char>, minMaxLocMaskCaller<char>,\r
+ minMaxLocMaskCaller<unsigned short>, minMaxLocMaskCaller<short>,\r
+ minMaxLocMaskCaller<int>, minMaxLocMaskCaller<float>, minMaxLocMaskCaller<double>\r
};\r
\r
+ CV_Assert(src.depth() <= CV_64F);\r
CV_Assert(src.channels() == 1);\r
-\r
CV_Assert(mask.empty() || (mask.type() == CV_8U && src.size() == mask.size()));\r
\r
+ if (src.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
double minVal_; if (!minVal) minVal = &minVal_;\r
double maxVal_; if (!maxVal) maxVal = &maxVal_;\r
int minLoc_[2];\r
int maxLoc_[2];\r
\r
Size valbuf_size, locbuf_size;\r
- getBufSizeRequired(src.cols, src.rows, static_cast<int>(src.elemSize()), valbuf_size.width, \r
+ getBufSizeRequired(src.cols, src.rows, static_cast<int>(src.elemSize()), valbuf_size.width,\r
valbuf_size.height, locbuf_size.width, locbuf_size.height);\r
ensureSizeIsEnough(valbuf_size, CV_8U, valBuf);\r
ensureSizeIsEnough(locbuf_size, CV_8U, locBuf);\r
callers = singlepass_callers;\r
\r
Caller caller = callers[src.type()];\r
- if (!caller) CV_Error(CV_StsBadArg, "minMaxLoc: unsupported type");\r
+ CV_Assert(caller != 0);\r
caller(src, minVal, maxVal, minLoc_, maxLoc_, valBuf, locBuf);\r
}\r
else\r
callers = masked_singlepass_callers;\r
\r
MaskedCaller caller = callers[src.type()];\r
- if (!caller) CV_Error(CV_StsBadArg, "minMaxLoc: unsupported type");\r
+ CV_Assert(caller != 0);\r
caller(src, mask, minVal, maxVal, minLoc_, maxLoc_, valBuf, locBuf);\r
}\r
\r
//////////////////////////////////////////////////////////////////////////////\r
// Count non-zero elements\r
\r
-namespace cv { namespace gpu { namespace device \r
+namespace cv { namespace gpu { namespace device\r
{\r
- namespace matrix_reductions \r
+ namespace matrix_reductions\r
{\r
- namespace countnonzero \r
+ namespace countnonzero\r
{\r
void getBufSizeRequired(int cols, int rows, int& bufcols, int& bufrows);\r
\r
- template <typename T> \r
+ template <typename T>\r
int countNonZeroCaller(const DevMem2Db src, PtrStepb buf);\r
\r
- template <typename T> \r
+ template <typename T>\r
int countNonZeroMultipassCaller(const DevMem2Db src, PtrStepb buf);\r
}\r
}\r
\r
typedef int (*Caller)(const DevMem2Db src, PtrStepb buf);\r
\r
- static Caller multipass_callers[7] = \r
+ static Caller multipass_callers[7] =\r
{\r
countNonZeroMultipassCaller<unsigned char>, countNonZeroMultipassCaller<char>,\r
countNonZeroMultipassCaller<unsigned short>, countNonZeroMultipassCaller<short>,\r
- countNonZeroMultipassCaller<int>, countNonZeroMultipassCaller<float>, 0 \r
+ countNonZeroMultipassCaller<int>, countNonZeroMultipassCaller<float>, 0\r
};\r
\r
- static Caller singlepass_callers[7] = \r
+ static Caller singlepass_callers[7] =\r
{\r
countNonZeroCaller<unsigned char>, countNonZeroCaller<char>,\r
countNonZeroCaller<unsigned short>, countNonZeroCaller<short>,\r
countNonZeroCaller<int>, countNonZeroCaller<float>, countNonZeroCaller<double> };\r
\r
+ CV_Assert(src.depth() <= CV_64F);\r
CV_Assert(src.channels() == 1);\r
\r
+ if (src.depth() == CV_64F)\r
+ {\r
+ if (!TargetArchs::builtWith(NATIVE_DOUBLE) || !DeviceInfo().supports(NATIVE_DOUBLE))\r
+ CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");\r
+ }\r
+\r
Size buf_size;\r
getBufSizeRequired(src.cols, src.rows, buf_size.width, buf_size.height);\r
ensureSizeIsEnough(buf_size, CV_8U, buf);\r
callers = singlepass_callers;\r
\r
Caller caller = callers[src.type()];\r
- if (!caller) CV_Error(CV_StsBadArg, "countNonZero: unsupported type");\r
+ CV_Assert(caller != 0);\r
return caller(src, buf);\r
}\r
\r
//////////////////////////////////////////////////////////////////////////////\r
// reduce\r
\r
-namespace cv { namespace gpu { namespace device \r
+namespace cv { namespace gpu { namespace device\r
{\r
- namespace matrix_reductions \r
+ namespace matrix_reductions\r
{\r
template <typename T, typename S, typename D> void reduceRows_gpu(const DevMem2Db& src, const DevMem2Db& dst, int reduceOp, cudaStream_t stream);\r
template <typename T, typename S, typename D> void reduceCols_gpu(const DevMem2Db& src, int cn, const DevMem2Db& dst, int reduceOp, cudaStream_t stream);\r
{\r
typedef void (*caller_t)(const DevMem2Db& src, const DevMem2Db& dst, int reduceOp, cudaStream_t stream);\r
\r
- static const caller_t callers[6][6] = \r
+ static const caller_t callers[6][6] =\r
{\r
{\r
reduceRows_gpu<unsigned char, int, unsigned char>,\r
};\r
\r
const caller_t func = callers[src.depth()][dst.depth()];\r
+\r
if (!func)\r
CV_Error(CV_StsUnsupportedFormat, "Unsupported combination of input and output array formats");\r
\r
{\r
typedef void (*caller_t)(const DevMem2Db& src, int cn, const DevMem2Db& dst, int reduceOp, cudaStream_t stream);\r
\r
- static const caller_t callers[6][6] = \r
+ static const caller_t callers[6][6] =\r
{\r
{\r
reduceCols_gpu<unsigned char, int, unsigned char>,\r
};\r
\r
const caller_t func = callers[src.depth()][dst.depth()];\r
+\r
if (!func)\r
CV_Error(CV_StsUnsupportedFormat, "Unsupported combination of input and output array formats");\r
\r
- func(src, src.channels(), dst, reduceOp, StreamAccessor::getStream(stream)); \r
+ func(src, src.channels(), dst, reduceOp, StreamAccessor::getStream(stream));\r
}\r
}\r
\r
#include "cuda.h"\r
#include "cuda_runtime_api.h"\r
#include "npp.h"\r
- \r
+\r
#ifdef HAVE_CUFFT\r
#include "cufft.h"\r
#endif\r
\r
#include "internal_shared.hpp"\r
#include "opencv2/gpu/stream_accessor.hpp"\r
- \r
+\r
#include "nvidia/core/NCV.hpp"\r
#include "nvidia/NPP_staging/NPP_staging.hpp"\r
#include "nvidia/NCVHaarObjectDetection.hpp"\r
#error "OpenCV GPU module doesn't support NVIDIA compute capability 1.0"\r
#endif\r
\r
- static inline void throw_nogpu() { CV_Error(CV_GpuNotSupported, "The called functionality is disabled for current build or platform"); }\r
+ static inline void throw_nogpu() { CV_Error(CV_StsNotImplemented, "The called functionality is disabled for current build or platform"); }\r
\r
#else /* defined(HAVE_CUDA) */\r
\r
cv::Mat src1 = randomMat(size, depth);\r
cv::Mat src2 = randomMat(size, depth);\r
\r
- cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
- cv::gpu::absdiff(loadMat(src1, useRoi), loadMat(src2, useRoi), dst);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::absdiff(loadMat(src1), loadMat(src2), dst);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
+ cv::gpu::absdiff(loadMat(src1, useRoi), loadMat(src2, useRoi), dst);\r
\r
- cv::Mat dst_gold;\r
- cv::absdiff(src1, src2, dst_gold);\r
+ cv::Mat dst_gold;\r
+ cv::absdiff(src1, src2, dst_gold);\r
\r
- EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ }\r
}\r
\r
TEST_P(AbsDiff, Scalar)\r
cv::Mat src = randomMat(size, depth);\r
cv::Scalar val = randomScalar(0.0, 255.0);\r
\r
- cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
- cv::gpu::absdiff(loadMat(src, useRoi), val, dst);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::absdiff(loadMat(src), val, dst);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
+ cv::gpu::absdiff(loadMat(src, useRoi), val, dst);\r
\r
- cv::Mat dst_gold;\r
- cv::absdiff(src, val, dst_gold);\r
+ cv::Mat dst_gold;\r
+ cv::absdiff(src, val, dst_gold);\r
\r
- EXPECT_MAT_NEAR(dst_gold, dst, depth <= CV_32F ? 1.0 : 1e-5);\r
+ EXPECT_MAT_NEAR(dst_gold, dst, depth <= CV_32F ? 1.0 : 1e-5);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, AbsDiff, testing::Combine(\r
////////////////////////////////////////////////////////////////////////////////\r
// Exp\r
\r
+template <typename T> void expImpl(const cv::Mat& src, cv::Mat& dst)\r
+{\r
+ dst.create(src.size(), src.type());\r
+\r
+ for (int y = 0; y < src.rows; ++y)\r
+ {\r
+ for (int x = 0; x < src.cols; ++x)\r
+ dst.at<T>(y, x) = cv::saturate_cast<T>(static_cast<int>(std::exp(static_cast<float>(src.at<T>(y, x)))));\r
+ }\r
+}\r
+void expImpl_float(const cv::Mat& src, cv::Mat& dst)\r
+{\r
+ dst.create(src.size(), src.type());\r
+\r
+ for (int y = 0; y < src.rows; ++y)\r
+ {\r
+ for (int x = 0; x < src.cols; ++x)\r
+ dst.at<float>(y, x) = std::exp(static_cast<float>(src.at<float>(y, x)));\r
+ }\r
+}\r
+\r
+void expGold(const cv::Mat& src, cv::Mat& dst)\r
+{\r
+ typedef void (*func_t)(const cv::Mat& src, cv::Mat& dst);\r
+\r
+ const func_t funcs[] =\r
+ {\r
+ expImpl<uchar>, expImpl<schar>, expImpl<ushort>, expImpl<short>,\r
+ expImpl<int>, expImpl_float\r
+ };\r
+\r
+ funcs[src.depth()](src, dst);\r
+}\r
+\r
PARAM_TEST_CASE(Exp, cv::gpu::DeviceInfo, cv::Size, MatType, UseRoi)\r
{\r
cv::gpu::DeviceInfo devInfo;\r
cv::gpu::exp(loadMat(src, useRoi), dst);\r
\r
cv::Mat dst_gold;\r
- cv::exp(src, dst_gold);\r
+ expGold(src, dst_gold);\r
\r
EXPECT_MAT_NEAR(dst_gold, dst, 1e-2);\r
}\r
INSTANTIATE_TEST_CASE_P(GPU_Core, Exp, testing::Combine(\r
ALL_DEVICES,\r
DIFFERENT_SIZES,\r
- testing::Values(MatType(CV_32FC1)),\r
+ testing::Values(MatType(CV_8UC1),\r
+ MatType(CV_16UC1),\r
+ MatType(CV_16SC1),\r
+ MatType(CV_32FC1)),\r
WHOLE_SUBMAT));\r
\r
////////////////////////////////////////////////////////////////////////////////\r
cv::Mat src1 = randomMat(size, depth);\r
cv::Mat src2 = randomMat(size, depth);\r
\r
- cv::gpu::GpuMat dst = createMat(size, CV_8UC1, useRoi);\r
- cv::gpu::compare(loadMat(src1, useRoi), loadMat(src2, useRoi), dst, cmp_code);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::compare(loadMat(src1), loadMat(src2), dst, cmp_code);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, CV_8UC1, useRoi);\r
+ cv::gpu::compare(loadMat(src1, useRoi), loadMat(src2, useRoi), dst, cmp_code);\r
\r
- cv::Mat dst_gold;\r
- cv::compare(src1, src2, dst_gold, cmp_code);\r
+ cv::Mat dst_gold;\r
+ cv::compare(src1, src2, dst_gold, cmp_code);\r
\r
- EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, Compare, testing::Combine(\r
}\r
};\r
\r
-TEST_P(Min, Accuracy)\r
+TEST_P(Min, Array)\r
{\r
cv::Mat src1 = randomMat(size, depth);\r
cv::Mat src2 = randomMat(size, depth);\r
\r
- cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
- cv::gpu::min(loadMat(src1, useRoi), loadMat(src2, useRoi), dst);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::min(loadMat(src1), loadMat(src2), dst);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
+ cv::gpu::min(loadMat(src1, useRoi), loadMat(src2, useRoi), dst);\r
\r
- cv::Mat dst_gold = cv::min(src1, src2);\r
+ cv::Mat dst_gold = cv::min(src1, src2);\r
\r
- EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ }\r
+}\r
+\r
+TEST_P(Min, Scalar)\r
+{\r
+ cv::Mat src = randomMat(size, depth);\r
+ double val = randomDouble(0.0, 255.0);\r
+\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::min(loadMat(src), val, dst);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
+ cv::gpu::min(loadMat(src, useRoi), val, dst);\r
+\r
+ cv::Mat dst_gold = cv::min(src, val);\r
+\r
+ EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, Min, testing::Combine(\r
}\r
};\r
\r
-TEST_P(Max, Accuracy)\r
+TEST_P(Max, Array)\r
{\r
cv::Mat src1 = randomMat(size, depth);\r
cv::Mat src2 = randomMat(size, depth);\r
\r
- cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
- cv::gpu::max(loadMat(src1, useRoi), loadMat(src2, useRoi), dst);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::max(loadMat(src1), loadMat(src2), dst);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
+ cv::gpu::max(loadMat(src1, useRoi), loadMat(src2, useRoi), dst);\r
+\r
+ cv::Mat dst_gold = cv::max(src1, src2);\r
+\r
+ EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ }\r
+}\r
\r
- cv::Mat dst_gold = cv::max(src1, src2);\r
+TEST_P(Max, Scalar)\r
+{\r
+ cv::Mat src = randomMat(size, depth);\r
+ double val = randomDouble(0.0, 255.0);\r
\r
- EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::max(loadMat(src), val, dst);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
+ cv::gpu::max(loadMat(src, useRoi), val, dst);\r
+\r
+ cv::Mat dst_gold = cv::max(src, val);\r
+\r
+ EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, Max, testing::Combine(\r
if (src.depth() < CV_32F)\r
power = static_cast<int>(power);\r
\r
- cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
- cv::gpu::pow(loadMat(src, useRoi), power, dst);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::pow(loadMat(src), power, dst);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, depth, useRoi);\r
+ cv::gpu::pow(loadMat(src, useRoi), power, dst);\r
\r
- cv::Mat dst_gold;\r
- cv::pow(src, power, dst_gold);\r
+ cv::Mat dst_gold;\r
+ cv::pow(src, power, dst_gold);\r
\r
- EXPECT_MAT_NEAR(dst_gold, dst, depth < CV_32F ? 0.0 : 1e-1);\r
+ EXPECT_MAT_NEAR(dst_gold, dst, depth < CV_32F ? 0.0 : 1e-1);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, Pow, testing::Combine(\r
int dst_depth;\r
bool useRoi;\r
\r
-\r
virtual void SetUp()\r
{\r
devInfo = GET_PARAM(0);\r
double beta = randomDouble(-10.0, 10.0);\r
double gamma = randomDouble(-10.0, 10.0);\r
\r
- cv::gpu::GpuMat dst = createMat(size, dst_depth, useRoi);\r
- cv::gpu::addWeighted(loadMat(src1, useRoi), alpha, loadMat(src2, useRoi), beta, gamma, dst, dst_depth);\r
+ if ((depth1 == CV_64F || depth2 == CV_64F || dst_depth == CV_64F) && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::addWeighted(loadMat(src1), alpha, loadMat(src2), beta, gamma, dst, dst_depth);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, dst_depth, useRoi);\r
+ cv::gpu::addWeighted(loadMat(src1, useRoi), alpha, loadMat(src2, useRoi), beta, gamma, dst, dst_depth);\r
\r
- cv::Mat dst_gold;\r
- cv::addWeighted(src1, alpha, src2, beta, gamma, dst_gold, dst_depth);\r
+ cv::Mat dst_gold;\r
+ cv::addWeighted(src1, alpha, src2, beta, gamma, dst_gold, dst_depth);\r
\r
- EXPECT_MAT_NEAR(dst_gold, dst, dst_depth < CV_32F ? 1.0 : 1e-12);\r
+ EXPECT_MAT_NEAR(dst_gold, dst, dst_depth < CV_32F ? 1.0 : 1e-12);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, AddWeighted, testing::Combine(\r
double alpha = randomDouble(-10.0, 10.0);\r
double beta = randomDouble(-10.0, 10.0);\r
\r
- cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
- cv::gpu::gemm(loadMat(src1, useRoi), loadMat(src2, useRoi), alpha, loadMat(src3, useRoi), beta, dst, flags);\r
+#ifndef HAVE_CUBLAS\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::gemm(loadMat(src1), loadMat(src2), alpha, loadMat(src3), beta, dst, flags);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+ }\r
+#else\r
+ if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::gemm(loadMat(src1), loadMat(src2), alpha, loadMat(src3), beta, dst, flags);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else if (type == CV_64FC2 && flags != 0)\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::gemm(loadMat(src1), loadMat(src2), alpha, loadMat(src3), beta, dst, flags);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsNotImplemented, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
+ cv::gpu::gemm(loadMat(src1, useRoi), loadMat(src2, useRoi), alpha, loadMat(src3, useRoi), beta, dst, flags);\r
\r
- cv::Mat dst_gold;\r
- cv::gemm(src1, src2, alpha, src3, beta, dst_gold, flags);\r
+ cv::Mat dst_gold;\r
+ cv::gemm(src1, src2, alpha, src3, beta, dst_gold, flags);\r
\r
- EXPECT_MAT_NEAR(dst_gold, dst, CV_MAT_DEPTH(type) == CV_32F ? 1e-1 : 1e-10);\r
+ EXPECT_MAT_NEAR(dst_gold, dst, CV_MAT_DEPTH(type) == CV_32F ? 1e-1 : 1e-10);\r
+ }\r
+#endif\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, GEMM, testing::Combine(\r
{\r
cv::Mat src = randomMat(size, type);\r
\r
- cv::gpu::GpuMat dst = createMat(cv::Size(size.height, size.width), type, useRoi);\r
- cv::gpu::transpose(loadMat(src, useRoi), dst);\r
+ if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::transpose(loadMat(src), dst);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ cv::gpu::GpuMat dst = createMat(cv::Size(size.height, size.width), type, useRoi);\r
+ cv::gpu::transpose(loadMat(src, useRoi), dst);\r
\r
- cv::Mat dst_gold;\r
- cv::transpose(src, dst_gold);\r
+ cv::Mat dst_gold;\r
+ cv::transpose(src, dst_gold);\r
\r
- EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, Transpose, testing::Combine(\r
{\r
cv::Mat src = randomMat(size, depth);\r
\r
- double minVal, maxVal;\r
- cv::gpu::minMax(loadMat(src, useRoi), &minVal, &maxVal);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ double minVal, maxVal;\r
+ cv::gpu::minMax(loadMat(src), &minVal, &maxVal);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ double minVal, maxVal;\r
+ cv::gpu::minMax(loadMat(src, useRoi), &minVal, &maxVal);\r
\r
- double minVal_gold, maxVal_gold;\r
- minMaxLocGold(src, &minVal_gold, &maxVal_gold);\r
+ double minVal_gold, maxVal_gold;\r
+ minMaxLocGold(src, &minVal_gold, &maxVal_gold);\r
\r
- EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
- EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
+ EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
+ EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
+ }\r
}\r
\r
TEST_P(MinMax, WithMask)\r
cv::Mat src = randomMat(size, depth);\r
cv::Mat mask = randomMat(size, CV_8UC1, 0.0, 2.0);\r
\r
- double minVal, maxVal;\r
- cv::gpu::minMax(loadMat(src, useRoi), &minVal, &maxVal, loadMat(mask, useRoi));\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ double minVal, maxVal;\r
+ cv::gpu::minMax(loadMat(src), &minVal, &maxVal, loadMat(mask));\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ double minVal, maxVal;\r
+ cv::gpu::minMax(loadMat(src, useRoi), &minVal, &maxVal, loadMat(mask, useRoi));\r
\r
- double minVal_gold, maxVal_gold;\r
- minMaxLocGold(src, &minVal_gold, &maxVal_gold, 0, 0, mask);\r
+ double minVal_gold, maxVal_gold;\r
+ minMaxLocGold(src, &minVal_gold, &maxVal_gold, 0, 0, mask);\r
\r
- EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
- EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
+ EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
+ EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
+ }\r
}\r
\r
TEST_P(MinMax, NullPtr)\r
{\r
cv::Mat src = randomMat(size, depth);\r
\r
- cv::gpu::minMax(loadMat(src, useRoi), 0, 0);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ double minVal, maxVal;\r
+ cv::gpu::minMax(loadMat(src), &minVal, 0);\r
+ cv::gpu::minMax(loadMat(src), 0, &maxVal);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ double minVal, maxVal;\r
+ cv::gpu::minMax(loadMat(src, useRoi), &minVal, 0);\r
+ cv::gpu::minMax(loadMat(src, useRoi), 0, &maxVal);\r
+\r
+ double minVal_gold, maxVal_gold;\r
+ minMaxLocGold(src, &minVal_gold, &maxVal_gold, 0, 0);\r
+\r
+ EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
+ EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, MinMax, testing::Combine(\r
{\r
cv::Mat src = randomMat(size, depth);\r
\r
- double minVal, maxVal;\r
- cv::Point minLoc, maxLoc;\r
- cv::gpu::minMaxLoc(loadMat(src, useRoi), &minVal, &maxVal, &minLoc, &maxLoc);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ double minVal, maxVal;\r
+ cv::Point minLoc, maxLoc;\r
+ cv::gpu::minMaxLoc(loadMat(src), &minVal, &maxVal, &minLoc, &maxLoc);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ double minVal, maxVal;\r
+ cv::Point minLoc, maxLoc;\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), &minVal, &maxVal, &minLoc, &maxLoc);\r
\r
- double minVal_gold, maxVal_gold;\r
- cv::Point minLoc_gold, maxLoc_gold;\r
- minMaxLocGold(src, &minVal_gold, &maxVal_gold, &minLoc_gold, &maxLoc_gold);\r
+ double minVal_gold, maxVal_gold;\r
+ cv::Point minLoc_gold, maxLoc_gold;\r
+ minMaxLocGold(src, &minVal_gold, &maxVal_gold, &minLoc_gold, &maxLoc_gold);\r
\r
- EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
- EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
+ EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
+ EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
\r
- expectEqual(src, minLoc_gold, minLoc);\r
- expectEqual(src, maxLoc_gold, maxLoc);\r
+ expectEqual(src, minLoc_gold, minLoc);\r
+ expectEqual(src, maxLoc_gold, maxLoc);\r
+ }\r
}\r
\r
TEST_P(MinMaxLoc, WithMask)\r
cv::Mat src = randomMat(size, depth);\r
cv::Mat mask = randomMat(size, CV_8UC1, 0.0, 2.0);\r
\r
- double minVal, maxVal;\r
- cv::Point minLoc, maxLoc;\r
- cv::gpu::minMaxLoc(loadMat(src, useRoi), &minVal, &maxVal, &minLoc, &maxLoc, loadMat(mask, useRoi));\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ double minVal, maxVal;\r
+ cv::Point minLoc, maxLoc;\r
+ cv::gpu::minMaxLoc(loadMat(src), &minVal, &maxVal, &minLoc, &maxLoc, loadMat(mask));\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ double minVal, maxVal;\r
+ cv::Point minLoc, maxLoc;\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), &minVal, &maxVal, &minLoc, &maxLoc, loadMat(mask, useRoi));\r
\r
- double minVal_gold, maxVal_gold;\r
- cv::Point minLoc_gold, maxLoc_gold;\r
- minMaxLocGold(src, &minVal_gold, &maxVal_gold, &minLoc_gold, &maxLoc_gold, mask);\r
+ double minVal_gold, maxVal_gold;\r
+ cv::Point minLoc_gold, maxLoc_gold;\r
+ minMaxLocGold(src, &minVal_gold, &maxVal_gold, &minLoc_gold, &maxLoc_gold, mask);\r
\r
- EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
- EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
+ EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
+ EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
\r
- expectEqual(src, minLoc_gold, minLoc);\r
- expectEqual(src, maxLoc_gold, maxLoc);\r
+ expectEqual(src, minLoc_gold, minLoc);\r
+ expectEqual(src, maxLoc_gold, maxLoc);\r
+ }\r
}\r
\r
TEST_P(MinMaxLoc, NullPtr)\r
{\r
cv::Mat src = randomMat(size, depth);\r
\r
- cv::gpu::minMaxLoc(loadMat(src, useRoi), 0, 0, 0, 0);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ double minVal, maxVal;\r
+ cv::Point minLoc, maxLoc;\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), &minVal, 0, 0, 0);\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), 0, &maxVal, 0, 0);\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), 0, 0, &minLoc, 0);\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), 0, 0, 0, &maxLoc);\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ double minVal, maxVal;\r
+ cv::Point minLoc, maxLoc;\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), &minVal, 0, 0, 0);\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), 0, &maxVal, 0, 0);\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), 0, 0, &minLoc, 0);\r
+ cv::gpu::minMaxLoc(loadMat(src, useRoi), 0, 0, 0, &maxLoc);\r
+\r
+ double minVal_gold, maxVal_gold;\r
+ cv::Point minLoc_gold, maxLoc_gold;\r
+ minMaxLocGold(src, &minVal_gold, &maxVal_gold, &minLoc_gold, &maxLoc_gold);\r
+\r
+ EXPECT_DOUBLE_EQ(minVal_gold, minVal);\r
+ EXPECT_DOUBLE_EQ(maxVal_gold, maxVal);\r
+\r
+ expectEqual(src, minLoc_gold, minLoc);\r
+ expectEqual(src, maxLoc_gold, maxLoc);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, MinMaxLoc, testing::Combine(\r
cv::Mat src;\r
srcBase.convertTo(src, depth);\r
\r
- int val = cv::gpu::countNonZero(loadMat(src, useRoi));\r
-\r
- int val_gold = cv::countNonZero(src);\r
+ if (depth == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+ {\r
+ try\r
+ {\r
+ cv::gpu::countNonZero(loadMat(src));\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ ASSERT_EQ(CV_StsUnsupportedFormat, e.code);\r
+ }\r
+ }\r
+ else\r
+ {\r
+ int val = cv::gpu::countNonZero(loadMat(src, useRoi));\r
\r
+ int val_gold = cv::countNonZero(src);\r
\r
- ASSERT_EQ(val_gold, val);\r
+ ASSERT_EQ(val_gold, val);\r
+ }\r
}\r
\r
INSTANTIATE_TEST_CASE_P(GPU_Core, CountNonZero, testing::Combine(\r
projects / profile / ivi / opencv.git / commitdiff