From d72bd6ec45e39da862532fb84bb117c208bb4f18 Mon Sep 17 00:00:00 2001 From: vbeckham Date: Tue, 28 Jul 2026 11:26:42 +0100 Subject: [PATCH 1/5] Original text. --- adoc/chapters/architecture.adoc | 25 +++++++++++++++++++++---- 1 file changed, 21 insertions(+), 4 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index cb45c8852..b05ff60e6 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1086,6 +1086,24 @@ values: the memory allocation containing the referenced object, as defined by the capabilities of <> and <>. +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~_. + +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#). @@ -1106,10 +1124,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 <>. +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 <>. 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. ==== From 5e79627e54c284f521698a869356a4b3771b2183 Mon Sep 17 00:00:00 2001 From: vbeckham Date: Tue, 28 Jul 2026 11:58:13 +0100 Subject: [PATCH 2/5] Rewording from OpenCL. --- adoc/chapters/architecture.adoc | 37 ++++++++++++++++++--------------- 1 file changed, 20 insertions(+), 17 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index b05ff60e6..0fb5b21d0 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1086,23 +1086,26 @@ values: the memory allocation containing the referenced object, as defined by the capabilities of <> and <>. -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~_. - -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~_. +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 + 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. The memory scopes are listed above from narrowest ([code]#memory_scope::work_item#) to widest ([code]#memory_scope::system#). From f35afa36ae72b412cfdb9333dacfe53bb3616d42 Mon Sep 17 00:00:00 2001 From: vbeckham Date: Tue, 4 Aug 2026 10:08:08 +0100 Subject: [PATCH 3/5] Remove reference to OpenCL spec. --- adoc/chapters/architecture.adoc | 5 ----- adoc/chapters/references.adoc | 5 ----- 2 files changed, 10 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index 0fb5b21d0..ee23a6891 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1127,11 +1127,6 @@ supplied. ==== The addition of memory scopes to the {cpp} memory model modifies the definition of some concepts from the {cpp} core language. -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 <>. -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 diff --git a/adoc/chapters/references.adoc b/adoc/chapters/references.adoc index 542d14f12..4d74bde4d 100644 --- a/adoc/chapters/references.adoc +++ b/adoc/chapters/references.adoc @@ -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++ ", From bfd9fddba7b69a8d6a96b2c56ef0da2019e4614a Mon Sep 17 00:00:00 2001 From: vbeckham Date: Tue, 4 Aug 2026 10:12:14 +0100 Subject: [PATCH 4/5] Remove new section down. --- adoc/chapters/architecture.adoc | 32 ++++++++++++++++---------------- 1 file changed, 16 insertions(+), 16 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index ee23a6891..2dcd37693 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1086,6 +1086,22 @@ values: the memory allocation containing the referenced object, as defined by the capabilities of <> and <>. +The memory scopes are listed above from narrowest +([code]#memory_scope::work_item#) to widest ([code]#memory_scope::system#). + +The complete set of memory scopes is not guaranteed to be supported by every +device. +The set of supported memory scopes can be queried via the information +descriptors for the [code]#sycl::device# and [code]#sycl::context# classes. + +The widest scope that can be applied to an atomic operation corresponds to the +set of work-items which can access the associated memory location. +For example, the widest scope that can be applied to atomic operations in +work-group local memory is [code]#sycl::memory_scope::work_group#. +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. + 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: @@ -1107,22 +1123,6 @@ 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. -The memory scopes are listed above from narrowest -([code]#memory_scope::work_item#) to widest ([code]#memory_scope::system#). - -The complete set of memory scopes is not guaranteed to be supported by every -device. -The set of supported memory scopes can be queried via the information -descriptors for the [code]#sycl::device# and [code]#sycl::context# classes. - -The widest scope that can be applied to an atomic operation corresponds to the -set of work-items which can access the associated memory location. -For example, the widest scope that can be applied to atomic operations in -work-group local memory is [code]#sycl::memory_scope::work_group#. -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 From c02e1823a8e616a5c5d87e67bc29759cb2126010 Mon Sep 17 00:00:00 2001 From: vbeckham Date: Tue, 11 Aug 2026 12:50:14 +0100 Subject: [PATCH 5/5] Remove the rest of the note. --- adoc/chapters/architecture.adoc | 6 ------ 1 file changed, 6 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index 2dcd37693..ddfc876cc 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1123,12 +1123,6 @@ 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. -[NOTE] -==== -The addition of memory scopes to the {cpp} memory model modifies the definition -of some concepts from the {cpp} core language. -==== - ==== Atomic operations Atomic operations can be performed on memory in buffers and USM.