Merge branch 'pub/main' into gmlueck/reformat-context

adoc/chapters/architecture.adoc

@@ -1389,7 +1389,7 @@ implementation-defined.
 
 
 [[sec:coordination]]
-=== Coordination and Synchronization
+=== Coordination and synchronization
 
 Coordination between the host and any devices can be expressed in the host SYCL
 application using calls into the SYCL runtime.
@@ -1403,7 +1403,7 @@ Such functions can be used to ensure that the host and any devices do not access
 data concurrently, and/or to reason about the ordering of operations across the
 host and any devices.
 
-==== Host-Device Coordination
+==== Host-device coordination
 
 The following operations can be used to coordinate host and device(s):
 
@@ -1456,7 +1456,7 @@ So it is up to the programmer to use a member function to wait for completion in
 some cases if this does not fit the goal.
 See <<sec:managing-object-lifetimes>> for more information on object life time.
 
-==== Work-item Coordination
+==== Work-item coordination
 
 A <<group-barrier>> provides a mechanism to coordinate all work-items in the
 same group.
@@ -1500,6 +1500,7 @@ Any error reported by a <<backend>> must derive from the base
 When a user wishes to capture specifically an error thrown by a <<backend>>, she
 must include the <<backend>>-specific headers for said <<backend>>.
 
+[[sec::fallback-mechanism]]
 === Fallback mechanism
 
 A <<command-group-function-object>> can be submitted either to a single queue to

adoc/chapters/device_compiler.adoc

@@ -205,6 +205,16 @@ Amongst other things, this restriction makes it illegal for a
     implementation defines the [code]#SYCL_EXTERNAL# macro as described in
     <<subsec:syclexternal>>.
 
+Inside a <<discarded-statement>> or in the case of a
+<<manifestly-constant-evaluated>>, any code accepted by the C++ standard in this
+case is also accepted in a SYCL <<device-function>>.
+
+[NOTE]
+====
+The restriction waiver in <<discarded-statement>> or
+<<manifestly-constant-evaluated>> allows any kind of meta-programming in a
+<<device-function>>.
+====
 
 [[subsec:scalartypes]]
 == Built-in scalar data types
@@ -377,8 +387,17 @@ in <<sec:backends>> must be defined by all conformant implementations.
 
 A number of kernel features defined by this SYCL specification are optional;
 they may be supported on some devices but not on other devices.
+
+As stated in <<sec:language.restrictions.kernels>>, the restrictions for
+optional kernel features do not apply to discarded statements or to manifestly
+constant-evaluated expressions or conversions in device code.
+Device code may use optional features in <<discarded-statement>> or
+<<manifestly-constant-evaluated>> even if the device does not support the
+optional feature.
+
 As described in <<sec:device-aspects>>, an application can test whether a device
-supports these features by testing whether the device has an associated aspect.
+supports an optional feature by testing whether the device has an associated
+aspect.
 The following aspects are those that correspond to optional kernel features:
 
   * [code]#fp16#

adoc/chapters/glossary.adoc

@@ -188,6 +188,12 @@ For the full description please refer to <<subsec:buffers>>.
     One of the device with the highest non-negative value is selected.
     See <<sec:device-selector>> for more details.
 
+[[discarded-statement]]discarded statement::
+    ISO C++ +[stmt.if]+ describes a discarded statement as the branch statement
+    of an [code]#if constexpr# which is not instantiated because of the boolean
+    condition.
+    For more context, see <<sec:language.restrictions.kernels>>.
+
 [[event]]event::
     A SYCL object that represents the status of an operation that is being
     executed by the SYCL runtime.
@@ -339,6 +345,19 @@ For the full description please refer to <<subsec:buffers>>.
     Local memory is a memory region associated with a <<work-group>> and
     accessible only by <<work-item,work-items>> in that <<work-group>>.
 
+[[manifestly-constant-evaluated]]manifestly constant-evaluated expression or conversion::
+    ISO C++ +[expr.const]+ describes manifestly constant-evaluated expression or
+    conversion like constant expressions, condition of an +if constexpr+, an
+    immediate invocation, used in template parameters, in constant
+    initialization, etc.
+    This is evaluated at compile-time by the compiler.
+    For more context, see <<sec:language.restrictions.kernels>>.
+
+[[mem-fence]]mem-fence::
+    A memory fence provides control over re-ordering of memory load and store
+    operations when coupled with an atomic operation.
+    See the definition of the [code]#sycl::atomic_fence# function.
+
 [[native-backend-object]]native backend object::
     An opaque object defined by a specific backend that represents a high-level
     SYCL object on said backend.
@@ -348,7 +367,6 @@ For the full description please refer to <<subsec:buffers>>.
     A <<specialization-constant>> in a device image whose value can be used by
     an online compiler as an immediate value during the compilation.
 
-
 [[nd-item]]nd-item::
     A unique identifier representing a single <<work-item>> and <<work-group>>
     within the index space of a SYCL kernel execution.
@@ -370,11 +388,6 @@ For the full description please refer to <<subsec:buffers>>.
     In the SYCL interface an <<nd-range>> is represented by the [code]#nd_range#
     class (see <<subsubsec:nd-range-class>>).
 
-[[mem-fence]]mem-fence::
-    A memory fence provides control over re-ordering of memory load and store
-    operations when coupled with an atomic operation.
-    See the definition of the [code]#sycl::atomic_fence# function.
-
 [[object]]object::
     A state which a <<kernel-bundle>> can be in, representing
     <<sycl-kernel-function,SYCL kernel functions>> as a non-executable object.
@@ -396,7 +409,7 @@ For the full description please refer to <<subsec:buffers>>.
     SYCL provides a heterogeneous platform integration using device queue, which
     is the minimum requirement for a SYCL application to run on a SYCL
     <<device>>.
-    For the full description please refer to <<sec:interface.queue.class>>.
+    For the full description please refer to <<sec:queue-class>>.
 
 [[range]]range::
     A representation of a number of <<work-item,work-items>> or

adoc/chapters/information_descriptors.adoc

@@ -33,7 +33,7 @@ include::{header_dir}/contextInfo.h[lines=4..-1]
 == Device information descriptors
 
 The following interface includes all the information descriptors for the
-[code]#device# class as described in <<table.device.info>>.
+[code]#device# class.
 [source,,linenums]
 ----
 include::{header_dir}/deviceInfo.h[lines=4..-1]
@@ -44,7 +44,7 @@ include::{header_dir}/deviceInfo.h[lines=4..-1]
 == Queue information descriptors
 
 The following interface includes all the information descriptors for the
-[code]#queue# class as described in <<table.queue.info>>.
+[code]#queue# class.
 [source,,linenums]
 ----
 include::{header_dir}/queueInfo.h[lines=4..-1]