Clarify group synchronization behavior #499
Merged
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
The previous description of group_barrier required implementations to execute certain fences, which should be an implementation detail. This commit adds a "Synchronization" section to group_barrier and describes behavior in terms of a happens-before relationship. Implementations may choose to implement the barrier using any synchronization operations which satisfy this requirement.
State explicitly that: - Group algorithms block until the algorithm executes. - Group algorithms have similar happens-before implications to barriers. Closes internal issue 576.
C++ just says "blocks" rather than "blocks this thread", so we can drop "blocks this work-item".
Explains that a group function may define operation(s) that take place only when all work-items have arried at the group function's synchronization point. The intent is that group algorithms can then be considered a specific instance of these group functions, where the operation(s) constitute the algorithm that should be executed.
These were accidentally overlooked during the previous round of changes.
TApplencourt
approved these changes
Nov 29, 2023
gmlueck
reviewed
Nov 30, 2023
Co-authored-by: Greg Lueck <gregory.m.lueck@intel.com>
gmlueck
approved these changes
Nov 30, 2023
|
This needs one more reviewer. Can someone volunteer? |
keryell
approved these changes
Dec 7, 2023
I agree that there are probably some opportunities to reduce the amount of duplicated information. I'd rather we do that in a follow-up PR, though, if everybody is okay with that. Ideally, I'd like us to agree that the specification for each overload is correct, and then we can look for opportunities to refactor things. |
keryell
pushed a commit
that referenced
this pull request
Sep 10, 2024
Clarify group synchronization behavior
gmlueck
added a commit
that referenced
this pull request
Nov 7, 2024
Clarify group synchronization behavior (cherry picked from commit fdc8a5e)
gmlueck
added a commit
that referenced
this pull request
Nov 7, 2024
Clarify group synchronization behavior (cherry picked from commit fdc8a5e)
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.
These changes build on #484 to clarify the expected synchronization behavior of group barriers and algorithms.
The wording proposed here deliberately does not specify how a SYCL implementation must satisfy certain guarantees. Previously, barriers were described in terms of fences and atomics, and we had discussed (offline) the possibility of defining algorithms similarly (or in terms of barriers). The approach taken by
std::barrierin ISO C++ is much simpler, giving implementations the freedom to use any combination of synchronization operations that guarantee certain events to happen before other events.For example, it should be legal to implement
sycl::group_barrierusing any of the following:std::atomic_thread_fence+ relaxedstd::atomic_ref, release atomics paired with acquire atomics, a call tostd::barrier, or any implementation-defined (backend-specific) functions that provide the desired memory ordering semantics.There's a lot of duplication here now, but there are two reasons for that:
I wanted some feedback on the wording, and specifically whether we would like the "Effects" for each algorithm to mention the algorithm by name.
I wanted to separate the review of the clarifications from any later formatting changes. The new table of overloads format that @gmlueck has been working on would allow us to combine some of these definitions.
Closes internal issue 576.