-
Notifications
You must be signed in to change notification settings - Fork 724
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
[SYCL][Reduction] Remove atomic64 check for float
reductions
#6434
[SYCL][Reduction] Remove atomic64 check for float
reductions
#6434
Conversation
This patch refers to intel/llvm#6434
@aelovikov-intel - can you please review this patch as the only active developer of reduction these days? |
@v-klochkov do we require to run any performance tests before commit? |
26bedb6
to
6739e29
Compare
Tag @steffenlarsen , in case you'd like to have a look as well. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Looks good! @pgorlani - Please resolve the conflicts.
Co-authored-by: aelovikov-intel <andrei.elovikov@intel.com>
934454b
6739e29
to
934454b
Compare
Thanks for your reviews, @steffenlarsen @aelovikov-intel! I rebased this PR onto the last refactoring of reductions, it should be fine. |
Thanks! |
intel#6434 enabled treating "float" as suitable for Reduction::has_fast_atomics implementation but that is slower than the one available under Reduction::has_fast_reduce. Make sure to check for the latter first.
#6434 enabled treating "float" as suitable for Reduction::has_fast_atomics implementation but that is slower than the one available under Reduction::has_fast_reduce. Make sure to check for the latter first.
This patch
removes the check on
aspect::atomic64
for reductions using atomic operations of single precision floating-point values. This solves sycl::reduction for float not working on AMD gpu #6054. However, it's not strictly necessary since also [SYCL][HIP] Implement PI_DEVICE_INFO_ATOMIC_64 for HIP #6429 solves it.adds the possibility of using
min
andmax
atomic operations for reductions ofdouble
values.