diff --git a/modules/gpufilters/include/opencv2/gpufilters.hpp b/modules/gpufilters/include/opencv2/gpufilters.hpp index 7595bd5439..c5a61a06e9 100644 --- a/modules/gpufilters/include/opencv2/gpufilters.hpp +++ b/modules/gpufilters/include/opencv2/gpufilters.hpp @@ -241,7 +241,18 @@ inline void morphologyEx(InputArray src, OutputArray dst, int op, InputArray ker f->apply(src, dst, stream); } +//////////////////////////////////////////////////////////////////////////////////////////////////// +// Image Rank Filter +//! Result pixel value is the maximum of pixel values under the rectangular mask region +CV_EXPORTS Ptr createBoxMaxFilter(int srcType, Size ksize, + Point anchor = Point(-1, -1), + int borderMode = BORDER_DEFAULT, Scalar borderVal = Scalar::all(0)); + +//! Result pixel value is the maximum of pixel values under the rectangular mask region +CV_EXPORTS Ptr createBoxMinFilter(int srcType, Size ksize, + Point anchor = Point(-1, -1), + int borderMode = BORDER_DEFAULT, Scalar borderVal = Scalar::all(0)); @@ -326,12 +337,6 @@ CV_EXPORTS Ptr getColumnSumFilter_GPU(int sumType, int dst -//! returns maximum filter -CV_EXPORTS Ptr getMaxFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1,-1)); - -//! returns minimum filter -CV_EXPORTS Ptr getMinFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1,-1)); - diff --git a/modules/gpufilters/src/filtering.cpp b/modules/gpufilters/src/filtering.cpp index 145e049842..a8ec83a74c 100644 --- a/modules/gpufilters/src/filtering.cpp +++ b/modules/gpufilters/src/filtering.cpp @@ -63,6 +63,8 @@ Ptr cv::gpu::createGaussianFilter(int, int, Size, double, double, int, i Ptr cv::gpu::createMorphologyFilter(int, int, InputArray, Point, int) { throw_no_cuda(); return Ptr(); } +Ptr cv::gpu::createBoxMaxFilter(int, Size, Point, int, Scalar) { throw_no_cuda(); return Ptr(); } +Ptr cv::gpu::createBoxMinFilter(int, Size, Point, int, Scalar) { throw_no_cuda(); return Ptr(); } @@ -70,8 +72,6 @@ Ptr cv::gpu::createMorphologyFilter(int, int, InputArray, Point, int) { Ptr cv::gpu::getRowSumFilter_GPU(int, int, int, int) { throw_no_cuda(); return Ptr(0); } Ptr cv::gpu::getColumnSumFilter_GPU(int, int, int, int) { throw_no_cuda(); return Ptr(0); } -Ptr cv::gpu::getMaxFilter_GPU(int, int, const Size&, Point) { throw_no_cuda(); return Ptr(0); } -Ptr cv::gpu::getMinFilter_GPU(int, int, const Size&, Point) { throw_no_cuda(); return Ptr(0); } @@ -783,6 +783,99 @@ Ptr cv::gpu::createMorphologyFilter(int op, int srcType, InputArray kern } } +//////////////////////////////////////////////////////////////////////////////////////////////////// +// Image Rank Filter + +namespace +{ + enum + { + RANK_MAX, + RANK_MIN + }; + + class NPPRankFilter : public Filter + { + public: + NPPRankFilter(int op, int srcType, Size ksize, Point anchor, int borderMode, Scalar borderVal); + + void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); + + private: + typedef NppStatus (*nppFilterRank_t)(const Npp8u* pSrc, Npp32s nSrcStep, Npp8u* pDst, Npp32s nDstStep, NppiSize oSizeROI, + NppiSize oMaskSize, NppiPoint oAnchor); + + int type_; + Size ksize_; + Point anchor_; + int borderMode_; + Scalar borderVal_; + nppFilterRank_t func_; + + GpuMat srcBorder_; + }; + + NPPRankFilter::NPPRankFilter(int op, int srcType, Size ksize, Point anchor, int borderMode, Scalar borderVal) : + type_(srcType), ksize_(ksize), anchor_(anchor), borderMode_(borderMode), borderVal_(borderVal) + { + static const nppFilterRank_t maxFuncs[] = {0, nppiFilterMax_8u_C1R, 0, 0, nppiFilterMax_8u_C4R}; + static const nppFilterRank_t minFuncs[] = {0, nppiFilterMin_8u_C1R, 0, 0, nppiFilterMin_8u_C4R}; + + CV_Assert( srcType == CV_8UC1 || srcType == CV_8UC4 ); + + normalizeAnchor(anchor_, ksize_); + + if (op == RANK_MAX) + func_ = maxFuncs[CV_MAT_CN(srcType)]; + else + func_ = minFuncs[CV_MAT_CN(srcType)]; + } + + void NPPRankFilter::apply(InputArray _src, OutputArray _dst, Stream& _stream) + { + GpuMat src = _src.getGpuMat(); + CV_Assert( src.type() == type_ ); + + gpu::copyMakeBorder(src, srcBorder_, ksize_.height, ksize_.height, ksize_.width, ksize_.width, borderMode_, borderVal_, _stream); + + _dst.create(src.size(), src.type()); + GpuMat dst = _dst.getGpuMat(); + + GpuMat srcRoi = srcBorder_(Rect(ksize_.width, ksize_.height, src.cols, src.rows)); + + cudaStream_t stream = StreamAccessor::getStream(_stream); + NppStreamHandler h(stream); + + NppiSize oSizeROI; + oSizeROI.width = src.cols; + oSizeROI.height = src.rows; + + NppiSize oMaskSize; + oMaskSize.height = ksize_.height; + oMaskSize.width = ksize_.width; + + NppiPoint oAnchor; + oAnchor.x = anchor_.x; + oAnchor.y = anchor_.y; + + nppSafeCall( func_(srcRoi.ptr(), static_cast(srcRoi.step), dst.ptr(), static_cast(dst.step), + oSizeROI, oMaskSize, oAnchor) ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } +} + +Ptr cv::gpu::createBoxMaxFilter(int srcType, Size ksize, Point anchor, int borderMode, Scalar borderVal) +{ + return new NPPRankFilter(RANK_MAX, srcType, ksize, anchor, borderMode, borderVal); +} + +Ptr cv::gpu::createBoxMinFilter(int srcType, Size ksize, Point anchor, int borderMode, Scalar borderVal) +{ + return new NPPRankFilter(RANK_MIN, srcType, ksize, anchor, borderMode, borderVal); +} + @@ -924,64 +1017,4 @@ Ptr cv::gpu::getColumnSumFilter_GPU(int sumType, int dstTy -//////////////////////////////////////////////////////////////////////////////////////////////////// -// Image Rank Filter - -namespace -{ - typedef NppStatus (*nppFilterRank_t)(const Npp8u * pSrc, Npp32s nSrcStep, Npp8u * pDst, Npp32s nDstStep, NppiSize oSizeROI, - NppiSize oMaskSize, NppiPoint oAnchor); - - struct NPPRankFilter : public BaseFilter_GPU - { - NPPRankFilter(const Size& ksize_, const Point& anchor_, nppFilterRank_t func_) : BaseFilter_GPU(ksize_, anchor_), func(func_) {} - - virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& s = Stream::Null()) - { - NppiSize sz; - sz.width = src.cols; - sz.height = src.rows; - NppiSize oKernelSize; - oKernelSize.height = ksize.height; - oKernelSize.width = ksize.width; - NppiPoint oAnchor; - oAnchor.x = anchor.x; - oAnchor.y = anchor.y; - - cudaStream_t stream = StreamAccessor::getStream(s); - - NppStreamHandler h(stream); - - nppSafeCall( func(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz, oKernelSize, oAnchor) ); - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } - - nppFilterRank_t func; - }; -} - -Ptr cv::gpu::getMaxFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor) -{ - static const nppFilterRank_t nppFilterRank_callers[] = {0, nppiFilterMax_8u_C1R, 0, 0, nppiFilterMax_8u_C4R}; - - CV_Assert((srcType == CV_8UC1 || srcType == CV_8UC4) && dstType == srcType); - - normalizeAnchor(anchor, ksize); - - return Ptr(new NPPRankFilter(ksize, anchor, nppFilterRank_callers[CV_MAT_CN(srcType)])); -} - -Ptr cv::gpu::getMinFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor) -{ - static const nppFilterRank_t nppFilterRank_callers[] = {0, nppiFilterMin_8u_C1R, 0, 0, nppiFilterMin_8u_C4R}; - - CV_Assert((srcType == CV_8UC1 || srcType == CV_8UC4) && dstType == srcType); - - normalizeAnchor(anchor, ksize); - - return Ptr(new NPPRankFilter(ksize, anchor, nppFilterRank_callers[CV_MAT_CN(srcType)])); -} - #endif