Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fix memcheck errors found in REDUCTION_TEST #13574

Merged
merged 2 commits into from
Jun 20, 2023

Conversation

davidwendt
Copy link
Contributor

Description

The nightly tests found memcheck errors in REDUCTION_TEST for any/all aggregation types on dictionary columns.
The error is caused by the internal atomicOr() function which reads the target value as 32-bit even if the type is bool (8-bit).
The fix works around this limitation by using atomicOr() on a 32-bit device scalar and then converting that to a bool scalar on return.

All other tests in REDUCTION_TEST now pass memcheck successfully.

Closes #13572

Checklist

  • I am familiar with the Contributing Guidelines.
  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

@davidwendt davidwendt added bug Something isn't working 3 - Ready for Review Ready for review by team libcudf Affects libcudf (C++/CUDA) code. non-breaking Non-breaking change labels Jun 14, 2023
@davidwendt davidwendt self-assigned this Jun 14, 2023
@davidwendt
Copy link
Contributor Author

The gtest was recoded slightly for clarity and to make this problem easier to debug. It looks larger only because of the reformatting.

@davidwendt davidwendt changed the title Fix memcheck errors found in REDUCTION_TEST Fix memcheck errors found in REDUCTION_TEST Jun 14, 2023
@davidwendt davidwendt marked this pull request as ready for review June 14, 2023 19:27
@davidwendt davidwendt requested a review from a team as a code owner June 14, 2023 19:27
@davidwendt davidwendt changed the title Fix memcheck errors found in REDUCTION_TEST Fix memcheck errors found in REDUCTION_TEST Jun 14, 2023
Copy link
Member

@PointKernel PointKernel left a comment

Choose a reason for hiding this comment

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

LGTM

@davidwendt
Copy link
Contributor Author

/merge

@rapids-bot rapids-bot bot merged commit 7e0ae51 into rapidsai:branch-23.08 Jun 20, 2023
@davidwendt davidwendt deleted the bug-reduction-memcheck branch June 20, 2023 14:38
rapids-bot bot pushed a commit that referenced this pull request Jun 22, 2023
Fixes a memcheck error found in `STRINGS_TEST` where an `atomicOr` was used on a boolean device scalar. The workaround uses a `cub::WarpReduce` to compute the result in the warp-per-string kernel.

Reference #13574

Authors:
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - Nghia Truong (https://github.com/ttnghia)

URL: #13578
rapids-bot bot pushed a commit that referenced this pull request Jul 10, 2023
Contributes to #13575

Depends on #13574, #13578

This PR cleans up custom atomic implementations in libcudf by using `cuda::atomic_ref` when possible. It removes atomic bitwise operations like `and`, `or` and `xor` since libcudac++ already provides proper replacements.

Authors:
  - Yunsong Wang (https://github.com/PointKernel)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - David Wendt (https://github.com/davidwendt)

URL: #13583
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
3 - Ready for Review Ready for review by team bug Something isn't working libcudf Affects libcudf (C++/CUDA) code. non-breaking Non-breaking change
Projects
None yet
Development

Successfully merging this pull request may close these issues.

[BUG] Invalid memory read in cudf::reduction::detail::reduce with aggregation ANY/ALL on dictionary column
4 participants