From 29e1f911ae883bb35c3ab8a675d116c29527b339 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Thu, 27 Aug 2026 11:10:57 -0400 Subject: [PATCH 1/2] Clarify atomics memory model Cherry pick #1032 from main (cherry picked from commit 4dc10b05a2f9e72b4c10c9316c279bad91b86732) --- adoc/chapters/architecture.adoc | 39 +++++++++++++++++++++++---------- adoc/chapters/references.adoc | 5 ----- 2 files changed, 28 insertions(+), 16 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index 1378d9339..3d0f3dec3 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1129,17 +1129,34 @@ 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. + +{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.{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++ ", From 6184a01ed82b58f27908ac9ee782e33551214520 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Fri, 28 Aug 2026 08:16:08 -0400 Subject: [PATCH 2/2] Edit out SYCL-Next information The previous commit was cherry picked from the SYCL-Next branch, and it contains some information specific to SYCL-Next. Edit out this information, leaving only the parts relevant to SYCL 2020. --- adoc/chapters/architecture.adoc | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/adoc/chapters/architecture.adoc b/adoc/chapters/architecture.adoc index 3d0f3dec3..44fc8211c 100644 --- a/adoc/chapters/architecture.adoc +++ b/adoc/chapters/architecture.adoc @@ -1148,15 +1148,13 @@ Atomic operations _A_ and _B_ are defined to have an inclusive scope if: 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 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.{endnote} +Implementations of SYCL 2020 using an OpenCL backend may use the older OpenCL +semantic.{endnote} ==== Atomic operations