From 1c8ec2d1a6d4c2ba2adbfce4bfebb7f66e4fad33 Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Wed, 25 Sep 2024 17:33:42 +0200 Subject: [PATCH 1/6] Get CUDA code to compile with clang CUDA and without CUDA --- modules/cudaarithm/src/cuda/polar_cart.cu | 10 ++++---- modules/cudacodec/src/cuda/nv12_to_rgb.cu | 4 ++-- modules/cudaimgproc/src/color.cpp | 2 +- modules/cudaimgproc/src/cuda/canny.cu | 6 ++--- .../src/cuda/connectedcomponents.cu | 10 ++++---- .../cudaimgproc/src/cuda/generalized_hough.cu | 4 ++-- modules/cudaimgproc/src/cuda/hough_lines.cu | 2 +- .../cudaimgproc/src/cuda/hough_segments.cu | 2 +- modules/cudaimgproc/src/cuda/moments.cu | 6 ++--- modules/cudaimgproc/src/moments.cpp | 11 +++++---- modules/cudaoptflow/src/cuda/farneback.cu | 24 +++++++++---------- .../cudaoptflow/src/cuda/nvidiaOpticalFlow.cu | 2 +- modules/cudaoptflow/src/farneback.cpp | 2 +- modules/hfs/src/cuda/gslic_seg_engine_gpu.cu | 18 +++++++------- modules/hfs/src/cuda/magnitude.cu | 4 ++-- 15 files changed, 54 insertions(+), 53 deletions(-) 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/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/cuda/canny.cu b/modules/cudaimgproc/src/cuda/canny.cu index 46b9624fb68..369024a2dd4 100644 --- a/modules/cudaimgproc/src/cuda/canny.cu +++ b/modules/cudaimgproc/src/cuda/canny.cu @@ -368,7 +368,7 @@ namespace canny while (s_counter > 0 && s_counter <= stack_size - blockDim.x) { const int subTaskIdx = threadIdx.x >> 3; - const int portion = ::min(s_counter, blockDim.x >> 3); + const int portion = std::min(s_counter, (int)(blockDim.x >> 3)); if (subTaskIdx < portion) pos = s_st[s_counter - 1 - subTaskIdx]; @@ -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_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/moments.cpp b/modules/cudaimgproc/src/moments.cpp index 3c2e62c4b90..a9c8beb9f83 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) { @@ -36,6 +31,12 @@ cv::Moments cv::cuda::convertSpatialMoments(Mat spatialMoments, const MomentsOrd void 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/farneback.cu b/modules/cudaoptflow/src/cuda/farneback.cu index 7c902c8d226..b8a51ae8899 100644 --- a/modules/cudaoptflow/src/cuda/farneback.cu +++ b/modules/cudaoptflow/src/cuda/farneback.cu @@ -74,7 +74,7 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback { extern __shared__ float smem[]; volatile float *row = smem + tx; - int xWarped = ::min(::max(x, 0), width - 1); + int xWarped = std::min(std::max(x, 0), width - 1); row[0] = src(y, xWarped) * c_g[0]; row[bdx] = 0.f; @@ -82,8 +82,8 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback for (int k = 1; k <= polyN; ++k) { - float t0 = src(::max(y - k, 0), xWarped); - float t1 = src(::min(y + k, height - 1), xWarped); + float t0 = src(std::max(y - k, 0), xWarped); + float t1 = src(std::min(y + k, height - 1), xWarped); row[0] += c_g[k] * (t0 + t1); row[bdx] += c_xg[k] * (t1 - t0); @@ -224,10 +224,10 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback r3 += r6*dy + r5*dx; float scale = - c_border[::min(x, BORDER_SIZE)] * - c_border[::min(y, BORDER_SIZE)] * - c_border[::min(width - x - 1, BORDER_SIZE)] * - c_border[::min(height - y - 1, BORDER_SIZE)]; + c_border[std::min(x, BORDER_SIZE)] * + c_border[std::min(y, BORDER_SIZE)] * + c_border[std::min(width - x - 1, BORDER_SIZE)] * + c_border[std::min(height - y - 1, BORDER_SIZE)]; r2 *= scale; r3 *= scale; r4 *= scale; r5 *= scale; r6 *= scale; @@ -316,11 +316,11 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx) { int xExt = int(bx * bdx) + i - ksizeHalf; - xExt = ::min(::max(xExt, 0), width - 1); + xExt = std::min(std::max(xExt, 0), width - 1); row[i] = src(y, xExt); for (int j = 1; j <= ksizeHalf; ++j) - row[i] += src(::max(y - j, 0), xExt) + src(::min(y + j, height - 1), xExt); + row[i] += src(std::max(y - j, 0), xExt) + src(std::min(y + j, height - 1), xExt); } if (x < width) @@ -372,7 +372,7 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx) { int xExt = int(bx * bdx) + i - ksizeHalf; - xExt = ::min(::max(xExt, 0), width - 1); + xExt = std::min(std::max(xExt, 0), width - 1); #pragma unroll for (int k = 0; k < 5; ++k) @@ -382,8 +382,8 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback #pragma unroll for (int k = 0; k < 5; ++k) row[k*smw + i] += - src(k*height + ::max(y - j, 0), xExt) + - src(k*height + ::min(y + j, height - 1), xExt); + src(k*height + std::max(y - j, 0), xExt) + + src(k*height + std::min(y + j, height - 1), xExt); } if (x < width) 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/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); } From 02f037d1ed0c1c5a91fc86f08100d660b25b7cd4 Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Thu, 26 Sep 2024 09:41:41 +0200 Subject: [PATCH 2/6] Fix canny.cu --- modules/cudaimgproc/src/cuda/canny.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/cudaimgproc/src/cuda/canny.cu b/modules/cudaimgproc/src/cuda/canny.cu index 369024a2dd4..4c6de962a00 100644 --- a/modules/cudaimgproc/src/cuda/canny.cu +++ b/modules/cudaimgproc/src/cuda/canny.cu @@ -368,7 +368,7 @@ namespace canny while (s_counter > 0 && s_counter <= stack_size - blockDim.x) { const int subTaskIdx = threadIdx.x >> 3; - const int portion = std::min(s_counter, (int)(blockDim.x >> 3)); + const int portion = ::min(s_counter, blockDim.x >> 3); if (subTaskIdx < portion) pos = s_st[s_counter - 1 - subTaskIdx]; From 319eddc0119a0e99f3d3454fb8c8b28e4c672aae Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Thu, 26 Sep 2024 10:19:34 +0200 Subject: [PATCH 3/6] farneback.cu did not need to be modified --- modules/cudaoptflow/src/cuda/farneback.cu | 24 +++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/modules/cudaoptflow/src/cuda/farneback.cu b/modules/cudaoptflow/src/cuda/farneback.cu index b8a51ae8899..7c902c8d226 100644 --- a/modules/cudaoptflow/src/cuda/farneback.cu +++ b/modules/cudaoptflow/src/cuda/farneback.cu @@ -74,7 +74,7 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback { extern __shared__ float smem[]; volatile float *row = smem + tx; - int xWarped = std::min(std::max(x, 0), width - 1); + int xWarped = ::min(::max(x, 0), width - 1); row[0] = src(y, xWarped) * c_g[0]; row[bdx] = 0.f; @@ -82,8 +82,8 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback for (int k = 1; k <= polyN; ++k) { - float t0 = src(std::max(y - k, 0), xWarped); - float t1 = src(std::min(y + k, height - 1), xWarped); + float t0 = src(::max(y - k, 0), xWarped); + float t1 = src(::min(y + k, height - 1), xWarped); row[0] += c_g[k] * (t0 + t1); row[bdx] += c_xg[k] * (t1 - t0); @@ -224,10 +224,10 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback r3 += r6*dy + r5*dx; float scale = - c_border[std::min(x, BORDER_SIZE)] * - c_border[std::min(y, BORDER_SIZE)] * - c_border[std::min(width - x - 1, BORDER_SIZE)] * - c_border[std::min(height - y - 1, BORDER_SIZE)]; + c_border[::min(x, BORDER_SIZE)] * + c_border[::min(y, BORDER_SIZE)] * + c_border[::min(width - x - 1, BORDER_SIZE)] * + c_border[::min(height - y - 1, BORDER_SIZE)]; r2 *= scale; r3 *= scale; r4 *= scale; r5 *= scale; r6 *= scale; @@ -316,11 +316,11 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx) { int xExt = int(bx * bdx) + i - ksizeHalf; - xExt = std::min(std::max(xExt, 0), width - 1); + xExt = ::min(::max(xExt, 0), width - 1); row[i] = src(y, xExt); for (int j = 1; j <= ksizeHalf; ++j) - row[i] += src(std::max(y - j, 0), xExt) + src(std::min(y + j, height - 1), xExt); + row[i] += src(::max(y - j, 0), xExt) + src(::min(y + j, height - 1), xExt); } if (x < width) @@ -372,7 +372,7 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx) { int xExt = int(bx * bdx) + i - ksizeHalf; - xExt = std::min(std::max(xExt, 0), width - 1); + xExt = ::min(::max(xExt, 0), width - 1); #pragma unroll for (int k = 0; k < 5; ++k) @@ -382,8 +382,8 @@ namespace cv { namespace cuda { namespace device { namespace optflow_farneback #pragma unroll for (int k = 0; k < 5; ++k) row[k*smw + i] += - src(k*height + std::max(y - j, 0), xExt) + - src(k*height + std::min(y + j, height - 1), xExt); + src(k*height + ::max(y - j, 0), xExt) + + src(k*height + ::min(y + j, height - 1), xExt); } if (x < width) From b7580470971b39f4ff3ba68833eae26785ac8c04 Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Thu, 26 Sep 2024 11:11:24 +0200 Subject: [PATCH 4/6] Forgot a few more files --- modules/cudaimgproc/src/cuda/hough_circles.cu | 2 +- modules/cudaoptflow/src/precomp.hpp | 2 ++ modules/cudev/include/opencv2/cudev/grid/detail/minmaxloc.hpp | 4 ++-- 3 files changed, 5 insertions(+), 3 deletions(-) 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/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 From 63734663a9a706c3de972b9a0b2212cd581e3b98 Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Wed, 2 Oct 2024 12:05:29 +0200 Subject: [PATCH 5/6] Add missing calcHist --- modules/cudaimgproc/src/histogram.cpp | 1 + 1 file changed, 1 insertion(+) 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(); } From 821de70182f9045a92dc2f704aa5e72ce3754da9 Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Wed, 2 Oct 2024 15:04:47 +0200 Subject: [PATCH 6/6] Add missing implementations. --- modules/cudaarithm/src/arithm.cpp | 2 ++ modules/cudaarithm/src/element_operations.cpp | 5 +++++ modules/cudaarithm/src/reductions.cpp | 4 +++- modules/cudaimgproc/src/connectedcomponents.cpp | 4 ++-- modules/cudaimgproc/src/moments.cpp | 3 ++- 5 files changed, 14 insertions(+), 4 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/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/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/moments.cpp b/modules/cudaimgproc/src/moments.cpp index a9c8beb9f83..40ab9414ad7 100644 --- a/modules/cudaimgproc/src/moments.cpp +++ b/modules/cudaimgproc/src/moments.cpp @@ -27,8 +27,9 @@ 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"