From 5720eaf35472b6b40985792c7a9c338b74ed1652 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 29 Apr 2013 16:03:16 +0400 Subject: [PATCH] refactored Morphology Filters --- modules/gpubgsegm/src/fgd.cpp | 16 +- .../gpufilters/include/opencv2/gpufilters.hpp | 61 +- modules/gpufilters/perf/perf_filters.cpp | 23 +- modules/gpufilters/src/filtering.cpp | 559 +++++++++--------- modules/gpufilters/test/test_filters.cpp | 21 +- samples/gpu/morphology.cpp | 226 ++++--- samples/gpu/performance/tests.cpp | 6 +- 7 files changed, 503 insertions(+), 409 deletions(-) diff --git a/modules/gpubgsegm/src/fgd.cpp b/modules/gpubgsegm/src/fgd.cpp index 1b4038304a..fb14ff172a 100644 --- a/modules/gpubgsegm/src/fgd.cpp +++ b/modules/gpubgsegm/src/fgd.cpp @@ -228,11 +228,10 @@ private: cv::gpu::GpuMat countBuf_; cv::gpu::GpuMat buf_; - cv::gpu::GpuMat filterBuf_; cv::gpu::GpuMat filterBrd_; - cv::Ptr dilateFilter_; - cv::Ptr erodeFilter_; + cv::Ptr dilateFilter_; + cv::Ptr erodeFilter_; CvMemStorage* storage_; }; @@ -305,8 +304,8 @@ void cv::gpu::FGDStatModel::Impl::create(const cv::gpu::GpuMat& firstFrame, cons cv::Mat kernel = cv::getStructuringElement(cv::MORPH_RECT, cv::Size(1 + params_.perform_morphing * 2, 1 + params_.perform_morphing * 2)); cv::Point anchor(params_.perform_morphing, params_.perform_morphing); - dilateFilter_ = cv::gpu::createMorphologyFilter_GPU(cv::MORPH_DILATE, CV_8UC1, kernel, filterBuf_, anchor); - erodeFilter_ = cv::gpu::createMorphologyFilter_GPU(cv::MORPH_ERODE, CV_8UC1, kernel, filterBuf_, anchor); + dilateFilter_ = cv::gpu::createMorphologyFilter(cv::MORPH_DILATE, CV_8UC1, kernel, anchor); + erodeFilter_ = cv::gpu::createMorphologyFilter(cv::MORPH_ERODE, CV_8UC1, kernel, anchor); } } @@ -326,7 +325,6 @@ void cv::gpu::FGDStatModel::Impl::release() countBuf_.release(); buf_.release(); - filterBuf_.release(); filterBrd_.release(); } @@ -488,14 +486,14 @@ namespace namespace { - void morphology(const cv::gpu::GpuMat& src, cv::gpu::GpuMat& dst, cv::gpu::GpuMat& filterBrd, int brd, cv::Ptr& filter, cv::Scalar brdVal) + void morphology(const cv::gpu::GpuMat& src, cv::gpu::GpuMat& dst, cv::gpu::GpuMat& filterBrd, int brd, cv::Ptr& filter, cv::Scalar brdVal) { cv::gpu::copyMakeBorder(src, filterBrd, brd, brd, brd, brd, cv::BORDER_CONSTANT, brdVal); - filter->apply(filterBrd(cv::Rect(brd, brd, src.cols, src.rows)), dst, cv::Rect(0, 0, src.cols, src.rows)); + filter->apply(filterBrd(cv::Rect(brd, brd, src.cols, src.rows)), dst); } void smoothForeground(cv::gpu::GpuMat& foreground, cv::gpu::GpuMat& filterBrd, cv::gpu::GpuMat& buf, - cv::Ptr& erodeFilter, cv::Ptr& dilateFilter, + cv::Ptr& erodeFilter, cv::Ptr& dilateFilter, const cv::gpu::FGDStatModel::Params& params) { const int brd = params.perform_morphing; diff --git a/modules/gpufilters/include/opencv2/gpufilters.hpp b/modules/gpufilters/include/opencv2/gpufilters.hpp index bf7a8e928b..7595bd5439 100644 --- a/modules/gpufilters/include/opencv2/gpufilters.hpp +++ b/modules/gpufilters/include/opencv2/gpufilters.hpp @@ -48,6 +48,7 @@ #endif #include "opencv2/core/gpu.hpp" +#include "opencv2/imgproc.hpp" #if defined __GNUC__ #define __OPENCV_GPUFILTERS_DEPR_BEFORE__ @@ -203,8 +204,42 @@ inline void GaussianBlur(InputArray src, OutputArray dst, Size ksize, double sig f->apply(src, dst, stream); } +//////////////////////////////////////////////////////////////////////////////////////////////////// +// Morphology Filter +//! returns 2D morphological filter +//! supports CV_8UC1 and CV_8UC4 types +CV_EXPORTS Ptr createMorphologyFilter(int op, int srcType, InputArray kernel, Point anchor = Point(-1, -1), int iterations = 1); +__OPENCV_GPUFILTERS_DEPR_BEFORE__ void erode(InputArray src, OutputArray dst, InputArray kernel, + Point anchor = Point(-1, -1), int iterations = 1, + Stream& stream = Stream::Null()) __OPENCV_GPUFILTERS_DEPR_AFTER__; + +inline void erode(InputArray src, OutputArray dst, InputArray kernel, Point anchor, int iterations, Stream& stream) +{ + Ptr f = gpu::createMorphologyFilter(MORPH_ERODE, src.type(), kernel, anchor, iterations); + f->apply(src, dst, stream); +} + +__OPENCV_GPUFILTERS_DEPR_BEFORE__ void dilate(InputArray src, OutputArray dst, InputArray kernel, + Point anchor = Point(-1, -1), int iterations = 1, + Stream& stream = Stream::Null()) __OPENCV_GPUFILTERS_DEPR_AFTER__; + +inline void dilate(InputArray src, OutputArray dst, InputArray kernel, Point anchor, int iterations, Stream& stream) +{ + Ptr f = gpu::createMorphologyFilter(MORPH_DILATE, src.type(), kernel, anchor, iterations); + f->apply(src, dst, stream); +} + +__OPENCV_GPUFILTERS_DEPR_BEFORE__ void morphologyEx(InputArray src, OutputArray dst, int op, + InputArray kernel, Point anchor = Point(-1, -1), int iterations = 1, + Stream& stream = Stream::Null()) __OPENCV_GPUFILTERS_DEPR_AFTER__; + +inline void morphologyEx(InputArray src, OutputArray dst, int op, InputArray kernel, Point anchor, int iterations, Stream& stream) +{ + Ptr f = gpu::createMorphologyFilter(op, src.type(), kernel, anchor, iterations); + f->apply(src, dst, stream); +} @@ -285,18 +320,7 @@ CV_EXPORTS Ptr getColumnSumFilter_GPU(int sumType, int dst -//! returns 2D morphological filter -//! only MORPH_ERODE and MORPH_DILATE are supported -//! supports CV_8UC1 and CV_8UC4 types -//! kernel must have CV_8UC1 type, one rows and cols == ksize.width * ksize.height -CV_EXPORTS Ptr getMorphologyFilter_GPU(int op, int type, const Mat& kernel, const Size& ksize, - Point anchor=Point(-1,-1)); -//! returns morphological filter engine. Only MORPH_ERODE and MORPH_DILATE are supported. -CV_EXPORTS Ptr createMorphologyFilter_GPU(int op, int type, const Mat& kernel, - const Point& anchor = Point(-1,-1), int iterations = 1); -CV_EXPORTS Ptr createMorphologyFilter_GPU(int op, int type, const Mat& kernel, GpuMat& buf, - const Point& anchor = Point(-1,-1), int iterations = 1); @@ -310,22 +334,7 @@ CV_EXPORTS Ptr getMinFilter_GPU(int srcType, int dstType, const -//! erodes the image (applies the local minimum operator) -CV_EXPORTS void erode(const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor = Point(-1, -1), int iterations = 1); -CV_EXPORTS void erode(const GpuMat& src, GpuMat& dst, const Mat& kernel, GpuMat& buf, - Point anchor = Point(-1, -1), int iterations = 1, - Stream& stream = Stream::Null()); -//! dilates the image (applies the local maximum operator) -CV_EXPORTS void dilate(const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor = Point(-1, -1), int iterations = 1); -CV_EXPORTS void dilate(const GpuMat& src, GpuMat& dst, const Mat& kernel, GpuMat& buf, - Point anchor = Point(-1, -1), int iterations = 1, - Stream& stream = Stream::Null()); - -//! applies an advanced morphological operation to the image -CV_EXPORTS void morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, Point anchor = Point(-1, -1), int iterations = 1); -CV_EXPORTS void morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, GpuMat& buf1, GpuMat& buf2, - Point anchor = Point(-1, -1), int iterations = 1, Stream& stream = Stream::Null()); diff --git a/modules/gpufilters/perf/perf_filters.cpp b/modules/gpufilters/perf/perf_filters.cpp index 9c7dcb8b43..6ad0998a5b 100644 --- a/modules/gpufilters/perf/perf_filters.cpp +++ b/modules/gpufilters/perf/perf_filters.cpp @@ -263,13 +263,6 @@ PERF_TEST_P(Sz_Type_KernelSz, GaussianBlur, Combine(GPU_TYPICAL_MAT_SIZES, Value } } - - - - - - - ////////////////////////////////////////////////////////////////////// // Erode @@ -289,9 +282,10 @@ PERF_TEST_P(Sz_Type, Erode, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8U { const cv::gpu::GpuMat d_src(src); cv::gpu::GpuMat dst; - cv::gpu::GpuMat d_buf; - TEST_CYCLE() cv::gpu::erode(d_src, dst, ker, d_buf); + cv::Ptr erode = cv::gpu::createMorphologyFilter(cv::MORPH_ERODE, src.type(), ker); + + TEST_CYCLE() erode->apply(d_src, dst); GPU_SANITY_CHECK(dst); } @@ -324,9 +318,10 @@ PERF_TEST_P(Sz_Type, Dilate, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8 { const cv::gpu::GpuMat d_src(src); cv::gpu::GpuMat dst; - cv::gpu::GpuMat d_buf; - TEST_CYCLE() cv::gpu::dilate(d_src, dst, ker, d_buf); + cv::Ptr dilate = cv::gpu::createMorphologyFilter(cv::MORPH_DILATE, src.type(), ker); + + TEST_CYCLE() dilate->apply(d_src, dst); GPU_SANITY_CHECK(dst); } @@ -364,10 +359,10 @@ PERF_TEST_P(Sz_Type_Op, MorphologyEx, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8 { const cv::gpu::GpuMat d_src(src); cv::gpu::GpuMat dst; - cv::gpu::GpuMat d_buf1; - cv::gpu::GpuMat d_buf2; - TEST_CYCLE() cv::gpu::morphologyEx(d_src, dst, morphOp, ker, d_buf1, d_buf2); + cv::Ptr morph = cv::gpu::createMorphologyFilter(morphOp, src.type(), ker); + + TEST_CYCLE() morph->apply(d_src, dst); GPU_SANITY_CHECK(dst); } diff --git a/modules/gpufilters/src/filtering.cpp b/modules/gpufilters/src/filtering.cpp index c4de574408..145e049842 100644 --- a/modules/gpufilters/src/filtering.cpp +++ b/modules/gpufilters/src/filtering.cpp @@ -61,6 +61,8 @@ Ptr cv::gpu::createScharrFilter(int, int, int, int, double, int, int) { Ptr cv::gpu::createGaussianFilter(int, int, Size, double, double, int, int) { throw_no_cuda(); return Ptr(); } +Ptr cv::gpu::createMorphologyFilter(int, int, InputArray, Point, int) { throw_no_cuda(); return Ptr(); } + @@ -68,18 +70,9 @@ Ptr cv::gpu::createGaussianFilter(int, int, Size, double, double, int, i 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::getMorphologyFilter_GPU(int, int, const Mat&, const Size&, Point) { throw_no_cuda(); return Ptr(0); } -Ptr cv::gpu::createMorphologyFilter_GPU(int, int, const Mat&, const Point&, int) { throw_no_cuda(); return Ptr(0); } -Ptr cv::gpu::createMorphologyFilter_GPU(int, int, const Mat&, GpuMat&, const Point&, 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); } -void cv::gpu::erode(const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_no_cuda(); } -void cv::gpu::erode(const GpuMat&, GpuMat&, const Mat&, GpuMat&, Point, int, Stream&) { throw_no_cuda(); } -void cv::gpu::dilate(const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_no_cuda(); } -void cv::gpu::dilate(const GpuMat&, GpuMat&, const Mat&, GpuMat&, Point, int, Stream&) { throw_no_cuda(); } -void cv::gpu::morphologyEx(const GpuMat&, GpuMat&, int, const Mat&, Point, int) { throw_no_cuda(); } -void cv::gpu::morphologyEx(const GpuMat&, GpuMat&, int, const Mat&, GpuMat&, GpuMat&, Point, int, Stream&) { throw_no_cuda(); } #else @@ -506,6 +499,297 @@ Ptr cv::gpu::createGaussianFilter(int srcType, int dstType, Size ksize, return createSeparableLinearFilter(srcType, dstType, kx, ky, Point(-1,-1), rowBorderMode, columnBorderMode); } +//////////////////////////////////////////////////////////////////////////////////////////////////// +// Morphology Filter + +namespace +{ + class MorphologyFilter : public Filter + { + public: + MorphologyFilter(int op, int srcType, InputArray kernel, Point anchor, int iterations); + + void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); + + private: + typedef NppStatus (*nppMorfFilter_t)(const Npp8u* pSrc, Npp32s nSrcStep, Npp8u* pDst, Npp32s nDstStep, NppiSize oSizeROI, + const Npp8u* pMask, NppiSize oMaskSize, NppiPoint oAnchor); + + int type_; + GpuMat kernel_; + Point anchor_; + int iters_; + nppMorfFilter_t func_; + + GpuMat srcBorder_; + GpuMat buf_; + }; + + MorphologyFilter::MorphologyFilter(int op, int srcType, InputArray _kernel, Point anchor, int iterations) : + type_(srcType), anchor_(anchor), iters_(iterations) + { + static const nppMorfFilter_t funcs[2][5] = + { + {0, nppiErode_8u_C1R, 0, 0, nppiErode_8u_C4R }, + {0, nppiDilate_8u_C1R, 0, 0, nppiDilate_8u_C4R } + }; + + CV_Assert( op == MORPH_ERODE || op == MORPH_DILATE ); + CV_Assert( srcType == CV_8UC1 || srcType == CV_8UC4 ); + + Mat kernel = _kernel.getMat(); + Size ksize = !kernel.empty() ? _kernel.size() : Size(3, 3); + + normalizeAnchor(anchor_, ksize); + + if (kernel.empty()) + { + kernel = getStructuringElement(MORPH_RECT, Size(1 + iters_ * 2, 1 + iters_ * 2)); + anchor_ = Point(iters_, iters_); + iters_ = 1; + } + else if (iters_ > 1 && countNonZero(kernel) == (int) kernel.total()) + { + anchor_ = Point(anchor_.x * iters_, anchor_.y * iters_); + kernel = getStructuringElement(MORPH_RECT, + Size(ksize.width + (iters_ - 1) * (ksize.width - 1), + ksize.height + (iters_ - 1) * (ksize.height - 1)), + anchor_); + iters_ = 1; + } + + CV_Assert( kernel.channels() == 1 ); + + Mat kernel8U; + kernel.convertTo(kernel8U, CV_8U); + + kernel_ = gpu::createContinuous(kernel.size(), CV_8UC1); + kernel_.upload(kernel8U); + + func_ = funcs[op][CV_MAT_CN(srcType)]; + } + + void MorphologyFilter::apply(InputArray _src, OutputArray _dst, Stream& _stream) + { + GpuMat src = _src.getGpuMat(); + CV_Assert( src.type() == type_ ); + + Size ksize = kernel_.size(); + gpu::copyMakeBorder(src, srcBorder_, ksize.height, ksize.height, ksize.width, ksize.width, BORDER_DEFAULT, Scalar(), _stream); + + GpuMat srcRoi = srcBorder_(Rect(ksize.width, ksize.height, src.cols, src.rows)); + + GpuMat bufRoi; + if (iters_ > 1) + { + ensureSizeIsEnough(srcBorder_.size(), type_, buf_); + buf_.setTo(Scalar::all(0), _stream); + bufRoi = buf_(Rect(ksize.width, ksize.height, src.cols, src.rows)); + } + + _dst.create(src.size(), src.type()); + GpuMat dst = _dst.getGpuMat(); + + 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, kernel_.ptr(), oMaskSize, oAnchor) ); + + for(int i = 1; i < iters_; ++i) + { + dst.copyTo(bufRoi, _stream); + + nppSafeCall( func_(bufRoi.ptr(), static_cast(bufRoi.step), dst.ptr(), static_cast(dst.step), + oSizeROI, kernel_.ptr(), oMaskSize, oAnchor) ); + } + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } +} + +namespace +{ + class MorphologyExFilter : public Filter + { + public: + MorphologyExFilter(int srcType, InputArray kernel, Point anchor, int iterations); + + protected: + Ptr erodeFilter_, dilateFilter_; + GpuMat buf_; + }; + + MorphologyExFilter::MorphologyExFilter(int srcType, InputArray kernel, Point anchor, int iterations) + { + erodeFilter_ = gpu::createMorphologyFilter(MORPH_ERODE, srcType, kernel, anchor, iterations); + dilateFilter_ = gpu::createMorphologyFilter(MORPH_DILATE, srcType, kernel, anchor, iterations); + } + + // MORPH_OPEN + + class MorphologyOpenFilter : public MorphologyExFilter + { + public: + MorphologyOpenFilter(int srcType, InputArray kernel, Point anchor, int iterations); + + void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); + }; + + MorphologyOpenFilter::MorphologyOpenFilter(int srcType, InputArray kernel, Point anchor, int iterations) : + MorphologyExFilter(srcType, kernel, anchor, iterations) + { + } + + void MorphologyOpenFilter::apply(InputArray src, OutputArray dst, Stream& stream) + { + erodeFilter_->apply(src, buf_, stream); + dilateFilter_->apply(buf_, dst, stream); + } + + // MORPH_CLOSE + + class MorphologyCloseFilter : public MorphologyExFilter + { + public: + MorphologyCloseFilter(int srcType, InputArray kernel, Point anchor, int iterations); + + void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); + }; + + MorphologyCloseFilter::MorphologyCloseFilter(int srcType, InputArray kernel, Point anchor, int iterations) : + MorphologyExFilter(srcType, kernel, anchor, iterations) + { + } + + void MorphologyCloseFilter::apply(InputArray src, OutputArray dst, Stream& stream) + { + dilateFilter_->apply(src, buf_, stream); + erodeFilter_->apply(buf_, dst, stream); + } + + // MORPH_GRADIENT + + class MorphologyGradientFilter : public MorphologyExFilter + { + public: + MorphologyGradientFilter(int srcType, InputArray kernel, Point anchor, int iterations); + + void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); + }; + + MorphologyGradientFilter::MorphologyGradientFilter(int srcType, InputArray kernel, Point anchor, int iterations) : + MorphologyExFilter(srcType, kernel, anchor, iterations) + { + } + + void MorphologyGradientFilter::apply(InputArray src, OutputArray dst, Stream& stream) + { + erodeFilter_->apply(src, buf_, stream); + dilateFilter_->apply(src, dst, stream); + gpu::subtract(dst, buf_, dst, noArray(), -1, stream); + } + + // MORPH_TOPHAT + + class MorphologyTophatFilter : public MorphologyExFilter + { + public: + MorphologyTophatFilter(int srcType, InputArray kernel, Point anchor, int iterations); + + void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); + }; + + MorphologyTophatFilter::MorphologyTophatFilter(int srcType, InputArray kernel, Point anchor, int iterations) : + MorphologyExFilter(srcType, kernel, anchor, iterations) + { + } + + void MorphologyTophatFilter::apply(InputArray src, OutputArray dst, Stream& stream) + { + erodeFilter_->apply(src, dst, stream); + dilateFilter_->apply(dst, buf_, stream); + gpu::subtract(src, buf_, dst, noArray(), -1, stream); + } + + // MORPH_BLACKHAT + + class MorphologyBlackhatFilter : public MorphologyExFilter + { + public: + MorphologyBlackhatFilter(int srcType, InputArray kernel, Point anchor, int iterations); + + void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); + }; + + MorphologyBlackhatFilter::MorphologyBlackhatFilter(int srcType, InputArray kernel, Point anchor, int iterations) : + MorphologyExFilter(srcType, kernel, anchor, iterations) + { + } + + void MorphologyBlackhatFilter::apply(InputArray src, OutputArray dst, Stream& stream) + { + dilateFilter_->apply(src, dst, stream); + erodeFilter_->apply(dst, buf_, stream); + gpu::subtract(buf_, src, dst, noArray(), -1, stream); + } +} + +Ptr cv::gpu::createMorphologyFilter(int op, int srcType, InputArray kernel, Point anchor, int iterations) +{ + switch( op ) + { + case MORPH_ERODE: + case MORPH_DILATE: + return new MorphologyFilter(op, srcType, kernel, anchor, iterations); + break; + + case MORPH_OPEN: + return new MorphologyOpenFilter(srcType, kernel, anchor, iterations); + break; + + case MORPH_CLOSE: + return new MorphologyCloseFilter(srcType, kernel, anchor, iterations); + break; + + case MORPH_GRADIENT: + return new MorphologyGradientFilter(srcType, kernel, anchor, iterations); + break; + + case MORPH_TOPHAT: + return new MorphologyTophatFilter(srcType, kernel, anchor, iterations); + break; + + case MORPH_BLACKHAT: + return new MorphologyBlackhatFilter(srcType, kernel, anchor, iterations); + break; + + default: + CV_Error(Error::StsBadArg, "Unknown morphological operation"); + return Ptr(); + } +} + + + + + + + + @@ -638,264 +922,7 @@ Ptr cv::gpu::getColumnSumFilter_GPU(int sumType, int dstTy return Ptr(new NppColumnSumFilter(ksize, anchor)); } -//////////////////////////////////////////////////////////////////////////////////////////////////// -// Morphology Filter -namespace -{ - typedef NppStatus (*nppMorfFilter_t)(const Npp8u*, Npp32s, Npp8u*, Npp32s, NppiSize, const Npp8u*, NppiSize, NppiPoint); - - struct NPPMorphFilter : public BaseFilter_GPU - { - NPPMorphFilter(const Size& ksize_, const Point& anchor_, const GpuMat& kernel_, nppMorfFilter_t func_) : - BaseFilter_GPU(ksize_, anchor_), kernel(kernel_), 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, kernel.ptr(), oKernelSize, oAnchor) ); - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } - - GpuMat kernel; - nppMorfFilter_t func; - }; -} - -Ptr cv::gpu::getMorphologyFilter_GPU(int op, int type, const Mat& kernel, const Size& ksize, Point anchor) -{ - static const nppMorfFilter_t nppMorfFilter_callers[2][5] = - { - {0, nppiErode_8u_C1R, 0, 0, nppiErode_8u_C4R }, - {0, nppiDilate_8u_C1R, 0, 0, nppiDilate_8u_C4R } - }; - - CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE); - CV_Assert(type == CV_8UC1 || type == CV_8UC4); - - GpuMat gpu_krnl; - normalizeKernel(kernel, gpu_krnl); - normalizeAnchor(anchor, ksize); - - return Ptr(new NPPMorphFilter(ksize, anchor, gpu_krnl, nppMorfFilter_callers[op][CV_MAT_CN(type)])); -} - -namespace -{ - struct MorphologyFilterEngine_GPU : public FilterEngine_GPU - { - MorphologyFilterEngine_GPU(const Ptr& filter2D_, int type_, int iters_) : - filter2D(filter2D_), type(type_), iters(iters_) - { - pbuf = &buf; - } - - MorphologyFilterEngine_GPU(const Ptr& filter2D_, int type_, int iters_, GpuMat& buf_) : - filter2D(filter2D_), type(type_), iters(iters_) - { - pbuf = &buf_; - } - - virtual void apply(const GpuMat& src, GpuMat& dst, Rect roi = Rect(0,0,-1,-1), Stream& stream = Stream::Null()) - { - CV_Assert(src.type() == type); - - Size src_size = src.size(); - - dst.create(src_size, type); - - if (roi.size() != src_size) - { - dst.setTo(Scalar::all(0), stream); - } - - normalizeROI(roi, filter2D->ksize, filter2D->anchor, src_size); - - if (iters > 1) - pbuf->create(src_size, type); - - GpuMat srcROI = src(roi); - GpuMat dstROI = dst(roi); - - (*filter2D)(srcROI, dstROI, stream); - - for(int i = 1; i < iters; ++i) - { - dst.swap((*pbuf)); - - dstROI = dst(roi); - GpuMat bufROI = (*pbuf)(roi); - - (*filter2D)(bufROI, dstROI, stream); - } - } - - Ptr filter2D; - - int type; - int iters; - - GpuMat buf; - GpuMat* pbuf; - }; -} - -Ptr cv::gpu::createMorphologyFilter_GPU(int op, int type, const Mat& kernel, const Point& anchor, int iterations) -{ - CV_Assert(iterations > 0); - - Size ksize = kernel.size(); - - Ptr filter2D = getMorphologyFilter_GPU(op, type, kernel, ksize, anchor); - - return Ptr(new MorphologyFilterEngine_GPU(filter2D, type, iterations)); -} - -Ptr cv::gpu::createMorphologyFilter_GPU(int op, int type, const Mat& kernel, GpuMat& buf, const Point& anchor, int iterations) -{ - CV_Assert(iterations > 0); - - Size ksize = kernel.size(); - - Ptr filter2D = getMorphologyFilter_GPU(op, type, kernel, ksize, anchor); - - return Ptr(new MorphologyFilterEngine_GPU(filter2D, type, iterations, buf)); -} - -namespace -{ - void morphOp(int op, const GpuMat& src, GpuMat& dst, const Mat& _kernel, GpuMat& buf, Point anchor, int iterations, Stream& stream = Stream::Null()) - { - Mat kernel; - Size ksize = _kernel.data ? _kernel.size() : Size(3, 3); - - normalizeAnchor(anchor, ksize); - - if (iterations == 0 || _kernel.rows * _kernel.cols == 1) - { - src.copyTo(dst, stream); - return; - } - - dst.create(src.size(), src.type()); - - if (!_kernel.data) - { - kernel = getStructuringElement(MORPH_RECT, Size(1 + iterations * 2, 1 + iterations * 2)); - anchor = Point(iterations, iterations); - iterations = 1; - } - else if (iterations > 1 && countNonZero(_kernel) == _kernel.rows * _kernel.cols) - { - anchor = Point(anchor.x * iterations, anchor.y * iterations); - kernel = getStructuringElement(MORPH_RECT, - Size(ksize.width + (iterations - 1) * (ksize.width - 1), - ksize.height + (iterations - 1) * (ksize.height - 1)), - anchor); - iterations = 1; - } - else - kernel = _kernel; - - Ptr f = createMorphologyFilter_GPU(op, src.type(), kernel, buf, anchor, iterations); - - f->apply(src, dst, Rect(0,0,-1,-1), stream); - } - - void morphOp(int op, const GpuMat& src, GpuMat& dst, const Mat& _kernel, Point anchor, int iterations) - { - GpuMat buf; - morphOp(op, src, dst, _kernel, buf, anchor, iterations); - } -} - -void cv::gpu::erode( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations) -{ - morphOp(MORPH_ERODE, src, dst, kernel, anchor, iterations); -} - -void cv::gpu::erode( const GpuMat& src, GpuMat& dst, const Mat& kernel, GpuMat& buf, Point anchor, int iterations, Stream& stream) -{ - morphOp(MORPH_ERODE, src, dst, kernel, buf, anchor, iterations, stream); -} - -void cv::gpu::dilate( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations) -{ - morphOp(MORPH_DILATE, src, dst, kernel, anchor, iterations); -} - -void cv::gpu::dilate( const GpuMat& src, GpuMat& dst, const Mat& kernel, GpuMat& buf, Point anchor, int iterations, Stream& stream) -{ - morphOp(MORPH_DILATE, src, dst, kernel, buf, anchor, iterations, stream); -} - -void cv::gpu::morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, Point anchor, int iterations) -{ - GpuMat buf1; - GpuMat buf2; - morphologyEx(src, dst, op, kernel, buf1, buf2, anchor, iterations); -} - -void cv::gpu::morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, GpuMat& buf1, GpuMat& buf2, Point anchor, int iterations, Stream& stream) -{ - switch( op ) - { - case MORPH_ERODE: - erode(src, dst, kernel, buf1, anchor, iterations, stream); - break; - - case MORPH_DILATE: - dilate(src, dst, kernel, buf1, anchor, iterations, stream); - break; - - case MORPH_OPEN: - erode(src, buf2, kernel, buf1, anchor, iterations, stream); - dilate(buf2, dst, kernel, buf1, anchor, iterations, stream); - break; - - case MORPH_CLOSE: - dilate(src, buf2, kernel, buf1, anchor, iterations, stream); - erode(buf2, dst, kernel, buf1, anchor, iterations, stream); - break; - - case MORPH_GRADIENT: - erode(src, buf2, kernel, buf1, anchor, iterations, stream); - dilate(src, dst, kernel, buf1, anchor, iterations, stream); - gpu::subtract(dst, buf2, dst, GpuMat(), -1, stream); - break; - - case MORPH_TOPHAT: - erode(src, dst, kernel, buf1, anchor, iterations, stream); - dilate(dst, buf2, kernel, buf1, anchor, iterations, stream); - gpu::subtract(src, buf2, dst, GpuMat(), -1, stream); - break; - - case MORPH_BLACKHAT: - dilate(src, dst, kernel, buf1, anchor, iterations, stream); - erode(dst, buf2, kernel, buf1, anchor, iterations, stream); - gpu::subtract(buf2, src, dst, GpuMat(), -1, stream); - break; - - default: - CV_Error(cv::Error::StsBadArg, "unknown morphological operation"); - } -} //////////////////////////////////////////////////////////////////////////////////////////////////// // Image Rank Filter diff --git a/modules/gpufilters/test/test_filters.cpp b/modules/gpufilters/test/test_filters.cpp index e052ede6cb..03bea05e6d 100644 --- a/modules/gpufilters/test/test_filters.cpp +++ b/modules/gpufilters/test/test_filters.cpp @@ -489,15 +489,6 @@ INSTANTIATE_TEST_CASE_P(GPU_Filters, GaussianBlur, testing::Combine( BorderType(cv::BORDER_REFLECT)), WHOLE_SUBMAT)); - - - - - - - - - ///////////////////////////////////////////////////////////////////////////////////////////////// // Erode @@ -528,8 +519,10 @@ GPU_TEST_P(Erode, Accuracy) cv::Mat src = randomMat(size, type); cv::Mat kernel = cv::Mat::ones(3, 3, CV_8U); + cv::Ptr erode = cv::gpu::createMorphologyFilter(cv::MORPH_ERODE, src.type(), kernel, anchor, iterations); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); - cv::gpu::erode(loadMat(src, useRoi), dst, kernel, anchor, iterations); + erode->apply(loadMat(src, useRoi), dst); cv::Mat dst_gold; cv::erode(src, dst_gold, kernel, anchor, iterations); @@ -577,8 +570,10 @@ GPU_TEST_P(Dilate, Accuracy) cv::Mat src = randomMat(size, type); cv::Mat kernel = cv::Mat::ones(3, 3, CV_8U); + cv::Ptr dilate = cv::gpu::createMorphologyFilter(cv::MORPH_DILATE, src.type(), kernel, anchor, iterations); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); - cv::gpu::dilate(loadMat(src, useRoi), dst, kernel, anchor, iterations); + dilate->apply(loadMat(src, useRoi), dst); cv::Mat dst_gold; cv::dilate(src, dst_gold, kernel, anchor, iterations); @@ -630,8 +625,10 @@ GPU_TEST_P(MorphEx, Accuracy) cv::Mat src = randomMat(size, type); cv::Mat kernel = cv::Mat::ones(3, 3, CV_8U); + cv::Ptr morph = cv::gpu::createMorphologyFilter(morphOp, src.type(), kernel, anchor, iterations); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); - cv::gpu::morphologyEx(loadMat(src, useRoi), dst, morphOp, kernel, anchor, iterations); + morph->apply(loadMat(src, useRoi), dst); cv::Mat dst_gold; cv::morphologyEx(src, dst_gold, morphOp, kernel, anchor, iterations); diff --git a/samples/gpu/morphology.cpp b/samples/gpu/morphology.cpp index 1ed8f96dc9..abc6c980b0 100644 --- a/samples/gpu/morphology.cpp +++ b/samples/gpu/morphology.cpp @@ -1,120 +1,186 @@ +#include -#include "opencv2/imgproc/imgproc.hpp" -#include "opencv2/highgui/highgui.hpp" -#include "opencv2/gpu/gpu.hpp" -#include -#include +#include "opencv2/imgproc.hpp" +#include "opencv2/highgui.hpp" +#include "opencv2/gpufilters.hpp" +#include "opencv2/gpuimgproc.hpp" using namespace std; using namespace cv; -using namespace cv::gpu; -static void help() +class App { +public: + App(int argc, const char* argv[]); -printf("\nShow off image morphology: erosion, dialation, open and close\n" - "Call:\n morphology2 [image]\n" - "This program also shows use of rect, elipse and cross kernels\n\n"); -printf( "Hot keys: \n" - "\tESC - quit the program\n" - "\tr - use rectangle structuring element\n" - "\te - use elliptic structuring element\n" - "\tc - use cross-shaped structuring element\n" - "\tSPACE - loop through all the options\n" ); -} + int run(); -GpuMat src, dst; +private: + void help(); -int element_shape = MORPH_RECT; + void OpenClose(); + void ErodeDilate(); -//the address of variable which receives trackbar position update -int max_iters = 10; -int open_close_pos = 0; -int erode_dilate_pos = 0; + static void OpenCloseCallback(int, void*); + static void ErodeDilateCallback(int, void*); -// callback function for open/close trackbar -static void OpenClose(int, void*) + gpu::GpuMat src, dst; + + int element_shape; + + int max_iters; + int open_close_pos; + int erode_dilate_pos; +}; + +App::App(int argc, const char* argv[]) { - int n = open_close_pos - max_iters; - int an = n > 0 ? n : -n; - Mat element = getStructuringElement(element_shape, Size(an*2+1, an*2+1), Point(an, an) ); - if( n < 0 ) - cv::gpu::morphologyEx(src, dst, MORPH_OPEN, element); - else - cv::gpu::morphologyEx(src, dst, MORPH_CLOSE, element); - imshow("Open/Close",(Mat)dst); -} + element_shape = MORPH_RECT; + open_close_pos = erode_dilate_pos = max_iters = 10; -// callback function for erode/dilate trackbar -static void ErodeDilate(int, void*) -{ - int n = erode_dilate_pos - max_iters; - int an = n > 0 ? n : -n; - Mat element = getStructuringElement(element_shape, Size(an*2+1, an*2+1), Point(an, an) ); - if( n < 0 ) - cv::gpu::erode(src, dst, element); - else - cv::gpu::dilate(src, dst, element); - imshow("Erode/Dilate",(Mat)dst); -} - - -int main( int argc, char** argv ) -{ - char* filename = argc == 2 ? argv[1] : (char*)"baboon.jpg"; - if (string(argv[1]) == "--help") + if (argc == 2 && String(argv[1]) == "--help") { help(); - return -1; + exit(0); } - src.upload(imread(filename, 1)); - if (src.empty()) + String filename = argc == 2 ? argv[1] : "baboon.jpg"; + + Mat img = imread(filename); + if (img.empty()) { - help(); - return -1; + cerr << "Can't open image " << filename.c_str() << endl; + exit(-1); } - cv::gpu::printShortCudaDeviceInfo(cv::gpu::getDevice()); - - help(); - - + src.upload(img); if (src.channels() == 3) { // gpu support only 4th channel images - GpuMat src4ch; - cv::gpu::cvtColor(src, src4ch, COLOR_BGR2BGRA); + gpu::GpuMat src4ch; + gpu::cvtColor(src, src4ch, COLOR_BGR2BGRA); src = src4ch; } - //create windows for output images - namedWindow("Open/Close",1); - namedWindow("Erode/Dilate",1); + help(); - open_close_pos = erode_dilate_pos = max_iters; - createTrackbar("iterations", "Open/Close",&open_close_pos,max_iters*2+1,OpenClose); - createTrackbar("iterations", "Erode/Dilate",&erode_dilate_pos,max_iters*2+1,ErodeDilate); + gpu::printShortCudaDeviceInfo(gpu::getDevice()); +} + +int App::run() +{ + // create windows for output images + namedWindow("Open/Close"); + namedWindow("Erode/Dilate"); + + createTrackbar("iterations", "Open/Close", &open_close_pos, max_iters * 2 + 1, OpenCloseCallback, this); + createTrackbar("iterations", "Erode/Dilate", &erode_dilate_pos, max_iters * 2 + 1, ErodeDilateCallback, this); for(;;) { - int c; + OpenClose(); + ErodeDilate(); - OpenClose(open_close_pos, 0); - ErodeDilate(erode_dilate_pos, 0); - c = waitKey(); + char c = (char) waitKey(); - if( (char)c == 27 ) + switch (c) + { + case 27: + return 0; break; - if( (char)c == 'e' ) + + case 'e': element_shape = MORPH_ELLIPSE; - else if( (char)c == 'r' ) + break; + + case 'r': element_shape = MORPH_RECT; - else if( (char)c == 'c' ) + break; + + case 'c': element_shape = MORPH_CROSS; - else if( (char)c == ' ' ) + break; + + case ' ': element_shape = (element_shape + 1) % 3; + break; + } + } +} + +void App::help() +{ + cout << "Show off image morphology: erosion, dialation, open and close \n"; + cout << "Call: \n"; + cout << " gpu-example-morphology [image] \n"; + cout << "This program also shows use of rect, elipse and cross kernels \n" << endl; + + cout << "Hot keys: \n"; + cout << "\tESC - quit the program \n"; + cout << "\tr - use rectangle structuring element \n"; + cout << "\te - use elliptic structuring element \n"; + cout << "\tc - use cross-shaped structuring element \n"; + cout << "\tSPACE - loop through all the options \n" << endl; +} + +void App::OpenClose() +{ + int n = open_close_pos - max_iters; + int an = n > 0 ? n : -n; + + Mat element = getStructuringElement(element_shape, Size(an*2+1, an*2+1), Point(an, an)); + + if (n < 0) + { + Ptr openFilter = gpu::createMorphologyFilter(MORPH_OPEN, src.type(), element); + openFilter->apply(src, dst); + } + else + { + Ptr closeFilter = gpu::createMorphologyFilter(MORPH_CLOSE, src.type(), element); + closeFilter->apply(src, dst); } - return 0; + Mat h_dst(dst); + imshow("Open/Close", h_dst); +} + +void App::ErodeDilate() +{ + int n = erode_dilate_pos - max_iters; + int an = n > 0 ? n : -n; + + Mat element = getStructuringElement(element_shape, Size(an*2+1, an*2+1), Point(an, an)); + + if (n < 0) + { + Ptr erodeFilter = gpu::createMorphologyFilter(MORPH_ERODE, src.type(), element); + erodeFilter->apply(src, dst); + } + else + { + Ptr dilateFilter = gpu::createMorphologyFilter(MORPH_DILATE, src.type(), element); + dilateFilter->apply(src, dst); + } + + Mat h_dst(dst); + imshow("Erode/Dilate", h_dst); +} + +void App::OpenCloseCallback(int, void* data) +{ + App* thiz = (App*) data; + thiz->OpenClose(); +} + +void App::ErodeDilateCallback(int, void* data) +{ + App* thiz = (App*) data; + thiz->ErodeDilate(); +} + +int main(int argc, const char* argv[]) +{ + App app(argc, argv); + return app.run(); } diff --git a/samples/gpu/performance/tests.cpp b/samples/gpu/performance/tests.cpp index cbe50a40f6..4333b76257 100644 --- a/samples/gpu/performance/tests.cpp +++ b/samples/gpu/performance/tests.cpp @@ -746,10 +746,12 @@ TEST(erode) d_src.upload(src); - gpu::erode(d_src, d_dst, ker, d_buf); + Ptr erode = gpu::createMorphologyFilter(MORPH_ERODE, d_src.type(), ker); + + erode->apply(d_src, d_dst); GPU_ON; - gpu::erode(d_src, d_dst, ker, d_buf); + erode->apply(d_src, d_dst); GPU_OFF; } }