From 84de6ce0363b6e894428272d2099ae16b5c42337 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Thu, 18 Apr 2013 10:26:59 +0400 Subject: [PATCH] gpufilters module for image filtering --- modules/gpu/CMakeLists.txt | 2 +- modules/gpu/doc/gpu.rst | 1 - modules/gpu/include/opencv2/gpu.hpp | 215 +--------------- modules/gpu/src/cuda/imgproc.cu | 106 -------- modules/gpufilters/CMakeLists.txt | 9 + .../doc/filtering.rst} | 0 modules/gpufilters/doc/gpufilters.rst | 8 + modules/gpufilters/include/opencv2/gpufilters.hpp | 269 +++++++++++++++++++++ modules/{gpu => gpufilters}/perf/perf_filters.cpp | 18 +- modules/gpufilters/perf/perf_main.cpp | 47 ++++ modules/gpufilters/perf/perf_precomp.cpp | 43 ++++ modules/gpufilters/perf/perf_precomp.hpp | 64 +++++ .../src/cuda/column_filter.16sc1.cu} | 2 +- .../src/cuda/column_filter.16sc3.cu} | 2 +- .../src/cuda/column_filter.16sc4.cu} | 2 +- .../src/cuda/column_filter.16uc1.cu} | 2 +- .../src/cuda/column_filter.16uc3.cu} | 2 +- .../src/cuda/column_filter.16uc4.cu} | 2 +- .../src/cuda/column_filter.32fc1.cu} | 2 +- .../src/cuda/column_filter.32fc3.cu} | 2 +- .../src/cuda/column_filter.32fc4.cu} | 2 +- .../src/cuda/column_filter.32sc1.cu} | 2 +- .../src/cuda/column_filter.32sc3.cu} | 2 +- .../src/cuda/column_filter.32sc4.cu} | 2 +- .../src/cuda/column_filter.8uc1.cu} | 2 +- .../src/cuda/column_filter.8uc3.cu} | 2 +- .../src/cuda/column_filter.8uc4.cu} | 2 +- .../src/cuda/column_filter.hpp} | 0 modules/gpufilters/src/cuda/filter2d.cu | 158 ++++++++++++ .../src/cuda/row_filter.16sc1.cu} | 2 +- .../src/cuda/row_filter.16sc3.cu} | 2 +- .../src/cuda/row_filter.16sc4.cu} | 2 +- .../src/cuda/row_filter.16uc1.cu} | 2 +- .../src/cuda/row_filter.16uc3.cu} | 2 +- .../src/cuda/row_filter.16uc4.cu} | 2 +- .../src/cuda/row_filter.32fc1.cu} | 2 +- .../src/cuda/row_filter.32fc3.cu} | 2 +- .../src/cuda/row_filter.32fc4.cu} | 2 +- .../src/cuda/row_filter.32sc1.cu} | 2 +- .../src/cuda/row_filter.32sc3.cu} | 2 +- .../src/cuda/row_filter.32sc4.cu} | 2 +- .../src/cuda/row_filter.8uc1.cu} | 2 +- .../src/cuda/row_filter.8uc3.cu} | 2 +- .../src/cuda/row_filter.8uc4.cu} | 2 +- .../src/cuda/row_filter.hpp} | 0 modules/{gpu => gpufilters}/src/filtering.cpp | 24 +- modules/gpufilters/src/precomp.cpp | 43 ++++ modules/gpufilters/src/precomp.hpp | 59 +++++ modules/{gpu => gpufilters}/test/test_filters.cpp | 18 +- modules/gpufilters/test/test_main.cpp | 120 +++++++++ modules/gpufilters/test/test_precomp.cpp | 43 ++++ modules/gpufilters/test/test_precomp.hpp | 60 +++++ modules/stitching/CMakeLists.txt | 2 +- modules/superres/CMakeLists.txt | 2 +- samples/cpp/CMakeLists.txt | 1 + samples/gpu/CMakeLists.txt | 2 +- 56 files changed, 995 insertions(+), 379 deletions(-) create mode 100644 modules/gpufilters/CMakeLists.txt rename modules/{gpu/doc/image_filtering.rst => gpufilters/doc/filtering.rst} (100%) create mode 100644 modules/gpufilters/doc/gpufilters.rst create mode 100644 modules/gpufilters/include/opencv2/gpufilters.hpp rename modules/{gpu => gpufilters}/perf/perf_filters.cpp (89%) create mode 100644 modules/gpufilters/perf/perf_main.cpp create mode 100644 modules/gpufilters/perf/perf_precomp.cpp create mode 100644 modules/gpufilters/perf/perf_precomp.hpp rename modules/{gpu/src/cuda/column_filter.8.cu => gpufilters/src/cuda/column_filter.16sc1.cu} (98%) rename modules/{gpu/src/cuda/column_filter.3.cu => gpufilters/src/cuda/column_filter.16sc3.cu} (98%) rename modules/{gpu/src/cuda/column_filter.9.cu => gpufilters/src/cuda/column_filter.16sc4.cu} (98%) rename modules/{gpu/src/cuda/column_filter.10.cu => gpufilters/src/cuda/column_filter.16uc1.cu} (98%) rename modules/{gpu/src/cuda/column_filter.11.cu => gpufilters/src/cuda/column_filter.16uc3.cu} (98%) rename modules/{gpu/src/cuda/column_filter.12.cu => gpufilters/src/cuda/column_filter.16uc4.cu} (98%) rename modules/{gpu/src/cuda/column_filter.5.cu => gpufilters/src/cuda/column_filter.32fc1.cu} (98%) rename modules/{gpu/src/cuda/column_filter.6.cu => gpufilters/src/cuda/column_filter.32fc3.cu} (98%) rename modules/{gpu/src/cuda/column_filter.7.cu => gpufilters/src/cuda/column_filter.32fc4.cu} (98%) rename modules/{gpu/src/cuda/column_filter.4.cu => gpufilters/src/cuda/column_filter.32sc1.cu} (98%) rename modules/{gpu/src/cuda/column_filter.13.cu => gpufilters/src/cuda/column_filter.32sc3.cu} (98%) rename modules/{gpu/src/cuda/column_filter.14.cu => gpufilters/src/cuda/column_filter.32sc4.cu} (98%) rename modules/{gpu/src/cuda/column_filter.0.cu => gpufilters/src/cuda/column_filter.8uc1.cu} (98%) rename modules/{gpu/src/cuda/column_filter.1.cu => gpufilters/src/cuda/column_filter.8uc3.cu} (98%) rename modules/{gpu/src/cuda/column_filter.2.cu => gpufilters/src/cuda/column_filter.8uc4.cu} (98%) rename modules/{gpu/src/cuda/column_filter.h => gpufilters/src/cuda/column_filter.hpp} (100%) create mode 100644 modules/gpufilters/src/cuda/filter2d.cu rename modules/{gpu/src/cuda/row_filter.8.cu => gpufilters/src/cuda/row_filter.16sc1.cu} (98%) rename modules/{gpu/src/cuda/row_filter.3.cu => gpufilters/src/cuda/row_filter.16sc3.cu} (98%) rename modules/{gpu/src/cuda/row_filter.9.cu => gpufilters/src/cuda/row_filter.16sc4.cu} (98%) rename modules/{gpu/src/cuda/row_filter.10.cu => gpufilters/src/cuda/row_filter.16uc1.cu} (98%) rename modules/{gpu/src/cuda/row_filter.11.cu => gpufilters/src/cuda/row_filter.16uc3.cu} (98%) rename modules/{gpu/src/cuda/row_filter.12.cu => gpufilters/src/cuda/row_filter.16uc4.cu} (98%) rename modules/{gpu/src/cuda/row_filter.5.cu => gpufilters/src/cuda/row_filter.32fc1.cu} (98%) rename modules/{gpu/src/cuda/row_filter.6.cu => gpufilters/src/cuda/row_filter.32fc3.cu} (98%) rename modules/{gpu/src/cuda/row_filter.7.cu => gpufilters/src/cuda/row_filter.32fc4.cu} (98%) rename modules/{gpu/src/cuda/row_filter.4.cu => gpufilters/src/cuda/row_filter.32sc1.cu} (98%) rename modules/{gpu/src/cuda/row_filter.13.cu => gpufilters/src/cuda/row_filter.32sc3.cu} (98%) rename modules/{gpu/src/cuda/row_filter.14.cu => gpufilters/src/cuda/row_filter.32sc4.cu} (98%) rename modules/{gpu/src/cuda/row_filter.0.cu => gpufilters/src/cuda/row_filter.8uc1.cu} (98%) rename modules/{gpu/src/cuda/row_filter.1.cu => gpufilters/src/cuda/row_filter.8uc3.cu} (98%) rename modules/{gpu/src/cuda/row_filter.2.cu => gpufilters/src/cuda/row_filter.8uc4.cu} (98%) rename modules/{gpu/src/cuda/row_filter.h => gpufilters/src/cuda/row_filter.hpp} (100%) rename modules/{gpu => gpufilters}/src/filtering.cpp (99%) create mode 100644 modules/gpufilters/src/precomp.cpp create mode 100644 modules/gpufilters/src/precomp.hpp rename modules/{gpu => gpufilters}/test/test_filters.cpp (96%) create mode 100644 modules/gpufilters/test/test_main.cpp create mode 100644 modules/gpufilters/test/test_precomp.cpp create mode 100644 modules/gpufilters/test/test_precomp.hpp diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt index 2f884b3..1ed9806 100644 --- a/modules/gpu/CMakeLists.txt +++ b/modules/gpu/CMakeLists.txt @@ -3,7 +3,7 @@ if(ANDROID OR IOS) endif() set(the_description "GPU-accelerated Computer Vision") -ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video opencv_photo opencv_legacy opencv_gpuarithm) +ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video opencv_photo opencv_legacy opencv_gpuarithm opencv_gpufilters) ocv_module_include_directories("${CMAKE_CURRENT_SOURCE_DIR}/src/cuda") diff --git a/modules/gpu/doc/gpu.rst b/modules/gpu/doc/gpu.rst index f17ed70..de52cea 100644 --- a/modules/gpu/doc/gpu.rst +++ b/modules/gpu/doc/gpu.rst @@ -11,6 +11,5 @@ gpu. GPU-accelerated Computer Vision image_processing object_detection feature_detection_and_description - image_filtering camera_calibration_and_3d_reconstruction video diff --git a/modules/gpu/include/opencv2/gpu.hpp b/modules/gpu/include/opencv2/gpu.hpp index cfad817..8f837da 100644 --- a/modules/gpu/include/opencv2/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu.hpp @@ -51,225 +51,12 @@ #include "opencv2/core/gpumat.hpp" #include "opencv2/gpuarithm.hpp" +#include "opencv2/gpufilters.hpp" #include "opencv2/imgproc.hpp" #include "opencv2/objdetect.hpp" #include "opencv2/features2d.hpp" namespace cv { namespace gpu { -//////////////////////////////// Filter Engine //////////////////////////////// - -/*! -The Base Class for 1D or Row-wise Filters - -This is the base class for linear or non-linear filters that process 1D data. -In particular, such filters are used for the "horizontal" filtering parts in separable filters. -*/ -class CV_EXPORTS BaseRowFilter_GPU -{ -public: - BaseRowFilter_GPU(int ksize_, int anchor_) : ksize(ksize_), anchor(anchor_) {} - virtual ~BaseRowFilter_GPU() {} - virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) = 0; - int ksize, anchor; -}; - -/*! -The Base Class for Column-wise Filters - -This is the base class for linear or non-linear filters that process columns of 2D arrays. -Such filters are used for the "vertical" filtering parts in separable filters. -*/ -class CV_EXPORTS BaseColumnFilter_GPU -{ -public: - BaseColumnFilter_GPU(int ksize_, int anchor_) : ksize(ksize_), anchor(anchor_) {} - virtual ~BaseColumnFilter_GPU() {} - virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) = 0; - int ksize, anchor; -}; - -/*! -The Base Class for Non-Separable 2D Filters. - -This is the base class for linear or non-linear 2D filters. -*/ -class CV_EXPORTS BaseFilter_GPU -{ -public: - BaseFilter_GPU(const Size& ksize_, const Point& anchor_) : ksize(ksize_), anchor(anchor_) {} - virtual ~BaseFilter_GPU() {} - virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) = 0; - Size ksize; - Point anchor; -}; - -/*! -The Base Class for Filter Engine. - -The class can be used to apply an arbitrary filtering operation to an image. -It contains all the necessary intermediate buffers. -*/ -class CV_EXPORTS FilterEngine_GPU -{ -public: - virtual ~FilterEngine_GPU() {} - - virtual void apply(const GpuMat& src, GpuMat& dst, Rect roi = Rect(0,0,-1,-1), Stream& stream = Stream::Null()) = 0; -}; - -//! returns the non-separable filter engine with the specified filter -CV_EXPORTS Ptr createFilter2D_GPU(const Ptr& filter2D, int srcType, int dstType); - -//! returns the separable filter engine with the specified filters -CV_EXPORTS Ptr createSeparableFilter_GPU(const Ptr& rowFilter, - const Ptr& columnFilter, int srcType, int bufType, int dstType); -CV_EXPORTS Ptr createSeparableFilter_GPU(const Ptr& rowFilter, - const Ptr& columnFilter, int srcType, int bufType, int dstType, GpuMat& buf); - -//! returns horizontal 1D box filter -//! supports only CV_8UC1 source type and CV_32FC1 sum type -CV_EXPORTS Ptr getRowSumFilter_GPU(int srcType, int sumType, int ksize, int anchor = -1); - -//! returns vertical 1D box filter -//! 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 -//! 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); - -//! returns 2D filter with the specified kernel -//! supports CV_8U, CV_16U and CV_32F one and four channel image -CV_EXPORTS Ptr getLinearFilter_GPU(int srcType, int dstType, const Mat& kernel, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT); - -//! returns the non-separable linear filter engine -CV_EXPORTS Ptr createLinearFilter_GPU(int srcType, int dstType, const Mat& kernel, - Point anchor = Point(-1,-1), int borderType = BORDER_DEFAULT); - -//! returns the primitive row filter with the specified kernel. -//! supports only CV_8UC1, CV_8UC4, CV_16SC1, CV_16SC2, CV_32SC1, CV_32FC1 source type. -//! there are two version of algorithm: NPP and OpenCV. -//! NPP calls when srcType == CV_8UC1 or srcType == CV_8UC4 and bufType == srcType, -//! otherwise calls OpenCV version. -//! NPP supports only BORDER_CONSTANT border type. -//! OpenCV version supports only CV_32F as buffer depth and -//! BORDER_REFLECT101, BORDER_REPLICATE and BORDER_CONSTANT border types. -CV_EXPORTS Ptr getLinearRowFilter_GPU(int srcType, int bufType, const Mat& rowKernel, - int anchor = -1, int borderType = BORDER_DEFAULT); - -//! returns the primitive column filter with the specified kernel. -//! supports only CV_8UC1, CV_8UC4, CV_16SC1, CV_16SC2, CV_32SC1, CV_32FC1 dst type. -//! there are two version of algorithm: NPP and OpenCV. -//! NPP calls when dstType == CV_8UC1 or dstType == CV_8UC4 and bufType == dstType, -//! otherwise calls OpenCV version. -//! NPP supports only BORDER_CONSTANT border type. -//! OpenCV version supports only CV_32F as buffer depth and -//! BORDER_REFLECT101, BORDER_REPLICATE and BORDER_CONSTANT border types. -CV_EXPORTS Ptr getLinearColumnFilter_GPU(int bufType, int dstType, const Mat& columnKernel, - int anchor = -1, int borderType = BORDER_DEFAULT); - -//! returns the separable linear filter engine -CV_EXPORTS Ptr createSeparableLinearFilter_GPU(int srcType, int dstType, const Mat& rowKernel, - const Mat& columnKernel, const Point& anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, - int columnBorderType = -1); -CV_EXPORTS Ptr createSeparableLinearFilter_GPU(int srcType, int dstType, const Mat& rowKernel, - const Mat& columnKernel, GpuMat& buf, const Point& anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, - int columnBorderType = -1); - -//! returns filter engine for the generalized Sobel operator -CV_EXPORTS Ptr createDerivFilter_GPU(int srcType, int dstType, int dx, int dy, int ksize, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); -CV_EXPORTS Ptr createDerivFilter_GPU(int srcType, int dstType, int dx, int dy, int ksize, GpuMat& buf, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); - -//! returns the Gaussian filter engine -CV_EXPORTS Ptr createGaussianFilter_GPU(int type, Size ksize, double sigma1, double sigma2 = 0, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); -CV_EXPORTS Ptr createGaussianFilter_GPU(int type, Size ksize, GpuMat& buf, double sigma1, double sigma2 = 0, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); - -//! returns maximum filter -CV_EXPORTS Ptr getMaxFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1,-1)); - -//! returns minimum filter -CV_EXPORTS Ptr getMinFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1,-1)); - -//! 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); -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()); - -//! applies non-separable 2D linear filter to the image -CV_EXPORTS void filter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& kernel, Point anchor=Point(-1,-1), int borderType = BORDER_DEFAULT, Stream& stream = Stream::Null()); - -//! applies separable 2D linear filter to the image -CV_EXPORTS void sepFilter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& kernelX, const Mat& kernelY, - Point anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); -CV_EXPORTS void sepFilter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& kernelX, const Mat& kernelY, GpuMat& buf, - Point anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, - Stream& stream = Stream::Null()); - -//! applies generalized Sobel operator to the image -CV_EXPORTS void Sobel(const GpuMat& src, GpuMat& dst, int ddepth, int dx, int dy, int ksize = 3, double scale = 1, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); -CV_EXPORTS void Sobel(const GpuMat& src, GpuMat& dst, int ddepth, int dx, int dy, GpuMat& buf, int ksize = 3, double scale = 1, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, Stream& stream = Stream::Null()); - -//! applies the vertical or horizontal Scharr operator to the image -CV_EXPORTS void Scharr(const GpuMat& src, GpuMat& dst, int ddepth, int dx, int dy, double scale = 1, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); -CV_EXPORTS void Scharr(const GpuMat& src, GpuMat& dst, int ddepth, int dx, int dy, GpuMat& buf, double scale = 1, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, Stream& stream = Stream::Null()); - -//! smooths the image using Gaussian filter. -CV_EXPORTS void GaussianBlur(const GpuMat& src, GpuMat& dst, Size ksize, double sigma1, double sigma2 = 0, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); -CV_EXPORTS void GaussianBlur(const GpuMat& src, GpuMat& dst, Size ksize, GpuMat& buf, double sigma1, double sigma2 = 0, - int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, Stream& stream = Stream::Null()); - -//! applies Laplacian operator to the image -//! supports only ksize = 1 and ksize = 3 -CV_EXPORTS void Laplacian(const GpuMat& src, GpuMat& dst, int ddepth, int ksize = 1, double scale = 1, int borderType = BORDER_DEFAULT, Stream& stream = Stream::Null()); - ////////////////////////////// Image processing ////////////////////////////// diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index fc27ec1..71f5e87 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -895,112 +895,6 @@ namespace cv { namespace gpu { namespace cudev if (stream == 0) cudaSafeCall(cudaDeviceSynchronize()); } - - ////////////////////////////////////////////////////////////////////////// - // filter2D - - #define FILTER2D_MAX_KERNEL_SIZE 16 - - __constant__ float c_filter2DKernel[FILTER2D_MAX_KERNEL_SIZE * FILTER2D_MAX_KERNEL_SIZE]; - - template - __global__ void filter2D(const SrcT src, PtrStepSz dst, const int kWidth, const int kHeight, const int anchorX, const int anchorY) - { - typedef typename TypeVec::cn>::vec_type sum_t; - - const int x = blockIdx.x * blockDim.x + threadIdx.x; - const int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= dst.cols || y >= dst.rows) - return; - - sum_t res = VecTraits::all(0); - int kInd = 0; - - for (int i = 0; i < kHeight; ++i) - { - for (int j = 0; j < kWidth; ++j) - res = res + src(y - anchorY + i, x - anchorX + j) * c_filter2DKernel[kInd++]; - } - - dst(y, x) = saturate_cast(res); - } - - template class Brd> struct Filter2DCaller; - - #define IMPLEMENT_FILTER2D_TEX_READER(type) \ - texture< type , cudaTextureType2D, cudaReadModeElementType> tex_filter2D_ ## type (0, cudaFilterModePoint, cudaAddressModeClamp); \ - struct tex_filter2D_ ## type ## _reader \ - { \ - typedef type elem_type; \ - typedef int index_type; \ - const int xoff; \ - const int yoff; \ - tex_filter2D_ ## type ## _reader (int xoff_, int yoff_) : xoff(xoff_), yoff(yoff_) {} \ - __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const \ - { \ - return tex2D(tex_filter2D_ ## type , x + xoff, y + yoff); \ - } \ - }; \ - template class Brd> struct Filter2DCaller< type , D, Brd> \ - { \ - static void call(const PtrStepSz< type > srcWhole, int xoff, int yoff, PtrStepSz dst, \ - int kWidth, int kHeight, int anchorX, int anchorY, const float* borderValue, cudaStream_t stream) \ - { \ - typedef typename TypeVec::cn>::vec_type work_type; \ - dim3 block(16, 16); \ - dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \ - bindTexture(&tex_filter2D_ ## type , srcWhole); \ - tex_filter2D_ ## type ##_reader texSrc(xoff, yoff); \ - Brd brd(dst.rows, dst.cols, VecTraits::make(borderValue)); \ - BorderReader< tex_filter2D_ ## type ##_reader, Brd > brdSrc(texSrc, brd); \ - filter2D<<>>(brdSrc, dst, kWidth, kHeight, anchorX, anchorY); \ - cudaSafeCall( cudaGetLastError() ); \ - if (stream == 0) \ - cudaSafeCall( cudaDeviceSynchronize() ); \ - } \ - }; - - IMPLEMENT_FILTER2D_TEX_READER(uchar); - IMPLEMENT_FILTER2D_TEX_READER(uchar4); - - IMPLEMENT_FILTER2D_TEX_READER(ushort); - IMPLEMENT_FILTER2D_TEX_READER(ushort4); - - IMPLEMENT_FILTER2D_TEX_READER(float); - IMPLEMENT_FILTER2D_TEX_READER(float4); - - #undef IMPLEMENT_FILTER2D_TEX_READER - - template - void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, - int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, - int borderMode, const float* borderValue, cudaStream_t stream) - { - typedef void (*func_t)(const PtrStepSz srcWhole, int xoff, int yoff, PtrStepSz dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* borderValue, cudaStream_t stream); - static const func_t funcs[] = - { - Filter2DCaller::call, - Filter2DCaller::call, - Filter2DCaller::call, - Filter2DCaller::call, - Filter2DCaller::call - }; - - if (stream == 0) - cudaSafeCall( cudaMemcpyToSymbol(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); - else - cudaSafeCall( cudaMemcpyToSymbolAsync(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) ); - - funcs[borderMode](static_cast< PtrStepSz >(srcWhole), ofsX, ofsY, static_cast< PtrStepSz >(dst), kWidth, kHeight, anchorX, anchorY, borderValue, stream); - } - - template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); - template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); - template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); - template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); - template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); - template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); } // namespace imgproc }}} // namespace cv { namespace gpu { namespace cudev { diff --git a/modules/gpufilters/CMakeLists.txt b/modules/gpufilters/CMakeLists.txt new file mode 100644 index 0000000..18f6d7f --- /dev/null +++ b/modules/gpufilters/CMakeLists.txt @@ -0,0 +1,9 @@ +if(ANDROID OR IOS) + ocv_module_disable(gpufilters) +endif() + +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) diff --git a/modules/gpu/doc/image_filtering.rst b/modules/gpufilters/doc/filtering.rst similarity index 100% rename from modules/gpu/doc/image_filtering.rst rename to modules/gpufilters/doc/filtering.rst diff --git a/modules/gpufilters/doc/gpufilters.rst b/modules/gpufilters/doc/gpufilters.rst new file mode 100644 index 0000000..778554b --- /dev/null +++ b/modules/gpufilters/doc/gpufilters.rst @@ -0,0 +1,8 @@ +******************************************* +gpufilters. GPU-accelerated Image Filtering +******************************************* + +.. toctree:: + :maxdepth: 1 + + filtering diff --git a/modules/gpufilters/include/opencv2/gpufilters.hpp b/modules/gpufilters/include/opencv2/gpufilters.hpp new file mode 100644 index 0000000..853755d --- /dev/null +++ b/modules/gpufilters/include/opencv2/gpufilters.hpp @@ -0,0 +1,269 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __OPENCV_GPUFILTERS_HPP__ +#define __OPENCV_GPUFILTERS_HPP__ + +#ifndef __cplusplus +# error gpufilters.hpp header must be compiled as C++ +#endif + +#include "opencv2/core/gpumat.hpp" +#include "opencv2/core/base.hpp" + +namespace cv { namespace gpu { + +/*! +The Base Class for 1D or Row-wise Filters + +This is the base class for linear or non-linear filters that process 1D data. +In particular, such filters are used for the "horizontal" filtering parts in separable filters. +*/ +class CV_EXPORTS BaseRowFilter_GPU +{ +public: + BaseRowFilter_GPU(int ksize_, int anchor_) : ksize(ksize_), anchor(anchor_) {} + virtual ~BaseRowFilter_GPU() {} + virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) = 0; + int ksize, anchor; +}; + +/*! +The Base Class for Column-wise Filters + +This is the base class for linear or non-linear filters that process columns of 2D arrays. +Such filters are used for the "vertical" filtering parts in separable filters. +*/ +class CV_EXPORTS BaseColumnFilter_GPU +{ +public: + BaseColumnFilter_GPU(int ksize_, int anchor_) : ksize(ksize_), anchor(anchor_) {} + virtual ~BaseColumnFilter_GPU() {} + virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) = 0; + int ksize, anchor; +}; + +/*! +The Base Class for Non-Separable 2D Filters. + +This is the base class for linear or non-linear 2D filters. +*/ +class CV_EXPORTS BaseFilter_GPU +{ +public: + BaseFilter_GPU(const Size& ksize_, const Point& anchor_) : ksize(ksize_), anchor(anchor_) {} + virtual ~BaseFilter_GPU() {} + virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) = 0; + Size ksize; + Point anchor; +}; + +/*! +The Base Class for Filter Engine. + +The class can be used to apply an arbitrary filtering operation to an image. +It contains all the necessary intermediate buffers. +*/ +class CV_EXPORTS FilterEngine_GPU +{ +public: + virtual ~FilterEngine_GPU() {} + + virtual void apply(const GpuMat& src, GpuMat& dst, Rect roi = Rect(0,0,-1,-1), Stream& stream = Stream::Null()) = 0; +}; + +//! returns the non-separable filter engine with the specified filter +CV_EXPORTS Ptr createFilter2D_GPU(const Ptr& filter2D, int srcType, int dstType); + +//! returns the separable filter engine with the specified filters +CV_EXPORTS Ptr createSeparableFilter_GPU(const Ptr& rowFilter, + const Ptr& columnFilter, int srcType, int bufType, int dstType); +CV_EXPORTS Ptr createSeparableFilter_GPU(const Ptr& rowFilter, + const Ptr& columnFilter, int srcType, int bufType, int dstType, GpuMat& buf); + +//! returns horizontal 1D box filter +//! supports only CV_8UC1 source type and CV_32FC1 sum type +CV_EXPORTS Ptr getRowSumFilter_GPU(int srcType, int sumType, int ksize, int anchor = -1); + +//! returns vertical 1D box filter +//! 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 +//! 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); + +//! returns 2D filter with the specified kernel +//! supports CV_8U, CV_16U and CV_32F one and four channel image +CV_EXPORTS Ptr getLinearFilter_GPU(int srcType, int dstType, const Mat& kernel, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT); + +//! returns the non-separable linear filter engine +CV_EXPORTS Ptr createLinearFilter_GPU(int srcType, int dstType, const Mat& kernel, + Point anchor = Point(-1,-1), int borderType = BORDER_DEFAULT); + +//! returns the primitive row filter with the specified kernel. +//! supports only CV_8UC1, CV_8UC4, CV_16SC1, CV_16SC2, CV_32SC1, CV_32FC1 source type. +//! there are two version of algorithm: NPP and OpenCV. +//! NPP calls when srcType == CV_8UC1 or srcType == CV_8UC4 and bufType == srcType, +//! otherwise calls OpenCV version. +//! NPP supports only BORDER_CONSTANT border type. +//! OpenCV version supports only CV_32F as buffer depth and +//! BORDER_REFLECT101, BORDER_REPLICATE and BORDER_CONSTANT border types. +CV_EXPORTS Ptr getLinearRowFilter_GPU(int srcType, int bufType, const Mat& rowKernel, + int anchor = -1, int borderType = BORDER_DEFAULT); + +//! returns the primitive column filter with the specified kernel. +//! supports only CV_8UC1, CV_8UC4, CV_16SC1, CV_16SC2, CV_32SC1, CV_32FC1 dst type. +//! there are two version of algorithm: NPP and OpenCV. +//! NPP calls when dstType == CV_8UC1 or dstType == CV_8UC4 and bufType == dstType, +//! otherwise calls OpenCV version. +//! NPP supports only BORDER_CONSTANT border type. +//! OpenCV version supports only CV_32F as buffer depth and +//! BORDER_REFLECT101, BORDER_REPLICATE and BORDER_CONSTANT border types. +CV_EXPORTS Ptr getLinearColumnFilter_GPU(int bufType, int dstType, const Mat& columnKernel, + int anchor = -1, int borderType = BORDER_DEFAULT); + +//! returns the separable linear filter engine +CV_EXPORTS Ptr createSeparableLinearFilter_GPU(int srcType, int dstType, const Mat& rowKernel, + const Mat& columnKernel, const Point& anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, + int columnBorderType = -1); +CV_EXPORTS Ptr createSeparableLinearFilter_GPU(int srcType, int dstType, const Mat& rowKernel, + const Mat& columnKernel, GpuMat& buf, const Point& anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, + int columnBorderType = -1); + +//! returns filter engine for the generalized Sobel operator +CV_EXPORTS Ptr createDerivFilter_GPU(int srcType, int dstType, int dx, int dy, int ksize, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); +CV_EXPORTS Ptr createDerivFilter_GPU(int srcType, int dstType, int dx, int dy, int ksize, GpuMat& buf, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); + +//! returns the Gaussian filter engine +CV_EXPORTS Ptr createGaussianFilter_GPU(int type, Size ksize, double sigma1, double sigma2 = 0, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); +CV_EXPORTS Ptr createGaussianFilter_GPU(int type, Size ksize, GpuMat& buf, double sigma1, double sigma2 = 0, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); + +//! returns maximum filter +CV_EXPORTS Ptr getMaxFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1,-1)); + +//! returns minimum filter +CV_EXPORTS Ptr getMinFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1,-1)); + +//! 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); +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()); + +//! applies non-separable 2D linear filter to the image +CV_EXPORTS void filter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& kernel, Point anchor=Point(-1,-1), int borderType = BORDER_DEFAULT, Stream& stream = Stream::Null()); + +//! applies separable 2D linear filter to the image +CV_EXPORTS void sepFilter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& kernelX, const Mat& kernelY, + Point anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); +CV_EXPORTS void sepFilter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& kernelX, const Mat& kernelY, GpuMat& buf, + Point anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, + Stream& stream = Stream::Null()); + +//! applies generalized Sobel operator to the image +CV_EXPORTS void Sobel(const GpuMat& src, GpuMat& dst, int ddepth, int dx, int dy, int ksize = 3, double scale = 1, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); +CV_EXPORTS void Sobel(const GpuMat& src, GpuMat& dst, int ddepth, int dx, int dy, GpuMat& buf, int ksize = 3, double scale = 1, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, Stream& stream = Stream::Null()); + +//! applies the vertical or horizontal Scharr operator to the image +CV_EXPORTS void Scharr(const GpuMat& src, GpuMat& dst, int ddepth, int dx, int dy, double scale = 1, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); +CV_EXPORTS void Scharr(const GpuMat& src, GpuMat& dst, int ddepth, int dx, int dy, GpuMat& buf, double scale = 1, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, Stream& stream = Stream::Null()); + +//! smooths the image using Gaussian filter. +CV_EXPORTS void GaussianBlur(const GpuMat& src, GpuMat& dst, Size ksize, double sigma1, double sigma2 = 0, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); +CV_EXPORTS void GaussianBlur(const GpuMat& src, GpuMat& dst, Size ksize, GpuMat& buf, double sigma1, double sigma2 = 0, + int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, Stream& stream = Stream::Null()); + +//! applies Laplacian operator to the image +//! supports only ksize = 1 and ksize = 3 +CV_EXPORTS void Laplacian(const GpuMat& src, GpuMat& dst, int ddepth, int ksize = 1, double scale = 1, int borderType = BORDER_DEFAULT, Stream& stream = Stream::Null()); + +}} // namespace cv { namespace gpu { + +#endif /* __OPENCV_GPUFILTERS_HPP__ */ diff --git a/modules/gpu/perf/perf_filters.cpp b/modules/gpufilters/perf/perf_filters.cpp similarity index 89% rename from modules/gpu/perf/perf_filters.cpp rename to modules/gpufilters/perf/perf_filters.cpp index 40d88aa..0dc506b 100644 --- a/modules/gpu/perf/perf_filters.cpp +++ b/modules/gpufilters/perf/perf_filters.cpp @@ -51,7 +51,7 @@ using namespace perf; DEF_PARAM_TEST(Sz_Type_KernelSz, cv::Size, MatType, int); -PERF_TEST_P(Sz_Type_KernelSz, Filters_Blur, +PERF_TEST_P(Sz_Type_KernelSz, Blur, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4), Values(3, 5, 7))) @@ -87,7 +87,7 @@ PERF_TEST_P(Sz_Type_KernelSz, Filters_Blur, ////////////////////////////////////////////////////////////////////// // Sobel -PERF_TEST_P(Sz_Type_KernelSz, Filters_Sobel, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1), Values(3, 5, 7, 9, 11, 13, 15))) +PERF_TEST_P(Sz_Type_KernelSz, Sobel, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1), Values(3, 5, 7, 9, 11, 13, 15))) { declare.time(20.0); @@ -121,7 +121,7 @@ PERF_TEST_P(Sz_Type_KernelSz, Filters_Sobel, Combine(GPU_TYPICAL_MAT_SIZES, Valu ////////////////////////////////////////////////////////////////////// // Scharr -PERF_TEST_P(Sz_Type, Filters_Scharr, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1))) +PERF_TEST_P(Sz_Type, Scharr, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1))) { declare.time(20.0); @@ -154,7 +154,7 @@ PERF_TEST_P(Sz_Type, Filters_Scharr, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8U ////////////////////////////////////////////////////////////////////// // GaussianBlur -PERF_TEST_P(Sz_Type_KernelSz, Filters_GaussianBlur, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1), Values(3, 5, 7, 9, 11, 13, 15))) +PERF_TEST_P(Sz_Type_KernelSz, GaussianBlur, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1), Values(3, 5, 7, 9, 11, 13, 15))) { declare.time(20.0); @@ -188,7 +188,7 @@ PERF_TEST_P(Sz_Type_KernelSz, Filters_GaussianBlur, Combine(GPU_TYPICAL_MAT_SIZE ////////////////////////////////////////////////////////////////////// // Laplacian -PERF_TEST_P(Sz_Type_KernelSz, Filters_Laplacian, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(1, 3))) +PERF_TEST_P(Sz_Type_KernelSz, Laplacian, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(1, 3))) { declare.time(20.0); @@ -221,7 +221,7 @@ PERF_TEST_P(Sz_Type_KernelSz, Filters_Laplacian, Combine(GPU_TYPICAL_MAT_SIZES, ////////////////////////////////////////////////////////////////////// // Erode -PERF_TEST_P(Sz_Type, Filters_Erode, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4))) +PERF_TEST_P(Sz_Type, Erode, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4))) { declare.time(20.0); @@ -256,7 +256,7 @@ PERF_TEST_P(Sz_Type, Filters_Erode, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC ////////////////////////////////////////////////////////////////////// // Dilate -PERF_TEST_P(Sz_Type, Filters_Dilate, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4))) +PERF_TEST_P(Sz_Type, Dilate, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4))) { declare.time(20.0); @@ -295,7 +295,7 @@ CV_ENUM(MorphOp, MORPH_OPEN, MORPH_CLOSE, MORPH_GRADIENT, MORPH_TOPHAT, MORPH_BL DEF_PARAM_TEST(Sz_Type_Op, cv::Size, MatType, MorphOp); -PERF_TEST_P(Sz_Type_Op, Filters_MorphologyEx, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4), MorphOp::all())) +PERF_TEST_P(Sz_Type_Op, MorphologyEx, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4), MorphOp::all())) { declare.time(20.0); @@ -332,7 +332,7 @@ PERF_TEST_P(Sz_Type_Op, Filters_MorphologyEx, Combine(GPU_TYPICAL_MAT_SIZES, Val ////////////////////////////////////////////////////////////////////// // Filter2D -PERF_TEST_P(Sz_Type_KernelSz, Filters_Filter2D, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(3, 5, 7, 9, 11, 13, 15))) +PERF_TEST_P(Sz_Type_KernelSz, Filter2D, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(3, 5, 7, 9, 11, 13, 15))) { declare.time(20.0); diff --git a/modules/gpufilters/perf/perf_main.cpp b/modules/gpufilters/perf/perf_main.cpp new file mode 100644 index 0000000..b5a3eda --- /dev/null +++ b/modules/gpufilters/perf/perf_main.cpp @@ -0,0 +1,47 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "perf_precomp.hpp" + +using namespace perf; + +CV_PERF_TEST_MAIN(gpufilters, printCudaInfo()) diff --git a/modules/gpufilters/perf/perf_precomp.cpp b/modules/gpufilters/perf/perf_precomp.cpp new file mode 100644 index 0000000..81f16e8 --- /dev/null +++ b/modules/gpufilters/perf/perf_precomp.cpp @@ -0,0 +1,43 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "perf_precomp.hpp" diff --git a/modules/gpufilters/perf/perf_precomp.hpp b/modules/gpufilters/perf/perf_precomp.hpp new file mode 100644 index 0000000..02ca5ce --- /dev/null +++ b/modules/gpufilters/perf/perf_precomp.hpp @@ -0,0 +1,64 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifdef __GNUC__ +# pragma GCC diagnostic ignored "-Wmissing-declarations" +# if defined __clang__ || defined __APPLE__ +# pragma GCC diagnostic ignored "-Wmissing-prototypes" +# pragma GCC diagnostic ignored "-Wextra" +# endif +#endif + +#ifndef __OPENCV_PERF_PRECOMP_HPP__ +#define __OPENCV_PERF_PRECOMP_HPP__ + +#include "opencv2/ts.hpp" +#include "opencv2/ts/gpu_perf.hpp" + +#include "opencv2/gpufilters.hpp" +#include "opencv2/imgproc.hpp" + +#ifdef GTEST_CREATE_SHARED_LIBRARY +#error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined +#endif + +#endif diff --git a/modules/gpu/src/cuda/column_filter.8.cu b/modules/gpufilters/src/cuda/column_filter.16sc1.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.8.cu rename to modules/gpufilters/src/cuda/column_filter.16sc1.cu index 0a63a1d..d4c6d19 100644 --- a/modules/gpu/src/cuda/column_filter.8.cu +++ b/modules/gpufilters/src/cuda/column_filter.16sc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.3.cu b/modules/gpufilters/src/cuda/column_filter.16sc3.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.3.cu rename to modules/gpufilters/src/cuda/column_filter.16sc3.cu index 7304565..419fdea 100644 --- a/modules/gpu/src/cuda/column_filter.3.cu +++ b/modules/gpufilters/src/cuda/column_filter.16sc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.9.cu b/modules/gpufilters/src/cuda/column_filter.16sc4.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.9.cu rename to modules/gpufilters/src/cuda/column_filter.16sc4.cu index 758d928..1caeb87 100644 --- a/modules/gpu/src/cuda/column_filter.9.cu +++ b/modules/gpufilters/src/cuda/column_filter.16sc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.10.cu b/modules/gpufilters/src/cuda/column_filter.16uc1.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.10.cu rename to modules/gpufilters/src/cuda/column_filter.16uc1.cu index b71e252..dc68b71 100644 --- a/modules/gpu/src/cuda/column_filter.10.cu +++ b/modules/gpufilters/src/cuda/column_filter.16uc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.11.cu b/modules/gpufilters/src/cuda/column_filter.16uc3.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.11.cu rename to modules/gpufilters/src/cuda/column_filter.16uc3.cu index ccfbf8e..f0a07d6 100644 --- a/modules/gpu/src/cuda/column_filter.11.cu +++ b/modules/gpufilters/src/cuda/column_filter.16uc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.12.cu b/modules/gpufilters/src/cuda/column_filter.16uc4.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.12.cu rename to modules/gpufilters/src/cuda/column_filter.16uc4.cu index a38f93b..638ef79 100644 --- a/modules/gpu/src/cuda/column_filter.12.cu +++ b/modules/gpufilters/src/cuda/column_filter.16uc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.5.cu b/modules/gpufilters/src/cuda/column_filter.32fc1.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.5.cu rename to modules/gpufilters/src/cuda/column_filter.32fc1.cu index a192660..aa30933 100644 --- a/modules/gpu/src/cuda/column_filter.5.cu +++ b/modules/gpufilters/src/cuda/column_filter.32fc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.6.cu b/modules/gpufilters/src/cuda/column_filter.32fc3.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.6.cu rename to modules/gpufilters/src/cuda/column_filter.32fc3.cu index f4f7c4f..c0ed3ac 100644 --- a/modules/gpu/src/cuda/column_filter.6.cu +++ b/modules/gpufilters/src/cuda/column_filter.32fc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.7.cu b/modules/gpufilters/src/cuda/column_filter.32fc4.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.7.cu rename to modules/gpufilters/src/cuda/column_filter.32fc4.cu index 9f94bed..f37f717 100644 --- a/modules/gpu/src/cuda/column_filter.7.cu +++ b/modules/gpufilters/src/cuda/column_filter.32fc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.4.cu b/modules/gpufilters/src/cuda/column_filter.32sc1.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.4.cu rename to modules/gpufilters/src/cuda/column_filter.32sc1.cu index 8c9db69..ee05205 100644 --- a/modules/gpu/src/cuda/column_filter.4.cu +++ b/modules/gpufilters/src/cuda/column_filter.32sc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.13.cu b/modules/gpufilters/src/cuda/column_filter.32sc3.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.13.cu rename to modules/gpufilters/src/cuda/column_filter.32sc3.cu index 40eec7a..b921d96 100644 --- a/modules/gpu/src/cuda/column_filter.13.cu +++ b/modules/gpufilters/src/cuda/column_filter.32sc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.14.cu b/modules/gpufilters/src/cuda/column_filter.32sc4.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.14.cu rename to modules/gpufilters/src/cuda/column_filter.32sc4.cu index 08151ac..dd21524 100644 --- a/modules/gpu/src/cuda/column_filter.14.cu +++ b/modules/gpufilters/src/cuda/column_filter.32sc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.0.cu b/modules/gpufilters/src/cuda/column_filter.8uc1.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.0.cu rename to modules/gpufilters/src/cuda/column_filter.8uc1.cu index 339fb80..470f3ee 100644 --- a/modules/gpu/src/cuda/column_filter.0.cu +++ b/modules/gpufilters/src/cuda/column_filter.8uc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.1.cu b/modules/gpufilters/src/cuda/column_filter.8uc3.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.1.cu rename to modules/gpufilters/src/cuda/column_filter.8uc3.cu index 53914a2..5d5be58 100644 --- a/modules/gpu/src/cuda/column_filter.1.cu +++ b/modules/gpufilters/src/cuda/column_filter.8uc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.2.cu b/modules/gpufilters/src/cuda/column_filter.8uc4.cu similarity index 98% rename from modules/gpu/src/cuda/column_filter.2.cu rename to modules/gpufilters/src/cuda/column_filter.8uc4.cu index a615944..8a322f2 100644 --- a/modules/gpu/src/cuda/column_filter.2.cu +++ b/modules/gpufilters/src/cuda/column_filter.8uc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "column_filter.h" +#include "column_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/column_filter.h b/modules/gpufilters/src/cuda/column_filter.hpp similarity index 100% rename from modules/gpu/src/cuda/column_filter.h rename to modules/gpufilters/src/cuda/column_filter.hpp diff --git a/modules/gpufilters/src/cuda/filter2d.cu b/modules/gpufilters/src/cuda/filter2d.cu new file mode 100644 index 0000000..0bb5fcd --- /dev/null +++ b/modules/gpufilters/src/cuda/filter2d.cu @@ -0,0 +1,158 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#if !defined CUDA_DISABLER + +#include "opencv2/core/cuda/common.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" + +namespace cv { namespace gpu { namespace cudev +{ + namespace imgproc + { + #define FILTER2D_MAX_KERNEL_SIZE 16 + + __constant__ float c_filter2DKernel[FILTER2D_MAX_KERNEL_SIZE * FILTER2D_MAX_KERNEL_SIZE]; + + template + __global__ void filter2D(const SrcT src, PtrStepSz dst, const int kWidth, const int kHeight, const int anchorX, const int anchorY) + { + typedef typename TypeVec::cn>::vec_type sum_t; + + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x >= dst.cols || y >= dst.rows) + return; + + sum_t res = VecTraits::all(0); + int kInd = 0; + + for (int i = 0; i < kHeight; ++i) + { + for (int j = 0; j < kWidth; ++j) + res = res + src(y - anchorY + i, x - anchorX + j) * c_filter2DKernel[kInd++]; + } + + dst(y, x) = saturate_cast(res); + } + + template class Brd> struct Filter2DCaller; + + #define IMPLEMENT_FILTER2D_TEX_READER(type) \ + texture< type , cudaTextureType2D, cudaReadModeElementType> tex_filter2D_ ## type (0, cudaFilterModePoint, cudaAddressModeClamp); \ + struct tex_filter2D_ ## type ## _reader \ + { \ + typedef type elem_type; \ + typedef int index_type; \ + const int xoff; \ + const int yoff; \ + tex_filter2D_ ## type ## _reader (int xoff_, int yoff_) : xoff(xoff_), yoff(yoff_) {} \ + __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const \ + { \ + return tex2D(tex_filter2D_ ## type , x + xoff, y + yoff); \ + } \ + }; \ + template class Brd> struct Filter2DCaller< type , D, Brd> \ + { \ + static void call(const PtrStepSz< type > srcWhole, int xoff, int yoff, PtrStepSz dst, \ + int kWidth, int kHeight, int anchorX, int anchorY, const float* borderValue, cudaStream_t stream) \ + { \ + typedef typename TypeVec::cn>::vec_type work_type; \ + dim3 block(16, 16); \ + dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \ + bindTexture(&tex_filter2D_ ## type , srcWhole); \ + tex_filter2D_ ## type ##_reader texSrc(xoff, yoff); \ + Brd brd(dst.rows, dst.cols, VecTraits::make(borderValue)); \ + BorderReader< tex_filter2D_ ## type ##_reader, Brd > brdSrc(texSrc, brd); \ + filter2D<<>>(brdSrc, dst, kWidth, kHeight, anchorX, anchorY); \ + cudaSafeCall( cudaGetLastError() ); \ + if (stream == 0) \ + cudaSafeCall( cudaDeviceSynchronize() ); \ + } \ + }; + + IMPLEMENT_FILTER2D_TEX_READER(uchar); + IMPLEMENT_FILTER2D_TEX_READER(uchar4); + + IMPLEMENT_FILTER2D_TEX_READER(ushort); + IMPLEMENT_FILTER2D_TEX_READER(ushort4); + + IMPLEMENT_FILTER2D_TEX_READER(float); + IMPLEMENT_FILTER2D_TEX_READER(float4); + + #undef IMPLEMENT_FILTER2D_TEX_READER + + template + void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, + int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, + int borderMode, const float* borderValue, cudaStream_t stream) + { + typedef void (*func_t)(const PtrStepSz srcWhole, int xoff, int yoff, PtrStepSz dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* borderValue, cudaStream_t stream); + static const func_t funcs[] = + { + Filter2DCaller::call, + Filter2DCaller::call, + Filter2DCaller::call, + Filter2DCaller::call, + Filter2DCaller::call + }; + + if (stream == 0) + cudaSafeCall( cudaMemcpyToSymbol(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice) ); + else + cudaSafeCall( cudaMemcpyToSymbolAsync(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) ); + + funcs[borderMode](static_cast< PtrStepSz >(srcWhole), ofsX, ofsY, static_cast< PtrStepSz >(dst), kWidth, kHeight, anchorX, anchorY, borderValue, stream); + } + + template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); + template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); + template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); + template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); + template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); + template void filter2D_gpu(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream); + } +}}} + +#endif // CUDA_DISABLER diff --git a/modules/gpu/src/cuda/row_filter.8.cu b/modules/gpufilters/src/cuda/row_filter.16sc1.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.8.cu rename to modules/gpufilters/src/cuda/row_filter.16sc1.cu index b899e87..59ebb9f 100644 --- a/modules/gpu/src/cuda/row_filter.8.cu +++ b/modules/gpufilters/src/cuda/row_filter.16sc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.3.cu b/modules/gpufilters/src/cuda/row_filter.16sc3.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.3.cu rename to modules/gpufilters/src/cuda/row_filter.16sc3.cu index fe84666..fcf40d8 100644 --- a/modules/gpu/src/cuda/row_filter.3.cu +++ b/modules/gpufilters/src/cuda/row_filter.16sc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.9.cu b/modules/gpufilters/src/cuda/row_filter.16sc4.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.9.cu rename to modules/gpufilters/src/cuda/row_filter.16sc4.cu index 516dd8f..c5d4726 100644 --- a/modules/gpu/src/cuda/row_filter.9.cu +++ b/modules/gpufilters/src/cuda/row_filter.16sc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.10.cu b/modules/gpufilters/src/cuda/row_filter.16uc1.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.10.cu rename to modules/gpufilters/src/cuda/row_filter.16uc1.cu index 7d93ee3..02e125a 100644 --- a/modules/gpu/src/cuda/row_filter.10.cu +++ b/modules/gpufilters/src/cuda/row_filter.16uc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.11.cu b/modules/gpufilters/src/cuda/row_filter.16uc3.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.11.cu rename to modules/gpufilters/src/cuda/row_filter.16uc3.cu index 31bccc4..494c604 100644 --- a/modules/gpu/src/cuda/row_filter.11.cu +++ b/modules/gpufilters/src/cuda/row_filter.16uc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.12.cu b/modules/gpufilters/src/cuda/row_filter.16uc4.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.12.cu rename to modules/gpufilters/src/cuda/row_filter.16uc4.cu index 7be543f..1eb1ac2 100644 --- a/modules/gpu/src/cuda/row_filter.12.cu +++ b/modules/gpufilters/src/cuda/row_filter.16uc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.5.cu b/modules/gpufilters/src/cuda/row_filter.32fc1.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.5.cu rename to modules/gpufilters/src/cuda/row_filter.32fc1.cu index 975aea4..bf577c6 100644 --- a/modules/gpu/src/cuda/row_filter.5.cu +++ b/modules/gpufilters/src/cuda/row_filter.32fc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.6.cu b/modules/gpufilters/src/cuda/row_filter.32fc3.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.6.cu rename to modules/gpufilters/src/cuda/row_filter.32fc3.cu index d589445..594fc04 100644 --- a/modules/gpu/src/cuda/row_filter.6.cu +++ b/modules/gpufilters/src/cuda/row_filter.32fc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.7.cu b/modules/gpufilters/src/cuda/row_filter.32fc4.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.7.cu rename to modules/gpufilters/src/cuda/row_filter.32fc4.cu index ac3fcc1..5f2812b 100644 --- a/modules/gpu/src/cuda/row_filter.7.cu +++ b/modules/gpufilters/src/cuda/row_filter.32fc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.4.cu b/modules/gpufilters/src/cuda/row_filter.32sc1.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.4.cu rename to modules/gpufilters/src/cuda/row_filter.32sc1.cu index 050f7af..67f3fb0 100644 --- a/modules/gpu/src/cuda/row_filter.4.cu +++ b/modules/gpufilters/src/cuda/row_filter.32sc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.13.cu b/modules/gpufilters/src/cuda/row_filter.32sc3.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.13.cu rename to modules/gpufilters/src/cuda/row_filter.32sc3.cu index bd700b1..8e881a2 100644 --- a/modules/gpu/src/cuda/row_filter.13.cu +++ b/modules/gpufilters/src/cuda/row_filter.32sc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.14.cu b/modules/gpufilters/src/cuda/row_filter.32sc4.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.14.cu rename to modules/gpufilters/src/cuda/row_filter.32sc4.cu index 97df2f1..66f00cf 100644 --- a/modules/gpu/src/cuda/row_filter.14.cu +++ b/modules/gpufilters/src/cuda/row_filter.32sc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.0.cu b/modules/gpufilters/src/cuda/row_filter.8uc1.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.0.cu rename to modules/gpufilters/src/cuda/row_filter.8uc1.cu index a4b4239..c94b39f 100644 --- a/modules/gpu/src/cuda/row_filter.0.cu +++ b/modules/gpufilters/src/cuda/row_filter.8uc1.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.1.cu b/modules/gpufilters/src/cuda/row_filter.8uc3.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.1.cu rename to modules/gpufilters/src/cuda/row_filter.8uc3.cu index ac4724f..1c924c1 100644 --- a/modules/gpu/src/cuda/row_filter.1.cu +++ b/modules/gpufilters/src/cuda/row_filter.8uc3.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.2.cu b/modules/gpufilters/src/cuda/row_filter.8uc4.cu similarity index 98% rename from modules/gpu/src/cuda/row_filter.2.cu rename to modules/gpufilters/src/cuda/row_filter.8uc4.cu index d630b6f..1ae9651 100644 --- a/modules/gpu/src/cuda/row_filter.2.cu +++ b/modules/gpufilters/src/cuda/row_filter.8uc4.cu @@ -42,7 +42,7 @@ #if !defined CUDA_DISABLER -#include "row_filter.h" +#include "row_filter.hpp" namespace filter { diff --git a/modules/gpu/src/cuda/row_filter.h b/modules/gpufilters/src/cuda/row_filter.hpp similarity index 100% rename from modules/gpu/src/cuda/row_filter.h rename to modules/gpufilters/src/cuda/row_filter.hpp diff --git a/modules/gpu/src/filtering.cpp b/modules/gpufilters/src/filtering.cpp similarity index 99% rename from modules/gpu/src/filtering.cpp rename to modules/gpufilters/src/filtering.cpp index 33d179b..6416325 100644 --- a/modules/gpu/src/filtering.cpp +++ b/modules/gpufilters/src/filtering.cpp @@ -45,7 +45,6 @@ using namespace cv; using namespace cv::gpu; - #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) Ptr cv::gpu::createFilter2D_GPU(const Ptr&, int, int) { throw_no_cuda(); return Ptr(0); } @@ -628,31 +627,44 @@ void cv::gpu::morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& ke { 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_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; + +#ifdef HAVE_OPENCV_GPUARITHM case MORPH_GRADIENT: erode(src, buf2, kernel, buf1, anchor, iterations, stream); dilate(src, dst, kernel, buf1, anchor, iterations, stream); - subtract(dst, buf2, dst, GpuMat(), -1, 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); - subtract(src, buf2, dst, GpuMat(), -1, 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); - subtract(buf2, src, dst, GpuMat(), -1, 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.cpp b/modules/gpufilters/src/precomp.cpp new file mode 100644 index 0000000..3c01a25 --- /dev/null +++ b/modules/gpufilters/src/precomp.cpp @@ -0,0 +1,43 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "precomp.hpp" diff --git a/modules/gpufilters/src/precomp.hpp b/modules/gpufilters/src/precomp.hpp new file mode 100644 index 0000000..e8e46ff --- /dev/null +++ b/modules/gpufilters/src/precomp.hpp @@ -0,0 +1,59 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __OPENCV_PRECOMP_H__ +#define __OPENCV_PRECOMP_H__ + +#include + +#include "opencv2/gpufilters.hpp" +#include "opencv2/imgproc.hpp" + +#include "opencv2/core/gpu_private.hpp" + +#include "opencv2/opencv_modules.hpp" + +#ifdef HAVE_OPENCV_GPUARITHM +# include "opencv2/gpuarithm.hpp" +#endif + +#endif /* __OPENCV_PRECOMP_H__ */ diff --git a/modules/gpu/test/test_filters.cpp b/modules/gpufilters/test/test_filters.cpp similarity index 96% rename from modules/gpu/test/test_filters.cpp rename to modules/gpufilters/test/test_filters.cpp index 1c0679e..5adcd87 100644 --- a/modules/gpu/test/test_filters.cpp +++ b/modules/gpufilters/test/test_filters.cpp @@ -105,7 +105,7 @@ GPU_TEST_P(Blur, Accuracy) EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 1.0); } -INSTANTIATE_TEST_CASE_P(GPU_Filter, Blur, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Filters, Blur, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), @@ -164,7 +164,7 @@ GPU_TEST_P(Sobel, Accuracy) EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), CV_MAT_DEPTH(type) < CV_32F ? 0.0 : 0.1); } -INSTANTIATE_TEST_CASE_P(GPU_Filter, Sobel, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Filters, Sobel, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_16S), MatDepth(CV_32F)), @@ -227,7 +227,7 @@ GPU_TEST_P(Scharr, Accuracy) EXPECT_MAT_NEAR(getInnerROI(dst_gold, cv::Size(3, 3)), getInnerROI(dst, cv::Size(3, 3)), CV_MAT_DEPTH(type) < CV_32F ? 0.0 : 0.1); } -INSTANTIATE_TEST_CASE_P(GPU_Filter, Scharr, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Filters, Scharr, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_16S), MatDepth(CV_32F)), @@ -301,7 +301,7 @@ GPU_TEST_P(GaussianBlur, Accuracy) } } -INSTANTIATE_TEST_CASE_P(GPU_Filter, GaussianBlur, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Filters, GaussianBlur, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_16S), MatDepth(CV_32F)), @@ -363,7 +363,7 @@ GPU_TEST_P(Laplacian, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, src.depth() < CV_32F ? 0.0 : 1e-3); } -INSTANTIATE_TEST_CASE_P(GPU_Filter, Laplacian, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Filters, Laplacian, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)), @@ -411,7 +411,7 @@ GPU_TEST_P(Erode, Accuracy) EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0); } -INSTANTIATE_TEST_CASE_P(GPU_Filter, Erode, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Filters, Erode, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), @@ -460,7 +460,7 @@ GPU_TEST_P(Dilate, Accuracy) EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0); } -INSTANTIATE_TEST_CASE_P(GPU_Filter, Dilate, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Filters, Dilate, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), @@ -513,7 +513,7 @@ GPU_TEST_P(MorphEx, Accuracy) EXPECT_MAT_NEAR(getInnerROI(dst_gold, border), getInnerROI(dst, border), 0.0); } -INSTANTIATE_TEST_CASE_P(GPU_Filter, MorphEx, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Filters, MorphEx, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), @@ -565,7 +565,7 @@ GPU_TEST_P(Filter2D, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, CV_MAT_DEPTH(type) == CV_32F ? 1e-1 : 1.0); } -INSTANTIATE_TEST_CASE_P(GPU_Filter, Filter2D, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Filters, Filter2D, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_16UC1), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC4)), diff --git a/modules/gpufilters/test/test_main.cpp b/modules/gpufilters/test/test_main.cpp new file mode 100644 index 0000000..c37a85c --- /dev/null +++ b/modules/gpufilters/test/test_main.cpp @@ -0,0 +1,120 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "test_precomp.hpp" + +#ifdef HAVE_CUDA + +using namespace std; +using namespace cv; +using namespace cv::gpu; +using namespace cvtest; +using namespace testing; + +int main(int argc, char** argv) +{ + try + { + const std::string keys = + "{ h help ? | | Print help}" + "{ i info | | Print information about system and exit }" + "{ device | -1 | Device on which tests will be executed (-1 means all devices) }" + ; + + CommandLineParser cmd(argc, (const char**)argv, keys); + + if (cmd.has("help")) + { + cmd.printMessage(); + return 0; + } + + printCudaInfo(); + + if (cmd.has("info")) + { + return 0; + } + + int device = cmd.get("device"); + if (device < 0) + { + DeviceManager::instance().loadAll(); + + cout << "Run tests on all supported devices \n" << endl; + } + else + { + DeviceManager::instance().load(device); + + DeviceInfo info(device); + cout << "Run tests on device " << device << " [" << info.name() << "] \n" << endl; + } + + TS::ptr()->init("gpu"); + InitGoogleTest(&argc, argv); + + return RUN_ALL_TESTS(); + } + catch (const exception& e) + { + cerr << e.what() << endl; + return -1; + } + catch (...) + { + cerr << "Unknown error" << endl; + return -1; + } + + return 0; +} + +#else // HAVE_CUDA + +int main() +{ + printf("OpenCV was built without CUDA support\n"); + return 0; +} + +#endif // HAVE_CUDA diff --git a/modules/gpufilters/test/test_precomp.cpp b/modules/gpufilters/test/test_precomp.cpp new file mode 100644 index 0000000..0fb6521 --- /dev/null +++ b/modules/gpufilters/test/test_precomp.cpp @@ -0,0 +1,43 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "test_precomp.hpp" diff --git a/modules/gpufilters/test/test_precomp.hpp b/modules/gpufilters/test/test_precomp.hpp new file mode 100644 index 0000000..9598492 --- /dev/null +++ b/modules/gpufilters/test/test_precomp.hpp @@ -0,0 +1,60 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifdef __GNUC__ +# pragma GCC diagnostic ignored "-Wmissing-declarations" +# if defined __clang__ || defined __APPLE__ +# pragma GCC diagnostic ignored "-Wmissing-prototypes" +# pragma GCC diagnostic ignored "-Wextra" +# endif +#endif + +#ifndef __OPENCV_TEST_PRECOMP_HPP__ +#define __OPENCV_TEST_PRECOMP_HPP__ + +#include "opencv2/ts.hpp" +#include "opencv2/ts/gpu_test.hpp" + +#include "opencv2/gpufilters.hpp" +#include "opencv2/imgproc.hpp" + +#endif diff --git a/modules/stitching/CMakeLists.txt b/modules/stitching/CMakeLists.txt index 2282d81..647d8b1 100644 --- a/modules/stitching/CMakeLists.txt +++ b/modules/stitching/CMakeLists.txt @@ -1,3 +1,3 @@ set(the_description "Images stitching") -ocv_define_module(stitching opencv_imgproc opencv_features2d opencv_calib3d opencv_objdetect OPTIONAL opencv_gpu opencv_gpuarithm opencv_nonfree) +ocv_define_module(stitching opencv_imgproc opencv_features2d opencv_calib3d opencv_objdetect OPTIONAL opencv_gpu opencv_gpuarithm opencv_gpufilters opencv_nonfree) diff --git a/modules/superres/CMakeLists.txt b/modules/superres/CMakeLists.txt index fae24e1..378a2a9 100644 --- a/modules/superres/CMakeLists.txt +++ b/modules/superres/CMakeLists.txt @@ -4,4 +4,4 @@ endif() set(the_description "Super Resolution") ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 -Wundef) -ocv_define_module(superres opencv_imgproc opencv_video OPTIONAL opencv_highgui opencv_gpu opencv_gpucodec opencv_gpuarithm) +ocv_define_module(superres opencv_imgproc opencv_video OPTIONAL opencv_highgui opencv_gpu opencv_gpuarithm opencv_gpufilters opencv_gpucodec) diff --git a/samples/cpp/CMakeLists.txt b/samples/cpp/CMakeLists.txt index 92f7dc7..4678532 100644 --- a/samples/cpp/CMakeLists.txt +++ b/samples/cpp/CMakeLists.txt @@ -18,6 +18,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND) if(HAVE_opencv_gpu) ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpuarithm/include") + ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpufilters/include") ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpu/include") endif() diff --git a/samples/gpu/CMakeLists.txt b/samples/gpu/CMakeLists.txt index efff5f9..760bc26 100644 --- a/samples/gpu/CMakeLists.txt +++ b/samples/gpu/CMakeLists.txt @@ -2,7 +2,7 @@ SET(OPENCV_GPU_SAMPLES_REQUIRED_DEPS opencv_core opencv_flann opencv_imgproc ope opencv_ml opencv_video opencv_objdetect opencv_features2d opencv_calib3d opencv_legacy opencv_contrib opencv_gpu opencv_nonfree opencv_softcascade opencv_superres - opencv_gpucodec opencv_gpuarithm) + opencv_gpucodec opencv_gpuarithm opencv_gpufilters) ocv_check_dependencies(${OPENCV_GPU_SAMPLES_REQUIRED_DEPS}) -- 2.7.4