diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index 1378d9339..44fc8211c 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1129,17 +1129,32 @@ 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 <>. -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 + 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. + +{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. +Implementations of SYCL 2020 using an OpenCL backend may use the older OpenCL +semantic.{endnote} ==== 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++ ",