1
0
mirror of https://github.com/opencv/opencv_contrib.git synced 2025-10-17 15:26:00 +08:00
Commit Graph

41 Commits

Author SHA1 Message Date
Troels Ynddal
84024547e3 Merge pull request #3970 from troelsy:4.x
Fix race condition in Otsu's method #3970

I found that in some cases, Otsu's method can have two or more equal thresholds. In case of the CUDA implementation, this would create a race-condition where the last thread to write the output gets to choose which Otsu threshold to use. The resulting image from cv::cuda::threshold would be the same, but the returned threshold would vary for each run.

I solve this by first doing a poll to check if there are multiple matches with `__syncthreads_count`. If not, we can return the threshold as before. If there are more than one, the kernel does a minimum reduction to find the lowest threshold that matches the otsu score. It doesn't matter which threshold you choose as long as it is consistent, so I choose the smallest one.

I'm unsure when `__syncthreads_count` as introduced, but it is compatible with CC≥2.0. CUDA Toolkit's archive only goes back to version 8.0, but it was documented back then (https://docs.nvidia.com/cuda/archive)


### 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
- [ ] The PR is proposed to the proper branch
- [ ] There is a reference to the original bug report and related work
- [ ] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [ ] The feature is well documented and sample code can be built with the project CMake
2025-07-17 09:36:33 +03:00
Alexander Smorkalov
b72527fbca Fixed out-of-bound access in CUDA Otsu threshold implementation. 2025-06-27 17:04:40 +03:00
Troels Ynddal
6329974d4d Merge pull request #3943 from troelsy:4.x
Add Otsu's method to cv::cuda::threshold #3943

I implemented Otsu's method in CUDA for a separate project and want to add it to cv::cuda::threshold

I have made an effort to use existing OpenCV functions in my code, but I had some trouble with `ThresholdTypes` and `cv::cuda::calcHist`. I couldn't figure out how to include `precomp.hpp` to get the definition of `ThresholdTypes`. For `cv::cuda::calcHist` I tried adding `opencv_cudaimgproc`, but it creates a circular dependency on `cudaarithm`. I have include a simple implementation of `calcHist` so the code runs, but I would like input on how to use `cv::cuda::calcHist` instead. 

### 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
- [ ] The PR is proposed to the proper branch
- [ ] There is a reference to the original bug report and related work
- [ ] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [ ] The feature is well documented and sample code can be built with the project CMake
2025-06-19 16:38:53 +03:00
Alexander Smorkalov
c76792f4f7 Merge pull request #3803 from cudawarped:cuda_update_to_npp_stream_ctx
cuda - update npp calls to use the new NppStreamContext API if available
2024-11-05 09:07:31 +03:00
Alexander Smorkalov
5741f22f56 Merge pull request #3815 from cudawarped:cuda_fix_python_binaryop_with_scalar
`cudaarithm`: fix python bindings for binary ops involving scalars
2024-10-31 16:02:31 +03:00
Vincent Rabaud
843b6ede6a 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
2024-10-29 17:31:10 +03:00
cudawarped
06ddccebed cudaarithm: fix python bindings for binary ops involving scalars 2024-10-29 15:42:50 +02:00
cudawarped
06f764abc1 cuda - update npp calls to use the new NppStreamContext API if available 2024-10-03 15:38:54 +03:00
Alexander Smorkalov
09eb618804 Workaround for CUDA 12.6 tuple_size issue #3773. 2024-09-03 13:47:48 +03:00
Pierre Chatelier
667a66ee0e Merge pull request #3607 from chacha21:cuda_phase_interleaved
Add interleaved versions of phase/cartToPolar/polarToCart #3607

This PR is for performance only (at the cost of more template code and increased GPU code size) The additional variants can help the caller skip the creation of temporary GPU mats (where memory is more likely to be a critical resource), and can even allow in-place processing. magnitude/angles/x/y are often already interleaved when dealing with DFTs.

### 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
- [ ] There is a reference to the original bug report and related work
- [X] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [X] The feature is well documented and sample code can be built with the project CMake
2024-07-22 16:00:24 +03:00
Gregor Burger
14e34187c5 also link to CUDA::cufft_static in case of BUILD_SHARED_LIBS=OFF 2024-07-09 08:42:51 +02:00
Alexander Smorkalov
1ed3dd2c53 Merge pull request #3744 from asmorkalov:as/variadic_tuple
Added CUDA 12.4+ support #3744

Tries to fix https://github.com/opencv/opencv_contrib/issues/3690 for CUDA 12.4+
Related patch to main repo: https://github.com/opencv/opencv/pull/25658

Changes:
- Added branches to support new variadic implementation of thrust::tuple
- Added branch with std::array instead of std::tuple in split-merge and grid operations. The new branch got rid of namespace clash: cv::cuda in OpenCV and ::cuda in CUDA standard library (injected by Thrust). Old tuple branches presumed for compatibility with old code and CUDA versions before 12.4.

### 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
- [x] There is a reference to the original bug report and related work
- [ ] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [ ] The feature is well documented and sample code can be built with the project CMake
2024-05-30 15:48:48 +03:00
sdy623
4e766a039e cudaarithm: fix the compile faiure of CUDA 12.4.x .
A slight API change of NPP nppiMeanStdDevGetBufferHostSize_8u_C1R
The type of bufSize is size_t instead of int in CUDA 12.4.x
2024-04-23 21:43:40 +09:00
James Bowley
1ae5f03a07 build: add first class cuda support 2023-12-14 18:15:34 +02:00
Alexander Smorkalov
51de5807e8 Set of build fixes for ## 3580, 3574. 2023-10-26 20:19:11 +03:00
Alexander Smorkalov
8db3e627fb Merge pull request #3378 from cudawarped:replace_texture_ref_with_texture_obj
Fix CUDA texture bugs and replace all instances of CUDA texture references with texture objects
2022-12-20 15:07:09 +03:00
James Bowley
8a6ea82ed0 Replace all instances of texture references/objects with texture objects using the existing updated cv::cudev::Texture class.
Fixes bugs in cv::cuda::demosaicing, cv::cuda::resize and cv::cuda::HoughSegmentDetector.
2022-12-15 12:36:59 +02:00
David Geldreich
53ab7c3efa keep cufftPlan2d across ConvolveImpl::convolve calls
on some CUDA versions creating/destroying cufftPlan2d is very time consuming
we now create them in ConvolveImpl::create() and destroy them in the dtor

this solves issue #3385
2022-12-01 16:41:36 +00:00
cudawarped
ea450e6378 Fix cuda::minMax and cuda::minMaxLoc python bindings to correctly output the results.
Added test case.
2022-03-22 18:41:53 +00:00
Alexander Alekhin
dd8be233ad Merge pull request #3191 from asmorkalov:as/cuda_meanstddev 2022-03-11 17:53:49 +00:00
Alexander Smorkalov
df101fd7e3 Added 32FC1 type support and mask to cuda::meanStdDev implementation. 2022-03-11 08:33:03 +03:00
cudawarped
55cbe767ac Add locateROI test case to check GpuMat::dataend calculation. 2022-02-02 14:26:56 +00:00
cudawarped
fdb53d633d Merge pull request #3144 from cudawarped:test_cuda_event
Test cv::cuda::event
2022-01-11 23:58:30 +03:00
Suleyman TURKMEN
180775f5bb fix legacy contants 2021-12-01 02:21:38 +03:00
Alexander Alekhin
4b20db546b [moved from opencv] cuda: fix inplace condition in cv::cuda::flip
original commit: e4b0251e9d
2021-04-01 02:26:59 +00:00
Aaron Miller
f1c0b5eae6 Implement cv::cuda::inRange (Fixes OpenCV #6295) 2021-01-18 21:50:41 -05:00
Rob Timpe
5b640a53f1 [moved from opencv] Fix errors when building with cuda stubs
Fixes two errors when building with the options WITH_CUDA=ON and BUILD_CUDA_STUBS=ON on a machine without CUDA.

In the cudaarithm module, make sure cuda_runtime.h only gets included when CUDA is installed.

In the stitching module, don't assume that cuda is present just because cudaarithm and cudawarping are present (as is the case when building with the above options).

original commit: 22ee5c0c4d
2020-10-21 15:51:46 -07:00
Tomoaki Teshima
46dba0f855 [moved from opencv] brush up by following the comments
original commit: 234117800f
2020-09-25 23:57:15 +09:00
Tomoaki Teshima
2ae9011019 [moved from opencv] use only even number for inplace flip
original commit: a61546680b
2020-09-16 15:45:03 +09:00
Namgoo Lee
b364cfed1d [moved from opencv] test code
original commit: 11ac26bfb4
2020-07-30 01:42:44 +09:00
Tomoaki Teshima
ccc5f3293a [moved from opencv] let the test pass on Jetson
original commit: 19646ad049
2020-07-25 23:45:51 +09:00
Namgoo Lee
3ba40a5316 Port opencv/#17863 2020-07-22 16:30:03 +09:00
cudawarped
d7d6360fce Merge pull request #2396 from cudawarped:fix_python_cudawarping_cudaarithm
Add python bindings to cudaobjdetect, cudawarping and cudaarithm

* Overload cudawarping functions to generate correct python bindings.
Add python wrapper to convolution funciton.

* Added shift and hog.

* Moved cuda python tests to this repo and added python bindings to SURF.

* Fix SURF documentation and allow meanshiftsegmention to create GpuMat internaly if not passed for python bindings consistency.

* Add correct cuda SURF test case.

* Fix python mog and mog2 python bindings, add tests and  correct cudawarping documentation.

* Updated KeyPoints in cuda::ORB::Convert python wrapper to be an output argument.

* Add changes suggested by alalek

* Added changes suggested by asmorkalov
2020-01-29 12:54:42 +03:00
cudawarped
454434bdff Cudaarithm python bindings for merge and split are inconsistent with the format of existing cuda python functions. Split does not return its output and merg's documentation is incorrect. 2019-12-16 18:34:42 +00:00
Ahmed Ashour
55f389e8fa [moved from opencv] java: generated code to have javadoc
original commit: 5c56b8ce92
2019-06-05 12:44:03 +02:00
Namgoo Lee
90d1627194 [moved from opencv] Move Ptr-related code from lut.cu to lut.cpp
original commit: a54affeb8d
2019-03-13 20:25:28 +09:00
Antonio Borondo
c7c7036121 [moved from opencv] Fix documentation of cv::cuda::compare
original commit: 28e55dc5f3
2018-10-01 13:23:18 +01:00
Hamdi Sahloul
4b1e26a9ca [moved from opencv] cuda::polarToCart: update documentation
original commit: 5db13fe2a7
2018-09-22 07:02:43 +09:00
Hamdi Sahloul
96320cacfc [moved from opencv] cuda::polarToCart: test double precision and tune tolerance
original commit: 9932612373
2018-09-22 07:02:43 +09:00
Hamdi Sahloul
4457f88244 [moved from opencv] cuda::polarToCart: Support double precision
original commit: 6203c95d31
2018-09-22 07:02:43 +09:00
Alexander Alekhin
82733fe56b cuda: move CUDA modules to opencv_contrib
OpenCV 4.0+
2018-09-19 18:57:43 +03:00