Skip to content

cudafilters: remove danger race condition#19534

Merged
opencv-pushbot merged 1 commit intoopencv:3.4from
tomoaki0705:fixCudaFiltersRace
Feb 18, 2021
Merged

cudafilters: remove danger race condition#19534
opencv-pushbot merged 1 commit intoopencv:3.4from
tomoaki0705:fixCudaFiltersRace

Conversation

@tomoaki0705
Copy link
Copy Markdown
Contributor

@tomoaki0705 tomoaki0705 commented Feb 15, 2021 •
edited by alalek
Loading

Follow up of #12721

There are still some race condition in median_filter.cu

opencv/modules/cudafilters/src/cuda/median_filter.cu

Lines 130 to 137 in fa8684b

if (1 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-1];
__syncthreads();
if (2 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-2];
__syncthreads();
if (4 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-4];

Though I inserted __syncthreads() between the lines, the operation += does both read and write.
This causes a race condition

Pull Request Readiness Checklist

See details at https://github.com/opencv/opencv/wiki/How_to_contribute#making-a-good-pull-request

  • I agree to contribute to the project under Apache 2 License.
  • To the best of my knowledge, the proposed patch is not based on a code under GPL or other license that is incompatible with OpenCV
  • The PR is proposed to proper branch
  • There is reference to 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
force_builders=Custom
buildworker:Custom=linux-4,linux-6
build_image:Custom=ubuntu-cuda:16.04

@asmorkalov
Copy link
Copy Markdown
Contributor

asmorkalov commented Feb 15, 2021

cc @nglee

@nglee
Copy link
Copy Markdown
Contributor

nglee commented Feb 16, 2021 •
edited
Loading

I have not tested this fix myself on a machine yet, but I agree that it is safer with this PR.

alalek
alalek approved these changes Feb 18, 2021
View reviewed changes
Copy link
Copy Markdown
Member

@alalek alalek left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you 👍

@opencv-pushbot opencv-pushbot merged commit 0421c39 into opencv:3.4 Feb 18, 2021
@tomoaki0705 tomoaki0705 deleted the fixCudaFiltersRace branch February 18, 2021 20:51
This was referenced Feb 21, 2021
Merged
Merged
@alalek alalek mentioned this pull request Apr 9, 2021
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@alalek alalek alalek approved these changes

Assignees

@alalek alalek

Labels

bug category: gpu/cuda (contrib) OpenCV 4.0+: moved to opencv_contrib

Projects

None yet

Milestone

3.4.14

Development

Successfully merging this pull request may close these issues.

5 participants

@tomoaki0705 @asmorkalov @nglee @alalek @opencv-pushbot