Skip to content

Clarify work-item scope atomics (and memory model in general) #849

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.

Sign up for GitHub

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

Open
Pennycook wants to merge 6 commits into
base: main
Choose a base branch
from
Open

Clarify work-item scope atomics (and memory model in general) #849

Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
9576b31
Add note about work-item scope atomics
Pennycook Jun 25, 2025
cb816e2
Clarify behavior of atomic_ref DefaultScope
Pennycook Jun 25, 2025
1141414
Clarify that using different scopes is a data race
Pennycook Jun 25, 2025
0665b69
Clarify which atomic operations synchronize
Pennycook Jun 25, 2025
eb0b73a
Fix a bug in description of data races with scopes
Pennycook Jun 26, 2025
04a8e80
Remove requirement on DefaultScope
gmlueck Sep 4, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
32 changes: 28 additions & 4 deletions adoc/chapters/architecture.adoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1089,6 +1089,31 @@ values:
the memory allocation containing the referenced object, as defined by the
capabilities of <<buffer,buffers>> and <<usm>>.

{note}An atomic operation with work-item scope is effectively the same as a
non-atomic operation.
[code]#sycl::memory_scope::work_item# is primarily intended to simplify generic
programming and to provide a meaningful way to describe the behavior of
<<group,groups>> containing a single work-item.
{endnote}

Potentially concurrent conflicting actions with different memory scopes may lead
to a data race, resulting in undefined behavior.
An atomic operation _A_ with scope _S~1~_ operating on the same memory location
as atomic operation _B_ with scope _S~2~_ is a data race if:

* The work-items which executed _A_ and _B_ are not both in the same group of
work-items associated with scope _S~1~_; or
* The work-items which executed _A_ and _B_ are not both in the same group of
work-items associated with scope _S~2~_.
Copy link
Contributor

@gmlueck gmlueck Sep 3, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

John and I were having a side discussion about this part of the PR before he left. The question is whether operations A and B need to have the same scope, or whether it is sufficient for the scopes to include both work-items. To illustrate, consider the following example:

// Work-item A
sycl::atomic_ref<int, memory_order::release, memory_scope::work_group> a(mem);
a.store(1);

// Work-item B (in the same work-group a A)
sycl::atomic_ref<int, memory_order::acquire, memory_scope::device> a(mem);
int x = a.load();

Note that the two operations have different scopes, but each scope includes both A and B.

The question is whether SYCL should guarantee that these operations are atomic even though the scopes are different. My first question to John was about Intel hardware. We think that Intel hardware is guaranteed to be atomic in this scenario, so we have no concerns from the standpoint of our own ability to implement the proposed SYCL wording.

However, then we realized that the OpenCL specification seems to not guarantee atomicity in this case. Instead, the OpenCL wording seems to require both operations to have the same scope in order to guarantee atomicity. There is some debate, though, about whether the OpenCL wording should be changed. There is an open internal issue against the OpenCL specification on this point:

https://gitlab.khronos.org/opencl/OpenCL-Docs/-/issues/367

The SYCL WG should consider whether we want to adopt the wording that John proposes in this PR even though it guarantees atomicity in a case that OpenCL does not guarantee. Or, whether we should adopt the same language about atomicity that is currently in the OpenCL spec,


An atomic operation _A_ with scope _S~1~_ can only synchronize with another
atomic operation _B_ with scope _S~2~_ if:

* The work-items which executed _A_ and _B_ are both in the same group of
work-items associated with scope _S~1~_; and
* The work-items which executed _A_ and _B_ are both in the same group of
work-items associated with scope _S~2~_.

The memory scopes are listed above from narrowest
([code]#memory_scope::work_item#) to widest ([code]#memory_scope::system#).

Expand All @@ -1109,10 +1134,9 @@ supplied.
====
The addition of memory scopes to the {cpp} memory model modifies the definition
of some concepts from the {cpp} core language.
For example: data races, the synchronizes-with relationship and sequential
consistency must be defined in a way that accounts for atomic operations with
differing (but compatible) scopes, in a manner similar to the <<opencl20, OpenCL
2.0 specification>>.
For example: sequential consistency must be defined in a way that accounts for
atomic operations with differing (but compatible) scopes, in a manner similar to
the <<opencl20, OpenCL 2.0 specification>>.
Efforts to formalize the memory model of SYCL are ongoing, and a formal memory
model will be included in a future version of the SYCL specification.
====
Expand Down