Skip to content

Clarify atomics memory model #1032

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
VerenaBeckham wants to merge 5 commits into
base: main
Choose a base branch
from
+20 −16
Open

Clarify atomics memory model #1032

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
d72bd6e
Original text.
VerenaBeckham Jul 28, 2026
5e79627
Rewording from OpenCL.
VerenaBeckham Jul 28, 2026
f35afa3
Remove reference to OpenCL spec.
VerenaBeckham Aug 4, 2026
bfd9fdd
Remove new section down.
VerenaBeckham Aug 4, 2026
c02e182
Remove the rest of the note.
VerenaBeckham Aug 11, 2026
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
31 changes: 20 additions & 11 deletions adoc/chapters/architecture.adoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1102,17 +1102,26 @@ If a wider scope is supplied, the behavior is as-if the narrowest scope
containing all work-items which can access the associated memory location was
supplied.

[NOTE]
====
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>>.
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.
====
Atomic operations on the same memory location must have inclusive scope,
otherwise the behaviour is undefined.
Atomic operations _A_ and _B_ are defined to have an inclusive scope if:

* _A_ and _B_ are executed by work-items within the same sub-group and the
memory scope of _A_ and _B_ is [code]#sycl::memory_scope::sub_group# or
wider.
* _A_ and _B_ are executed by work-items within the same work-group and the
memory scope of _A_ and _B_ is [code]#sycl::memory_scope::work_group# or
wider.
* _A_ and _B_ are executed by work-items in the same kernel instance, and the
Comment thread
VerenaBeckham marked this conversation as resolved.
memory scope of _A_ and _B_ is [code]#sycl::memory_scope::device# or wider.
* _A_ and _B_ are executed by work-items on the same device, and the memory
scope of _A_ and _B_ is [code]#sycl::memory_scope::device# or wider.
* _A_ and _B_ are executed by a host thread or by work-items, and the memory
scope of _A_ and _B_ is [code]#sycl::memory_scope::system#.

In SYCL 2020 the memory scope of _A_ and _B_ must additionally be the same for
the operations to have an inclusive scope.
From SYCL {SYCL_VERSION} the memory scope of _A_ and _B_ may be different.

@gmlueck gmlueck Aug 11, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

I wonder if we should add a non-normative note here to explain the rationale. Something like:

[Note: The SYCL definition of inclusive scope follows the OpenCL definition. In OpenCL versions prior to 3.1, the memory scopes of A and B must be the same for their scopes to be inclusive, but this was relaxed in OpenCL version 3.1. As a result, implementations of SYCL 2020 using an OpenCL backend may use the older OpenCL semantic, but implementations of {SYCL_VERSION} must use the OpenCL 3.1 semantic. -- end note]


==== Atomic operations

Expand Down
5 changes: 0 additions & 5 deletions adoc/chapters/references.adoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,11 +21,6 @@ Khronos OpenCL Working Group.
_The OpenCL Specification, Version 1.2.19_ (11/14/12).
https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf .

[[opencl20]]
Khronos OpenCL Working Group.
_The OpenCL Specification, Version 2.0.29_ (July 21, 2015).
https://www.khronos.org/registry/OpenCL/specs/opencl-2.0.pdf .

[[cpp20]]
International Organization for Standardization (ISO).
" Programming Languages — {cpp}, Langages de programmation — C++ ",
Expand Down
Loading