diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt index de132cf9fb..578957037d 100644 --- a/modules/gpu/CMakeLists.txt +++ b/modules/gpu/CMakeLists.txt @@ -6,7 +6,7 @@ set(the_description "GPU-accelerated Computer Vision") ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations -Wshadow -Wunused-parameter) -ocv_define_module(gpu opencv_gpuarithm opencv_gpufilters opencv_gpuimgproc +ocv_define_module(gpu opencv_gpuarithm opencv_gpufilters opencv_gpuwarping opencv_gpuimgproc opencv_gpufeatures2d opencv_gpuvideo opencv_gpucalib3d opencv_gpuobjdetect) if(HAVE_CUDA) diff --git a/modules/gpu/include/opencv2/gpu.hpp b/modules/gpu/include/opencv2/gpu.hpp index e2f747806d..10fbbd7d81 100644 --- a/modules/gpu/include/opencv2/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu.hpp @@ -46,6 +46,7 @@ #include "opencv2/core/gpumat.hpp" #include "opencv2/gpuarithm.hpp" #include "opencv2/gpufilters.hpp" +#include "opencv2/gpuwarping.hpp" #include "opencv2/gpuimgproc.hpp" #include "opencv2/gpufeatures2d.hpp" #include "opencv2/gpuvideo.hpp" diff --git a/modules/gpuarithm/doc/operations_on_matrices.rst b/modules/gpuarithm/doc/operations_on_matrices.rst index d1762f442a..a25100728a 100644 --- a/modules/gpuarithm/doc/operations_on_matrices.rst +++ b/modules/gpuarithm/doc/operations_on_matrices.rst @@ -272,3 +272,166 @@ Normalizes the norm or value range of an array. :param cvt_buf: Optional buffer to avoid extra memory allocations. It is resized automatically. .. seealso:: :ocv:func:`normalize` + + + +gpu::mulSpectrums +--------------------- +Performs a per-element multiplication of two Fourier spectrums. + +.. ocv:function:: void gpu::mulSpectrums( const GpuMat& a, const GpuMat& b, GpuMat& c, int flags, bool conjB=false, Stream& stream=Stream::Null() ) + + :param a: First spectrum. + + :param b: Second spectrum with the same size and type as ``a`` . + + :param c: Destination spectrum. + + :param flags: Mock parameter used for CPU/GPU interfaces similarity. + + :param conjB: Optional flag to specify if the second spectrum needs to be conjugated before the multiplication. + + Only full (not packed) ``CV_32FC2`` complex spectrums in the interleaved format are supported for now. + +.. seealso:: :ocv:func:`mulSpectrums` + + + +gpu::mulAndScaleSpectrums +----------------------------- +Performs a per-element multiplication of two Fourier spectrums and scales the result. + +.. ocv:function:: void gpu::mulAndScaleSpectrums( const GpuMat& a, const GpuMat& b, GpuMat& c, int flags, float scale, bool conjB=false, Stream& stream=Stream::Null() ) + + :param a: First spectrum. + + :param b: Second spectrum with the same size and type as ``a`` . + + :param c: Destination spectrum. + + :param flags: Mock parameter used for CPU/GPU interfaces similarity. + + :param scale: Scale constant. + + :param conjB: Optional flag to specify if the second spectrum needs to be conjugated before the multiplication. + + Only full (not packed) ``CV_32FC2`` complex spectrums in the interleaved format are supported for now. + +.. seealso:: :ocv:func:`mulSpectrums` + + + +gpu::dft +------------ +Performs a forward or inverse discrete Fourier transform (1D or 2D) of the floating point matrix. + +.. ocv:function:: void gpu::dft( const GpuMat& src, GpuMat& dst, Size dft_size, int flags=0, Stream& stream=Stream::Null() ) + + :param src: Source matrix (real or complex). + + :param dst: Destination matrix (real or complex). + + :param dft_size: Size of a discrete Fourier transform. + + :param flags: Optional flags: + + * **DFT_ROWS** transforms each individual row of the source matrix. + + * **DFT_SCALE** scales the result: divide it by the number of elements in the transform (obtained from ``dft_size`` ). + + * **DFT_INVERSE** inverts DFT. Use for complex-complex cases (real-complex and complex-real cases are always forward and inverse, respectively). + + * **DFT_REAL_OUTPUT** specifies the output as real. The source matrix is the result of real-complex transform, so the destination matrix must be real. + +Use to handle real matrices ( ``CV32FC1`` ) and complex matrices in the interleaved format ( ``CV32FC2`` ). + +The source matrix should be continuous, otherwise reallocation and data copying is performed. The function chooses an operation mode depending on the flags, size, and channel count of the source matrix: + + * If the source matrix is complex and the output is not specified as real, the destination matrix is complex and has the ``dft_size`` size and ``CV_32FC2`` type. The destination matrix contains a full result of the DFT (forward or inverse). + + * If the source matrix is complex and the output is specified as real, the function assumes that its input is the result of the forward transform (see the next item). The destination matrix has the ``dft_size`` size and ``CV_32FC1`` type. It contains the result of the inverse DFT. + + * If the source matrix is real (its type is ``CV_32FC1`` ), forward DFT is performed. The result of the DFT is packed into complex ( ``CV_32FC2`` ) matrix. So, the width of the destination matrix is ``dft_size.width / 2 + 1`` . But if the source is a single column, the height is reduced instead of the width. + +.. seealso:: :ocv:func:`dft` + + +gpu::ConvolveBuf +---------------- +.. ocv:struct:: gpu::ConvolveBuf + +Class providing a memory buffer for :ocv:func:`gpu::convolve` function, plus it allows to adjust some specific parameters. :: + + struct CV_EXPORTS ConvolveBuf + { + Size result_size; + Size block_size; + Size user_block_size; + Size dft_size; + int spect_len; + + GpuMat image_spect, templ_spect, result_spect; + GpuMat image_block, templ_block, result_data; + + void create(Size image_size, Size templ_size); + static Size estimateBlockSize(Size result_size, Size templ_size); + }; + +You can use field `user_block_size` to set specific block size for :ocv:func:`gpu::convolve` function. If you leave its default value `Size(0,0)` then automatic estimation of block size will be used (which is optimized for speed). By varying `user_block_size` you can reduce memory requirements at the cost of speed. + +gpu::ConvolveBuf::create +------------------------ +.. ocv:function:: gpu::ConvolveBuf::create(Size image_size, Size templ_size) + +Constructs a buffer for :ocv:func:`gpu::convolve` function with respective arguments. + + +gpu::convolve +----------------- +Computes a convolution (or cross-correlation) of two images. + +.. ocv:function:: void gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, bool ccorr=false) + +.. ocv:function:: void gpu::convolve( const GpuMat& image, const GpuMat& templ, GpuMat& result, bool ccorr, ConvolveBuf& buf, Stream& stream=Stream::Null() ) + + :param image: Source image. Only ``CV_32FC1`` images are supported for now. + + :param templ: Template image. The size is not greater than the ``image`` size. The type is the same as ``image`` . + + :param result: Result image. If ``image`` is *W x H* and ``templ`` is *w x h*, then ``result`` must be *W-w+1 x H-h+1*. + + :param ccorr: Flags to evaluate cross-correlation instead of convolution. + + :param buf: Optional buffer to avoid extra memory allocations and to adjust some specific parameters. See :ocv:struct:`gpu::ConvolveBuf`. + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`gpu::filter2D` + + + +gpu::copyMakeBorder +----------------------- +Forms a border around an image. + +.. ocv:function:: void gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value = Scalar(), Stream& stream = Stream::Null()) + + :param src: Source image. ``CV_8UC1`` , ``CV_8UC4`` , ``CV_32SC1`` , and ``CV_32FC1`` types are supported. + + :param dst: Destination image with the same type as ``src``. The size is ``Size(src.cols+left+right, src.rows+top+bottom)`` . + + :param top: + + :param bottom: + + :param left: + + :param right: Number of pixels in each direction from the source image rectangle to extrapolate. For example: ``top=1, bottom=1, left=1, right=1`` mean that 1 pixel-wide border needs to be built. + + :param borderType: Border type. See :ocv:func:`borderInterpolate` for details. ``BORDER_REFLECT101`` , ``BORDER_REPLICATE`` , ``BORDER_CONSTANT`` , ``BORDER_REFLECT`` and ``BORDER_WRAP`` are supported for now. + + :param value: Border value. + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`copyMakeBorder` diff --git a/modules/gpuarithm/perf/perf_core.cpp b/modules/gpuarithm/perf/perf_core.cpp index 8957d06d5e..bb1a89d03c 100644 --- a/modules/gpuarithm/perf/perf_core.cpp +++ b/modules/gpuarithm/perf/perf_core.cpp @@ -2411,7 +2411,7 @@ CV_ENUM(ThreshOp, cv::THRESH_BINARY, cv::THRESH_BINARY_INV, cv::THRESH_TRUNC, cv DEF_PARAM_TEST(Sz_Depth_Op, cv::Size, MatDepth, ThreshOp); -PERF_TEST_P(Sz_Depth_Op, Threshold, +PERF_TEST_P(Sz_Depth_Op, ImgProc_Threshold, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8U, CV_16U, CV_32F, CV_64F), ALL_THRESH_OPS)) diff --git a/modules/gpufeatures2d/CMakeLists.txt b/modules/gpufeatures2d/CMakeLists.txt index 4a93be34af..7162fae158 100644 --- a/modules/gpufeatures2d/CMakeLists.txt +++ b/modules/gpufeatures2d/CMakeLists.txt @@ -6,4 +6,4 @@ set(the_description "GPU-accelerated Feature Detection and Description") ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4100 /wd4324 /wd4512 -Wundef -Wmissing-declarations -Wshadow -Wunused-parameter) -ocv_define_module(gpufeatures2d opencv_features2d opencv_gpufilters opencv_gpuimgproc) +ocv_define_module(gpufeatures2d opencv_features2d opencv_gpufilters opencv_gpuwarping) diff --git a/modules/gpufeatures2d/src/precomp.hpp b/modules/gpufeatures2d/src/precomp.hpp index d3936264b1..9fbecc71ec 100644 --- a/modules/gpufeatures2d/src/precomp.hpp +++ b/modules/gpufeatures2d/src/precomp.hpp @@ -49,8 +49,7 @@ #include "opencv2/gpufeatures2d.hpp" #include "opencv2/gpuarithm.hpp" -#include "opencv2/gpuimgproc.hpp" - +#include "opencv2/gpuwarping.hpp" #include "opencv2/features2d.hpp" #include "opencv2/core/gpu_private.hpp" diff --git a/modules/gpuimgproc/doc/image_processing.rst b/modules/gpuimgproc/doc/image_processing.rst index 69e5003743..52d7b92adc 100644 --- a/modules/gpuimgproc/doc/image_processing.rst +++ b/modules/gpuimgproc/doc/image_processing.rst @@ -97,18 +97,6 @@ Computes a squared integral image. -gpu::columnSum ------------------- -Computes a vertical (column) sum. - -.. ocv:function:: void gpu::columnSum(const GpuMat& src, GpuMat& sum) - - :param src: Source image. Only ``CV_32FC1`` images are supported for now. - - :param sum: Destination image of the ``CV_32FC1`` type. - - - gpu::cornerHarris --------------------- Computes the Harris cornerness criteria at each image pixel. @@ -155,139 +143,6 @@ Computes the minimum eigen value of a 2x2 derivative covariation matrix at each -gpu::mulSpectrums ---------------------- -Performs a per-element multiplication of two Fourier spectrums. - -.. ocv:function:: void gpu::mulSpectrums( const GpuMat& a, const GpuMat& b, GpuMat& c, int flags, bool conjB=false, Stream& stream=Stream::Null() ) - - :param a: First spectrum. - - :param b: Second spectrum with the same size and type as ``a`` . - - :param c: Destination spectrum. - - :param flags: Mock parameter used for CPU/GPU interfaces similarity. - - :param conjB: Optional flag to specify if the second spectrum needs to be conjugated before the multiplication. - - Only full (not packed) ``CV_32FC2`` complex spectrums in the interleaved format are supported for now. - -.. seealso:: :ocv:func:`mulSpectrums` - - - -gpu::mulAndScaleSpectrums ------------------------------ -Performs a per-element multiplication of two Fourier spectrums and scales the result. - -.. ocv:function:: void gpu::mulAndScaleSpectrums( const GpuMat& a, const GpuMat& b, GpuMat& c, int flags, float scale, bool conjB=false, Stream& stream=Stream::Null() ) - - :param a: First spectrum. - - :param b: Second spectrum with the same size and type as ``a`` . - - :param c: Destination spectrum. - - :param flags: Mock parameter used for CPU/GPU interfaces similarity. - - :param scale: Scale constant. - - :param conjB: Optional flag to specify if the second spectrum needs to be conjugated before the multiplication. - - Only full (not packed) ``CV_32FC2`` complex spectrums in the interleaved format are supported for now. - -.. seealso:: :ocv:func:`mulSpectrums` - - - -gpu::dft ------------- -Performs a forward or inverse discrete Fourier transform (1D or 2D) of the floating point matrix. - -.. ocv:function:: void gpu::dft( const GpuMat& src, GpuMat& dst, Size dft_size, int flags=0, Stream& stream=Stream::Null() ) - - :param src: Source matrix (real or complex). - - :param dst: Destination matrix (real or complex). - - :param dft_size: Size of a discrete Fourier transform. - - :param flags: Optional flags: - - * **DFT_ROWS** transforms each individual row of the source matrix. - - * **DFT_SCALE** scales the result: divide it by the number of elements in the transform (obtained from ``dft_size`` ). - - * **DFT_INVERSE** inverts DFT. Use for complex-complex cases (real-complex and complex-real cases are always forward and inverse, respectively). - - * **DFT_REAL_OUTPUT** specifies the output as real. The source matrix is the result of real-complex transform, so the destination matrix must be real. - -Use to handle real matrices ( ``CV32FC1`` ) and complex matrices in the interleaved format ( ``CV32FC2`` ). - -The source matrix should be continuous, otherwise reallocation and data copying is performed. The function chooses an operation mode depending on the flags, size, and channel count of the source matrix: - - * If the source matrix is complex and the output is not specified as real, the destination matrix is complex and has the ``dft_size`` size and ``CV_32FC2`` type. The destination matrix contains a full result of the DFT (forward or inverse). - - * If the source matrix is complex and the output is specified as real, the function assumes that its input is the result of the forward transform (see the next item). The destination matrix has the ``dft_size`` size and ``CV_32FC1`` type. It contains the result of the inverse DFT. - - * If the source matrix is real (its type is ``CV_32FC1`` ), forward DFT is performed. The result of the DFT is packed into complex ( ``CV_32FC2`` ) matrix. So, the width of the destination matrix is ``dft_size.width / 2 + 1`` . But if the source is a single column, the height is reduced instead of the width. - -.. seealso:: :ocv:func:`dft` - - -gpu::ConvolveBuf ----------------- -.. ocv:struct:: gpu::ConvolveBuf - -Class providing a memory buffer for :ocv:func:`gpu::convolve` function, plus it allows to adjust some specific parameters. :: - - struct CV_EXPORTS ConvolveBuf - { - Size result_size; - Size block_size; - Size user_block_size; - Size dft_size; - int spect_len; - - GpuMat image_spect, templ_spect, result_spect; - GpuMat image_block, templ_block, result_data; - - void create(Size image_size, Size templ_size); - static Size estimateBlockSize(Size result_size, Size templ_size); - }; - -You can use field `user_block_size` to set specific block size for :ocv:func:`gpu::convolve` function. If you leave its default value `Size(0,0)` then automatic estimation of block size will be used (which is optimized for speed). By varying `user_block_size` you can reduce memory requirements at the cost of speed. - -gpu::ConvolveBuf::create ------------------------- -.. ocv:function:: gpu::ConvolveBuf::create(Size image_size, Size templ_size) - -Constructs a buffer for :ocv:func:`gpu::convolve` function with respective arguments. - - -gpu::convolve ------------------ -Computes a convolution (or cross-correlation) of two images. - -.. ocv:function:: void gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, bool ccorr=false) - -.. ocv:function:: void gpu::convolve( const GpuMat& image, const GpuMat& templ, GpuMat& result, bool ccorr, ConvolveBuf& buf, Stream& stream=Stream::Null() ) - - :param image: Source image. Only ``CV_32FC1`` images are supported for now. - - :param templ: Template image. The size is not greater than the ``image`` size. The type is the same as ``image`` . - - :param result: Result image. If ``image`` is *W x H* and ``templ`` is *w x h*, then ``result`` must be *W-w+1 x H-h+1*. - - :param ccorr: Flags to evaluate cross-correlation instead of convolution. - - :param buf: Optional buffer to avoid extra memory allocations and to adjust some specific parameters. See :ocv:struct:`gpu::ConvolveBuf`. - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`gpu::filter2D` - gpu::MatchTemplateBuf --------------------- .. ocv:struct:: gpu::MatchTemplateBuf @@ -305,6 +160,8 @@ Class providing memory buffers for :ocv:func:`gpu::matchTemplate` function, plus You can use field `user_block_size` to set specific block size for :ocv:func:`gpu::matchTemplate` function. If you leave its default value `Size(0,0)` then automatic estimation of block size will be used (which is optimized for speed). By varying `user_block_size` you can reduce memory requirements at the cost of speed. + + gpu::matchTemplate ---------------------- Computes a proximity map for a raster template and an image where the template is searched for. @@ -342,39 +199,6 @@ Computes a proximity map for a raster template and an image where the template i .. seealso:: :ocv:func:`matchTemplate` -gpu::remap --------------- -Applies a generic geometrical transformation to an image. - -.. ocv:function:: void gpu::remap( const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap, int interpolation, int borderMode=BORDER_CONSTANT, Scalar borderValue=Scalar(), Stream& stream=Stream::Null() ) - - :param src: Source image. - - :param dst: Destination image with the size the same as ``xmap`` and the type the same as ``src`` . - - :param xmap: X values. Only ``CV_32FC1`` type is supported. - - :param ymap: Y values. Only ``CV_32FC1`` type is supported. - - :param interpolation: Interpolation method (see :ocv:func:`resize` ). ``INTER_NEAREST`` , ``INTER_LINEAR`` and ``INTER_CUBIC`` are supported for now. - - :param borderMode: Pixel extrapolation method (see :ocv:func:`borderInterpolate` ). ``BORDER_REFLECT101`` , ``BORDER_REPLICATE`` , ``BORDER_CONSTANT`` , ``BORDER_REFLECT`` and ``BORDER_WRAP`` are supported for now. - - :param borderValue: Value used in case of a constant border. By default, it is 0. - - :param stream: Stream for the asynchronous version. - -The function transforms the source image using the specified map: - -.. math:: - - \texttt{dst} (x,y) = \texttt{src} (xmap(x,y), ymap(x,y)) - -Values of pixels with non-integer coordinates are computed using the bilinear interpolation. - -.. seealso:: :ocv:func:`remap` - - gpu::cvtColor ----------------- @@ -414,185 +238,6 @@ The methods support arbitrary permutations of the original channels, including r -gpu::resize ---------------- -Resizes an image. - -.. ocv:function:: void gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, double fy=0, int interpolation = INTER_LINEAR, Stream& stream = Stream::Null()) - - :param src: Source image. - - :param dst: Destination image with the same type as ``src`` . The size is ``dsize`` (when it is non-zero) or the size is computed from ``src.size()`` , ``fx`` , and ``fy`` . - - :param dsize: Destination image size. If it is zero, it is computed as: - - .. math:: - \texttt{dsize = Size(round(fx*src.cols), round(fy*src.rows))} - - Either ``dsize`` or both ``fx`` and ``fy`` must be non-zero. - - :param fx: Scale factor along the horizontal axis. If it is zero, it is computed as: - - .. math:: - - \texttt{(double)dsize.width/src.cols} - - :param fy: Scale factor along the vertical axis. If it is zero, it is computed as: - - .. math:: - - \texttt{(double)dsize.height/src.rows} - - :param interpolation: Interpolation method. ``INTER_NEAREST`` , ``INTER_LINEAR`` and ``INTER_CUBIC`` are supported for now. - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`resize` - - - -gpu::warpAffine -------------------- -Applies an affine transformation to an image. - -.. ocv:function:: void gpu::warpAffine( const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags=INTER_LINEAR, int borderMode=BORDER_CONSTANT, Scalar borderValue=Scalar(), Stream& stream=Stream::Null() ) - - :param src: Source image. ``CV_8U`` , ``CV_16U`` , ``CV_32S`` , or ``CV_32F`` depth and 1, 3, or 4 channels are supported. - - :param dst: Destination image with the same type as ``src`` . The size is ``dsize`` . - - :param M: *2x3* transformation matrix. - - :param dsize: Size of the destination image. - - :param flags: Combination of interpolation methods (see :ocv:func:`resize`) and the optional flag ``WARP_INVERSE_MAP`` specifying that ``M`` is an inverse transformation ( ``dst=>src`` ). Only ``INTER_NEAREST`` , ``INTER_LINEAR`` , and ``INTER_CUBIC`` interpolation methods are supported. - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`warpAffine` - - - -gpu::buildWarpAffineMaps ------------------------- -Builds transformation maps for affine transformation. - -.. ocv:function:: void gpu::buildWarpAffineMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream = Stream::Null()) - - :param M: *2x3* transformation matrix. - - :param inverse: Flag specifying that ``M`` is an inverse transformation ( ``dst=>src`` ). - - :param dsize: Size of the destination image. - - :param xmap: X values with ``CV_32FC1`` type. - - :param ymap: Y values with ``CV_32FC1`` type. - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`gpu::warpAffine` , :ocv:func:`gpu::remap` - - - -gpu::warpPerspective ------------------------- -Applies a perspective transformation to an image. - -.. ocv:function:: void gpu::warpPerspective( const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags=INTER_LINEAR, int borderMode=BORDER_CONSTANT, Scalar borderValue=Scalar(), Stream& stream=Stream::Null() ) - - :param src: Source image. ``CV_8U`` , ``CV_16U`` , ``CV_32S`` , or ``CV_32F`` depth and 1, 3, or 4 channels are supported. - - :param dst: Destination image with the same type as ``src`` . The size is ``dsize`` . - - :param M: *3x3* transformation matrix. - - :param dsize: Size of the destination image. - - :param flags: Combination of interpolation methods (see :ocv:func:`resize` ) and the optional flag ``WARP_INVERSE_MAP`` specifying that ``M`` is the inverse transformation ( ``dst => src`` ). Only ``INTER_NEAREST`` , ``INTER_LINEAR`` , and ``INTER_CUBIC`` interpolation methods are supported. - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`warpPerspective` - - - -gpu::buildWarpPerspectiveMaps ------------------------------ -Builds transformation maps for perspective transformation. - -.. ocv:function:: void gpu::buildWarpAffineMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream = Stream::Null()) - - :param M: *3x3* transformation matrix. - - :param inverse: Flag specifying that ``M`` is an inverse transformation ( ``dst=>src`` ). - - :param dsize: Size of the destination image. - - :param xmap: X values with ``CV_32FC1`` type. - - :param ymap: Y values with ``CV_32FC1`` type. - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`gpu::warpPerspective` , :ocv:func:`gpu::remap` - - - -gpu::rotate ---------------- -Rotates an image around the origin (0,0) and then shifts it. - -.. ocv:function:: void gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift = 0, double yShift = 0, int interpolation = INTER_LINEAR, Stream& stream = Stream::Null()) - - :param src: Source image. Supports 1, 3 or 4 channels images with ``CV_8U`` , ``CV_16U`` or ``CV_32F`` depth. - - :param dst: Destination image with the same type as ``src`` . The size is ``dsize`` . - - :param dsize: Size of the destination image. - - :param angle: Angle of rotation in degrees. - - :param xShift: Shift along the horizontal axis. - - :param yShift: Shift along the vertical axis. - - :param interpolation: Interpolation method. Only ``INTER_NEAREST`` , ``INTER_LINEAR`` , and ``INTER_CUBIC`` are supported. - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`gpu::warpAffine` - - - -gpu::copyMakeBorder ------------------------ -Forms a border around an image. - -.. ocv:function:: void gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value = Scalar(), Stream& stream = Stream::Null()) - - :param src: Source image. ``CV_8UC1`` , ``CV_8UC4`` , ``CV_32SC1`` , and ``CV_32FC1`` types are supported. - - :param dst: Destination image with the same type as ``src``. The size is ``Size(src.cols+left+right, src.rows+top+bottom)`` . - - :param top: - - :param bottom: - - :param left: - - :param right: Number of pixels in each direction from the source image rectangle to extrapolate. For example: ``top=1, bottom=1, left=1, right=1`` mean that 1 pixel-wide border needs to be built. - - :param borderType: Border type. See :ocv:func:`borderInterpolate` for details. ``BORDER_REFLECT101`` , ``BORDER_REPLICATE`` , ``BORDER_CONSTANT`` , ``BORDER_REFLECT`` and ``BORDER_WRAP`` are supported for now. - - :param value: Border value. - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`copyMakeBorder` - - - gpu::rectStdDev ------------------- Computes a standard deviation of integral images. @@ -711,68 +356,6 @@ Equalizes the histogram of a grayscale image. -gpu::buildWarpPlaneMaps ------------------------ -Builds plane warping maps. - -.. ocv:function:: void gpu::buildWarpPlaneMaps( Size src_size, Rect dst_roi, const Mat & K, const Mat& R, const Mat & T, float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream=Stream::Null() ) - - :param stream: Stream for the asynchronous version. - - - -gpu::buildWarpCylindricalMaps ------------------------------ -Builds cylindrical warping maps. - -.. ocv:function:: void gpu::buildWarpCylindricalMaps( Size src_size, Rect dst_roi, const Mat & K, const Mat& R, float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream=Stream::Null() ) - - :param stream: Stream for the asynchronous version. - - - -gpu::buildWarpSphericalMaps ---------------------------- -Builds spherical warping maps. - -.. ocv:function:: void gpu::buildWarpSphericalMaps( Size src_size, Rect dst_roi, const Mat & K, const Mat& R, float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream=Stream::Null() ) - - :param stream: Stream for the asynchronous version. - - - -gpu::pyrDown -------------------- -Smoothes an image and downsamples it. - -.. ocv:function:: void gpu::pyrDown(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) - - :param src: Source image. - - :param dst: Destination image. Will have ``Size((src.cols+1)/2, (src.rows+1)/2)`` size and the same type as ``src`` . - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`pyrDown` - - - -gpu::pyrUp -------------------- -Upsamples an image and then smoothes it. - -.. ocv:function:: void gpu::pyrUp(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) - - :param src: Source image. - - :param dst: Destination image. Will have ``Size(src.cols*2, src.rows*2)`` size and the same type as ``src`` . - - :param stream: Stream for the asynchronous version. - -.. seealso:: :ocv:func:`pyrUp` - - - gpu::blendLinear ------------------- Performs linear blending of two images. @@ -841,6 +424,8 @@ Performs pure non local means denoising without any simplification, and thus it :ocv:func:`fastNlMeansDenoising` + + gpu::FastNonLocalMeansDenoising ------------------------------- .. ocv:class:: gpu::FastNonLocalMeansDenoising @@ -858,6 +443,8 @@ gpu::FastNonLocalMeansDenoising The class implements fast approximate Non Local Means Denoising algorithm. + + gpu::FastNonLocalMeansDenoising::simpleMethod() ----------------------------------------------- Perform image denoising using Non-local Means Denoising algorithm http://www.ipol.im/pub/algo/bcm_non_local_means_denoising with several computational optimizations. Noise expected to be a gaussian white noise @@ -882,6 +469,8 @@ This function expected to be applied to grayscale images. For colored images loo :ocv:func:`fastNlMeansDenoising` + + gpu::FastNonLocalMeansDenoising::labMethod() -------------------------------------------- Modification of ``FastNonLocalMeansDenoising::simpleMethod`` for color images @@ -908,6 +497,8 @@ The function converts image to CIELAB colorspace and then separately denoise L a :ocv:func:`fastNlMeansDenoisingColored` + + gpu::alphaComp ------------------- Composites two images using alpha opacity values contained in each image. diff --git a/modules/gpuimgproc/include/opencv2/gpuimgproc.hpp b/modules/gpuimgproc/include/opencv2/gpuimgproc.hpp index a0b1e3094f..00ef0a3e9e 100644 --- a/modules/gpuimgproc/include/opencv2/gpuimgproc.hpp +++ b/modules/gpuimgproc/include/opencv2/gpuimgproc.hpp @@ -60,12 +60,6 @@ enum { ALPHA_OVER, ALPHA_IN, ALPHA_OUT, ALPHA_ATOP, ALPHA_XOR, ALPHA_PLUS, ALPHA //! Supports CV_8UC4, CV_16UC4, CV_32SC4 and CV_32FC4 types CV_EXPORTS void alphaComp(const GpuMat& img1, const GpuMat& img2, GpuMat& dst, int alpha_op, Stream& stream = Stream::Null()); -//! DST[x,y] = SRC[xmap[x,y],ymap[x,y]] -//! supports only CV_32FC1 map type -CV_EXPORTS void remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap, - int interpolation, int borderMode = BORDER_CONSTANT, Scalar borderValue = Scalar(), - Stream& stream = Stream::Null()); - //! Does mean shift filtering on GPU. CV_EXPORTS void meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1), @@ -113,42 +107,6 @@ CV_EXPORTS void swapChannels(GpuMat& image, const int dstOrder[4], Stream& strea //! Routines for correcting image color gamma CV_EXPORTS void gammaCorrection(const GpuMat& src, GpuMat& dst, bool forward = true, Stream& stream = Stream::Null()); -//! resizes the image -//! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_AREA -CV_EXPORTS void resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, double fy=0, int interpolation = INTER_LINEAR, Stream& stream = Stream::Null()); - -//! warps the image using affine transformation -//! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC -CV_EXPORTS void warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR, - int borderMode = BORDER_CONSTANT, Scalar borderValue = Scalar(), Stream& stream = Stream::Null()); - -CV_EXPORTS void buildWarpAffineMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream = Stream::Null()); - -//! warps the image using perspective transformation -//! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC -CV_EXPORTS void warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR, - int borderMode = BORDER_CONSTANT, Scalar borderValue = Scalar(), Stream& stream = Stream::Null()); - -CV_EXPORTS void buildWarpPerspectiveMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream = Stream::Null()); - -//! builds plane warping maps -CV_EXPORTS void buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, const Mat &T, float scale, - GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); - -//! builds cylindrical warping maps -CV_EXPORTS void buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, - GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); - -//! builds spherical warping maps -CV_EXPORTS void buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, - GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); - -//! rotates an image around the origin (0,0) and then shifts it -//! supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC -//! supports 1, 3 or 4 channels images with CV_8U, CV_16U or CV_32F depth -CV_EXPORTS void rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift = 0, double yShift = 0, - int interpolation = INTER_LINEAR, Stream& stream = Stream::Null()); - //! computes Harris cornerness criteria at each image pixel CV_EXPORTS void cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, double k, int borderType = BORDER_REFLECT101); CV_EXPORTS void cornerHarris(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, double k, int borderType = BORDER_REFLECT101); @@ -176,12 +134,6 @@ CV_EXPORTS void matchTemplate(const GpuMat& image, const GpuMat& templ, GpuMat& //! computes the proximity map for the raster template and the image where the template is searched for CV_EXPORTS void matchTemplate(const GpuMat& image, const GpuMat& templ, GpuMat& result, int method, MatchTemplateBuf &buf, Stream& stream = Stream::Null()); -//! smoothes the source image and downsamples it -CV_EXPORTS void pyrDown(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()); - -//! upsamples the source image and then smoothes it -CV_EXPORTS void pyrUp(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()); - //! performs linear blending of two images //! to avoid accuracy errors sum of weigths shouldn't be very close to zero CV_EXPORTS void blendLinear(const GpuMat& img1, const GpuMat& img2, const GpuMat& weights1, const GpuMat& weights2, @@ -227,32 +179,6 @@ CV_EXPORTS void Canny(const GpuMat& image, CannyBuf& buf, GpuMat& edges, double CV_EXPORTS void Canny(const GpuMat& dx, const GpuMat& dy, GpuMat& edges, double low_thresh, double high_thresh, bool L2gradient = false); CV_EXPORTS void Canny(const GpuMat& dx, const GpuMat& dy, CannyBuf& buf, GpuMat& edges, double low_thresh, double high_thresh, bool L2gradient = false); -class CV_EXPORTS ImagePyramid -{ -public: - inline ImagePyramid() : nLayers_(0) {} - inline ImagePyramid(const GpuMat& img, int nLayers, Stream& stream = Stream::Null()) - { - build(img, nLayers, stream); - } - - void build(const GpuMat& img, int nLayers, Stream& stream = Stream::Null()); - - void getLayer(GpuMat& outImg, Size outRoi, Stream& stream = Stream::Null()) const; - - inline void release() - { - layer0_.release(); - pyramid_.clear(); - nLayers_ = 0; - } - -private: - GpuMat layer0_; - std::vector pyramid_; - int nLayers_; -}; - //! HoughLines struct HoughLinesBuf diff --git a/modules/gpuimgproc/perf/perf_imgproc.cpp b/modules/gpuimgproc/perf/perf_imgproc.cpp index 349dcc825d..ff19e14edb 100644 --- a/modules/gpuimgproc/perf/perf_imgproc.cpp +++ b/modules/gpuimgproc/perf/perf_imgproc.cpp @@ -46,323 +46,6 @@ using namespace std; using namespace testing; using namespace perf; -////////////////////////////////////////////////////////////////////// -// Remap - -enum { HALF_SIZE=0, UPSIDE_DOWN, REFLECTION_X, REFLECTION_BOTH }; -CV_ENUM(RemapMode, HALF_SIZE, UPSIDE_DOWN, REFLECTION_X, REFLECTION_BOTH); - -void generateMap(cv::Mat& map_x, cv::Mat& map_y, int remapMode) -{ - for (int j = 0; j < map_x.rows; ++j) - { - for (int i = 0; i < map_x.cols; ++i) - { - switch (remapMode) - { - case HALF_SIZE: - if (i > map_x.cols*0.25 && i < map_x.cols*0.75 && j > map_x.rows*0.25 && j < map_x.rows*0.75) - { - map_x.at(j,i) = 2.f * (i - map_x.cols * 0.25f) + 0.5f; - map_y.at(j,i) = 2.f * (j - map_x.rows * 0.25f) + 0.5f; - } - else - { - map_x.at(j,i) = 0.f; - map_y.at(j,i) = 0.f; - } - break; - case UPSIDE_DOWN: - map_x.at(j,i) = static_cast(i); - map_y.at(j,i) = static_cast(map_x.rows - j); - break; - case REFLECTION_X: - map_x.at(j,i) = static_cast(map_x.cols - i); - map_y.at(j,i) = static_cast(j); - break; - case REFLECTION_BOTH: - map_x.at(j,i) = static_cast(map_x.cols - i); - map_y.at(j,i) = static_cast(map_x.rows - j); - break; - } // end of switch - } - } -} - -DEF_PARAM_TEST(Sz_Depth_Cn_Inter_Border_Mode, cv::Size, MatDepth, MatCn, Interpolation, BorderMode, RemapMode); - -PERF_TEST_P(Sz_Depth_Cn_Inter_Border_Mode, ImgProc_Remap, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4, - Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), - ALL_BORDER_MODES, - RemapMode::all())) -{ - declare.time(20.0); - - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - const int interpolation = GET_PARAM(3); - const int borderMode = GET_PARAM(4); - const int remapMode = GET_PARAM(5); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - cv::Mat xmap(size, CV_32FC1); - cv::Mat ymap(size, CV_32FC1); - generateMap(xmap, ymap, remapMode); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - const cv::gpu::GpuMat d_xmap(xmap); - const cv::gpu::GpuMat d_ymap(ymap); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::remap(d_src, dst, d_xmap, d_ymap, interpolation, borderMode); - - GPU_SANITY_CHECK(dst); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::remap(src, dst, xmap, ymap, interpolation, borderMode); - - CPU_SANITY_CHECK(dst); - } -} - -////////////////////////////////////////////////////////////////////// -// Resize - -DEF_PARAM_TEST(Sz_Depth_Cn_Inter_Scale, cv::Size, MatDepth, MatCn, Interpolation, double); - -PERF_TEST_P(Sz_Depth_Cn_Inter_Scale, ImgProc_Resize, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4, - Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), - Values(0.5, 0.3, 2.0))) -{ - declare.time(20.0); - - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - const int interpolation = GET_PARAM(3); - const double f = GET_PARAM(4); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::resize(d_src, dst, cv::Size(), f, f, interpolation); - - GPU_SANITY_CHECK(dst, 1e-3, ERROR_RELATIVE); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::resize(src, dst, cv::Size(), f, f, interpolation); - - CPU_SANITY_CHECK(dst); - } -} - -////////////////////////////////////////////////////////////////////// -// ResizeArea - -DEF_PARAM_TEST(Sz_Depth_Cn_Scale, cv::Size, MatDepth, MatCn, double); - -PERF_TEST_P(Sz_Depth_Cn_Scale, ImgProc_ResizeArea, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4, - Values(0.2, 0.1, 0.05))) -{ - declare.time(1.0); - - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - const int interpolation = cv::INTER_AREA; - const double f = GET_PARAM(3); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::resize(d_src, dst, cv::Size(), f, f, interpolation); - - GPU_SANITY_CHECK(dst); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::resize(src, dst, cv::Size(), f, f, interpolation); - - CPU_SANITY_CHECK(dst); - } -} - -////////////////////////////////////////////////////////////////////// -// WarpAffine - -DEF_PARAM_TEST(Sz_Depth_Cn_Inter_Border, cv::Size, MatDepth, MatCn, Interpolation, BorderMode); - -PERF_TEST_P(Sz_Depth_Cn_Inter_Border, ImgProc_WarpAffine, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4, - Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), - ALL_BORDER_MODES)) -{ - declare.time(20.0); - - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - const int interpolation = GET_PARAM(3); - const int borderMode = GET_PARAM(4); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - const double aplha = CV_PI / 4; - const double mat[2 * 3] = - { - std::cos(aplha), -std::sin(aplha), src.cols / 2, - std::sin(aplha), std::cos(aplha), 0 - }; - const cv::Mat M(2, 3, CV_64F, (void*) mat); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::warpAffine(d_src, dst, M, size, interpolation, borderMode); - - GPU_SANITY_CHECK(dst, 1); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::warpAffine(src, dst, M, size, interpolation, borderMode); - - CPU_SANITY_CHECK(dst); - } -} - -////////////////////////////////////////////////////////////////////// -// WarpPerspective - -PERF_TEST_P(Sz_Depth_Cn_Inter_Border, ImgProc_WarpPerspective, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4, - Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), - ALL_BORDER_MODES)) -{ - declare.time(20.0); - - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - const int interpolation = GET_PARAM(3); - const int borderMode = GET_PARAM(4); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - const double aplha = CV_PI / 4; - double mat[3][3] = { {std::cos(aplha), -std::sin(aplha), src.cols / 2}, - {std::sin(aplha), std::cos(aplha), 0}, - {0.0, 0.0, 1.0}}; - const cv::Mat M(3, 3, CV_64F, (void*) mat); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::warpPerspective(d_src, dst, M, size, interpolation, borderMode); - - GPU_SANITY_CHECK(dst, 1); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::warpPerspective(src, dst, M, size, interpolation, borderMode); - - CPU_SANITY_CHECK(dst); - } -} - -////////////////////////////////////////////////////////////////////// -// Threshold - -CV_ENUM(ThreshOp, THRESH_BINARY, THRESH_BINARY_INV, THRESH_TRUNC, THRESH_TOZERO, THRESH_TOZERO_INV) - -DEF_PARAM_TEST(Sz_Depth_Op, cv::Size, MatDepth, ThreshOp); - -PERF_TEST_P(Sz_Depth_Op, ImgProc_Threshold, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F, CV_64F), - ThreshOp::all())) -{ - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int threshOp = GET_PARAM(2); - - cv::Mat src(size, depth); - declare.in(src, WARMUP_RNG); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::threshold(d_src, dst, 100.0, 255.0, threshOp); - - GPU_SANITY_CHECK(dst, 1e-10); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::threshold(src, dst, 100.0, 255.0, threshOp); - - CPU_SANITY_CHECK(dst); - } -} - ////////////////////////////////////////////////////////////////////// // HistEvenC1 @@ -892,196 +575,6 @@ PERF_TEST_P(Image_Type_Border_BlockSz_ApertureSz, CornerMinEigenVal, } } -////////////////////////////////////////////////////////////////////// -// BuildWarpPlaneMaps - -PERF_TEST_P(Sz, ImgProc_BuildWarpPlaneMaps, - GPU_TYPICAL_MAT_SIZES) -{ - const cv::Size size = GetParam(); - - const cv::Mat K = cv::Mat::eye(3, 3, CV_32FC1); - const cv::Mat R = cv::Mat::ones(3, 3, CV_32FC1); - const cv::Mat T = cv::Mat::zeros(1, 3, CV_32F); - - if (PERF_RUN_GPU()) - { - cv::gpu::GpuMat map_x; - cv::gpu::GpuMat map_y; - - TEST_CYCLE() cv::gpu::buildWarpPlaneMaps(size, cv::Rect(0, 0, size.width, size.height), K, R, T, 1.0, map_x, map_y); - - GPU_SANITY_CHECK(map_x); - GPU_SANITY_CHECK(map_y); - } - else - { - FAIL_NO_CPU(); - } -} - -////////////////////////////////////////////////////////////////////// -// BuildWarpCylindricalMaps - -PERF_TEST_P(Sz, ImgProc_BuildWarpCylindricalMaps, - GPU_TYPICAL_MAT_SIZES) -{ - const cv::Size size = GetParam(); - - const cv::Mat K = cv::Mat::eye(3, 3, CV_32FC1); - const cv::Mat R = cv::Mat::ones(3, 3, CV_32FC1); - - if (PERF_RUN_GPU()) - { - cv::gpu::GpuMat map_x; - cv::gpu::GpuMat map_y; - - TEST_CYCLE() cv::gpu::buildWarpCylindricalMaps(size, cv::Rect(0, 0, size.width, size.height), K, R, 1.0, map_x, map_y); - - GPU_SANITY_CHECK(map_x); - GPU_SANITY_CHECK(map_y); - } - else - { - FAIL_NO_CPU(); - } -} - -////////////////////////////////////////////////////////////////////// -// BuildWarpSphericalMaps - -PERF_TEST_P(Sz, ImgProc_BuildWarpSphericalMaps, - GPU_TYPICAL_MAT_SIZES) -{ - const cv::Size size = GetParam(); - - const cv::Mat K = cv::Mat::eye(3, 3, CV_32FC1); - const cv::Mat R = cv::Mat::ones(3, 3, CV_32FC1); - - if (PERF_RUN_GPU()) - { - cv::gpu::GpuMat map_x; - cv::gpu::GpuMat map_y; - - TEST_CYCLE() cv::gpu::buildWarpSphericalMaps(size, cv::Rect(0, 0, size.width, size.height), K, R, 1.0, map_x, map_y); - - GPU_SANITY_CHECK(map_x); - GPU_SANITY_CHECK(map_y); - } - else - { - FAIL_NO_CPU(); - } -} - -////////////////////////////////////////////////////////////////////// -// Rotate - -DEF_PARAM_TEST(Sz_Depth_Cn_Inter, cv::Size, MatDepth, MatCn, Interpolation); - -PERF_TEST_P(Sz_Depth_Cn_Inter, ImgProc_Rotate, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4, - Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)))) -{ - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - const int interpolation = GET_PARAM(3); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::rotate(d_src, dst, size, 30.0, 0, 0, interpolation); - - GPU_SANITY_CHECK(dst, 1e-3, ERROR_RELATIVE); - } - else - { - FAIL_NO_CPU(); - } -} - -////////////////////////////////////////////////////////////////////// -// PyrDown - -PERF_TEST_P(Sz_Depth_Cn, ImgProc_PyrDown, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4)) -{ - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::pyrDown(d_src, dst); - - GPU_SANITY_CHECK(dst); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::pyrDown(src, dst); - - CPU_SANITY_CHECK(dst); - } -} - -////////////////////////////////////////////////////////////////////// -// PyrUp - -PERF_TEST_P(Sz_Depth_Cn, ImgProc_PyrUp, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4)) -{ - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::pyrUp(d_src, dst); - - GPU_SANITY_CHECK(dst); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::pyrUp(src, dst); - - CPU_SANITY_CHECK(dst); - } -} - ////////////////////////////////////////////////////////////////////// // CvtColor @@ -1284,82 +777,6 @@ PERF_TEST_P(Sz_Type_Op, AlphaComp, } } -////////////////////////////////////////////////////////////////////// -// ImagePyramidBuild - -PERF_TEST_P(Sz_Depth_Cn, ImgProc_ImagePyramidBuild, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4)) -{ - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - const int nLayers = 5; - const cv::Size dstSize(size.width / 2 + 10, size.height / 2 + 10); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - - cv::gpu::ImagePyramid d_pyr; - - TEST_CYCLE() d_pyr.build(d_src, nLayers); - - cv::gpu::GpuMat dst; - d_pyr.getLayer(dst, dstSize); - - GPU_SANITY_CHECK(dst); - } - else - { - FAIL_NO_CPU(); - } -} - -////////////////////////////////////////////////////////////////////// -// ImagePyramidGetLayer - -PERF_TEST_P(Sz_Depth_Cn, ImgProc_ImagePyramidGetLayer, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4)) -{ - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - const int nLayers = 3; - const cv::Size dstSize(size.width / 2 + 10, size.height / 2 + 10); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - cv::gpu::ImagePyramid d_pyr(d_src, nLayers); - - TEST_CYCLE() d_pyr.getLayer(dst, dstSize); - - GPU_SANITY_CHECK(dst); - } - else - { - FAIL_NO_CPU(); - } -} - ////////////////////////////////////////////////////////////////////// // HoughLines diff --git a/modules/gpuimgproc/src/cuda/imgproc.cu b/modules/gpuimgproc/src/cuda/imgproc.cu index d2d0d0f3c5..c47076f44d 100644 --- a/modules/gpuimgproc/src/cuda/imgproc.cu +++ b/modules/gpuimgproc/src/cuda/imgproc.cu @@ -399,172 +399,6 @@ namespace cv { namespace gpu { namespace cudev if (stream == 0) cudaSafeCall(cudaDeviceSynchronize()); } - - ////////////////////////////////////////////////////////////////////////// - // buildWarpMaps - - // TODO use intrinsics like __sinf and so on - - namespace build_warp_maps - { - - __constant__ float ck_rinv[9]; - __constant__ float cr_kinv[9]; - __constant__ float ct[3]; - __constant__ float cscale; - } - - - class PlaneMapper - { - public: - static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y) - { - using namespace build_warp_maps; - - float x_ = u / cscale - ct[0]; - float y_ = v / cscale - ct[1]; - - float z; - x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * (1 - ct[2]); - y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * (1 - ct[2]); - z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * (1 - ct[2]); - - x /= z; - y /= z; - } - }; - - - class CylindricalMapper - { - public: - static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y) - { - using namespace build_warp_maps; - - u /= cscale; - float x_ = ::sinf(u); - float y_ = v / cscale; - float z_ = ::cosf(u); - - float z; - x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; - y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; - z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; - - if (z > 0) { x /= z; y /= z; } - else x = y = -1; - } - }; - - - class SphericalMapper - { - public: - static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y) - { - using namespace build_warp_maps; - - v /= cscale; - u /= cscale; - - float sinv = ::sinf(v); - float x_ = sinv * ::sinf(u); - float y_ = -::cosf(v); - float z_ = sinv * ::cosf(u); - - float z; - x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; - y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; - z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; - - if (z > 0) { x /= z; y /= z; } - else x = y = -1; - } - }; - - - template - __global__ void buildWarpMapsKernel(int tl_u, int tl_v, int cols, int rows, - PtrStepf map_x, PtrStepf map_y) - { - int du = blockIdx.x * blockDim.x + threadIdx.x; - int dv = blockIdx.y * blockDim.y + threadIdx.y; - if (du < cols && dv < rows) - { - float u = tl_u + du; - float v = tl_v + dv; - float x, y; - Mapper::mapBackward(u, v, x, y); - map_x.ptr(dv)[du] = x; - map_y.ptr(dv)[du] = y; - } - } - - - void buildWarpPlaneMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, - const float k_rinv[9], const float r_kinv[9], const float t[3], - float scale, cudaStream_t stream) - { - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ct, t, 3*sizeof(float))); - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); - - int cols = map_x.cols; - int rows = map_x.rows; - - dim3 threads(32, 8); - dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); - - buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); - cudaSafeCall(cudaGetLastError()); - if (stream == 0) - cudaSafeCall(cudaDeviceSynchronize()); - } - - - void buildWarpCylindricalMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, - const float k_rinv[9], const float r_kinv[9], float scale, - cudaStream_t stream) - { - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); - - int cols = map_x.cols; - int rows = map_x.rows; - - dim3 threads(32, 8); - dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); - - buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); - cudaSafeCall(cudaGetLastError()); - if (stream == 0) - cudaSafeCall(cudaDeviceSynchronize()); - } - - - void buildWarpSphericalMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, - const float k_rinv[9], const float r_kinv[9], float scale, - cudaStream_t stream) - { - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); - cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); - - int cols = map_x.cols; - int rows = map_x.rows; - - dim3 threads(32, 8); - dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); - - buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); - cudaSafeCall(cudaGetLastError()); - if (stream == 0) - cudaSafeCall(cudaDeviceSynchronize()); - } } // namespace imgproc }}} // namespace cv { namespace gpu { namespace cudev { diff --git a/modules/gpuimgproc/src/imgproc.cpp b/modules/gpuimgproc/src/imgproc.cpp index dabf054b60..939b149370 100644 --- a/modules/gpuimgproc/src/imgproc.cpp +++ b/modules/gpuimgproc/src/imgproc.cpp @@ -49,10 +49,6 @@ using namespace cv::gpu; void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria, Stream&) { throw_no_cuda(); } void cv::gpu::meanShiftProc(const GpuMat&, GpuMat&, GpuMat&, int, int, TermCriteria, Stream&) { throw_no_cuda(); } -void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } -void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } -void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } -void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int, Stream&) { throw_no_cuda(); } void cv::gpu::evenLevels(GpuMat&, int, int, int) { throw_no_cuda(); } void cv::gpu::histEven(const GpuMat&, GpuMat&, int, int, int, Stream&) { throw_no_cuda(); } void cv::gpu::histEven(const GpuMat&, GpuMat&, GpuMat&, int, int, int, Stream&) { throw_no_cuda(); } @@ -155,184 +151,6 @@ void cv::gpu::meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int meanShiftProc_gpu(src, dstr, dstsp, sp, sr, maxIter, eps, StreamAccessor::getStream(stream)); } -////////////////////////////////////////////////////////////////////////////// -// buildWarpPlaneMaps - -namespace cv { namespace gpu { namespace cudev -{ - namespace imgproc - { - void buildWarpPlaneMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, - const float k_rinv[9], const float r_kinv[9], const float t[3], float scale, - cudaStream_t stream); - } -}}} - -void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, const Mat &T, - float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream) -{ - (void)src_size; - using namespace ::cv::gpu::cudev::imgproc; - - CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); - CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); - CV_Assert((T.size() == Size(3,1) || T.size() == Size(1,3)) && T.type() == CV_32F && T.isContinuous()); - - Mat K_Rinv = K * R.t(); - Mat R_Kinv = R * K.inv(); - CV_Assert(K_Rinv.isContinuous()); - CV_Assert(R_Kinv.isContinuous()); - - map_x.create(dst_roi.size(), CV_32F); - map_y.create(dst_roi.size(), CV_32F); - cudev::imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), - T.ptr(), scale, StreamAccessor::getStream(stream)); -} - -////////////////////////////////////////////////////////////////////////////// -// buildWarpCylyndricalMaps - -namespace cv { namespace gpu { namespace cudev -{ - namespace imgproc - { - void buildWarpCylindricalMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, - const float k_rinv[9], const float r_kinv[9], float scale, - cudaStream_t stream); - } -}}} - -void cv::gpu::buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, - GpuMat& map_x, GpuMat& map_y, Stream& stream) -{ - (void)src_size; - using namespace ::cv::gpu::cudev::imgproc; - - CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); - CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); - - Mat K_Rinv = K * R.t(); - Mat R_Kinv = R * K.inv(); - CV_Assert(K_Rinv.isContinuous()); - CV_Assert(R_Kinv.isContinuous()); - - map_x.create(dst_roi.size(), CV_32F); - map_y.create(dst_roi.size(), CV_32F); - cudev::imgproc::buildWarpCylindricalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), scale, StreamAccessor::getStream(stream)); -} - - -////////////////////////////////////////////////////////////////////////////// -// buildWarpSphericalMaps - -namespace cv { namespace gpu { namespace cudev -{ - namespace imgproc - { - void buildWarpSphericalMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, - const float k_rinv[9], const float r_kinv[9], float scale, - cudaStream_t stream); - } -}}} - -void cv::gpu::buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, - GpuMat& map_x, GpuMat& map_y, Stream& stream) -{ - (void)src_size; - using namespace ::cv::gpu::cudev::imgproc; - - CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); - CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); - - Mat K_Rinv = K * R.t(); - Mat R_Kinv = R * K.inv(); - CV_Assert(K_Rinv.isContinuous()); - CV_Assert(R_Kinv.isContinuous()); - - map_x.create(dst_roi.size(), CV_32F); - map_y.create(dst_roi.size(), CV_32F); - cudev::imgproc::buildWarpSphericalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), scale, StreamAccessor::getStream(stream)); -} - -//////////////////////////////////////////////////////////////////////// -// rotate - -namespace -{ - template struct NppTypeTraits; - template<> struct NppTypeTraits { typedef Npp8u npp_t; }; - template<> struct NppTypeTraits { typedef Npp8s npp_t; }; - template<> struct NppTypeTraits { typedef Npp16u npp_t; }; - template<> struct NppTypeTraits { typedef Npp16s npp_t; }; - template<> struct NppTypeTraits { typedef Npp32s npp_t; }; - template<> struct NppTypeTraits { typedef Npp32f npp_t; }; - template<> struct NppTypeTraits { typedef Npp64f npp_t; }; - - template struct NppRotateFunc - { - typedef typename NppTypeTraits::npp_t npp_t; - - typedef NppStatus (*func_t)(const npp_t* pSrc, NppiSize oSrcSize, int nSrcStep, NppiRect oSrcROI, - npp_t* pDst, int nDstStep, NppiRect oDstROI, - double nAngle, double nShiftX, double nShiftY, int eInterpolation); - }; - - template ::func_t func> struct NppRotate - { - typedef typename NppRotateFunc::npp_t npp_t; - - static void call(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, cudaStream_t stream) - { - (void)dsize; - static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC}; - - NppStreamHandler h(stream); - - NppiSize srcsz; - srcsz.height = src.rows; - srcsz.width = src.cols; - NppiRect srcroi; - srcroi.x = srcroi.y = 0; - srcroi.height = src.rows; - srcroi.width = src.cols; - NppiRect dstroi; - dstroi.x = dstroi.y = 0; - dstroi.height = dst.rows; - dstroi.width = dst.cols; - - nppSafeCall( func(src.ptr(), srcsz, static_cast(src.step), srcroi, - dst.ptr(), static_cast(dst.step), dstroi, angle, xShift, yShift, npp_inter[interpolation]) ); - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } - }; -} - -void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, Stream& stream) -{ - typedef void (*func_t)(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, cudaStream_t stream); - - static const func_t funcs[6][4] = - { - {NppRotate::call, 0, NppRotate::call, NppRotate::call}, - {0,0,0,0}, - {NppRotate::call, 0, NppRotate::call, NppRotate::call}, - {0,0,0,0}, - {0,0,0,0}, - {NppRotate::call, 0, NppRotate::call, NppRotate::call} - }; - - CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); - CV_Assert(src.channels() == 1 || src.channels() == 3 || src.channels() == 4); - CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC); - - dst.create(dsize, src.type()); - dst.setTo(Scalar::all(0)); - - funcs[src.depth()][src.channels() - 1](src, dst, dsize, angle, xShift, yShift, interpolation, StreamAccessor::getStream(stream)); -} - //////////////////////////////////////////////////////////////////////// // Histogram @@ -344,14 +162,14 @@ namespace template struct NppHistogramEvenFuncC1 { - typedef typename NppTypeTraits::npp_t src_t; + typedef typename NPPTypeTraits::npp_type src_t; typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s * pHist, int nLevels, Npp32s nLowerLevel, Npp32s nUpperLevel, Npp8u * pBuffer); }; template struct NppHistogramEvenFuncC4 { - typedef typename NppTypeTraits::npp_t src_t; + typedef typename NPPTypeTraits::npp_type src_t; typedef NppStatus (*func_ptr)(const src_t* pSrc, int nSrcStep, NppiSize oSizeROI, Npp32s * pHist[4], int nLevels[4], Npp32s nLowerLevel[4], Npp32s nUpperLevel[4], Npp8u * pBuffer); @@ -420,7 +238,7 @@ namespace template struct NppHistogramRangeFuncC1 { - typedef typename NppTypeTraits::npp_t src_t; + typedef typename NPPTypeTraits::npp_type src_t; typedef Npp32s level_t; enum {LEVEL_TYPE_CODE=CV_32SC1}; @@ -438,7 +256,7 @@ namespace }; template struct NppHistogramRangeFuncC4 { - typedef typename NppTypeTraits::npp_t src_t; + typedef typename NPPTypeTraits::npp_type src_t; typedef Npp32s level_t; enum {LEVEL_TYPE_CODE=CV_32SC1}; @@ -1042,14 +860,14 @@ namespace { template struct NppAlphaCompFunc { - typedef typename NppTypeTraits::npp_t npp_t; + typedef typename NPPTypeTraits::npp_type npp_t; typedef NppStatus (*func_t)(const npp_t* pSrc1, int nSrc1Step, const npp_t* pSrc2, int nSrc2Step, npp_t* pDst, int nDstStep, NppiSize oSizeROI, NppiAlphaOp eAlphaOp); }; template ::func_t func> struct NppAlphaComp { - typedef typename NppTypeTraits::npp_t npp_t; + typedef typename NPPTypeTraits::npp_type npp_t; static void call(const GpuMat& img1, const GpuMat& img2, GpuMat& dst, NppiAlphaOp eAlphaOp, cudaStream_t stream) { diff --git a/modules/gpuimgproc/src/precomp.hpp b/modules/gpuimgproc/src/precomp.hpp index 7df02aadd9..93bcf3a28b 100644 --- a/modules/gpuimgproc/src/precomp.hpp +++ b/modules/gpuimgproc/src/precomp.hpp @@ -43,9 +43,9 @@ #ifndef __OPENCV_PRECOMP_H__ #define __OPENCV_PRECOMP_H__ +#include "opencv2/gpuimgproc.hpp" #include "opencv2/gpufilters.hpp" #include "opencv2/gpuarithm.hpp" -#include "opencv2/gpuimgproc.hpp" #include "opencv2/core/private.hpp" #include "opencv2/core/gpu_private.hpp" diff --git a/modules/gpuimgproc/test/test_denoising.cpp b/modules/gpuimgproc/test/test_denoising.cpp index 2f1a93be1c..cb4ea1ecce 100644 --- a/modules/gpuimgproc/test/test_denoising.cpp +++ b/modules/gpuimgproc/test/test_denoising.cpp @@ -46,53 +46,6 @@ using namespace cvtest; -//////////////////////////////////////////////////////// -// BilateralFilter - -PARAM_TEST_CASE(BilateralFilter, cv::gpu::DeviceInfo, cv::Size, MatType) -{ - cv::gpu::DeviceInfo devInfo; - cv::Size size; - int type; - int kernel_size; - float sigma_color; - float sigma_spatial; - - virtual void SetUp() - { - devInfo = GET_PARAM(0); - size = GET_PARAM(1); - type = GET_PARAM(2); - - kernel_size = 5; - sigma_color = 10.f; - sigma_spatial = 3.5f; - - cv::gpu::setDevice(devInfo.deviceID()); - } -}; - -GPU_TEST_P(BilateralFilter, Accuracy) -{ - cv::Mat src = randomMat(size, type); - - src.convertTo(src, type); - cv::gpu::GpuMat dst; - - cv::gpu::bilateralFilter(loadMat(src), dst, kernel_size, sigma_color, sigma_spatial); - - cv::Mat dst_gold; - cv::bilateralFilter(src, dst_gold, kernel_size, sigma_color, sigma_spatial); - - EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-3 : 1.0); -} - -INSTANTIATE_TEST_CASE_P(GPU_Denoising, BilateralFilter, testing::Combine( - ALL_DEVICES, - testing::Values(cv::Size(128, 128), cv::Size(113, 113), cv::Size(639, 481)), - testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_32FC1), MatType(CV_32FC3)) - )); - //////////////////////////////////////////////////////// // Brute Force Non local means diff --git a/modules/gpuimgproc/test/test_imgproc.cpp b/modules/gpuimgproc/test/test_imgproc.cpp index 6957f54375..0e66fe03dd 100644 --- a/modules/gpuimgproc/test/test_imgproc.cpp +++ b/modules/gpuimgproc/test/test_imgproc.cpp @@ -840,4 +840,51 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CornerMinEigen, testing::Combine( testing::Values(BlockSize(3), BlockSize(5), BlockSize(7)), testing::Values(ApertureSize(0), ApertureSize(3), ApertureSize(5), ApertureSize(7)))); +//////////////////////////////////////////////////////// +// BilateralFilter + +PARAM_TEST_CASE(BilateralFilter, cv::gpu::DeviceInfo, cv::Size, MatType) +{ + cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; + int kernel_size; + float sigma_color; + float sigma_spatial; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + type = GET_PARAM(2); + + kernel_size = 5; + sigma_color = 10.f; + sigma_spatial = 3.5f; + + cv::gpu::setDevice(devInfo.deviceID()); + } +}; + +GPU_TEST_P(BilateralFilter, Accuracy) +{ + cv::Mat src = randomMat(size, type); + + src.convertTo(src, type); + cv::gpu::GpuMat dst; + + cv::gpu::bilateralFilter(loadMat(src), dst, kernel_size, sigma_color, sigma_spatial); + + cv::Mat dst_gold; + cv::bilateralFilter(src, dst_gold, kernel_size, sigma_color, sigma_spatial); + + EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-3 : 1.0); +} + +INSTANTIATE_TEST_CASE_P(GPU_Denoising, BilateralFilter, testing::Combine( + ALL_DEVICES, + testing::Values(cv::Size(128, 128), cv::Size(113, 113), cv::Size(639, 481)), + testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_32FC1), MatType(CV_32FC3)) + )); + #endif // HAVE_CUDA diff --git a/modules/gpuimgproc/test/test_precomp.hpp b/modules/gpuimgproc/test/test_precomp.hpp index a80f5e5f44..4196aa9fe1 100644 --- a/modules/gpuimgproc/test/test_precomp.hpp +++ b/modules/gpuimgproc/test/test_precomp.hpp @@ -58,6 +58,4 @@ #include "opencv2/gpuarithm.hpp" #include "opencv2/imgproc.hpp" -#include "interpolation.hpp" - #endif diff --git a/modules/gpuobjdetect/CMakeLists.txt b/modules/gpuobjdetect/CMakeLists.txt index 5bce4d283a..bcc2242c5d 100644 --- a/modules/gpuobjdetect/CMakeLists.txt +++ b/modules/gpuobjdetect/CMakeLists.txt @@ -6,4 +6,4 @@ set(the_description "GPU-accelerated Object Detection") ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations) -ocv_define_module(gpuobjdetect opencv_objdetect opencv_gpuimgproc OPTIONAL opencv_gpulegacy) +ocv_define_module(gpuobjdetect opencv_objdetect opencv_gpuwarping opencv_gpuimgproc OPTIONAL opencv_gpulegacy) diff --git a/modules/gpuobjdetect/src/precomp.hpp b/modules/gpuobjdetect/src/precomp.hpp index 003df2eca0..40fd6c46d7 100644 --- a/modules/gpuobjdetect/src/precomp.hpp +++ b/modules/gpuobjdetect/src/precomp.hpp @@ -44,6 +44,7 @@ #define __OPENCV_PRECOMP_H__ #include "opencv2/gpuobjdetect.hpp" +#include "opencv2/gpuwarping.hpp" #include "opencv2/gpuimgproc.hpp" #include "opencv2/gpuarithm.hpp" diff --git a/modules/gpuvideo/CMakeLists.txt b/modules/gpuvideo/CMakeLists.txt index 4e6264aed5..6c15bd147a 100644 --- a/modules/gpuvideo/CMakeLists.txt +++ b/modules/gpuvideo/CMakeLists.txt @@ -6,4 +6,4 @@ set(the_description "GPU-accelerated Video Analysis") ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations) -ocv_define_module(gpuvideo opencv_video opencv_legacy opencv_gpufilters opencv_gpuimgproc OPTIONAL opencv_gpulegacy) +ocv_define_module(gpuvideo opencv_video opencv_legacy opencv_gpufilters opencv_gpuwarping opencv_gpuimgproc OPTIONAL opencv_gpulegacy) diff --git a/modules/gpuvideo/src/precomp.hpp b/modules/gpuvideo/src/precomp.hpp index 1e37cf77de..a1e1b47d54 100644 --- a/modules/gpuvideo/src/precomp.hpp +++ b/modules/gpuvideo/src/precomp.hpp @@ -49,6 +49,7 @@ #include "opencv2/gpuarithm.hpp" #include "opencv2/gpufilters.hpp" +#include "opencv2/gpuwarping.hpp" #include "opencv2/gpuimgproc.hpp" #include "opencv2/video.hpp" diff --git a/modules/gpuwarping/CMakeLists.txt b/modules/gpuwarping/CMakeLists.txt new file mode 100644 index 0000000000..0c4ca7b12e --- /dev/null +++ b/modules/gpuwarping/CMakeLists.txt @@ -0,0 +1,9 @@ +if(ANDROID OR IOS) + ocv_module_disable(gpuwarping) +endif() + +set(the_description "GPU-accelerated Image Warping") + +ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4324 /wd4512 -Wundef -Wmissing-declarations) + +ocv_define_module(gpuwarping opencv_imgproc OPTIONAL opencv_gpulegacy) diff --git a/modules/gpuwarping/doc/gpuwarping.rst b/modules/gpuwarping/doc/gpuwarping.rst new file mode 100644 index 0000000000..4bdaa8d4eb --- /dev/null +++ b/modules/gpuwarping/doc/gpuwarping.rst @@ -0,0 +1,8 @@ +***************************************** +gpuwarping. GPU-accelerated Image Warping +***************************************** + +.. toctree:: + :maxdepth: 1 + + warping diff --git a/modules/gpuwarping/doc/warping.rst b/modules/gpuwarping/doc/warping.rst new file mode 100644 index 0000000000..b2c95e2b02 --- /dev/null +++ b/modules/gpuwarping/doc/warping.rst @@ -0,0 +1,251 @@ +Image Warping +============= + +.. highlight:: cpp + + + +gpu::remap +-------------- +Applies a generic geometrical transformation to an image. + +.. ocv:function:: void gpu::remap( const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap, int interpolation, int borderMode=BORDER_CONSTANT, Scalar borderValue=Scalar(), Stream& stream=Stream::Null() ) + + :param src: Source image. + + :param dst: Destination image with the size the same as ``xmap`` and the type the same as ``src`` . + + :param xmap: X values. Only ``CV_32FC1`` type is supported. + + :param ymap: Y values. Only ``CV_32FC1`` type is supported. + + :param interpolation: Interpolation method (see :ocv:func:`resize` ). ``INTER_NEAREST`` , ``INTER_LINEAR`` and ``INTER_CUBIC`` are supported for now. + + :param borderMode: Pixel extrapolation method (see :ocv:func:`borderInterpolate` ). ``BORDER_REFLECT101`` , ``BORDER_REPLICATE`` , ``BORDER_CONSTANT`` , ``BORDER_REFLECT`` and ``BORDER_WRAP`` are supported for now. + + :param borderValue: Value used in case of a constant border. By default, it is 0. + + :param stream: Stream for the asynchronous version. + +The function transforms the source image using the specified map: + +.. math:: + + \texttt{dst} (x,y) = \texttt{src} (xmap(x,y), ymap(x,y)) + +Values of pixels with non-integer coordinates are computed using the bilinear interpolation. + +.. seealso:: :ocv:func:`remap` + + + +gpu::resize +--------------- +Resizes an image. + +.. ocv:function:: void gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, double fy=0, int interpolation = INTER_LINEAR, Stream& stream = Stream::Null()) + + :param src: Source image. + + :param dst: Destination image with the same type as ``src`` . The size is ``dsize`` (when it is non-zero) or the size is computed from ``src.size()`` , ``fx`` , and ``fy`` . + + :param dsize: Destination image size. If it is zero, it is computed as: + + .. math:: + \texttt{dsize = Size(round(fx*src.cols), round(fy*src.rows))} + + Either ``dsize`` or both ``fx`` and ``fy`` must be non-zero. + + :param fx: Scale factor along the horizontal axis. If it is zero, it is computed as: + + .. math:: + + \texttt{(double)dsize.width/src.cols} + + :param fy: Scale factor along the vertical axis. If it is zero, it is computed as: + + .. math:: + + \texttt{(double)dsize.height/src.rows} + + :param interpolation: Interpolation method. ``INTER_NEAREST`` , ``INTER_LINEAR`` and ``INTER_CUBIC`` are supported for now. + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`resize` + + + +gpu::warpAffine +------------------- +Applies an affine transformation to an image. + +.. ocv:function:: void gpu::warpAffine( const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags=INTER_LINEAR, int borderMode=BORDER_CONSTANT, Scalar borderValue=Scalar(), Stream& stream=Stream::Null() ) + + :param src: Source image. ``CV_8U`` , ``CV_16U`` , ``CV_32S`` , or ``CV_32F`` depth and 1, 3, or 4 channels are supported. + + :param dst: Destination image with the same type as ``src`` . The size is ``dsize`` . + + :param M: *2x3* transformation matrix. + + :param dsize: Size of the destination image. + + :param flags: Combination of interpolation methods (see :ocv:func:`resize`) and the optional flag ``WARP_INVERSE_MAP`` specifying that ``M`` is an inverse transformation ( ``dst=>src`` ). Only ``INTER_NEAREST`` , ``INTER_LINEAR`` , and ``INTER_CUBIC`` interpolation methods are supported. + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`warpAffine` + + + +gpu::buildWarpAffineMaps +------------------------ +Builds transformation maps for affine transformation. + +.. ocv:function:: void gpu::buildWarpAffineMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream = Stream::Null()) + + :param M: *2x3* transformation matrix. + + :param inverse: Flag specifying that ``M`` is an inverse transformation ( ``dst=>src`` ). + + :param dsize: Size of the destination image. + + :param xmap: X values with ``CV_32FC1`` type. + + :param ymap: Y values with ``CV_32FC1`` type. + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`gpu::warpAffine` , :ocv:func:`gpu::remap` + + + +gpu::warpPerspective +------------------------ +Applies a perspective transformation to an image. + +.. ocv:function:: void gpu::warpPerspective( const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags=INTER_LINEAR, int borderMode=BORDER_CONSTANT, Scalar borderValue=Scalar(), Stream& stream=Stream::Null() ) + + :param src: Source image. ``CV_8U`` , ``CV_16U`` , ``CV_32S`` , or ``CV_32F`` depth and 1, 3, or 4 channels are supported. + + :param dst: Destination image with the same type as ``src`` . The size is ``dsize`` . + + :param M: *3x3* transformation matrix. + + :param dsize: Size of the destination image. + + :param flags: Combination of interpolation methods (see :ocv:func:`resize` ) and the optional flag ``WARP_INVERSE_MAP`` specifying that ``M`` is the inverse transformation ( ``dst => src`` ). Only ``INTER_NEAREST`` , ``INTER_LINEAR`` , and ``INTER_CUBIC`` interpolation methods are supported. + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`warpPerspective` + + + +gpu::buildWarpPerspectiveMaps +----------------------------- +Builds transformation maps for perspective transformation. + +.. ocv:function:: void gpu::buildWarpAffineMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream = Stream::Null()) + + :param M: *3x3* transformation matrix. + + :param inverse: Flag specifying that ``M`` is an inverse transformation ( ``dst=>src`` ). + + :param dsize: Size of the destination image. + + :param xmap: X values with ``CV_32FC1`` type. + + :param ymap: Y values with ``CV_32FC1`` type. + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`gpu::warpPerspective` , :ocv:func:`gpu::remap` + + + +gpu::rotate +--------------- +Rotates an image around the origin (0,0) and then shifts it. + +.. ocv:function:: void gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift = 0, double yShift = 0, int interpolation = INTER_LINEAR, Stream& stream = Stream::Null()) + + :param src: Source image. Supports 1, 3 or 4 channels images with ``CV_8U`` , ``CV_16U`` or ``CV_32F`` depth. + + :param dst: Destination image with the same type as ``src`` . The size is ``dsize`` . + + :param dsize: Size of the destination image. + + :param angle: Angle of rotation in degrees. + + :param xShift: Shift along the horizontal axis. + + :param yShift: Shift along the vertical axis. + + :param interpolation: Interpolation method. Only ``INTER_NEAREST`` , ``INTER_LINEAR`` , and ``INTER_CUBIC`` are supported. + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`gpu::warpAffine` + + + +gpu::buildWarpPlaneMaps +----------------------- +Builds plane warping maps. + +.. ocv:function:: void gpu::buildWarpPlaneMaps( Size src_size, Rect dst_roi, const Mat & K, const Mat& R, const Mat & T, float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream=Stream::Null() ) + + :param stream: Stream for the asynchronous version. + + + +gpu::buildWarpCylindricalMaps +----------------------------- +Builds cylindrical warping maps. + +.. ocv:function:: void gpu::buildWarpCylindricalMaps( Size src_size, Rect dst_roi, const Mat & K, const Mat& R, float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream=Stream::Null() ) + + :param stream: Stream for the asynchronous version. + + + +gpu::buildWarpSphericalMaps +--------------------------- +Builds spherical warping maps. + +.. ocv:function:: void gpu::buildWarpSphericalMaps( Size src_size, Rect dst_roi, const Mat & K, const Mat& R, float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream=Stream::Null() ) + + :param stream: Stream for the asynchronous version. + + + +gpu::pyrDown +------------------- +Smoothes an image and downsamples it. + +.. ocv:function:: void gpu::pyrDown(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) + + :param src: Source image. + + :param dst: Destination image. Will have ``Size((src.cols+1)/2, (src.rows+1)/2)`` size and the same type as ``src`` . + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`pyrDown` + + + +gpu::pyrUp +------------------- +Upsamples an image and then smoothes it. + +.. ocv:function:: void gpu::pyrUp(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()) + + :param src: Source image. + + :param dst: Destination image. Will have ``Size(src.cols*2, src.rows*2)`` size and the same type as ``src`` . + + :param stream: Stream for the asynchronous version. + +.. seealso:: :ocv:func:`pyrUp` diff --git a/modules/gpuwarping/include/opencv2/gpuwarping.hpp b/modules/gpuwarping/include/opencv2/gpuwarping.hpp new file mode 100644 index 0000000000..ada180d9f9 --- /dev/null +++ b/modules/gpuwarping/include/opencv2/gpuwarping.hpp @@ -0,0 +1,131 @@ +/*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_GPUWARPING_HPP__ +#define __OPENCV_GPUWARPING_HPP__ + +#ifndef __cplusplus +# error gpuwarping.hpp header must be compiled as C++ +#endif + +#include "opencv2/core/gpumat.hpp" +#include "opencv2/imgproc.hpp" + +namespace cv { namespace gpu { + +//! DST[x,y] = SRC[xmap[x,y],ymap[x,y]] +//! supports only CV_32FC1 map type +CV_EXPORTS void remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap, + int interpolation, int borderMode = BORDER_CONSTANT, Scalar borderValue = Scalar(), + Stream& stream = Stream::Null()); + +//! resizes the image +//! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_AREA +CV_EXPORTS void resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, double fy=0, int interpolation = INTER_LINEAR, Stream& stream = Stream::Null()); + +//! warps the image using affine transformation +//! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC +CV_EXPORTS void warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR, + int borderMode = BORDER_CONSTANT, Scalar borderValue = Scalar(), Stream& stream = Stream::Null()); + +CV_EXPORTS void buildWarpAffineMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream = Stream::Null()); + +//! warps the image using perspective transformation +//! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC +CV_EXPORTS void warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR, + int borderMode = BORDER_CONSTANT, Scalar borderValue = Scalar(), Stream& stream = Stream::Null()); + +CV_EXPORTS void buildWarpPerspectiveMaps(const Mat& M, bool inverse, Size dsize, GpuMat& xmap, GpuMat& ymap, Stream& stream = Stream::Null()); + +//! builds plane warping maps +CV_EXPORTS void buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, const Mat &T, float scale, + GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); + +//! builds cylindrical warping maps +CV_EXPORTS void buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, + GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); + +//! builds spherical warping maps +CV_EXPORTS void buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, + GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); + +//! rotates an image around the origin (0,0) and then shifts it +//! supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC +//! supports 1, 3 or 4 channels images with CV_8U, CV_16U or CV_32F depth +CV_EXPORTS void rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift = 0, double yShift = 0, + int interpolation = INTER_LINEAR, Stream& stream = Stream::Null()); + +//! smoothes the source image and downsamples it +CV_EXPORTS void pyrDown(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()); + +//! upsamples the source image and then smoothes it +CV_EXPORTS void pyrUp(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null()); + +class CV_EXPORTS ImagePyramid +{ +public: + inline ImagePyramid() : nLayers_(0) {} + inline ImagePyramid(const GpuMat& img, int nLayers, Stream& stream = Stream::Null()) + { + build(img, nLayers, stream); + } + + void build(const GpuMat& img, int nLayers, Stream& stream = Stream::Null()); + + void getLayer(GpuMat& outImg, Size outRoi, Stream& stream = Stream::Null()) const; + + inline void release() + { + layer0_.release(); + pyramid_.clear(); + nLayers_ = 0; + } + +private: + GpuMat layer0_; + std::vector pyramid_; + int nLayers_; +}; + +}} // namespace cv { namespace gpu { + +#endif /* __OPENCV_GPUWARPING_HPP__ */ diff --git a/modules/gpuwarping/perf/perf_main.cpp b/modules/gpuwarping/perf/perf_main.cpp new file mode 100644 index 0000000000..a7c1d5c85c --- /dev/null +++ b/modules/gpuwarping/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(gpuwarping, printCudaInfo()) diff --git a/modules/gpuwarping/perf/perf_precomp.cpp b/modules/gpuwarping/perf/perf_precomp.cpp new file mode 100644 index 0000000000..81f16e8f14 --- /dev/null +++ b/modules/gpuwarping/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/gpuwarping/perf/perf_precomp.hpp b/modules/gpuwarping/perf/perf_precomp.hpp new file mode 100644 index 0000000000..fc46b26474 --- /dev/null +++ b/modules/gpuwarping/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/gpuwarping.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/gpuwarping/perf/perf_warping.cpp b/modules/gpuwarping/perf/perf_warping.cpp new file mode 100644 index 0000000000..fd555cbe56 --- /dev/null +++ b/modules/gpuwarping/perf/perf_warping.cpp @@ -0,0 +1,592 @@ +/*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 std; +using namespace testing; +using namespace perf; + +////////////////////////////////////////////////////////////////////// +// Remap + +enum { HALF_SIZE=0, UPSIDE_DOWN, REFLECTION_X, REFLECTION_BOTH }; +CV_ENUM(RemapMode, HALF_SIZE, UPSIDE_DOWN, REFLECTION_X, REFLECTION_BOTH) + +void generateMap(cv::Mat& map_x, cv::Mat& map_y, int remapMode) +{ + for (int j = 0; j < map_x.rows; ++j) + { + for (int i = 0; i < map_x.cols; ++i) + { + switch (remapMode) + { + case HALF_SIZE: + if (i > map_x.cols*0.25 && i < map_x.cols*0.75 && j > map_x.rows*0.25 && j < map_x.rows*0.75) + { + map_x.at(j,i) = 2.f * (i - map_x.cols * 0.25f) + 0.5f; + map_y.at(j,i) = 2.f * (j - map_x.rows * 0.25f) + 0.5f; + } + else + { + map_x.at(j,i) = 0.f; + map_y.at(j,i) = 0.f; + } + break; + case UPSIDE_DOWN: + map_x.at(j,i) = static_cast(i); + map_y.at(j,i) = static_cast(map_x.rows - j); + break; + case REFLECTION_X: + map_x.at(j,i) = static_cast(map_x.cols - i); + map_y.at(j,i) = static_cast(j); + break; + case REFLECTION_BOTH: + map_x.at(j,i) = static_cast(map_x.cols - i); + map_y.at(j,i) = static_cast(map_x.rows - j); + break; + } // end of switch + } + } +} + +DEF_PARAM_TEST(Sz_Depth_Cn_Inter_Border_Mode, cv::Size, MatDepth, MatCn, Interpolation, BorderMode, RemapMode); + +PERF_TEST_P(Sz_Depth_Cn_Inter_Border_Mode, Remap, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4, + Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), + ALL_BORDER_MODES, + RemapMode::all())) +{ + declare.time(20.0); + + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + const int interpolation = GET_PARAM(3); + const int borderMode = GET_PARAM(4); + const int remapMode = GET_PARAM(5); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + cv::Mat xmap(size, CV_32FC1); + cv::Mat ymap(size, CV_32FC1); + generateMap(xmap, ymap, remapMode); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + const cv::gpu::GpuMat d_xmap(xmap); + const cv::gpu::GpuMat d_ymap(ymap); + cv::gpu::GpuMat dst; + + TEST_CYCLE() cv::gpu::remap(d_src, dst, d_xmap, d_ymap, interpolation, borderMode); + + GPU_SANITY_CHECK(dst); + } + else + { + cv::Mat dst; + + TEST_CYCLE() cv::remap(src, dst, xmap, ymap, interpolation, borderMode); + + CPU_SANITY_CHECK(dst); + } +} + +////////////////////////////////////////////////////////////////////// +// Resize + +DEF_PARAM_TEST(Sz_Depth_Cn_Inter_Scale, cv::Size, MatDepth, MatCn, Interpolation, double); + +PERF_TEST_P(Sz_Depth_Cn_Inter_Scale, Resize, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4, + Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), + Values(0.5, 0.3, 2.0))) +{ + declare.time(20.0); + + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + const int interpolation = GET_PARAM(3); + const double f = GET_PARAM(4); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + cv::gpu::GpuMat dst; + + TEST_CYCLE() cv::gpu::resize(d_src, dst, cv::Size(), f, f, interpolation); + + GPU_SANITY_CHECK(dst, 1e-3, ERROR_RELATIVE); + } + else + { + cv::Mat dst; + + TEST_CYCLE() cv::resize(src, dst, cv::Size(), f, f, interpolation); + + CPU_SANITY_CHECK(dst); + } +} + +////////////////////////////////////////////////////////////////////// +// ResizeArea + +DEF_PARAM_TEST(Sz_Depth_Cn_Scale, cv::Size, MatDepth, MatCn, double); + +PERF_TEST_P(Sz_Depth_Cn_Scale, ResizeArea, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4, + Values(0.2, 0.1, 0.05))) +{ + declare.time(1.0); + + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + const int interpolation = cv::INTER_AREA; + const double f = GET_PARAM(3); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + cv::gpu::GpuMat dst; + + TEST_CYCLE() cv::gpu::resize(d_src, dst, cv::Size(), f, f, interpolation); + + GPU_SANITY_CHECK(dst); + } + else + { + cv::Mat dst; + + TEST_CYCLE() cv::resize(src, dst, cv::Size(), f, f, interpolation); + + CPU_SANITY_CHECK(dst); + } +} + +////////////////////////////////////////////////////////////////////// +// WarpAffine + +DEF_PARAM_TEST(Sz_Depth_Cn_Inter_Border, cv::Size, MatDepth, MatCn, Interpolation, BorderMode); + +PERF_TEST_P(Sz_Depth_Cn_Inter_Border, WarpAffine, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4, + Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), + ALL_BORDER_MODES)) +{ + declare.time(20.0); + + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + const int interpolation = GET_PARAM(3); + const int borderMode = GET_PARAM(4); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + const double aplha = CV_PI / 4; + const double mat[2 * 3] = + { + std::cos(aplha), -std::sin(aplha), src.cols / 2, + std::sin(aplha), std::cos(aplha), 0 + }; + const cv::Mat M(2, 3, CV_64F, (void*) mat); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + cv::gpu::GpuMat dst; + + TEST_CYCLE() cv::gpu::warpAffine(d_src, dst, M, size, interpolation, borderMode); + + GPU_SANITY_CHECK(dst, 1); + } + else + { + cv::Mat dst; + + TEST_CYCLE() cv::warpAffine(src, dst, M, size, interpolation, borderMode); + + CPU_SANITY_CHECK(dst); + } +} + +////////////////////////////////////////////////////////////////////// +// WarpPerspective + +PERF_TEST_P(Sz_Depth_Cn_Inter_Border, WarpPerspective, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4, + Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), + ALL_BORDER_MODES)) +{ + declare.time(20.0); + + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + const int interpolation = GET_PARAM(3); + const int borderMode = GET_PARAM(4); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + const double aplha = CV_PI / 4; + double mat[3][3] = { {std::cos(aplha), -std::sin(aplha), src.cols / 2}, + {std::sin(aplha), std::cos(aplha), 0}, + {0.0, 0.0, 1.0}}; + const cv::Mat M(3, 3, CV_64F, (void*) mat); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + cv::gpu::GpuMat dst; + + TEST_CYCLE() cv::gpu::warpPerspective(d_src, dst, M, size, interpolation, borderMode); + + GPU_SANITY_CHECK(dst, 1); + } + else + { + cv::Mat dst; + + TEST_CYCLE() cv::warpPerspective(src, dst, M, size, interpolation, borderMode); + + CPU_SANITY_CHECK(dst); + } +} + +////////////////////////////////////////////////////////////////////// +// BuildWarpPlaneMaps + +PERF_TEST_P(Sz, BuildWarpPlaneMaps, + GPU_TYPICAL_MAT_SIZES) +{ + const cv::Size size = GetParam(); + + const cv::Mat K = cv::Mat::eye(3, 3, CV_32FC1); + const cv::Mat R = cv::Mat::ones(3, 3, CV_32FC1); + const cv::Mat T = cv::Mat::zeros(1, 3, CV_32F); + + if (PERF_RUN_GPU()) + { + cv::gpu::GpuMat map_x; + cv::gpu::GpuMat map_y; + + TEST_CYCLE() cv::gpu::buildWarpPlaneMaps(size, cv::Rect(0, 0, size.width, size.height), K, R, T, 1.0, map_x, map_y); + + GPU_SANITY_CHECK(map_x); + GPU_SANITY_CHECK(map_y); + } + else + { + FAIL_NO_CPU(); + } +} + +////////////////////////////////////////////////////////////////////// +// BuildWarpCylindricalMaps + +PERF_TEST_P(Sz, BuildWarpCylindricalMaps, + GPU_TYPICAL_MAT_SIZES) +{ + const cv::Size size = GetParam(); + + const cv::Mat K = cv::Mat::eye(3, 3, CV_32FC1); + const cv::Mat R = cv::Mat::ones(3, 3, CV_32FC1); + + if (PERF_RUN_GPU()) + { + cv::gpu::GpuMat map_x; + cv::gpu::GpuMat map_y; + + TEST_CYCLE() cv::gpu::buildWarpCylindricalMaps(size, cv::Rect(0, 0, size.width, size.height), K, R, 1.0, map_x, map_y); + + GPU_SANITY_CHECK(map_x); + GPU_SANITY_CHECK(map_y); + } + else + { + FAIL_NO_CPU(); + } +} + +////////////////////////////////////////////////////////////////////// +// BuildWarpSphericalMaps + +PERF_TEST_P(Sz, BuildWarpSphericalMaps, + GPU_TYPICAL_MAT_SIZES) +{ + const cv::Size size = GetParam(); + + const cv::Mat K = cv::Mat::eye(3, 3, CV_32FC1); + const cv::Mat R = cv::Mat::ones(3, 3, CV_32FC1); + + if (PERF_RUN_GPU()) + { + cv::gpu::GpuMat map_x; + cv::gpu::GpuMat map_y; + + TEST_CYCLE() cv::gpu::buildWarpSphericalMaps(size, cv::Rect(0, 0, size.width, size.height), K, R, 1.0, map_x, map_y); + + GPU_SANITY_CHECK(map_x); + GPU_SANITY_CHECK(map_y); + } + else + { + FAIL_NO_CPU(); + } +} + +////////////////////////////////////////////////////////////////////// +// Rotate + +DEF_PARAM_TEST(Sz_Depth_Cn_Inter, cv::Size, MatDepth, MatCn, Interpolation); + +PERF_TEST_P(Sz_Depth_Cn_Inter, Rotate, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4, + Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)))) +{ + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + const int interpolation = GET_PARAM(3); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + cv::gpu::GpuMat dst; + + TEST_CYCLE() cv::gpu::rotate(d_src, dst, size, 30.0, 0, 0, interpolation); + + GPU_SANITY_CHECK(dst, 1e-3, ERROR_RELATIVE); + } + else + { + FAIL_NO_CPU(); + } +} + +////////////////////////////////////////////////////////////////////// +// PyrDown + +PERF_TEST_P(Sz_Depth_Cn, PyrDown, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4)) +{ + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + cv::gpu::GpuMat dst; + + TEST_CYCLE() cv::gpu::pyrDown(d_src, dst); + + GPU_SANITY_CHECK(dst); + } + else + { + cv::Mat dst; + + TEST_CYCLE() cv::pyrDown(src, dst); + + CPU_SANITY_CHECK(dst); + } +} + +////////////////////////////////////////////////////////////////////// +// PyrUp + +PERF_TEST_P(Sz_Depth_Cn, PyrUp, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4)) +{ + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + cv::gpu::GpuMat dst; + + TEST_CYCLE() cv::gpu::pyrUp(d_src, dst); + + GPU_SANITY_CHECK(dst); + } + else + { + cv::Mat dst; + + TEST_CYCLE() cv::pyrUp(src, dst); + + CPU_SANITY_CHECK(dst); + } +} + +////////////////////////////////////////////////////////////////////// +// ImagePyramidBuild + +PERF_TEST_P(Sz_Depth_Cn, ImagePyramidBuild, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4)) +{ + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + const int nLayers = 5; + const cv::Size dstSize(size.width / 2 + 10, size.height / 2 + 10); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + + cv::gpu::ImagePyramid d_pyr; + + TEST_CYCLE() d_pyr.build(d_src, nLayers); + + cv::gpu::GpuMat dst; + d_pyr.getLayer(dst, dstSize); + + GPU_SANITY_CHECK(dst); + } + else + { + FAIL_NO_CPU(); + } +} + +////////////////////////////////////////////////////////////////////// +// ImagePyramidGetLayer + +PERF_TEST_P(Sz_Depth_Cn, ImagePyramidGetLayer, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4)) +{ + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + const int nLayers = 3; + const cv::Size dstSize(size.width / 2 + 10, size.height / 2 + 10); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + cv::gpu::GpuMat dst; + + cv::gpu::ImagePyramid d_pyr(d_src, nLayers); + + TEST_CYCLE() d_pyr.getLayer(dst, dstSize); + + GPU_SANITY_CHECK(dst); + } + else + { + FAIL_NO_CPU(); + } +} diff --git a/modules/gpuwarping/src/cuda/build_warp_maps.cu b/modules/gpuwarping/src/cuda/build_warp_maps.cu new file mode 100644 index 0000000000..6bd4e335bb --- /dev/null +++ b/modules/gpuwarping/src/cuda/build_warp_maps.cu @@ -0,0 +1,221 @@ +/*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/vec_traits.hpp" +#include "opencv2/core/cuda/vec_math.hpp" +#include "opencv2/core/cuda/saturate_cast.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" + +namespace cv { namespace gpu { namespace cudev +{ + namespace imgproc + { + // TODO use intrinsics like __sinf and so on + + namespace build_warp_maps + { + + __constant__ float ck_rinv[9]; + __constant__ float cr_kinv[9]; + __constant__ float ct[3]; + __constant__ float cscale; + } + + + class PlaneMapper + { + public: + static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y) + { + using namespace build_warp_maps; + + float x_ = u / cscale - ct[0]; + float y_ = v / cscale - ct[1]; + + float z; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * (1 - ct[2]); + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * (1 - ct[2]); + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * (1 - ct[2]); + + x /= z; + y /= z; + } + }; + + + class CylindricalMapper + { + public: + static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y) + { + using namespace build_warp_maps; + + u /= cscale; + float x_ = ::sinf(u); + float y_ = v / cscale; + float z_ = ::cosf(u); + + float z; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; + + if (z > 0) { x /= z; y /= z; } + else x = y = -1; + } + }; + + + class SphericalMapper + { + public: + static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y) + { + using namespace build_warp_maps; + + v /= cscale; + u /= cscale; + + float sinv = ::sinf(v); + float x_ = sinv * ::sinf(u); + float y_ = -::cosf(v); + float z_ = sinv * ::cosf(u); + + float z; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * z_; + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * z_; + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * z_; + + if (z > 0) { x /= z; y /= z; } + else x = y = -1; + } + }; + + + template + __global__ void buildWarpMapsKernel(int tl_u, int tl_v, int cols, int rows, + PtrStepf map_x, PtrStepf map_y) + { + int du = blockIdx.x * blockDim.x + threadIdx.x; + int dv = blockIdx.y * blockDim.y + threadIdx.y; + if (du < cols && dv < rows) + { + float u = tl_u + du; + float v = tl_v + dv; + float x, y; + Mapper::mapBackward(u, v, x, y); + map_x.ptr(dv)[du] = x; + map_y.ptr(dv)[du] = y; + } + } + + + void buildWarpPlaneMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, + const float k_rinv[9], const float r_kinv[9], const float t[3], + float scale, cudaStream_t stream) + { + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ct, t, 3*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); + + int cols = map_x.cols; + int rows = map_x.rows; + + dim3 threads(32, 8); + dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); + + buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); + cudaSafeCall(cudaGetLastError()); + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + + + void buildWarpCylindricalMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, + const float k_rinv[9], const float r_kinv[9], float scale, + cudaStream_t stream) + { + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); + + int cols = map_x.cols; + int rows = map_x.rows; + + dim3 threads(32, 8); + dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); + + buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); + cudaSafeCall(cudaGetLastError()); + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + + + void buildWarpSphericalMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, + const float k_rinv[9], const float r_kinv[9], float scale, + cudaStream_t stream) + { + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); + + int cols = map_x.cols; + int rows = map_x.rows; + + dim3 threads(32, 8); + dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); + + buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); + cudaSafeCall(cudaGetLastError()); + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + } // namespace imgproc +}}} // namespace cv { namespace gpu { namespace cudev { + + +#endif /* CUDA_DISABLER */ diff --git a/modules/gpuimgproc/src/cuda/pyr_down.cu b/modules/gpuwarping/src/cuda/pyr_down.cu similarity index 100% rename from modules/gpuimgproc/src/cuda/pyr_down.cu rename to modules/gpuwarping/src/cuda/pyr_down.cu diff --git a/modules/gpuimgproc/src/cuda/pyr_up.cu b/modules/gpuwarping/src/cuda/pyr_up.cu similarity index 100% rename from modules/gpuimgproc/src/cuda/pyr_up.cu rename to modules/gpuwarping/src/cuda/pyr_up.cu diff --git a/modules/gpuimgproc/src/cuda/remap.cu b/modules/gpuwarping/src/cuda/remap.cu similarity index 100% rename from modules/gpuimgproc/src/cuda/remap.cu rename to modules/gpuwarping/src/cuda/remap.cu diff --git a/modules/gpuimgproc/src/cuda/resize.cu b/modules/gpuwarping/src/cuda/resize.cu similarity index 100% rename from modules/gpuimgproc/src/cuda/resize.cu rename to modules/gpuwarping/src/cuda/resize.cu diff --git a/modules/gpuimgproc/src/cuda/warp.cu b/modules/gpuwarping/src/cuda/warp.cu similarity index 100% rename from modules/gpuimgproc/src/cuda/warp.cu rename to modules/gpuwarping/src/cuda/warp.cu diff --git a/modules/gpuwarping/src/precomp.cpp b/modules/gpuwarping/src/precomp.cpp new file mode 100644 index 0000000000..3c01a2596d --- /dev/null +++ b/modules/gpuwarping/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/gpuwarping/src/precomp.hpp b/modules/gpuwarping/src/precomp.hpp new file mode 100644 index 0000000000..52cc69a346 --- /dev/null +++ b/modules/gpuwarping/src/precomp.hpp @@ -0,0 +1,57 @@ +/*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 "opencv2/gpuwarping.hpp" + +#include "opencv2/core/gpu_private.hpp" + +#include "opencv2/opencv_modules.hpp" + +#ifdef HAVE_OPENCV_GPULEGACY +# include "opencv2/gpulegacy.hpp" +# include "opencv2/gpulegacy/private.hpp" +#endif + +#endif /* __OPENCV_PRECOMP_H__ */ diff --git a/modules/gpuimgproc/src/pyramids.cpp b/modules/gpuwarping/src/pyramids.cpp similarity index 97% rename from modules/gpuimgproc/src/pyramids.cpp rename to modules/gpuwarping/src/pyramids.cpp index 9e9fbe3437..91b568d706 100644 --- a/modules/gpuimgproc/src/pyramids.cpp +++ b/modules/gpuwarping/src/pyramids.cpp @@ -45,7 +45,9 @@ #if !defined HAVE_CUDA || defined(CUDA_DISABLER) void cv::gpu::pyrDown(const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } + void cv::gpu::pyrUp(const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } + void cv::gpu::ImagePyramid::build(const GpuMat&, int, Stream&) { throw_no_cuda(); } void cv::gpu::ImagePyramid::getLayer(GpuMat&, Size, Stream&) const { throw_no_cuda(); } @@ -130,17 +132,14 @@ void cv::gpu::pyrUp(const GpuMat& src, GpuMat& dst, Stream& stream) ////////////////////////////////////////////////////////////////////////////// // ImagePyramid -namespace cv { namespace gpu { namespace cudev -{ - namespace pyramid - { - template void kernelDownsampleX2_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void kernelInterpolateFrom1_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - } -}}} - void cv::gpu::ImagePyramid::build(const GpuMat& img, int numLayers, Stream& stream) { +#ifndef HAVE_OPENCV_GPULEGACY + (void) img; + (void) numLayers; + (void) stream; + throw_no_cuda(); +#else using namespace cv::gpu::cudev::pyramid; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -185,10 +184,17 @@ void cv::gpu::ImagePyramid::build(const GpuMat& img, int numLayers, Stream& stre szLastLayer = szCurLayer; } +#endif } void cv::gpu::ImagePyramid::getLayer(GpuMat& outImg, Size outRoi, Stream& stream) const { +#ifndef HAVE_OPENCV_GPULEGACY + (void) outImg; + (void) outRoi; + (void) stream; + throw_no_cuda(); +#else using namespace cv::gpu::cudev::pyramid; typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -244,6 +250,7 @@ void cv::gpu::ImagePyramid::getLayer(GpuMat& outImg, Size outRoi, Stream& stream } func(lastLayer, outImg, StreamAccessor::getStream(stream)); +#endif } #endif // HAVE_CUDA diff --git a/modules/gpuimgproc/src/remap.cpp b/modules/gpuwarping/src/remap.cpp similarity index 100% rename from modules/gpuimgproc/src/remap.cpp rename to modules/gpuwarping/src/remap.cpp diff --git a/modules/gpuimgproc/src/resize.cpp b/modules/gpuwarping/src/resize.cpp similarity index 96% rename from modules/gpuimgproc/src/resize.cpp rename to modules/gpuwarping/src/resize.cpp index 32afa54de9..68708b41a0 100644 --- a/modules/gpuimgproc/src/resize.cpp +++ b/modules/gpuwarping/src/resize.cpp @@ -44,18 +44,7 @@ #if !defined HAVE_CUDA || defined(CUDA_DISABLER) -void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, double fy, int interpolation, Stream& s) -{ - (void)src; - (void)dst; - (void)dsize; - (void)fx; - (void)fy; - (void)interpolation; - (void)s; - - throw_no_cuda(); -} +void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int, Stream&) { throw_no_cuda(); } #else // HAVE_CUDA diff --git a/modules/gpuimgproc/src/warp.cpp b/modules/gpuwarping/src/warp.cpp similarity index 70% rename from modules/gpuimgproc/src/warp.cpp rename to modules/gpuwarping/src/warp.cpp index 007091e6a3..e15c11b748 100644 --- a/modules/gpuimgproc/src/warp.cpp +++ b/modules/gpuwarping/src/warp.cpp @@ -42,8 +42,10 @@ #include "precomp.hpp" -#if !defined HAVE_CUDA || defined(CUDA_DISABLER) +using namespace cv; +using namespace cv::gpu; +#if !defined HAVE_CUDA || defined(CUDA_DISABLER) void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int, int, Scalar, Stream&) { throw_no_cuda(); } void cv::gpu::buildWarpAffineMaps(const Mat&, bool, Size, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } @@ -51,6 +53,12 @@ void cv::gpu::buildWarpAffineMaps(const Mat&, bool, Size, GpuMat&, GpuMat&, Stre void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int, int, Scalar, Stream&) { throw_no_cuda(); } void cv::gpu::buildWarpPerspectiveMaps(const Mat&, bool, Size, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } +void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } + +void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int, Stream&) { throw_no_cuda(); } + #else // HAVE_CUDA namespace cv { namespace gpu { namespace cudev @@ -121,27 +129,18 @@ void cv::gpu::buildWarpPerspectiveMaps(const Mat& M, bool inverse, Size dsize, G namespace { - template struct NppTypeTraits; - template<> struct NppTypeTraits { typedef Npp8u npp_t; }; - template<> struct NppTypeTraits { typedef Npp8s npp_t; }; - template<> struct NppTypeTraits { typedef Npp16u npp_t; }; - template<> struct NppTypeTraits { typedef Npp16s npp_t; typedef Npp16sc npp_complex_type; }; - template<> struct NppTypeTraits { typedef Npp32s npp_t; typedef Npp32sc npp_complex_type; }; - template<> struct NppTypeTraits { typedef Npp32f npp_t; typedef Npp32fc npp_complex_type; }; - template<> struct NppTypeTraits { typedef Npp64f npp_t; typedef Npp64fc npp_complex_type; }; - template struct NppWarpFunc { - typedef typename NppTypeTraits::npp_t npp_t; + typedef typename NPPTypeTraits::npp_type npp_type; - typedef NppStatus (*func_t)(const npp_t* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, npp_t* pDst, + typedef NppStatus (*func_t)(const npp_type* pSrc, NppiSize srcSize, int srcStep, NppiRect srcRoi, npp_type* pDst, int dstStep, NppiRect dstRoi, const double coeffs[][3], int interpolation); }; template ::func_t func> struct NppWarp { - typedef typename NppWarpFunc::npp_t npp_t; + typedef typename NppWarpFunc::npp_type npp_type; static void call(const cv::gpu::GpuMat& src, cv::gpu::GpuMat& dst, double coeffs[][3], int interpolation, cudaStream_t stream) { @@ -165,8 +164,8 @@ namespace cv::gpu::NppStreamHandler h(stream); - nppSafeCall( func(src.ptr(), srcsz, static_cast(src.step), srcroi, - dst.ptr(), static_cast(dst.step), dstroi, + nppSafeCall( func(src.ptr(), srcsz, static_cast(src.step), srcroi, + dst.ptr(), static_cast(dst.step), dstroi, coeffs, npp_inter[interpolation]) ); if (stream == 0) @@ -451,4 +450,173 @@ void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size } } +////////////////////////////////////////////////////////////////////////////// +// buildWarpPlaneMaps + +namespace cv { namespace gpu { namespace cudev +{ + namespace imgproc + { + void buildWarpPlaneMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, + const float k_rinv[9], const float r_kinv[9], const float t[3], float scale, + cudaStream_t stream); + } +}}} + +void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, const Mat &T, + float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream) +{ + (void)src_size; + using namespace ::cv::gpu::cudev::imgproc; + + CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); + CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); + CV_Assert((T.size() == Size(3,1) || T.size() == Size(1,3)) && T.type() == CV_32F && T.isContinuous()); + + Mat K_Rinv = K * R.t(); + Mat R_Kinv = R * K.inv(); + CV_Assert(K_Rinv.isContinuous()); + CV_Assert(R_Kinv.isContinuous()); + + map_x.create(dst_roi.size(), CV_32F); + map_y.create(dst_roi.size(), CV_32F); + cudev::imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), + T.ptr(), scale, StreamAccessor::getStream(stream)); +} + +////////////////////////////////////////////////////////////////////////////// +// buildWarpCylyndricalMaps + +namespace cv { namespace gpu { namespace cudev +{ + namespace imgproc + { + void buildWarpCylindricalMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, + const float k_rinv[9], const float r_kinv[9], float scale, + cudaStream_t stream); + } +}}} + +void cv::gpu::buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, + GpuMat& map_x, GpuMat& map_y, Stream& stream) +{ + (void)src_size; + using namespace ::cv::gpu::cudev::imgproc; + + CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); + CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); + + Mat K_Rinv = K * R.t(); + Mat R_Kinv = R * K.inv(); + CV_Assert(K_Rinv.isContinuous()); + CV_Assert(R_Kinv.isContinuous()); + + map_x.create(dst_roi.size(), CV_32F); + map_y.create(dst_roi.size(), CV_32F); + cudev::imgproc::buildWarpCylindricalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), scale, StreamAccessor::getStream(stream)); +} + + +////////////////////////////////////////////////////////////////////////////// +// buildWarpSphericalMaps + +namespace cv { namespace gpu { namespace cudev +{ + namespace imgproc + { + void buildWarpSphericalMaps(int tl_u, int tl_v, PtrStepSzf map_x, PtrStepSzf map_y, + const float k_rinv[9], const float r_kinv[9], float scale, + cudaStream_t stream); + } +}}} + +void cv::gpu::buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, + GpuMat& map_x, GpuMat& map_y, Stream& stream) +{ + (void)src_size; + using namespace ::cv::gpu::cudev::imgproc; + + CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); + CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); + + Mat K_Rinv = K * R.t(); + Mat R_Kinv = R * K.inv(); + CV_Assert(K_Rinv.isContinuous()); + CV_Assert(R_Kinv.isContinuous()); + + map_x.create(dst_roi.size(), CV_32F); + map_y.create(dst_roi.size(), CV_32F); + cudev::imgproc::buildWarpSphericalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), scale, StreamAccessor::getStream(stream)); +} + +//////////////////////////////////////////////////////////////////////// +// rotate + +namespace +{ + template struct NppRotateFunc + { + typedef typename NPPTypeTraits::npp_type npp_type; + + typedef NppStatus (*func_t)(const npp_type* pSrc, NppiSize oSrcSize, int nSrcStep, NppiRect oSrcROI, + npp_type* pDst, int nDstStep, NppiRect oDstROI, + double nAngle, double nShiftX, double nShiftY, int eInterpolation); + }; + + template ::func_t func> struct NppRotate + { + typedef typename NppRotateFunc::npp_type npp_type; + + static void call(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, cudaStream_t stream) + { + (void)dsize; + static const int npp_inter[] = {NPPI_INTER_NN, NPPI_INTER_LINEAR, NPPI_INTER_CUBIC}; + + NppStreamHandler h(stream); + + NppiSize srcsz; + srcsz.height = src.rows; + srcsz.width = src.cols; + NppiRect srcroi; + srcroi.x = srcroi.y = 0; + srcroi.height = src.rows; + srcroi.width = src.cols; + NppiRect dstroi; + dstroi.x = dstroi.y = 0; + dstroi.height = dst.rows; + dstroi.width = dst.cols; + + nppSafeCall( func(src.ptr(), srcsz, static_cast(src.step), srcroi, + dst.ptr(), static_cast(dst.step), dstroi, angle, xShift, yShift, npp_inter[interpolation]) ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } + }; +} + +void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, Stream& stream) +{ + typedef void (*func_t)(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift, double yShift, int interpolation, cudaStream_t stream); + + static const func_t funcs[6][4] = + { + {NppRotate::call, 0, NppRotate::call, NppRotate::call}, + {0,0,0,0}, + {NppRotate::call, 0, NppRotate::call, NppRotate::call}, + {0,0,0,0}, + {0,0,0,0}, + {NppRotate::call, 0, NppRotate::call, NppRotate::call} + }; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 1 || src.channels() == 3 || src.channels() == 4); + CV_Assert(interpolation == INTER_NEAREST || interpolation == INTER_LINEAR || interpolation == INTER_CUBIC); + + dst.create(dsize, src.type()); + dst.setTo(Scalar::all(0)); + + funcs[src.depth()][src.channels() - 1](src, dst, dsize, angle, xShift, yShift, interpolation, StreamAccessor::getStream(stream)); +} + #endif // HAVE_CUDA diff --git a/modules/gpuimgproc/test/interpolation.hpp b/modules/gpuwarping/test/interpolation.hpp similarity index 100% rename from modules/gpuimgproc/test/interpolation.hpp rename to modules/gpuwarping/test/interpolation.hpp diff --git a/modules/gpuwarping/test/test_main.cpp b/modules/gpuwarping/test/test_main.cpp new file mode 100644 index 0000000000..eea3d7c008 --- /dev/null +++ b/modules/gpuwarping/test/test_main.cpp @@ -0,0 +1,45 @@ +/*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" + +CV_GPU_TEST_MAIN("gpu") diff --git a/modules/gpuwarping/test/test_precomp.cpp b/modules/gpuwarping/test/test_precomp.cpp new file mode 100644 index 0000000000..0fb6521809 --- /dev/null +++ b/modules/gpuwarping/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/gpuwarping/test/test_precomp.hpp b/modules/gpuwarping/test/test_precomp.hpp new file mode 100644 index 0000000000..90a28fe1a7 --- /dev/null +++ b/modules/gpuwarping/test/test_precomp.hpp @@ -0,0 +1,62 @@ +/*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/gpuwarping.hpp" +#include "opencv2/imgproc.hpp" + +#include "interpolation.hpp" + +#endif diff --git a/modules/gpuimgproc/test/test_pyramids.cpp b/modules/gpuwarping/test/test_pyramids.cpp similarity index 97% rename from modules/gpuimgproc/test/test_pyramids.cpp rename to modules/gpuwarping/test/test_pyramids.cpp index 6b0540fc10..f296b7d6f2 100644 --- a/modules/gpuimgproc/test/test_pyramids.cpp +++ b/modules/gpuwarping/test/test_pyramids.cpp @@ -80,7 +80,7 @@ GPU_TEST_P(PyrDown, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-4 : 1.0); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, PyrDown, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, PyrDown, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), @@ -120,7 +120,7 @@ GPU_TEST_P(PyrUp, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-4 : 1.0); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, PyrUp, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, PyrUp, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), diff --git a/modules/gpuimgproc/test/test_remap.cpp b/modules/gpuwarping/test/test_remap.cpp similarity index 99% rename from modules/gpuimgproc/test/test_remap.cpp rename to modules/gpuwarping/test/test_remap.cpp index eb4b9ece85..c1899ff611 100644 --- a/modules/gpuimgproc/test/test_remap.cpp +++ b/modules/gpuwarping/test/test_remap.cpp @@ -169,7 +169,7 @@ GPU_TEST_P(Remap, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-3 : 1.0); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, Remap, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, Remap, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), diff --git a/modules/gpuimgproc/test/test_resize.cpp b/modules/gpuwarping/test/test_resize.cpp similarity index 97% rename from modules/gpuimgproc/test/test_resize.cpp rename to modules/gpuwarping/test/test_resize.cpp index 593c891e6a..27289834a9 100644 --- a/modules/gpuimgproc/test/test_resize.cpp +++ b/modules/gpuwarping/test/test_resize.cpp @@ -152,7 +152,7 @@ GPU_TEST_P(Resize, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-2 : 1.0); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, Resize, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, Resize, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC3), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), @@ -198,7 +198,7 @@ GPU_TEST_P(ResizeSameAsHost, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-2 : 1.0); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, ResizeSameAsHost, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, ResizeSameAsHost, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC3), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), @@ -241,7 +241,7 @@ GPU_TEST_P(ResizeNPP, Accuracy) EXPECT_MAT_SIMILAR(dst_gold, dst, 1e-1); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, ResizeNPP, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, ResizeNPP, testing::Combine( ALL_DEVICES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), testing::Values(0.3, 0.5, 1.5, 2.0), diff --git a/modules/gpuimgproc/test/test_warp_affine.cpp b/modules/gpuwarping/test/test_warp_affine.cpp similarity index 98% rename from modules/gpuimgproc/test/test_warp_affine.cpp rename to modules/gpuwarping/test/test_warp_affine.cpp index 43bf0f6d9e..206446c4d0 100644 --- a/modules/gpuimgproc/test/test_warp_affine.cpp +++ b/modules/gpuwarping/test/test_warp_affine.cpp @@ -101,7 +101,7 @@ GPU_TEST_P(BuildWarpAffineMaps, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, 0.0); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, BuildWarpAffineMaps, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, BuildWarpAffineMaps, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, DIRECT_INVERSE)); @@ -222,7 +222,7 @@ GPU_TEST_P(WarpAffine, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-1 : 1.0); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, WarpAffine, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, WarpAffine, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), @@ -271,7 +271,7 @@ GPU_TEST_P(WarpAffineNPP, Accuracy) EXPECT_MAT_SIMILAR(dst_gold, dst, 2e-2); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, WarpAffineNPP, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, WarpAffineNPP, testing::Combine( ALL_DEVICES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), DIRECT_INVERSE, diff --git a/modules/gpuimgproc/test/test_warp_perspective.cpp b/modules/gpuwarping/test/test_warp_perspective.cpp similarity index 98% rename from modules/gpuimgproc/test/test_warp_perspective.cpp rename to modules/gpuwarping/test/test_warp_perspective.cpp index d225e58b66..49f844c3f6 100644 --- a/modules/gpuimgproc/test/test_warp_perspective.cpp +++ b/modules/gpuwarping/test/test_warp_perspective.cpp @@ -102,7 +102,7 @@ GPU_TEST_P(BuildWarpPerspectiveMaps, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, 0.0); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, BuildWarpPerspectiveMaps, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, BuildWarpPerspectiveMaps, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, DIRECT_INVERSE)); @@ -225,7 +225,7 @@ GPU_TEST_P(WarpPerspective, Accuracy) EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-1 : 1.0); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, WarpPerspective, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, WarpPerspective, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), @@ -274,7 +274,7 @@ GPU_TEST_P(WarpPerspectiveNPP, Accuracy) EXPECT_MAT_SIMILAR(dst_gold, dst, 2e-2); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, WarpPerspectiveNPP, testing::Combine( +INSTANTIATE_TEST_CASE_P(GPU_Warping, WarpPerspectiveNPP, testing::Combine( ALL_DEVICES, testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), DIRECT_INVERSE, diff --git a/samples/cpp/CMakeLists.txt b/samples/cpp/CMakeLists.txt index e90bcb6acd..d1807e9290 100644 --- a/samples/cpp/CMakeLists.txt +++ b/samples/cpp/CMakeLists.txt @@ -19,6 +19,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/gpuwarping/include") ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpuimgproc/include") ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpufeatures2d/include") ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpuvideo/include") diff --git a/samples/gpu/CMakeLists.txt b/samples/gpu/CMakeLists.txt index d69ccdc2ea..670c71c9e2 100644 --- a/samples/gpu/CMakeLists.txt +++ b/samples/gpu/CMakeLists.txt @@ -2,8 +2,9 @@ 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_gpufilters opencv_gpulegacy opencv_gpuimgproc opencv_gpufeatures2d opencv_gpuvideo opencv_gpuobjdetect - opencv_gpucalib3d) + opencv_gpucodec opencv_gpuarithm opencv_gpufilters opencv_gpuwarping opencv_gpuimgproc + opencv_gpufeatures2d opencv_gpuvideo opencv_gpuobjdetect + opencv_gpucalib3d opencv_gpulegacy) ocv_check_dependencies(${OPENCV_GPU_SAMPLES_REQUIRED_DEPS})