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