From a2f8817df1a336f3d69332e40d753386a8cfda13 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 20 Sep 2010 10:34:46 +0000 Subject: [PATCH] minor refactoring of GPU module and GPU tests added gpu compare version for CMP_NE operation --- modules/gpu/include/opencv2/gpu/gpu.hpp | 2 +- modules/gpu/src/arithm.cpp | 263 ++++++++++++---------- modules/gpu/src/cuda/matrix_operations.cu | 42 +++- modules/gpu/src/filtering_npp.cpp | 6 +- tests/gpu/src/meanshift.cpp | 42 ++-- tests/gpu/src/morf_filters.cpp | 6 +- tests/gpu/src/npp_image_arithm.cpp | 257 +++++++++++++-------- tests/gpu/src/operator_async_call.cpp | 13 +- tests/gpu/src/operator_convert_to.cpp | 6 +- tests/gpu/src/operator_copy_to.cpp | 19 +- tests/gpu/src/operator_set_to.cpp | 17 +- tests/gpu/src/stereo_bm.cpp | 32 ++- tests/gpu/src/stereo_bm_async.cpp | 39 ++-- tests/gpu/src/stereo_bp.cpp | 35 ++- tests/gpu/src/stereo_csbp.cpp | 36 ++- 15 files changed, 529 insertions(+), 286 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 3986cbba73..51ce18c0ea 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -364,7 +364,7 @@ namespace cv //! 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, double maxVal, int thresholdType); + CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh); //! compares elements of two arrays (c = a b) //! Now doesn't support CMP_NE. diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index e4b85b93a8..5fbebf40ed 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -57,7 +57,7 @@ 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, double, int) { throw_nogpu(); return 0.0; } +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(); } @@ -85,27 +85,14 @@ void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int) #else /* !defined (HAVE_CUDA) */ 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); - - typedef NppStatus (*npp_binary_func_8u_scale_t)(const Npp8u* pSrc1, int nSrc1Step, const Npp8u* pSrc2, int nSrc2Step, Npp8u* pDst, int nDstStep, - NppiSize oSizeROI, int nScaleFactor); - typedef NppStatus (*npp_binary_func_32f_t)(const Npp32f* pSrc1, int nSrc1Step, const Npp32f* pSrc2, int nSrc2Step, Npp32f* pDst, - int nDstStep, NppiSize oSizeROI); +{ + typedef NppStatus (*npp_arithm_8u_t)(const Npp8u* pSrc1, int nSrc1Step, const Npp8u* pSrc2, int nSrc2Step, Npp8u* pDst, int nDstStep, + NppiSize oSizeROI, int nScaleFactor); + typedef NppStatus (*npp_arithm_32f_t)(const Npp32f* pSrc1, int nSrc1Step, const Npp32f* pSrc2, int nSrc2Step, Npp32f* pDst, + int nDstStep, NppiSize oSizeROI); void nppFuncCaller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, - npp_binary_func_8u_scale_t npp_func_8uc1, npp_binary_func_8u_scale_t npp_func_8uc4, npp_binary_func_32f_t npp_func_32fc1) + npp_arithm_8u_t npp_func_8uc1, npp_arithm_8u_t npp_func_8uc4, npp_arithm_32f_t npp_func_32fc1) { CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type()); @@ -117,27 +104,26 @@ namespace sz.width = src1.cols; sz.height = src1.rows; - if (src1.depth() == CV_8U) - { - if (src1.channels() == 1) - { - nppSafeCall( npp_func_8uc1((const Npp8u*)src1.ptr(), src1.step, - (const Npp8u*)src2.ptr(), src2.step, - (Npp8u*)dst.ptr(), dst.step, sz, 0) ); - } - else - { - nppSafeCall( npp_func_8uc4((const Npp8u*)src1.ptr(), src1.step, - (const Npp8u*)src2.ptr(), src2.step, - (Npp8u*)dst.ptr(), dst.step, sz, 0) ); - } - } - else //if (src1.depth() == CV_32F) - { - nppSafeCall( npp_func_32fc1((const Npp32f*)src1.ptr(), src1.step, - (const Npp32f*)src2.ptr(), src2.step, - (Npp32f*)dst.ptr(), dst.step, sz) ); - } + switch (src1.type()) + { + case CV_8UC1: + nppSafeCall( npp_func_8uc1(src1.ptr(), src1.step, + src2.ptr(), src2.step, + dst.ptr(), dst.step, sz, 0) ); + break; + case CV_8UC4: + nppSafeCall( npp_func_8uc4(src1.ptr(), src1.step, + src2.ptr(), src2.step, + dst.ptr(), dst.step, sz, 0) ); + break; + case CV_32FC1: + nppSafeCall( npp_func_32fc1(src1.ptr(), src1.step, + src2.ptr(), src2.step, + dst.ptr(), dst.step, sz) ); + break; + default: + CV_Assert(!"Unsupported source type"); + } } } @@ -171,14 +157,14 @@ void cv::gpu::transpose(const GpuMat& src, GpuMat& dst) sz.width = src.cols; sz.height = src.rows; - nppSafeCall( nppiTranspose_8u_C1R((const Npp8u*)src.ptr(), src.step, (Npp8u*)dst.ptr(), dst.step, sz) ); + nppSafeCall( nppiTranspose_8u_C1R(src.ptr(), src.step, dst.ptr(), dst.step, sz) ); } void cv::gpu::absdiff(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) { CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type()); - CV_Assert((src1.depth() == CV_8U || src1.depth() == CV_32F) && src1.channels() == 1); + CV_Assert(src1.type() == CV_8UC1 || src1.type() == CV_32FC1); dst.create( src1.size(), src1.type() ); @@ -186,23 +172,23 @@ void cv::gpu::absdiff(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) sz.width = src1.cols; sz.height = src1.rows; - if (src1.depth() == CV_8U) + if (src1.type() == CV_8UC1) { - nppSafeCall( nppiAbsDiff_8u_C1R((const Npp8u*)src1.ptr(), src1.step, - (const Npp8u*)src2.ptr(), src2.step, - (Npp8u*)dst.ptr(), dst.step, sz) ); + nppSafeCall( nppiAbsDiff_8u_C1R(src1.ptr(), src1.step, + src2.ptr(), src2.step, + dst.ptr(), dst.step, sz) ); } - else //if (src1.depth() == CV_32F) + else { - nppSafeCall( nppiAbsDiff_32f_C1R((const Npp32f*)src1.ptr(), src1.step, - (const Npp32f*)src2.ptr(), src2.step, - (Npp32f*)dst.ptr(), dst.step, sz) ); + nppSafeCall( nppiAbsDiff_32f_C1R(src1.ptr(), src1.step, + src2.ptr(), src2.step, + dst.ptr(), dst.step, sz) ); } } -double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh, double /*maxVal*/, int thresholdType) +double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh) { - CV_Assert(src.type() == CV_32FC1 && thresholdType == THRESH_TRUNC); + CV_Assert(src.type() == CV_32FC1) dst.create( src.size(), src.type() ); @@ -210,17 +196,23 @@ double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh, double sz.width = src.cols; sz.height = src.rows; - nppSafeCall( nppiThreshold_32f_C1R((const Npp32f*)src.ptr(), src.step, - (Npp32f*)dst.ptr(), dst.step, sz, (Npp32f)thresh, NPP_CMP_GREATER) ); + nppSafeCall( nppiThreshold_32f_C1R(src.ptr(), src.step, + dst.ptr(), dst.step, sz, static_cast(thresh), NPP_CMP_GREATER) ); return thresh; } +namespace cv { namespace gpu { namespace matrix_operations +{ + void compare_ne_8u(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst); + void compare_ne_32f(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst); +}}} + void cv::gpu::compare(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, int cmpop) { CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type()); - CV_Assert((src1.type() == CV_8UC4 || src1.type() == CV_32FC1) && cmpop != CMP_NE); + CV_Assert(src1.type() == CV_8UC4 || src1.type() == CV_32FC1); dst.create( src1.size(), CV_8UC1 ); @@ -230,17 +222,31 @@ void cv::gpu::compare(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, int c sz.width = src1.cols; sz.height = src1.rows; - if (src1.depth() == CV_8U) + if (src1.type() == CV_8UC4) { - nppSafeCall( nppiCompare_8u_C4R((const Npp8u*)src1.ptr(), src1.step, - (const Npp8u*)src2.ptr(), src2.step, - (Npp8u*)dst.ptr(), dst.step, sz, nppCmpOp[cmpop]) ); + if (cmpop != CMP_NE) + { + nppSafeCall( nppiCompare_8u_C4R(src1.ptr(), src1.step, + src2.ptr(), src2.step, + dst.ptr(), dst.step, sz, nppCmpOp[cmpop]) ); + } + else + { + matrix_operations::compare_ne_8u(src1, src2, dst); + } } - else //if (src1.depth() == CV_32F) + else { - nppSafeCall( nppiCompare_32f_C1R((const Npp32f*)src1.ptr(), src1.step, - (const Npp32f*)src2.ptr(), src2.step, - (Npp8u*)dst.ptr(), dst.step, sz, nppCmpOp[cmpop]) ); + if (cmpop != CMP_NE) + { + nppSafeCall( nppiCompare_32f_C1R(src1.ptr(), src1.step, + src2.ptr(), src2.step, + dst.ptr(), dst.step, sz, nppCmpOp[cmpop]) ); + } + else + { + matrix_operations::compare_ne_32f(src1, src2, dst); + } } } @@ -252,7 +258,7 @@ void cv::gpu::meanStdDev(const GpuMat& src, Scalar& mean, Scalar& stddev) sz.width = src.cols; sz.height = src.rows; - nppSafeCall( nppiMean_StdDev_8u_C1R((const Npp8u*)src.ptr(), src.step, sz, mean.val, stddev.val) ); + nppSafeCall( nppiMean_StdDev_8u_C1R(src.ptr(), src.step, sz, mean.val, stddev.val) ); } double cv::gpu::norm(const GpuMat& src1, int normType) @@ -264,7 +270,8 @@ double cv::gpu::norm(const GpuMat& src1, const GpuMat& src2, int normType) { CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type()); - CV_Assert((src1.type() == CV_8UC1) && (normType == NORM_INF || normType == NORM_L1 || normType == NORM_L2)); + CV_Assert(src1.type() == CV_8UC1); + CV_Assert(normType == NORM_INF || normType == NORM_L1 || normType == NORM_L2); typedef NppStatus (*npp_norm_diff_func_t)(const Npp8u* pSrc1, int nSrcStep1, const Npp8u* pSrc2, int nSrcStep2, NppiSize oSizeROI, Npp64f* pRetVal); @@ -278,8 +285,8 @@ double cv::gpu::norm(const GpuMat& src1, const GpuMat& src2, int normType) int funcIdx = normType >> 1; Scalar retVal; - nppSafeCall( npp_norm_diff_func[funcIdx]((const Npp8u*)src1.ptr(), src1.step, - (const Npp8u*)src2.ptr(), src2.step, + nppSafeCall( npp_norm_diff_func[funcIdx](src1.ptr(), src1.step, + src2.ptr(), src2.step, sz, retVal.val) ); return retVal[0]; @@ -295,16 +302,16 @@ void cv::gpu::flip(const GpuMat& src, GpuMat& dst, int flipCode) sz.width = src.cols; sz.height = src.rows; - if (src.channels() == 1) + if (src.type() == CV_8UC1) { - nppSafeCall( nppiMirror_8u_C1R((const Npp8u*)src.ptr(), src.step, - (Npp8u*)dst.ptr(), dst.step, sz, + nppSafeCall( nppiMirror_8u_C1R(src.ptr(), src.step, + dst.ptr(), dst.step, sz, (flipCode == 0 ? NPP_HORIZONTAL_AXIS : (flipCode > 0 ? NPP_VERTICAL_AXIS : NPP_BOTH_AXIS))) ); } else { - nppSafeCall( nppiMirror_8u_C4R((const Npp8u*)src.ptr(), src.step, - (Npp8u*)dst.ptr(), dst.step, sz, + nppSafeCall( nppiMirror_8u_C4R(src.ptr(), src.step, + dst.ptr(), dst.step, sz, (flipCode == 0 ? NPP_HORIZONTAL_AXIS : (flipCode > 0 ? NPP_VERTICAL_AXIS : NPP_BOTH_AXIS))) ); } } @@ -313,11 +320,12 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub { 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) && - (interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC || interpolation == INTER_LANCZOS4)); + 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)); @@ -327,6 +335,7 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub fx = (double)dsize.width / src.cols; fy = (double)dsize.height / src.rows; } + dst.create(dsize, src.type()); NppiSize srcsz; @@ -340,15 +349,15 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub dstsz.width = dst.cols; dstsz.height = dst.rows; - if (src.channels() == 1) + if (src.type() == CV_8UC1) { - nppSafeCall( nppiResize_8u_C1R((const Npp8u*)src.ptr(), srcsz, src.step, srcrect, - (Npp8u*)dst.ptr(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) ); + 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((const Npp8u*)src.ptr(), srcsz, src.step, srcrect, - (Npp8u*)dst.ptr(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) ); + nppSafeCall( nppiResize_8u_C4R(src.ptr(), srcsz, src.step, srcrect, + dst.ptr(), dst.step, dstsz, fx, fy, npp_inter[interpolation]) ); } } @@ -362,13 +371,13 @@ Scalar cv::gpu::sum(const GpuMat& src) sz.width = src.cols; sz.height = src.rows; - if (src.channels() == 1) + if (src.type() == CV_8UC1) { - nppSafeCall( nppiSum_8u_C1R((const Npp8u*)src.ptr(), src.step, sz, res.val) ); + nppSafeCall( nppiSum_8u_C1R(src.ptr(), src.step, sz, res.val) ); } else { - nppSafeCall( nppiSum_8u_C4R((const Npp8u*)src.ptr(), src.step, sz, res.val) ); + nppSafeCall( nppiSum_8u_C4R(src.ptr(), src.step, sz, res.val) ); } return res; @@ -384,7 +393,7 @@ void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal) Npp8u min_res, max_res; - nppSafeCall( nppiMinMax_8u_C1R((const Npp8u*)src.ptr(), src.step, sz, &min_res, &max_res) ); + nppSafeCall( nppiMinMax_8u_C1R(src.ptr(), src.step, sz, &min_res, &max_res) ); if (minVal) *minVal = min_res; @@ -406,31 +415,49 @@ void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom dstsz.width = dst.cols; dstsz.height = dst.rows; - if (src.depth() == CV_8U) - { - if (src.channels() == 1) + switch (src.type()) + { + case CV_8UC1: { - Npp8u nVal = (Npp8u)value[0]; - nppSafeCall( nppiCopyConstBorder_8u_C1R((const Npp8u*)src.ptr(), src.step, srcsz, - (Npp8u*)dst.ptr(), dst.step, dstsz, top, left, nVal) ); + Npp8u nVal = static_cast(value[0]); + nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr(), src.step, srcsz, + dst.ptr(), dst.step, dstsz, top, left, nVal) ); + break; } - else + case CV_8UC4: { - Npp8u nVal[] = {(Npp8u)value[0], (Npp8u)value[1], (Npp8u)value[2], (Npp8u)value[3]}; - nppSafeCall( nppiCopyConstBorder_8u_C4R((const Npp8u*)src.ptr(), src.step, srcsz, - (Npp8u*)dst.ptr(), dst.step, dstsz, top, left, nVal) ); - } - } - else //if (src.depth() == CV_32S) - { - Npp32s nVal = (Npp32s)value[0]; - nppSafeCall( nppiCopyConstBorder_32s_C1R((const Npp32s*)src.ptr(), src.step, srcsz, - (Npp32s*)dst.ptr(), dst.step, dstsz, top, left, nVal) ); - } + 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]) @@ -461,20 +488,20 @@ namespace switch (src.depth()) { case CV_8U: - nppSafeCall( npp_warp_8u[src.channels()][warpInd]((const Npp8u*)src.ptr(), srcsz, src.step, srcroi, - (Npp8u*)dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); + 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]((const Npp16u*)src.ptr(), srcsz, src.step, srcroi, - (Npp16u*)dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); + 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_32SC1: - nppSafeCall( npp_warp_32s[src.channels()][warpInd]((const Npp32s*)src.ptr(), srcsz, src.step, srcroi, - (Npp32s*)dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); + 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_32FC1: - nppSafeCall( npp_warp_32f[src.channels()][warpInd]((const Npp32f*)src.ptr(), srcsz, src.step, srcroi, - (Npp32f*)dst.ptr(), dst.step, dstroi, coeffs, npp_inter[interpolation]) ); + 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"); @@ -591,15 +618,15 @@ void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, d dstroi.height = dst.rows; dstroi.width = dst.cols; - if (src.channels() == 1) + if (src.type() == CV_8UC1) { - nppSafeCall( nppiRotate_8u_C1R((const Npp8u*)src.ptr(), srcsz, src.step, srcroi, - (Npp8u*)dst.ptr(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) ); + 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((const Npp8u*)src.ptr(), srcsz, src.step, srcroi, - (Npp8u*)dst.ptr(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) ); + nppSafeCall( nppiRotate_8u_C4R(src.ptr(), srcsz, src.step, srcroi, + dst.ptr(), dst.step, dstroi, angle, xShift, yShift, npp_inter[interpolation]) ); } } diff --git a/modules/gpu/src/cuda/matrix_operations.cu b/modules/gpu/src/cuda/matrix_operations.cu index e16e5533b2..0b791fa72d 100644 --- a/modules/gpu/src/cuda/matrix_operations.cu +++ b/modules/gpu/src/cuda/matrix_operations.cu @@ -255,6 +255,24 @@ namespace mat_operators } } + /////////////////////////////////////////////////////////////////////////// + /////////////////////////////// compare_ne //////////////////////////////// + /////////////////////////////////////////////////////////////////////////// + + template + __global__ void kernel_compare_ne(uchar* src1, size_t src1_step, uchar* src2, size_t src2_step, uchar* dst, size_t dst_step, int cols, int rows) + { + const size_t x = threadIdx.x + blockIdx.x * blockDim.x; + const size_t y = threadIdx.y + blockIdx.y * blockDim.y; + + if (x < cols && y < rows) + { + T src1_pix = ((T*)(src1 + y * src1_step))[x]; + T src2_pix = ((T*)(src2 + y * src2_step))[x]; + uchar res = (uchar)(src1_pix != src2_pix) * 255; + ((dst + y * dst_step))[x] = res; + } + } } // namespace mat_operators namespace cv @@ -460,6 +478,28 @@ namespace cv cv::gpu::error("Unsupported convert operation", __FILE__, __LINE__); func(src, dst, src.cols * channels, src.rows, alpha, beta, stream); } - } // namespace impl + + /////////////////////////////////////////////////////////////////////////// + /////////////////////////////// compare_ne //////////////////////////////// + /////////////////////////////////////////////////////////////////////////// + + void compare_ne_8u(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst) + { + dim3 block(32, 8); + dim3 grid(divUp(src1.cols, block.x), divUp(src1.rows, block.y)); + + mat_operators::kernel_compare_ne<<>>(src1.ptr, src1.step, src2.ptr, src2.step, dst.ptr, dst.step, src1.cols, src1.rows); + cudaSafeCall( cudaThreadSynchronize() ); + } + + void compare_ne_32f(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst) + { + dim3 block(32, 8); + dim3 grid(divUp(src1.cols, block.x), divUp(src1.rows, block.y)); + + mat_operators::kernel_compare_ne<<>>(src1.ptr, src1.step, src2.ptr, src2.step, dst.ptr, dst.step, src1.cols, src1.rows); + cudaSafeCall( cudaThreadSynchronize() ); + } + } // namespace matrix_operations } // namespace gpu } // namespace cv diff --git a/modules/gpu/src/filtering_npp.cpp b/modules/gpu/src/filtering_npp.cpp index 869e3307bc..6c9bb1cea7 100644 --- a/modules/gpu/src/filtering_npp.cpp +++ b/modules/gpu/src/filtering_npp.cpp @@ -48,9 +48,9 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) -void cv::gpu::erode( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations) { throw_nogpu(); } -void cv::gpu::dilate( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations) { throw_nogpu(); } -void morphologyEx( const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, Point anchor, int iterations) { throw_nogpu(); } +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(); } #else diff --git a/tests/gpu/src/meanshift.cpp b/tests/gpu/src/meanshift.cpp index e5443c7db3..7e7b4b81c1 100644 --- a/tests/gpu/src/meanshift.cpp +++ b/tests/gpu/src/meanshift.cpp @@ -60,21 +60,23 @@ CV_GpuMeanShiftTest::CV_GpuMeanShiftTest(): CvTest( "GPU-MeanShift", "MeanShift" void CV_GpuMeanShiftTest::run(int) { - int spatialRad = 30; - int colorRad = 30; + int spatialRad = 30; + int colorRad = 30; - cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png"); - cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result.png"); + cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png"); + cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result.png"); - if (img.empty() || img_template.empty()) - { - ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA); - return; - } + if (img.empty() || img_template.empty()) + { + ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA); + return; + } - cv::Mat rgba; - cvtColor(img, rgba, CV_BGR2BGRA); + cv::Mat rgba; + cvtColor(img, rgba, CV_BGR2BGRA); + try + { cv::gpu::GpuMat res; cv::gpu::meanShiftFiltering( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad ); if (res.type() != CV_8UC4) @@ -98,15 +100,27 @@ void CV_GpuMeanShiftTest::run(int) { const uchar& ch1 = res_line[result.channels()*i + k]; const uchar& ch2 = ref_line[img_template.channels()*i + k]; - uchar diff = abs(ch1 - ch2); + uchar diff = static_cast(abs(ch1 - ch2)); if (maxDiff < diff) maxDiff = diff; } } } if (maxDiff > 0) + { ts->printf(CvTS::CONSOLE, "\nMeanShift maxDiff = %d\n", maxDiff); - ts->set_failed_test_info((maxDiff == 0) ? CvTS::OK : CvTS::FAIL_GENERIC); + ts->set_failed_test_info(CvTS::FAIL_GENERIC); + return; + } + } + catch(const cv::Exception& e) + { + if (!check_and_treat_gpu_exception(e, ts)) + throw; + return; + } + + ts->set_failed_test_info(CvTS::OK); } -CV_GpuMeanShiftTest CV_GpuMeanShift_test; \ No newline at end of file +CV_GpuMeanShiftTest CV_GpuMeanShift_test; diff --git a/tests/gpu/src/morf_filters.cpp b/tests/gpu/src/morf_filters.cpp index a14ac49236..5e33dff66e 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::LOG, "\nNorm: %f\n", res); + ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", res); return CvTS::FAIL_GENERIC; } }; @@ -116,7 +116,8 @@ void CV_GpuNppMorphogyTest::run( int ) catch(const cv::Exception& e) { if (!check_and_treat_gpu_exception(e, ts)) - throw; + throw; + return; } ts->set_failed_test_info(CvTS::OK); @@ -174,7 +175,6 @@ protected: CV_GpuDilateTest CV_GpuDilate_test; - //////////////////////////////////////////////////////////////////////////////// // Dilate class CV_GpuMorphExTest : public CV_GpuNppMorphogyTest diff --git a/tests/gpu/src/npp_image_arithm.cpp b/tests/gpu/src/npp_image_arithm.cpp index ec28361932..11dd3d40c4 100644 --- a/tests/gpu/src/npp_image_arithm.cpp +++ b/tests/gpu/src/npp_image_arithm.cpp @@ -132,7 +132,7 @@ int CV_GpuNppImageArithmTest::CheckNorm(const Mat& m1, const Mat& m2) } else { - ts->printf(CvTS::LOG, "\nNorm: %f\n", ret); + ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret); return CvTS::FAIL_GENERIC; } } @@ -154,7 +154,7 @@ int CV_GpuNppImageArithmTest::CheckNorm(double d1, double d2) } else { - ts->printf(CvTS::LOG, "\nNorm: %f\n", ret); + ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", ret); return CvTS::FAIL_GENERIC; } } @@ -165,8 +165,14 @@ void CV_GpuNppImageArithmTest::run( int ) //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::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()) { @@ -174,32 +180,41 @@ void CV_GpuNppImageArithmTest::run( int ) return; } - //run tests - int testResult = test8UC1(img_l, img_r); - if (testResult != CvTS::OK) + try { - ts->set_failed_test_info(testResult); - return; - } + //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 = 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 = 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) + testResult = test32FC1(img_l, img_r); + if (testResult != CvTS::OK) + { + ts->set_failed_test_info(testResult); + return; + } + } + catch(const cv::Exception& e) { - ts->set_failed_test_info(testResult); + if (!check_and_treat_gpu_exception(e, ts)) + throw; return; } @@ -423,15 +438,15 @@ int CV_GpuNppImageThresholdTest::test( const Mat& cpu1, const Mat& ) if (cpu1.type() != CV_32FC1) return CvTS::OK; - const double thresh = 0.5; - const double maxval = 0.0; + cv::RNG rng(*ts->get_rng()); + const double thresh = rng; cv::Mat cpuRes; - cv::threshold(cpu1, cpuRes, thresh, maxval, THRESH_TRUNC); + cv::threshold(cpu1, cpuRes, thresh, 0.0, THRESH_TRUNC); GpuMat gpu1(cpu1); GpuMat gpuRes; - cv::gpu::threshold(gpu1, gpuRes, thresh, maxval, THRESH_TRUNC); + cv::gpu::threshold(gpu1, gpuRes, thresh); return CheckNorm(cpuRes, gpuRes); } @@ -458,15 +473,30 @@ int CV_GpuNppImageCompareTest::test( const Mat& cpu1, const Mat& cpu2 ) if (cpu1.type() != CV_32FC1) return CvTS::OK; - cv::Mat cpuRes; - cv::compare(cpu1, cpu2, cpuRes, CMP_GT); + 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); - GpuMat gpu1(cpu1); - GpuMat gpu2(cpu2); - GpuMat gpuRes; - cv::gpu::compare(gpu1, gpu2, gpuRes, CMP_GT); + int test_res = CvTS::OK; - return CheckNorm(cpuRes, gpuRes); + 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; @@ -525,19 +555,28 @@ int CV_GpuNppImageNormTest::test( const Mat& cpu1, const Mat& cpu2 ) if (cpu1.type() != CV_8UC1) return CvTS::OK; - double cpu_norm_inf = cv::norm(cpu1, cpu2, NORM_INF); - double cpu_norm_L1 = cv::norm(cpu1, cpu2, NORM_L1); - double cpu_norm_L2 = cv::norm(cpu1, cpu2, NORM_L2); + 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); - GpuMat gpu1(cpu1); - GpuMat gpu2(cpu2); - double gpu_norm_inf = cv::gpu::norm(gpu1, gpu2, NORM_INF); - double gpu_norm_L1 = cv::gpu::norm(gpu1, gpu2, NORM_L1); - double gpu_norm_L2 = cv::gpu::norm(gpu1, gpu2, NORM_L2); + int test_res = CvTS::OK; - return (CheckNorm(cpu_norm_inf, gpu_norm_inf) == CvTS::OK - && CheckNorm(cpu_norm_L1, gpu_norm_L1) == CvTS::OK - && CheckNorm(cpu_norm_L2, gpu_norm_L2) == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC; + 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; @@ -562,20 +601,29 @@ int CV_GpuNppImageFlipTest::test( const Mat& cpu1, const Mat& ) if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4) return CvTS::OK; - Mat cpux, cpuy, cpub; - cv::flip(cpu1, cpux, 0); - cv::flip(cpu1, cpuy, 1); - cv::flip(cpu1, cpub, -1); + int flip_codes[] = {0, 1, -1}; + const char* flip_axis[] = {"X", "Y", "Both"}; + int flip_codes_num = sizeof(flip_codes) / sizeof(int); - GpuMat gpu1(cpu1); - GpuMat gpux, gpuy, gpub; - cv::gpu::flip(gpu1, gpux, 0); - cv::gpu::flip(gpu1, gpuy, 1); - cv::gpu::flip(gpu1, gpub, -1); + int test_res = CvTS::OK; - return (CheckNorm(cpux, gpux) == CvTS::OK && - CheckNorm(cpuy, gpuy) == CvTS::OK && - CheckNorm(cpub, gpub) == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC; + 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; @@ -600,25 +648,28 @@ int CV_GpuNppImageResizeTest::test( const Mat& cpu1, const Mat& ) if (cpu1.type() != CV_8UC1 && cpu1.type() != CV_8UC4) return CvTS::OK; - Mat cpunn, cpulin, cpucub, cpulanc; - cv::resize(cpu1, cpunn, Size(), 0.5, 0.5, INTER_NEAREST); - cv::resize(cpu1, cpulin, Size(), 0.5, 0.5, INTER_LINEAR); - cv::resize(cpu1, cpucub, Size(), 0.5, 0.5, INTER_CUBIC); - cv::resize(cpu1, cpulanc, Size(), 0.5, 0.5, INTER_LANCZOS4); + 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); - GpuMat gpu1(cpu1); - GpuMat gpunn, gpulin, gpucub, gpulanc; - cv::gpu::resize(gpu1, gpunn, Size(), 0.5, 0.5, INTER_NEAREST); - cv::gpu::resize(gpu1, gpulin, Size(), 0.5, 0.5, INTER_LINEAR); - cv::gpu::resize(gpu1, gpucub, Size(), 0.5, 0.5, INTER_CUBIC); - cv::gpu::resize(gpu1, gpulanc, Size(), 0.5, 0.5, INTER_LANCZOS4); + int test_res = CvTS::OK; - int nnres =CheckNorm(cpunn, gpunn); - int linres = CheckNorm(cpulin, gpulin); - int cubres = CheckNorm(cpucub, gpucub); - int lancres = CheckNorm(cpulanc, gpulanc); + for (int i = 0; i < interpolations_num; ++i) + { + Mat cpu_res; + cv::resize(cpu1, cpu_res, Size(), 0.5, 0.5, interpolations[i]); - return (nnres == CvTS::OK && linres == CvTS::OK && cubres == CvTS::OK && lancres == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC; + 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; @@ -744,14 +795,29 @@ int CV_GpuNppImageWarpAffineTest::test( const Mat& cpu1, const Mat& ) if (cpu1.type() == CV_32SC1) return CvTS::OK; - Mat cpudst; - cv::warpAffine(cpu1, cpudst, M, cpu1.size(), INTER_CUBIC | WARP_INVERSE_MAP); + 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); - GpuMat gpu1(cpu1); - GpuMat gpudst; - cv::gpu::warpAffine(gpu1, gpudst, M, gpu1.size(), INTER_CUBIC | WARP_INVERSE_MAP); + int test_res = CvTS::OK; - return CheckNorm(cpudst, gpudst); + 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; @@ -784,14 +850,29 @@ int CV_GpuNppImageWarpPerspectiveTest::test( const Mat& cpu1, const Mat& ) if (cpu1.type() == CV_32SC1) return CvTS::OK; - Mat cpudst; - cv::warpPerspective(cpu1, cpudst, M, cpu1.size(), INTER_CUBIC | WARP_INVERSE_MAP); + 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); - GpuMat gpu1(cpu1); - GpuMat gpudst; - cv::gpu::warpPerspective(gpu1, gpudst, M, gpu1.size(), INTER_CUBIC | WARP_INVERSE_MAP); + int test_res = CvTS::OK; - return CheckNorm(cpudst, gpudst); + 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; \ No newline at end of file +CV_GpuNppImageWarpPerspectiveTest CV_GpuNppImageWarpPerspective_test; diff --git a/tests/gpu/src/operator_async_call.cpp b/tests/gpu/src/operator_async_call.cpp index 1ecddc671d..2a9cf4f0ef 100644 --- a/tests/gpu/src/operator_async_call.cpp +++ b/tests/gpu/src/operator_async_call.cpp @@ -143,7 +143,16 @@ void CV_GpuMatAsyncCallTest::run( int /* start_from */) Mat cpumat(rows, cols, CV_8U); cpumat.setTo(Scalar::all(127)); - is_test_good &= compare_matrix(cpumat); + try + { + is_test_good &= compare_matrix(cpumat); + } + catch(cv::Exception& e) + { + if (!check_and_treat_gpu_exception(e, ts)) + throw; + return; + } if (is_test_good == true) ts->set_failed_test_info(CvTS::OK); @@ -151,4 +160,4 @@ void CV_GpuMatAsyncCallTest::run( int /* start_from */) ts->set_failed_test_info(CvTS::FAIL_GENERIC); } -CV_GpuMatAsyncCallTest CV_GpuMatAsyncCall_test; +//CV_GpuMatAsyncCallTest CV_GpuMatAsyncCall_test; diff --git a/tests/gpu/src/operator_convert_to.cpp b/tests/gpu/src/operator_convert_to.cpp index 47414a20d2..d9482e59b2 100644 --- a/tests/gpu/src/operator_convert_to.cpp +++ b/tests/gpu/src/operator_convert_to.cpp @@ -115,10 +115,12 @@ void CV_GpuMatOpConvertToTest::run(int /* start_from */) } catch(cv::Exception& e) { - ts->printf(CvTS::CONSOLE, "\nERROR: %s\n", e.what()); + if (!check_and_treat_gpu_exception(e, ts)) + throw; + return; } + ts->set_failed_test_info(passed ? CvTS::OK : CvTS::FAIL_GENERIC); } CV_GpuMatOpConvertToTest CV_GpuMatOpConvertToTest_test; - diff --git a/tests/gpu/src/operator_copy_to.cpp b/tests/gpu/src/operator_copy_to.cpp index 74d8d60bde..9ada88b488 100644 --- a/tests/gpu/src/operator_copy_to.cpp +++ b/tests/gpu/src/operator_copy_to.cpp @@ -136,14 +136,23 @@ void CV_GpuMatOpCopyToTest::run( int /* start_from */) { bool is_test_good = true; - for (int i = 0 ; i < 7; i++) + try { - Mat cpumat(rows, cols, i); - cpumat.setTo(Scalar::all(127)); + for (int i = 0 ; i < 7; i++) + { + Mat cpumat(rows, cols, i); + cpumat.setTo(Scalar::all(127)); - GpuMat gpumat(cpumat); + GpuMat gpumat(cpumat); - is_test_good &= compare_matrix(cpumat, gpumat); + is_test_good &= compare_matrix(cpumat, gpumat); + } + } + catch(const cv::Exception& e) + { + if (!check_and_treat_gpu_exception(e, ts)) + throw; + return; } if (is_test_good == true) diff --git a/tests/gpu/src/operator_set_to.cpp b/tests/gpu/src/operator_set_to.cpp index 71db0d2b38..6075c22344 100644 --- a/tests/gpu/src/operator_set_to.cpp +++ b/tests/gpu/src/operator_set_to.cpp @@ -132,11 +132,20 @@ void CV_GpuMatOpSetToTest::run( int /* start_from */) { bool is_test_good = true; - for (int i = 0; i < 7; i++) + try { - Mat cpumat(rows, cols, i, Scalar::all(0)); - GpuMat gpumat(cpumat); - is_test_good &= compare_matrix(cpumat, gpumat); + for (int i = 0; i < 7; i++) + { + Mat cpumat(rows, cols, i, Scalar::all(0)); + GpuMat gpumat(cpumat); + is_test_good &= compare_matrix(cpumat, gpumat); + } + } + catch(const cv::Exception& e) + { + if (!check_and_treat_gpu_exception(e, ts)) + throw; + return; } if (is_test_good == true) diff --git a/tests/gpu/src/stereo_bm.cpp b/tests/gpu/src/stereo_bm.cpp index 8631c571b7..2af46f4fce 100644 --- a/tests/gpu/src/stereo_bm.cpp +++ b/tests/gpu/src/stereo_bm.cpp @@ -70,18 +70,30 @@ void CV_GpuStereoBMTest::run(int ) return; } - cv::gpu::GpuMat disp; - cv::gpu::StereoBM_GPU bm(0, 128, 19); - bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp); + try + { + cv::gpu::GpuMat disp; + cv::gpu::StereoBM_GPU bm(0, 128, 19); + bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp); - disp.convertTo(disp, img_reference.type()); - double norm = cv::norm(disp, img_reference, cv::NORM_INF); + disp.convertTo(disp, img_reference.type()); + double norm = cv::norm(disp, img_reference, cv::NORM_INF); - if (norm >= 100) - ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm); - ts->set_failed_test_info((norm < 100) ? CvTS::OK : CvTS::FAIL_GENERIC); + if (norm >= 100) + { + ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm); + ts->set_failed_test_info(CvTS::FAIL_GENERIC); + return; + } + } + catch(const cv::Exception& e) + { + if (!check_and_treat_gpu_exception(e, ts)) + throw; + return; + } + + ts->set_failed_test_info(CvTS::OK); } - CV_GpuStereoBMTest CV_GpuStereoBM_test; - diff --git a/tests/gpu/src/stereo_bm_async.cpp b/tests/gpu/src/stereo_bm_async.cpp index fdbe9516f1..b178a0ec2b 100644 --- a/tests/gpu/src/stereo_bm_async.cpp +++ b/tests/gpu/src/stereo_bm_async.cpp @@ -74,24 +74,37 @@ void CV_GpuMatAsyncCallStereoBMTest::run( int /* start_from */) return; } - cv::gpu::GpuMat disp; - cv::gpu::StereoBM_GPU bm(0, 128, 19); + try + { + cv::gpu::GpuMat disp; + cv::gpu::StereoBM_GPU bm(0, 128, 19); - cv::gpu::Stream stream; + cv::gpu::Stream stream; - for (size_t i = 0; i < 50; i++) - { - bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp, stream); - } + for (size_t i = 0; i < 50; i++) + { + bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp, stream); + } - stream.waitForCompletion(); - disp.convertTo(disp, img_reference.type()); - double norm = cv::norm(disp, img_reference, cv::NORM_INF); + stream.waitForCompletion(); + disp.convertTo(disp, img_reference.type()); + double norm = cv::norm(disp, img_reference, cv::NORM_INF); - if (norm >= 100) - ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm); - ts->set_failed_test_info((norm < 100) ? CvTS::OK : CvTS::FAIL_GENERIC); + if (norm >= 100) + { + ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm); + ts->set_failed_test_info(CvTS::FAIL_GENERIC); + return; + } + } + catch(const cv::Exception& e) + { + if (!check_and_treat_gpu_exception(e, ts)) + throw; + return; + } + ts->set_failed_test_info(CvTS::OK); } CV_GpuMatAsyncCallStereoBMTest CV_GpuMatAsyncCallStereoBMTest_test; diff --git a/tests/gpu/src/stereo_bp.cpp b/tests/gpu/src/stereo_bp.cpp index f91898508d..f0b2e4a84e 100644 --- a/tests/gpu/src/stereo_bp.cpp +++ b/tests/gpu/src/stereo_bp.cpp @@ -69,20 +69,33 @@ void CV_GpuStereoBPTest::run(int ) return; } - cv::gpu::GpuMat disp; - cv::gpu::StereoBeliefPropagation bpm(64, 8, 2, 25, 0.1f, 15, 1, CV_16S); + try + { + cv::gpu::GpuMat disp; + cv::gpu::StereoBeliefPropagation bpm(64, 8, 2, 25, 0.1f, 15, 1, CV_16S); - bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp); + bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp); - //cv::imwrite(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", disp); + //cv::imwrite(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", disp); - disp.convertTo(disp, img_template.type()); + disp.convertTo(disp, img_template.type()); - double norm = cv::norm(disp, img_template, cv::NORM_INF); - if (norm >= 0.5) - ts->printf(CvTS::CONSOLE, "\nStereoBP norm = %f\n", norm); - ts->set_failed_test_info((norm < 0.5) ? CvTS::OK : CvTS::FAIL_GENERIC); + double norm = cv::norm(disp, img_template, cv::NORM_INF); + if (norm >= 0.5) + { + ts->printf(CvTS::CONSOLE, "\nStereoBP norm = %f\n", norm); + ts->set_failed_test_info(CvTS::FAIL_GENERIC); + return; + } + } + catch(const cv::Exception& e) + { + if (!check_and_treat_gpu_exception(e, ts)) + throw; + return; + } + + ts->set_failed_test_info(CvTS::OK); } - -CV_GpuStereoBPTest CV_GpuStereoBP_test; \ No newline at end of file +CV_GpuStereoBPTest CV_GpuStereoBP_test; diff --git a/tests/gpu/src/stereo_csbp.cpp b/tests/gpu/src/stereo_csbp.cpp index c6e9bf4d6e..e5613da697 100644 --- a/tests/gpu/src/stereo_csbp.cpp +++ b/tests/gpu/src/stereo_csbp.cpp @@ -59,16 +59,18 @@ CV_GpuStereoCSBPTest::CV_GpuStereoCSBPTest(): CvTest( "GPU-StereoCSBP", "Constan void CV_GpuStereoCSBPTest::run(int ) { - cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-L.png"); - cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-R.png"); - cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-disp.png", 0); + cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-L.png"); + cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-R.png"); + cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-disp.png", 0); - if (img_l.empty() || img_r.empty() || img_template.empty()) - { - ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA); - return; - } + if (img_l.empty() || img_r.empty() || img_template.empty()) + { + ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA); + return; + } + try + { cv::gpu::GpuMat disp; cv::gpu::StereoConstantSpaceBP bpm(128, 16, 4, 4); @@ -79,9 +81,21 @@ void CV_GpuStereoCSBPTest::run(int ) disp.convertTo(disp, img_template.type()); double norm = cv::norm(disp, img_template, cv::NORM_INF); - if (norm >= 0.5) + if (norm >= 0.5) + { ts->printf(CvTS::CONSOLE, "\nConstantSpaceStereoBP norm = %f\n", norm); - ts->set_failed_test_info((norm < 0.5) ? CvTS::OK : CvTS::FAIL_GENERIC); + ts->set_failed_test_info(CvTS::FAIL_GENERIC); + return; + } + } + catch(const cv::Exception& e) + { + if (!check_and_treat_gpu_exception(e, ts)) + throw; + return; + } + + ts->set_failed_test_info(CvTS::OK); } -CV_GpuStereoCSBPTest CV_GpuCSStereoBP_test; \ No newline at end of file +CV_GpuStereoCSBPTest CV_GpuCSStereoBP_test;