From: Vladislav Vinogradov Date: Wed, 22 Sep 2010 10:58:01 +0000 (+0000) Subject: added gpu version of LUT, integral, boxFilter and cvtColor (RGB <-> YCrCb), based... X-Git-Tag: accepted/2.0/20130307.220821~4467 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=4100cbd9973c1e563a53e1ed13a22d9992673c7f;p=profile%2Fivi%2Fopencv.git added gpu version of LUT, integral, boxFilter and cvtColor (RGB <-> YCrCb), based on NPP. minor refactoring of GPU module and GPU tests, split arithm and imgproc parts. --- diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 0a5d8ee..38edb9d 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -360,11 +360,7 @@ namespace cv CV_EXPORTS void transpose(const GpuMat& src1, GpuMat& dst); //! computes element-wise absolute difference of two arrays (c = abs(a - b)) - CV_EXPORTS void absdiff(const GpuMat& a, const GpuMat& b, GpuMat& c); - - //! applies fixed threshold to the image. - //! Now supports only THRESH_TRUNC threshold type and one channels float source. - CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh); + CV_EXPORTS void absdiff(const GpuMat& a, const GpuMat& b, GpuMat& c); //! compares elements of two arrays (c = a b) //! Now doesn't support CMP_NE. @@ -383,30 +379,17 @@ namespace cv //! reverses the order of the rows, columns or both in a matrix CV_EXPORTS void flip(const GpuMat& a, GpuMat& b, int flipCode); - //! resizes the image - //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_LANCZOS4 - CV_EXPORTS void resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, double fy=0, int interpolation = INTER_LINEAR); - //! computes sum of array elements CV_EXPORTS Scalar sum(const GpuMat& m); //! finds global minimum and maximum array elements and returns their values CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal = 0); - //! copies 2D array to a larger destination array and pads borders with user-specifiable constant - CV_EXPORTS void copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value = Scalar()); - - //! warps the image using affine transformation - //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC - CV_EXPORTS void warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR); - - //! warps the image using perspective transformation - //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC - CV_EXPORTS void warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR); - - //! rotate 8bit single or four channel image - //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC - CV_EXPORTS void rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift = 0, double yShift = 0, int interpolation = INTER_LINEAR); + //! transforms 8-bit unsigned integers using lookup table: dst(i)=lut(src(i)) + //! supports only single channels source + //! destination array will have the same type as source + //! lut must hase CV_32S depth and the same number of channels as in the source array + CV_EXPORTS void LUT(const GpuMat& src, const Mat& lut, GpuMat& dst); //! makes multi-channel array out of several single-channel arrays CV_EXPORTS void merge(const GpuMat* src, size_t n, GpuMat& dst); @@ -434,33 +417,69 @@ namespace cv ////////////////////////////// Image processing ////////////////////////////// - // DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation. - // xymap.type() == xymap.type() == CV_32FC1 + //! DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation. + //! xymap.type() == xymap.type() == CV_32FC1 CV_EXPORTS void remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap); - // Does mean shift filtering on GPU. + //! Does mean shift filtering on GPU. CV_EXPORTS void meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1)); - // Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV. - // Supported types of input disparity: CV_8U, CV_16S. - // Output disparity has CV_8UC4 type in BGRA format (alpha = 255). + //! Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV. + //! Supported types of input disparity: CV_8U, CV_16S. + //! Output disparity has CV_8UC4 type in BGRA format (alpha = 255). CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndisp); - // Acync version + //! Acync version CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndisp, const Stream& stream); - // Reprojects disparity image to 3D space. - // Supports CV_8U and CV_16S types of input disparity. - // The output is a 4-channel floating-point (CV_32FC4) matrix. - // Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map. - // Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify. + //! Reprojects disparity image to 3D space. + //! Supports CV_8U and CV_16S types of input disparity. + //! The output is a 4-channel floating-point (CV_32FC4) matrix. + //! Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map. + //! Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify. CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q); - // Acync version + //! Acync version CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream); + //! converts image from one color space to another CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn = 0); + //! Acync version CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream); + //! applies fixed threshold to the image. + //! Now supports only THRESH_TRUNC threshold type and one channels float source. + CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh); + + //! resizes the image + //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_LANCZOS4 + CV_EXPORTS void resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, double fy=0, int interpolation = INTER_LINEAR); + + //! warps the image using affine transformation + //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC + CV_EXPORTS void warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR); + + //! warps the image using perspective transformation + //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC + CV_EXPORTS void warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR); + + //! rotate 8bit single or four channel image + //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC + CV_EXPORTS void rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift = 0, double yShift = 0, int interpolation = INTER_LINEAR); + + //! copies 2D array to a larger destination array and pads borders with user-specifiable constant + CV_EXPORTS void copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value = Scalar()); + + //! computes the integral image and integral for the squared image + //! sum will have CV_32S type, sqsum - CV32F type + CV_EXPORTS void integral(GpuMat& src, GpuMat& sum, GpuMat& sqsum); + + //! smooths the image using the normalized box filter + CV_EXPORTS void boxFilter(const GpuMat& src, GpuMat& dst, Size ksize, Point anchor = Point(-1,-1)); + //! a synonym for normalized box filter + static inline void blur(const GpuMat& src, GpuMat& dst, Size ksize, Point anchor = Point(-1,-1)) + { + boxFilter(src, dst, ksize, anchor); + } //! erodes the image (applies the local minimum operator) CV_EXPORTS void erode( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations); diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index 5fbebf4..a21acbb 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -52,38 +52,22 @@ void cv::gpu::add(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::subtract(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::multiply(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::divide(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } - void cv::gpu::transpose(const GpuMat&, GpuMat&) { throw_nogpu(); } - void cv::gpu::absdiff(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } - -double cv::gpu::threshold(const GpuMat&, GpuMat&, double) { throw_nogpu(); return 0.0; } - void cv::gpu::compare(const GpuMat&, const GpuMat&, GpuMat&, int) { throw_nogpu(); } - void cv::gpu::meanStdDev(const GpuMat&, Scalar&, Scalar&) { throw_nogpu(); } - double cv::gpu::norm(const GpuMat&, int) { throw_nogpu(); return 0.0; } double cv::gpu::norm(const GpuMat&, const GpuMat&, int) { throw_nogpu(); return 0.0; } - void cv::gpu::flip(const GpuMat&, GpuMat&, int) { throw_nogpu(); } - -void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int) { throw_nogpu(); } - Scalar cv::gpu::sum(const GpuMat&) { throw_nogpu(); return Scalar(); } - void cv::gpu::minMax(const GpuMat&, double*, double*) { throw_nogpu(); } - -void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, const Scalar&) { throw_nogpu(); } - -void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); } - -void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); } - -void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int) { throw_nogpu(); } +void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst) { throw_nogpu(); } #else /* !defined (HAVE_CUDA) */ +//////////////////////////////////////////////////////////////////////// +// add subtract multiply divide + namespace { typedef NppStatus (*npp_arithm_8u_t)(const Npp8u* pSrc1, int nSrc1Step, const Npp8u* pSrc2, int nSrc2Step, Npp8u* pDst, int nDstStep, @@ -147,6 +131,9 @@ void cv::gpu::divide(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) nppFuncCaller(src2, src1, dst, nppiDiv_8u_C1RSfs, nppiDiv_8u_C4RSfs, nppiDiv_32f_C1R); } +//////////////////////////////////////////////////////////////////////// +// transpose + void cv::gpu::transpose(const GpuMat& src, GpuMat& dst) { CV_Assert(src.type() == CV_8UC1); @@ -160,6 +147,9 @@ void cv::gpu::transpose(const GpuMat& src, GpuMat& dst) nppSafeCall( nppiTranspose_8u_C1R(src.ptr(), src.step, dst.ptr(), dst.step, sz) ); } +//////////////////////////////////////////////////////////////////////// +// absdiff + void cv::gpu::absdiff(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) { CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type()); @@ -186,21 +176,8 @@ void cv::gpu::absdiff(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) } } -double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh) -{ - CV_Assert(src.type() == CV_32FC1) - - dst.create( src.size(), src.type() ); - - NppiSize sz; - sz.width = src.cols; - sz.height = src.rows; - - nppSafeCall( nppiThreshold_32f_C1R(src.ptr(), src.step, - dst.ptr(), dst.step, sz, static_cast(thresh), NPP_CMP_GREATER) ); - - return thresh; -} +//////////////////////////////////////////////////////////////////////// +// compare namespace cv { namespace gpu { namespace matrix_operations { @@ -250,6 +227,9 @@ void cv::gpu::compare(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, int c } } +//////////////////////////////////////////////////////////////////////// +// meanStdDev + void cv::gpu::meanStdDev(const GpuMat& src, Scalar& mean, Scalar& stddev) { CV_Assert(src.type() == CV_8UC1); @@ -261,6 +241,9 @@ void cv::gpu::meanStdDev(const GpuMat& src, Scalar& mean, Scalar& stddev) nppSafeCall( nppiMean_StdDev_8u_C1R(src.ptr(), src.step, sz, mean.val, stddev.val) ); } +//////////////////////////////////////////////////////////////////////// +// norm + double cv::gpu::norm(const GpuMat& src1, int normType) { return norm(src1, GpuMat(src1.size(), src1.type(), Scalar::all(0.0)), normType); @@ -292,6 +275,9 @@ double cv::gpu::norm(const GpuMat& src1, const GpuMat& src2, int normType) return retVal[0]; } +//////////////////////////////////////////////////////////////////////// +// flip + void cv::gpu::flip(const GpuMat& src, GpuMat& dst, int flipCode) { CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4); @@ -316,50 +302,8 @@ void cv::gpu::flip(const GpuMat& src, GpuMat& dst, int flipCode) } } -void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, double fy, int interpolation) -{ - static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC, 0, NPPI_INTER_LANCZOS}; - - CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4); - CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC || interpolation == INTER_LANCZOS4); - - CV_Assert( src.size().area() > 0 ); - CV_Assert( !(dsize == Size()) || (fx > 0 && fy > 0) ); - - if( dsize == Size() ) - { - dsize = Size(saturate_cast(src.cols * fx), saturate_cast(src.rows * fy)); - } - else - { - fx = (double)dsize.width / src.cols; - fy = (double)dsize.height / src.rows; - } - - dst.create(dsize, src.type()); - - NppiSize srcsz; - srcsz.width = src.cols; - srcsz.height = src.rows; - NppiRect srcrect; - srcrect.x = srcrect.y = 0; - srcrect.width = src.cols; - srcrect.height = src.rows; - NppiSize dstsz; - dstsz.width = dst.cols; - dstsz.height = dst.rows; - - if (src.type() == CV_8UC1) - { - nppSafeCall( nppiResize_8u_C1R(src.ptr(), srcsz, src.step, srcrect, - dst.ptr(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) ); - } - else - { - nppSafeCall( nppiResize_8u_C4R(src.ptr(), srcsz, src.step, srcrect, - dst.ptr(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) ); - } -} +//////////////////////////////////////////////////////////////////////// +// sum Scalar cv::gpu::sum(const GpuMat& src) { @@ -383,6 +327,9 @@ Scalar cv::gpu::sum(const GpuMat& src) return res; } +//////////////////////////////////////////////////////////////////////// +// minMax + void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal) { CV_Assert(src.type() == CV_8UC1); @@ -402,232 +349,37 @@ void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal) *maxVal = max_res; } -void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value) -{ - CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1); - - dst.create(src.rows + top + bottom, src.cols + left + right, src.type()); - - NppiSize srcsz; - srcsz.width = src.cols; - srcsz.height = src.rows; - NppiSize dstsz; - dstsz.width = dst.cols; - dstsz.height = dst.rows; +//////////////////////////////////////////////////////////////////////// +// LUT - switch (src.type()) - { - case CV_8UC1: - { - Npp8u nVal = static_cast(value[0]); - nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr(), src.step, srcsz, - dst.ptr(), dst.step, dstsz, top, left, nVal) ); - break; - } - case CV_8UC4: - { - Npp8u nVal[] = {static_cast(value[0]), static_cast(value[1]), static_cast(value[2]), static_cast(value[3])}; - nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr(), src.step, srcsz, - dst.ptr(), dst.step, dstsz, top, left, nVal) ); - break; - } - case CV_32SC1: - { - Npp32s nVal = static_cast(value[0]); - nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr(), src.step, srcsz, - dst.ptr(), dst.step, dstsz, top, left, nVal) ); - break; - } - default: - CV_Assert(!"Unsupported source type"); - } -} - -namespace -{ - typedef NppStatus (*npp_warp_8u_t)(const Npp8u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp8u* pDst, - int dstStep, NppiRect dstRoi, const double coeffs[][3], - int interpolation); - typedef NppStatus (*npp_warp_16u_t)(const Npp16u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp16u* pDst, - int dstStep, NppiRect dstRoi, const double coeffs[][3], - int interpolation); - typedef NppStatus (*npp_warp_32s_t)(const Npp32s* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32s* pDst, - int dstStep, NppiRect dstRoi, const double coeffs[][3], - int interpolation); - typedef NppStatus (*npp_warp_32f_t)(const Npp32f* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32f* pDst, - int dstStep, NppiRect dstRoi, const double coeffs[][3], - int interpolation); - - void nppWarpCaller(const GpuMat& src, GpuMat& dst, double coeffs[][3], const Size& dsize, int flags, - npp_warp_8u_t npp_warp_8u[][2], npp_warp_16u_t npp_warp_16u[][2], - npp_warp_32s_t npp_warp_32s[][2], npp_warp_32f_t npp_warp_32f[][2]) +void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst) +{ + class LevelsInit { - static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC}; - - int interpolation = flags & INTER_MAX; - - CV_Assert((src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32S || src.depth() == CV_32F) && src.channels() != 2); - CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC); - - dst.create(dsize, src.type()); - - NppiSize srcsz; - srcsz.height = src.rows; - srcsz.width = src.cols; - NppiRect srcroi; - srcroi.x = srcroi.y = 0; - srcroi.height = src.rows; - srcroi.width = src.cols; - NppiRect dstroi; - dstroi.x = dstroi.y = 0; - dstroi.height = dst.rows; - dstroi.width = dst.cols; + public: + Npp32s pLevels[256]; - int warpInd = (flags & WARP_INVERSE_MAP) >> 4; - - switch (src.depth()) - { - case CV_8U: - nppSafeCall( npp_warp_8u[src.channels()][warpInd](src.ptr(), srcsz, src.step, srcroi, - dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); - break; - case CV_16U: - nppSafeCall( npp_warp_16u[src.channels()][warpInd](src.ptr(), srcsz, src.step, srcroi, - dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); - break; - case CV_32S: - nppSafeCall( npp_warp_32s[src.channels()][warpInd](src.ptr(), srcsz, src.step, srcroi, - dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); - break; - case CV_32F: - nppSafeCall( npp_warp_32f[src.channels()][warpInd](src.ptr(), srcsz, src.step, srcroi, - dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); - break; - default: - CV_Assert(!"Unsupported source type"); + LevelsInit() + { + for (int i = 0; i < 256; ++i) + pLevels[i] = i; } - } -} - -void cv::gpu::warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags) -{ - static npp_warp_8u_t npp_warpAffine_8u[][2] = - { - {0, 0}, - {nppiWarpAffine_8u_C1R, nppiWarpAffineBack_8u_C1R}, - {0, 0}, - {nppiWarpAffine_8u_C3R, nppiWarpAffineBack_8u_C3R}, - {nppiWarpAffine_8u_C4R, nppiWarpAffineBack_8u_C4R} - }; - static npp_warp_16u_t npp_warpAffine_16u[][2] = - { - {0, 0}, - {nppiWarpAffine_16u_C1R, nppiWarpAffineBack_16u_C1R}, - {0, 0}, - {nppiWarpAffine_16u_C3R, nppiWarpAffineBack_16u_C3R}, - {nppiWarpAffine_16u_C4R, nppiWarpAffineBack_16u_C4R} - }; - static npp_warp_32s_t npp_warpAffine_32s[][2] = - { - {0, 0}, - {nppiWarpAffine_32s_C1R, nppiWarpAffineBack_32s_C1R}, - {0, 0}, - {nppiWarpAffine_32s_C3R, nppiWarpAffineBack_32s_C3R}, - {nppiWarpAffine_32s_C4R, nppiWarpAffineBack_32s_C4R} - }; - static npp_warp_32f_t npp_warpAffine_32f[][2] = - { - {0, 0}, - {nppiWarpAffine_32f_C1R, nppiWarpAffineBack_32f_C1R}, - {0, 0}, - {nppiWarpAffine_32f_C3R, nppiWarpAffineBack_32f_C3R}, - {nppiWarpAffine_32f_C4R, nppiWarpAffineBack_32f_C4R} - }; - - CV_Assert(M.rows == 2 && M.cols == 3); - - double coeffs[2][3]; - Mat coeffsMat(2, 3, CV_64F, (void*)coeffs); - M.convertTo(coeffsMat, coeffsMat.type()); - - nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpAffine_8u, npp_warpAffine_16u, npp_warpAffine_32s, npp_warpAffine_32f); -} - -void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags) -{ - static npp_warp_8u_t npp_warpPerspective_8u[][2] = - { - {0, 0}, - {nppiWarpPerspective_8u_C1R, nppiWarpPerspectiveBack_8u_C1R}, - {0, 0}, - {nppiWarpPerspective_8u_C3R, nppiWarpPerspectiveBack_8u_C3R}, - {nppiWarpPerspective_8u_C4R, nppiWarpPerspectiveBack_8u_C4R} - }; - static npp_warp_16u_t npp_warpPerspective_16u[][2] = - { - {0, 0}, - {nppiWarpPerspective_16u_C1R, nppiWarpPerspectiveBack_16u_C1R}, - {0, 0}, - {nppiWarpPerspective_16u_C3R, nppiWarpPerspectiveBack_16u_C3R}, - {nppiWarpPerspective_16u_C4R, nppiWarpPerspectiveBack_16u_C4R} - }; - static npp_warp_32s_t npp_warpPerspective_32s[][2] = - { - {0, 0}, - {nppiWarpPerspective_32s_C1R, nppiWarpPerspectiveBack_32s_C1R}, - {0, 0}, - {nppiWarpPerspective_32s_C3R, nppiWarpPerspectiveBack_32s_C3R}, - {nppiWarpPerspective_32s_C4R, nppiWarpPerspectiveBack_32s_C4R} - }; - static npp_warp_32f_t npp_warpPerspective_32f[][2] = - { - {0, 0}, - {nppiWarpPerspective_32f_C1R, nppiWarpPerspectiveBack_32f_C1R}, - {0, 0}, - {nppiWarpPerspective_32f_C3R, nppiWarpPerspectiveBack_32f_C3R}, - {nppiWarpPerspective_32f_C4R, nppiWarpPerspectiveBack_32f_C4R} - }; + }; + static LevelsInit lvls; - CV_Assert(M.rows == 3 && M.cols == 3); + int cn = src.channels(); - double coeffs[3][3]; - Mat coeffsMat(3, 3, CV_64F, (void*)coeffs); - M.convertTo(coeffsMat, coeffsMat.type()); + CV_Assert(src.type() == CV_8UC1); + CV_Assert(lut.depth() == CV_32SC1 && lut.rows * lut.cols == 256 && lut.isContinuous()); - nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpPerspective_8u, npp_warpPerspective_16u, npp_warpPerspective_32s, npp_warpPerspective_32f); -} + dst.create(src.size(), src.type()); -void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation) -{ - static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC}; - - CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4); - CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC); - - dst.create(dsize, src.type()); - - NppiSize srcsz; - srcsz.height = src.rows; - srcsz.width = src.cols; - NppiRect srcroi; - srcroi.x = srcroi.y = 0; - srcroi.height = src.rows; - srcroi.width = src.cols; - NppiRect dstroi; - dstroi.x = dstroi.y = 0; - dstroi.height = dst.rows; - dstroi.width = dst.cols; + NppiSize sz; + sz.height = src.rows; + sz.width = src.cols; - if (src.type() == CV_8UC1) - { - nppSafeCall( nppiRotate_8u_C1R(src.ptr(), srcsz, src.step, srcroi, - dst.ptr(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) ); - } - else - { - nppSafeCall( nppiRotate_8u_C4R(src.ptr(), srcsz, src.step, srcroi, - dst.ptr(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) ); - } + nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr(), src.step, dst.ptr(), dst.step, sz, + lut.ptr(), lvls.pLevels, 256) ); } #endif /* !defined (HAVE_CUDA) */ \ No newline at end of file diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu index a53a4f8..6954d57 100644 --- a/modules/gpu/src/cuda/color.cu +++ b/modules/gpu/src/cuda/color.cu @@ -46,20 +46,30 @@ using namespace cv::gpu; #ifndef CV_DESCALE -#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n)) +#define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n)) #endif namespace imgproc { - template struct ColorChannel - { - }; + template struct TypeVec {}; + template<> struct TypeVec { typedef uchar1 vec_t; }; + template<> struct TypeVec { typedef uchar2 vec_t; }; + template<> struct TypeVec { typedef uchar3 vec_t; }; + template<> struct TypeVec { typedef uchar4 vec_t; }; + template<> struct TypeVec { typedef ushort1 vec_t; }; + template<> struct TypeVec { typedef ushort2 vec_t; }; + template<> struct TypeVec { typedef ushort3 vec_t; }; + template<> struct TypeVec { typedef ushort4 vec_t; }; + template<> struct TypeVec { typedef float1 vec_t; }; + template<> struct TypeVec { typedef float2 vec_t; }; + template<> struct TypeVec { typedef float3 vec_t; }; + template<> struct TypeVec { typedef float4 vec_t; }; + + template struct ColorChannel {}; template<> struct ColorChannel { typedef float worktype_f; - typedef uchar3 vec3_t; - typedef uchar4 vec4_t; static __device__ unsigned char max() { return UCHAR_MAX; } static __device__ unsigned char half() { return (unsigned char)(max()/2 + 1); } }; @@ -67,8 +77,6 @@ namespace imgproc template<> struct ColorChannel { typedef float worktype_f; - typedef ushort3 vec3_t; - typedef ushort4 vec4_t; static __device__ unsigned short max() { return USHRT_MAX; } static __device__ unsigned short half() { return (unsigned short)(max()/2 + 1); } }; @@ -76,94 +84,114 @@ namespace imgproc template<> struct ColorChannel { typedef float worktype_f; - typedef float3 vec3_t; - typedef float4 vec4_t; static __device__ float max() { return 1.f; } static __device__ float half() { return 0.5f; } - }; + }; } -////////////////// Various 3/4-channel to 3/4-channel RGB transformations ///////////////// +//////////////////////////////////////// SwapChannels ///////////////////////////////////// namespace imgproc { - template - __global__ void RGB2RGB_3_3(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx) - { - const int x = blockDim.x * blockIdx.x + threadIdx.x; - const int y = blockDim.y * blockIdx.y + threadIdx.y; + __constant__ int ccoeffs[4]; + + template + __global__ void swapChannels(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols) + { + typedef typename TypeVec::vec_t vec_t; + + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; if (y < rows && x < cols) { - const T* src = src_ + y * src_step + x * 3; - T* dst = dst_ + y * dst_step + x * 3; + vec_t src = *(const vec_t*)(src_ + y * src_step + x * CN); + vec_t dst; - T t0 = src[bidx], t1 = src[1], t2 = src[bidx ^ 2]; - dst[0] = t0; dst[1] = t1; dst[2] = t2; - } + const T* src_ptr = (const T*)(&src); + T* dst_ptr = (T*)(&dst); + + for (int i = 0; i < CN; ++i) + dst_ptr[i] = src_ptr[ccoeffs[i]]; + + *(vec_t*)(dst_ + y * dst_step + x * CN) = dst; } + } +} +namespace cv { namespace gpu { namespace improc +{ template - __global__ void RGB2RGB_4_3(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx) - { - typedef typename ColorChannel::vec4_t vec4_t; + void swapChannels_caller(const DevMem2D_& src, const DevMem2D_& dst, int cn, const int* coeffs, cudaStream_t stream) + { + dim3 threads(32, 8, 1); + dim3 grid(1, 1, 1); - const int x = blockDim.x * blockIdx.x + threadIdx.x; - const int y = blockDim.y * blockIdx.y + threadIdx.y; + grid.x = divUp(src.cols, threads.x); + grid.y = divUp(src.rows, threads.y); - if (y < rows && x < cols) - { - vec4_t src = *(vec4_t*)(src_ + y * src_step + (x << 2)); - T* dst = dst_ + y * dst_step + x * 3; + cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, cn * sizeof(int)) ); - T t0 = ((T*)(&src))[bidx], t1 = src.y, t2 = ((T*)(&src))[bidx ^ 2]; - dst[0] = t0; dst[1] = t1; dst[2] = t2; - } + switch (cn) + { + case 3: + imgproc::swapChannels<3><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); + break; + case 4: + imgproc::swapChannels<4><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); + break; + default: + cv::gpu::error("Unsupported channels count", __FILE__, __LINE__); + break; } - template - __global__ void RGB2RGB_3_4(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx) - { - typedef typename ColorChannel::vec4_t vec4_t; + if (stream == 0) + cudaSafeCall( cudaThreadSynchronize() ); + } - const int x = blockDim.x * blockIdx.x + threadIdx.x; - const int y = blockDim.y * blockIdx.y + threadIdx.y; + void swapChannels_gpu(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream) + { + swapChannels_caller(src, dst, cn, coeffs, stream); + } - if (y < rows && x < cols) - { - const T* src = src_ + y * src_step + x * 3; + void swapChannels_gpu(const DevMem2D_& src, const DevMem2D_& dst, int cn, const int* coeffs, cudaStream_t stream) + { + swapChannels_caller(src, dst, cn, coeffs, stream); + } - vec4_t dst; + void swapChannels_gpu(const DevMem2Df& src, const DevMem2Df& dst, int cn, const int* coeffs, cudaStream_t stream) + { + swapChannels_caller(src, dst, cn, coeffs, stream); + } +}}} - dst.x = src[bidx]; - dst.y = src[1]; - dst.z = src[bidx ^ 2]; - dst.w = ColorChannel::max(); - *(vec4_t*)(dst_ + y * dst_step + (x << 2)) = dst; - } - } +////////////////// Various 3/4-channel to 3/4-channel RGB transformations ///////////////// - template - __global__ void RGB2RGB_4_4(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx) - { - typedef typename ColorChannel::vec4_t vec4_t; +namespace imgproc +{ + template + __global__ void RGB2RGB(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx) + { + typedef typename TypeVec::vec_t src_t; + typedef typename TypeVec::vec_t dst_t; const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; if (y < rows && x < cols) { - vec4_t src = *(const vec4_t*)(src_ + y * src_step + (x << 2)); - vec4_t dst; + src_t src = *(const src_t*)(src_ + y * src_step + x * SRCCN); + dst_t dst; - dst.x = ((T*)(&src))[bidx]; + dst.x = ((const T*)(&src))[bidx]; dst.y = src.y; - dst.z = ((T*)(&src))[bidx ^ 2]; - dst.w = src.w; - - *(vec4_t*)(dst_ + y * dst_step + (x << 2)) = dst; - } + dst.z = ((const T*)(&src))[bidx ^ 2]; + if (DSTCN == 4) + ((T*)(&dst))[3] = ColorChannel::max(); + + *(dst_t*)(dst_ + y * dst_step + x * DSTCN) = dst; } + } } namespace cv { namespace gpu { namespace improc @@ -183,12 +211,15 @@ namespace cv { namespace gpu { namespace improc switch (srccn) { case 3: - imgproc::RGB2RGB_3_3<<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), - src.rows, src.cols, bidx); + { + int coeffs[] = {2, 1, 0}; + cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, 3 * sizeof(int)) ); + imgproc::swapChannels<3><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); break; + } case 4: - imgproc::RGB2RGB_4_3<<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), - src.rows, src.cols, bidx); + imgproc::RGB2RGB<4, 3><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), + src.rows, src.cols, bidx); break; default: cv::gpu::error("Unsupported channels count", __FILE__, __LINE__); @@ -199,13 +230,16 @@ namespace cv { namespace gpu { namespace improc switch (srccn) { case 3: - imgproc::RGB2RGB_3_4<<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), - src.rows, src.cols, bidx); + imgproc::RGB2RGB<3, 4><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), + src.rows, src.cols, bidx); break; case 4: - imgproc::RGB2RGB_4_4<<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), - src.rows, src.cols, bidx); + { + int coeffs[] = {2, 1, 0, 3}; + cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, 4 * sizeof(int)) ); + imgproc::swapChannels<4><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); break; + } default: cv::gpu::error("Unsupported channels count", __FILE__, __LINE__); break; @@ -319,8 +353,8 @@ namespace imgproc template __global__ void Gray2RGB_3(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols) { - const int x = blockDim.x * blockIdx.x + threadIdx.x; - const int y = blockDim.y * blockIdx.y + threadIdx.y; + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; if (y < rows && x < cols) { @@ -335,7 +369,7 @@ namespace imgproc template __global__ void Gray2RGB_4(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols) { - typedef typename ColorChannel::vec4_t vec4_t; + typedef typename TypeVec::vec_t vec4_t; const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; diff --git a/modules/gpu/src/filtering_npp.cpp b/modules/gpu/src/filtering_npp.cpp index 6c9bb1c..7a914f8 100644 --- a/modules/gpu/src/filtering_npp.cpp +++ b/modules/gpu/src/filtering_npp.cpp @@ -50,7 +50,7 @@ using namespace cv::gpu; void cv::gpu::erode( const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_nogpu(); } void cv::gpu::dilate( const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_nogpu(); } -void morphologyEx( const GpuMat&, GpuMat&, int, const Mat&, Point, int) { throw_nogpu(); } +void cv::gpu::morphologyEx( const GpuMat&, GpuMat&, int, const Mat&, Point, int) { throw_nogpu(); } #else @@ -132,7 +132,6 @@ void cv::gpu::morphologyEx( const GpuMat& src, GpuMat& dst, int op, const Mat& k temp = dst; dilate( src, temp, kernel, anchor, iterations); erode( temp, temp, kernel, anchor, iterations); - dst = temp - src; subtract(temp, src, dst); break; default: diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index d8685fa..8a207f2 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/modules/gpu/src/imgproc_gpu.cpp @@ -55,6 +55,14 @@ void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&) { throw_nog void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, const Stream&) { throw_nogpu(); } void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); } void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); } +double cv::gpu::threshold(const GpuMat&, GpuMat&, double) { throw_nogpu(); return 0.0; } +void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int) { throw_nogpu(); } +void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, const Scalar&) { throw_nogpu(); } +void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); } +void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); } +void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int) { throw_nogpu(); } +void cv::gpu::integral(GpuMat&, GpuMat&, GpuMat&) { throw_nogpu(); } +void cv::gpu::boxFilter(const GpuMat&, GpuMat&, Size, Point) { throw_nogpu(); } #else /* !defined (HAVE_CUDA) */ @@ -73,6 +81,10 @@ namespace cv { namespace gpu void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream); void reprojectImageTo3D_gpu(const DevMem2D_& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream); + void swapChannels_gpu(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream); + void swapChannels_gpu(const DevMem2D_& src, const DevMem2D_& dst, int cn, const int* coeffs, cudaStream_t stream); + void swapChannels_gpu(const DevMem2Df& src, const DevMem2Df& dst, int cn, const int* coeffs, cudaStream_t stream); + void RGB2RGB_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream); void RGB2RGB_gpu(const DevMem2D_& src, int srccn, const DevMem2D_& dst, int dstcn, int bidx, cudaStream_t stream); void RGB2RGB_gpu(const DevMem2Df& src, int srccn, const DevMem2Df& dst, int dstcn, int bidx, cudaStream_t stream); @@ -218,6 +230,10 @@ namespace if (dst.data != src.data) out = dst; + NppiSize nppsz; + nppsz.height = src.rows; + nppsz.width = src.cols; + switch (code) { case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR: @@ -305,6 +321,31 @@ namespace // CvtColorLoop(src, dst, Gray2RGB5x5(code == CV_GRAY2BGR565 ? 6 : 5)); // break; + case CV_RGB2YCrCb: + CV_Assert(scn == 3 && depth == CV_8U); + + out.create(sz, CV_MAKETYPE(depth, 3)); + + nppSafeCall( nppiRGBToYCbCr_8u_C3R(src.ptr(), src.step, out.ptr(), out.step, nppsz) ); + { + static int coeffs[] = {0, 2, 1}; + improc::swapChannels_gpu((DevMem2D)out, (DevMem2D)out, 3, coeffs, 0); + } + break; + + case CV_YCrCb2RGB: + CV_Assert(scn == 3 && depth == CV_8U); + + out.create(sz, CV_MAKETYPE(depth, 3)); + + { + static int coeffs[] = {0, 2, 1}; + GpuMat src1(src.size(), src.type()); + improc::swapChannels_gpu((DevMem2D)src, (DevMem2D)src1, 3, coeffs, 0); + nppSafeCall( nppiYCbCrToRGB_8u_C3R(src1.ptr(), src1.step, out.ptr(), out.step, nppsz) ); + } + break; + //case CV_BGR2YCrCb: case CV_RGB2YCrCb: //case CV_BGR2YUV: case CV_RGB2YUV: // { @@ -526,4 +567,366 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream)); } +//////////////////////////////////////////////////////////////////////// +// threshold + +double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh) +{ + CV_Assert(src.type() == CV_32FC1) + + dst.create( src.size(), src.type() ); + + NppiSize sz; + sz.width = src.cols; + sz.height = src.rows; + + nppSafeCall( nppiThreshold_32f_C1R(src.ptr(), src.step, + dst.ptr(), dst.step, sz, static_cast(thresh), NPP_CMP_GREATER) ); + + return thresh; +} + +//////////////////////////////////////////////////////////////////////// +// resize + +void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, double fy, int interpolation) +{ + static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC, 0, NPPI_INTER_LANCZOS}; + + CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4); + CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC || interpolation == INTER_LANCZOS4); + + CV_Assert( src.size().area() > 0 ); + CV_Assert( !(dsize == Size()) || (fx > 0 && fy > 0) ); + + if( dsize == Size() ) + { + dsize = Size(saturate_cast(src.cols * fx), saturate_cast(src.rows * fy)); + } + else + { + fx = (double)dsize.width / src.cols; + fy = (double)dsize.height / src.rows; + } + + dst.create(dsize, src.type()); + + NppiSize srcsz; + srcsz.width = src.cols; + srcsz.height = src.rows; + NppiRect srcrect; + srcrect.x = srcrect.y = 0; + srcrect.width = src.cols; + srcrect.height = src.rows; + NppiSize dstsz; + dstsz.width = dst.cols; + dstsz.height = dst.rows; + + if (src.type() == CV_8UC1) + { + nppSafeCall( nppiResize_8u_C1R(src.ptr(), srcsz, src.step, srcrect, + dst.ptr(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) ); + } + else + { + nppSafeCall( nppiResize_8u_C4R(src.ptr(), srcsz, src.step, srcrect, + dst.ptr(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) ); + } +} + +//////////////////////////////////////////////////////////////////////// +// copyMakeBorder + +void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value) +{ + CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1); + + dst.create(src.rows + top + bottom, src.cols + left + right, src.type()); + + NppiSize srcsz; + srcsz.width = src.cols; + srcsz.height = src.rows; + NppiSize dstsz; + dstsz.width = dst.cols; + dstsz.height = dst.rows; + + switch (src.type()) + { + case CV_8UC1: + { + Npp8u nVal = static_cast(value[0]); + nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr(), src.step, srcsz, + dst.ptr(), dst.step, dstsz, top, left, nVal) ); + break; + } + case CV_8UC4: + { + Npp8u nVal[] = {static_cast(value[0]), static_cast(value[1]), static_cast(value[2]), static_cast(value[3])}; + nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr(), src.step, srcsz, + dst.ptr(), dst.step, dstsz, top, left, nVal) ); + break; + } + case CV_32SC1: + { + Npp32s nVal = static_cast(value[0]); + nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr(), src.step, srcsz, + dst.ptr(), dst.step, dstsz, top, left, nVal) ); + break; + } + default: + CV_Assert(!"Unsupported source type"); + } +} + +//////////////////////////////////////////////////////////////////////// +// warp + +namespace +{ + typedef NppStatus (*npp_warp_8u_t)(const Npp8u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp8u* pDst, + int dstStep, NppiRect dstRoi, const double coeffs[][3], + int interpolation); + typedef NppStatus (*npp_warp_16u_t)(const Npp16u* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp16u* pDst, + int dstStep, NppiRect dstRoi, const double coeffs[][3], + int interpolation); + typedef NppStatus (*npp_warp_32s_t)(const Npp32s* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32s* pDst, + int dstStep, NppiRect dstRoi, const double coeffs[][3], + int interpolation); + typedef NppStatus (*npp_warp_32f_t)(const Npp32f* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, Npp32f* pDst, + int dstStep, NppiRect dstRoi, const double coeffs[][3], + int interpolation); + + void nppWarpCaller(const GpuMat& src, GpuMat& dst, double coeffs[][3], const Size& dsize, int flags, + npp_warp_8u_t npp_warp_8u[][2], npp_warp_16u_t npp_warp_16u[][2], + npp_warp_32s_t npp_warp_32s[][2], npp_warp_32f_t npp_warp_32f[][2]) + { + static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC}; + + int interpolation = flags & INTER_MAX; + + CV_Assert((src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32S || src.depth() == CV_32F) && src.channels() != 2); + CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC); + + dst.create(dsize, src.type()); + + NppiSize srcsz; + srcsz.height = src.rows; + srcsz.width = src.cols; + NppiRect srcroi; + srcroi.x = srcroi.y = 0; + srcroi.height = src.rows; + srcroi.width = src.cols; + NppiRect dstroi; + dstroi.x = dstroi.y = 0; + dstroi.height = dst.rows; + dstroi.width = dst.cols; + + int warpInd = (flags & WARP_INVERSE_MAP) >> 4; + + switch (src.depth()) + { + case CV_8U: + nppSafeCall( npp_warp_8u[src.channels()][warpInd](src.ptr(), srcsz, src.step, srcroi, + dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); + break; + case CV_16U: + nppSafeCall( npp_warp_16u[src.channels()][warpInd](src.ptr(), srcsz, src.step, srcroi, + dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); + break; + case CV_32S: + nppSafeCall( npp_warp_32s[src.channels()][warpInd](src.ptr(), srcsz, src.step, srcroi, + dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); + break; + case CV_32F: + nppSafeCall( npp_warp_32f[src.channels()][warpInd](src.ptr(), srcsz, src.step, srcroi, + dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); + break; + default: + CV_Assert(!"Unsupported source type"); + } + } +} + +void cv::gpu::warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags) +{ + static npp_warp_8u_t npp_warpAffine_8u[][2] = + { + {0, 0}, + {nppiWarpAffine_8u_C1R, nppiWarpAffineBack_8u_C1R}, + {0, 0}, + {nppiWarpAffine_8u_C3R, nppiWarpAffineBack_8u_C3R}, + {nppiWarpAffine_8u_C4R, nppiWarpAffineBack_8u_C4R} + }; + static npp_warp_16u_t npp_warpAffine_16u[][2] = + { + {0, 0}, + {nppiWarpAffine_16u_C1R, nppiWarpAffineBack_16u_C1R}, + {0, 0}, + {nppiWarpAffine_16u_C3R, nppiWarpAffineBack_16u_C3R}, + {nppiWarpAffine_16u_C4R, nppiWarpAffineBack_16u_C4R} + }; + static npp_warp_32s_t npp_warpAffine_32s[][2] = + { + {0, 0}, + {nppiWarpAffine_32s_C1R, nppiWarpAffineBack_32s_C1R}, + {0, 0}, + {nppiWarpAffine_32s_C3R, nppiWarpAffineBack_32s_C3R}, + {nppiWarpAffine_32s_C4R, nppiWarpAffineBack_32s_C4R} + }; + static npp_warp_32f_t npp_warpAffine_32f[][2] = + { + {0, 0}, + {nppiWarpAffine_32f_C1R, nppiWarpAffineBack_32f_C1R}, + {0, 0}, + {nppiWarpAffine_32f_C3R, nppiWarpAffineBack_32f_C3R}, + {nppiWarpAffine_32f_C4R, nppiWarpAffineBack_32f_C4R} + }; + + CV_Assert(M.rows == 2 && M.cols == 3); + + double coeffs[2][3]; + Mat coeffsMat(2, 3, CV_64F, (void*)coeffs); + M.convertTo(coeffsMat, coeffsMat.type()); + + nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpAffine_8u, npp_warpAffine_16u, npp_warpAffine_32s, npp_warpAffine_32f); +} + +void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags) +{ + static npp_warp_8u_t npp_warpPerspective_8u[][2] = + { + {0, 0}, + {nppiWarpPerspective_8u_C1R, nppiWarpPerspectiveBack_8u_C1R}, + {0, 0}, + {nppiWarpPerspective_8u_C3R, nppiWarpPerspectiveBack_8u_C3R}, + {nppiWarpPerspective_8u_C4R, nppiWarpPerspectiveBack_8u_C4R} + }; + static npp_warp_16u_t npp_warpPerspective_16u[][2] = + { + {0, 0}, + {nppiWarpPerspective_16u_C1R, nppiWarpPerspectiveBack_16u_C1R}, + {0, 0}, + {nppiWarpPerspective_16u_C3R, nppiWarpPerspectiveBack_16u_C3R}, + {nppiWarpPerspective_16u_C4R, nppiWarpPerspectiveBack_16u_C4R} + }; + static npp_warp_32s_t npp_warpPerspective_32s[][2] = + { + {0, 0}, + {nppiWarpPerspective_32s_C1R, nppiWarpPerspectiveBack_32s_C1R}, + {0, 0}, + {nppiWarpPerspective_32s_C3R, nppiWarpPerspectiveBack_32s_C3R}, + {nppiWarpPerspective_32s_C4R, nppiWarpPerspectiveBack_32s_C4R} + }; + static npp_warp_32f_t npp_warpPerspective_32f[][2] = + { + {0, 0}, + {nppiWarpPerspective_32f_C1R, nppiWarpPerspectiveBack_32f_C1R}, + {0, 0}, + {nppiWarpPerspective_32f_C3R, nppiWarpPerspectiveBack_32f_C3R}, + {nppiWarpPerspective_32f_C4R, nppiWarpPerspectiveBack_32f_C4R} + }; + + CV_Assert(M.rows == 3 && M.cols == 3); + + double coeffs[3][3]; + Mat coeffsMat(3, 3, CV_64F, (void*)coeffs); + M.convertTo(coeffsMat, coeffsMat.type()); + + nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpPerspective_8u, npp_warpPerspective_16u, npp_warpPerspective_32s, npp_warpPerspective_32f); +} + +//////////////////////////////////////////////////////////////////////// +// rotate + +void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation) +{ + static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC}; + + CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4); + CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC); + + dst.create(dsize, src.type()); + + NppiSize srcsz; + srcsz.height = src.rows; + srcsz.width = src.cols; + NppiRect srcroi; + srcroi.x = srcroi.y = 0; + srcroi.height = src.rows; + srcroi.width = src.cols; + NppiRect dstroi; + dstroi.x = dstroi.y = 0; + dstroi.height = dst.rows; + dstroi.width = dst.cols; + + if (src.type() == CV_8UC1) + { + nppSafeCall( nppiRotate_8u_C1R(src.ptr(), srcsz, src.step, srcroi, + dst.ptr(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) ); + } + else + { + nppSafeCall( nppiRotate_8u_C4R(src.ptr(), srcsz, src.step, srcroi, + dst.ptr(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) ); + } +} + +//////////////////////////////////////////////////////////////////////// +// integral + +void cv::gpu::integral(GpuMat& src, GpuMat& sum, GpuMat& sqsum) +{ + CV_Assert(src.type() == CV_8UC1); + + int w = src.cols + 1, h = src.rows + 1; + + sum.create(h, w, CV_32S); + sqsum.create(h, w, CV_32F); + + NppiSize sz; + sz.width = src.cols; + sz.height = src.rows; + + nppSafeCall( nppiSqrIntegral_8u32s32f_C1R(src.ptr(), src.step, sum.ptr(), + sum.step, sqsum.ptr(), sqsum.step, sz, 0, 0.0f, h) ); +} + +//////////////////////////////////////////////////////////////////////// +// boxFilter + +void cv::gpu::boxFilter(const GpuMat& src, GpuMat& dst, Size ksize, Point anchor) +{ + CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4); + CV_Assert(ksize.height == 3 || ksize.height == 5 || ksize.height == 7); + CV_Assert(ksize.height == ksize.width); + + if (anchor.x == -1) + anchor.x = 0; + if (anchor.y == -1) + anchor.y = 0; + + CV_Assert(anchor.x == 0 && anchor.y == 0); + + dst.create(src.size(), src.type()); + + NppiSize srcsz; + srcsz.height = src.rows; + srcsz.width = src.cols; + NppiSize masksz; + masksz.height = ksize.height; + masksz.width = ksize.width; + NppiPoint anc; + anc.x = anchor.x; + anc.y = anchor.y; + + if (src.type() == CV_8UC1) + { + nppSafeCall( nppiFilterBox_8u_C1R(src.ptr(), src.step, dst.ptr(), dst.step, srcsz, masksz, anc) ); + } + else + { + nppSafeCall( nppiFilterBox_8u_C4R(src.ptr(), src.step, dst.ptr(), dst.step, srcsz, masksz, anc) ); + } +} + #endif /* !defined (HAVE_CUDA) */ diff --git a/tests/gpu/src/arithm.cpp b/tests/gpu/src/arithm.cpp new file mode 100644 index 0000000..d4511a4 --- /dev/null +++ b/tests/gpu/src/arithm.cpp @@ -0,0 +1,696 @@ +/*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 +#include +#include +#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::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::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; diff --git a/tests/gpu/src/gputest.hpp b/tests/gpu/src/gputest.hpp index 0800e8f..57afd49 100644 --- a/tests/gpu/src/gputest.hpp +++ b/tests/gpu/src/gputest.hpp @@ -68,15 +68,15 @@ static inline bool check_and_treat_gpu_exception(const cv::Exception& e, CvTS* t 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: diff --git a/tests/gpu/src/imgproc_gpu.cpp b/tests/gpu/src/imgproc_gpu.cpp new file mode 100644 index 0000000..74076cf --- /dev/null +++ b/tests/gpu/src/imgproc_gpu.cpp @@ -0,0 +1,613 @@ +/*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 +#include +#include +#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::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::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; diff --git a/tests/gpu/src/meanshift.cpp b/tests/gpu/src/meanshift.cpp index 7e7b4b8..224db58 100644 --- a/tests/gpu/src/meanshift.cpp +++ b/tests/gpu/src/meanshift.cpp @@ -108,7 +108,7 @@ void CV_GpuMeanShiftTest::run(int) } if (maxDiff > 0) { - ts->printf(CvTS::CONSOLE, "\nMeanShift maxDiff = %d\n", maxDiff); + ts->printf(CvTS::LOG, "\nMeanShift maxDiff = %d\n", maxDiff); ts->set_failed_test_info(CvTS::FAIL_GENERIC); return; } diff --git a/tests/gpu/src/morf_filters.cpp b/tests/gpu/src/morf_filters.cpp index 5e33dff..d2d4d12 100644 --- a/tests/gpu/src/morf_filters.cpp +++ b/tests/gpu/src/morf_filters.cpp @@ -81,7 +81,7 @@ protected: if (res < std::numeric_limits::epsilon()) return CvTS::OK; - ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", res); + ts->printf(CvTS::LOG, "\nNorm: %f\n", res); return CvTS::FAIL_GENERIC; } }; diff --git a/tests/gpu/src/npp_image_arithm.cpp b/tests/gpu/src/npp_image_arithm.cpp deleted file mode 100644 index 11dd3d4..0000000 --- a/tests/gpu/src/npp_image_arithm.cpp +++ /dev/null @@ -1,878 +0,0 @@ -/*M/////////////////////////////////////////////////////////////////////////////////////// -// -// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. -// -// By downloading, copying, installing or using the software you agree to this license. -// If you do not agree to this license, do not download, install, -// copy or use the software. -// -// -// 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 -#include -#include -#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::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::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; diff --git a/tests/gpu/src/operator_async_call.cpp b/tests/gpu/src/operator_async_call.cpp index 2a9cf4f..821e367 100644 --- a/tests/gpu/src/operator_async_call.cpp +++ b/tests/gpu/src/operator_async_call.cpp @@ -131,7 +131,7 @@ bool CV_GpuMatAsyncCallTest::compare_matrix(cv::Mat & cpumat) return true; else { - ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret); + ts->printf(CvTS::LOG, "\nNorm: %f\n", ret); return false; } } diff --git a/tests/gpu/src/operator_convert_to.cpp b/tests/gpu/src/operator_convert_to.cpp index d9482e5..ab7b015 100644 --- a/tests/gpu/src/operator_convert_to.cpp +++ b/tests/gpu/src/operator_convert_to.cpp @@ -104,7 +104,7 @@ void CV_GpuMatOpConvertToTest::run(int /* start_from */) 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; diff --git a/tests/gpu/src/operator_copy_to.cpp b/tests/gpu/src/operator_copy_to.cpp index 9ada88b..f6ecf7d 100644 --- a/tests/gpu/src/operator_copy_to.cpp +++ b/tests/gpu/src/operator_copy_to.cpp @@ -127,7 +127,7 @@ bool CV_GpuMatOpCopyToTest::compare_matrix(cv::Mat & cpumat, gpu::GpuMat & gpuma return true; else { - ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret); + ts->printf(CvTS::LOG, "\nNorm: %f\n", ret); return false; } } diff --git a/tests/gpu/src/operator_set_to.cpp b/tests/gpu/src/operator_set_to.cpp index 6075c22..c208ea9 100644 --- a/tests/gpu/src/operator_set_to.cpp +++ b/tests/gpu/src/operator_set_to.cpp @@ -123,7 +123,7 @@ bool CV_GpuMatOpSetToTest::compare_matrix(cv::Mat & cpumat, gpu::GpuMat & gpumat return true; else { - ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret); + ts->printf(CvTS::LOG, "\nNorm: %f\n", ret); return false; } } diff --git a/tests/gpu/src/stereo_bm.cpp b/tests/gpu/src/stereo_bm.cpp index 2af46f4..682c7d1 100644 --- a/tests/gpu/src/stereo_bm.cpp +++ b/tests/gpu/src/stereo_bm.cpp @@ -81,7 +81,7 @@ void CV_GpuStereoBMTest::run(int ) 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; } diff --git a/tests/gpu/src/stereo_bm_async.cpp b/tests/gpu/src/stereo_bm_async.cpp index b178a0e..5d1255c 100644 --- a/tests/gpu/src/stereo_bm_async.cpp +++ b/tests/gpu/src/stereo_bm_async.cpp @@ -92,7 +92,7 @@ void CV_GpuMatAsyncCallStereoBMTest::run( int /* start_from */) 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; } diff --git a/tests/gpu/src/stereo_bp.cpp b/tests/gpu/src/stereo_bp.cpp index f0b2e4a..a6baf90 100644 --- a/tests/gpu/src/stereo_bp.cpp +++ b/tests/gpu/src/stereo_bp.cpp @@ -83,7 +83,7 @@ void CV_GpuStereoBPTest::run(int ) 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; } diff --git a/tests/gpu/src/stereo_csbp.cpp b/tests/gpu/src/stereo_csbp.cpp index e5613da..3d80c6d 100644 --- a/tests/gpu/src/stereo_csbp.cpp +++ b/tests/gpu/src/stereo_csbp.cpp @@ -83,7 +83,7 @@ void CV_GpuStereoCSBPTest::run(int ) 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; }