[Draft] Refactor GPU reductions and add unsafe atomic tunings #247
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Use atomic fucntions directly in Base and Lambda variants instead of using RAJA atomics. Add a variety of util functions.
Use single macro for all single reduction kernels instead of duplicating the code
Now gpu kernels using atomics have a default tuning called atomic and tunings with unsafe atomics have a tuning called unsafeAtomics.
This makes it clear where atomics are being used as part of a reduction or not
rhornung67
approved these changes
Jun 3, 2022
|
I kept finding things to change... |
CRobeck
reviewed
Jun 7, 2022
| #endif | ||
| template < size_t block_size > | ||
| __launch_bounds__(block_size) | ||
| __global__ void reduce_sum_unsafe(Real_ptr x, Real_ptr dsum, Real_type sum_init, |
There was a problem hiding this comment.
If we start to add more tests with atomics we might not want to split these up into separate kernels for compile time concerns but should be OK for now.
CRobeck
reviewed
Jun 7, 2022
| return devProp.gcnArchName; | ||
| } | ||
|
|
||
| #if defined(__gfx90a__) |
There was a problem hiding this comment.
Will be able to do away with gfx90a arch check in future Rocm release.
MrBurmark
changed the title
|
I'm closing this PR for a number of reasons. Its worth differentiating between safe and unsafe atomics as that is a temporary issue. The reducer implementation is difficult to put in a macro and not identically duplicated, so its not a huge gain to abstract it. The lambda variants are not necessary. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Refactor GPU reductions and add unsafe atomic tunings
This adds a "unsafeAtomic" tuning of each of the kernels with a hip variant using atomics.
This also refactors gpu reductions so the implementation is not duplicated in each kernel with a reduction.
This also adds lambda variants of gpu reduction kernels.