From ade7394e77e46e1f7bd45329865530965733083a Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Wed, 14 Mar 2012 15:54:17 +0000 Subject: [PATCH] refactored and fixed bugs in gpu warp functions (remap, resize, warpAffine, warpPerspective) wrote more complicated tests for them implemented own version of warpAffine and warpPerspective for different border interpolation types refactored some gpu tests --- modules/core/src/opengl_interop.cpp | 2 +- modules/gpu/include/opencv2/gpu/gpu.hpp | 128 +- modules/gpu/perf/perf_imgproc.cpp | 132 +- modules/gpu/src/color.cpp | 319 +- modules/gpu/src/cuda/remap.cu | 56 +- modules/gpu/src/cuda/resize.cu | 121 +- modules/gpu/src/cuda/warp.cu | 380 + modules/gpu/src/imgproc.cpp | 331 - .../gpu/src/opencv2/gpu/device/filters.hpp | 30 +- modules/gpu/src/remap.cpp | 105 + modules/gpu/src/resize.cpp | 150 + modules/gpu/src/warp.cpp | 463 ++ modules/gpu/test/interpolation.hpp | 120 + modules/gpu/test/{test_main.cpp => main.cpp} | 2 +- .../test/{test_precomp.cpp => precomp.cpp} | 2 +- .../test/{test_precomp.hpp => precomp.hpp} | 7 +- modules/gpu/test/test_arithm.cpp | 58 +- modules/gpu/test/test_calib3d.cpp | 2 +- modules/gpu/test/test_copy_make_border.cpp | 90 + modules/gpu/test/test_features2d.cpp | 2 +- modules/gpu/test/test_filters.cpp | 20 +- modules/gpu/test/test_hog.cpp | 2 +- modules/gpu/test/test_imgproc.cpp | 7080 ++++++++--------- modules/gpu/test/test_matop.cpp | 12 +- modules/gpu/test/test_nvidia.cpp | 2 +- modules/gpu/test/test_remap.cpp | 177 + modules/gpu/test/test_resize.cpp | 202 + modules/gpu/test/test_threshold.cpp | 88 + modules/gpu/test/test_video.cpp | 2 +- modules/gpu/test/test_warp_affine.cpp | 275 + modules/gpu/test/test_warp_perspective.cpp | 275 + .../test/{test_gpu_base.cpp => utility.cpp} | 102 +- .../test/{test_gpu_base.hpp => utility.hpp} | 49 +- 33 files changed, 6242 insertions(+), 4544 deletions(-) create mode 100644 modules/gpu/src/cuda/warp.cu create mode 100644 modules/gpu/src/remap.cpp create mode 100644 modules/gpu/src/resize.cpp create mode 100644 modules/gpu/src/warp.cpp create mode 100644 modules/gpu/test/interpolation.hpp rename modules/gpu/test/{test_main.cpp => main.cpp} (96%) rename modules/gpu/test/{test_precomp.cpp => precomp.cpp} (96%) rename modules/gpu/test/{test_precomp.hpp => precomp.hpp} (95%) create mode 100644 modules/gpu/test/test_copy_make_border.cpp create mode 100644 modules/gpu/test/test_remap.cpp create mode 100644 modules/gpu/test/test_resize.cpp create mode 100644 modules/gpu/test/test_threshold.cpp create mode 100644 modules/gpu/test/test_warp_affine.cpp create mode 100644 modules/gpu/test/test_warp_perspective.cpp rename modules/gpu/test/{test_gpu_base.cpp => utility.cpp} (65%) rename modules/gpu/test/{test_gpu_base.hpp => utility.hpp} (78%) diff --git a/modules/core/src/opengl_interop.cpp b/modules/core/src/opengl_interop.cpp index a9c7819d00..e30096ff03 100644 --- a/modules/core/src/opengl_interop.cpp +++ b/modules/core/src/opengl_interop.cpp @@ -1251,7 +1251,7 @@ cv::GlFont::GlFont(const string& family, int height, Weight weight, Style style) base_ = glGenLists(256); CV_CheckGlError(); - glFuncTab()->generateBitmapFont(family, height, weight, style & STYLE_ITALIC, style & STYLE_UNDERLINE, 0, 256, base_); + glFuncTab()->generateBitmapFont(family, height, weight, (style & STYLE_ITALIC) != 0, (style & STYLE_UNDERLINE) != 0, 0, 256, base_); #endif } diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index cc6d607509..244f336d1d 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -64,7 +64,7 @@ CV_EXPORTS int getCudaEnabledDeviceCount(); CV_EXPORTS void setDevice(int device); CV_EXPORTS int getDevice(); -//! Explicitly destroys and cleans up all resources associated with the current device in the current process. +//! Explicitly destroys and cleans up all resources associated with the current device in the current process. //! Any subsequent API call to this device will reinitialize the device. CV_EXPORTS void resetDevice(); @@ -81,7 +81,7 @@ enum FeatureSet NATIVE_DOUBLE = FEATURE_SET_COMPUTE_13 }; -// Gives information about what GPU archs this OpenCV GPU module was +// Gives information about what GPU archs this OpenCV GPU module was // compiled for class CV_EXPORTS TargetArchs { @@ -266,10 +266,10 @@ private: Impl *impl; friend struct StreamAccessor; - + explicit Stream(Impl* impl); }; - + //////////////////////////////// Filter Engine //////////////////////////////// @@ -432,26 +432,26 @@ CV_EXPORTS Ptr getMinFilter_GPU(int srcType, int dstType, const CV_EXPORTS void boxFilter(const GpuMat& src, GpuMat& dst, int ddepth, Size ksize, Point anchor = Point(-1,-1), Stream& stream = Stream::Null()); //! a synonym for normalized box filter -static inline void blur(const GpuMat& src, GpuMat& dst, Size ksize, Point anchor = Point(-1,-1), Stream& stream = Stream::Null()) -{ - boxFilter(src, dst, -1, ksize, anchor, stream); +static inline void blur(const GpuMat& src, GpuMat& dst, Size ksize, Point anchor = Point(-1,-1), Stream& stream = Stream::Null()) +{ + boxFilter(src, dst, -1, ksize, anchor, stream); } //! erodes the image (applies the local minimum operator) CV_EXPORTS void erode(const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor = Point(-1, -1), int iterations = 1); -CV_EXPORTS void erode(const GpuMat& src, GpuMat& dst, const Mat& kernel, GpuMat& buf, - Point anchor = Point(-1, -1), int iterations = 1, +CV_EXPORTS void erode(const GpuMat& src, GpuMat& dst, const Mat& kernel, GpuMat& buf, + Point anchor = Point(-1, -1), int iterations = 1, Stream& stream = Stream::Null()); //! dilates the image (applies the local maximum operator) CV_EXPORTS void dilate(const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor = Point(-1, -1), int iterations = 1); -CV_EXPORTS void dilate(const GpuMat& src, GpuMat& dst, const Mat& kernel, GpuMat& buf, - Point anchor = Point(-1, -1), int iterations = 1, +CV_EXPORTS void dilate(const GpuMat& src, GpuMat& dst, const Mat& kernel, GpuMat& buf, + Point anchor = Point(-1, -1), int iterations = 1, Stream& stream = Stream::Null()); //! applies an advanced morphological operation to the image CV_EXPORTS void morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, Point anchor = Point(-1, -1), int iterations = 1); -CV_EXPORTS void morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, GpuMat& buf1, GpuMat& buf2, +CV_EXPORTS void morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, GpuMat& buf1, GpuMat& buf2, Point anchor = Point(-1, -1), int iterations = 1, Stream& stream = Stream::Null()); //! applies non-separable 2D linear filter to the image @@ -461,7 +461,7 @@ CV_EXPORTS void filter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& CV_EXPORTS void sepFilter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& kernelX, const Mat& kernelY, Point anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1); CV_EXPORTS void sepFilter2D(const GpuMat& src, GpuMat& dst, int ddepth, const Mat& kernelX, const Mat& kernelY, GpuMat& buf, - Point anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, + Point anchor = Point(-1,-1), int rowBorderType = BORDER_DEFAULT, int columnBorderType = -1, Stream& stream = Stream::Null()); //! applies generalized Sobel operator to the image @@ -490,7 +490,7 @@ CV_EXPORTS void Laplacian(const GpuMat& src, GpuMat& dst, int ddepth, int ksize ////////////////////////////// Arithmetics /////////////////////////////////// //! implements generalized matrix product algorithm GEMM from BLAS -CV_EXPORTS void gemm(const GpuMat& src1, const GpuMat& src2, double alpha, +CV_EXPORTS void gemm(const GpuMat& src1, const GpuMat& src2, double alpha, const GpuMat& src3, double beta, GpuMat& dst, int flags = 0, Stream& stream = Stream::Null()); //! transposes the matrix @@ -572,7 +572,7 @@ CV_EXPORTS void divide(const GpuMat& a, const Scalar& sc, GpuMat& c, double scal CV_EXPORTS void divide(double scale, const GpuMat& src2, GpuMat& dst, int dtype = -1, Stream& stream = Stream::Null()); //! computes the weighted sum of two arrays (dst = alpha*src1 + beta*src2 + gamma) -CV_EXPORTS void addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, +CV_EXPORTS void addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, int dtype = -1, Stream& stream = Stream::Null()); //! adds scaled array to another one (dst = alpha*src1 + src2) @@ -669,17 +669,17 @@ CV_EXPORTS void alphaComp(const GpuMat& img1, const GpuMat& img2, GpuMat& dst, i //! 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, const Scalar& borderValue = Scalar(), + 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), + TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1), Stream& stream = Stream::Null()); //! Does mean shift procedure on GPU. CV_EXPORTS void meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int sp, int sr, - TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1), + TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1), Stream& stream = Stream::Null()); //! Does mean shift segmentation with elimination of small regions. @@ -717,11 +717,17 @@ CV_EXPORTS void resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx=0, //! 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, Stream& stream = Stream::Null()); +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, Stream& stream = Stream::Null()); +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, @@ -738,11 +744,11 @@ CV_EXPORTS void buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat &K //! 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, +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()); //! copies 2D array to a larger destination array and pads borders with user-specifiable constant -CV_EXPORTS void copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, +CV_EXPORTS void copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value = Scalar(), Stream& stream = Stream::Null()); //! computes the integral image @@ -768,13 +774,13 @@ CV_EXPORTS void rectStdDev(const GpuMat& src, const GpuMat& sqr, GpuMat& dst, co //! 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); -CV_EXPORTS void cornerHarris(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, GpuMat& buf, int blockSize, int ksize, double k, +CV_EXPORTS void cornerHarris(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, GpuMat& buf, int blockSize, int ksize, double k, int borderType = BORDER_REFLECT101, Stream& stream = Stream::Null()); //! computes minimum eigen value of 2x2 derivative covariation matrix at each pixel - the cornerness criteria CV_EXPORTS void cornerMinEigenVal(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, int borderType=BORDER_REFLECT101); CV_EXPORTS void cornerMinEigenVal(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, int blockSize, int ksize, int borderType=BORDER_REFLECT101); -CV_EXPORTS void cornerMinEigenVal(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, GpuMat& buf, int blockSize, int ksize, +CV_EXPORTS void cornerMinEigenVal(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& Dy, GpuMat& buf, int blockSize, int ksize, int borderType=BORDER_REFLECT101, Stream& stream = Stream::Null()); //! performs per-element multiplication of two full (not packed) Fourier spectrums @@ -787,7 +793,7 @@ CV_EXPORTS void mulAndScaleSpectrums(const GpuMat& a, const GpuMat& b, GpuMat& c //! Performs a forward or inverse discrete Fourier transform (1D or 2D) of floating point matrix. //! Param dft_size is the size of DFT transform. -//! +//! //! If the source matrix is not continous, then additional copy will be done, //! so to avoid copying ensure the source matrix is continous one. If you want to use //! preallocated output ensure it is continuous too, otherwise it will be reallocated. @@ -808,7 +814,7 @@ CV_EXPORTS void convolve(const GpuMat& image, const GpuMat& templ, GpuMat& resul struct CV_EXPORTS ConvolveBuf { ConvolveBuf() {} - ConvolveBuf(Size image_size, Size templ_size) + ConvolveBuf(Size image_size, Size templ_size) { create(image_size, templ_size); } void create(Size image_size, Size templ_size); void create(Size image_size, Size templ_size, Size block_size); @@ -837,10 +843,10 @@ CV_EXPORTS void pyrUp(const GpuMat& src, GpuMat& dst, int borderType = BORDER_DE //! 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, +CV_EXPORTS void blendLinear(const GpuMat& img1, const GpuMat& img2, const GpuMat& weights1, const GpuMat& weights2, GpuMat& result, Stream& stream = Stream::Null()); - + struct CV_EXPORTS CannyBuf; CV_EXPORTS void Canny(const GpuMat& image, GpuMat& edges, double low_thresh, double high_thresh, int apperture_size = 3, bool L2gradient = false); @@ -855,7 +861,7 @@ struct CV_EXPORTS CannyBuf CannyBuf(const GpuMat& dx_, const GpuMat& dy_); void create(const Size& image_size, int apperture_size = 3); - + void release(); GpuMat dx, dy; @@ -968,24 +974,24 @@ CV_EXPORTS void transformPoints(const GpuMat& src, const Mat& rvec, const Mat& t GpuMat& dst, Stream& stream = Stream::Null()); CV_EXPORTS void projectPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec, - const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst, + const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst, Stream& stream = Stream::Null()); CV_EXPORTS void solvePnPRansac(const Mat& object, const Mat& image, const Mat& camera_mat, const Mat& dist_coef, Mat& rvec, Mat& tvec, bool use_extrinsic_guess=false, - int num_iters=100, float max_dist=8.0, int min_inlier_count=100, + int num_iters=100, float max_dist=8.0, int min_inlier_count=100, std::vector* inliers=NULL); //////////////////////////////// Image Labeling //////////////////////////////// -//!performs labeling via graph cuts of a 2D regular 4-connected graph. -CV_EXPORTS void graphcut(GpuMat& terminals, GpuMat& leftTransp, GpuMat& rightTransp, GpuMat& top, GpuMat& bottom, GpuMat& labels, +//!performs labeling via graph cuts of a 2D regular 4-connected graph. +CV_EXPORTS void graphcut(GpuMat& terminals, GpuMat& leftTransp, GpuMat& rightTransp, GpuMat& top, GpuMat& bottom, GpuMat& labels, GpuMat& buf, Stream& stream = Stream::Null()); -//!performs labeling via graph cuts of a 2D regular 8-connected graph. -CV_EXPORTS void graphcut(GpuMat& terminals, GpuMat& leftTransp, GpuMat& rightTransp, GpuMat& top, GpuMat& topLeft, GpuMat& topRight, +//!performs labeling via graph cuts of a 2D regular 8-connected graph. +CV_EXPORTS void graphcut(GpuMat& terminals, GpuMat& leftTransp, GpuMat& rightTransp, GpuMat& top, GpuMat& topLeft, GpuMat& topRight, GpuMat& bottom, GpuMat& bottomLeft, GpuMat& bottomRight, - GpuMat& labels, + GpuMat& labels, GpuMat& buf, Stream& stream = Stream::Null()); ////////////////////////////////// Histograms ////////////////////////////////// @@ -1243,16 +1249,16 @@ struct CV_EXPORTS HOGDescriptor static vector getPeopleDetector48x96(); static vector getPeopleDetector64x128(); - void detect(const GpuMat& img, vector& found_locations, - double hit_threshold=0, Size win_stride=Size(), + void detect(const GpuMat& img, vector& found_locations, + double hit_threshold=0, Size win_stride=Size(), Size padding=Size()); void detectMultiScale(const GpuMat& img, vector& found_locations, - double hit_threshold=0, Size win_stride=Size(), - Size padding=Size(), double scale0=1.05, + double hit_threshold=0, Size win_stride=Size(), + Size padding=Size(), double scale0=1.05, int group_threshold=2); - void getDescriptors(const GpuMat& img, Size win_stride, + void getDescriptors(const GpuMat& img, Size win_stride, GpuMat& descriptors, int descr_format=DESCR_FORMAT_COL_BY_COL); @@ -1290,11 +1296,11 @@ protected: // Gradients conputation results GpuMat grad, qangle, grad_buf, qangle_buf; - // returns subbuffer with required size, reallocates buffer if nessesary. - static GpuMat getBuffer(const Size& sz, int type, GpuMat& buf); - static GpuMat getBuffer(int rows, int cols, int type, GpuMat& buf); + // returns subbuffer with required size, reallocates buffer if nessesary. + static GpuMat getBuffer(const Size& sz, int type, GpuMat& buf); + static GpuMat getBuffer(int rows, int cols, int type, GpuMat& buf); - std::vector image_scales; + std::vector image_scales; }; @@ -1323,8 +1329,8 @@ public: bool isMaskSupported() const; // Find one best match for each query descriptor - void matchSingle(const GpuMat& query, const GpuMat& train, - GpuMat& trainIdx, GpuMat& distance, + void matchSingle(const GpuMat& query, const GpuMat& train, + GpuMat& trainIdx, GpuMat& distance, const GpuMat& mask = GpuMat(), Stream& stream = Stream::Null()); // Download trainIdx and distance and convert it to CPU vector with DMatch @@ -1339,7 +1345,7 @@ public: void makeGpuCollection(GpuMat& trainCollection, GpuMat& maskCollection, const std::vector& masks = std::vector()); // Find one best match from train collection for each query descriptor - void matchCollection(const GpuMat& query, const GpuMat& trainCollection, + void matchCollection(const GpuMat& query, const GpuMat& trainCollection, GpuMat& trainIdx, GpuMat& imgIdx, GpuMat& distance, const GpuMat& masks = GpuMat(), Stream& stream = Stream::Null()); @@ -1508,7 +1514,7 @@ private: class CV_EXPORTS SURF_GPU : public CvSURFParams { public: - enum KeypointLayout + enum KeypointLayout { SF_X = 0, SF_Y, @@ -1535,7 +1541,7 @@ public: //! download descriptors from device to host memory void downloadDescriptors(const GpuMat& descriptorsGPU, vector& descriptors); - + //! finds the keypoints using fast hessian detector used in SURF //! supports CV_8UC1 images //! keypoints will have nFeature cols and 6 rows @@ -1546,16 +1552,16 @@ public: //! keypoints.ptr(SF_DIR)[i] will contain orientation of i'th feature //! keypoints.ptr(SF_HESSIAN)[i] will contain response of i'th feature void operator()(const GpuMat& img, const GpuMat& mask, GpuMat& keypoints); - //! finds the keypoints and computes their descriptors. + //! finds the keypoints and computes their descriptors. //! Optionally it can compute descriptors for the user-provided keypoints and recompute keypoints direction - void operator()(const GpuMat& img, const GpuMat& mask, GpuMat& keypoints, GpuMat& descriptors, + void operator()(const GpuMat& img, const GpuMat& mask, GpuMat& keypoints, GpuMat& descriptors, bool useProvidedKeypoints = false); void operator()(const GpuMat& img, const GpuMat& mask, std::vector& keypoints); - void operator()(const GpuMat& img, const GpuMat& mask, std::vector& keypoints, GpuMat& descriptors, + void operator()(const GpuMat& img, const GpuMat& mask, std::vector& keypoints, GpuMat& descriptors, bool useProvidedKeypoints = false); - void operator()(const GpuMat& img, const GpuMat& mask, std::vector& keypoints, std::vector& descriptors, + void operator()(const GpuMat& img, const GpuMat& mask, std::vector& keypoints, std::vector& descriptors, bool useProvidedKeypoints = false); void releaseMemory(); @@ -1589,7 +1595,7 @@ public: //! finds the keypoints using FAST detector //! supports only CV_8UC1 images - void operator ()(const GpuMat& image, const GpuMat& mask, GpuMat& keypoints); + void operator ()(const GpuMat& image, const GpuMat& mask, GpuMat& keypoints); void operator ()(const GpuMat& image, const GpuMat& mask, std::vector& keypoints); //! download keypoints from device to host memory @@ -1709,7 +1715,7 @@ private: GpuMat pattern_; std::vector imagePyr_; - std::vector maskPyr_; + std::vector maskPyr_; GpuMat buf_; @@ -1729,7 +1735,7 @@ class CV_EXPORTS BroxOpticalFlow { public: BroxOpticalFlow(float alpha_, float gamma_, float scale_factor_, int inner_iterations_, int outer_iterations_, int solver_iterations_) : - alpha(alpha_), gamma(gamma_), scale_factor(scale_factor_), + alpha(alpha_), gamma(gamma_), scale_factor(scale_factor_), inner_iterations(inner_iterations_), outer_iterations(outer_iterations_), solver_iterations(solver_iterations_) { } @@ -1857,7 +1863,7 @@ private: GpuMat dy_calcBuf_; vector prevPyr_; - vector nextPyr_; + vector nextPyr_; GpuMat dx_buf_; GpuMat dy_buf_; @@ -1943,10 +1949,10 @@ private: //! occlusion masks 0, occlusion masks 1, //! interpolated forward flow 0, interpolated forward flow 1, //! interpolated backward flow 0, interpolated backward flow 1 -//! -CV_EXPORTS void interpolateFrames(const GpuMat& frame0, const GpuMat& frame1, +//! +CV_EXPORTS void interpolateFrames(const GpuMat& frame0, const GpuMat& frame1, const GpuMat& fu, const GpuMat& fv, - const GpuMat& bu, const GpuMat& bv, + const GpuMat& bu, const GpuMat& bv, float pos, GpuMat& newFrame, GpuMat& buf, Stream& stream = Stream::Null()); diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp index 8c31d83fbe..5472acf85f 100644 --- a/modules/gpu/perf/perf_imgproc.cpp +++ b/modules/gpu/perf/perf_imgproc.cpp @@ -35,9 +35,9 @@ GPU_PERF_TEST(Remap, cv::gpu::DeviceInfo, cv::Size, perf::MatType, Interpolation } INSTANTIATE_TEST_CASE_P(ImgProc, Remap, testing::Combine( - ALL_DEVICES, - GPU_TYPICAL_MAT_SIZES, - testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_32FC1), + ALL_DEVICES, + GPU_TYPICAL_MAT_SIZES, + testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_32FC1), testing::Values((int) cv::INTER_NEAREST, (int) cv::INTER_LINEAR, (int) cv::INTER_CUBIC), testing::Values((int) cv::BORDER_REFLECT101, (int) cv::BORDER_REPLICATE, (int) cv::BORDER_CONSTANT))); @@ -52,7 +52,7 @@ GPU_PERF_TEST_1(MeanShiftFiltering, cv::gpu::DeviceInfo) cv::Mat img = readImage("gpu/meanshift/cones.png"); ASSERT_FALSE(img.empty()); - + cv::Mat rgba; cv::cvtColor(img, rgba, cv::COLOR_BGR2BGRA); @@ -80,7 +80,7 @@ GPU_PERF_TEST_1(MeanShiftProc, cv::gpu::DeviceInfo) cv::Mat img = readImage("gpu/meanshift/cones.png"); ASSERT_FALSE(img.empty()); - + cv::Mat rgba; cv::cvtColor(img, rgba, cv::COLOR_BGR2BGRA); @@ -109,7 +109,7 @@ GPU_PERF_TEST_1(MeanShiftSegmentation, cv::gpu::DeviceInfo) cv::Mat img = readImage("gpu/meanshift/cones.png"); ASSERT_FALSE(img.empty()); - + cv::Mat rgba; cv::cvtColor(img, rgba, cv::COLOR_BGR2BGRA); @@ -151,8 +151,8 @@ GPU_PERF_TEST(DrawColorDisp, cv::gpu::DeviceInfo, cv::Size, perf::MatType) } INSTANTIATE_TEST_CASE_P(ImgProc, DrawColorDisp, testing::Combine( - ALL_DEVICES, - GPU_TYPICAL_MAT_SIZES, + ALL_DEVICES, + GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_16SC1))); ////////////////////////////////////////////////////////////////////// @@ -180,8 +180,8 @@ GPU_PERF_TEST(ReprojectImageTo3D, cv::gpu::DeviceInfo, cv::Size, perf::MatType) } INSTANTIATE_TEST_CASE_P(ImgProc, ReprojectImageTo3D, testing::Combine( - ALL_DEVICES, - GPU_TYPICAL_MAT_SIZES, + ALL_DEVICES, + GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_16SC1))); ////////////////////////////////////////////////////////////////////// @@ -210,12 +210,12 @@ GPU_PERF_TEST(CvtColor, cv::gpu::DeviceInfo, cv::Size, perf::MatType, CvtColorIn } INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor, testing::Combine( - ALL_DEVICES, - GPU_TYPICAL_MAT_SIZES, + ALL_DEVICES, + GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_16UC1, CV_32FC1), testing::Values( - CvtColorInfo(4, 4, cv::COLOR_RGBA2BGRA), CvtColorInfo(4, 1, cv::COLOR_BGRA2GRAY), CvtColorInfo(1, 4, cv::COLOR_GRAY2BGRA), - CvtColorInfo(4, 4, cv::COLOR_BGR2XYZ), CvtColorInfo(4, 4, cv::COLOR_BGR2YCrCb), CvtColorInfo(4, 4, cv::COLOR_YCrCb2BGR), + CvtColorInfo(4, 4, cv::COLOR_RGBA2BGRA), CvtColorInfo(4, 1, cv::COLOR_BGRA2GRAY), CvtColorInfo(1, 4, cv::COLOR_GRAY2BGRA), + CvtColorInfo(4, 4, cv::COLOR_BGR2XYZ), CvtColorInfo(4, 4, cv::COLOR_BGR2YCrCb), CvtColorInfo(4, 4, cv::COLOR_YCrCb2BGR), CvtColorInfo(4, 4, cv::COLOR_BGR2HSV), CvtColorInfo(4, 4, cv::COLOR_HSV2BGR)))); ////////////////////////////////////////////////////////////////////// @@ -269,8 +269,8 @@ GPU_PERF_TEST(Threshold, cv::gpu::DeviceInfo, cv::Size, perf::MatType) } INSTANTIATE_TEST_CASE_P(ImgProc, Threshold, testing::Combine( - ALL_DEVICES, - GPU_TYPICAL_MAT_SIZES, + ALL_DEVICES, + GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_16UC1, CV_32FC1))); ////////////////////////////////////////////////////////////////////// @@ -302,8 +302,8 @@ GPU_PERF_TEST(Resize, cv::gpu::DeviceInfo, cv::Size, perf::MatType, Interpolatio } INSTANTIATE_TEST_CASE_P(ImgProc, Resize, testing::Combine( - ALL_DEVICES, - testing::Values(perf::szSXGA, perf::sz1080p), + ALL_DEVICES, + testing::Values(perf::szSXGA, perf::sz1080p), testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_32FC1), testing::Values((int) cv::INTER_NEAREST, (int) cv::INTER_LINEAR, (int) cv::INTER_CUBIC), testing::Values(0.5, 2.0))); @@ -327,22 +327,21 @@ GPU_PERF_TEST(WarpAffine, cv::gpu::DeviceInfo, cv::Size, perf::MatType, Interpol cv::gpu::GpuMat src(src_host); cv::gpu::GpuMat dst; - double reflect[2][3] = { {-1, 0, 0}, - { 0, -1, 0}}; - reflect[0][2] = size.width; - reflect[1][2] = size.height; - cv::Mat M(2, 3, CV_64F, (void*) reflect); + const double aplha = CV_PI / 4; + double mat[2][3] = { {std::cos(aplha), -std::sin(aplha), src.cols / 2}, + {std::sin(aplha), std::cos(aplha), 0}}; + cv::Mat M(2, 3, CV_64F, (void*) mat); TEST_CYCLE() { - cv::gpu::warpAffine(src, dst, M, size, interpolation); + cv::gpu::warpAffine(src, dst, M, size, interpolation, cv::BORDER_CONSTANT, cv::Scalar()); } } INSTANTIATE_TEST_CASE_P(ImgProc, WarpAffine, testing::Combine( - ALL_DEVICES, - GPU_TYPICAL_MAT_SIZES, - testing::Values(CV_8UC1, CV_8UC4, CV_32FC1), + ALL_DEVICES, + GPU_TYPICAL_MAT_SIZES, + testing::Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_16UC1, CV_16UC3, CV_16UC4, CV_32FC1, CV_32FC3, CV_32FC4), testing::Values((int) cv::INTER_NEAREST, (int) cv::INTER_LINEAR, (int) cv::INTER_CUBIC))); ////////////////////////////////////////////////////////////////////// @@ -364,23 +363,22 @@ GPU_PERF_TEST(WarpPerspective, cv::gpu::DeviceInfo, cv::Size, perf::MatType, Int cv::gpu::GpuMat src(src_host); cv::gpu::GpuMat dst; - double reflect[3][3] = { {-1, 0, 0}, - { 0, -1, 0}, - { 0, 0, 1}}; - reflect[0][2] = size.width; - reflect[1][2] = size.height; - cv::Mat M(3, 3, CV_64F, (void*)reflect); + 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}}; + cv::Mat M(3, 3, CV_64F, (void*) mat); TEST_CYCLE() { - cv::gpu::warpPerspective(src, dst, M, size, interpolation); + cv::gpu::warpPerspective(src, dst, M, size, interpolation, cv::BORDER_CONSTANT, cv::Scalar()); } } INSTANTIATE_TEST_CASE_P(ImgProc, WarpPerspective, testing::Combine( - ALL_DEVICES, - GPU_TYPICAL_MAT_SIZES, - testing::Values(CV_8UC1, CV_8UC4, CV_32FC1), + ALL_DEVICES, + GPU_TYPICAL_MAT_SIZES, + testing::Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_16UC1, CV_16UC3, CV_16UC4, CV_32FC1, CV_32FC3, CV_32FC4), testing::Values((int) cv::INTER_NEAREST, (int) cv::INTER_LINEAR, (int) cv::INTER_CUBIC))); ////////////////////////////////////////////////////////////////////// @@ -398,13 +396,13 @@ GPU_PERF_TEST(BuildWarpPlaneMaps, cv::gpu::DeviceInfo, cv::Size) TEST_CYCLE() { - cv::gpu::buildWarpPlaneMaps(size, cv::Rect(0, 0, size.width, size.height), cv::Mat::eye(3, 3, CV_32FC1), + cv::gpu::buildWarpPlaneMaps(size, cv::Rect(0, 0, size.width, size.height), cv::Mat::eye(3, 3, CV_32FC1), cv::Mat::ones(3, 3, CV_32FC1), cv::Mat::zeros(1, 3, CV_32F), 1.0, map_x, map_y); } } INSTANTIATE_TEST_CASE_P(ImgProc, BuildWarpPlaneMaps, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -428,7 +426,7 @@ GPU_PERF_TEST(BuildWarpCylindricalMaps, cv::gpu::DeviceInfo, cv::Size) } INSTANTIATE_TEST_CASE_P(ImgProc, BuildWarpCylindricalMaps, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -452,7 +450,7 @@ GPU_PERF_TEST(BuildWarpSphericalMaps, cv::gpu::DeviceInfo, cv::Size) } INSTANTIATE_TEST_CASE_P(ImgProc, BuildWarpSphericalMaps, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -481,8 +479,8 @@ GPU_PERF_TEST(Rotate, cv::gpu::DeviceInfo, cv::Size, perf::MatType, Interpolatio } INSTANTIATE_TEST_CASE_P(ImgProc, Rotate, testing::Combine( - ALL_DEVICES, - GPU_TYPICAL_MAT_SIZES, + ALL_DEVICES, + GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_16UC1, CV_16UC3, CV_16UC4, CV_32FC1, CV_32FC3, CV_32FC4), testing::Values((int) cv::INTER_NEAREST, (int) cv::INTER_LINEAR, (int) cv::INTER_CUBIC))); @@ -512,8 +510,8 @@ GPU_PERF_TEST(CopyMakeBorder, cv::gpu::DeviceInfo, cv::Size, perf::MatType, Bord } INSTANTIATE_TEST_CASE_P(ImgProc, CopyMakeBorder, testing::Combine( - ALL_DEVICES, - GPU_TYPICAL_MAT_SIZES, + ALL_DEVICES, + GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_8UC4, CV_32FC1), testing::Values((int) cv::BORDER_REPLICATE, (int) cv::BORDER_REFLECT, (int) cv::BORDER_WRAP, (int) cv::BORDER_CONSTANT))); @@ -542,7 +540,7 @@ GPU_PERF_TEST(Integral, cv::gpu::DeviceInfo, cv::Size) } INSTANTIATE_TEST_CASE_P(ImgProc, Integral, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -569,7 +567,7 @@ GPU_PERF_TEST(IntegralSqr, cv::gpu::DeviceInfo, cv::Size) } INSTANTIATE_TEST_CASE_P(ImgProc, IntegralSqr, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -596,7 +594,7 @@ GPU_PERF_TEST(ColumnSum, cv::gpu::DeviceInfo, cv::Size) } INSTANTIATE_TEST_CASE_P(ImgProc, ColumnSum, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -608,7 +606,7 @@ GPU_PERF_TEST(CornerHarris, cv::gpu::DeviceInfo, perf::MatType) int type = GET_PARAM(1); cv::gpu::setDevice(devInfo.deviceID()); - + cv::Mat img = readImage("gpu/stereobm/aloe-L.png", cv::IMREAD_GRAYSCALE); ASSERT_FALSE(img.empty()); @@ -620,7 +618,7 @@ GPU_PERF_TEST(CornerHarris, cv::gpu::DeviceInfo, perf::MatType) cv::gpu::GpuMat Dy; int blockSize = 3; - int ksize = 7; + int ksize = 7; double k = 0.5; TEST_CYCLE() @@ -630,7 +628,7 @@ GPU_PERF_TEST(CornerHarris, cv::gpu::DeviceInfo, perf::MatType) } INSTANTIATE_TEST_CASE_P(ImgProc, CornerHarris, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, testing::Values(CV_8UC1, CV_32FC1))); ////////////////////////////////////////////////////////////////////// @@ -642,7 +640,7 @@ GPU_PERF_TEST(CornerMinEigenVal, cv::gpu::DeviceInfo, perf::MatType) int type = GET_PARAM(1); cv::gpu::setDevice(devInfo.deviceID()); - + cv::Mat img = readImage("gpu/stereobm/aloe-L.png", cv::IMREAD_GRAYSCALE); ASSERT_FALSE(img.empty()); @@ -654,7 +652,7 @@ GPU_PERF_TEST(CornerMinEigenVal, cv::gpu::DeviceInfo, perf::MatType) cv::gpu::GpuMat Dy; int blockSize = 3; - int ksize = 7; + int ksize = 7; TEST_CYCLE() { @@ -663,7 +661,7 @@ GPU_PERF_TEST(CornerMinEigenVal, cv::gpu::DeviceInfo, perf::MatType) } INSTANTIATE_TEST_CASE_P(ImgProc, CornerMinEigenVal, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, testing::Values(CV_8UC1, CV_32FC1))); ////////////////////////////////////////////////////////////////////// @@ -692,7 +690,7 @@ GPU_PERF_TEST(MulSpectrums, cv::gpu::DeviceInfo, cv::Size) } INSTANTIATE_TEST_CASE_P(ImgProc, MulSpectrums, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -721,7 +719,7 @@ GPU_PERF_TEST(Dft, cv::gpu::DeviceInfo, cv::Size) } INSTANTIATE_TEST_CASE_P(ImgProc, Dft, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -754,7 +752,7 @@ GPU_PERF_TEST(Convolve, cv::gpu::DeviceInfo, cv::Size, int, bool) } INSTANTIATE_TEST_CASE_P(ImgProc, Convolve, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES, testing::Values(3, 9, 27, 32, 64), testing::Bool())); @@ -784,7 +782,7 @@ GPU_PERF_TEST(PyrDown, cv::gpu::DeviceInfo, cv::Size, perf::MatType) } INSTANTIATE_TEST_CASE_P(ImgProc, PyrDown, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_8UC4, CV_16SC3, CV_32FC1))); @@ -813,7 +811,7 @@ GPU_PERF_TEST(PyrUp, cv::gpu::DeviceInfo, cv::Size, perf::MatType) } INSTANTIATE_TEST_CASE_P(ImgProc, PyrUp, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_8UC4, CV_16SC3, CV_32FC1))); @@ -846,7 +844,7 @@ GPU_PERF_TEST(BlendLinear, cv::gpu::DeviceInfo, cv::Size, perf::MatType) } INSTANTIATE_TEST_CASE_P(ImgProc, BlendLinear, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_32FC1))); @@ -878,7 +876,7 @@ GPU_PERF_TEST(AlphaComp, cv::gpu::DeviceInfo, cv::Size, perf::MatType, AlphaOp) } INSTANTIATE_TEST_CASE_P(ImgProc, AlphaComp, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC4, CV_16UC4, CV_32SC4, CV_32FC4), testing::Values((int)cv::gpu::ALPHA_OVER, (int)cv::gpu::ALPHA_IN, (int)cv::gpu::ALPHA_OUT, (int)cv::gpu::ALPHA_ATOP, (int)cv::gpu::ALPHA_XOR, (int)cv::gpu::ALPHA_PLUS, (int)cv::gpu::ALPHA_OVER_PREMUL, (int)cv::gpu::ALPHA_IN_PREMUL, (int)cv::gpu::ALPHA_OUT_PREMUL, (int)cv::gpu::ALPHA_ATOP_PREMUL, (int)cv::gpu::ALPHA_XOR_PREMUL, (int)cv::gpu::ALPHA_PLUS_PREMUL, (int)cv::gpu::ALPHA_PREMUL))); @@ -932,7 +930,7 @@ GPU_PERF_TEST(CalcHist, cv::gpu::DeviceInfo, cv::Size) } INSTANTIATE_TEST_CASE_P(ImgProc, CalcHist, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -961,7 +959,7 @@ GPU_PERF_TEST(EqualizeHist, cv::gpu::DeviceInfo, cv::Size) } INSTANTIATE_TEST_CASE_P(ImgProc, EqualizeHist, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES)); ////////////////////////////////////////////////////////////////////// @@ -982,7 +980,7 @@ GPU_PERF_TEST(ImagePyramid_build, cv::gpu::DeviceInfo, cv::Size, perf::MatType) cv::gpu::GpuMat src(src_host); cv::gpu::ImagePyramid pyr; - + TEST_CYCLE() { pyr.build(src, 5); @@ -990,7 +988,7 @@ GPU_PERF_TEST(ImagePyramid_build, cv::gpu::DeviceInfo, cv::Size, perf::MatType) } INSTANTIATE_TEST_CASE_P(ImgProc, ImagePyramid_build, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_16UC1, CV_16UC3, CV_16UC4, CV_32FC1, CV_32FC3, CV_32FC4))); @@ -1010,7 +1008,7 @@ GPU_PERF_TEST(ImagePyramid_getLayer, cv::gpu::DeviceInfo, cv::Size, perf::MatTyp cv::gpu::GpuMat dst; cv::gpu::ImagePyramid pyr(src, 3); - + TEST_CYCLE() { pyr.getLayer(dst, cv::Size(size.width / 2 + 10, size.height / 2 + 10)); @@ -1018,7 +1016,7 @@ GPU_PERF_TEST(ImagePyramid_getLayer, cv::gpu::DeviceInfo, cv::Size, perf::MatTyp } INSTANTIATE_TEST_CASE_P(ImgProc, ImagePyramid_getLayer, testing::Combine( - ALL_DEVICES, + ALL_DEVICES, GPU_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_16UC1, CV_16UC3, CV_16UC4, CV_32FC1, CV_32FC3, CV_32FC4))); diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp index 89b44b9a9b..5fcd076edc 100644 --- a/modules/gpu/src/color.cpp +++ b/modules/gpu/src/color.cpp @@ -52,7 +52,7 @@ void cv::gpu::swapChannels(GpuMat&, const int[], Stream&) { throw_nogpu(); } #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace device { #define OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name) \ void name(const DevMem2Db& src, const DevMem2Db& dst, cudaStream_t stream); @@ -213,11 +213,11 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {bgr_to_rgb_8u, 0, bgr_to_rgb_16u, 0, 0, bgr_to_rgb_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -226,11 +226,11 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {bgr_to_bgra_8u, 0, bgr_to_bgra_16u, 0, 0, bgr_to_bgra_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -239,11 +239,11 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {bgr_to_rgba_8u, 0, bgr_to_rgba_16u, 0, 0, bgr_to_rgba_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -252,11 +252,11 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {bgra_to_bgr_8u, 0, bgra_to_bgr_16u, 0, 0, bgra_to_bgr_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -265,11 +265,11 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {bgra_to_rgb_8u, 0, bgra_to_rgb_16u, 0, 0, bgra_to_rgb_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -278,171 +278,171 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {bgra_to_rgba_8u, 0, bgra_to_rgba_16u, 0, 0, bgra_to_rgba_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } void bgr_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 3); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void bgr_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 3); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void rgb_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 3); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void rgb_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 3); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void bgra_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 4); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void bgra_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 4); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void rgba_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 4); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void rgba_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 4); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC3); + dst.create(src.size(), CV_8UC3); device::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC3); + dst.create(src.size(), CV_8UC3); device::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC3); + dst.create(src.size(), CV_8UC3); device::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC3); + dst.create(src.size(), CV_8UC3); device::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC4); + dst.create(src.size(), CV_8UC4); device::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC4); + dst.create(src.size(), CV_8UC4); device::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC4); + dst.create(src.size(), CV_8UC4); device::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC4); + dst.create(src.size(), CV_8UC4); device::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream)); } @@ -451,11 +451,11 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {gray_to_bgr_8u, 0, gray_to_bgr_16u, 0, 0, gray_to_bgr_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 1); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -464,51 +464,51 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {gray_to_bgra_8u, 0, gray_to_bgra_16u, 0, 0, gray_to_bgra_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 1); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } void gray_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 1); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream)); } void gray_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 1); - dst.create(src.size(), CV_8UC2); + dst.create(src.size(), CV_8UC2); device::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream)); } void bgr555_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC1); + dst.create(src.size(), CV_8UC1); device::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream)); } void bgr565_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) - { + { CV_Assert(src.depth() == CV_8U); CV_Assert(src.channels() == 2); - dst.create(src.size(), CV_8UC1); + dst.create(src.size(), CV_8UC1); device::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream)); } @@ -517,11 +517,11 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {rgb_to_gray_8u, 0, rgb_to_gray_16u, 0, 0, rgb_to_gray_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -530,11 +530,11 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {bgr_to_gray_8u, 0, bgr_to_gray_16u, 0, 0, bgr_to_gray_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -543,11 +543,11 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {rgba_to_gray_8u, 0, rgba_to_gray_16u, 0, 0, rgba_to_gray_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -556,19 +556,19 @@ namespace { using namespace cv::gpu::device; static const gpu_func_t funcs[] = {bgra_to_gray_8u, 0, bgra_to_gray_16u, 0, 0, bgra_to_gray_32f}; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); } - + void rgb_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {rgb_to_yuv_8u, 0, rgb_to_yuv_16u, 0, 0, rgb_to_yuv_32f}, @@ -581,12 +581,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -594,7 +594,7 @@ namespace void bgr_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {bgr_to_yuv_8u, 0, bgr_to_yuv_16u, 0, 0, bgr_to_yuv_32f}, @@ -607,12 +607,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -620,7 +620,7 @@ namespace void yuv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {yuv_to_rgb_8u, 0, yuv_to_rgb_16u, 0, 0, yuv_to_rgb_32f}, @@ -633,12 +633,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -646,7 +646,7 @@ namespace void yuv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {yuv_to_bgr_8u, 0, yuv_to_bgr_16u, 0, 0, yuv_to_bgr_32f}, @@ -659,20 +659,20 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } - + void rgb_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {rgb_to_YCrCb_8u, 0, rgb_to_YCrCb_16u, 0, 0, rgb_to_YCrCb_32f}, @@ -685,12 +685,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -698,7 +698,7 @@ namespace void bgr_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {bgr_to_YCrCb_8u, 0, bgr_to_YCrCb_16u, 0, 0, bgr_to_YCrCb_32f}, @@ -711,12 +711,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -724,7 +724,7 @@ namespace void YCrCb_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {YCrCb_to_rgb_8u, 0, YCrCb_to_rgb_16u, 0, 0, YCrCb_to_rgb_32f}, @@ -737,12 +737,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -750,7 +750,7 @@ namespace void YCrCb_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {YCrCb_to_bgr_8u, 0, YCrCb_to_bgr_16u, 0, 0, YCrCb_to_bgr_32f}, @@ -763,12 +763,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -776,7 +776,7 @@ namespace void rgb_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {rgb_to_xyz_8u, 0, rgb_to_xyz_16u, 0, 0, rgb_to_xyz_32f}, @@ -789,12 +789,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -802,7 +802,7 @@ namespace void bgr_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {bgr_to_xyz_8u, 0, bgr_to_xyz_16u, 0, 0, bgr_to_xyz_32f}, @@ -815,12 +815,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -828,7 +828,7 @@ namespace void xyz_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {xyz_to_rgb_8u, 0, xyz_to_rgb_16u, 0, 0, xyz_to_rgb_32f}, @@ -841,12 +841,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -854,7 +854,7 @@ namespace void xyz_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {xyz_to_bgr_8u, 0, xyz_to_bgr_16u, 0, 0, xyz_to_bgr_32f}, @@ -867,12 +867,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -880,7 +880,7 @@ namespace void rgb_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {rgb_to_hsv_8u, 0, 0, 0, 0, rgb_to_hsv_32f}, @@ -893,12 +893,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -906,7 +906,7 @@ namespace void bgr_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {bgr_to_hsv_8u, 0, 0, 0, 0, bgr_to_hsv_32f}, @@ -919,12 +919,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -932,7 +932,7 @@ namespace void hsv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {hsv_to_rgb_8u, 0, 0, 0, 0, hsv_to_rgb_32f}, @@ -945,12 +945,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -958,7 +958,7 @@ namespace void hsv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {hsv_to_bgr_8u, 0, 0, 0, 0, hsv_to_bgr_32f}, @@ -971,20 +971,20 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); - } + } void rgb_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {rgb_to_hls_8u, 0, 0, 0, 0, rgb_to_hls_32f}, @@ -997,12 +997,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1010,7 +1010,7 @@ namespace void bgr_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {bgr_to_hls_8u, 0, 0, 0, 0, bgr_to_hls_32f}, @@ -1023,12 +1023,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1036,7 +1036,7 @@ namespace void hls_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {hls_to_rgb_8u, 0, 0, 0, 0, hls_to_rgb_32f}, @@ -1049,12 +1049,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1062,7 +1062,7 @@ namespace void hls_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {hls_to_bgr_8u, 0, 0, 0, 0, hls_to_bgr_32f}, @@ -1075,20 +1075,20 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); - } + } void rgb_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {rgb_to_hsv_full_8u, 0, 0, 0, 0, rgb_to_hsv_full_32f}, @@ -1101,12 +1101,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1114,7 +1114,7 @@ namespace void bgr_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {bgr_to_hsv_full_8u, 0, 0, 0, 0, bgr_to_hsv_full_32f}, @@ -1127,12 +1127,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1140,7 +1140,7 @@ namespace void hsv_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {hsv_to_rgb_full_8u, 0, 0, 0, 0, hsv_to_rgb_full_32f}, @@ -1153,12 +1153,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1166,7 +1166,7 @@ namespace void hsv_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {hsv_to_bgr_full_8u, 0, 0, 0, 0, hsv_to_bgr_full_32f}, @@ -1179,20 +1179,20 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); - } + } void rgb_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {rgb_to_hls_full_8u, 0, 0, 0, 0, rgb_to_hls_full_32f}, @@ -1205,12 +1205,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1218,7 +1218,7 @@ namespace void bgr_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {bgr_to_hls_full_8u, 0, 0, 0, 0, bgr_to_hls_full_32f}, @@ -1231,12 +1231,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1244,7 +1244,7 @@ namespace void hls_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {hls_to_rgb_full_8u, 0, 0, 0, 0, hls_to_rgb_full_32f}, @@ -1257,12 +1257,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1270,7 +1270,7 @@ namespace void hls_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { using namespace cv::gpu::device; - static const gpu_func_t funcs[2][2][6] = + static const gpu_func_t funcs[2][2][6] = { { {hls_to_bgr_full_8u, 0, 0, 0, 0, hls_to_bgr_full_32f}, @@ -1283,12 +1283,12 @@ namespace }; if (dcn <= 0) dcn = 3; - + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); CV_Assert(src.channels() == 3 || src.channels() == 4); CV_Assert(dcn == 3 || dcn == 4); - dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } @@ -1297,7 +1297,7 @@ namespace void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream& stream) { typedef void (*func_t)(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream); - static const func_t funcs[] = + static const func_t funcs[] = { bgr_to_bgra, // CV_BGR2BGRA =0 bgra_to_bgr, // CV_BGRA2BGR =1 @@ -1353,7 +1353,7 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream 0, // =42 0, // =43 - 0, // CV_BGR2Lab =44 + 0, // CV_BGR2Lab =44 0, // CV_RGB2Lab =45 0, // CV_BayerBG2BGR =46 @@ -1374,7 +1374,7 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream 0, // CV_Lab2RGB =57 0, // CV_Luv2BGR =58 0, // CV_Luv2RGB =59 - + hls_to_bgr, // CV_HLS2BGR =60 hls_to_rgb, // CV_HLS2RGB =61 @@ -1436,7 +1436,10 @@ void cv::gpu::swapChannels(GpuMat& image, const int dstOrder[4], Stream& s) sz.width = image.cols; sz.height = image.rows; - nppSafeCall( nppiSwapChannels_8u_C4IR(image.ptr(), static_cast(image.step), sz, dstOrder) ); + nppSafeCall( nppiSwapChannels_8u_C4IR(image.ptr(), static_cast(image.step), sz, dstOrder) ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); } #endif /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/src/cuda/remap.cu b/modules/gpu/src/cuda/remap.cu index f77adee924..7eb2da3acc 100644 --- a/modules/gpu/src/cuda/remap.cu +++ b/modules/gpu/src/cuda/remap.cu @@ -47,10 +47,10 @@ #include "opencv2/gpu/device/saturate_cast.hpp" #include "opencv2/gpu/device/filters.hpp" -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace device { - namespace imgproc - { + namespace imgproc + { template __global__ void remap(const Ptr2D src, const PtrStepf mapx, const PtrStepf mapy, DevMem2D_ dst) { const int x = blockDim.x * blockIdx.x + threadIdx.x; @@ -67,11 +67,10 @@ namespace cv { namespace gpu { namespace device template