minor refactoring of GPU module and GPU tests, split arithm and imgproc parts.
CV_EXPORTS void transpose(const GpuMat& src1, GpuMat& dst);\r
\r
//! computes element-wise absolute difference of two arrays (c = abs(a - b))\r
- CV_EXPORTS void absdiff(const GpuMat& a, const GpuMat& b, GpuMat& c);\r
-\r
- //! applies fixed threshold to the image. \r
- //! Now supports only THRESH_TRUNC threshold type and one channels float source.\r
- CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh);\r
+ CV_EXPORTS void absdiff(const GpuMat& a, const GpuMat& b, GpuMat& c); \r
\r
//! compares elements of two arrays (c = a <cmpop> b)\r
//! Now doesn't support CMP_NE.\r
//! reverses the order of the rows, columns or both in a matrix\r
CV_EXPORTS void flip(const GpuMat& a, GpuMat& b, int flipCode);\r
\r
- //! resizes the image\r
- //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_LANCZOS4\r
- CV_EXPORTS void resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, double fy=0, int interpolation = INTER_LINEAR);\r
-\r
//! computes sum of array elements\r
CV_EXPORTS Scalar sum(const GpuMat& m);\r
\r
//! finds global minimum and maximum array elements and returns their values\r
CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal = 0);\r
\r
- //! copies 2D array to a larger destination array and pads borders with user-specifiable constant\r
- CV_EXPORTS void copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value = Scalar());\r
-\r
- //! warps the image using affine transformation\r
- //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC\r
- CV_EXPORTS void warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR);\r
-\r
- //! warps the image using perspective transformation\r
- //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC\r
- CV_EXPORTS void warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR);\r
-\r
- //! rotate 8bit single or four channel image\r
- //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC\r
- CV_EXPORTS void rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift = 0, double yShift = 0, int interpolation = INTER_LINEAR);\r
+ //! transforms 8-bit unsigned integers using lookup table: dst(i)=lut(src(i))\r
+ //! supports only single channels source\r
+ //! destination array will have the same type as source\r
+ //! lut must hase CV_32S depth and the same number of channels as in the source array\r
+ CV_EXPORTS void LUT(const GpuMat& src, const Mat& lut, GpuMat& dst);\r
\r
//! makes multi-channel array out of several single-channel arrays\r
CV_EXPORTS void merge(const GpuMat* src, size_t n, GpuMat& dst);\r
\r
////////////////////////////// Image processing //////////////////////////////\r
\r
- // DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation.\r
- // xymap.type() == xymap.type() == CV_32FC1\r
+ //! DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation.\r
+ //! xymap.type() == xymap.type() == CV_32FC1\r
CV_EXPORTS void remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap);\r
\r
- // Does mean shift filtering on GPU.\r
+ //! Does mean shift filtering on GPU.\r
CV_EXPORTS void meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, \r
TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1));\r
\r
- // Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV.\r
- // Supported types of input disparity: CV_8U, CV_16S.\r
- // Output disparity has CV_8UC4 type in BGRA format (alpha = 255).\r
+ //! Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV.\r
+ //! Supported types of input disparity: CV_8U, CV_16S.\r
+ //! Output disparity has CV_8UC4 type in BGRA format (alpha = 255).\r
CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndisp);\r
- // Acync version\r
+ //! Acync version\r
CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndisp, const Stream& stream);\r
\r
- // Reprojects disparity image to 3D space. \r
- // Supports CV_8U and CV_16S types of input disparity.\r
- // The output is a 4-channel floating-point (CV_32FC4) matrix. \r
- // Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map.\r
- // Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify.\r
+ //! Reprojects disparity image to 3D space. \r
+ //! Supports CV_8U and CV_16S types of input disparity.\r
+ //! The output is a 4-channel floating-point (CV_32FC4) matrix. \r
+ //! Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map.\r
+ //! Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify.\r
CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q);\r
- // Acync version\r
+ //! Acync version\r
CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream);\r
\r
+ //! converts image from one color space to another\r
CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn = 0);\r
+ //! Acync version\r
CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream);\r
\r
+ //! applies fixed threshold to the image. \r
+ //! Now supports only THRESH_TRUNC threshold type and one channels float source.\r
+ CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh);\r
+\r
+ //! resizes the image\r
+ //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_LANCZOS4\r
+ CV_EXPORTS void resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, double fy=0, int interpolation = INTER_LINEAR);\r
+ \r
+ //! warps the image using affine transformation\r
+ //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC\r
+ CV_EXPORTS void warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR);\r
+\r
+ //! warps the image using perspective transformation\r
+ //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC\r
+ CV_EXPORTS void warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR);\r
+ \r
+ //! rotate 8bit single or four channel image\r
+ //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC\r
+ CV_EXPORTS void rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift = 0, double yShift = 0, int interpolation = INTER_LINEAR);\r
+ \r
+ //! copies 2D array to a larger destination array and pads borders with user-specifiable constant\r
+ CV_EXPORTS void copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value = Scalar());\r
+ \r
+ //! computes the integral image and integral for the squared image\r
+ //! sum will have CV_32S type, sqsum - CV32F type\r
+ CV_EXPORTS void integral(GpuMat& src, GpuMat& sum, GpuMat& sqsum);\r
+\r
+ //! smooths the image using the normalized box filter\r
+ CV_EXPORTS void boxFilter(const GpuMat& src, GpuMat& dst, Size ksize, Point anchor = Point(-1,-1));\r
+ //! a synonym for normalized box filter\r
+ static inline void blur(const GpuMat& src, GpuMat& dst, Size ksize, Point anchor = Point(-1,-1))\r
+ {\r
+ boxFilter(src, dst, ksize, anchor);\r
+ }\r
\r
//! erodes the image (applies the local minimum operator)\r
CV_EXPORTS void erode( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations);\r
void cv::gpu::subtract(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
void cv::gpu::multiply(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
void cv::gpu::divide(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
-\r
void cv::gpu::transpose(const GpuMat&, GpuMat&) { throw_nogpu(); }\r
-\r
void cv::gpu::absdiff(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
-\r
-double cv::gpu::threshold(const GpuMat&, GpuMat&, double) { throw_nogpu(); return 0.0; }\r
-\r
void cv::gpu::compare(const GpuMat&, const GpuMat&, GpuMat&, int) { throw_nogpu(); }\r
-\r
void cv::gpu::meanStdDev(const GpuMat&, Scalar&, Scalar&) { throw_nogpu(); }\r
-\r
double cv::gpu::norm(const GpuMat&, int) { throw_nogpu(); return 0.0; }\r
double cv::gpu::norm(const GpuMat&, const GpuMat&, int) { throw_nogpu(); return 0.0; }\r
-\r
void cv::gpu::flip(const GpuMat&, GpuMat&, int) { throw_nogpu(); }\r
-\r
-void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int) { throw_nogpu(); }\r
-\r
Scalar cv::gpu::sum(const GpuMat&) { throw_nogpu(); return Scalar(); }\r
-\r
void cv::gpu::minMax(const GpuMat&, double*, double*) { throw_nogpu(); }\r
-\r
-void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, const Scalar&) { throw_nogpu(); }\r
-\r
-void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); }\r
-\r
-void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); }\r
-\r
-void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int) { throw_nogpu(); }\r
+void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst) { throw_nogpu(); }\r
\r
#else /* !defined (HAVE_CUDA) */\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// add subtract multiply divide\r
+\r
namespace\r
{\r
typedef NppStatus (*npp_arithm_8u_t)(const Npp8u* pSrc1, int nSrc1Step, const Npp8u* pSrc2, int nSrc2Step, Npp8u* pDst, int nDstStep, \r
nppFuncCaller(src2, src1, dst, nppiDiv_8u_C1RSfs, nppiDiv_8u_C4RSfs, nppiDiv_32f_C1R);\r
}\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// transpose\r
+\r
void cv::gpu::transpose(const GpuMat& src, GpuMat& dst)\r
{\r
CV_Assert(src.type() == CV_8UC1);\r
nppSafeCall( nppiTranspose_8u_C1R(src.ptr<Npp8u>(), src.step, dst.ptr<Npp8u>(), dst.step, sz) );\r
}\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// absdiff\r
+\r
void cv::gpu::absdiff(const GpuMat& src1, const GpuMat& src2, GpuMat& dst)\r
{\r
CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type());\r
}\r
}\r
\r
-double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh) \r
-{ \r
- CV_Assert(src.type() == CV_32FC1)\r
-\r
- dst.create( src.size(), src.type() );\r
-\r
- NppiSize sz;\r
- sz.width = src.cols;\r
- sz.height = src.rows;\r
-\r
- nppSafeCall( nppiThreshold_32f_C1R(src.ptr<Npp32f>(), src.step, \r
- dst.ptr<Npp32f>(), dst.step, sz, static_cast<Npp32f>(thresh), NPP_CMP_GREATER) );\r
-\r
- return thresh;\r
-}\r
+////////////////////////////////////////////////////////////////////////\r
+// compare\r
\r
namespace cv { namespace gpu { namespace matrix_operations\r
{\r
}\r
}\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// meanStdDev\r
+\r
void cv::gpu::meanStdDev(const GpuMat& src, Scalar& mean, Scalar& stddev) \r
{\r
CV_Assert(src.type() == CV_8UC1);\r
nppSafeCall( nppiMean_StdDev_8u_C1R(src.ptr<Npp8u>(), src.step, sz, mean.val, stddev.val) );\r
}\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// norm\r
+\r
double cv::gpu::norm(const GpuMat& src1, int normType) \r
{\r
return norm(src1, GpuMat(src1.size(), src1.type(), Scalar::all(0.0)), normType);\r
return retVal[0];\r
}\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// flip\r
+\r
void cv::gpu::flip(const GpuMat& src, GpuMat& dst, int flipCode)\r
{\r
CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
}\r
}\r
\r
-void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, double fy, int interpolation)\r
-{\r
- static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC, 0, NPPI_INTER_LANCZOS};\r
-\r
- CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
- CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC || interpolation == INTER_LANCZOS4);\r
-\r
- CV_Assert( src.size().area() > 0 );\r
- CV_Assert( !(dsize == Size()) || (fx > 0 && fy > 0) );\r
-\r
- if( dsize == Size() )\r
- {\r
- dsize = Size(saturate_cast<int>(src.cols * fx), saturate_cast<int>(src.rows * fy));\r
- }\r
- else\r
- {\r
- fx = (double)dsize.width / src.cols;\r
- fy = (double)dsize.height / src.rows;\r
- }\r
-\r
- dst.create(dsize, src.type());\r
-\r
- NppiSize srcsz;\r
- srcsz.width = src.cols;\r
- srcsz.height = src.rows;\r
- NppiRect srcrect;\r
- srcrect.x = srcrect.y = 0;\r
- srcrect.width = src.cols;\r
- srcrect.height = src.rows;\r
- NppiSize dstsz;\r
- dstsz.width = dst.cols;\r
- dstsz.height = dst.rows;\r
-\r
- if (src.type() == CV_8UC1)\r
- {\r
- nppSafeCall( nppiResize_8u_C1R(src.ptr<Npp8u>(), srcsz, src.step, srcrect,\r
- dst.ptr<Npp8u>(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) );\r
- }\r
- else\r
- {\r
- nppSafeCall( nppiResize_8u_C4R(src.ptr<Npp8u>(), srcsz, src.step, srcrect,\r
- dst.ptr<Npp8u>(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) );\r
- }\r
-}\r
+////////////////////////////////////////////////////////////////////////\r
+// sum\r
\r
Scalar cv::gpu::sum(const GpuMat& src)\r
{\r
return res;\r
}\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// minMax\r
+\r
void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal) \r
{\r
CV_Assert(src.type() == CV_8UC1);\r
*maxVal = max_res;\r
}\r
\r
-void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value) \r
-{\r
- CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1);\r
-\r
- dst.create(src.rows + top + bottom, src.cols + left + right, src.type());\r
-\r
- NppiSize srcsz;\r
- srcsz.width = src.cols;\r
- srcsz.height = src.rows;\r
- NppiSize dstsz;\r
- dstsz.width = dst.cols;\r
- dstsz.height = dst.rows;\r
+////////////////////////////////////////////////////////////////////////\r
+// LUT\r
\r
- switch (src.type())\r
- {\r
- case CV_8UC1:\r
- {\r
- Npp8u nVal = static_cast<Npp8u>(value[0]);\r
- nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr<Npp8u>(), src.step, srcsz, \r
- dst.ptr<Npp8u>(), dst.step, dstsz, top, left, nVal) );\r
- break;\r
- }\r
- case CV_8UC4:\r
- {\r
- Npp8u nVal[] = {static_cast<Npp8u>(value[0]), static_cast<Npp8u>(value[1]), static_cast<Npp8u>(value[2]), static_cast<Npp8u>(value[3])};\r
- nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr<Npp8u>(), src.step, srcsz, \r
- dst.ptr<Npp8u>(), dst.step, dstsz, top, left, nVal) );\r
- break;\r
- }\r
- case CV_32SC1:\r
- {\r
- Npp32s nVal = static_cast<Npp32s>(value[0]);\r
- nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr<Npp32s>(), src.step, srcsz, \r
- dst.ptr<Npp32s>(), dst.step, dstsz, top, left, nVal) );\r
- break;\r
- }\r
- default:\r
- CV_Assert(!"Unsupported source type");\r
- }\r
-}\r
-\r
-namespace\r
-{ \r
- typedef NppStatus (*npp_warp_8u_t)(const Npp8u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp8u* pDst, \r
- int dstStep, NppiRect dstRoi, const double coeffs[][3], \r
- int interpolation);\r
- typedef NppStatus (*npp_warp_16u_t)(const Npp16u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp16u* pDst, \r
- int dstStep, NppiRect dstRoi, const double coeffs[][3], \r
- int interpolation);\r
- typedef NppStatus (*npp_warp_32s_t)(const Npp32s* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32s* pDst, \r
- int dstStep, NppiRect dstRoi, const double coeffs[][3], \r
- int interpolation);\r
- typedef NppStatus (*npp_warp_32f_t)(const Npp32f* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32f* pDst, \r
- int dstStep, NppiRect dstRoi, const double coeffs[][3], \r
- int interpolation);\r
-\r
- void nppWarpCaller(const GpuMat& src, GpuMat& dst, double coeffs[][3], const Size& dsize, int flags, \r
- npp_warp_8u_t npp_warp_8u[][2], npp_warp_16u_t npp_warp_16u[][2], \r
- npp_warp_32s_t npp_warp_32s[][2], npp_warp_32f_t npp_warp_32f[][2]) \r
+void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst)\r
+{\r
+ class LevelsInit\r
{\r
- static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC};\r
- \r
- int interpolation = flags & INTER_MAX;\r
-\r
- CV_Assert((src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32S || src.depth() == CV_32F) && src.channels() != 2);\r
- CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC);\r
-\r
- dst.create(dsize, src.type());\r
-\r
- NppiSize srcsz;\r
- srcsz.height = src.rows;\r
- srcsz.width = src.cols;\r
- NppiRect srcroi;\r
- srcroi.x = srcroi.y = 0;\r
- srcroi.height = src.rows;\r
- srcroi.width = src.cols;\r
- NppiRect dstroi;\r
- dstroi.x = dstroi.y = 0;\r
- dstroi.height = dst.rows;\r
- dstroi.width = dst.cols;\r
+ public:\r
+ Npp32s pLevels[256];\r
\r
- int warpInd = (flags & WARP_INVERSE_MAP) >> 4;\r
-\r
- switch (src.depth())\r
- {\r
- case CV_8U:\r
- nppSafeCall( npp_warp_8u[src.channels()][warpInd](src.ptr<Npp8u>(), srcsz, src.step, srcroi, \r
- dst.ptr<Npp8u>(), dst.step, dstroi, coeffs, npp_inter[interpolation]) );\r
- break;\r
- case CV_16U:\r
- nppSafeCall( npp_warp_16u[src.channels()][warpInd](src.ptr<Npp16u>(), srcsz, src.step, srcroi, \r
- dst.ptr<Npp16u>(), dst.step, dstroi, coeffs, npp_inter[interpolation]) );\r
- break;\r
- case CV_32S:\r
- nppSafeCall( npp_warp_32s[src.channels()][warpInd](src.ptr<Npp32s>(), srcsz, src.step, srcroi, \r
- dst.ptr<Npp32s>(), dst.step, dstroi, coeffs, npp_inter[interpolation]) );\r
- break;\r
- case CV_32F:\r
- nppSafeCall( npp_warp_32f[src.channels()][warpInd](src.ptr<Npp32f>(), srcsz, src.step, srcroi, \r
- dst.ptr<Npp32f>(), dst.step, dstroi, coeffs, npp_inter[interpolation]) );\r
- break;\r
- default:\r
- CV_Assert(!"Unsupported source type");\r
+ LevelsInit()\r
+ { \r
+ for (int i = 0; i < 256; ++i)\r
+ pLevels[i] = i;\r
}\r
- }\r
-}\r
-\r
-void cv::gpu::warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags) \r
-{\r
- static npp_warp_8u_t npp_warpAffine_8u[][2] = \r
- {\r
- {0, 0}, \r
- {nppiWarpAffine_8u_C1R, nppiWarpAffineBack_8u_C1R}, \r
- {0, 0}, \r
- {nppiWarpAffine_8u_C3R, nppiWarpAffineBack_8u_C3R}, \r
- {nppiWarpAffine_8u_C4R, nppiWarpAffineBack_8u_C4R}\r
- };\r
- static npp_warp_16u_t npp_warpAffine_16u[][2] = \r
- {\r
- {0, 0}, \r
- {nppiWarpAffine_16u_C1R, nppiWarpAffineBack_16u_C1R}, \r
- {0, 0}, \r
- {nppiWarpAffine_16u_C3R, nppiWarpAffineBack_16u_C3R}, \r
- {nppiWarpAffine_16u_C4R, nppiWarpAffineBack_16u_C4R}\r
- };\r
- static npp_warp_32s_t npp_warpAffine_32s[][2] = \r
- {\r
- {0, 0}, \r
- {nppiWarpAffine_32s_C1R, nppiWarpAffineBack_32s_C1R}, \r
- {0, 0}, \r
- {nppiWarpAffine_32s_C3R, nppiWarpAffineBack_32s_C3R}, \r
- {nppiWarpAffine_32s_C4R, nppiWarpAffineBack_32s_C4R}\r
- };\r
- static npp_warp_32f_t npp_warpAffine_32f[][2] = \r
- {\r
- {0, 0}, \r
- {nppiWarpAffine_32f_C1R, nppiWarpAffineBack_32f_C1R}, \r
- {0, 0}, \r
- {nppiWarpAffine_32f_C3R, nppiWarpAffineBack_32f_C3R}, \r
- {nppiWarpAffine_32f_C4R, nppiWarpAffineBack_32f_C4R}\r
- };\r
-\r
- CV_Assert(M.rows == 2 && M.cols == 3);\r
-\r
- double coeffs[2][3];\r
- Mat coeffsMat(2, 3, CV_64F, (void*)coeffs);\r
- M.convertTo(coeffsMat, coeffsMat.type());\r
-\r
- nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpAffine_8u, npp_warpAffine_16u, npp_warpAffine_32s, npp_warpAffine_32f);\r
-}\r
-\r
-void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags)\r
-{\r
- static npp_warp_8u_t npp_warpPerspective_8u[][2] = \r
- {\r
- {0, 0}, \r
- {nppiWarpPerspective_8u_C1R, nppiWarpPerspectiveBack_8u_C1R}, \r
- {0, 0}, \r
- {nppiWarpPerspective_8u_C3R, nppiWarpPerspectiveBack_8u_C3R}, \r
- {nppiWarpPerspective_8u_C4R, nppiWarpPerspectiveBack_8u_C4R}\r
- };\r
- static npp_warp_16u_t npp_warpPerspective_16u[][2] = \r
- {\r
- {0, 0}, \r
- {nppiWarpPerspective_16u_C1R, nppiWarpPerspectiveBack_16u_C1R}, \r
- {0, 0}, \r
- {nppiWarpPerspective_16u_C3R, nppiWarpPerspectiveBack_16u_C3R}, \r
- {nppiWarpPerspective_16u_C4R, nppiWarpPerspectiveBack_16u_C4R}\r
- };\r
- static npp_warp_32s_t npp_warpPerspective_32s[][2] = \r
- {\r
- {0, 0}, \r
- {nppiWarpPerspective_32s_C1R, nppiWarpPerspectiveBack_32s_C1R}, \r
- {0, 0}, \r
- {nppiWarpPerspective_32s_C3R, nppiWarpPerspectiveBack_32s_C3R}, \r
- {nppiWarpPerspective_32s_C4R, nppiWarpPerspectiveBack_32s_C4R}\r
- };\r
- static npp_warp_32f_t npp_warpPerspective_32f[][2] = \r
- {\r
- {0, 0}, \r
- {nppiWarpPerspective_32f_C1R, nppiWarpPerspectiveBack_32f_C1R}, \r
- {0, 0}, \r
- {nppiWarpPerspective_32f_C3R, nppiWarpPerspectiveBack_32f_C3R}, \r
- {nppiWarpPerspective_32f_C4R, nppiWarpPerspectiveBack_32f_C4R}\r
- };\r
+ };\r
+ static LevelsInit lvls;\r
\r
- CV_Assert(M.rows == 3 && M.cols == 3);\r
+ int cn = src.channels();\r
\r
- double coeffs[3][3];\r
- Mat coeffsMat(3, 3, CV_64F, (void*)coeffs);\r
- M.convertTo(coeffsMat, coeffsMat.type());\r
+ CV_Assert(src.type() == CV_8UC1);\r
+ CV_Assert(lut.depth() == CV_32SC1 && lut.rows * lut.cols == 256 && lut.isContinuous());\r
\r
- nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpPerspective_8u, npp_warpPerspective_16u, npp_warpPerspective_32s, npp_warpPerspective_32f);\r
-}\r
+ dst.create(src.size(), src.type());\r
\r
-void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation)\r
-{\r
- static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC};\r
- \r
- CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
- CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC);\r
-\r
- dst.create(dsize, src.type());\r
-\r
- NppiSize srcsz;\r
- srcsz.height = src.rows;\r
- srcsz.width = src.cols;\r
- NppiRect srcroi;\r
- srcroi.x = srcroi.y = 0;\r
- srcroi.height = src.rows;\r
- srcroi.width = src.cols;\r
- NppiRect dstroi;\r
- dstroi.x = dstroi.y = 0;\r
- dstroi.height = dst.rows;\r
- dstroi.width = dst.cols;\r
+ NppiSize sz;\r
+ sz.height = src.rows;\r
+ sz.width = src.cols;\r
\r
- if (src.type() == CV_8UC1)\r
- {\r
- nppSafeCall( nppiRotate_8u_C1R(src.ptr<Npp8u>(), srcsz, src.step, srcroi, \r
- dst.ptr<Npp8u>(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) );\r
- }\r
- else\r
- {\r
- nppSafeCall( nppiRotate_8u_C4R(src.ptr<Npp8u>(), srcsz, src.step, srcroi, \r
- dst.ptr<Npp8u>(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) );\r
- }\r
+ nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr<Npp8u>(), src.step, dst.ptr<Npp8u>(), dst.step, sz, \r
+ lut.ptr<Npp32s>(), lvls.pLevels, 256) );\r
}\r
\r
#endif /* !defined (HAVE_CUDA) */
\ No newline at end of file
using namespace cv::gpu;\r
\r
#ifndef CV_DESCALE\r
-#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))\r
+#define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n))\r
#endif\r
\r
namespace imgproc\r
{\r
- template<typename _Tp> struct ColorChannel\r
- {\r
- };\r
+ template<typename T, int N> struct TypeVec {};\r
+ template<> struct TypeVec<uchar, 1> { typedef uchar1 vec_t; };\r
+ template<> struct TypeVec<uchar, 2> { typedef uchar2 vec_t; };\r
+ template<> struct TypeVec<uchar, 3> { typedef uchar3 vec_t; };\r
+ template<> struct TypeVec<uchar, 4> { typedef uchar4 vec_t; };\r
+ template<> struct TypeVec<unsigned short, 1> { typedef ushort1 vec_t; };\r
+ template<> struct TypeVec<unsigned short, 2> { typedef ushort2 vec_t; };\r
+ template<> struct TypeVec<unsigned short, 3> { typedef ushort3 vec_t; };\r
+ template<> struct TypeVec<unsigned short, 4> { typedef ushort4 vec_t; };\r
+ template<> struct TypeVec<float, 1> { typedef float1 vec_t; };\r
+ template<> struct TypeVec<float, 2> { typedef float2 vec_t; };\r
+ template<> struct TypeVec<float, 3> { typedef float3 vec_t; };\r
+ template<> struct TypeVec<float, 4> { typedef float4 vec_t; };\r
+\r
+ template<typename _Tp> struct ColorChannel {};\r
\r
template<> struct ColorChannel<uchar>\r
{\r
typedef float worktype_f;\r
- typedef uchar3 vec3_t;\r
- typedef uchar4 vec4_t;\r
static __device__ unsigned char max() { return UCHAR_MAX; }\r
static __device__ unsigned char half() { return (unsigned char)(max()/2 + 1); }\r
};\r
template<> struct ColorChannel<unsigned short>\r
{\r
typedef float worktype_f;\r
- typedef ushort3 vec3_t;\r
- typedef ushort4 vec4_t;\r
static __device__ unsigned short max() { return USHRT_MAX; }\r
static __device__ unsigned short half() { return (unsigned short)(max()/2 + 1); }\r
};\r
template<> struct ColorChannel<float>\r
{\r
typedef float worktype_f;\r
- typedef float3 vec3_t;\r
- typedef float4 vec4_t;\r
static __device__ float max() { return 1.f; }\r
static __device__ float half() { return 0.5f; }\r
- };\r
+ }; \r
}\r
\r
-////////////////// Various 3/4-channel to 3/4-channel RGB transformations /////////////////\r
+//////////////////////////////////////// SwapChannels /////////////////////////////////////\r
\r
namespace imgproc\r
{\r
- template <typename T>\r
- __global__ void RGB2RGB_3_3(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx)\r
- {\r
- const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
- const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+ __constant__ int ccoeffs[4];\r
+\r
+ template <int CN, typename T>\r
+ __global__ void swapChannels(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols)\r
+ {\r
+ typedef typename TypeVec<T, CN>::vec_t vec_t;\r
+\r
+ const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
+ const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
\r
if (y < rows && x < cols)\r
{\r
- const T* src = src_ + y * src_step + x * 3;\r
- T* dst = dst_ + y * dst_step + x * 3;\r
+ vec_t src = *(const vec_t*)(src_ + y * src_step + x * CN);\r
+ vec_t dst;\r
\r
- T t0 = src[bidx], t1 = src[1], t2 = src[bidx ^ 2];\r
- dst[0] = t0; dst[1] = t1; dst[2] = t2;\r
- }\r
+ const T* src_ptr = (const T*)(&src);\r
+ T* dst_ptr = (T*)(&dst);\r
+\r
+ for (int i = 0; i < CN; ++i)\r
+ dst_ptr[i] = src_ptr[ccoeffs[i]];\r
+\r
+ *(vec_t*)(dst_ + y * dst_step + x * CN) = dst;\r
}\r
+ }\r
+}\r
\r
+namespace cv { namespace gpu { namespace improc\r
+{\r
template <typename T>\r
- __global__ void RGB2RGB_4_3(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx)\r
- {\r
- typedef typename ColorChannel<T>::vec4_t vec4_t;\r
+ void swapChannels_caller(const DevMem2D_<T>& src, const DevMem2D_<T>& dst, int cn, const int* coeffs, cudaStream_t stream)\r
+ {\r
+ dim3 threads(32, 8, 1);\r
+ dim3 grid(1, 1, 1);\r
\r
- const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
- const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+ grid.x = divUp(src.cols, threads.x);\r
+ grid.y = divUp(src.rows, threads.y);\r
\r
- if (y < rows && x < cols)\r
- {\r
- vec4_t src = *(vec4_t*)(src_ + y * src_step + (x << 2));\r
- T* dst = dst_ + y * dst_step + x * 3;\r
+ cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, cn * sizeof(int)) );\r
\r
- T t0 = ((T*)(&src))[bidx], t1 = src.y, t2 = ((T*)(&src))[bidx ^ 2];\r
- dst[0] = t0; dst[1] = t1; dst[2] = t2;\r
- }\r
+ switch (cn)\r
+ {\r
+ case 3:\r
+ imgproc::swapChannels<3><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
+ break;\r
+ case 4:\r
+ imgproc::swapChannels<4><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
+ break;\r
+ default:\r
+ cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
+ break;\r
}\r
\r
- template <typename T>\r
- __global__ void RGB2RGB_3_4(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx)\r
- {\r
- typedef typename ColorChannel<T>::vec4_t vec4_t;\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaThreadSynchronize() );\r
+ }\r
\r
- const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
- const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+ void swapChannels_gpu(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream)\r
+ {\r
+ swapChannels_caller(src, dst, cn, coeffs, stream);\r
+ }\r
\r
- if (y < rows && x < cols)\r
- {\r
- const T* src = src_ + y * src_step + x * 3;\r
+ void swapChannels_gpu(const DevMem2D_<unsigned short>& src, const DevMem2D_<unsigned short>& dst, int cn, const int* coeffs, cudaStream_t stream)\r
+ {\r
+ swapChannels_caller(src, dst, cn, coeffs, stream);\r
+ }\r
\r
- vec4_t dst;\r
+ void swapChannels_gpu(const DevMem2Df& src, const DevMem2Df& dst, int cn, const int* coeffs, cudaStream_t stream)\r
+ {\r
+ swapChannels_caller(src, dst, cn, coeffs, stream);\r
+ }\r
+}}}\r
\r
- dst.x = src[bidx];\r
- dst.y = src[1];\r
- dst.z = src[bidx ^ 2];\r
- dst.w = ColorChannel<T>::max();\r
- *(vec4_t*)(dst_ + y * dst_step + (x << 2)) = dst;\r
- }\r
- }\r
+////////////////// Various 3/4-channel to 3/4-channel RGB transformations /////////////////\r
\r
- template <typename T>\r
- __global__ void RGB2RGB_4_4(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx)\r
- {\r
- typedef typename ColorChannel<T>::vec4_t vec4_t;\r
+namespace imgproc\r
+{\r
+ template <int SRCCN, int DSTCN, typename T>\r
+ __global__ void RGB2RGB(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx)\r
+ {\r
+ typedef typename TypeVec<T, SRCCN>::vec_t src_t;\r
+ typedef typename TypeVec<T, DSTCN>::vec_t dst_t;\r
\r
const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
\r
if (y < rows && x < cols)\r
{\r
- vec4_t src = *(const vec4_t*)(src_ + y * src_step + (x << 2));\r
- vec4_t dst;\r
+ src_t src = *(const src_t*)(src_ + y * src_step + x * SRCCN);\r
+ dst_t dst;\r
\r
- dst.x = ((T*)(&src))[bidx];\r
+ dst.x = ((const T*)(&src))[bidx];\r
dst.y = src.y;\r
- dst.z = ((T*)(&src))[bidx ^ 2];\r
- dst.w = src.w;\r
-\r
- *(vec4_t*)(dst_ + y * dst_step + (x << 2)) = dst;\r
- }\r
+ dst.z = ((const T*)(&src))[bidx ^ 2];\r
+ if (DSTCN == 4)\r
+ ((T*)(&dst))[3] = ColorChannel<T>::max();\r
+ \r
+ *(dst_t*)(dst_ + y * dst_step + x * DSTCN) = dst;\r
}\r
+ }\r
}\r
\r
namespace cv { namespace gpu { namespace improc\r
switch (srccn)\r
{\r
case 3:\r
- imgproc::RGB2RGB_3_3<<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T),\r
- src.rows, src.cols, bidx);\r
+ {\r
+ int coeffs[] = {2, 1, 0};\r
+ cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, 3 * sizeof(int)) );\r
+ imgproc::swapChannels<3><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
break;\r
+ }\r
case 4:\r
- imgproc::RGB2RGB_4_3<<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T),\r
- src.rows, src.cols, bidx);\r
+ imgproc::RGB2RGB<4, 3><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T),\r
+ src.rows, src.cols, bidx);\r
break;\r
default:\r
cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
switch (srccn)\r
{\r
case 3:\r
- imgproc::RGB2RGB_3_4<<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T),\r
- src.rows, src.cols, bidx);\r
+ imgproc::RGB2RGB<3, 4><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T),\r
+ src.rows, src.cols, bidx);\r
break;\r
case 4:\r
- imgproc::RGB2RGB_4_4<<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T),\r
- src.rows, src.cols, bidx);\r
+ {\r
+ int coeffs[] = {2, 1, 0, 3};\r
+ cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, 4 * sizeof(int)) );\r
+ imgproc::swapChannels<4><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
break;\r
+ }\r
default:\r
cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
break;\r
template <typename T>\r
__global__ void Gray2RGB_3(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols)\r
{\r
- const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
- const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+ const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
+ const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
\r
if (y < rows && x < cols)\r
{\r
template <typename T>\r
__global__ void Gray2RGB_4(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols)\r
{\r
- typedef typename ColorChannel<T>::vec4_t vec4_t;\r
+ typedef typename TypeVec<T, 4>::vec_t vec4_t;\r
\r
const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
\r
void cv::gpu::erode( const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_nogpu(); }\r
void cv::gpu::dilate( const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_nogpu(); }\r
-void morphologyEx( const GpuMat&, GpuMat&, int, const Mat&, Point, int) { throw_nogpu(); }\r
+void cv::gpu::morphologyEx( const GpuMat&, GpuMat&, int, const Mat&, Point, int) { throw_nogpu(); }\r
\r
#else\r
\r
temp = dst;\r
dilate( src, temp, kernel, anchor, iterations);\r
erode( temp, temp, kernel, anchor, iterations);\r
- dst = temp - src;\r
subtract(temp, src, dst);\r
break;\r
default:\r
void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, const Stream&) { throw_nogpu(); }\r
void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); }\r
void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); }\r
+double cv::gpu::threshold(const GpuMat&, GpuMat&, double) { throw_nogpu(); return 0.0; }\r
+void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int) { throw_nogpu(); }\r
+void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, const Scalar&) { throw_nogpu(); }\r
+void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); }\r
+void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); }\r
+void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int) { throw_nogpu(); }\r
+void cv::gpu::integral(GpuMat&, GpuMat&, GpuMat&) { throw_nogpu(); }\r
+void cv::gpu::boxFilter(const GpuMat&, GpuMat&, Size, Point) { throw_nogpu(); }\r
\r
#else /* !defined (HAVE_CUDA) */\r
\r
void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);\r
void reprojectImageTo3D_gpu(const DevMem2D_<short>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);\r
\r
+ void swapChannels_gpu(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream);\r
+ void swapChannels_gpu(const DevMem2D_<ushort>& src, const DevMem2D_<ushort>& dst, int cn, const int* coeffs, cudaStream_t stream);\r
+ void swapChannels_gpu(const DevMem2Df& src, const DevMem2Df& dst, int cn, const int* coeffs, cudaStream_t stream);\r
+\r
void RGB2RGB_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
void RGB2RGB_gpu(const DevMem2D_<ushort>& src, int srccn, const DevMem2D_<ushort>& dst, int dstcn, int bidx, cudaStream_t stream);\r
void RGB2RGB_gpu(const DevMem2Df& src, int srccn, const DevMem2Df& dst, int dstcn, int bidx, cudaStream_t stream);\r
if (dst.data != src.data)\r
out = dst;\r
\r
+ NppiSize nppsz;\r
+ nppsz.height = src.rows;\r
+ nppsz.width = src.cols;\r
+\r
switch (code)\r
{\r
case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR:\r
// CvtColorLoop(src, dst, Gray2RGB5x5(code == CV_GRAY2BGR565 ? 6 : 5));\r
// break;\r
\r
+ case CV_RGB2YCrCb:\r
+ CV_Assert(scn == 3 && depth == CV_8U);\r
+ \r
+ out.create(sz, CV_MAKETYPE(depth, 3));\r
+\r
+ nppSafeCall( nppiRGBToYCbCr_8u_C3R(src.ptr<Npp8u>(), src.step, out.ptr<Npp8u>(), out.step, nppsz) );\r
+ {\r
+ static int coeffs[] = {0, 2, 1};\r
+ improc::swapChannels_gpu((DevMem2D)out, (DevMem2D)out, 3, coeffs, 0);\r
+ }\r
+ break;\r
+\r
+ case CV_YCrCb2RGB:\r
+ CV_Assert(scn == 3 && depth == CV_8U);\r
+ \r
+ out.create(sz, CV_MAKETYPE(depth, 3));\r
+\r
+ {\r
+ static int coeffs[] = {0, 2, 1};\r
+ GpuMat src1(src.size(), src.type());\r
+ improc::swapChannels_gpu((DevMem2D)src, (DevMem2D)src1, 3, coeffs, 0);\r
+ nppSafeCall( nppiYCbCrToRGB_8u_C3R(src1.ptr<Npp8u>(), src1.step, out.ptr<Npp8u>(), out.step, nppsz) ); \r
+ } \r
+ break;\r
+\r
//case CV_BGR2YCrCb: case CV_RGB2YCrCb:\r
//case CV_BGR2YUV: case CV_RGB2YUV:\r
// {\r
cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream));\r
}\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// threshold\r
+\r
+double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh) \r
+{ \r
+ CV_Assert(src.type() == CV_32FC1)\r
+\r
+ dst.create( src.size(), src.type() );\r
+\r
+ NppiSize sz;\r
+ sz.width = src.cols;\r
+ sz.height = src.rows;\r
+\r
+ nppSafeCall( nppiThreshold_32f_C1R(src.ptr<Npp32f>(), src.step, \r
+ dst.ptr<Npp32f>(), dst.step, sz, static_cast<Npp32f>(thresh), NPP_CMP_GREATER) );\r
+\r
+ return thresh;\r
+}\r
+\r
+////////////////////////////////////////////////////////////////////////\r
+// resize\r
+\r
+void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, double fy, int interpolation)\r
+{\r
+ static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC, 0, NPPI_INTER_LANCZOS};\r
+\r
+ CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
+ CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC || interpolation == INTER_LANCZOS4);\r
+\r
+ CV_Assert( src.size().area() > 0 );\r
+ CV_Assert( !(dsize == Size()) || (fx > 0 && fy > 0) );\r
+\r
+ if( dsize == Size() )\r
+ {\r
+ dsize = Size(saturate_cast<int>(src.cols * fx), saturate_cast<int>(src.rows * fy));\r
+ }\r
+ else\r
+ {\r
+ fx = (double)dsize.width / src.cols;\r
+ fy = (double)dsize.height / src.rows;\r
+ }\r
+\r
+ dst.create(dsize, src.type());\r
+\r
+ NppiSize srcsz;\r
+ srcsz.width = src.cols;\r
+ srcsz.height = src.rows;\r
+ NppiRect srcrect;\r
+ srcrect.x = srcrect.y = 0;\r
+ srcrect.width = src.cols;\r
+ srcrect.height = src.rows;\r
+ NppiSize dstsz;\r
+ dstsz.width = dst.cols;\r
+ dstsz.height = dst.rows;\r
+\r
+ if (src.type() == CV_8UC1)\r
+ {\r
+ nppSafeCall( nppiResize_8u_C1R(src.ptr<Npp8u>(), srcsz, src.step, srcrect,\r
+ dst.ptr<Npp8u>(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) );\r
+ }\r
+ else\r
+ {\r
+ nppSafeCall( nppiResize_8u_C4R(src.ptr<Npp8u>(), srcsz, src.step, srcrect,\r
+ dst.ptr<Npp8u>(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) );\r
+ }\r
+}\r
+\r
+////////////////////////////////////////////////////////////////////////\r
+// copyMakeBorder\r
+\r
+void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value) \r
+{\r
+ CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1);\r
+\r
+ dst.create(src.rows + top + bottom, src.cols + left + right, src.type());\r
+\r
+ NppiSize srcsz;\r
+ srcsz.width = src.cols;\r
+ srcsz.height = src.rows;\r
+ NppiSize dstsz;\r
+ dstsz.width = dst.cols;\r
+ dstsz.height = dst.rows;\r
+\r
+ switch (src.type())\r
+ {\r
+ case CV_8UC1:\r
+ {\r
+ Npp8u nVal = static_cast<Npp8u>(value[0]);\r
+ nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr<Npp8u>(), src.step, srcsz, \r
+ dst.ptr<Npp8u>(), dst.step, dstsz, top, left, nVal) );\r
+ break;\r
+ }\r
+ case CV_8UC4:\r
+ {\r
+ Npp8u nVal[] = {static_cast<Npp8u>(value[0]), static_cast<Npp8u>(value[1]), static_cast<Npp8u>(value[2]), static_cast<Npp8u>(value[3])};\r
+ nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr<Npp8u>(), src.step, srcsz, \r
+ dst.ptr<Npp8u>(), dst.step, dstsz, top, left, nVal) );\r
+ break;\r
+ }\r
+ case CV_32SC1:\r
+ {\r
+ Npp32s nVal = static_cast<Npp32s>(value[0]);\r
+ nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr<Npp32s>(), src.step, srcsz, \r
+ dst.ptr<Npp32s>(), dst.step, dstsz, top, left, nVal) );\r
+ break;\r
+ }\r
+ default:\r
+ CV_Assert(!"Unsupported source type");\r
+ }\r
+}\r
+\r
+////////////////////////////////////////////////////////////////////////\r
+// warp\r
+\r
+namespace\r
+{ \r
+ typedef NppStatus (*npp_warp_8u_t)(const Npp8u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp8u* pDst, \r
+ int dstStep, NppiRect dstRoi, const double coeffs[][3], \r
+ int interpolation);\r
+ typedef NppStatus (*npp_warp_16u_t)(const Npp16u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp16u* pDst, \r
+ int dstStep, NppiRect dstRoi, const double coeffs[][3], \r
+ int interpolation);\r
+ typedef NppStatus (*npp_warp_32s_t)(const Npp32s* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32s* pDst, \r
+ int dstStep, NppiRect dstRoi, const double coeffs[][3], \r
+ int interpolation);\r
+ typedef NppStatus (*npp_warp_32f_t)(const Npp32f* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32f* pDst, \r
+ int dstStep, NppiRect dstRoi, const double coeffs[][3], \r
+ int interpolation);\r
+\r
+ void nppWarpCaller(const GpuMat& src, GpuMat& dst, double coeffs[][3], const Size& dsize, int flags, \r
+ npp_warp_8u_t npp_warp_8u[][2], npp_warp_16u_t npp_warp_16u[][2], \r
+ npp_warp_32s_t npp_warp_32s[][2], npp_warp_32f_t npp_warp_32f[][2]) \r
+ {\r
+ static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC};\r
+ \r
+ int interpolation = flags & INTER_MAX;\r
+\r
+ CV_Assert((src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32S || src.depth() == CV_32F) && src.channels() != 2);\r
+ CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC);\r
+\r
+ dst.create(dsize, src.type());\r
+\r
+ NppiSize srcsz;\r
+ srcsz.height = src.rows;\r
+ srcsz.width = src.cols;\r
+ NppiRect srcroi;\r
+ srcroi.x = srcroi.y = 0;\r
+ srcroi.height = src.rows;\r
+ srcroi.width = src.cols;\r
+ NppiRect dstroi;\r
+ dstroi.x = dstroi.y = 0;\r
+ dstroi.height = dst.rows;\r
+ dstroi.width = dst.cols;\r
+\r
+ int warpInd = (flags & WARP_INVERSE_MAP) >> 4;\r
+\r
+ switch (src.depth())\r
+ {\r
+ case CV_8U:\r
+ nppSafeCall( npp_warp_8u[src.channels()][warpInd](src.ptr<Npp8u>(), srcsz, src.step, srcroi, \r
+ dst.ptr<Npp8u>(), dst.step, dstroi, coeffs, npp_inter[interpolation]) );\r
+ break;\r
+ case CV_16U:\r
+ nppSafeCall( npp_warp_16u[src.channels()][warpInd](src.ptr<Npp16u>(), srcsz, src.step, srcroi, \r
+ dst.ptr<Npp16u>(), dst.step, dstroi, coeffs, npp_inter[interpolation]) );\r
+ break;\r
+ case CV_32S:\r
+ nppSafeCall( npp_warp_32s[src.channels()][warpInd](src.ptr<Npp32s>(), srcsz, src.step, srcroi, \r
+ dst.ptr<Npp32s>(), dst.step, dstroi, coeffs, npp_inter[interpolation]) );\r
+ break;\r
+ case CV_32F:\r
+ nppSafeCall( npp_warp_32f[src.channels()][warpInd](src.ptr<Npp32f>(), srcsz, src.step, srcroi, \r
+ dst.ptr<Npp32f>(), dst.step, dstroi, coeffs, npp_inter[interpolation]) );\r
+ break;\r
+ default:\r
+ CV_Assert(!"Unsupported source type");\r
+ }\r
+ }\r
+}\r
+\r
+void cv::gpu::warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags) \r
+{\r
+ static npp_warp_8u_t npp_warpAffine_8u[][2] = \r
+ {\r
+ {0, 0}, \r
+ {nppiWarpAffine_8u_C1R, nppiWarpAffineBack_8u_C1R}, \r
+ {0, 0}, \r
+ {nppiWarpAffine_8u_C3R, nppiWarpAffineBack_8u_C3R}, \r
+ {nppiWarpAffine_8u_C4R, nppiWarpAffineBack_8u_C4R}\r
+ };\r
+ static npp_warp_16u_t npp_warpAffine_16u[][2] = \r
+ {\r
+ {0, 0}, \r
+ {nppiWarpAffine_16u_C1R, nppiWarpAffineBack_16u_C1R}, \r
+ {0, 0}, \r
+ {nppiWarpAffine_16u_C3R, nppiWarpAffineBack_16u_C3R}, \r
+ {nppiWarpAffine_16u_C4R, nppiWarpAffineBack_16u_C4R}\r
+ };\r
+ static npp_warp_32s_t npp_warpAffine_32s[][2] = \r
+ {\r
+ {0, 0}, \r
+ {nppiWarpAffine_32s_C1R, nppiWarpAffineBack_32s_C1R}, \r
+ {0, 0}, \r
+ {nppiWarpAffine_32s_C3R, nppiWarpAffineBack_32s_C3R}, \r
+ {nppiWarpAffine_32s_C4R, nppiWarpAffineBack_32s_C4R}\r
+ };\r
+ static npp_warp_32f_t npp_warpAffine_32f[][2] = \r
+ {\r
+ {0, 0}, \r
+ {nppiWarpAffine_32f_C1R, nppiWarpAffineBack_32f_C1R}, \r
+ {0, 0}, \r
+ {nppiWarpAffine_32f_C3R, nppiWarpAffineBack_32f_C3R}, \r
+ {nppiWarpAffine_32f_C4R, nppiWarpAffineBack_32f_C4R}\r
+ };\r
+\r
+ CV_Assert(M.rows == 2 && M.cols == 3);\r
+\r
+ double coeffs[2][3];\r
+ Mat coeffsMat(2, 3, CV_64F, (void*)coeffs);\r
+ M.convertTo(coeffsMat, coeffsMat.type());\r
+\r
+ nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpAffine_8u, npp_warpAffine_16u, npp_warpAffine_32s, npp_warpAffine_32f);\r
+}\r
+\r
+void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags)\r
+{\r
+ static npp_warp_8u_t npp_warpPerspective_8u[][2] = \r
+ {\r
+ {0, 0}, \r
+ {nppiWarpPerspective_8u_C1R, nppiWarpPerspectiveBack_8u_C1R}, \r
+ {0, 0}, \r
+ {nppiWarpPerspective_8u_C3R, nppiWarpPerspectiveBack_8u_C3R}, \r
+ {nppiWarpPerspective_8u_C4R, nppiWarpPerspectiveBack_8u_C4R}\r
+ };\r
+ static npp_warp_16u_t npp_warpPerspective_16u[][2] = \r
+ {\r
+ {0, 0}, \r
+ {nppiWarpPerspective_16u_C1R, nppiWarpPerspectiveBack_16u_C1R}, \r
+ {0, 0}, \r
+ {nppiWarpPerspective_16u_C3R, nppiWarpPerspectiveBack_16u_C3R}, \r
+ {nppiWarpPerspective_16u_C4R, nppiWarpPerspectiveBack_16u_C4R}\r
+ };\r
+ static npp_warp_32s_t npp_warpPerspective_32s[][2] = \r
+ {\r
+ {0, 0}, \r
+ {nppiWarpPerspective_32s_C1R, nppiWarpPerspectiveBack_32s_C1R}, \r
+ {0, 0}, \r
+ {nppiWarpPerspective_32s_C3R, nppiWarpPerspectiveBack_32s_C3R}, \r
+ {nppiWarpPerspective_32s_C4R, nppiWarpPerspectiveBack_32s_C4R}\r
+ };\r
+ static npp_warp_32f_t npp_warpPerspective_32f[][2] = \r
+ {\r
+ {0, 0}, \r
+ {nppiWarpPerspective_32f_C1R, nppiWarpPerspectiveBack_32f_C1R}, \r
+ {0, 0}, \r
+ {nppiWarpPerspective_32f_C3R, nppiWarpPerspectiveBack_32f_C3R}, \r
+ {nppiWarpPerspective_32f_C4R, nppiWarpPerspectiveBack_32f_C4R}\r
+ };\r
+\r
+ CV_Assert(M.rows == 3 && M.cols == 3);\r
+\r
+ double coeffs[3][3];\r
+ Mat coeffsMat(3, 3, CV_64F, (void*)coeffs);\r
+ M.convertTo(coeffsMat, coeffsMat.type());\r
+\r
+ nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpPerspective_8u, npp_warpPerspective_16u, npp_warpPerspective_32s, npp_warpPerspective_32f);\r
+}\r
+\r
+////////////////////////////////////////////////////////////////////////\r
+// rotate\r
+\r
+void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation)\r
+{\r
+ static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC};\r
+ \r
+ CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
+ CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC);\r
+\r
+ dst.create(dsize, src.type());\r
+\r
+ NppiSize srcsz;\r
+ srcsz.height = src.rows;\r
+ srcsz.width = src.cols;\r
+ NppiRect srcroi;\r
+ srcroi.x = srcroi.y = 0;\r
+ srcroi.height = src.rows;\r
+ srcroi.width = src.cols;\r
+ NppiRect dstroi;\r
+ dstroi.x = dstroi.y = 0;\r
+ dstroi.height = dst.rows;\r
+ dstroi.width = dst.cols;\r
+\r
+ if (src.type() == CV_8UC1)\r
+ {\r
+ nppSafeCall( nppiRotate_8u_C1R(src.ptr<Npp8u>(), srcsz, src.step, srcroi, \r
+ dst.ptr<Npp8u>(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) );\r
+ }\r
+ else\r
+ {\r
+ nppSafeCall( nppiRotate_8u_C4R(src.ptr<Npp8u>(), srcsz, src.step, srcroi, \r
+ dst.ptr<Npp8u>(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) );\r
+ }\r
+}\r
+\r
+////////////////////////////////////////////////////////////////////////\r
+// integral\r
+\r
+void cv::gpu::integral(GpuMat& src, GpuMat& sum, GpuMat& sqsum)\r
+{\r
+ CV_Assert(src.type() == CV_8UC1);\r
+ \r
+ int w = src.cols + 1, h = src.rows + 1;\r
+\r
+ sum.create(h, w, CV_32S);\r
+ sqsum.create(h, w, CV_32F);\r
+\r
+ NppiSize sz;\r
+ sz.width = src.cols;\r
+ sz.height = src.rows;\r
+\r
+ nppSafeCall( nppiSqrIntegral_8u32s32f_C1R(src.ptr<Npp8u>(), src.step, sum.ptr<Npp32s>(), \r
+ sum.step, sqsum.ptr<Npp32f>(), sqsum.step, sz, 0, 0.0f, h) );\r
+}\r
+\r
+////////////////////////////////////////////////////////////////////////\r
+// boxFilter\r
+\r
+void cv::gpu::boxFilter(const GpuMat& src, GpuMat& dst, Size ksize, Point anchor)\r
+{\r
+ CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
+ CV_Assert(ksize.height == 3 || ksize.height == 5 || ksize.height == 7);\r
+ CV_Assert(ksize.height == ksize.width);\r
+\r
+ if (anchor.x == -1)\r
+ anchor.x = 0;\r
+ if (anchor.y == -1)\r
+ anchor.y = 0;\r
+\r
+ CV_Assert(anchor.x == 0 && anchor.y == 0);\r
+\r
+ dst.create(src.size(), src.type());\r
+\r
+ NppiSize srcsz;\r
+ srcsz.height = src.rows;\r
+ srcsz.width = src.cols;\r
+ NppiSize masksz;\r
+ masksz.height = ksize.height;\r
+ masksz.width = ksize.width;\r
+ NppiPoint anc;\r
+ anc.x = anchor.x;\r
+ anc.y = anchor.y;\r
+\r
+ if (src.type() == CV_8UC1)\r
+ {\r
+ nppSafeCall( nppiFilterBox_8u_C1R(src.ptr<Npp8u>(), src.step, dst.ptr<Npp8u>(), dst.step, srcsz, masksz, anc) );\r
+ }\r
+ else\r
+ {\r
+ nppSafeCall( nppiFilterBox_8u_C4R(src.ptr<Npp8u>(), src.step, dst.ptr<Npp8u>(), dst.step, srcsz, masksz, anc) );\r
+ }\r
+}\r
+\r
#endif /* !defined (HAVE_CUDA) */\r
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+// By downloading, copying, installing or using the software you agree to this license.
+// If you do not agree to this license, do not download, install,
+// copy or use the software.
+//
+//
+// Intel License Agreement
+// For Open Source Computer Vision Library
+//
+// Copyright (C) 2000, Intel Corporation, all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+// * Redistribution's of source code must retain the above copyright notice,
+// this list of conditions and the following disclaimer.
+//
+// * Redistribution's in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// * The name of Intel Corporation may not be used to endorse or promote products
+// derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include <iostream>
+#include <cmath>
+#include <limits>
+#include "gputest.hpp"
+#include "opencv2/core/core.hpp"
+#include "opencv2/imgproc/imgproc.hpp"
+#include "opencv2/highgui/highgui.hpp"
+
+using namespace cv;
+using namespace std;
+using namespace gpu;
+
+class CV_GpuArithmTest : public CvTest
+{
+public:
+ CV_GpuArithmTest(const char* test_name, const char* test_funcs);
+ virtual ~CV_GpuArithmTest();
+
+protected:
+ void run(int);
+
+ int test(int type);
+
+ virtual int test(const Mat& mat1, const Mat& mat2) = 0;
+
+ int CheckNorm(const Mat& m1, const Mat& m2);
+ int CheckNorm(const Scalar& s1, const Scalar& s2);
+ int CheckNorm(double d1, double d2);
+};
+
+CV_GpuArithmTest::CV_GpuArithmTest(const char* test_name, const char* test_funcs): CvTest(test_name, test_funcs)
+{
+}
+
+CV_GpuArithmTest::~CV_GpuArithmTest() {}
+
+int CV_GpuArithmTest::test(int type)
+{
+ cv::Size sz(200, 200);
+ cv::Mat mat1(sz, type), mat2(sz, type);
+ cv::RNG rng(*ts->get_rng());
+ rng.fill(mat1, cv::RNG::UNIFORM, cv::Scalar::all(10), cv::Scalar::all(100));
+ rng.fill(mat2, cv::RNG::UNIFORM, cv::Scalar::all(10), cv::Scalar::all(100));
+
+ return test(mat1, mat2);
+}
+
+int CV_GpuArithmTest::CheckNorm(const Mat& m1, const Mat& m2)
+{
+ double ret = norm(m1, m2, NORM_INF);
+
+ if (ret < std::numeric_limits<double>::epsilon())
+ {
+ return CvTS::OK;
+ }
+ else
+ {
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
+ return CvTS::FAIL_GENERIC;
+ }
+}
+
+int CV_GpuArithmTest::CheckNorm(const Scalar& s1, const Scalar& s2)
+{
+ double ret0 = CheckNorm(s1[0], s2[0]), ret1 = CheckNorm(s1[1], s2[1]), ret2 = CheckNorm(s1[2], s2[2]), ret3 = CheckNorm(s1[3], s2[3]);
+
+ return (ret0 == CvTS::OK && ret1 == CvTS::OK && ret2 == CvTS::OK && ret3 == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC;
+}
+
+int CV_GpuArithmTest::CheckNorm(double d1, double d2)
+{
+ double ret = ::fabs(d1 - d2);
+
+ if (ret < std::numeric_limits<double>::epsilon())
+ {
+ return CvTS::OK;
+ }
+ else
+ {
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
+ return CvTS::FAIL_GENERIC;
+ }
+}
+
+void CV_GpuArithmTest::run( int )
+{
+ int testResult = CvTS::OK;
+ try
+ {
+ //run tests
+ ts->printf(CvTS::LOG, "\n========Start test 8UC1========\n");
+ if (test(CV_8UC1) == CvTS::OK)
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");
+ else
+ {
+ ts->printf(CvTS::LOG, "\nFAIL\n");
+ testResult = CvTS::FAIL_GENERIC;
+ }
+
+ ts->printf(CvTS::LOG, "\n========Start test 8UC3========\n");
+ if (test(CV_8UC3) == CvTS::OK)
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");
+ else
+ {
+ ts->printf(CvTS::LOG, "\nFAIL\n");
+ testResult = CvTS::FAIL_GENERIC;
+ }
+
+ ts->printf(CvTS::LOG, "\n========Start test 8UC4========\n");
+ if (test(CV_8UC4) == CvTS::OK)
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");
+ else
+ {
+ ts->printf(CvTS::LOG, "\nFAIL\n");
+ testResult = CvTS::FAIL_GENERIC;
+ }
+
+ ts->printf(CvTS::LOG, "\n========Start test 32FC1========\n");
+ if (test(CV_32FC1) == CvTS::OK)
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");
+ else
+ {
+ ts->printf(CvTS::LOG, "\nFAIL\n");
+ testResult = CvTS::FAIL_GENERIC;
+ }
+ }
+ catch(const cv::Exception& e)
+ {
+ if (!check_and_treat_gpu_exception(e, ts))
+ throw;
+ return;
+ }
+
+ ts->set_failed_test_info(testResult);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Add
+class CV_GpuNppImageAddTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageAddTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageAddTest::CV_GpuNppImageAddTest(): CV_GpuArithmTest( "GPU-NppImageAdd", "add" )
+{
+}
+
+int CV_GpuNppImageAddTest::test( const Mat& mat1, const Mat& mat2 )
+{
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ cv::Mat cpuRes;
+ cv::add(mat1, mat2, cpuRes);
+
+ GpuMat gpu1(mat1);
+ GpuMat gpu2(mat2);
+ GpuMat gpuRes;
+ cv::gpu::add(gpu1, gpu2, gpuRes);
+
+ return CheckNorm(cpuRes, gpuRes);
+}
+
+CV_GpuNppImageAddTest CV_GpuNppImageAdd_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// Sub
+class CV_GpuNppImageSubtractTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageSubtractTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageSubtractTest::CV_GpuNppImageSubtractTest(): CV_GpuArithmTest( "GPU-NppImageSubtract", "subtract" )
+{
+}
+
+int CV_GpuNppImageSubtractTest::test( const Mat& mat1, const Mat& mat2 )
+{
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ cv::Mat cpuRes;
+ cv::subtract(mat1, mat2, cpuRes);
+
+ GpuMat gpu1(mat1);
+ GpuMat gpu2(mat2);
+ GpuMat gpuRes;
+ cv::gpu::subtract(gpu1, gpu2, gpuRes);
+
+ return CheckNorm(cpuRes, gpuRes);
+}
+
+CV_GpuNppImageSubtractTest CV_GpuNppImageSubtract_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// multiply
+class CV_GpuNppImageMultiplyTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageMultiplyTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageMultiplyTest::CV_GpuNppImageMultiplyTest(): CV_GpuArithmTest( "GPU-NppImageMultiply", "multiply" )
+{
+}
+
+int CV_GpuNppImageMultiplyTest::test( const Mat& mat1, const Mat& mat2 )
+{
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ cv::Mat cpuRes;
+ cv::multiply(mat1, mat2, cpuRes);
+
+ GpuMat gpu1(mat1);
+ GpuMat gpu2(mat2);
+ GpuMat gpuRes;
+ cv::gpu::multiply(gpu1, gpu2, gpuRes);
+
+ return CheckNorm(cpuRes, gpuRes);
+}
+
+CV_GpuNppImageMultiplyTest CV_GpuNppImageMultiply_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// divide
+class CV_GpuNppImageDivideTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageDivideTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageDivideTest::CV_GpuNppImageDivideTest(): CV_GpuArithmTest( "GPU-NppImageDivide", "divide" )
+{
+}
+
+int CV_GpuNppImageDivideTest::test( const Mat& mat1, const Mat& mat2 )
+{
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ cv::Mat cpuRes;
+ cv::divide(mat1, mat2, cpuRes);
+
+ GpuMat gpu1(mat1);
+ GpuMat gpu2(mat2);
+ GpuMat gpuRes;
+ cv::gpu::divide(gpu1, gpu2, gpuRes);
+
+ return CheckNorm(cpuRes, gpuRes);
+}
+
+CV_GpuNppImageDivideTest CV_GpuNppImageDivide_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// transpose
+class CV_GpuNppImageTransposeTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageTransposeTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageTransposeTest::CV_GpuNppImageTransposeTest(): CV_GpuArithmTest( "GPU-NppImageTranspose", "transpose" )
+{
+}
+
+int CV_GpuNppImageTransposeTest::test( const Mat& mat1, const Mat& )
+{
+ if (mat1.type() != CV_8UC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ cv::Mat cpuRes;
+ cv::transpose(mat1, cpuRes);
+
+ GpuMat gpu1(mat1);
+ GpuMat gpuRes;
+ cv::gpu::transpose(gpu1, gpuRes);
+
+ return CheckNorm(cpuRes, gpuRes);
+}
+
+CV_GpuNppImageTransposeTest CV_GpuNppImageTranspose_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// absdiff
+class CV_GpuNppImageAbsdiffTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageAbsdiffTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageAbsdiffTest::CV_GpuNppImageAbsdiffTest(): CV_GpuArithmTest( "GPU-NppImageAbsdiff", "absdiff" )
+{
+}
+
+int CV_GpuNppImageAbsdiffTest::test( const Mat& mat1, const Mat& mat2 )
+{
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_32FC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ cv::Mat cpuRes;
+ cv::absdiff(mat1, mat2, cpuRes);
+
+ GpuMat gpu1(mat1);
+ GpuMat gpu2(mat2);
+ GpuMat gpuRes;
+ cv::gpu::absdiff(gpu1, gpu2, gpuRes);
+
+ return CheckNorm(cpuRes, gpuRes);
+}
+
+CV_GpuNppImageAbsdiffTest CV_GpuNppImageAbsdiff_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// compare
+class CV_GpuNppImageCompareTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageCompareTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageCompareTest::CV_GpuNppImageCompareTest(): CV_GpuArithmTest( "GPU-NppImageCompare", "compare" )
+{
+}
+
+int CV_GpuNppImageCompareTest::test( const Mat& mat1, const Mat& mat2 )
+{
+ if (mat1.type() != CV_32FC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ int cmp_codes[] = {CMP_EQ, CMP_GT, CMP_GE, CMP_LT, CMP_LE, CMP_NE};
+ const char* cmp_str[] = {"CMP_EQ", "CMP_GT", "CMP_GE", "CMP_LT", "CMP_LE", "CMP_NE"};
+ int cmp_num = sizeof(cmp_codes) / sizeof(int);
+
+ int test_res = CvTS::OK;
+
+ for (int i = 0; i < cmp_num; ++i)
+ {
+ ts->printf(CvTS::LOG, "\nCompare operation: %s\n", cmp_str[i]);
+
+ cv::Mat cpuRes;
+ cv::compare(mat1, mat2, cpuRes, cmp_codes[i]);
+
+ GpuMat gpu1(mat1);
+ GpuMat gpu2(mat2);
+ GpuMat gpuRes;
+ cv::gpu::compare(gpu1, gpu2, gpuRes, cmp_codes[i]);
+
+ if (CheckNorm(cpuRes, gpuRes) != CvTS::OK)
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ return test_res;
+}
+
+CV_GpuNppImageCompareTest CV_GpuNppImageCompare_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// meanStdDev
+class CV_GpuNppImageMeanStdDevTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageMeanStdDevTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageMeanStdDevTest::CV_GpuNppImageMeanStdDevTest(): CV_GpuArithmTest( "GPU-NppImageMeanStdDev", "meanStdDev" )
+{
+}
+
+int CV_GpuNppImageMeanStdDevTest::test( const Mat& mat1, const Mat& )
+{
+ if (mat1.type() != CV_8UC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ Scalar cpumean;
+ Scalar cpustddev;
+ cv::meanStdDev(mat1, cpumean, cpustddev);
+
+ GpuMat gpu1(mat1);
+ Scalar gpumean;
+ Scalar gpustddev;
+ cv::gpu::meanStdDev(gpu1, gpumean, gpustddev);
+
+ int test_res = CvTS::OK;
+
+ if (CheckNorm(cpumean, gpumean) != CvTS::OK)
+ {
+ ts->printf(CvTS::LOG, "\nMean FAILED\n");
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ if (CheckNorm(cpustddev, gpustddev) != CvTS::OK)
+ {
+ ts->printf(CvTS::LOG, "\nStdDev FAILED\n");
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ return test_res;
+}
+
+CV_GpuNppImageMeanStdDevTest CV_GpuNppImageMeanStdDev_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// norm
+class CV_GpuNppImageNormTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageNormTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageNormTest::CV_GpuNppImageNormTest(): CV_GpuArithmTest( "GPU-NppImageNorm", "norm" )
+{
+}
+
+int CV_GpuNppImageNormTest::test( const Mat& mat1, const Mat& mat2 )
+{
+ if (mat1.type() != CV_8UC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ int norms[] = {NORM_INF, NORM_L1, NORM_L2};
+ const char* norms_str[] = {"NORM_INF", "NORM_L1", "NORM_L2"};
+ int norms_num = sizeof(norms) / sizeof(int);
+
+ int test_res = CvTS::OK;
+
+ for (int i = 0; i < norms_num; ++i)
+ {
+ ts->printf(CvTS::LOG, "\nNorm type: %s\n", norms_str[i]);
+
+ double cpu_norm = cv::norm(mat1, mat2, norms[i]);
+
+ GpuMat gpu1(mat1);
+ GpuMat gpu2(mat2);
+ double gpu_norm = cv::gpu::norm(gpu1, gpu2, norms[i]);
+
+ if (CheckNorm(cpu_norm, gpu_norm) != CvTS::OK)
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ return test_res;
+}
+
+CV_GpuNppImageNormTest CV_GpuNppImageNorm_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// flip
+class CV_GpuNppImageFlipTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageFlipTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageFlipTest::CV_GpuNppImageFlipTest(): CV_GpuArithmTest( "GPU-NppImageFlip", "flip" )
+{
+}
+
+int CV_GpuNppImageFlipTest::test( const Mat& mat1, const Mat& )
+{
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ int flip_codes[] = {0, 1, -1};
+ const char* flip_axis[] = {"X", "Y", "Both"};
+ int flip_codes_num = sizeof(flip_codes) / sizeof(int);
+
+ int test_res = CvTS::OK;
+
+ for (int i = 0; i < flip_codes_num; ++i)
+ {
+ ts->printf(CvTS::LOG, "\nFlip Axis: %s\n", flip_axis[i]);
+
+ Mat cpu_res;
+ cv::flip(mat1, cpu_res, flip_codes[i]);
+
+ GpuMat gpu1(mat1);
+ GpuMat gpu_res;
+ cv::gpu::flip(gpu1, gpu_res, flip_codes[i]);
+
+ if (CheckNorm(cpu_res, gpu_res) != CvTS::OK)
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ return test_res;
+}
+
+CV_GpuNppImageFlipTest CV_GpuNppImageFlip_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// sum
+class CV_GpuNppImageSumTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageSumTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageSumTest::CV_GpuNppImageSumTest(): CV_GpuArithmTest( "GPU-NppImageSum", "sum" )
+{
+}
+
+int CV_GpuNppImageSumTest::test( const Mat& mat1, const Mat& )
+{
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ Scalar cpures = cv::sum(mat1);
+
+ GpuMat gpu1(mat1);
+ Scalar gpures = cv::gpu::sum(gpu1);
+
+ return CheckNorm(cpures, gpures);
+}
+
+CV_GpuNppImageSumTest CV_GpuNppImageSum_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// minNax
+class CV_GpuNppImageMinNaxTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageMinNaxTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageMinNaxTest::CV_GpuNppImageMinNaxTest(): CV_GpuArithmTest( "GPU-NppImageMinNax", "minNax" )
+{
+}
+
+int CV_GpuNppImageMinNaxTest::test( const Mat& mat1, const Mat& )
+{
+ if (mat1.type() != CV_8UC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ double cpumin, cpumax;
+ cv::minMaxLoc(mat1, &cpumin, &cpumax);
+
+ GpuMat gpu1(mat1);
+ double gpumin, gpumax;
+ cv::gpu::minMax(gpu1, &gpumin, &gpumax);
+
+ return (CheckNorm(cpumin, gpumin) == CvTS::OK && CheckNorm(cpumax, gpumax) == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC;
+}
+
+CV_GpuNppImageMinNaxTest CV_GpuNppImageMinNax_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// LUT
+class CV_GpuNppImageLUTTest : public CV_GpuArithmTest
+{
+public:
+ CV_GpuNppImageLUTTest();
+
+protected:
+ virtual int test(const Mat& mat1, const Mat& mat2);
+};
+
+CV_GpuNppImageLUTTest::CV_GpuNppImageLUTTest(): CV_GpuArithmTest( "GPU-NppImageLUT", "LUT" )
+{
+}
+
+int CV_GpuNppImageLUTTest::test( const Mat& mat1, const Mat& )
+{
+ if (mat1.type() != CV_8UC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ cv::Mat lut(1, 256, CV_32SC1);
+ cv::RNG rng(*ts->get_rng());
+ rng.fill(lut, cv::RNG::UNIFORM, cv::Scalar::all(100), cv::Scalar::all(200));
+
+ cv::Mat cpuRes;
+ cv::LUT(mat1, lut, cpuRes);
+ cpuRes.convertTo(cpuRes, CV_8U);
+
+ cv::gpu::GpuMat gpuRes;
+ cv::gpu::LUT(GpuMat(mat1), lut, gpuRes);
+
+ return CheckNorm(cpuRes, gpuRes);
+}
+
+CV_GpuNppImageLUTTest CV_GpuNppImageLUT_test;
switch (e.code)
{
case CV_GpuNotFound:
- ts->printf(CvTS::CONSOLE, "\nGpu not found");
+ ts->printf(CvTS::LOG, "\nGpu not found");
break;
case CV_GpuApiCallError:
- ts->printf(CvTS::CONSOLE, "\nGPU Error: %s", e.what());
+ ts->printf(CvTS::LOG, "\nGPU Error: %s", e.what());
break;
case CV_GpuNppCallError:
- ts->printf(CvTS::CONSOLE, "\nNPP Error: %s", e.what());
+ ts->printf(CvTS::LOG, "\nNPP Error: %s", e.what());
break;
default:
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+// By downloading, copying, installing or using the software you agree to this license.
+// If you do not agree to this license, do not download, install,
+// copy or use the software.
+//
+//
+// Intel License Agreement
+// For Open Source Computer Vision Library
+//
+// Copyright (C) 2000, Intel Corporation, all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+// * Redistribution's of source code must retain the above copyright notice,
+// this list of conditions and the following disclaimer.
+//
+// * Redistribution's in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// * The name of Intel Corporation may not be used to endorse or promote products
+// derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include <iostream>
+#include <cmath>
+#include <limits>
+#include "gputest.hpp"
+#include "opencv2/core/core.hpp"
+#include "opencv2/imgproc/imgproc.hpp"
+#include "opencv2/highgui/highgui.hpp"
+
+using namespace cv;
+using namespace std;
+using namespace gpu;
+
+class CV_GpuImageProcTest : public CvTest
+{
+public:
+ CV_GpuImageProcTest(const char* test_name, const char* test_funcs);
+ virtual ~CV_GpuImageProcTest();
+
+protected:
+ void run(int);
+
+ int test8UC1 (const Mat& img);
+ int test8UC4 (const Mat& img);
+ int test32SC1(const Mat& img);
+ int test32FC1(const Mat& img);
+
+ virtual int test(const Mat& img) = 0;
+
+ int CheckNorm(const Mat& m1, const Mat& m2);
+};
+
+CV_GpuImageProcTest::CV_GpuImageProcTest(const char* test_name, const char* test_funcs): CvTest(test_name, test_funcs)
+{
+}
+
+CV_GpuImageProcTest::~CV_GpuImageProcTest() {}
+
+int CV_GpuImageProcTest::test8UC1(const Mat& img)
+{
+ cv::Mat img_C1;
+ cvtColor(img, img_C1, CV_BGR2GRAY);
+
+ return test(img_C1);
+}
+
+int CV_GpuImageProcTest::test8UC4(const Mat& img)
+{
+ cv::Mat img_C4;
+ cvtColor(img, img_C4, CV_BGR2BGRA);
+
+ return test(img_C4);
+}
+
+int CV_GpuImageProcTest::test32SC1(const Mat& img)
+{
+ cv::Mat img_C1;
+ cvtColor(img, img_C1, CV_BGR2GRAY);
+ img_C1.convertTo(img_C1, CV_32S);
+
+ return test(img_C1);
+}
+
+int CV_GpuImageProcTest::test32FC1(const Mat& img)
+{
+ cv::Mat temp, img_C1;
+ img.convertTo(temp, CV_32F);
+ cvtColor(temp, img_C1, CV_BGR2GRAY);
+
+ return test(img_C1);
+}
+
+int CV_GpuImageProcTest::CheckNorm(const Mat& m1, const Mat& m2)
+{
+ double ret = norm(m1, m2, NORM_INF);
+
+ if (ret < std::numeric_limits<double>::epsilon())
+ {
+ return CvTS::OK;
+ }
+ else
+ {
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
+ return CvTS::FAIL_GENERIC;
+ }
+}
+
+void CV_GpuImageProcTest::run( int )
+{
+ //load image
+ cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
+
+ if (img.empty())
+ {
+ ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
+ return;
+ }
+
+ int testResult = CvTS::OK;
+ try
+ {
+ //run tests
+ ts->printf(CvTS::LOG, "\n========Start test 8UC1========\n");
+ if (test8UC1(img) == CvTS::OK)
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");
+ else
+ {
+ ts->printf(CvTS::LOG, "\nFAIL\n");
+ testResult = CvTS::FAIL_GENERIC;
+ }
+
+ ts->printf(CvTS::LOG, "\n========Start test 8UC4========\n");
+ if (test8UC4(img) == CvTS::OK)
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");
+ else
+ {
+ ts->printf(CvTS::LOG, "\nFAIL\n");
+ testResult = CvTS::FAIL_GENERIC;
+ }
+
+ ts->printf(CvTS::LOG, "\n========Start test 32SC1========\n");
+ if (test32SC1(img) == CvTS::OK)
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");
+ else
+ {
+ ts->printf(CvTS::LOG, "\nFAIL\n");
+ testResult = CvTS::FAIL_GENERIC;
+ }
+
+ ts->printf(CvTS::LOG, "\n========Start test 32FC1========\n");
+ if (test32FC1(img) == CvTS::OK)
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");
+ else
+ {
+ ts->printf(CvTS::LOG, "\nFAIL\n");
+ testResult = CvTS::FAIL_GENERIC;
+ }
+ }
+ catch(const cv::Exception& e)
+ {
+ if (!check_and_treat_gpu_exception(e, ts))
+ throw;
+ return;
+ }
+
+ ts->set_failed_test_info(testResult);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// threshold
+class CV_GpuNppImageThresholdTest : public CV_GpuImageProcTest
+{
+public:
+ CV_GpuNppImageThresholdTest();
+
+protected:
+ virtual int test(const Mat& img);
+};
+
+CV_GpuNppImageThresholdTest::CV_GpuNppImageThresholdTest(): CV_GpuImageProcTest( "GPU-NppImageThreshold", "threshold" )
+{
+}
+
+int CV_GpuNppImageThresholdTest::test(const Mat& img)
+{
+ if (img.type() != CV_32FC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ cv::RNG rng(*ts->get_rng());
+ const double thresh = rng;
+
+ cv::Mat cpuRes;
+ cv::threshold(img, cpuRes, thresh, 0.0, THRESH_TRUNC);
+
+ GpuMat gpu1(img);
+ GpuMat gpuRes;
+ cv::gpu::threshold(gpu1, gpuRes, thresh);
+
+ return CheckNorm(cpuRes, gpuRes);
+}
+
+CV_GpuNppImageThresholdTest CV_GpuNppImageThreshold_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// resize
+class CV_GpuNppImageResizeTest : public CV_GpuImageProcTest
+{
+public:
+ CV_GpuNppImageResizeTest();
+
+protected:
+ virtual int test(const Mat& img);
+};
+
+CV_GpuNppImageResizeTest::CV_GpuNppImageResizeTest(): CV_GpuImageProcTest( "GPU-NppImageResize", "resize" )
+{
+}
+
+int CV_GpuNppImageResizeTest::test(const Mat& img)
+{
+ if (img.type() != CV_8UC1 && img.type() != CV_8UC4)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ int interpolations[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_LANCZOS4};
+ const char* interpolations_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_LANCZOS4"};
+ int interpolations_num = sizeof(interpolations) / sizeof(int);
+
+ int test_res = CvTS::OK;
+
+ for (int i = 0; i < interpolations_num; ++i)
+ {
+ ts->printf(CvTS::LOG, "\nInterpolation type: %s\n", interpolations_str[i]);
+
+ Mat cpu_res;
+ cv::resize(img, cpu_res, Size(), 0.5, 0.5, interpolations[i]);
+
+ GpuMat gpu1(img), gpu_res;
+ cv::gpu::resize(gpu1, gpu_res, Size(), 0.5, 0.5, interpolations[i]);
+
+ if (CheckNorm(cpu_res, gpu_res) != CvTS::OK)
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ return test_res;
+}
+
+//CV_GpuNppImageResizeTest CV_GpuNppImageResize_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// copyMakeBorder
+class CV_GpuNppImageCopyMakeBorderTest : public CV_GpuImageProcTest
+{
+public:
+ CV_GpuNppImageCopyMakeBorderTest();
+
+protected:
+ virtual int test(const Mat& img);
+};
+
+CV_GpuNppImageCopyMakeBorderTest::CV_GpuNppImageCopyMakeBorderTest(): CV_GpuImageProcTest( "GPU-NppImageCopyMakeBorder", "copyMakeBorder" )
+{
+}
+
+int CV_GpuNppImageCopyMakeBorderTest::test(const Mat& img)
+{
+ if (img.type() != CV_8UC1 && img.type() != CV_8UC4 && img.type() != CV_32SC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ cv::RNG rng(*ts->get_rng());
+ int top = rng.uniform(1, 10);
+ int botton = rng.uniform(1, 10);
+ int left = rng.uniform(1, 10);
+ int right = rng.uniform(1, 10);
+ cv::Scalar val(rng.uniform(0, 255), rng.uniform(0, 255), rng.uniform(0, 255), rng.uniform(0, 255));
+
+ Mat cpudst;
+ cv::copyMakeBorder(img, cpudst, top, botton, left, right, BORDER_CONSTANT, val);
+
+ GpuMat gpu1(img);
+ GpuMat gpudst;
+ cv::gpu::copyMakeBorder(gpu1, gpudst, top, botton, left, right, val);
+
+ return CheckNorm(cpudst, gpudst);
+}
+
+CV_GpuNppImageCopyMakeBorderTest CV_GpuNppImageCopyMakeBorder_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// warpAffine
+class CV_GpuNppImageWarpAffineTest : public CV_GpuImageProcTest
+{
+public:
+ CV_GpuNppImageWarpAffineTest();
+
+protected:
+ virtual int test(const Mat& img);
+};
+
+CV_GpuNppImageWarpAffineTest::CV_GpuNppImageWarpAffineTest(): CV_GpuImageProcTest( "GPU-NppImageWarpAffine", "warpAffine" )
+{
+}
+
+int CV_GpuNppImageWarpAffineTest::test(const Mat& img)
+{
+ if (img.type() == CV_32SC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ static const double coeffs[2][3] =
+ {
+ {cos(3.14 / 6), -sin(3.14 / 6), 100.0},
+ {sin(3.14 / 6), cos(3.14 / 6), -100.0}
+ };
+ Mat M(2, 3, CV_64F, (void*)coeffs);
+
+ int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};
+ const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};
+ int flags_num = sizeof(flags) / sizeof(int);
+
+ int test_res = CvTS::OK;
+
+ for (int i = 0; i < flags_num; ++i)
+ {
+ ts->printf(CvTS::LOG, "\nFlags: %s\n", flags_str[i]);
+
+ Mat cpudst;
+ cv::warpAffine(img, cpudst, M, img.size(), flags[i]);
+
+ GpuMat gpu1(img);
+ GpuMat gpudst;
+ cv::gpu::warpAffine(gpu1, gpudst, M, gpu1.size(), flags[i]);
+
+ if (CheckNorm(cpudst, gpudst) != CvTS::OK)
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ return test_res;
+}
+
+//CV_GpuNppImageWarpAffineTest CV_GpuNppImageWarpAffine_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// warpPerspective
+class CV_GpuNppImageWarpPerspectiveTest : public CV_GpuImageProcTest
+{
+public:
+ CV_GpuNppImageWarpPerspectiveTest();
+
+protected:
+ virtual int test(const Mat& img);
+};
+
+CV_GpuNppImageWarpPerspectiveTest::CV_GpuNppImageWarpPerspectiveTest(): CV_GpuImageProcTest( "GPU-NppImageWarpPerspective", "warpPerspective" )
+{
+}
+
+int CV_GpuNppImageWarpPerspectiveTest::test(const Mat& img)
+{
+ if (img.type() == CV_32SC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ static const double coeffs[3][3] =
+ {
+ {cos(3.14 / 6), -sin(3.14 / 6), 100.0},
+ {sin(3.14 / 6), cos(3.14 / 6), -100.0},
+ {0.0, 0.0, 1.0}
+ };
+ Mat M(3, 3, CV_64F, (void*)coeffs);
+
+ int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};
+ const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};
+ int flags_num = sizeof(flags) / sizeof(int);
+
+ int test_res = CvTS::OK;
+
+ for (int i = 0; i < flags_num; ++i)
+ {
+ ts->printf(CvTS::LOG, "\nFlags: %s\n", flags_str[i]);
+
+ Mat cpudst;
+ cv::warpPerspective(img, cpudst, M, img.size(), flags[i]);
+
+ GpuMat gpu1(img);
+ GpuMat gpudst;
+ cv::gpu::warpPerspective(gpu1, gpudst, M, gpu1.size(), flags[i]);
+
+ if (CheckNorm(cpudst, gpudst) != CvTS::OK)
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ return test_res;
+}
+
+//CV_GpuNppImageWarpPerspectiveTest CV_GpuNppImageWarpPerspective_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// integral
+class CV_GpuNppImageIntegralTest : public CV_GpuImageProcTest
+{
+public:
+ CV_GpuNppImageIntegralTest();
+
+protected:
+ virtual int test(const Mat& img);
+};
+
+CV_GpuNppImageIntegralTest::CV_GpuNppImageIntegralTest(): CV_GpuImageProcTest( "GPU-NppImageIntegral", "integral" )
+{
+}
+
+int CV_GpuNppImageIntegralTest::test(const Mat& img)
+{
+ if (img.type() != CV_8UC1)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ Mat cpusum, cpusqsum;
+ cv::integral(img, cpusum, cpusqsum, CV_32S);
+
+ GpuMat gpu1(img);
+ GpuMat gpusum, gpusqsum;
+ cv::gpu::integral(gpu1, gpusum, gpusqsum);
+
+ gpusqsum.convertTo(gpusqsum, CV_64F);
+
+ int test_res = CvTS::OK;
+
+ if (CheckNorm(cpusum, gpusum) != CvTS::OK)
+ {
+ ts->printf(CvTS::LOG, "\nSum failed\n");
+ test_res = CvTS::FAIL_GENERIC;
+ }
+ if (CheckNorm(cpusqsum, gpusqsum) != CvTS::OK)
+ {
+ ts->printf(CvTS::LOG, "\nSquared sum failed\n");
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ return test_res;
+}
+
+CV_GpuNppImageIntegralTest CV_GpuNppImageIntegral_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// blur
+class CV_GpuNppImageBlurTest : public CV_GpuImageProcTest
+{
+public:
+ CV_GpuNppImageBlurTest();
+
+protected:
+ virtual int test(const Mat& img);
+};
+
+CV_GpuNppImageBlurTest::CV_GpuNppImageBlurTest(): CV_GpuImageProcTest( "GPU-NppImageBlur", "blur" )
+{
+}
+
+int CV_GpuNppImageBlurTest::test(const Mat& img)
+{
+ if (img.type() != CV_8UC1 && img.type() != CV_8UC4)
+ {
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");
+ return CvTS::OK;
+ }
+
+ int ksizes[] = {3, 5, 7};
+ int ksizes_num = sizeof(ksizes) / sizeof(int);
+
+ int test_res = CvTS::OK;
+
+ for (int i = 0; i < ksizes_num; ++i)
+ {
+ ts->printf(CvTS::LOG, "\nksize = %d\n", ksizes[i]);
+
+ Mat cpudst;
+ cv::blur(img, cpudst, Size(ksizes[i], ksizes[i]));
+
+ GpuMat gpu1(img);
+ GpuMat gpudst;
+ cv::gpu::blur(gpu1, gpudst, Size(ksizes[i], ksizes[i]));
+
+ cv::Mat c;
+ cv::absdiff(cpudst, gpudst, c);
+
+ if (CheckNorm(cpudst, gpudst) != CvTS::OK)
+ test_res = CvTS::FAIL_GENERIC;
+ }
+
+ return test_res;
+}
+
+//CV_GpuNppImageBlurTest CV_GpuNppImageBlur_test;
+
+////////////////////////////////////////////////////////////////////////////////
+// cvtColor
+class CV_GpuCvtColorTest : public CvTest
+{
+public:
+ CV_GpuCvtColorTest();
+
+protected:
+ void run(int);
+
+ int CheckNorm(const Mat& m1, const Mat& m2);
+};
+
+CV_GpuCvtColorTest::CV_GpuCvtColorTest(): CvTest("GPU-NppCvtColor", "cvtColor")
+{
+}
+
+int CV_GpuCvtColorTest::CheckNorm(const Mat& m1, const Mat& m2)
+{
+ double ret = norm(m1, m2, NORM_INF);
+
+ if (ret < std::numeric_limits<double>::epsilon())
+ {
+ return CvTS::OK;
+ }
+ else
+ {
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
+ return CvTS::FAIL_GENERIC;
+ }
+}
+
+void CV_GpuCvtColorTest::run( int )
+{
+ //load image
+ cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
+
+ if (img.empty())
+ {
+ ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
+ return;
+ }
+
+ int testResult = CvTS::OK;
+ cv::Mat cpuRes;
+ cv::gpu::GpuMat gpuImg(img), gpuRes;
+ try
+ {
+ //run tests
+ int codes[] = {CV_BGR2RGB, CV_RGB2YCrCb, CV_YCrCb2RGB, CV_RGB2RGBA, CV_RGBA2BGRA, CV_BGRA2GRAY, CV_GRAY2RGB};
+ const char* codes_str[] = {"CV_BGR2RGB", "CV_RGB2YCrCb", "CV_YCrCb2RGB", "CV_RGB2RGBA", "CV_RGBA2BGRA", "CV_BGRA2GRAY", "CV_GRAY2RGB"};
+ int codes_num = sizeof(codes) / sizeof(int);
+
+ for (int i = 0; i < codes_num; ++i)
+ {
+ ts->printf(CvTS::LOG, "\n%s\n", codes_str[i]);
+
+ cv::cvtColor(img, cpuRes, codes[i]);
+ cv::gpu::cvtColor(gpuImg, gpuRes, codes[i]);
+
+ if (CheckNorm(cpuRes, gpuRes) == CvTS::OK)
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");
+ else
+ {
+ ts->printf(CvTS::LOG, "\nFAIL\n");
+ testResult = CvTS::FAIL_GENERIC;
+ }
+
+ img = cpuRes;
+ gpuImg = gpuRes;
+ }
+ }
+ catch(const cv::Exception& e)
+ {
+ if (!check_and_treat_gpu_exception(e, ts))
+ throw;
+ return;
+ }
+
+ ts->set_failed_test_info(testResult);
+}
+
+CV_GpuCvtColorTest CV_GpuCvtColor_test;
}\r
if (maxDiff > 0) \r
{\r
- ts->printf(CvTS::CONSOLE, "\nMeanShift maxDiff = %d\n", maxDiff);\r
+ ts->printf(CvTS::LOG, "\nMeanShift maxDiff = %d\n", maxDiff);\r
ts->set_failed_test_info(CvTS::FAIL_GENERIC);\r
return;\r
}\r
if (res < std::numeric_limits<double>::epsilon())
return CvTS::OK;
- ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", res);
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", res);
return CvTS::FAIL_GENERIC;
}
};
+++ /dev/null
-/*M///////////////////////////////////////////////////////////////////////////////////////
-//
-// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
-//
-// By downloading, copying, installing or using the software you agree to this license.
-// If you do not agree to this license, do not download, install,
-// copy or use the software.
-//
-//
-// Intel License Agreement
-// For Open Source Computer Vision Library
-//
-// Copyright (C) 2000, Intel Corporation, all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// Redistribution and use in source and binary forms, with or without modification,
-// are permitted provided that the following conditions are met:
-//
-// * Redistribution's of source code must retain the above copyright notice,
-// this list of conditions and the following disclaimer.
-//
-// * Redistribution's in binary form must reproduce the above copyright notice,
-// this list of conditions and the following disclaimer in the documentation
-// and/or other materials provided with the distribution.
-//
-// * The name of Intel Corporation may not be used to endorse or promote products
-// derived from this software without specific prior written permission.
-//
-// This software is provided by the copyright holders and contributors "as is" and
-// any express or implied warranties, including, but not limited to, the implied
-// warranties of merchantability and fitness for a particular purpose are disclaimed.
-// In no event shall the Intel Corporation or contributors be liable for any direct,
-// indirect, incidental, special, exemplary, or consequential damages
-// (including, but not limited to, procurement of substitute goods or services;
-// loss of use, data, or profits; or business interruption) however caused
-// and on any theory of liability, whether in contract, strict liability,
-// or tort (including negligence or otherwise) arising in any way out of
-// the use of this software, even if advised of the possibility of such damage.
-//
-//M*/
-
-#include <iostream>
-#include <cmath>
-#include <limits>
-#include "gputest.hpp"
-#include "opencv2/imgproc/imgproc.hpp"
-#include "opencv2/highgui/highgui.hpp"
-
-using namespace cv;
-using namespace std;
-using namespace gpu;
-
-class CV_GpuNppImageArithmTest : public CvTest
-{
-public:
- CV_GpuNppImageArithmTest(const char* test_name, const char* test_funcs);
- virtual ~CV_GpuNppImageArithmTest();
-
-protected:
- void run(int);
-
- int test8UC1(const Mat& cpu1, const Mat& cpu2);
- int test8UC4(const Mat& cpu1, const Mat& cpu2);
- int test32SC1(const Mat& cpu1, const Mat& cpu2);
- int test32FC1(const Mat& cpu1, const Mat& cpu2);
-
- virtual int test(const Mat& cpu1, const Mat& cpu2) = 0;
- int CheckNorm(const Mat& m1, const Mat& m2);
- int CheckNorm(const Scalar& s1, const Scalar& s2);
- int CheckNorm(double d1, double d2);
-};
-
-CV_GpuNppImageArithmTest::CV_GpuNppImageArithmTest(const char* test_name, const char* test_funcs): CvTest(test_name, test_funcs)
-{
-}
-
-CV_GpuNppImageArithmTest::~CV_GpuNppImageArithmTest() {}
-
-int CV_GpuNppImageArithmTest::test8UC1(const Mat& cpu1, const Mat& cpu2)
-{
- cv::Mat imgL_C1;
- cv::Mat imgR_C1;
- cvtColor(cpu1, imgL_C1, CV_BGR2GRAY);
- cvtColor(cpu2, imgR_C1, CV_BGR2GRAY);
-
- return test(imgL_C1, imgR_C1);
-}
-
-int CV_GpuNppImageArithmTest::test8UC4(const Mat& cpu1, const Mat& cpu2)
-{
- cv::Mat imgL_C4;
- cv::Mat imgR_C4;
- cvtColor(cpu1, imgL_C4, CV_BGR2BGRA);
- cvtColor(cpu2, imgR_C4, CV_BGR2BGRA);
-
- return test(imgL_C4, imgR_C4);
-}
-
-int CV_GpuNppImageArithmTest::test32SC1( const Mat& cpu1, const Mat& cpu2 )
-{
- cv::Mat imgL_C1;
- cv::Mat imgR_C1;
- cvtColor(cpu1, imgL_C1, CV_BGR2GRAY);
- cvtColor(cpu2, imgR_C1, CV_BGR2GRAY);
-
- imgL_C1.convertTo(imgL_C1, CV_32S);
- imgR_C1.convertTo(imgR_C1, CV_32S);
-
- return test(imgL_C1, imgR_C1);
-}
-
-int CV_GpuNppImageArithmTest::test32FC1( const Mat& cpu1, const Mat& cpu2 )
-{
- cv::Mat imgL_C1;
- cv::Mat imgR_C1;
- cvtColor(cpu1, imgL_C1, CV_BGR2GRAY);
- cvtColor(cpu2, imgR_C1, CV_BGR2GRAY);
-
- imgL_C1.convertTo(imgL_C1, CV_32F);
- imgR_C1.convertTo(imgR_C1, CV_32F);
-
- return test(imgL_C1, imgR_C1);
-}
-
-int CV_GpuNppImageArithmTest::CheckNorm(const Mat& m1, const Mat& m2)
-{
- double ret = norm(m1, m2, NORM_INF);
-
- if (ret < std::numeric_limits<double>::epsilon())
- {
- return CvTS::OK;
- }
- else
- {
- ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret);
- return CvTS::FAIL_GENERIC;
- }
-}
-
-int CV_GpuNppImageArithmTest::CheckNorm(const Scalar& s1, const Scalar& s2)
-{
- double ret0 = CheckNorm(s1[0], s2[0]), ret1 = CheckNorm(s1[1], s2[1]), ret2 = CheckNorm(s1[2], s2[2]), ret3 = CheckNorm(s1[3], s2[3]);
-
- return (ret0 == CvTS::OK && ret1 == CvTS::OK && ret2 == CvTS::OK && ret3 == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC;
-}
-
-int CV_GpuNppImageArithmTest::CheckNorm(double d1, double d2)
-{
- double ret = ::fabs(d1 - d2);
-
- if (ret < std::numeric_limits<double>::epsilon())
- {
- return CvTS::OK;
- }
- else
- {
- ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret);
- return CvTS::FAIL_GENERIC;
- }
-}
-
-void CV_GpuNppImageArithmTest::run( int )
-{
- //load images
- //cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-L.png");
- //cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-R.png");
-
- //cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
- //cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-R.png");
-
- cv::RNG rng(*ts->get_rng());
- cv::Size sz(200, 200);
- cv::Mat img_l(sz, CV_8UC3), img_r(sz, CV_8UC3);
- rng.fill(img_l, cv::RNG::UNIFORM, cv::Scalar::all(10), cv::Scalar::all(100));
- rng.fill(img_r, cv::RNG::UNIFORM, cv::Scalar::all(10), cv::Scalar::all(100));
-
- if (img_l.empty() || img_r.empty())
- {
- ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
- return;
- }
-
- try
- {
- //run tests
- int testResult = test8UC1(img_l, img_r);
- if (testResult != CvTS::OK)
- {
- ts->set_failed_test_info(testResult);
- return;
- }
-
- testResult = test8UC4(img_l, img_r);
- if (testResult != CvTS::OK)
- {
- ts->set_failed_test_info(testResult);
- return;
- }
-
- testResult = test32SC1(img_l, img_r);
- if (testResult != CvTS::OK)
- {
- ts->set_failed_test_info(testResult);
- return;
- }
-
- testResult = test32FC1(img_l, img_r);
- if (testResult != CvTS::OK)
- {
- ts->set_failed_test_info(testResult);
- return;
- }
- }
- catch(const cv::Exception& e)
- {
- if (!check_and_treat_gpu_exception(e, ts))
- throw;
- return;
- }
-
- ts->set_failed_test_info(CvTS::OK);
-}
-
-////////////////////////////////////////////////////////////////////////////////
-// Add
-class CV_GpuNppImageAddTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageAddTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageAddTest::CV_GpuNppImageAddTest(): CV_GpuNppImageArithmTest( "GPU-NppImageAdd", "add" )
-{
-}
-
-int CV_GpuNppImageAddTest::test( const Mat& cpu1, const Mat& cpu2 )
-{
- if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4 && cpu1.type() != CV_32FC1)
- return CvTS::OK;
-
- cv::Mat cpuRes;
- cv::add(cpu1, cpu2, cpuRes);
-
- GpuMat gpu1(cpu1);
- GpuMat gpu2(cpu2);
- GpuMat gpuRes;
- cv::gpu::add(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
-}
-
-CV_GpuNppImageAddTest CV_GpuNppImageAdd_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// Sub
-class CV_GpuNppImageSubtractTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageSubtractTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageSubtractTest::CV_GpuNppImageSubtractTest(): CV_GpuNppImageArithmTest( "GPU-NppImageSubtract", "subtract" )
-{
-}
-
-int CV_GpuNppImageSubtractTest::test( const Mat& cpu1, const Mat& cpu2 )
-{
- if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4 && cpu1.type() != CV_32FC1)
- return CvTS::OK;
-
- cv::Mat cpuRes;
- cv::subtract(cpu1, cpu2, cpuRes);
-
- GpuMat gpu1(cpu1);
- GpuMat gpu2(cpu2);
- GpuMat gpuRes;
- cv::gpu::subtract(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
-}
-
-CV_GpuNppImageSubtractTest CV_GpuNppImageSubtract_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// multiply
-class CV_GpuNppImageMultiplyTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageMultiplyTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageMultiplyTest::CV_GpuNppImageMultiplyTest(): CV_GpuNppImageArithmTest( "GPU-NppImageMultiply", "multiply" )
-{
-}
-
-int CV_GpuNppImageMultiplyTest::test( const Mat& cpu1, const Mat& cpu2 )
-{
- if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4 && cpu1.type() != CV_32FC1)
- return CvTS::OK;
-
- cv::Mat cpuRes;
- cv::multiply(cpu1, cpu2, cpuRes);
-
- GpuMat gpu1(cpu1);
- GpuMat gpu2(cpu2);
- GpuMat gpuRes;
- cv::gpu::multiply(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
-}
-
-CV_GpuNppImageMultiplyTest CV_GpuNppImageMultiply_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// divide
-class CV_GpuNppImageDivideTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageDivideTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageDivideTest::CV_GpuNppImageDivideTest(): CV_GpuNppImageArithmTest( "GPU-NppImageDivide", "divide" )
-{
-}
-
-int CV_GpuNppImageDivideTest::test( const Mat& cpu1, const Mat& cpu2 )
-{
- if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4 && cpu1.type() != CV_32FC1)
- return CvTS::OK;
-
- cv::Mat cpuRes;
- cv::divide(cpu1, cpu2, cpuRes);
-
- GpuMat gpu1(cpu1);
- GpuMat gpu2(cpu2);
- GpuMat gpuRes;
- cv::gpu::divide(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
-}
-
-CV_GpuNppImageDivideTest CV_GpuNppImageDivide_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// transpose
-class CV_GpuNppImageTransposeTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageTransposeTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageTransposeTest::CV_GpuNppImageTransposeTest(): CV_GpuNppImageArithmTest( "GPU-NppImageTranspose", "transpose" )
-{
-}
-
-int CV_GpuNppImageTransposeTest::test( const Mat& cpu1, const Mat& )
-{
- if (cpu1.type() != CV_8UC1)
- return CvTS::OK;
-
- cv::Mat cpuRes;
- cv::transpose(cpu1, cpuRes);
-
- GpuMat gpu1(cpu1);
- GpuMat gpuRes;
- cv::gpu::transpose(gpu1, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
-}
-
-CV_GpuNppImageTransposeTest CV_GpuNppImageTranspose_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// absdiff
-class CV_GpuNppImageAbsdiffTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageAbsdiffTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageAbsdiffTest::CV_GpuNppImageAbsdiffTest(): CV_GpuNppImageArithmTest( "GPU-NppImageAbsdiff", "absdiff" )
-{
-}
-
-int CV_GpuNppImageAbsdiffTest::test( const Mat& cpu1, const Mat& cpu2 )
-{
- if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_32FC1)
- return CvTS::OK;
-
- cv::Mat cpuRes;
- cv::absdiff(cpu1, cpu2, cpuRes);
-
- GpuMat gpu1(cpu1);
- GpuMat gpu2(cpu2);
- GpuMat gpuRes;
- cv::gpu::absdiff(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
-}
-
-CV_GpuNppImageAbsdiffTest CV_GpuNppImageAbsdiff_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// threshold
-class CV_GpuNppImageThresholdTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageThresholdTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageThresholdTest::CV_GpuNppImageThresholdTest(): CV_GpuNppImageArithmTest( "GPU-NppImageThreshold", "threshold" )
-{
-}
-
-int CV_GpuNppImageThresholdTest::test( const Mat& cpu1, const Mat& )
-{
- if (cpu1.type() != CV_32FC1)
- return CvTS::OK;
-
- cv::RNG rng(*ts->get_rng());
- const double thresh = rng;
-
- cv::Mat cpuRes;
- cv::threshold(cpu1, cpuRes, thresh, 0.0, THRESH_TRUNC);
-
- GpuMat gpu1(cpu1);
- GpuMat gpuRes;
- cv::gpu::threshold(gpu1, gpuRes, thresh);
-
- return CheckNorm(cpuRes, gpuRes);
-}
-
-CV_GpuNppImageThresholdTest CV_GpuNppImageThreshold_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// compare
-class CV_GpuNppImageCompareTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageCompareTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageCompareTest::CV_GpuNppImageCompareTest(): CV_GpuNppImageArithmTest( "GPU-NppImageCompare", "compare" )
-{
-}
-
-int CV_GpuNppImageCompareTest::test( const Mat& cpu1, const Mat& cpu2 )
-{
- if (cpu1.type() != CV_32FC1)
- return CvTS::OK;
-
- int cmp_codes[] = {CMP_EQ, CMP_GT, CMP_GE, CMP_LT, CMP_LE, CMP_NE};
- const char* cmp_str[] = {"CMP_EQ", "CMP_GT", "CMP_GE", "CMP_LT", "CMP_LE", "CMP_NE"};
- int cmp_num = sizeof(cmp_codes) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < cmp_num; ++i)
- {
- cv::Mat cpuRes;
- cv::compare(cpu1, cpu2, cpuRes, cmp_codes[i]);
-
- GpuMat gpu1(cpu1);
- GpuMat gpu2(cpu2);
- GpuMat gpuRes;
- cv::gpu::compare(gpu1, gpu2, gpuRes, cmp_codes[i]);
-
- if (CheckNorm(cpuRes, gpuRes) != CvTS::OK)
- {
- ts->printf(CvTS::CONSOLE, "\nCompare operation: %s\n", cmp_str[i]);
- test_res = CvTS::FAIL_GENERIC;
- }
- }
-
- return test_res;
-}
-
-CV_GpuNppImageCompareTest CV_GpuNppImageCompare_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// meanStdDev
-class CV_GpuNppImageMeanStdDevTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageMeanStdDevTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageMeanStdDevTest::CV_GpuNppImageMeanStdDevTest(): CV_GpuNppImageArithmTest( "GPU-NppImageMeanStdDev", "meanStdDev" )
-{
-}
-
-int CV_GpuNppImageMeanStdDevTest::test( const Mat& cpu1, const Mat& )
-{
- if (cpu1.type() != CV_8UC1)
- return CvTS::OK;
-
- Scalar cpumean;
- Scalar cpustddev;
- cv::meanStdDev(cpu1, cpumean, cpustddev);
-
- GpuMat gpu1(cpu1);
- Scalar gpumean;
- Scalar gpustddev;
- cv::gpu::meanStdDev(gpu1, gpumean, gpustddev);
-
- return (CheckNorm(cpumean, gpumean) == CvTS::OK && CheckNorm(cpustddev, gpustddev) == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC;
-}
-
-CV_GpuNppImageMeanStdDevTest CV_GpuNppImageMeanStdDev_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// norm
-class CV_GpuNppImageNormTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageNormTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageNormTest::CV_GpuNppImageNormTest(): CV_GpuNppImageArithmTest( "GPU-NppImageNorm", "norm" )
-{
-}
-
-int CV_GpuNppImageNormTest::test( const Mat& cpu1, const Mat& cpu2 )
-{
- if (cpu1.type() != CV_8UC1)
- return CvTS::OK;
-
- int norms[] = {NORM_INF, NORM_L1, NORM_L2};
- const char* norms_str[] = {"NORM_INF", "NORM_L1", "NORM_L2"};
- int norms_num = sizeof(norms) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < norms_num; ++i)
- {
- double cpu_norm = cv::norm(cpu1, cpu2, norms[i]);
-
- GpuMat gpu1(cpu1);
- GpuMat gpu2(cpu2);
- double gpu_norm = cv::gpu::norm(gpu1, gpu2, norms[i]);
-
- if (CheckNorm(cpu_norm, gpu_norm) != CvTS::OK)
- {
- ts->printf(CvTS::CONSOLE, "\nNorm type: %s\n", norms_str[i]);
- test_res = CvTS::FAIL_GENERIC;
- }
- }
-
- return test_res;
-}
-
-CV_GpuNppImageNormTest CV_GpuNppImageNorm_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// flip
-class CV_GpuNppImageFlipTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageFlipTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageFlipTest::CV_GpuNppImageFlipTest(): CV_GpuNppImageArithmTest( "GPU-NppImageFlip", "flip" )
-{
-}
-
-int CV_GpuNppImageFlipTest::test( const Mat& cpu1, const Mat& )
-{
- if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4)
- return CvTS::OK;
-
- int flip_codes[] = {0, 1, -1};
- const char* flip_axis[] = {"X", "Y", "Both"};
- int flip_codes_num = sizeof(flip_codes) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < flip_codes_num; ++i)
- {
- Mat cpu_res;
- cv::flip(cpu1, cpu_res, flip_codes[i]);
-
- GpuMat gpu1(cpu1);
- GpuMat gpu_res;
- cv::gpu::flip(gpu1, gpu_res, flip_codes[i]);
-
- if (CheckNorm(cpu_res, gpu_res) != CvTS::OK)
- {
- ts->printf(CvTS::CONSOLE, "\nFlip Axis: %s\n", flip_axis[i]);
- test_res = CvTS::FAIL_GENERIC;
- }
- }
-
- return test_res;
-}
-
-CV_GpuNppImageFlipTest CV_GpuNppImageFlip_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// resize
-class CV_GpuNppImageResizeTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageResizeTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageResizeTest::CV_GpuNppImageResizeTest(): CV_GpuNppImageArithmTest( "GPU-NppImageResize", "resize" )
-{
-}
-
-int CV_GpuNppImageResizeTest::test( const Mat& cpu1, const Mat& )
-{
- if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4)
- return CvTS::OK;
-
- int interpolations[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_LANCZOS4};
- const char* interpolations_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_LANCZOS4"};
- int interpolations_num = sizeof(interpolations) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < interpolations_num; ++i)
- {
- Mat cpu_res;
- cv::resize(cpu1, cpu_res, Size(), 0.5, 0.5, interpolations[i]);
-
- GpuMat gpu1(cpu1), gpu_res;
- cv::gpu::resize(gpu1, gpu_res, Size(), 0.5, 0.5, interpolations[i]);
-
- if (CheckNorm(cpu_res, gpu_res) != CvTS::OK)
- {
- ts->printf(CvTS::CONSOLE, "\nInterpolation type: %s\n", interpolations_str[i]);
- test_res = CvTS::FAIL_GENERIC;
- }
- }
-
- return test_res;
-}
-
-CV_GpuNppImageResizeTest CV_GpuNppImageResize_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// sum
-class CV_GpuNppImageSumTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageSumTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageSumTest::CV_GpuNppImageSumTest(): CV_GpuNppImageArithmTest( "GPU-NppImageSum", "sum" )
-{
-}
-
-int CV_GpuNppImageSumTest::test( const Mat& cpu1, const Mat& )
-{
- if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4)
- return CvTS::OK;
-
- Scalar cpures = cv::sum(cpu1);
-
- GpuMat gpu1(cpu1);
- Scalar gpures = cv::gpu::sum(gpu1);
-
- return CheckNorm(cpures, gpures);
-}
-
-CV_GpuNppImageSumTest CV_GpuNppImageSum_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// minNax
-class CV_GpuNppImageMinNaxTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageMinNaxTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageMinNaxTest::CV_GpuNppImageMinNaxTest(): CV_GpuNppImageArithmTest( "GPU-NppImageMinNax", "minNax" )
-{
-}
-
-int CV_GpuNppImageMinNaxTest::test( const Mat& cpu1, const Mat& )
-{
- if (cpu1.type() != CV_8UC1)
- return CvTS::OK;
-
- double cpumin, cpumax;
- cv::minMaxLoc(cpu1, &cpumin, &cpumax);
-
- GpuMat gpu1(cpu1);
- double gpumin, gpumax;
- cv::gpu::minMax(gpu1, &gpumin, &gpumax);
-
- return (CheckNorm(cpumin, gpumin) == CvTS::OK && CheckNorm(cpumax, gpumax) == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC;
-}
-
-CV_GpuNppImageMinNaxTest CV_GpuNppImageMinNax_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// copyConstBorder
-class CV_GpuNppImageCopyMakeBorderTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageCopyMakeBorderTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageCopyMakeBorderTest::CV_GpuNppImageCopyMakeBorderTest(): CV_GpuNppImageArithmTest( "GPU-NppImageCopyMakeBorder", "copyMakeBorder" )
-{
-}
-
-int CV_GpuNppImageCopyMakeBorderTest::test( const Mat& cpu1, const Mat& )
-{
- if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4 && cpu1.type() != CV_32SC1)
- return CvTS::OK;
-
- Mat cpudst;
- cv::copyMakeBorder(cpu1, cpudst, 5, 5, 5, 5, BORDER_CONSTANT);
-
- GpuMat gpu1(cpu1);
- GpuMat gpudst;
- cv::gpu::copyMakeBorder(gpu1, gpudst, 5, 5, 5, 5);
-
- return CheckNorm(cpudst, gpudst);
-}
-
-CV_GpuNppImageCopyMakeBorderTest CV_GpuNppImageCopyMakeBorder_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// warpAffine
-class CV_GpuNppImageWarpAffineTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageWarpAffineTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageWarpAffineTest::CV_GpuNppImageWarpAffineTest(): CV_GpuNppImageArithmTest( "GPU-NppImageWarpAffine", "warpAffine" )
-{
-}
-
-int CV_GpuNppImageWarpAffineTest::test( const Mat& cpu1, const Mat& )
-{
- static const double coeffs[2][3] =
- {
- {cos(3.14 / 6), -sin(3.14 / 6), 100.0},
- {sin(3.14 / 6), cos(3.14 / 6), -100.0}
- };
- Mat M(2, 3, CV_64F, (void*)coeffs);
-
- if (cpu1.type() == CV_32SC1)
- return CvTS::OK;
-
- int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};
- const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};
- int flags_num = sizeof(flags) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < flags_num; ++i)
- {
- Mat cpudst;
- cv::warpAffine(cpu1, cpudst, M, cpu1.size(), flags[i]);
-
- GpuMat gpu1(cpu1);
- GpuMat gpudst;
- cv::gpu::warpAffine(gpu1, gpudst, M, gpu1.size(), flags[i]);
-
- if (CheckNorm(cpudst, gpudst) != CvTS::OK)
- {
- ts->printf(CvTS::CONSOLE, "\nFlags: %s\n", flags_str[i]);
- test_res = CvTS::FAIL_GENERIC;
- }
- }
-
- return test_res;
-}
-
-CV_GpuNppImageWarpAffineTest CV_GpuNppImageWarpAffine_test;
-
-////////////////////////////////////////////////////////////////////////////////
-// warpAffine
-class CV_GpuNppImageWarpPerspectiveTest : public CV_GpuNppImageArithmTest
-{
-public:
- CV_GpuNppImageWarpPerspectiveTest();
-
-protected:
- virtual int test(const Mat& cpu1, const Mat& cpu2);
-};
-
-CV_GpuNppImageWarpPerspectiveTest::CV_GpuNppImageWarpPerspectiveTest(): CV_GpuNppImageArithmTest( "GPU-NppImageWarpPerspective", "warpPerspective" )
-{
-}
-
-int CV_GpuNppImageWarpPerspectiveTest::test( const Mat& cpu1, const Mat& )
-{
- static const double coeffs[3][3] =
- {
- {cos(3.14 / 6), -sin(3.14 / 6), 100.0},
- {sin(3.14 / 6), cos(3.14 / 6), -100.0},
- {0.0, 0.0, 1.0}
- };
- Mat M(3, 3, CV_64F, (void*)coeffs);
-
- if (cpu1.type() == CV_32SC1)
- return CvTS::OK;
-
- int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};
- const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};
- int flags_num = sizeof(flags) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < flags_num; ++i)
- {
- Mat cpudst;
- cv::warpPerspective(cpu1, cpudst, M, cpu1.size(), flags[i]);
-
- GpuMat gpu1(cpu1);
- GpuMat gpudst;
- cv::gpu::warpPerspective(gpu1, gpudst, M, gpu1.size(), flags[i]);
-
- if (CheckNorm(cpudst, gpudst) != CvTS::OK)
- {
- ts->printf(CvTS::CONSOLE, "\nFlags: %s\n", flags_str[i]);
- test_res = CvTS::FAIL_GENERIC;
- }
- }
-
- return test_res;
-}
-
-CV_GpuNppImageWarpPerspectiveTest CV_GpuNppImageWarpPerspective_test;
return true;
else
{
- ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret);
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
return false;
}
}
double r = norm(cpumatdst, gpumatdst, NORM_INF);
if (r > 1)
{
- ts->printf(CvTS::CONSOLE,
+ ts->printf(CvTS::LOG,
"\nFAILED: SRC_TYPE=%sC%d DST_TYPE=%s NORM = %d\n",
types_str[i], c, types_str[j], r);
passed = false;
return true;
else
{
- ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret);
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
return false;
}
}
return true;
else
{
- ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret);
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
return false;
}
}
if (norm >= 100)
{
- ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm);
+ ts->printf(CvTS::LOG, "\nStereoBM norm = %f\n", norm);
ts->set_failed_test_info(CvTS::FAIL_GENERIC);
return;
}
if (norm >= 100)
{
- ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm);
+ ts->printf(CvTS::LOG, "\nStereoBM norm = %f\n", norm);
ts->set_failed_test_info(CvTS::FAIL_GENERIC);
return;
}
double norm = cv::norm(disp, img_template, cv::NORM_INF);
if (norm >= 0.5)
{
- ts->printf(CvTS::CONSOLE, "\nStereoBP norm = %f\n", norm);
+ ts->printf(CvTS::LOG, "\nStereoBP norm = %f\n", norm);
ts->set_failed_test_info(CvTS::FAIL_GENERIC);
return;
}
double norm = cv::norm(disp, img_template, cv::NORM_INF);
if (norm >= 0.5)
{
- ts->printf(CvTS::CONSOLE, "\nConstantSpaceStereoBP norm = %f\n", norm);
+ ts->printf(CvTS::LOG, "\nConstantSpaceStereoBP norm = %f\n", norm);
ts->set_failed_test_info(CvTS::FAIL_GENERIC);
return;
}
projects / profile / ivi / opencv.git / commitdiff