Skip to content

Commit

Permalink
Merge branch 4.x
Browse files Browse the repository at this point in the history
  • Loading branch information
@asmorkalov
asmorkalov committed Nov 2, 2024
2 parents ddfb5cf + 5741f22 commit bd3b734
Show file tree
Hide file tree
Showing 24 changed files with 311 additions and 62 deletions.
230 changes: 215 additions & 15 deletions modules/cudaarithm/include/opencv2/cudaarithm.hpp
View file

Large diffs are not rendered by default.

35 changes: 35 additions & 0 deletions modules/cudaarithm/misc/python/test/test_cudaarithm.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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()
Expand All @@ -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)))
Expand Down Expand Up @@ -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()
Expand All @@ -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)))

Expand All @@ -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)))

Expand Down
2 changes: 2 additions & 0 deletions modules/cudaarithm/src/arithm.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<DFT> cv::cuda::createDFT(Size, int) { throw_no_cuda(); return Ptr<DFT>(); }

Ptr<Convolution> cv::cuda::createConvolution(Size) { throw_no_cuda(); return Ptr<Convolution>(); }

#else /* !defined (HAVE_CUDA) */
Expand Down
10 changes: 5 additions & 5 deletions modules/cudaarithm/src/cuda/polar_cart.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -289,9 +289,9 @@ namespace
const T scale = angleInDegrees ? static_cast<T>(CV_PI / 180.0) : static_cast<T>(1.0);

if (mag.empty())
polarToCartImpl_<T, false> << <grid, block, 0, stream >> >(mag, angle, x, y, scale);
polarToCartImpl_<T, false> <<<grid, block, 0, stream >>>(mag, angle, x, y, scale);
else
polarToCartImpl_<T, true> << <grid, block, 0, stream >> >(mag, angle, x, y, scale);
polarToCartImpl_<T, true> <<<grid, block, 0, stream >>>(mag, angle, x, y, scale);
}

template <typename T>
Expand All @@ -305,9 +305,9 @@ namespace
const T scale = angleInDegrees ? static_cast<T>(CV_PI / 180.0) : static_cast<T>(1.0);

if (mag.empty())
polarToCartDstInterleavedImpl_<T, false> << <grid, block, 0, stream >> >(mag, angle, xy, scale);
polarToCartDstInterleavedImpl_<T, false> <<<grid, block, 0, stream >>>(mag, angle, xy, scale);
else
polarToCartDstInterleavedImpl_<T, true> << <grid, block, 0, stream >> >(mag, angle, xy, scale);
polarToCartDstInterleavedImpl_<T, true> <<<grid, block, 0, stream >>>(mag, angle, xy, scale);
}

template <typename T>
Expand All @@ -320,7 +320,7 @@ namespace

const T scale = angleInDegrees ? static_cast<T>(CV_PI / 180.0) : static_cast<T>(1.0);

polarToCartInterleavedImpl_<T> << <grid, block, 0, stream >> >(magAngle, xy, scale);
polarToCartInterleavedImpl_<T> <<<grid, block, 0, stream >>>(magAngle, xy, scale);
}
}

Expand Down
5 changes: 5 additions & 0 deletions modules/cudaarithm/src/element_operations.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
4 changes: 3 additions & 1 deletion modules/cudaarithm/src/reductions.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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(); }

Expand Down
4 changes: 2 additions & 2 deletions modules/cudacodec/src/cuda/nv12_to_rgb.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<true> << <grid, block, 0, stream >> > (decodedFrame.ptr<uchar>(), decodedFrame.step, outFrame.ptr<uint>(), outFrame.step, width, height);
NV12_to_BGRA<true> <<<grid, block, 0, stream >>> (decodedFrame.ptr<uchar>(), decodedFrame.step, outFrame.ptr<uint>(), outFrame.step, width, height);
else
NV12_to_BGRA<false> << <grid, block, 0, stream >> > (decodedFrame.ptr<uchar>(), decodedFrame.step, outFrame.ptr<uint>(), outFrame.step, width, height);
NV12_to_BGRA<false> <<<grid, block, 0, stream >>> (decodedFrame.ptr<uchar>(), decodedFrame.step, outFrame.ptr<uint>(), outFrame.step, width, height);
CV_CUDEV_SAFE_CALL(cudaGetLastError());
if (stream == 0)
CV_CUDEV_SAFE_CALL(cudaDeviceSynchronize());
Expand Down
2 changes: 1 addition & 1 deletion modules/cudaimgproc/src/color.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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(); }

Expand Down
4 changes: 2 additions & 2 deletions modules/cudaimgproc/src/connectedcomponents.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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) */

Expand Down
4 changes: 2 additions & 2 deletions modules/cudaimgproc/src/cuda/canny.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<<<grid, block, 0, stream>>>(map, st1, st2, d_counter, count);
cudaSafeCall( cudaGetLastError() );
Expand All @@ -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;
Expand Down
10 changes: 5 additions & 5 deletions modules/cudaimgproc/src/cuda/connectedcomponents.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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 << <grid_size, block_size >> > (img, labels, last_pixel);
InitLabeling <<<grid_size, block_size >>> (img, labels, last_pixel);
cudaSafeCall(cudaGetLastError());

Compression << <grid_size, block_size >> > (labels);
Compression <<<grid_size, block_size >>> (labels);
cudaSafeCall(cudaGetLastError());

Merge << <grid_size, block_size >> > (labels, last_pixel);
Merge <<<grid_size, block_size >>> (labels, last_pixel);
cudaSafeCall(cudaGetLastError());

Compression << <grid_size, block_size >> > (labels);
Compression <<<grid_size, block_size >>> (labels);
cudaSafeCall(cudaGetLastError());

FinalLabeling << <grid_size, block_size >> > (img, labels);
FinalLabeling <<<grid_size, block_size >>> (img, labels);
cudaSafeCall(cudaGetLastError());

if (last_pixel_allocated) {
Expand Down
4 changes: 2 additions & 2 deletions modules/cudaimgproc/src/cuda/generalized_hough.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
Expand Down Expand Up @@ -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;
}
Expand Down
2 changes: 1 addition & 1 deletion modules/cudaimgproc/src/cuda/hough_circles.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
Expand Down
2 changes: 1 addition & 1 deletion modules/cudaimgproc/src/cuda/hough_lines.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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)
{
Expand Down
2 changes: 1 addition & 1 deletion modules/cudaimgproc/src/cuda/hough_segments.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -241,7 +241,7 @@ namespace cv { namespace cuda { namespace device

cudaSafeCall( cudaStreamSynchronize(stream) );

totalCount = ::min(totalCount, maxSize);
totalCount = std::min(totalCount, maxSize);
return totalCount;
}
}
Expand Down
6 changes: 3 additions & 3 deletions modules/cudaimgproc/src/cuda/moments.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -139,7 +139,7 @@ template <typename TSrc, typename TMoments, int nMoments> struct momentsDispatch
static void call(const PtrStepSz<TSrc> src, PtrStepSz<TMoments> moments, const bool binary, const int offsetX, const cudaStream_t stream) {
dim3 blockSize(blockSizeX, blockSizeY);
dim3 gridSize = dim3(divUp(src.rows, blockSizeY));
spatialMoments<TSrc, TMoments, false, false, nMoments> << <gridSize, blockSize, 0, stream >> > (src, binary, moments.ptr());
spatialMoments<TSrc, TMoments, false, false, nMoments> <<<gridSize, blockSize, 0, stream >>> (src, binary, moments.ptr());
if (stream == 0)
cudaSafeCall(cudaStreamSynchronize(stream));
};
Expand All @@ -150,9 +150,9 @@ template <typename TSrc, int nMoments> struct momentsDispatcherChar {
dim3 blockSize(blockSizeX, blockSizeY);
dim3 gridSize = dim3(divUp(src.rows, blockSizeY));
if (offsetX)
spatialMoments<TSrc, float, true, false, nMoments> << <gridSize, blockSize, 0, stream >> > (src, binary, moments.ptr(), offsetX);
spatialMoments<TSrc, float, true, false, nMoments> <<<gridSize, blockSize, 0, stream >>> (src, binary, moments.ptr(), offsetX);
else
spatialMoments<TSrc, float, true, true, nMoments> << <gridSize, blockSize, 0, stream >> > (src, binary, moments.ptr());
spatialMoments<TSrc, float, true, true, nMoments> <<<gridSize, blockSize, 0, stream >>> (src, binary, moments.ptr());

if (stream == 0)
cudaSafeCall(cudaStreamSynchronize(stream));
Expand Down
1 change: 1 addition & 0 deletions modules/cudaimgproc/src/histogram.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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(); }

Expand Down
14 changes: 8 additions & 6 deletions modules/cudaimgproc/src/moments.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<typename T>
cv::Moments convertSpatialMomentsT(Mat spatialMoments, const MomentsOrder order) {
switch (order) {
Expand All @@ -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 <typename TSrc, typename TMoments>
void moments(const PtrStepSzb src, PtrStepSzb moments, const bool binary, const int order, const int offsetX, const cudaStream_t stream);
Expand Down
2 changes: 1 addition & 1 deletion modules/cudaoptflow/src/cuda/nvidiaOpticalFlow.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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 << <gridDim, blockDim >> > (0, srcDevPtr, nSrcWidth, nSrcPitch, nSrcHeight,
NearestNeighborFlowKernel <<<gridDim, blockDim >>> (0, srcDevPtr, nSrcWidth, nSrcPitch, nSrcHeight,
0, dstDevPtr, nDstWidth, nDstPitch, nDstHeight,
nScaleFactor);

Expand Down
2 changes: 1 addition & 1 deletion modules/cudaoptflow/src/farneback.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -47,7 +47,7 @@ using namespace cv::cuda;

#if !defined HAVE_CUDA || defined(CUDA_DISABLER)

Ptr<FarnebackOpticalFlow> cv::cuda::FarnebackOpticalFlow::create(int, double, bool, int, int, int, double, int) { throw_no_cuda(); return Ptr<FarnebackOpticalFlow>(); }
Ptr<cv::cuda::FarnebackOpticalFlow> cv::cuda::FarnebackOpticalFlow::create(int, double, bool, int, int, int, double, int) { throw_no_cuda(); return Ptr<FarnebackOpticalFlow>(); }

#else

Expand Down
2 changes: 2 additions & 0 deletions modules/cudaoptflow/src/precomp.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
Loading

0 comments on commit bd3b734

Please sign in to comment.