From: Vladislav Vinogradov Date: Wed, 13 Feb 2013 11:54:50 +0000 (+0400) Subject: added mask support to gpu norm and sum X-Git-Tag: accepted/tizen/6.0/unified/20201030.111113~1314^2~1511^2~7 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=61b54149b14cbf41122058fbde92b9025b4c0eb5;p=platform%2Fupstream%2Fopencv.git added mask support to gpu norm and sum --- diff --git a/modules/gpu/doc/matrix_reductions.rst b/modules/gpu/doc/matrix_reductions.rst index 538267e..e9229f8 100644 --- a/modules/gpu/doc/matrix_reductions.rst +++ b/modules/gpu/doc/matrix_reductions.rst @@ -32,6 +32,8 @@ Returns the norm of a matrix (or difference of two matrices). .. ocv:function:: double gpu::norm(const GpuMat& src1, int normType, GpuMat& buf) +.. ocv:function:: double gpu::norm(const GpuMat& src1, int normType, const GpuMat& mask, GpuMat& buf) + .. ocv:function:: double gpu::norm(const GpuMat& src1, const GpuMat& src2, int normType=NORM_L2) :param src1: Source matrix. Any matrices except 64F are supported. @@ -40,6 +42,8 @@ Returns the norm of a matrix (or difference of two matrices). :param normType: Norm type. ``NORM_L1`` , ``NORM_L2`` , and ``NORM_INF`` are supported for now. + :param mask: optional operation mask; it must have the same size as ``src1`` and ``CV_8UC1`` type. + :param buf: Optional buffer to avoid extra memory allocations. It is resized automatically. .. seealso:: :ocv:func:`norm` @@ -54,8 +58,12 @@ Returns the sum of matrix elements. .. ocv:function:: Scalar gpu::sum(const GpuMat& src, GpuMat& buf) +.. ocv:function:: Scalar gpu::sum(const GpuMat& src, const GpuMat& mask, GpuMat& buf) + :param src: Source image of any depth except for ``CV_64F`` . + :param mask: optional operation mask; it must have the same size as ``src1`` and ``CV_8UC1`` type. + :param buf: Optional buffer to avoid extra memory allocations. It is resized automatically. .. seealso:: :ocv:func:`sum` @@ -70,8 +78,12 @@ Returns the sum of absolute values for matrix elements. .. ocv:function:: Scalar gpu::absSum(const GpuMat& src, GpuMat& buf) +.. ocv:function:: Scalar gpu::absSum(const GpuMat& src, const GpuMat& mask, GpuMat& buf) + :param src: Source image of any depth except for ``CV_64F`` . + :param mask: optional operation mask; it must have the same size as ``src1`` and ``CV_8UC1`` type. + :param buf: Optional buffer to avoid extra memory allocations. It is resized automatically. @@ -84,8 +96,12 @@ Returns the squared sum of matrix elements. .. ocv:function:: Scalar gpu::sqrSum(const GpuMat& src, GpuMat& buf) +.. ocv:function:: Scalar gpu::sqrSum(const GpuMat& src, const GpuMat& mask, GpuMat& buf) + :param src: Source image of any depth except for ``CV_64F`` . + :param mask: optional operation mask; it must have the same size as ``src1`` and ``CV_8UC1`` type. + :param buf: Optional buffer to avoid extra memory allocations. It is resized automatically. diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 4eb339e..046339b 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -919,11 +919,8 @@ CV_EXPORTS void meanStdDev(const GpuMat& mtx, Scalar& mean, Scalar& stddev, GpuM //! supports NORM_INF, NORM_L1, NORM_L2 //! supports all matrices except 64F CV_EXPORTS double norm(const GpuMat& src1, int normType=NORM_L2); - -//! computes norm of array -//! supports NORM_INF, NORM_L1, NORM_L2 -//! supports all matrices except 64F CV_EXPORTS double norm(const GpuMat& src1, int normType, GpuMat& buf); +CV_EXPORTS double norm(const GpuMat& src1, int normType, const GpuMat& mask, GpuMat& buf); //! computes norm of the difference between two arrays //! supports NORM_INF, NORM_L1, NORM_L2 @@ -933,45 +930,33 @@ CV_EXPORTS double norm(const GpuMat& src1, const GpuMat& src2, int normType=NORM //! computes sum of array elements //! supports only single channel images CV_EXPORTS Scalar sum(const GpuMat& src); - -//! computes sum of array elements -//! supports only single channel images CV_EXPORTS Scalar sum(const GpuMat& src, GpuMat& buf); +CV_EXPORTS Scalar sum(const GpuMat& src, const GpuMat& mask, GpuMat& buf); //! computes sum of array elements absolute values //! supports only single channel images CV_EXPORTS Scalar absSum(const GpuMat& src); - -//! computes sum of array elements absolute values -//! supports only single channel images CV_EXPORTS Scalar absSum(const GpuMat& src, GpuMat& buf); +CV_EXPORTS Scalar absSum(const GpuMat& src, const GpuMat& mask, GpuMat& buf); //! computes squared sum of array elements //! supports only single channel images CV_EXPORTS Scalar sqrSum(const GpuMat& src); - -//! computes squared sum of array elements -//! supports only single channel images CV_EXPORTS Scalar sqrSum(const GpuMat& src, GpuMat& buf); +CV_EXPORTS Scalar sqrSum(const GpuMat& src, const GpuMat& mask, GpuMat& buf); //! finds global minimum and maximum array elements and returns their values CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal=0, const GpuMat& mask=GpuMat()); - -//! finds global minimum and maximum array elements and returns their values CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal, const GpuMat& mask, GpuMat& buf); //! finds global minimum and maximum array elements and returns their values with locations CV_EXPORTS void minMaxLoc(const GpuMat& src, double* minVal, double* maxVal=0, Point* minLoc=0, Point* maxLoc=0, const GpuMat& mask=GpuMat()); - -//! finds global minimum and maximum array elements and returns their values with locations CV_EXPORTS void minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point* minLoc, Point* maxLoc, const GpuMat& mask, GpuMat& valbuf, GpuMat& locbuf); //! counts non-zero array elements CV_EXPORTS int countNonZero(const GpuMat& src); - -//! counts non-zero array elements CV_EXPORTS int countNonZero(const GpuMat& src, GpuMat& buf); //! reduces a matrix to a vector diff --git a/modules/gpu/perf/perf_core.cpp b/modules/gpu/perf/perf_core.cpp index c78bfd6..6b407f4 100644 --- a/modules/gpu/perf/perf_core.cpp +++ b/modules/gpu/perf/perf_core.cpp @@ -1631,7 +1631,7 @@ PERF_TEST_P(Sz_Depth_Norm, Core_Norm, Combine( cv::gpu::GpuMat d_src(src); cv::gpu::GpuMat d_buf; - TEST_CYCLE() dst = cv::gpu::norm(d_src, normType, d_buf); + TEST_CYCLE() dst = cv::gpu::norm(d_src, normType, cv::gpu::GpuMat(), d_buf); } else { @@ -1701,7 +1701,7 @@ PERF_TEST_P(Sz_Depth_Cn, Core_Sum, Combine( cv::gpu::GpuMat d_src(src); cv::gpu::GpuMat d_buf; - TEST_CYCLE() dst = cv::gpu::sum(d_src, d_buf); + TEST_CYCLE() dst = cv::gpu::sum(d_src, cv::gpu::GpuMat(), d_buf); } else { @@ -1736,7 +1736,7 @@ PERF_TEST_P(Sz_Depth_Cn, Core_SumAbs, Combine( cv::gpu::GpuMat d_src(src); cv::gpu::GpuMat d_buf; - TEST_CYCLE() dst = cv::gpu::absSum(d_src, d_buf); + TEST_CYCLE() dst = cv::gpu::absSum(d_src, cv::gpu::GpuMat(), d_buf); SANITY_CHECK(dst, 1e-6); } @@ -1770,7 +1770,7 @@ PERF_TEST_P(Sz_Depth_Cn, Core_SumSqr, Combine( cv::gpu::GpuMat d_src(src); cv::gpu::GpuMat d_buf; - TEST_CYCLE() dst = cv::gpu::sqrSum(d_src, d_buf); + TEST_CYCLE() dst = cv::gpu::sqrSum(d_src, cv::gpu::GpuMat(), d_buf); SANITY_CHECK(dst, 1e-6); } diff --git a/modules/gpu/src/cuda/matrix_reductions.cu b/modules/gpu/src/cuda/matrix_reductions.cu index b48c47e..745daca 100644 --- a/modules/gpu/src/cuda/matrix_reductions.cu +++ b/modules/gpu/src/cuda/matrix_reductions.cu @@ -352,8 +352,8 @@ namespace sum } }; - template - __global__ void kernel(const PtrStepSz src, result_type* result, const Op op, const int twidth, const int theight) + template + __global__ void kernel(const PtrStepSz src, result_type* result, const Mask mask, const Op op, const int twidth, const int theight) { typedef typename VecTraits::elem_type T; typedef typename VecTraits::elem_type R; @@ -375,9 +375,11 @@ namespace sum for (int j = 0, x = x0; j < twidth && x < src.cols; ++j, x += blockDim.x) { - const src_type srcVal = ptr[x]; - - sum = sum + op(saturate_cast(srcVal)); + if (mask(y, x)) + { + const src_type srcVal = ptr[x]; + sum = sum + op(saturate_cast(srcVal)); + } } } @@ -410,7 +412,7 @@ namespace sum } template class Op> - void caller(PtrStepSzb src_, void* buf_, double* out) + void caller(PtrStepSzb src_, void* buf_, double* out, PtrStepSzb mask) { typedef typename TypeVec::vec_type src_type; typedef typename TypeVec::vec_type result_type; @@ -426,7 +428,10 @@ namespace sum Op op; - kernel<<>>(src, buf, op, twidth, theight); + if (mask.data) + kernel<<>>(src, buf, SingleMask(mask), op, twidth, theight); + else + kernel<<>>(src, buf, WithOutMask(), op, twidth, theight); cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaDeviceSynchronize() ); @@ -450,88 +455,88 @@ namespace sum template <> struct SumType { typedef double R; }; template - void run(PtrStepSzb src, void* buf, double* out) + void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask) { typedef typename SumType::R R; - caller(src, buf, out); + caller(src, buf, out, mask); } - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); - template void run(PtrStepSzb src, void* buf, double* out); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void run(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); template - void runAbs(PtrStepSzb src, void* buf, double* out) + void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask) { typedef typename SumType::R R; - caller(src, buf, out); + caller(src, buf, out, mask); } - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); - template void runAbs(PtrStepSzb src, void* buf, double* out); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runAbs(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); template struct Sqr : unary_function { @@ -542,45 +547,45 @@ namespace sum }; template - void runSqr(PtrStepSzb src, void* buf, double* out) + void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask) { - caller(src, buf, out); + caller(src, buf, out, mask); } - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); - template void runSqr(PtrStepSzb src, void* buf, double* out); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); + template void runSqr(PtrStepSzb src, void* buf, double* out, PtrStepSzb mask); } ///////////////////////////////////////////////////////////// diff --git a/modules/gpu/src/matrix_reductions.cpp b/modules/gpu/src/matrix_reductions.cpp index 4295644..67e65fc 100644 --- a/modules/gpu/src/matrix_reductions.cpp +++ b/modules/gpu/src/matrix_reductions.cpp @@ -51,13 +51,17 @@ void cv::gpu::meanStdDev(const GpuMat&, Scalar&, Scalar&) { throw_nogpu(); } void cv::gpu::meanStdDev(const GpuMat&, Scalar&, Scalar&, GpuMat&) { throw_nogpu(); } double cv::gpu::norm(const GpuMat&, int) { throw_nogpu(); return 0.0; } double cv::gpu::norm(const GpuMat&, int, GpuMat&) { throw_nogpu(); return 0.0; } +double cv::gpu::norm(const GpuMat&, int, const GpuMat&, GpuMat&) { throw_nogpu(); return 0.0; } double cv::gpu::norm(const GpuMat&, const GpuMat&, int) { throw_nogpu(); return 0.0; } Scalar cv::gpu::sum(const GpuMat&) { throw_nogpu(); return Scalar(); } Scalar cv::gpu::sum(const GpuMat&, GpuMat&) { throw_nogpu(); return Scalar(); } +Scalar cv::gpu::sum(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); return Scalar(); } Scalar cv::gpu::absSum(const GpuMat&) { throw_nogpu(); return Scalar(); } Scalar cv::gpu::absSum(const GpuMat&, GpuMat&) { throw_nogpu(); return Scalar(); } +Scalar cv::gpu::absSum(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); return Scalar(); } Scalar cv::gpu::sqrSum(const GpuMat&) { throw_nogpu(); return Scalar(); } Scalar cv::gpu::sqrSum(const GpuMat&, GpuMat&) { throw_nogpu(); return Scalar(); } +Scalar cv::gpu::sqrSum(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); return Scalar(); } void cv::gpu::minMax(const GpuMat&, double*, double*, const GpuMat&) { throw_nogpu(); } void cv::gpu::minMax(const GpuMat&, double*, double*, const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::minMaxLoc(const GpuMat&, double*, double*, Point*, Point*, const GpuMat&) { throw_nogpu(); } @@ -150,24 +154,30 @@ void cv::gpu::meanStdDev(const GpuMat& src, Scalar& mean, Scalar& stddev, GpuMat double cv::gpu::norm(const GpuMat& src, int normType) { GpuMat buf; - return norm(src, normType, buf); + return norm(src, normType, GpuMat(), buf); } double cv::gpu::norm(const GpuMat& src, int normType, GpuMat& buf) { + return norm(src, normType, GpuMat(), buf); +} + +double cv::gpu::norm(const GpuMat& src, int normType, const GpuMat& mask, GpuMat& buf) +{ CV_Assert(normType == NORM_INF || normType == NORM_L1 || normType == NORM_L2); + CV_Assert(mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src.size() && src.channels() == 1)); GpuMat src_single_channel = src.reshape(1); if (normType == NORM_L1) - return absSum(src_single_channel, buf)[0]; + return absSum(src_single_channel, mask, buf)[0]; if (normType == NORM_L2) - return std::sqrt(sqrSum(src_single_channel, buf)[0]); + return std::sqrt(sqrSum(src_single_channel, mask, buf)[0]); // NORM_INF double min_val, max_val; - minMax(src_single_channel, &min_val, &max_val, GpuMat(), buf); + minMax(src_single_channel, &min_val, &max_val, mask, buf); return std::max(std::abs(min_val), std::abs(max_val)); } @@ -209,24 +219,29 @@ namespace sum void getBufSize(int cols, int rows, int cn, int& bufcols, int& bufrows); template - void run(PtrStepSzb src, void* buf, double* sum); + void run(PtrStepSzb src, void* buf, double* sum, PtrStepSzb mask); template - void runAbs(PtrStepSzb src, void* buf, double* sum); + void runAbs(PtrStepSzb src, void* buf, double* sum, PtrStepSzb mask); template - void runSqr(PtrStepSzb src, void* buf, double* sum); + void runSqr(PtrStepSzb src, void* buf, double* sum, PtrStepSzb mask); } Scalar cv::gpu::sum(const GpuMat& src) { GpuMat buf; - return sum(src, buf); + return sum(src, GpuMat(), buf); } Scalar cv::gpu::sum(const GpuMat& src, GpuMat& buf) { - typedef void (*func_t)(PtrStepSzb src, void* buf, double* sum); + return sum(src, GpuMat(), buf); +} + +Scalar cv::gpu::sum(const GpuMat& src, const GpuMat& mask, GpuMat& buf) +{ + typedef void (*func_t)(PtrStepSzb src, void* buf, double* sum, PtrStepSzb mask); static const func_t funcs[7][5] = { {0, ::sum::run, ::sum::run, ::sum::run, ::sum::run}, @@ -238,6 +253,8 @@ Scalar cv::gpu::sum(const GpuMat& src, GpuMat& buf) {0, ::sum::run, ::sum::run, ::sum::run, ::sum::run} }; + CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src.size()) ); + if (src.depth() == CV_64F) { if (!deviceSupports(NATIVE_DOUBLE)) @@ -252,7 +269,7 @@ Scalar cv::gpu::sum(const GpuMat& src, GpuMat& buf) const func_t func = funcs[src.depth()][src.channels()]; double result[4]; - func(src, buf.data, result); + func(src, buf.data, result, mask); return Scalar(result[0], result[1], result[2], result[3]); } @@ -260,12 +277,17 @@ Scalar cv::gpu::sum(const GpuMat& src, GpuMat& buf) Scalar cv::gpu::absSum(const GpuMat& src) { GpuMat buf; - return absSum(src, buf); + return absSum(src, GpuMat(), buf); } Scalar cv::gpu::absSum(const GpuMat& src, GpuMat& buf) { - typedef void (*func_t)(PtrStepSzb src, void* buf, double* sum); + return absSum(src, GpuMat(), buf); +} + +Scalar cv::gpu::absSum(const GpuMat& src, const GpuMat& mask, GpuMat& buf) +{ + typedef void (*func_t)(PtrStepSzb src, void* buf, double* sum, PtrStepSzb mask); static const func_t funcs[7][5] = { {0, ::sum::runAbs, ::sum::runAbs, ::sum::runAbs, ::sum::runAbs}, @@ -277,6 +299,8 @@ Scalar cv::gpu::absSum(const GpuMat& src, GpuMat& buf) {0, ::sum::runAbs, ::sum::runAbs, ::sum::runAbs, ::sum::runAbs} }; + CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src.size()) ); + if (src.depth() == CV_64F) { if (!deviceSupports(NATIVE_DOUBLE)) @@ -291,7 +315,7 @@ Scalar cv::gpu::absSum(const GpuMat& src, GpuMat& buf) const func_t func = funcs[src.depth()][src.channels()]; double result[4]; - func(src, buf.data, result); + func(src, buf.data, result, mask); return Scalar(result[0], result[1], result[2], result[3]); } @@ -299,12 +323,17 @@ Scalar cv::gpu::absSum(const GpuMat& src, GpuMat& buf) Scalar cv::gpu::sqrSum(const GpuMat& src) { GpuMat buf; - return sqrSum(src, buf); + return sqrSum(src, GpuMat(), buf); } Scalar cv::gpu::sqrSum(const GpuMat& src, GpuMat& buf) { - typedef void (*func_t)(PtrStepSzb src, void* buf, double* sum); + return sqrSum(src, GpuMat(), buf); +} + +Scalar cv::gpu::sqrSum(const GpuMat& src, const GpuMat& mask, GpuMat& buf) +{ + typedef void (*func_t)(PtrStepSzb src, void* buf, double* sum, PtrStepSzb mask); static const func_t funcs[7][5] = { {0, ::sum::runSqr, ::sum::runSqr, ::sum::runSqr, ::sum::runSqr}, @@ -316,6 +345,8 @@ Scalar cv::gpu::sqrSum(const GpuMat& src, GpuMat& buf) {0, ::sum::runSqr, ::sum::runSqr, ::sum::runSqr, ::sum::runSqr} }; + CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src.size()) ); + if (src.depth() == CV_64F) { if (!deviceSupports(NATIVE_DOUBLE)) @@ -330,7 +361,7 @@ Scalar cv::gpu::sqrSum(const GpuMat& src, GpuMat& buf) const func_t func = funcs[src.depth()][src.channels()]; double result[4]; - func(src, buf.data, result); + func(src, buf.data, result, mask); return Scalar(result[0], result[1], result[2], result[3]); } diff --git a/modules/gpu/test/test_core.cpp b/modules/gpu/test/test_core.cpp index eeea41e..ad0cf53 100644 --- a/modules/gpu/test/test_core.cpp +++ b/modules/gpu/test/test_core.cpp @@ -2918,10 +2918,12 @@ PARAM_TEST_CASE(Norm, cv::gpu::DeviceInfo, cv::Size, MatDepth, NormCode, UseRoi) GPU_TEST_P(Norm, Accuracy) { cv::Mat src = randomMat(size, depth); + cv::Mat mask = randomMat(size, CV_8UC1, 0, 2); - double val = cv::gpu::norm(loadMat(src, useRoi), normCode); + cv::gpu::GpuMat d_buf; + double val = cv::gpu::norm(loadMat(src, useRoi), normCode, loadMat(mask, useRoi), d_buf); - double val_gold = cv::norm(src, normCode); + double val_gold = cv::norm(src, normCode, mask); EXPECT_NEAR(val_gold, val, depth < CV_32F ? 0.0 : 1.0); }