From: Vladislav Vinogradov Date: Mon, 29 Apr 2013 07:46:49 +0000 (+0400) Subject: refactored box filter X-Git-Tag: submit/tizen_ivi/20141117.190038~2^2~1027^2~8 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=fcfcd4cbced5e188ffd01758150e07b2210b25ab;p=profile%2Fivi%2Fopencv.git refactored box filter --- diff --git a/modules/gpubgsegm/include/opencv2/gpubgsegm.hpp b/modules/gpubgsegm/include/opencv2/gpubgsegm.hpp index e7a29b5..3fe62ec 100644 --- a/modules/gpubgsegm/include/opencv2/gpubgsegm.hpp +++ b/modules/gpubgsegm/include/opencv2/gpubgsegm.hpp @@ -321,7 +321,7 @@ private: GpuMat colors_; GpuMat weights_; - Ptr boxFilter_; + Ptr boxFilter_; GpuMat buf_; }; diff --git a/modules/gpubgsegm/src/gmg.cpp b/modules/gpubgsegm/src/gmg.cpp index a38cbff..b97f083 100644 --- a/modules/gpubgsegm/src/gmg.cpp +++ b/modules/gpubgsegm/src/gmg.cpp @@ -100,7 +100,7 @@ void cv::gpu::GMG_GPU::initialize(cv::Size frameSize, float min, float max) nfeatures_.setTo(cv::Scalar::all(0)); if (smoothingRadius > 0) - boxFilter_ = cv::gpu::createBoxFilter_GPU(CV_8UC1, CV_8UC1, cv::Size(smoothingRadius, smoothingRadius)); + boxFilter_ = cv::gpu::createBoxFilter(CV_8UC1, -1, cv::Size(smoothingRadius, smoothingRadius)); loadConstants(frameSize_.width, frameSize_.height, minVal_, maxVal_, quantizationLevels, backgroundPrior, decisionThreshold, maxFeatures, numInitializationFrames); } @@ -141,7 +141,7 @@ void cv::gpu::GMG_GPU::operator ()(const cv::gpu::GpuMat& frame, cv::gpu::GpuMat // medianBlur if (smoothingRadius > 0) { - boxFilter_->apply(fgmask, buf_, cv::Rect(0,0,-1,-1), stream); + boxFilter_->apply(fgmask, buf_, stream); int minCount = (smoothingRadius * smoothingRadius + 1) / 2; double thresh = 255.0 * minCount / (smoothingRadius * smoothingRadius); cv::gpu::threshold(buf_, fgmask, thresh, 255.0, cv::THRESH_BINARY, stream); diff --git a/modules/gpufilters/CMakeLists.txt b/modules/gpufilters/CMakeLists.txt index 18f6d7f..640de8c 100644 --- a/modules/gpufilters/CMakeLists.txt +++ b/modules/gpufilters/CMakeLists.txt @@ -6,4 +6,4 @@ set(the_description "GPU-accelerated Image Filtering") ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4324 /wd4512 -Wundef -Wmissing-declarations) -ocv_define_module(gpufilters opencv_imgproc OPTIONAL opencv_gpuarithm) +ocv_define_module(gpufilters opencv_imgproc opencv_gpuarithm) diff --git a/modules/gpufilters/include/opencv2/gpufilters.hpp b/modules/gpufilters/include/opencv2/gpufilters.hpp index 582c55d..5cc2ac4 100644 --- a/modules/gpufilters/include/opencv2/gpufilters.hpp +++ b/modules/gpufilters/include/opencv2/gpufilters.hpp @@ -48,10 +48,61 @@ #endif #include "opencv2/core/gpu.hpp" -#include "opencv2/core/base.hpp" + +#if defined __GNUC__ + #define __OPENCV_GPUFILTERS_DEPR_BEFORE__ + #define __OPENCV_GPUFILTERS_DEPR_AFTER__ __attribute__ ((deprecated)) +#elif (defined WIN32 || defined _WIN32) + #define __OPENCV_GPUFILTERS_DEPR_BEFORE__ __declspec(deprecated) + #define __OPENCV_GPUFILTERS_DEPR_AFTER__ +#else + #define __OPENCV_GPUFILTERS_DEPR_BEFORE__ + #define __OPENCV_GPUFILTERS_DEPR_AFTER__ +#endif namespace cv { namespace gpu { +class CV_EXPORTS Filter : public Algorithm +{ +public: + virtual void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null()) = 0; +}; + +//////////////////////////////////////////////////////////////////////////////////////////////////// +// Box Filter + +//! smooths the image using the normalized box filter +//! supports CV_8UC1, CV_8UC4 types +CV_EXPORTS Ptr createBoxFilter(int srcType, int dstType, Size ksize, Point anchor = Point(-1,-1), + int borderMode = BORDER_DEFAULT, Scalar borderVal = Scalar::all(0)); + +__OPENCV_GPUFILTERS_DEPR_BEFORE__ void boxFilter(InputArray src, OutputArray dst, int dstType, + Size ksize, Point anchor = Point(-1,-1), + Stream& stream = Stream::Null()) __OPENCV_GPUFILTERS_DEPR_AFTER__; + +inline void boxFilter(InputArray src, OutputArray dst, int dstType, Size ksize, Point anchor, Stream& stream) +{ + Ptr f = gpu::createBoxFilter(src.type(), dstType, ksize, anchor); + f->apply(src, dst, stream); +} + +__OPENCV_GPUFILTERS_DEPR_BEFORE__ void blur(InputArray src, OutputArray dst, Size ksize, + Point anchor = Point(-1,-1), + Stream& stream = Stream::Null()) __OPENCV_GPUFILTERS_DEPR_AFTER__; + +inline void blur(InputArray src, OutputArray dst, Size ksize, Point anchor, Stream& stream) +{ + Ptr f = gpu::createBoxFilter(src.type(), -1, ksize, anchor); + f->apply(src, dst, stream); +} + + + + + + + + /*! The Base Class for 1D or Row-wise Filters @@ -128,13 +179,7 @@ CV_EXPORTS Ptr getRowSumFilter_GPU(int srcType, int sumType, //! supports only CV_8UC1 sum type and CV_32FC1 dst type CV_EXPORTS Ptr getColumnSumFilter_GPU(int sumType, int dstType, int ksize, int anchor = -1); -//! returns 2D box filter -//! supports CV_8UC1 and CV_8UC4 source type, dst type must be the same as source type -CV_EXPORTS Ptr getBoxFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1, -1)); -//! returns box filter engine -CV_EXPORTS Ptr createBoxFilter_GPU(int srcType, int dstType, const Size& ksize, - const Point& anchor = Point(-1,-1)); //! returns 2D morphological filter //! only MORPH_ERODE and MORPH_DILATE are supported @@ -205,15 +250,7 @@ CV_EXPORTS Ptr getMaxFilter_GPU(int srcType, int dstType, const //! returns minimum filter CV_EXPORTS Ptr getMinFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1,-1)); -//! smooths the image using the normalized box filter -//! supports CV_8UC1, CV_8UC4 types -CV_EXPORTS void boxFilter(const GpuMat& src, GpuMat& dst, int ddepth, Size ksize, Point anchor = Point(-1,-1), Stream& stream = Stream::Null()); -//! a synonym for normalized box filter -static inline void blur(const GpuMat& src, GpuMat& dst, Size ksize, Point anchor = Point(-1,-1), Stream& stream = Stream::Null()) -{ - boxFilter(src, dst, -1, ksize, anchor, stream); -} //! 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); @@ -266,4 +303,7 @@ CV_EXPORTS void Laplacian(const GpuMat& src, GpuMat& dst, int ddepth, int ksize }} // namespace cv { namespace gpu { +#undef __OPENCV_GPUFILTERS_DEPR_BEFORE__ +#undef __OPENCV_GPUFILTERS_DEPR_AFTER__ + #endif /* __OPENCV_GPUFILTERS_HPP__ */ diff --git a/modules/gpufilters/perf/perf_filters.cpp b/modules/gpufilters/perf/perf_filters.cpp index 64cf4cc..35c4a94 100644 --- a/modules/gpufilters/perf/perf_filters.cpp +++ b/modules/gpufilters/perf/perf_filters.cpp @@ -70,7 +70,9 @@ PERF_TEST_P(Sz_Type_KernelSz, Blur, const cv::gpu::GpuMat d_src(src); cv::gpu::GpuMat dst; - TEST_CYCLE() cv::gpu::blur(d_src, dst, cv::Size(ksize, ksize)); + cv::Ptr blurFilter = cv::gpu::createBoxFilter(d_src.type(), -1, cv::Size(ksize, ksize)); + + TEST_CYCLE() blurFilter->apply(d_src, dst); GPU_SANITY_CHECK(dst, 1); } diff --git a/modules/gpufilters/src/filtering.cpp b/modules/gpufilters/src/filtering.cpp index d40293d..35df05e 100644 --- a/modules/gpufilters/src/filtering.cpp +++ b/modules/gpufilters/src/filtering.cpp @@ -47,13 +47,13 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) +Ptr cv::gpu::createBoxFilter(int, int, Size, Point, int, Scalar) { throw_no_cuda(); return Ptr(); } + Ptr cv::gpu::createFilter2D_GPU(const Ptr&, int, int) { throw_no_cuda(); return Ptr(0); } Ptr cv::gpu::createSeparableFilter_GPU(const Ptr&, const Ptr&, int, int, int) { throw_no_cuda(); return Ptr(0); } Ptr cv::gpu::createSeparableFilter_GPU(const Ptr&, const Ptr&, int, int, int, GpuMat&) { throw_no_cuda(); return Ptr(0); } 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::getBoxFilter_GPU(int, int, const Size&, Point) { throw_no_cuda(); return Ptr(0); } -Ptr cv::gpu::createBoxFilter_GPU(int, int, const Size&, const Point&) { 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); } @@ -70,7 +70,6 @@ Ptr cv::gpu::createGaussianFilter_GPU(int, Size, GpuMat&, doub 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::boxFilter(const GpuMat&, GpuMat&, int, Size, Point, Stream&) { throw_no_cuda(); } 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(); } @@ -92,20 +91,135 @@ void cv::gpu::Laplacian(const GpuMat&, GpuMat&, int, int, double, int, Stream&) namespace { - inline void normalizeAnchor(int& anchor, int ksize) + void normalizeAnchor(int& anchor, int ksize) { if (anchor < 0) anchor = ksize >> 1; - CV_Assert(0 <= anchor && anchor < ksize); + CV_Assert( 0 <= anchor && anchor < ksize ); } - inline void normalizeAnchor(Point& anchor, const Size& ksize) + void normalizeAnchor(Point& anchor, Size ksize) { normalizeAnchor(anchor.x, ksize.width); normalizeAnchor(anchor.y, ksize.height); } +} + +//////////////////////////////////////////////////////////////////////////////////////////////////// +// Box Filter + +namespace +{ + class NPPBoxFilter : public Filter + { + public: + NPPBoxFilter(int srcType, int dstType, Size ksize, Point anchor, int borderMode, Scalar borderVal); + + void apply(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); + + private: + typedef NppStatus (*nppFilterBox_t)(const Npp8u* pSrc, Npp32s nSrcStep, Npp8u* pDst, Npp32s nDstStep, + NppiSize oSizeROI, NppiSize oMaskSize, NppiPoint oAnchor); + + Size ksize_; + Point anchor_; + int type_; + nppFilterBox_t func_; + int borderMode_; + Scalar borderVal_; + GpuMat srcBorder_; + }; + + NPPBoxFilter::NPPBoxFilter(int srcType, int dstType, Size ksize, Point anchor, int borderMode, Scalar borderVal) : + ksize_(ksize), anchor_(anchor), type_(srcType), borderMode_(borderMode), borderVal_(borderVal) + { + static const nppFilterBox_t funcs[] = {0, nppiFilterBox_8u_C1R, 0, 0, nppiFilterBox_8u_C4R}; + + CV_Assert( srcType == CV_8UC1 || srcType == CV_8UC4 ); + CV_Assert( dstType == srcType ); + + normalizeAnchor(anchor_, ksize); + + func_ = funcs[CV_MAT_CN(srcType)]; + } + + void NPPBoxFilter::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::createBoxFilter(int srcType, int dstType, Size ksize, Point anchor, int borderMode, Scalar borderVal) +{ + if (dstType < 0) + dstType = srcType; + + return new NPPBoxFilter(srcType, dstType, ksize, anchor, borderMode, borderVal); +} + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +namespace +{ inline void normalizeROI(Rect& roi, const Size& ksize, const Point& anchor, const Size& src_size) { if (roi == Rect(0,0,-1,-1)) @@ -330,74 +444,6 @@ Ptr cv::gpu::getColumnSumFilter_GPU(int sumType, int dstTy } //////////////////////////////////////////////////////////////////////////////////////////////////// -// Box Filter - -namespace -{ - typedef NppStatus (*nppFilterBox_t)(const Npp8u * pSrc, Npp32s nSrcStep, Npp8u * pDst, Npp32s nDstStep, NppiSize oSizeROI, - NppiSize oMaskSize, NppiPoint oAnchor); - - struct NPPBoxFilter : public BaseFilter_GPU - { - NPPBoxFilter(const Size& ksize_, const Point& anchor_, nppFilterBox_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() ); - } - - nppFilterBox_t func; - }; -} - -Ptr cv::gpu::getBoxFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor) -{ - static const nppFilterBox_t nppFilterBox_callers[] = {0, nppiFilterBox_8u_C1R, 0, 0, nppiFilterBox_8u_C4R}; - - CV_Assert((srcType == CV_8UC1 || srcType == CV_8UC4) && dstType == srcType); - - normalizeAnchor(anchor, ksize); - - return Ptr(new NPPBoxFilter(ksize, anchor, nppFilterBox_callers[CV_MAT_CN(srcType)])); -} - -Ptr cv::gpu::createBoxFilter_GPU(int srcType, int dstType, const Size& ksize, const Point& anchor) -{ - Ptr boxFilter = getBoxFilter_GPU(srcType, dstType, ksize, anchor); - return createFilter2D_GPU(boxFilter, srcType, dstType); -} - -void cv::gpu::boxFilter(const GpuMat& src, GpuMat& dst, int ddepth, Size ksize, Point anchor, Stream& stream) -{ - int sdepth = src.depth(), cn = src.channels(); - if( ddepth < 0 ) - ddepth = sdepth; - - dst.create(src.size(), CV_MAKETYPE(ddepth, cn)); - - Ptr f = createBoxFilter_GPU(src.type(), dst.type(), ksize, anchor); - f->apply(src, dst, Rect(0,0,-1,-1), stream); -} - -//////////////////////////////////////////////////////////////////////////////////////////////////// // Morphology Filter namespace @@ -633,7 +679,6 @@ void cv::gpu::morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& ke erode(buf2, dst, kernel, buf1, anchor, iterations, stream); break; -#ifdef HAVE_OPENCV_GPUARITHM case MORPH_GRADIENT: erode(src, buf2, kernel, buf1, anchor, iterations, stream); dilate(src, dst, kernel, buf1, anchor, iterations, stream); @@ -651,7 +696,6 @@ void cv::gpu::morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& ke erode(dst, buf2, kernel, buf1, anchor, iterations, stream); gpu::subtract(buf2, src, dst, GpuMat(), -1, stream); break; -#endif default: CV_Error(cv::Error::StsBadArg, "unknown morphological operation"); diff --git a/modules/gpufilters/src/precomp.hpp b/modules/gpufilters/src/precomp.hpp index 3add0f2..c3d5e02 100644 --- a/modules/gpufilters/src/precomp.hpp +++ b/modules/gpufilters/src/precomp.hpp @@ -46,14 +46,9 @@ #include #include "opencv2/gpufilters.hpp" +#include "opencv2/gpuarithm.hpp" #include "opencv2/imgproc.hpp" #include "opencv2/core/private.gpu.hpp" -#include "opencv2/opencv_modules.hpp" - -#ifdef HAVE_OPENCV_GPUARITHM -# include "opencv2/gpuarithm.hpp" -#endif - #endif /* __OPENCV_PRECOMP_H__ */ diff --git a/modules/gpufilters/test/test_filters.cpp b/modules/gpufilters/test/test_filters.cpp index 5adcd87..a63d92b 100644 --- a/modules/gpufilters/test/test_filters.cpp +++ b/modules/gpufilters/test/test_filters.cpp @@ -70,13 +70,14 @@ namespace ///////////////////////////////////////////////////////////////////////////////////////////////// // Blur -PARAM_TEST_CASE(Blur, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Anchor, UseRoi) +PARAM_TEST_CASE(Blur, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Anchor, BorderType, UseRoi) { cv::gpu::DeviceInfo devInfo; cv::Size size; int type; cv::Size ksize; cv::Point anchor; + int borderType; bool useRoi; virtual void SetUp() @@ -86,7 +87,8 @@ PARAM_TEST_CASE(Blur, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Anchor, Use type = GET_PARAM(2); ksize = GET_PARAM(3); anchor = GET_PARAM(4); - useRoi = GET_PARAM(5); + borderType = GET_PARAM(5); + useRoi = GET_PARAM(6); cv::gpu::setDevice(devInfo.deviceID()); } @@ -96,13 +98,15 @@ GPU_TEST_P(Blur, Accuracy) { cv::Mat src = randomMat(size, type); + cv::Ptr blurFilter = cv::gpu::createBoxFilter(src.type(), -1, ksize, anchor, borderType); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); - cv::gpu::blur(loadMat(src, useRoi), dst, ksize, anchor); + blurFilter->apply(loadMat(src, useRoi), dst); cv::Mat dst_gold; - cv::blur(src, dst_gold, ksize, anchor); + cv::blur(src, dst_gold, ksize, anchor, borderType); - EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 1.0); + EXPECT_MAT_NEAR(dst_gold, dst, 1.0); } INSTANTIATE_TEST_CASE_P(GPU_Filters, Blur, testing::Combine( @@ -111,6 +115,7 @@ INSTANTIATE_TEST_CASE_P(GPU_Filters, Blur, testing::Combine( testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), testing::Values(KSize(cv::Size(3, 3)), KSize(cv::Size(5, 5)), KSize(cv::Size(7, 7))), testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))), + testing::Values(BorderType(cv::BORDER_REFLECT101), BorderType(cv::BORDER_REPLICATE), BorderType(cv::BORDER_CONSTANT), BorderType(cv::BORDER_REFLECT)), WHOLE_SUBMAT)); /////////////////////////////////////////////////////////////////////////////////////////////////