Chiebot-Mirror
/
opencv
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

refactored Morphology Filters

Browse Source
pull/1010/head
Vladislav Vinogradov 12 years ago
parent 12ae11e2ff
commit 5720eaf354
7 changed files with 500 additions and 406 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 16
      modules/gpubgsegm/src/fgd.cpp
  2. 61
      modules/gpufilters/include/opencv2/gpufilters.hpp
  3. 23
      modules/gpufilters/perf/perf_filters.cpp
  4. 557
      modules/gpufilters/src/filtering.cpp
  5. 21
      modules/gpufilters/test/test_filters.cpp
  6. 222
      samples/gpu/morphology.cpp
  7. 6
      samples/gpu/performance/tests.cpp

16
modules/gpubgsegm/src/fgd.cpp
Unescape View File

@ -228,11 +228,10 @@ private:
cv::gpu::GpuMat countBuf_;
cv::gpu::GpuMat buf_;
cv::gpu::GpuMat filterBuf_;
cv::gpu::GpuMat filterBrd_;
cv::Ptr<cv::gpu::FilterEngine_GPU> dilateFilter_;
cv::Ptr<cv::gpu::FilterEngine_GPU> erodeFilter_;
cv::Ptr<cv::gpu::Filter> dilateFilter_;
cv::Ptr<cv::gpu::Filter> 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<cv::gpu::FilterEngine_GPU>& filter, cv::Scalar brdVal)
void morphology(const cv::gpu::GpuMat& src, cv::gpu::GpuMat& dst, cv::gpu::GpuMat& filterBrd, int brd, cv::Ptr<cv::gpu::Filter>& 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<cv::gpu::FilterEngine_GPU>& erodeFilter, cv::Ptr<cv::gpu::FilterEngine_GPU>& dilateFilter,
cv::Ptr<cv::gpu::Filter>& erodeFilter, cv::Ptr<cv::gpu::Filter>& dilateFilter,
const cv::gpu::FGDStatModel::Params& params)
{
const int brd = params.perform_morphing;

61
modules/gpufilters/include/opencv2/gpufilters.hpp
Unescape View File

@ -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<Filter> 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<gpu::Filter> 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<gpu::Filter> 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<gpu::Filter> f = gpu::createMorphologyFilter(op, src.type(), kernel, anchor, iterations);
f->apply(src, dst, stream);
}
@ -285,18 +320,7 @@ CV_EXPORTS Ptr<BaseColumnFilter_GPU> 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<BaseFilter_GPU> 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<FilterEngine_GPU> createMorphologyFilter_GPU(int op, int type, const Mat& kernel,
const Point& anchor = Point(-1,-1), int iterations = 1);
CV_EXPORTS Ptr<FilterEngine_GPU> 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<BaseFilter_GPU> 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());

23
modules/gpufilters/perf/perf_filters.cpp
Unescape View File

@ -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<cv::gpu::Filter> 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<cv::gpu::Filter> 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<cv::gpu::Filter> morph = cv::gpu::createMorphologyFilter(morphOp, src.type(), ker);
TEST_CYCLE() morph->apply(d_src, dst);
GPU_SANITY_CHECK(dst);
}

557
modules/gpufilters/src/filtering.cpp
Unescape View File

@ -61,6 +61,8 @@ Ptr<Filter> cv::gpu::createScharrFilter(int, int, int, int, double, int, int) {
Ptr<Filter> cv::gpu::createGaussianFilter(int, int, Size, double, double, int, int) { throw_no_cuda(); return Ptr<Filter>(); }
Ptr<Filter> cv::gpu::createMorphologyFilter(int, int, InputArray, Point, int) { throw_no_cuda(); return Ptr<Filter>(); }
@ -68,18 +70,9 @@ Ptr<Filter> cv::gpu::createGaussianFilter(int, int, Size, double, double, int, i
Ptr<BaseRowFilter_GPU> cv::gpu::getRowSumFilter_GPU(int, int, int, int) { throw_no_cuda(); return Ptr<BaseRowFilter_GPU>(0); }
Ptr<BaseColumnFilter_GPU> cv::gpu::getColumnSumFilter_GPU(int, int, int, int) { throw_no_cuda(); return Ptr<BaseColumnFilter_GPU>(0); }
Ptr<BaseFilter_GPU> cv::gpu::getMorphologyFilter_GPU(int, int, const Mat&, const Size&, Point) { throw_no_cuda(); return Ptr<BaseFilter_GPU>(0); }
Ptr<FilterEngine_GPU> cv::gpu::createMorphologyFilter_GPU(int, int, const Mat&, const Point&, int) { throw_no_cuda(); return Ptr<FilterEngine_GPU>(0); }
Ptr<FilterEngine_GPU> cv::gpu::createMorphologyFilter_GPU(int, int, const Mat&, GpuMat&, const Point&, int) { throw_no_cuda(); return Ptr<FilterEngine_GPU>(0); }
Ptr<BaseFilter_GPU> cv::gpu::getMaxFilter_GPU(int, int, const Size&, Point) { throw_no_cuda(); return Ptr<BaseFilter_GPU>(0); }
Ptr<BaseFilter_GPU> cv::gpu::getMinFilter_GPU(int, int, const Size&, Point) { throw_no_cuda(); return Ptr<BaseFilter_GPU>(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,397 +499,431 @@ Ptr<Filter> 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<Npp8u>(), static_cast<int>(srcRoi.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step),
oSizeROI, kernel_.ptr<Npp8u>(), oMaskSize, oAnchor) );
for(int i = 1; i < iters_; ++i)
{
dst.copyTo(bufRoi, _stream);
nppSafeCall( func_(bufRoi.ptr<Npp8u>(), static_cast<int>(bufRoi.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step),
oSizeROI, kernel_.ptr<Npp8u>(), oMaskSize, oAnchor) );
}
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
}
namespace
{
inline void normalizeROI(Rect& roi, const Size& ksize, const Point& anchor, const Size& src_size)
class MorphologyExFilter : public Filter
{
if (roi == Rect(0,0,-1,-1))
roi = Rect(anchor.x, anchor.y, src_size.width - ksize.width, src_size.height - ksize.height);
public:
MorphologyExFilter(int srcType, InputArray kernel, Point anchor, int iterations);
CV_Assert(roi.x >= 0 && roi.y >= 0 && roi.width <= src_size.width && roi.height <= src_size.height);
protected:
Ptr<gpu::Filter> 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);
}
inline void normalizeKernel(const Mat& kernel, GpuMat& gpu_krnl, int type = CV_8U, int* nDivisor = 0, bool reverse = false)
// MORPH_OPEN
class MorphologyOpenFilter : public MorphologyExFilter
{
int scale = nDivisor && (kernel.depth() == CV_32F || kernel.depth() == CV_64F) ? 256 : 1;
if (nDivisor) *nDivisor = scale;
public:
MorphologyOpenFilter(int srcType, InputArray kernel, Point anchor, int iterations);
Mat temp(kernel.size(), type);
kernel.convertTo(temp, type, scale);
Mat cont_krnl = temp.reshape(1, 1);
void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null());
};
if (reverse)
{
int count = cont_krnl.cols >> 1;
for (int i = 0; i < count; ++i)
{
std::swap(cont_krnl.at<int>(0, i), cont_krnl.at<int>(0, cont_krnl.cols - 1 - i));
}
}
MorphologyOpenFilter::MorphologyOpenFilter(int srcType, InputArray kernel, Point anchor, int iterations) :
MorphologyExFilter(srcType, kernel, anchor, iterations)
{
}
gpu_krnl.upload(cont_krnl);
void MorphologyOpenFilter::apply(InputArray src, OutputArray dst, Stream& stream)
{
erodeFilter_->apply(src, buf_, stream);
dilateFilter_->apply(buf_, dst, stream);
}
}
////////////////////////////////////////////////////////////////////////////////////////////////////
// 1D Sum Filter
// MORPH_CLOSE
namespace
{
struct NppRowSumFilter : public BaseRowFilter_GPU
class MorphologyCloseFilter : public MorphologyExFilter
{
NppRowSumFilter(int ksize_, int anchor_) : BaseRowFilter_GPU(ksize_, anchor_) {}
public:
MorphologyCloseFilter(int srcType, InputArray kernel, Point anchor, int iterations);
virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& s = Stream::Null())
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null());
};
cudaStream_t stream = StreamAccessor::getStream(s);
MorphologyCloseFilter::MorphologyCloseFilter(int srcType, InputArray kernel, Point anchor, int iterations) :
MorphologyExFilter(srcType, kernel, anchor, iterations)
{
}
NppStreamHandler h(stream);
void MorphologyCloseFilter::apply(InputArray src, OutputArray dst, Stream& stream)
{
dilateFilter_->apply(src, buf_, stream);
erodeFilter_->apply(buf_, dst, stream);
}
nppSafeCall( nppiSumWindowRow_8u32f_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step),
dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz, ksize, anchor) );
// MORPH_GRADIENT
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
class MorphologyGradientFilter : public MorphologyExFilter
{
public:
MorphologyGradientFilter(int srcType, InputArray kernel, Point anchor, int iterations);
void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null());
};
}
Ptr<BaseRowFilter_GPU> cv::gpu::getRowSumFilter_GPU(int srcType, int sumType, int ksize, int anchor)
{
CV_Assert(srcType == CV_8UC1 && sumType == CV_32FC1);
MorphologyGradientFilter::MorphologyGradientFilter(int srcType, InputArray kernel, Point anchor, int iterations) :
MorphologyExFilter(srcType, kernel, anchor, iterations)
{
}
normalizeAnchor(anchor, ksize);
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);
}
return Ptr<BaseRowFilter_GPU>(new NppRowSumFilter(ksize, anchor));
}
// MORPH_TOPHAT
namespace
{
struct NppColumnSumFilter : public BaseColumnFilter_GPU
class MorphologyTophatFilter : public MorphologyExFilter
{
NppColumnSumFilter(int ksize_, int anchor_) : BaseColumnFilter_GPU(ksize_, anchor_) {}
public:
MorphologyTophatFilter(int srcType, InputArray kernel, Point anchor, int iterations);
virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& s = Stream::Null())
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null());
};
cudaStream_t stream = StreamAccessor::getStream(s);
MorphologyTophatFilter::MorphologyTophatFilter(int srcType, InputArray kernel, Point anchor, int iterations) :
MorphologyExFilter(srcType, kernel, anchor, iterations)
{
}
NppStreamHandler h(stream);
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);
}
nppSafeCall( nppiSumWindowColumn_8u32f_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step),
dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz, ksize, anchor) );
// MORPH_BLACKHAT
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
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<BaseColumnFilter_GPU> cv::gpu::getColumnSumFilter_GPU(int sumType, int dstType, int ksize, int anchor)
Ptr<Filter> cv::gpu::createMorphologyFilter(int op, int srcType, InputArray kernel, Point anchor, int iterations)
{
CV_Assert(sumType == CV_8UC1 && dstType == CV_32FC1);
switch( op )
{
case MORPH_ERODE:
case MORPH_DILATE:
return new MorphologyFilter(op, srcType, kernel, anchor, iterations);
break;
normalizeAnchor(anchor, ksize);
case MORPH_OPEN:
return new MorphologyOpenFilter(srcType, kernel, anchor, iterations);
break;
return Ptr<BaseColumnFilter_GPU>(new NppColumnSumFilter(ksize, anchor));
}
case MORPH_CLOSE:
return new MorphologyCloseFilter(srcType, kernel, anchor, iterations);
break;
////////////////////////////////////////////////////////////////////////////////////////////////////
// Morphology Filter
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<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<Npp8u>(), static_cast<int>(src.step),
dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz, kernel.ptr<Npp8u>(), oKernelSize, oAnchor) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
GpuMat kernel;
nppMorfFilter_t func;
};
}
Ptr<BaseFilter_GPU> 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<BaseFilter_GPU>(new NPPMorphFilter(ksize, anchor, gpu_krnl, nppMorfFilter_callers[op][CV_MAT_CN(type)]));
}
namespace
{
struct MorphologyFilterEngine_GPU : public FilterEngine_GPU
{
MorphologyFilterEngine_GPU(const Ptr<BaseFilter_GPU>& filter2D_, int type_, int iters_) :
filter2D(filter2D_), type(type_), iters(iters_)
{
pbuf = &buf;
}
MorphologyFilterEngine_GPU(const Ptr<BaseFilter_GPU>& 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<BaseFilter_GPU> filter2D;
int type;
int iters;
GpuMat buf;
GpuMat* pbuf;
};
}
Ptr<FilterEngine_GPU> 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<BaseFilter_GPU> filter2D = getMorphologyFilter_GPU(op, type, kernel, ksize, anchor);
return Ptr<FilterEngine_GPU>(new MorphologyFilterEngine_GPU(filter2D, type, iterations));
}
Ptr<FilterEngine_GPU> 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<BaseFilter_GPU> filter2D = getMorphologyFilter_GPU(op, type, kernel, ksize, anchor);
return Ptr<FilterEngine_GPU>(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())
inline void normalizeROI(Rect& roi, const Size& ksize, const Point& anchor, const Size& src_size)
{
Mat kernel;
Size ksize = _kernel.data ? _kernel.size() : Size(3, 3);
if (roi == Rect(0,0,-1,-1))
roi = Rect(anchor.x, anchor.y, src_size.width - ksize.width, src_size.height - ksize.height);
normalizeAnchor(anchor, ksize);
CV_Assert(roi.x >= 0 && roi.y >= 0 && roi.width <= src_size.width && roi.height <= src_size.height);
}
if (iterations == 0 || _kernel.rows * _kernel.cols == 1)
{
src.copyTo(dst, stream);
return;
}
inline void normalizeKernel(const Mat& kernel, GpuMat& gpu_krnl, int type = CV_8U, int* nDivisor = 0, bool reverse = false)
{
int scale = nDivisor && (kernel.depth() == CV_32F || kernel.depth() == CV_64F) ? 256 : 1;
if (nDivisor) *nDivisor = scale;
dst.create(src.size(), src.type());
Mat temp(kernel.size(), type);
kernel.convertTo(temp, type, scale);
Mat cont_krnl = temp.reshape(1, 1);
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)
if (reverse)
{
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;
int count = cont_krnl.cols >> 1;
for (int i = 0; i < count; ++i)
{
std::swap(cont_krnl.at<int>(0, i), cont_krnl.at<int>(0, cont_krnl.cols - 1 - i));
}
}
else
kernel = _kernel;
Ptr<FilterEngine_GPU> 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);
gpu_krnl.upload(cont_krnl);
}
}
void cv::gpu::erode( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations)
{
morphOp(MORPH_ERODE, src, dst, kernel, anchor, iterations);
}
////////////////////////////////////////////////////////////////////////////////////////////////////
// 1D Sum Filter
void cv::gpu::erode( const GpuMat& src, GpuMat& dst, const Mat& kernel, GpuMat& buf, Point anchor, int iterations, Stream& stream)
namespace
{
morphOp(MORPH_ERODE, src, dst, kernel, buf, anchor, iterations, stream);
}
struct NppRowSumFilter : public BaseRowFilter_GPU
{
NppRowSumFilter(int ksize_, int anchor_) : BaseRowFilter_GPU(ksize_, anchor_) {}
void cv::gpu::dilate( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations)
{
morphOp(MORPH_DILATE, src, dst, kernel, anchor, iterations);
}
virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& s = Stream::Null())
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
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);
cudaStream_t stream = StreamAccessor::getStream(s);
NppStreamHandler h(stream);
nppSafeCall( nppiSumWindowRow_8u32f_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step),
dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz, ksize, anchor) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
}
void cv::gpu::morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, Point anchor, int iterations)
Ptr<BaseRowFilter_GPU> cv::gpu::getRowSumFilter_GPU(int srcType, int sumType, int ksize, int anchor)
{
GpuMat buf1;
GpuMat buf2;
morphologyEx(src, dst, op, kernel, buf1, buf2, anchor, iterations);
CV_Assert(srcType == CV_8UC1 && sumType == CV_32FC1);
normalizeAnchor(anchor, ksize);
return Ptr<BaseRowFilter_GPU>(new NppRowSumFilter(ksize, anchor));
}
void cv::gpu::morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, GpuMat& buf1, GpuMat& buf2, Point anchor, int iterations, Stream& stream)
namespace
{
switch( op )
struct NppColumnSumFilter : public BaseColumnFilter_GPU
{
case MORPH_ERODE:
erode(src, dst, kernel, buf1, anchor, iterations, stream);
break;
NppColumnSumFilter(int ksize_, int anchor_) : BaseColumnFilter_GPU(ksize_, anchor_) {}
case MORPH_DILATE:
dilate(src, dst, kernel, buf1, anchor, iterations, stream);
break;
virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& s = Stream::Null())
{
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
case MORPH_OPEN:
erode(src, buf2, kernel, buf1, anchor, iterations, stream);
dilate(buf2, dst, kernel, buf1, anchor, iterations, stream);
break;
cudaStream_t stream = StreamAccessor::getStream(s);
case MORPH_CLOSE:
dilate(src, buf2, kernel, buf1, anchor, iterations, stream);
erode(buf2, dst, kernel, buf1, anchor, iterations, stream);
break;
NppStreamHandler h(stream);
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;
nppSafeCall( nppiSumWindowColumn_8u32f_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step),
dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz, ksize, anchor) );
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;
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
};
}
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;
Ptr<BaseColumnFilter_GPU> cv::gpu::getColumnSumFilter_GPU(int sumType, int dstType, int ksize, int anchor)
{
CV_Assert(sumType == CV_8UC1 && dstType == CV_32FC1);
default:
CV_Error(cv::Error::StsBadArg, "unknown morphological operation");
}
normalizeAnchor(anchor, ksize);
return Ptr<BaseColumnFilter_GPU>(new NppColumnSumFilter(ksize, anchor));
}
////////////////////////////////////////////////////////////////////////////////////////////////////
// Image Rank Filter

21
modules/gpufilters/test/test_filters.cpp
Unescape View File

@ -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<cv::gpu::Filter> 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<cv::gpu::Filter> 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<cv::gpu::Filter> 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);

222
samples/gpu/morphology.cpp
Unescape View File

@ -1,120 +1,186 @@
#include <iostream>
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/gpu/gpu.hpp"
#include <stdlib.h>
#include <stdio.h>
#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*)
{
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);
}
gpu::GpuMat src, dst;
// 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 element_shape;
int max_iters;
int open_close_pos;
int erode_dilate_pos;
};
int main( int argc, char** argv )
App::App(int argc, const char* argv[])
{
char* filename = argc == 2 ? argv[1] : (char*)"baboon.jpg";
if (string(argv[1]) == "--help")
{
help();
return -1;
}
element_shape = MORPH_RECT;
open_close_pos = erode_dilate_pos = max_iters = 10;
src.upload(imread(filename, 1));
if (src.empty())
if (argc == 2 && String(argv[1]) == "--help")
{
help();
return -1;
exit(0);
}
cv::gpu::printShortCudaDeviceInfo(cv::gpu::getDevice());
help();
String filename = argc == 2 ? argv[1] : "baboon.jpg";
Mat img = imread(filename);
if (img.empty())
{
cerr << "Can't open image " << filename.c_str() << endl;
exit(-1);
}
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<gpu::Filter> openFilter = gpu::createMorphologyFilter(MORPH_OPEN, src.type(), element);
openFilter->apply(src, dst);
}
else
{
Ptr<gpu::Filter> closeFilter = gpu::createMorphologyFilter(MORPH_CLOSE, src.type(), element);
closeFilter->apply(src, dst);
}
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<gpu::Filter> erodeFilter = gpu::createMorphologyFilter(MORPH_ERODE, src.type(), element);
erodeFilter->apply(src, dst);
}
else
{
Ptr<gpu::Filter> dilateFilter = gpu::createMorphologyFilter(MORPH_DILATE, src.type(), element);
dilateFilter->apply(src, dst);
}
return 0;
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();
}

6
samples/gpu/performance/tests.cpp
Unescape View File

@ -746,10 +746,12 @@ TEST(erode)
d_src.upload(src);
gpu::erode(d_src, d_dst, ker, d_buf);
Ptr<gpu::Filter> 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;
}
}

Write Preview
Loading…
Cancel
Save