Skip to content

Commit 843b6ed

Browse files
vrabaudvrabaud
authored
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
1 parent 80f1ca2 commit 843b6ed

File tree

22 files changed

+61
-47
lines changed

22 files changed

+61
-47
lines changed

‎modules/cudaarithm/src/arithm.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -54,6 +54,8 @@ void cv::cuda::mulAndScaleSpectrums(InputArray, InputArray, OutputArray, int, fl
5454

5555
void cv::cuda::dft(InputArray, OutputArray, Size, int, Stream&) { throw_no_cuda(); }
5656

57+
Ptr<DFT> cv::cuda::createDFT(Size, int) { throw_no_cuda(); return Ptr<DFT>(); }
58+
5759
Ptr<Convolution> cv::cuda::createConvolution(Size) { throw_no_cuda(); return Ptr<Convolution>(); }
5860

5961
#else /* !defined (HAVE_CUDA) */

‎modules/cudaarithm/src/cuda/polar_cart.cu

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -289,9 +289,9 @@ namespace
289289
const T scale = angleInDegrees ? static_cast<T>(CV_PI / 180.0) : static_cast<T>(1.0);
290290

291291
if (mag.empty())
292-
polarToCartImpl_<T, false> << <grid, block, 0, stream >> >(mag, angle, x, y, scale);
292+
polarToCartImpl_<T, false> <<<grid, block, 0, stream >>>(mag, angle, x, y, scale);
293293
else
294-
polarToCartImpl_<T, true> << <grid, block, 0, stream >> >(mag, angle, x, y, scale);
294+
polarToCartImpl_<T, true> <<<grid, block, 0, stream >>>(mag, angle, x, y, scale);
295295
}
296296

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

307307
if (mag.empty())
308-
polarToCartDstInterleavedImpl_<T, false> << <grid, block, 0, stream >> >(mag, angle, xy, scale);
308+
polarToCartDstInterleavedImpl_<T, false> <<<grid, block, 0, stream >>>(mag, angle, xy, scale);
309309
else
310-
polarToCartDstInterleavedImpl_<T, true> << <grid, block, 0, stream >> >(mag, angle, xy, scale);
310+
polarToCartDstInterleavedImpl_<T, true> <<<grid, block, 0, stream >>>(mag, angle, xy, scale);
311311
}
312312

313313
template <typename T>
@@ -320,7 +320,7 @@ namespace
320320

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

323-
polarToCartInterleavedImpl_<T> << <grid, block, 0, stream >> >(magAngle, xy, scale);
323+
polarToCartInterleavedImpl_<T> <<<grid, block, 0, stream >>>(magAngle, xy, scale);
324324
}
325325
}
326326

‎modules/cudaarithm/src/element_operations.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -84,8 +84,13 @@ void cv::cuda::magnitude(InputArray, InputArray, OutputArray, Stream&) { throw_n
8484
void cv::cuda::magnitudeSqr(InputArray, OutputArray, Stream&) { throw_no_cuda(); }
8585
void cv::cuda::magnitudeSqr(InputArray, InputArray, OutputArray, Stream&) { throw_no_cuda(); }
8686
void cv::cuda::phase(InputArray, InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
87+
void cv::cuda::phase(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
8788
void cv::cuda::cartToPolar(InputArray, InputArray, OutputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
89+
void cv::cuda::cartToPolar(InputArray, OutputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
90+
void cv::cuda::cartToPolar(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
8891
void cv::cuda::polarToCart(InputArray, InputArray, OutputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
92+
void cv::cuda::polarToCart(InputArray, InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
93+
void cv::cuda::polarToCart(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
8994

9095
#else
9196

‎modules/cudaarithm/src/reductions.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -69,8 +69,10 @@ void cv::cuda::countNonZero(InputArray, OutputArray, Stream&) { throw_no_cuda();
6969

7070
void cv::cuda::reduce(InputArray, OutputArray, int, int, int, Stream&) { throw_no_cuda(); }
7171

72-
void cv::cuda::meanStdDev(InputArray, Scalar&, Scalar&) { throw_no_cuda(); }
72+
void cv::cuda::meanStdDev(InputArray, OutputArray, InputArray, Stream&) { throw_no_cuda(); }
7373
void cv::cuda::meanStdDev(InputArray, OutputArray, Stream&) { throw_no_cuda(); }
74+
void cv::cuda::meanStdDev(InputArray, Scalar&, Scalar&, InputArray) { throw_no_cuda(); }
75+
void cv::cuda::meanStdDev(InputArray, Scalar&, Scalar&) { throw_no_cuda(); }
7476

7577
void cv::cuda::rectStdDev(InputArray, InputArray, OutputArray, Rect, Stream&) { throw_no_cuda(); }
7678

‎modules/cudacodec/src/cuda/nv12_to_rgb.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -179,9 +179,9 @@ void nv12ToBgra(const GpuMat& decodedFrame, GpuMat& outFrame, int width, int hei
179179
dim3 block(32, 8);
180180
dim3 grid(divUp(width, 2 * block.x), divUp(height, block.y));
181181
if (videoFullRangeFlag)
182-
NV12_to_BGRA<true> << <grid, block, 0, stream >> > (decodedFrame.ptr<uchar>(), decodedFrame.step, outFrame.ptr<uint>(), outFrame.step, width, height);
182+
NV12_to_BGRA<true> <<<grid, block, 0, stream >>> (decodedFrame.ptr<uchar>(), decodedFrame.step, outFrame.ptr<uint>(), outFrame.step, width, height);
183183
else
184-
NV12_to_BGRA<false> << <grid, block, 0, stream >> > (decodedFrame.ptr<uchar>(), decodedFrame.step, outFrame.ptr<uint>(), outFrame.step, width, height);
184+
NV12_to_BGRA<false> <<<grid, block, 0, stream >>> (decodedFrame.ptr<uchar>(), decodedFrame.step, outFrame.ptr<uint>(), outFrame.step, width, height);
185185
CV_CUDEV_SAFE_CALL(cudaGetLastError());
186186
if (stream == 0)
187187
CV_CUDEV_SAFE_CALL(cudaDeviceSynchronize());

‎modules/cudaimgproc/src/color.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -51,7 +51,7 @@ void cv::cuda::cvtColor(InputArray, OutputArray, int, int, Stream&) { throw_no_c
5151

5252
void cv::cuda::demosaicing(InputArray, OutputArray, int, int, Stream&) { throw_no_cuda(); }
5353

54-
void cv::cuda::swapChannels(InputOutputArray, const int[], Stream&) { throw_no_cuda(); }
54+
void cv::cuda::swapChannels(InputOutputArray, const int[4], Stream&) { throw_no_cuda(); }
5555

5656
void cv::cuda::gammaCorrection(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
5757

‎modules/cudaimgproc/src/connectedcomponents.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,8 @@ using namespace cv::cuda;
99

1010
#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
1111

12-
void cv::cuda::connectedComponents(InputArray img_, OutputArray labels_, int connectivity,
13-
int ltype, ConnectedComponentsAlgorithmsTypes ccltype) { throw_no_cuda(); }
12+
void cv::cuda::connectedComponents(InputArray, OutputArray, int, int, ConnectedComponentsAlgorithmsTypes) { throw_no_cuda(); }
13+
void cv::cuda::connectedComponents(InputArray, OutputArray, int, int) { throw_no_cuda(); }
1414

1515
#else /* !defined (HAVE_CUDA) */
1616

‎modules/cudaimgproc/src/cuda/canny.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -428,7 +428,7 @@ namespace canny
428428
cudaSafeCall( cudaMemsetAsync(d_counter, 0, sizeof(int), stream) );
429429

430430
const dim3 block(128);
431-
const dim3 grid(::min(count, 65535u), divUp(count, 65535), 1);
431+
const dim3 grid(std::min(count, 65535), divUp(count, 65535), 1);
432432

433433
edgesHysteresisGlobalKernel<<<grid, block, 0, stream>>>(map, st1, st2, d_counter, count);
434434
cudaSafeCall( cudaGetLastError() );
@@ -439,7 +439,7 @@ namespace canny
439439
cudaSafeCall( cudaMemcpyAsync(&count, d_counter, sizeof(int), cudaMemcpyDeviceToHost, stream) );
440440
cudaSafeCall( cudaStreamSynchronize(stream) );
441441

442-
count = min(count, map.cols * map.rows);
442+
count = std::min(count, map.cols * map.rows);
443443

444444
//std::swap(st1, st2);
445445
short2* tmp = st1;

‎modules/cudaimgproc/src/cuda/connectedcomponents.cu

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -317,19 +317,19 @@ void BlockBasedKomuraEquivalence(const cv::cuda::GpuMat& img, cv::cuda::GpuMat&
317317
grid_size = dim3((((img.cols + 1) / 2) - 1) / kblock_cols + 1, (((img.rows + 1) / 2) - 1) / kblock_rows + 1, 1);
318318
block_size = dim3(kblock_cols, kblock_rows, 1);
319319

320-
InitLabeling << <grid_size, block_size >> > (img, labels, last_pixel);
320+
InitLabeling <<<grid_size, block_size >>> (img, labels, last_pixel);
321321
cudaSafeCall(cudaGetLastError());
322322

323-
Compression << <grid_size, block_size >> > (labels);
323+
Compression <<<grid_size, block_size >>> (labels);
324324
cudaSafeCall(cudaGetLastError());
325325

326-
Merge << <grid_size, block_size >> > (labels, last_pixel);
326+
Merge <<<grid_size, block_size >>> (labels, last_pixel);
327327
cudaSafeCall(cudaGetLastError());
328328

329-
Compression << <grid_size, block_size >> > (labels);
329+
Compression <<<grid_size, block_size >>> (labels);
330330
cudaSafeCall(cudaGetLastError());
331331

332-
FinalLabeling << <grid_size, block_size >> > (img, labels);
332+
FinalLabeling <<<grid_size, block_size >>> (img, labels);
333333
cudaSafeCall(cudaGetLastError());
334334

335335
if (last_pixel_allocated) {

‎modules/cudaimgproc/src/cuda/generalized_hough.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -302,7 +302,7 @@ namespace cv { namespace cuda { namespace device
302302
int totalCount;
303303
cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
304304

305-
totalCount = ::min(totalCount, maxSize);
305+
totalCount = std::min(totalCount, maxSize);
306306

307307
return totalCount;
308308
}
@@ -812,7 +812,7 @@ namespace cv { namespace cuda { namespace device
812812
int totalCount;
813813
cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
814814

815-
totalCount = ::min(totalCount, maxSize);
815+
totalCount = std::min(totalCount, maxSize);
816816

817817
return totalCount;
818818
}

0 commit comments

Comments
 (0)