-
Notifications
You must be signed in to change notification settings - Fork 769
[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 <john.pennycook@intel.com>
Signed-off-by: John Pennycook <john.pennycook@intel.com>
Closer alignment with 2020 provisional. Signed-off-by: John Pennycook <john.pennycook@intel.com>
Signed-off-by: John Pennycook <john.pennycook@intel.com>
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 <john.pennycook@intel.com>
Signed-off-by: John Pennycook <john.pennycook@intel.com>
Signed-off-by: John Pennycook <john.pennycook@intel.com>
Fix L0 build issues for testing. Signed-off-by: John Pennycook <john.pennycook@intel.com>
If scope is too narrow, implementation will ignore it. Signed-off-by: John Pennycook <john.pennycook@intel.com>
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 john.pennycook@intel.com