diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index cb45c8852..ddfc876cc 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -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 <>. -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. +From SYCL {SYCL_VERSION} the memory scope of _A_ and _B_ may be different. ==== 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++ ",