-
Notifications
You must be signed in to change notification settings - Fork 787
[SYCL] Add implementation of sycl::intel::barrier #2198
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
Conversation
Exposes barrier as a free function, aligned with group_barrier from SYCL 2020 provisional. Naming is aligned with the existing DPC++ group algorithms, hence barrier in place of group_barrier. Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
Closer alignment with 2020 provisional. Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
Please fix test/build failures. |
I don't think they're related to this PR. The same test is failing across multiple PRs. |
Even better! |
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.
LGTM
sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc
Outdated
Show resolved
Hide resolved
Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
Fix L0 build issues for testing. Signed-off-by: John Pennycook <[email protected]>
If scope is too narrow, implementation will ignore it. Signed-off-by: John Pennycook <[email protected]>
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.
LGTM, thanks!
@mkinsner: Hi Mike, this PR requires your approval. |
@v-klochkov: I think @jbrodman's review counts here. |
This reverts commit 3df81ac.
…2261) This reverts commit 3df81ac. SYCL::barrier.cpp is failing for every run after merging the PR both on Windows and Linux (e.g. http://ci.llvm.intel.com:8010/#/builders/18/builds/3851)
A set of llvm.vector.reduce.* intrinsics doesn't have straight forward operation equivalent on the SPIRV side. The easiest solution to this problem is to use scalar operation on each pair of vector elements and repeat until there is only one value. Original commit: KhronosGroup/SPIRV-LLVM-Translator@fe088cd
Exposes barrier as a free function, aligned with group_barrier from SYCL 2020
provisional.
Naming is aligned with the existing DPC++ group algorithms, hence barrier in
place of group_barrier.
Signed-off-by: John Pennycook [email protected]