From 06ddccebed1b0ba69607f2e90326fff7a5f30306 Mon Sep 17 00:00:00 2001 From: cudawarped <12133430+cudawarped@users.noreply.github.com> Date: Wed, 23 Oct 2024 17:22:55 +0300 Subject: [PATCH 1/2] cudaarithm: fix python bindings for binary ops involving scalars --- .../cudaarithm/include/opencv2/cudaarithm.hpp | 230 ++++++++++++++++-- .../misc/python/test/test_cudaarithm.py | 35 +++ 2 files changed, 250 insertions(+), 15 deletions(-) diff --git a/modules/cudaarithm/include/opencv2/cudaarithm.hpp b/modules/cudaarithm/include/opencv2/cudaarithm.hpp index a16c271881e..e5ab2cf2575 100644 --- a/modules/cudaarithm/include/opencv2/cudaarithm.hpp +++ b/modules/cudaarithm/include/opencv2/cudaarithm.hpp @@ -75,61 +75,137 @@ namespace cv { namespace cuda { @param src1 First source matrix or scalar. @param src2 Second source matrix or scalar. Matrix should have the same size and type as src1 . @param dst Destination matrix that has the same size and number of channels as the input array(s). -The depth is defined by dtype or src1 depth. +The depth is defined by dtype or @p src1 depth. @param mask Optional operation mask, 8-bit single channel array, that specifies elements of the destination array to be changed. The mask can be used only with single channel images. @param dtype Optional depth of the output array. @param stream Stream for the asynchronous version. -@sa add +@warning In python both @p src1 and @p src2 have to be matrices, see @ref addWithScalar for scalar overload. + +@sa cv::add, addWithScalar */ CV_EXPORTS_W void add(InputArray src1, InputArray src2, OutputArray dst, InputArray mask = noArray(), int dtype = -1, Stream& stream = Stream::Null()); +/** @brief Computes a matrix-scalar sum. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and number of channels as the input array. +The depth is defined by dtype or @p src1 depth. +@param mask Optional operation mask, 8-bit single channel array, that specifies elements of the +destination array to be changed. The mask can be used only with single channel images. +@param dtype Optional depth of the output array. +@param stream Stream for the asynchronous version. + +@sa add + */ +CV_EXPORTS_W void inline addWithScalar(InputArray src1, Scalar src2, OutputArray dst, InputArray mask = noArray(), int dtype = -1, Stream& stream = Stream::Null()) { + add(src1, src2, dst, mask, dtype, stream); +} + /** @brief Computes a matrix-matrix or matrix-scalar difference. @param src1 First source matrix or scalar. -@param src2 Second source matrix or scalar. Matrix should have the same size and type as src1 . +@param src2 Second source matrix or scalar. Matrix should have the same size and type as @p src1. @param dst Destination matrix that has the same size and number of channels as the input array(s). -The depth is defined by dtype or src1 depth. +The depth is defined by dtype or @p src1 depth. @param mask Optional operation mask, 8-bit single channel array, that specifies elements of the destination array to be changed. The mask can be used only with single channel images. @param dtype Optional depth of the output array. @param stream Stream for the asynchronous version. -@sa subtract +@warning In python both @p src1 and @p src2 have to be matrices, see @ref subtractWithScalar for scalar overload. + +@sa cv::subtract, subtractWithScalar */ CV_EXPORTS_W void subtract(InputArray src1, InputArray src2, OutputArray dst, InputArray mask = noArray(), int dtype = -1, Stream& stream = Stream::Null()); +/** @brief Computes matrix-scalar difference. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and number of channels as the input array. +The depth is defined by dtype or @p src1 depth. +@param mask Optional operation mask, 8-bit single channel array, that specifies elements of the +destination array to be changed. The mask can be used only with single channel images. +@param dtype Optional depth of the output array. +@param stream Stream for the asynchronous version. + +@sa cv::subtract + */ +CV_EXPORTS_W void inline subtractWithScalar(InputArray src1, Scalar src2, OutputArray dst, InputArray mask = noArray(), int dtype = -1, Stream& stream = Stream::Null()) { + subtract(src1, src2, dst, mask, dtype, stream); +} + /** @brief Computes a matrix-matrix or matrix-scalar per-element product. @param src1 First source matrix or scalar. @param src2 Second source matrix or scalar. @param dst Destination matrix that has the same size and number of channels as the input array(s). -The depth is defined by dtype or src1 depth. +The depth is defined by dtype or @p src1 depth. @param scale Optional scale factor. @param dtype Optional depth of the output array. @param stream Stream for the asynchronous version. -@sa multiply +@warning In python both @p src1 and @p src2 have to be matrices, see @ref multiplyWithScalar for scalar overload. + +@sa cv::multiply, multiplyWithScalar */ CV_EXPORTS_W void multiply(InputArray src1, InputArray src2, OutputArray dst, double scale = 1, int dtype = -1, Stream& stream = Stream::Null()); +/** @brief Computes a matrix-scalar per-element product. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and number of channels as the input array. +The depth is defined by dtype or @p src1 depth. +@param scale Optional scale factor. +@param dtype Optional depth of the output array. +@param stream Stream for the asynchronous version. + +@sa multiply + */ +CV_EXPORTS_W void inline multiplyWithScalar(InputArray src1, Scalar src2, OutputArray dst, double scale = 1, int dtype = -1, Stream& stream = Stream::Null()) { + multiply(src1, src2, dst, scale, dtype, stream); +} + /** @brief Computes a matrix-matrix or matrix-scalar division. -@param src1 First source matrix or a scalar. +@param src1 First source matrix or scalar. @param src2 Second source matrix or scalar. @param dst Destination matrix that has the same size and number of channels as the input array(s). -The depth is defined by dtype or src1 depth. +The depth is defined by dtype or @p src1 depth. @param scale Optional scale factor. @param dtype Optional depth of the output array. @param stream Stream for the asynchronous version. This function, in contrast to divide, uses a round-down rounding mode. -@sa divide +@warning In python both @p src1 and @p src2 have to be matrices, see @ref divideWithScalar for scalar overload. + +@sa cv::divide, divideWithScalar */ CV_EXPORTS_W void divide(InputArray src1, InputArray src2, OutputArray dst, double scale = 1, int dtype = -1, Stream& stream = Stream::Null()); +/** @brief Computes a matrix-scalar division. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and number of channels as the input array. +The depth is defined by dtype or @p src1 depth. +@param scale Optional scale factor. +@param dtype Optional depth of the output array. +@param stream Stream for the asynchronous version. + +This function, in contrast to divide, uses a round-down rounding mode. + +@sa divide + */ +CV_EXPORTS_W void inline divideWithScalar(InputArray src1, Scalar src2, OutputArray dst, double scale = 1, int dtype = -1, Stream& stream = Stream::Null()) { + divide(src1, src2, dst, scale, dtype, stream); +} + /** @brief Computes per-element absolute difference of two matrices (or of a matrix and scalar). @param src1 First source matrix or scalar. @@ -137,10 +213,25 @@ CV_EXPORTS_W void divide(InputArray src1, InputArray src2, OutputArray dst, doub @param dst Destination matrix that has the same size and type as the input array(s). @param stream Stream for the asynchronous version. -@sa absdiff +@warning In python both @p src1 and @p src2 have to be matrices, see @ref absdiffWithScalar for scalar overload. + +@sa cv::absdiff, absdiffWithScalar */ CV_EXPORTS_W void absdiff(InputArray src1, InputArray src2, OutputArray dst, Stream& stream = Stream::Null()); +/** @brief Computes per-element absolute difference of a matrix and scalar. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and type as the input array. +@param stream Stream for the asynchronous version. + +@sa absdiff + */ +CV_EXPORTS_W void inline absdiffWithScalar(InputArray src1, Scalar src2, OutputArray dst, Stream& stream = Stream::Null()) { + absdiff(src1, src2, dst, stream); +} + /** @brief Computes an absolute value of each matrix element. @param src Source matrix. @@ -218,10 +309,32 @@ CV_EXPORTS_W void pow(InputArray src, double power, OutputArray dst, Stream& str - **CMP_NE:** a(.) != b(.) @param stream Stream for the asynchronous version. -@sa compare +@warning In python both @p src1 and @p src2 have to be matrices, see @ref compareWithScalar for scalar overload. + +@sa cv::compare, compareWithScalar */ CV_EXPORTS_W void compare(InputArray src1, InputArray src2, OutputArray dst, int cmpop, Stream& stream = Stream::Null()); +/** @brief Compares elements of a matrix and scalar. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size as the input array and type \ref CV_8U. +@param cmpop Flag specifying the relation between the elements to be checked: +- **CMP_EQ:** a(.) == b(.) +- **CMP_GT:** a(.) \> b(.) +- **CMP_GE:** a(.) \>= b(.) +- **CMP_LT:** a(.) \< b(.) +- **CMP_LE:** a(.) \<= b(.) +- **CMP_NE:** a(.) != b(.) +@param stream Stream for the asynchronous version. + +@sa compare + */ +CV_EXPORTS_W void inline compareWithScalar(InputArray src1, Scalar src2, OutputArray dst, int cmpop, Stream& stream = Stream::Null()) { + compare(src1, src2, dst, cmpop, stream); +} + /** @brief Performs a per-element bitwise inversion. @param src Source matrix. @@ -240,9 +353,28 @@ CV_EXPORTS_W void bitwise_not(InputArray src, OutputArray dst, InputArray mask = @param mask Optional operation mask, 8-bit single channel array, that specifies elements of the destination array to be changed. The mask can be used only with single channel images. @param stream Stream for the asynchronous version. + +@warning In python both @p src1 and @p src2 have to be matrices, see @ref bitwise_or_with_scalar for scalar overload. + +@sa cv::bitwise_or, bitwise_or_with_scalar */ CV_EXPORTS_W void bitwise_or(InputArray src1, InputArray src2, OutputArray dst, InputArray mask = noArray(), Stream& stream = Stream::Null()); +/** @brief Performs a per-element bitwise disjunction of a matrix and scalar. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and type as the input array. +@param mask Optional operation mask, 8-bit single channel array, that specifies elements of the +destination array to be changed. The mask can be used only with single channel images. +@param stream Stream for the asynchronous version. + +@sa bitwise_or + */ +CV_EXPORTS_W void inline bitwise_or_with_scalar(InputArray src1, Scalar src2, OutputArray dst, InputArray mask = noArray(), Stream& stream = Stream::Null()) { + bitwise_or(src1, src2, dst, mask, stream); +} + /** @brief Performs a per-element bitwise conjunction of two matrices (or of matrix and scalar). @param src1 First source matrix or scalar. @@ -251,20 +383,58 @@ CV_EXPORTS_W void bitwise_or(InputArray src1, InputArray src2, OutputArray dst, @param mask Optional operation mask, 8-bit single channel array, that specifies elements of the destination array to be changed. The mask can be used only with single channel images. @param stream Stream for the asynchronous version. + +@warning In python both @p src1 and @p src2 have to be matrices, see @ref bitwise_and_with_scalar for scalar overload. + +@sa bitwise_and_with_scalar */ CV_EXPORTS_W void bitwise_and(InputArray src1, InputArray src2, OutputArray dst, InputArray mask = noArray(), Stream& stream = Stream::Null()); +/** @brief Performs a per-element bitwise conjunction of a matrix and a scalar. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and type as the input array. +@param mask Optional operation mask, 8-bit single channel array, that specifies elements of the +destination array to be changed. The mask can be used only with single channel images. +@param stream Stream for the asynchronous version. + +@sa bitwise_and + */ +CV_EXPORTS_W void inline bitwise_and_with_scalar(InputArray src1, Scalar src2, OutputArray dst, InputArray mask = noArray(), Stream& stream = Stream::Null()) { + bitwise_and(src1, src2, dst, mask, stream); +} + /** @brief Performs a per-element bitwise exclusive or operation of two matrices (or of matrix and scalar). @param src1 First source matrix or scalar. @param src2 Second source matrix or scalar. -@param dst Destination matrix that has the same size and type as the input array(s). +@param dst Destination matrix that has the same size and type as the input array. @param mask Optional operation mask, 8-bit single channel array, that specifies elements of the destination array to be changed. The mask can be used only with single channel images. @param stream Stream for the asynchronous version. + +@warning In python both @p src1 and @p src2 have to be matrices, see @ref bitwise_xor_with_scalar for scalar overload. + +@sa cv::bitwise_xor, bitwise_xor_with_scalar */ CV_EXPORTS_W void bitwise_xor(InputArray src1, InputArray src2, OutputArray dst, InputArray mask = noArray(), Stream& stream = Stream::Null()); +/** @brief Performs a per-element bitwise exclusive or operation of a matrix and a scalar. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and type as the input array(s). +@param mask Optional operation mask, 8-bit single channel array, that specifies elements of the +destination array to be changed. The mask can be used only with single channel images. +@param stream Stream for the asynchronous version. + +@sa bitwise_xor + */ +CV_EXPORTS_W void inline bitwise_xor_with_scalar(InputArray src1, Scalar src2, OutputArray dst, InputArray mask = noArray(), Stream& stream = Stream::Null()) { + bitwise_xor(src1, src2, dst, mask, stream); +} + /** @brief Performs pixel by pixel right shift of an image by a constant value. @param src Source matrix. Supports 1, 3 and 4 channels images with integers elements. @@ -299,10 +469,25 @@ CV_WRAP inline void lshift(InputArray src, Scalar val, OutputArray dst, Stream& @param dst Destination matrix that has the same size and type as the input array(s). @param stream Stream for the asynchronous version. -@sa min +@warning In python both @p src1 and @p src2 have to be matrices, see @ref minWithScalar for scalar overload. + +@sa cv::min, minWithScalar */ CV_EXPORTS_W void min(InputArray src1, InputArray src2, OutputArray dst, Stream& stream = Stream::Null()); +/** @brief Computes the per-element minimum or a matrix and a scalar. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and type as the input array. +@param stream Stream for the asynchronous version. + +@sa min + */ +CV_EXPORTS_W void inline minWithScalar(InputArray src1, Scalar src2, OutputArray dst, Stream& stream = Stream::Null()) { + min(src1, src2, dst, stream); +} + /** @brief Computes the per-element maximum of two matrices (or a matrix and a scalar). @param src1 First source matrix or scalar. @@ -310,10 +495,25 @@ CV_EXPORTS_W void min(InputArray src1, InputArray src2, OutputArray dst, Stream& @param dst Destination matrix that has the same size and type as the input array(s). @param stream Stream for the asynchronous version. -@sa max +@warning In python both @p src1 and @p src2 have to be matrices, see @ref maxWithScalar for scalar overload. + +@sa cv::max, maxWithScalar */ CV_EXPORTS_W void max(InputArray src1, InputArray src2, OutputArray dst, Stream& stream = Stream::Null()); +/** @brief Computes the per-element maximum of a matrix and a scalar. + +@param src1 First source matrix. +@param src2 Second source scalar. +@param dst Destination matrix that has the same size and type as the input array. +@param stream Stream for the asynchronous version. + +@sa max + */ +CV_EXPORTS_W void inline maxWithScalar(InputArray src1, Scalar src2, OutputArray dst, Stream& stream = Stream::Null()) { + max(src1, src2, dst, stream); +} + /** @brief Computes the weighted sum of two arrays. @param src1 First source array. diff --git a/modules/cudaarithm/misc/python/test/test_cudaarithm.py b/modules/cudaarithm/misc/python/test/test_cudaarithm.py index 1d959d791f3..897d5961dea 100644 --- a/modules/cudaarithm/misc/python/test/test_cudaarithm.py +++ b/modules/cudaarithm/misc/python/test/test_cudaarithm.py @@ -38,6 +38,7 @@ def test_cudaarithm(self): def test_arithmetic(self): npMat1 = np.random.random((128, 128, 3)) - 0.5 npMat2 = np.random.random((128, 128, 3)) - 0.5 + scalar = np.random.random() cuMat1 = cv.cuda_GpuMat() cuMat2 = cv.cuda_GpuMat() @@ -48,36 +49,54 @@ def test_arithmetic(self): self.assertTrue(np.allclose(cv.cuda.add(cuMat1, cuMat2).download(), cv.add(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.addWithScalar(cuMat1, [scalar]*3).download(), + cv.add(npMat1, scalar))) + cv.cuda.add(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.add(npMat1, npMat2))) self.assertTrue(np.allclose(cv.cuda.subtract(cuMat1, cuMat2).download(), cv.subtract(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.subtractWithScalar(cuMat1, [scalar]*3).download(), + cv.subtract(npMat1, scalar))) + cv.cuda.subtract(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.subtract(npMat1, npMat2))) self.assertTrue(np.allclose(cv.cuda.multiply(cuMat1, cuMat2).download(), cv.multiply(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.multiplyWithScalar(cuMat1, [scalar]*3).download(), + cv.multiply(npMat1, scalar))) + cv.cuda.multiply(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.multiply(npMat1, npMat2))) self.assertTrue(np.allclose(cv.cuda.divide(cuMat1, cuMat2).download(), cv.divide(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.divideWithScalar(cuMat1, [scalar]*3).download(), + cv.divide(npMat1, scalar))) + cv.cuda.divide(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.divide(npMat1, npMat2))) self.assertTrue(np.allclose(cv.cuda.absdiff(cuMat1, cuMat2).download(), cv.absdiff(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.absdiffWithScalar(cuMat1, [scalar]*3).download(), + cv.absdiff(npMat1, scalar))) + cv.cuda.absdiff(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.absdiff(npMat1, npMat2))) self.assertTrue(np.allclose(cv.cuda.compare(cuMat1, cuMat2, cv.CMP_GE).download(), cv.compare(npMat1, npMat2, cv.CMP_GE))) + self.assertTrue(np.allclose(cv.cuda.compareWithScalar(cuMat1, [scalar]*3, cv.CMP_GE).download(), + cv.compare(npMat1, scalar, cv.CMP_GE))) + cuMatDst1 = cv.cuda_GpuMat(cuMat1.size(),cv.CV_8UC3) cv.cuda.compare(cuMat1, cuMat2, cv.CMP_GE, cuMatDst1) self.assertTrue(np.allclose(cuMatDst1.download(),cv.compare(npMat1, npMat2, cv.CMP_GE))) @@ -111,6 +130,7 @@ def test_arithmetic(self): def test_logical(self): npMat1 = (np.random.random((128, 128)) * 255).astype(np.uint8) npMat2 = (np.random.random((128, 128)) * 255).astype(np.uint8) + scalar = np.random.random() cuMat1 = cv.cuda_GpuMat() cuMat2 = cv.cuda_GpuMat() @@ -121,18 +141,27 @@ def test_logical(self): self.assertTrue(np.allclose(cv.cuda.bitwise_or(cuMat1, cuMat2).download(), cv.bitwise_or(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.bitwise_or_with_scalar(cuMat1, scalar).download(), + cv.bitwise_or(npMat1, scalar))) + cv.cuda.bitwise_or(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.bitwise_or(npMat1, npMat2))) self.assertTrue(np.allclose(cv.cuda.bitwise_and(cuMat1, cuMat2).download(), cv.bitwise_and(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.bitwise_and_with_scalar(cuMat1, scalar).download(), + cv.bitwise_and(npMat1, scalar))) + cv.cuda.bitwise_and(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.bitwise_and(npMat1, npMat2))) self.assertTrue(np.allclose(cv.cuda.bitwise_xor(cuMat1, cuMat2).download(), cv.bitwise_xor(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.bitwise_xor_with_scalar(cuMat1, scalar).download(), + cv.bitwise_xor(npMat1, scalar))) + cv.cuda.bitwise_xor(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.bitwise_xor(npMat1, npMat2))) @@ -145,12 +174,18 @@ def test_logical(self): self.assertTrue(np.allclose(cv.cuda.min(cuMat1, cuMat2).download(), cv.min(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.minWithScalar(cuMat1, scalar).download(), + cv.min(npMat1, scalar))) + cv.cuda.min(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.min(npMat1, npMat2))) self.assertTrue(np.allclose(cv.cuda.max(cuMat1, cuMat2).download(), cv.max(npMat1, npMat2))) + self.assertTrue(np.allclose(cv.cuda.maxWithScalar(cuMat1, scalar).download(), + cv.max(npMat1, scalar))) + cv.cuda.max(cuMat1, cuMat2, cuMatDst) self.assertTrue(np.allclose(cuMatDst.download(),cv.max(npMat1, npMat2))) From 843b6ede6a9c438f35bf35e1460d55e3dbd1af17 Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Tue, 29 Oct 2024 15:31:10 +0100 Subject: [PATCH 2/2] Merge pull request #3800 from vrabaud:cuda Get CUDA code to compile with clang CUDA and without CUDA #3800 Changelist: - there are some syntactic changes: `<< <` -> `<<<`. For some reason, I do not need to change all those in the code. - `::min` -> `std::min` in `__host__` code - `modules/cudaimgproc/src/moments.cpp` needs to have the CUDA code in the `#ifdef` - The signature of `cv::cuda::swapChannels` is not exactly the same as the C++ one in `modules/cudaimgproc/src/color.cpp` - `cv::cuda::FarnebackOpticalFlow::create` needs to be explicit about which FarnebackOpticalFlow it returns ### Pull Request Readiness Checklist See details at https://github.com/opencv/opencv/wiki/How_to_contribute#making-a-good-pull-request - [x] I agree to contribute to the project under Apache 2 License. - [x] To the best of my knowledge, the proposed patch is not based on a code under GPL or another license that is incompatible with OpenCV - [x] The PR is proposed to the proper branch --- modules/cudaarithm/src/arithm.cpp | 2 ++ modules/cudaarithm/src/cuda/polar_cart.cu | 10 +++++----- modules/cudaarithm/src/element_operations.cpp | 5 +++++ modules/cudaarithm/src/reductions.cpp | 4 +++- modules/cudacodec/src/cuda/nv12_to_rgb.cu | 4 ++-- modules/cudaimgproc/src/color.cpp | 2 +- .../cudaimgproc/src/connectedcomponents.cpp | 4 ++-- modules/cudaimgproc/src/cuda/canny.cu | 4 ++-- .../src/cuda/connectedcomponents.cu | 10 +++++----- .../cudaimgproc/src/cuda/generalized_hough.cu | 4 ++-- modules/cudaimgproc/src/cuda/hough_circles.cu | 2 +- modules/cudaimgproc/src/cuda/hough_lines.cu | 2 +- modules/cudaimgproc/src/cuda/hough_segments.cu | 2 +- modules/cudaimgproc/src/cuda/moments.cu | 6 +++--- modules/cudaimgproc/src/histogram.cpp | 1 + modules/cudaimgproc/src/moments.cpp | 14 ++++++++------ .../cudaoptflow/src/cuda/nvidiaOpticalFlow.cu | 2 +- modules/cudaoptflow/src/farneback.cpp | 2 +- modules/cudaoptflow/src/precomp.hpp | 2 ++ .../opencv2/cudev/grid/detail/minmaxloc.hpp | 4 ++-- modules/hfs/src/cuda/gslic_seg_engine_gpu.cu | 18 +++++++++--------- modules/hfs/src/cuda/magnitude.cu | 4 ++-- 22 files changed, 61 insertions(+), 47 deletions(-) diff --git a/modules/cudaarithm/src/arithm.cpp b/modules/cudaarithm/src/arithm.cpp index 30cf225e154..6d2a145d610 100644 --- a/modules/cudaarithm/src/arithm.cpp +++ b/modules/cudaarithm/src/arithm.cpp @@ -54,6 +54,8 @@ void cv::cuda::mulAndScaleSpectrums(InputArray, InputArray, OutputArray, int, fl void cv::cuda::dft(InputArray, OutputArray, Size, int, Stream&) { throw_no_cuda(); } +Ptr cv::cuda::createDFT(Size, int) { throw_no_cuda(); return Ptr(); } + Ptr cv::cuda::createConvolution(Size) { throw_no_cuda(); return Ptr(); } #else /* !defined (HAVE_CUDA) */ diff --git a/modules/cudaarithm/src/cuda/polar_cart.cu b/modules/cudaarithm/src/cuda/polar_cart.cu index c65b894bf60..e263ce34881 100644 --- a/modules/cudaarithm/src/cuda/polar_cart.cu +++ b/modules/cudaarithm/src/cuda/polar_cart.cu @@ -289,9 +289,9 @@ namespace const T scale = angleInDegrees ? static_cast(CV_PI / 180.0) : static_cast(1.0); if (mag.empty()) - polarToCartImpl_ << > >(mag, angle, x, y, scale); + polarToCartImpl_ <<>>(mag, angle, x, y, scale); else - polarToCartImpl_ << > >(mag, angle, x, y, scale); + polarToCartImpl_ <<>>(mag, angle, x, y, scale); } template @@ -305,9 +305,9 @@ namespace const T scale = angleInDegrees ? static_cast(CV_PI / 180.0) : static_cast(1.0); if (mag.empty()) - polarToCartDstInterleavedImpl_ << > >(mag, angle, xy, scale); + polarToCartDstInterleavedImpl_ <<>>(mag, angle, xy, scale); else - polarToCartDstInterleavedImpl_ << > >(mag, angle, xy, scale); + polarToCartDstInterleavedImpl_ <<>>(mag, angle, xy, scale); } template @@ -320,7 +320,7 @@ namespace const T scale = angleInDegrees ? static_cast(CV_PI / 180.0) : static_cast(1.0); - polarToCartInterleavedImpl_ << > >(magAngle, xy, scale); + polarToCartInterleavedImpl_ <<>>(magAngle, xy, scale); } } diff --git a/modules/cudaarithm/src/element_operations.cpp b/modules/cudaarithm/src/element_operations.cpp index 1ad3c17c40f..6f810357b13 100644 --- a/modules/cudaarithm/src/element_operations.cpp +++ b/modules/cudaarithm/src/element_operations.cpp @@ -84,8 +84,13 @@ void cv::cuda::magnitude(InputArray, InputArray, OutputArray, Stream&) { throw_n void cv::cuda::magnitudeSqr(InputArray, OutputArray, Stream&) { throw_no_cuda(); } void cv::cuda::magnitudeSqr(InputArray, InputArray, OutputArray, Stream&) { throw_no_cuda(); } void cv::cuda::phase(InputArray, InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } +void cv::cuda::phase(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } void cv::cuda::cartToPolar(InputArray, InputArray, OutputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } +void cv::cuda::cartToPolar(InputArray, OutputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } +void cv::cuda::cartToPolar(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } void cv::cuda::polarToCart(InputArray, InputArray, OutputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } +void cv::cuda::polarToCart(InputArray, InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } +void cv::cuda::polarToCart(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } #else diff --git a/modules/cudaarithm/src/reductions.cpp b/modules/cudaarithm/src/reductions.cpp index b70a128558f..0ceb3d26a14 100644 --- a/modules/cudaarithm/src/reductions.cpp +++ b/modules/cudaarithm/src/reductions.cpp @@ -69,8 +69,10 @@ void cv::cuda::countNonZero(InputArray, OutputArray, Stream&) { throw_no_cuda(); void cv::cuda::reduce(InputArray, OutputArray, int, int, int, Stream&) { throw_no_cuda(); } -void cv::cuda::meanStdDev(InputArray, Scalar&, Scalar&) { throw_no_cuda(); } +void cv::cuda::meanStdDev(InputArray, OutputArray, InputArray, Stream&) { throw_no_cuda(); } void cv::cuda::meanStdDev(InputArray, OutputArray, Stream&) { throw_no_cuda(); } +void cv::cuda::meanStdDev(InputArray, Scalar&, Scalar&, InputArray) { throw_no_cuda(); } +void cv::cuda::meanStdDev(InputArray, Scalar&, Scalar&) { throw_no_cuda(); } void cv::cuda::rectStdDev(InputArray, InputArray, OutputArray, Rect, Stream&) { throw_no_cuda(); } diff --git a/modules/cudacodec/src/cuda/nv12_to_rgb.cu b/modules/cudacodec/src/cuda/nv12_to_rgb.cu index 049041a7588..a9031e0ec9e 100644 --- a/modules/cudacodec/src/cuda/nv12_to_rgb.cu +++ b/modules/cudacodec/src/cuda/nv12_to_rgb.cu @@ -179,9 +179,9 @@ void nv12ToBgra(const GpuMat& decodedFrame, GpuMat& outFrame, int width, int hei dim3 block(32, 8); dim3 grid(divUp(width, 2 * block.x), divUp(height, block.y)); if (videoFullRangeFlag) - NV12_to_BGRA << > > (decodedFrame.ptr(), decodedFrame.step, outFrame.ptr(), outFrame.step, width, height); + NV12_to_BGRA <<>> (decodedFrame.ptr(), decodedFrame.step, outFrame.ptr(), outFrame.step, width, height); else - NV12_to_BGRA << > > (decodedFrame.ptr(), decodedFrame.step, outFrame.ptr(), outFrame.step, width, height); + NV12_to_BGRA <<>> (decodedFrame.ptr(), decodedFrame.step, outFrame.ptr(), outFrame.step, width, height); CV_CUDEV_SAFE_CALL(cudaGetLastError()); if (stream == 0) CV_CUDEV_SAFE_CALL(cudaDeviceSynchronize()); diff --git a/modules/cudaimgproc/src/color.cpp b/modules/cudaimgproc/src/color.cpp index 5adfa6cda6c..2393a50b5a8 100644 --- a/modules/cudaimgproc/src/color.cpp +++ b/modules/cudaimgproc/src/color.cpp @@ -51,7 +51,7 @@ void cv::cuda::cvtColor(InputArray, OutputArray, int, int, Stream&) { throw_no_c void cv::cuda::demosaicing(InputArray, OutputArray, int, int, Stream&) { throw_no_cuda(); } -void cv::cuda::swapChannels(InputOutputArray, const int[], Stream&) { throw_no_cuda(); } +void cv::cuda::swapChannels(InputOutputArray, const int[4], Stream&) { throw_no_cuda(); } void cv::cuda::gammaCorrection(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } diff --git a/modules/cudaimgproc/src/connectedcomponents.cpp b/modules/cudaimgproc/src/connectedcomponents.cpp index e2e2bde057d..3f6ac944491 100644 --- a/modules/cudaimgproc/src/connectedcomponents.cpp +++ b/modules/cudaimgproc/src/connectedcomponents.cpp @@ -9,8 +9,8 @@ using namespace cv::cuda; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) -void cv::cuda::connectedComponents(InputArray img_, OutputArray labels_, int connectivity, - int ltype, ConnectedComponentsAlgorithmsTypes ccltype) { throw_no_cuda(); } +void cv::cuda::connectedComponents(InputArray, OutputArray, int, int, ConnectedComponentsAlgorithmsTypes) { throw_no_cuda(); } +void cv::cuda::connectedComponents(InputArray, OutputArray, int, int) { throw_no_cuda(); } #else /* !defined (HAVE_CUDA) */ diff --git a/modules/cudaimgproc/src/cuda/canny.cu b/modules/cudaimgproc/src/cuda/canny.cu index 46b9624fb68..4c6de962a00 100644 --- a/modules/cudaimgproc/src/cuda/canny.cu +++ b/modules/cudaimgproc/src/cuda/canny.cu @@ -428,7 +428,7 @@ namespace canny cudaSafeCall( cudaMemsetAsync(d_counter, 0, sizeof(int), stream) ); const dim3 block(128); - const dim3 grid(::min(count, 65535u), divUp(count, 65535), 1); + const dim3 grid(std::min(count, 65535), divUp(count, 65535), 1); edgesHysteresisGlobalKernel<<>>(map, st1, st2, d_counter, count); cudaSafeCall( cudaGetLastError() ); @@ -439,7 +439,7 @@ namespace canny cudaSafeCall( cudaMemcpyAsync(&count, d_counter, sizeof(int), cudaMemcpyDeviceToHost, stream) ); cudaSafeCall( cudaStreamSynchronize(stream) ); - count = min(count, map.cols * map.rows); + count = std::min(count, map.cols * map.rows); //std::swap(st1, st2); short2* tmp = st1; diff --git a/modules/cudaimgproc/src/cuda/connectedcomponents.cu b/modules/cudaimgproc/src/cuda/connectedcomponents.cu index 4946fa26276..d6ce49f0941 100644 --- a/modules/cudaimgproc/src/cuda/connectedcomponents.cu +++ b/modules/cudaimgproc/src/cuda/connectedcomponents.cu @@ -317,19 +317,19 @@ void BlockBasedKomuraEquivalence(const cv::cuda::GpuMat& img, cv::cuda::GpuMat& grid_size = dim3((((img.cols + 1) / 2) - 1) / kblock_cols + 1, (((img.rows + 1) / 2) - 1) / kblock_rows + 1, 1); block_size = dim3(kblock_cols, kblock_rows, 1); - InitLabeling << > > (img, labels, last_pixel); + InitLabeling <<>> (img, labels, last_pixel); cudaSafeCall(cudaGetLastError()); - Compression << > > (labels); + Compression <<>> (labels); cudaSafeCall(cudaGetLastError()); - Merge << > > (labels, last_pixel); + Merge <<>> (labels, last_pixel); cudaSafeCall(cudaGetLastError()); - Compression << > > (labels); + Compression <<>> (labels); cudaSafeCall(cudaGetLastError()); - FinalLabeling << > > (img, labels); + FinalLabeling <<>> (img, labels); cudaSafeCall(cudaGetLastError()); if (last_pixel_allocated) { diff --git a/modules/cudaimgproc/src/cuda/generalized_hough.cu b/modules/cudaimgproc/src/cuda/generalized_hough.cu index 232c625f9f8..150c502bc5b 100644 --- a/modules/cudaimgproc/src/cuda/generalized_hough.cu +++ b/modules/cudaimgproc/src/cuda/generalized_hough.cu @@ -302,7 +302,7 @@ namespace cv { namespace cuda { namespace device int totalCount; cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); - totalCount = ::min(totalCount, maxSize); + totalCount = std::min(totalCount, maxSize); return totalCount; } @@ -812,7 +812,7 @@ namespace cv { namespace cuda { namespace device int totalCount; cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); - totalCount = ::min(totalCount, maxSize); + totalCount = std::min(totalCount, maxSize); return totalCount; } diff --git a/modules/cudaimgproc/src/cuda/hough_circles.cu b/modules/cudaimgproc/src/cuda/hough_circles.cu index fcbf8c4122b..da122cd8d18 100644 --- a/modules/cudaimgproc/src/cuda/hough_circles.cu +++ b/modules/cudaimgproc/src/cuda/hough_circles.cu @@ -238,7 +238,7 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaMemcpyAsync(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost, stream) ); cudaSafeCall( cudaStreamSynchronize(stream) ); - totalCount = ::min(totalCount, maxCircles); + totalCount = std::min(totalCount, maxCircles); return totalCount; } diff --git a/modules/cudaimgproc/src/cuda/hough_lines.cu b/modules/cudaimgproc/src/cuda/hough_lines.cu index 6fa55791107..7c46b594e13 100644 --- a/modules/cudaimgproc/src/cuda/hough_lines.cu +++ b/modules/cudaimgproc/src/cuda/hough_lines.cu @@ -189,7 +189,7 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaStreamSynchronize(stream) ); - totalCount = ::min(totalCount, maxSize); + totalCount = std::min(totalCount, maxSize); if (doSort && totalCount > 0) { diff --git a/modules/cudaimgproc/src/cuda/hough_segments.cu b/modules/cudaimgproc/src/cuda/hough_segments.cu index 2ad7f1f1da5..50da12f0d56 100644 --- a/modules/cudaimgproc/src/cuda/hough_segments.cu +++ b/modules/cudaimgproc/src/cuda/hough_segments.cu @@ -241,7 +241,7 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaStreamSynchronize(stream) ); - totalCount = ::min(totalCount, maxSize); + totalCount = std::min(totalCount, maxSize); return totalCount; } } diff --git a/modules/cudaimgproc/src/cuda/moments.cu b/modules/cudaimgproc/src/cuda/moments.cu index daf479a75f7..88a3d86692f 100644 --- a/modules/cudaimgproc/src/cuda/moments.cu +++ b/modules/cudaimgproc/src/cuda/moments.cu @@ -139,7 +139,7 @@ template struct momentsDispatch static void call(const PtrStepSz src, PtrStepSz moments, const bool binary, const int offsetX, const cudaStream_t stream) { dim3 blockSize(blockSizeX, blockSizeY); dim3 gridSize = dim3(divUp(src.rows, blockSizeY)); - spatialMoments << > > (src, binary, moments.ptr()); + spatialMoments <<>> (src, binary, moments.ptr()); if (stream == 0) cudaSafeCall(cudaStreamSynchronize(stream)); }; @@ -150,9 +150,9 @@ template struct momentsDispatcherChar { dim3 blockSize(blockSizeX, blockSizeY); dim3 gridSize = dim3(divUp(src.rows, blockSizeY)); if (offsetX) - spatialMoments << > > (src, binary, moments.ptr(), offsetX); + spatialMoments <<>> (src, binary, moments.ptr(), offsetX); else - spatialMoments << > > (src, binary, moments.ptr()); + spatialMoments <<>> (src, binary, moments.ptr()); if (stream == 0) cudaSafeCall(cudaStreamSynchronize(stream)); diff --git a/modules/cudaimgproc/src/histogram.cpp b/modules/cudaimgproc/src/histogram.cpp index 51a5ce1a83e..88a1305b117 100644 --- a/modules/cudaimgproc/src/histogram.cpp +++ b/modules/cudaimgproc/src/histogram.cpp @@ -48,6 +48,7 @@ using namespace cv::cuda; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) void cv::cuda::calcHist(InputArray, OutputArray, Stream&) { throw_no_cuda(); } +void cv::cuda::calcHist(InputArray, InputArray, OutputArray, Stream&) { throw_no_cuda(); } void cv::cuda::equalizeHist(InputArray, OutputArray, Stream&) { throw_no_cuda(); } diff --git a/modules/cudaimgproc/src/moments.cpp b/modules/cudaimgproc/src/moments.cpp index 3c2e62c4b90..40ab9414ad7 100644 --- a/modules/cudaimgproc/src/moments.cpp +++ b/modules/cudaimgproc/src/moments.cpp @@ -3,15 +3,10 @@ // of this distribution and at http://opencv.org/license.html. #include "precomp.hpp" -#include "cuda/moments.cuh" using namespace cv; using namespace cv::cuda; -int cv::cuda::numMoments(const MomentsOrder order) { - return order == MomentsOrder::FIRST_ORDER_MOMENTS ? device::imgproc::n1 : order == MomentsOrder::SECOND_ORDER_MOMENTS ? device::imgproc::n12 : device::imgproc::n123; -} - template cv::Moments convertSpatialMomentsT(Mat spatialMoments, const MomentsOrder order) { switch (order) { @@ -32,10 +27,17 @@ cv::Moments cv::cuda::convertSpatialMoments(Mat spatialMoments, const MomentsOrd } #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) + int cv::cuda::numMoments(MomentsOrder) { throw_no_cuda(); return 0; } Moments cv::cuda::moments(InputArray src, const bool binary, const MomentsOrder order, const int momentsType) { throw_no_cuda(); } - void spatialMoments(InputArray src, OutputArray moments, const bool binary, const MomentsOrder order, const int momentsType, Stream& stream) { throw_no_cuda(); } + void cv::cuda::spatialMoments(InputArray src, OutputArray moments, const bool binary, const MomentsOrder order, const int momentsType, Stream& stream) { throw_no_cuda(); } #else /* !defined (HAVE_CUDA) */ +#include "cuda/moments.cuh" + +int cv::cuda::numMoments(const MomentsOrder order) { + return order == MomentsOrder::FIRST_ORDER_MOMENTS ? device::imgproc::n1 : order == MomentsOrder::SECOND_ORDER_MOMENTS ? device::imgproc::n12 : device::imgproc::n123; +} + namespace cv { namespace cuda { namespace device { namespace imgproc { template void moments(const PtrStepSzb src, PtrStepSzb moments, const bool binary, const int order, const int offsetX, const cudaStream_t stream); diff --git a/modules/cudaoptflow/src/cuda/nvidiaOpticalFlow.cu b/modules/cudaoptflow/src/cuda/nvidiaOpticalFlow.cu index e73eb430081..5395360b187 100644 --- a/modules/cudaoptflow/src/cuda/nvidiaOpticalFlow.cu +++ b/modules/cudaoptflow/src/cuda/nvidiaOpticalFlow.cu @@ -90,7 +90,7 @@ void FlowUpsample(void* srcDevPtr, uint32_t nSrcWidth, uint32_t nSrcPitch, uint3 dim3 blockDim(BLOCKDIM_X, BLOCKDIM_Y); dim3 gridDim((nDstWidth + blockDim.x - 1) / blockDim.x, (nDstHeight + blockDim.y - 1) / blockDim.y); - NearestNeighborFlowKernel << > > (0, srcDevPtr, nSrcWidth, nSrcPitch, nSrcHeight, + NearestNeighborFlowKernel <<>> (0, srcDevPtr, nSrcWidth, nSrcPitch, nSrcHeight, 0, dstDevPtr, nDstWidth, nDstPitch, nDstHeight, nScaleFactor); diff --git a/modules/cudaoptflow/src/farneback.cpp b/modules/cudaoptflow/src/farneback.cpp index eb82d0c34e4..f99a4c978f6 100644 --- a/modules/cudaoptflow/src/farneback.cpp +++ b/modules/cudaoptflow/src/farneback.cpp @@ -47,7 +47,7 @@ using namespace cv::cuda; #if !defined HAVE_CUDA || defined(CUDA_DISABLER) -Ptr cv::cuda::FarnebackOpticalFlow::create(int, double, bool, int, int, int, double, int) { throw_no_cuda(); return Ptr(); } +Ptr cv::cuda::FarnebackOpticalFlow::create(int, double, bool, int, int, int, double, int) { throw_no_cuda(); return Ptr(); } #else diff --git a/modules/cudaoptflow/src/precomp.hpp b/modules/cudaoptflow/src/precomp.hpp index d5ac493342d..c29b8b130a5 100644 --- a/modules/cudaoptflow/src/precomp.hpp +++ b/modules/cudaoptflow/src/precomp.hpp @@ -52,7 +52,9 @@ #include "opencv2/video.hpp" #include "opencv2/core/private.cuda.hpp" +#if defined HAVE_CUDA #include "opencv2/core/cuda/vec_traits.hpp" +#endif #include "opencv2/opencv_modules.hpp" #ifdef HAVE_OPENCV_CUDALEGACY diff --git a/modules/cudev/include/opencv2/cudev/grid/detail/minmaxloc.hpp b/modules/cudev/include/opencv2/cudev/grid/detail/minmaxloc.hpp index 9e4f348e1a8..13f0ebf2513 100644 --- a/modules/cudev/include/opencv2/cudev/grid/detail/minmaxloc.hpp +++ b/modules/cudev/include/opencv2/cudev/grid/detail/minmaxloc.hpp @@ -148,8 +148,8 @@ namespace grid_minmaxloc_detail block = dim3(Policy::block_size_x, Policy::block_size_y); grid = dim3(divUp(cols, block.x * Policy::patch_size_x), divUp(rows, block.y * Policy::patch_size_y)); - grid.x = ::min(grid.x, block.x); - grid.y = ::min(grid.y, block.y); + grid.x = std::min(grid.x, block.x); + grid.y = std::min(grid.y, block.y); } template diff --git a/modules/hfs/src/cuda/gslic_seg_engine_gpu.cu b/modules/hfs/src/cuda/gslic_seg_engine_gpu.cu index 69c054cdd96..fee4412b48d 100644 --- a/modules/hfs/src/cuda/gslic_seg_engine_gpu.cu +++ b/modules/hfs/src/cuda/gslic_seg_engine_gpu.cu @@ -75,7 +75,7 @@ void SegEngineGPU::cvtImgSpace(Ptr inimg, Ptr outimg) dim3 blockSize(HFS_BLOCK_DIM, HFS_BLOCK_DIM); dim3 gridSize = getGridSize(img_size, blockSize); - cvtImgSpaceDevice << > >(inimg_ptr, img_size, outimg_ptr); + cvtImgSpaceDevice <<>>(inimg_ptr, img_size, outimg_ptr); } void SegEngineGPU::initClusterCenters() @@ -85,7 +85,7 @@ void SegEngineGPU::initClusterCenters() dim3 blockSize(HFS_BLOCK_DIM, HFS_BLOCK_DIM); dim3 gridSize = getGridSize(map_size, blockSize); - initClusterCentersDevice << > > + initClusterCentersDevice <<>> (img_ptr, map_size, img_size, spixel_size, spixel_list); } @@ -98,7 +98,7 @@ void SegEngineGPU::findCenterAssociation() dim3 blockSize(HFS_BLOCK_DIM, HFS_BLOCK_DIM); dim3 gridSize = getGridSize(img_size, blockSize); - findCenterAssociationDevice << > > + findCenterAssociationDevice <<>> (img_ptr, spixel_list, map_size, img_size, spixel_size, slic_settings.coh_weight, max_xy_dist, max_color_dist, idx_ptr); @@ -116,13 +116,13 @@ void SegEngineGPU::updateClusterCenter() dim3 blockSize(HFS_BLOCK_DIM, HFS_BLOCK_DIM); dim3 gridSize(map_size.x, map_size.y, no_grid_per_center); - updateClusterCenterDevice << > > + updateClusterCenterDevice <<>> (img_ptr, idx_ptr, map_size, img_size, spixel_size, no_blocks_per_line, accum_map_ptr); dim3 gridSize2(map_size.x, map_size.y); - finalizeReductionResultDevice << > > + finalizeReductionResultDevice <<>> (accum_map_ptr, map_size, no_grid_per_center, spixel_list_ptr); } @@ -134,13 +134,13 @@ void SegEngineGPU::enforceConnectivity() dim3 blockSize(HFS_BLOCK_DIM, HFS_BLOCK_DIM); dim3 gridSize = getGridSize(img_size, blockSize); - enforceConnectivityDevice << > > + enforceConnectivityDevice <<>> (idx_ptr, img_size, tmp_idx_ptr); - enforceConnectivityDevice << > > + enforceConnectivityDevice <<>> (tmp_idx_ptr, img_size, idx_ptr); - enforceConnectivityDevice1_2 << > > + enforceConnectivityDevice1_2 <<>> (idx_ptr, img_size, tmp_idx_ptr); - enforceConnectivityDevice1_2 << > > + enforceConnectivityDevice1_2 <<>> (tmp_idx_ptr, img_size, idx_ptr); } diff --git a/modules/hfs/src/cuda/magnitude.cu b/modules/hfs/src/cuda/magnitude.cu index 5fc556a8119..35d8377aab6 100644 --- a/modules/hfs/src/cuda/magnitude.cu +++ b/modules/hfs/src/cuda/magnitude.cu @@ -194,7 +194,7 @@ void Magnitude::derrivativeXYGpu() dim3 gridSize((int)ceil((float)img_size.x / (float)blockSize.x), (int)ceil((float)img_size.y / (float)blockSize.y)); - derrivativeXYDevice << > > + derrivativeXYDevice <<>> (gray_ptr, dx_ptr, dy_ptr, mag_ptr, img_size); } @@ -209,7 +209,7 @@ void Magnitude::nonMaxSuppGpu() dim3 gridSize((int)ceil((float)img_size.x / (float)blockSize.x), (int)ceil((float)img_size.y / (float)blockSize.y)); - nonMaxSuppDevice << > > + nonMaxSuppDevice <<>> (nms_ptr, dx_ptr, dy_ptr, mag_ptr, img_size); }