Remove local (in-group) atomic usage from __parallel_find_or (V1) #1668
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
… local atomic from __parallel_find_or
…view comment: remove duplicated comments Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com> (cherry picked from commit cbf9e7a)
(cherry picked from commit 7328925)
…FPGA_DEVICE macro Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
…view comment: rename __save_state_to for works with atomic to __save_state_to_atomic Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
…sycl::bit_or instead of __dpl_sycl::__plus Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
…- using sycl::bit_or instead of __dpl_sycl::__plus" - rolled back due very small or absent perf profit Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
… FoundLocalReduceOp to _LocalResultsReduceOp Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
SergeyKopienko
force-pushed
the
dev/skopienko/remove_local_atomic_from_find_or_V1
branch
from
d6dd3da
to
32a9d32
Compare
… and use __found_state in the struct __parallel_or_tag Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
…dpl_sycl::__bit_or in struct __parallel_or_tag Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
SergeyKopienko
force-pushed
the
dev/skopienko/remove_local_atomic_from_find_or_V1
branch
from
5b735ce
to
9201751
Compare
SergeyKopienko
commented
Jul 8, 2024
| @@ -139,6 +141,9 @@ using __maximum = sycl::ONEAPI::maximum<_T>; | |||
|
|
|||
| template <typename _T> | |||
| using __minimum = sycl::ONEAPI::minimum<_T>; | |||
|
|
|||
| template <typename _T = void> | |||
| using __bit_or = sycl::ONEAPI::bit_or<_T>; | |||
There was a problem hiding this comment.
SergeyKopienko
commented
Jul 8, 2024
| @@ -130,6 +130,8 @@ using __maximum = sycl::maximum<_T>; | |||
| template <typename _T = void> | |||
| using __minimum = sycl::minimum<_T>; | |||
|
|
|||
| template <typename _T = void> | |||
| using __bit_or = sycl::bit_or<_T>; | |||
There was a problem hiding this comment.
There was a problem hiding this comment.
Test evidence for complier icx 2021.1.2 (whene group operations isn't present) ): https://godbolt.org/z/PcnYT4vqe
julianmi
reviewed
Jul 8, 2024
julianmi
reviewed
Jul 8, 2024
…view comment: remove unexpected comment Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
|
@danhoeflinger, @julianmi is something that mandatory required to change in this PR? |
julianmi
reviewed
Jul 8, 2024
Co-authored-by: Julian Miller <julian.miller@intel.com>
julianmi
previously approved these changes
Jul 8, 2024
…NEDPL_FPGA_EMU instead of _ONEDPL_FPGA_DEVICE Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
danhoeflinger
approved these changes
Jul 8, 2024
julianmi
approved these changes
Jul 9, 2024
SergeyKopienko
deleted the
dev/skopienko/remove_local_atomic_from_find_or_V1
branch
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.
In this PR we remove local (in-group) atomic usage from
__parallel_find_orimplementation:__found_localnow is local variable (for each item);__found_localwithin one group through__dpl_sycl::__reduce_over_groupoperation :__found_local = __dpl_sycl::__reduce_over_group(__item_id.get_group(), __found_local, typename _BrickTag::_LocalResultsReduceOp{});where
_LocalResultsReduceOpis :__dpl_sycl::__minimum(for the__parallel_find_forward_tag);__dpl_sycl::__maximum(for the__parallel_find_backward_tag);__dpl_sycl::__bit_or(for the__parallel_or_tag).The new operation
__dpl_sycl::__bit_orhas been defined in the file sycl_defs.h :