From 9576b3101e2ae475850d56ad03825f628ce0d7d9 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 25 Jun 2025 15:42:03 +0100 Subject: [PATCH 1/6] Add note about work-item scope atomics In OpenCL, these atomics are only required to support a very specific use-case involving images, and are forbidden in all other contexts. In SYCL, we would like a work-item to be viewed as a degenerate case of a group containing a single work-item. Work-item scope atomics should thus be permitted, and their effect should be equivalent to non-atomic operations. --- adoc/chapters/architecture.adoc | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index 4b06f72b3..718eba22c 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1089,6 +1089,13 @@ values: the memory allocation containing the referenced object, as defined by the capabilities of <> and <>. +{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 +<> containing a single work-item. +{endnote} + The memory scopes are listed above from narrowest ([code]#memory_scope::work_item#) to widest ([code]#memory_scope::system#). From cb816e236cc978ad0b797964f3d61e6a633e8ca0 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 25 Jun 2025 16:03:06 +0100 Subject: [PATCH 2/6] Clarify behavior of atomic_ref DefaultScope An implementation of atomic_ref that was not lock-free needs to know which work-items may access the lock in order to decide where to allocate the lock. This additionally serves as a clarification of the behavior of work-item scope; using an atomic_ref with work-item scope at the same time as an atomic_ref with broader scope is invalid. --- adoc/chapters/programming_interface.adoc | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/adoc/chapters/programming_interface.adoc b/adoc/chapters/programming_interface.adoc index 5c9dd5453..7e1679884 100644 --- a/adoc/chapters/programming_interface.adoc +++ b/adoc/chapters/programming_interface.adoc @@ -20340,6 +20340,10 @@ scope for the atomic operations. Most member functions also provide an optional parameter that allows the application to override this default. +All accesses to an object referenced by an [code]#sycl::atomic_ref# must +exclusively occur through instances of an [code]#sycl::atomic_ref# with the same +[code]#DefaultScope#. + The [code]#sycl::atomic_ref# class also has a template parameter [code]#AddressSpace#, which allows the application to make an assertion about the address space of the object of type [code]#T# that it references. From 114141496aed84436d8d744e5939a90302c9abe8 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 25 Jun 2025 16:34:24 +0100 Subject: [PATCH 3/6] Clarify that using different scopes is a data race Even when using two atomic_ref objects with the same DefaultScope, it's possible to encounter a data race by overriding the scope parameter of individual operations. This is a general clean-up but was motivated by work-item scope atomics: any potentially concurrent use of work-item scope atomics and atomics with a different scope results in undefined behavior. --- adoc/chapters/architecture.adoc | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index 718eba22c..99e78b5e0 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1096,6 +1096,9 @@ programming and to provide a meaningful way to describe the behavior of <> containing a single work-item. {endnote} +Potentially concurrent conflicting actions with different memory scopes are +considered a data race, resulting in undefined behavior. + The memory scopes are listed above from narrowest ([code]#memory_scope::work_item#) to widest ([code]#memory_scope::system#). @@ -1116,10 +1119,10 @@ 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: 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 <>. 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 0665b69ee60dabf897344db512abb63c8abc8e2a Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 25 Jun 2025 17:02:22 +0100 Subject: [PATCH 4/6] Clarify which atomic operations synchronize The ISO C++ synchronizes-with relationship does not account for scopes. The scopes do not need to match exactly, but there are restrictions on which pairs of scopes are valid. This is the final part of the clarification for work-item scope atomics; a work-item scope atomic cannot sychronize with the atomic operations performed by other work-items, and so their effects are not guaranteed to be visible to other work-items without some other synchronization taking place. --- adoc/chapters/architecture.adoc | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index 99e78b5e0..45d70a9c2 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1099,6 +1099,12 @@ programming and to provide a meaningful way to describe the behavior of Potentially concurrent conflicting actions with different memory scopes are considered a data race, resulting in undefined behavior. +An atomic operation _A_ with scope _S~1~_ can only synchronize with another +atomic operation _B_ with scope _S~2~_ if: + +* The work-item which executed _A_ is in both _S~1~_ and _S~2~_; and +* The work-item which executed _B_ is in both _S~1~_ and _S~2~_. + The memory scopes are listed above from narrowest ([code]#memory_scope::work_item#) to widest ([code]#memory_scope::system#). @@ -1119,10 +1125,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: 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 eb0b73a3bbdf507e797069c35c9cd0c9968825c7 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 26 Jun 2025 06:56:50 +0100 Subject: [PATCH 5/6] Fix a bug in description of data races with scopes The previously proposed wording suggested that any difference in scopes would lead to undefined behavior, which was inconsistent with the paragraph immediately afterwards about which atomics synchronize-with each other. --- adoc/chapters/architecture.adoc | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index 45d70a9c2..4a2e3c9ee 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1096,14 +1096,23 @@ programming and to provide a meaningful way to describe the behavior of <> containing a single work-item. {endnote} -Potentially concurrent conflicting actions with different memory scopes are -considered a data race, resulting in undefined behavior. +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-item which executed _A_ is in both _S~1~_ and _S~2~_; and -* The work-item which executed _B_ is in both _S~1~_ and _S~2~_. +* 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#). From 04a8e80cf06e21ab654b2f6d9d594aa1351a50ab Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Thu, 4 Sep 2025 12:46:34 -0400 Subject: [PATCH 6/6] Remove requirement on DefaultScope We decided that this requirement doesn't actually help implementations. --- adoc/chapters/programming_interface.adoc | 4 ---- 1 file changed, 4 deletions(-) diff --git a/adoc/chapters/programming_interface.adoc b/adoc/chapters/programming_interface.adoc index 7e1679884..5c9dd5453 100644 --- a/adoc/chapters/programming_interface.adoc +++ b/adoc/chapters/programming_interface.adoc @@ -20340,10 +20340,6 @@ scope for the atomic operations. Most member functions also provide an optional parameter that allows the application to override this default. -All accesses to an object referenced by an [code]#sycl::atomic_ref# must -exclusively occur through instances of an [code]#sycl::atomic_ref# with the same -[code]#DefaultScope#. - The [code]#sycl::atomic_ref# class also has a template parameter [code]#AddressSpace#, which allows the application to make an assertion about the address space of the object of type [code]#T# that it references.