From 655d35b70e7ab079fa8ec1816ca3fa17423db284 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 24 Mar 2020 10:20:18 -0700 Subject: [PATCH 1/8] [SYCL][Doc] Update ExtendedAtomics documentation Several major changes to the extension: - No longer replaces cl::sycl::atomic - Closer alignment with C++20 atomic_ref - Supports additional memory orders Signed-off-by: John Pennycook --- sycl/doc/extensions/ExtendedAtomics/README.md | 2 +- .../SYCL_INTEL_extended_atomics.asciidoc | 566 +++++++++--------- 2 files changed, 283 insertions(+), 285 deletions(-) mode change 100755 => 100644 sycl/doc/extensions/ExtendedAtomics/SYCL_INTEL_extended_atomics.asciidoc diff --git a/sycl/doc/extensions/ExtendedAtomics/README.md b/sycl/doc/extensions/ExtendedAtomics/README.md index c9037fd37430b..66431c1c16dc1 100644 --- a/sycl/doc/extensions/ExtendedAtomics/README.md +++ b/sycl/doc/extensions/ExtendedAtomics/README.md @@ -1,3 +1,3 @@ # SYCL_INTEL_extended_atomics -Replaces the `cl::sycl::atomic` class with the `cl::sycl::intel::atomic` class, which exposes additional functionality aligned with the `std::atomic` class from C++11 and C++20. +Introduces the `cl::sycl::intel::atomic_ref` class, which exposes additional functionality aligned with the `std::atomic_ref` class from C++20. diff --git a/sycl/doc/extensions/ExtendedAtomics/SYCL_INTEL_extended_atomics.asciidoc b/sycl/doc/extensions/ExtendedAtomics/SYCL_INTEL_extended_atomics.asciidoc old mode 100755 new mode 100644 index 3bab2c843087d..51bd1243f1582 --- a/sycl/doc/extensions/ExtendedAtomics/SYCL_INTEL_extended_atomics.asciidoc +++ b/sycl/doc/extensions/ExtendedAtomics/SYCL_INTEL_extended_atomics.asciidoc @@ -27,7 +27,7 @@ NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are tradema NOTE: This document is better viewed when rendered as html with asciidoctor. GitHub does not render image icons. -This document describes an extension that replaces the +cl::sycl::atomic+ class with the `cl::sycl::intel::atomic` class, which exposes additional functionality aligned with the +std::atomic+ class from {cpp}11 and {cpp}20. +This document describes an extension that introduces the `cl::sycl::intel::atomic_ref` class, which exposes additional functionality aligned with the +std::atomic_ref+ class from {cpp}20. == Name Strings @@ -55,56 +55,58 @@ John Pennycook, Intel (john 'dot' pennycook 'at' intel 'dot' com) == Dependencies -This extension is written against the SYCL 1.2.1 specification, Revision v1.2.1-6. +This extension is written against the SYCL 1.2.1 specification, Revision 6. == Overview -The SYCL atomic library (+cl::sycl::atomic+) defined in SYCL 1.2.1 is based on the standard atomic libary (+std::atomic+) but has some differences. This extension introduces an alternative atomic class (+cl::sycl::intel::atomic+) including additional features from {cpp}11 and {cpp}20: +The SYCL atomic library (+cl::sycl::atomic+) defined in SYCL 1.2.1 is based on the standard atomic libary (+std::atomic+) but has some differences. This extension introduces an alternative atomic class (+cl::sycl::intel::atomic_ref+) including additional features from {cpp}20: - Overloaded operators to reduce the verbosity of using atomics +- Missing functions (e.g. `is_lock_free()`) - Support for floating-point types +- Support for additional memory orderings besides `relaxed` + +This extension deprecates the SYCL 1.2.1 +cl::sycl::atomic+ class and accessors created with mode +access::mode::atomic+. The extension can be enabled using the `-fsycl-extended-atomics` flag, and applications can check whether the extension is enabled using `__has_extension(sycl_extended_atomics)`. === Overloaded Operators -In SYCL 1.2.1, the +cl::sycl::atomic+ class provides atomic operations by way of member functions (e.g. +fetch_add+) without defining the corresponding operators (e.g. `+=`). This increases the verbosity of simple uses of atomics, and requires developers to change their kernel code when switching between +read_write+ and +atomic+ accessors: +In SYCL 1.2.1, the +cl::sycl::atomic+ class provides atomic operations by way of member functions (e.g. +fetch_add+) without defining the corresponding operators (e.g. `+=`). This increases the verbosity of simple uses of atomics. + +The operators defined by this extension match those defined for +std::atomic_ref+ in {cpp}20. The functionality of each operator is equivalent to calling a corresponding member function of the +atomic_ref+ class -- the operators do not expose any new functionality of the class, but act as shorthands for common use-cases. + +==== Operators for All Supported Types [source,c++] ---- -q.submit([&](handler& cgh) -{ - auto acc = buf.get_access(cgh); - cgh.parallel_for(..., [=](id<1> i) - { - acc[i].fetch_add(1); // developer cannot write acc[i]++ without this extension - }); -}); +operator T(); // equivalent to load() +T operator=(T desired); // equivalent to store(desired) +T operator+=(T operand); // equivalent to fetch_add(operand) +T operator-=(T operand); // equivalent to fetch_sub(operand) ---- -The operators defined by this extension match those defined for +std::atomic+ in {cpp}11. The functionality of each operator is equivalent to calling a corresponding member function of the atomic class -- the operators do not expose any new functionality of the class, but act as shorthands for common use-cases. - -==== Operators for All Supported Types +==== Operators for Integral Types [source,c++] ---- - operator T(); - T operator=(T desired); - T operator+=(T operand); - T operator-=(T operand); +T operator++(int); // equivalent to fetch_add(1) +T operator--(int); // equivalent to fetch_sub(1) +T operator++(); // equivalent to fetch_add(1) + 1 +T operator--(); // equivalent to fetch_sub(1) - 1 +T operator&=(T operand); // equivalent to fetch_and(operand) +T operator|=(T operand); // equivalent to fetch_or(operand) +T operator^=(T operand); // equivalent to fetch_xor(operand) ---- -==== Operators for Integral Types +==== Operators for Pointer Types [source,c++] ---- - T operator++(int operand); - T operator--(int operand); - T operator++(); - T operator--(); - T operator&=(T operand); - T operator|=(T operand); - T operator^=(T operand); +T operator++(int); // equivalent to fetch_add(1) +T operator--(int); // equivalent to fetch_sub(1) +T operator++(); // equivalent to fetch_add(1) + 1 +T operator--(); // equivalent to fetch_sub(1) - 1 ---- === Support for Floating-point Types @@ -113,389 +115,384 @@ In SYCL 1.2.1, support for floating-point types is limited to the +load+, +store This extension extends support for floating-point types to the +compare_exchange+, +fetch_add+ and +fetch_sub+ functions in line with {cpp}20, as well as the +fetch_min+ and +fetch_max+ functions. These new functions do not require dedicated floating-point atomic instructions and can be emulated using integer operations, giving compilers the freedom to choose the best implementation for the target device. -== Modifications of SYCL 1.2.1 Specification - -=== Modify Sentence in Section 3.5.1 - -==== From: +=== Support for Additional Memory Orderings -Atomic access can also be requested on an accessor which allows +cl::sycl::atomic+ classes to be used via the accessor. +The atomic operations in SYCL 1.2.1 default to +memory_order_relaxed+, which is inconsistent with the default of +memory_order_seq_cst+ used by the +std::atomic+ class. Defaulting to +memory_order_relaxed+ may improve the performance and portability of SYCL 1.2.1 code across multiple target devices, but may also lead to unexpected behavior when code is migrated between {cpp} and SYCL. Different users have different understandings of which memory orders are the most common or useful, and the performance difference between memory orders is also expected to vary between devices. This extension therefore makes the default memory order of +cl::sycl::intel::atomic_ref+ dependent upon a template argument that must be specified by the user. -==== To: +All devices must support +memory_order_relaxed+, and the host device must support all {cpp} memory orders. These changes bring the SYCL memory model in line with modern {cpp} while allowing a device/compiler to implement only a subset of {cpp} memory orders. Supporting the standard {cpp} memory model in SYCL requires that disjoint address spaces (e.g. local and global memory) are treated as though they are part of a single address space (i.e. there must be a single happens-before relationship for all addresses). -Atomic access can also be requested on an accessor which allows +cl::sycl::intel::atomic+ classes to be used via the accessor. +=== The +atomic_ref+ Class -=== Modify Paragraph in Section 3.5.2.3 +The +cl::sycl::intel::atomic_ref+ class is constructed from a reference, and enables atomic operations to the referenced object. If any non-atomic access to the referenced object is made during the lifetime of the +cl::sycl::intel::atomic_ref+ class then the behavior is undefined. No subobject of the object referenced by an +atomic_ref+ shall be concurrently referenced by any other +atomic_ref+ object. -==== From: +The static member +required_alignment+ describes the minimum required alignment in bytes of an object that can be referenced by an +atomic_ref+, which must be at least +alignof(T)+. -Atomic operations can be performed on memory in buffers. The range of atomic operations available on a specific OpenCL device is limited by the atomic capabilities of that device. The +cl::sycl::atomic+ must be used for elements of a buffer to provide safe atomic access to the buffer from device code. +The static member +is_always_lock_free+ is true if all atomic operations for type +T+ are always lock-free. A SYCL implementation is not guaranteed to support atomic operations that are not lock-free. -==== To: +The static members +default_read_order+, +default_write_order+ and +default_read_modify_write_order+ reflect the default memory order values for each type of atomic operation, consistent with the +DefaultOrder+ template. -Atomic operations can be performed on memory. The range of atomic operations available on a specific OpenCL device is limited by the atomic capabilities of that device. The +cl::sycl::intel::atomic+ class may be used to provide safe atomic access to any memory location, in host or device code. +The member functions below are common to atomic references for any type +T+: -=== Modify Paragraph in Section 4.2 +|=== +|Member Functions|Description -==== From: +| `atomic_ref(T& ref)` +| Constructs an instance of +atomic_ref+ which is associated with the reference _ref_. -Each of the following SYCL runtime classes: +accessor+, +sampler+, +stream+, +vec+, +multi_ptr+, +device_event+, +id+, +range+, +item+, +nd_item+, +h_item+, +group+ and +atomic+ must be available within a SYCL kernel function. +| `bool is_lock_free() const` +| Return +true+ if the atomic operations provided by this +atomic_ref+ are lock-free. -==== To: +| `void store(T operand, memory_order order = default_write_order) const` +| Atomically stores _operand_ to the object referenced by this +atomic_ref+. The memory order of this atomic operation must be +memory_order::relaxed+, +memory_order::release+ or +memory_order::seq_cst+. -Each of the following SYCL runtime classes: +accessor+, +sampler+, +stream+, +vec+, +multi_ptr+, +device_event+, +id+, +range+, +item+, +nd_item+, +h_item+, +group+ and +intel::atomic+ must be available within a SYCL kernel function. +| `T operator=(T desired) const` +| Equivalent to +store(desired)+. Returns _desired_. -=== Modify Paragraph in Section 4.7.6.5 +| `T load(memory_order order = default_read_order) const` +| Atomically loads the value of the object referenced by this +atomic_ref+. The memory order of this atomic operation must be +memory_order::relaxed+, +memory_order::acquire+, or +memory_order::seq_cst+. -==== From: +| `operator T() const` +| Equivalent to +load()+. -A buffer accessor with access target +access::target::global_buffer+ can optionally provide atomic access to a SYCL buffer, using the access mode +access::mode::atomic+, in which case all operators which return an element of the SYCL buffer return an instance of the SYCL atomic class. +| `T exchange(T operand, memory_order order = default_read_modify_write_order)` +| Atomically replaces the value of the object referenced by this +atomic_ref+ with _operand_ and returns the original value of the referenced object. -==== To: +| `bool compare_exchange_weak(T &expected, T desired, memory_order success = default_read_modify_write_order, memory_order failure = default_read_order) const` +| Atomically compares the value of the object referenced by this +atomic_ref+ against the value of _expected_. If the values are equal attempts to replace the value of the referenced object with the value of +desired+, otherwise assigns the original value of the referenced object to _expected_. Returns +true+ if the comparison operation and replacement operation were successful. The _failure_ memory order of this atomic operation must be +memory_order::relaxed+, +memory_order::acquire+ or +memory_order::seq_cst+. -A buffer accessor with access target +access::target::global_buffer+ can optionally provide atomic access to a SYCL buffer, using the access mode +access::mode::atomic+, in which case all operators which return an element of the SYCL buffer return an instance of the +cl::sycl::intel::atomic+ class. +| `bool compare_exchange_weak(T &expected, T desired, memory_order order = default_read_modify_write_order) const` +| Equivalent to +compare_exchange_weak(expected, desired, order, order)+. -=== Modify Listing 4.1 +| `bool compare_exchange_strong(T &expected, T desired, memory_order success = default_read_modify_write_order, memory_order failure = default_read_order) const` +| Atomically compares the value of the object referenced by this +atomic_ref+ against the value of _expected_. If the values are equal replaces the value of the referenced object with the value of +desired+, otherwise assigns the original value of the referenced object to _expected_. Returns +true+ if the comparison operation was successful. The _failure_ memory order of this atomic operation must be +memory_order::relaxed+, +memory_order::acquire+ or +memory_order::seq_cst+. -==== From: +| `bool compare_exchange_strong(T &expected, T desired, memory_order order = default_read_modify_write_order) const` +| Equivalent to +compare_exchange_strong(expected, desired, order, order)+. -[source,c++] ----- -/* Available only when: accessMode == access::mode::atomic && dimensions == 0 */ -operator atomic () const; +|=== -/* Available only when: accessMode == access::mode::atomic && dimensions > 0 */ -atomic operator[](id index) const; +The additional member functions below are available for atomic references to integral types: -/* Available only when: accessMode == access::mode::atomic && dimensions == 1 */ -atomic operator[](size_t index) const; ----- +|=== +|Member Functions|Description -==== To: +| `T fetch_add(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically adds _operand_ to the value of the object referenced by this +atomic_ref+ and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -[source,c++] ----- -/* Available only when: accessMode == access::mode::atomic && dimensions == 0 */ -operator intel::atomic () const; +| `T operator+=(T operand) const` +| Equivalent to +fetch_add(operand)+. -/* Available only when: accessMode == access::mode::atomic && dimensions > 0 */ -intel::atomic operator[](id index) const; +| `T operator++(int) const` +| Equivalent to +fetch_add(1)+. -/* Available only when: accessMode == access::mode::atomic && dimensions == 1 */ -intel::atomic operator[](size_t index) const; ----- +| `T operator++() const` +| Equivalent to +fetch_add(1) + 1+. -=== Modify Table 4.46 +| `T fetch_sub(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically subtracts _operand_ from the value of the object referenced by this +atomic_ref+ and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -==== Replace each instance of: +| `T operator-=(T operand) const` +| Equivalent to +fetch_sub(operand)+. -+atomic+ +| `T operator--(int) const` +| Equivalent to +fetch_sub(1)+. -==== With: +| `T operator--() const` +| Equivalent to +fetch_sub(1) + 1+. -+intel::atomic+ +| `T fetch_and(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically performs a bitwise AND between _operand_ and the value of the object referenced by this +atomic_ref+, and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -=== Modify Paragraph in Section 4.7.6.7 +| `T operator&=(T operand) const` +| Equivalent to +fetch_and(operand)+. -==== From: +| `T fetch_or(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically performs a bitwise OR between _operand_ and the value of the object referenced by this +atomic_ref+, and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -A local accessor can optionally provide atomic access to allocated memory, using the access mode +access::mode::atomic+, in which case all operators which return an element of the allocated memory return an instance of the SYCL atomic class. +| `T operator\|=(T operand) const` +| Equivalent to +fetch_or(operand)+. -==== To: +| `T fetch_xor(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically performs a bitwise XOR between the value +operand+ and the value of the object referenced by this +atomic_ref+, and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -A local accessor can optionally provide atomic access to allocated memory, using the access mode +access::mode::atomic+, in which case all operators which return an element of the allocated memory return an instance of the +cl::sycl::intel::atomic+ class. +| `T operator^=(T operand) const` +| Equivalent to +fetch_xor(operand)+. -=== Modify Listing 4.2 +| `T fetch_min(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically computes the minimum of _operand_ and the value of the object referenced by this +atomic_ref+, and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -==== From: +| `T fetch_max(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically computes the maximum of _operand_ and the value of the object referenced by this +atomic_ref+, and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -[source,c++] ----- -/* Available only when: accessMode == access::mode::atomic && dimensions == 0 */ -operator atomic () const; +|=== -/* Available only when: accessMode == access::mode::atomic && dimensions > 0 */ -atomic operator[](id index) const; +The additional member functions below are available for atomic references to floating-point types: -/* Available only when: accessMode == access::mode::atomic && dimensions == 1 */ -atomic operator[](size_t index) const; ----- +|=== +| Member Function | Description -===== To: +| `T fetch_add(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically adds _operand_ to the value of the object referenced by this +atomic_ref+ and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -[source,c++] ----- -/* Available only when: accessMode == access::mode::atomic && dimensions == 0 */ -operator intel::atomic () const; +| `T operator+=(T operand) const` +| Equivalent to +fetch_add(operand)+. -/* Available only when: accessMode == access::mode::atomic && dimensions > 0 */ -intel::atomic operator[](id index) const; +| `T fetch_sub(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically subtracts _operand_ from the value of the object referenced by this +atomic_ref+ and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -/* Available only when: accessMode == access::mode::atomic && dimensions == 1 */ -intel::atomic operator[](size_t index) const; ----- +| `T operator-=(T operand) const` +| Equivalent to +fetch_sub(operand)+. -=== Modify Table 4.49 +| `T fetch_min(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically computes the minimum of _operand_ and the value of the object referenced by this +atomic_ref+, and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -==== Replace each instance of: +| `T fetch_max(T operand, memory_order order = default_read_modify_write_order) const` +| Atomically computes the maximum of _operand_ and the value of the object referenced by this +atomic_ref+, and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -+atomic+ +|=== -==== With: +The additional member functions below are available for atomic references to pointer types: -+intel::atomic+ +|=== +| Member Function | Description -=== Modify Section 4.11 +| `T* fetch_add(ptrdiff_t operand, memory_order order = default_read_modify_write_order) const` +| Atomically adds _operand_ to the value of the object referenced by this +atomic_ref+ and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -==== From: +| `T* operator+=(ptrdiff_t operand) const` +| Equivalent to +fetch_add(operand)+. -The SYCL specification provides atomic operations based on the {cpp}11 library syntax. The only available ordering, due to constraints of the OpenCL 1.2 memory model, is +memory_order_relaxed+. No default order is supported because a default order would imply sequential consistency. The SYCL atomic library may map directly to the underlying {cpp}11 library in host code, and must interact safely with the host {cpp}11 atomic library when used in host code. The SYCL library must be used in device code to ensure that only the limited subset of functionality is available. SYCL 1.2.1 device compilers should give a compilation error on use of the +std::atomic+ classes and functions in device code. +| `T* operator++(int) const` +| Equivalent to +fetch_add(1)+. -The template parameter +addressSpace+ is permitted to be +access::address_space::global_space+ or +access::address_space::local_space+. +| `T* operator++() const` +| Equivalent to +fetch_add(1) + 1+. -The data type +T+ is permitted to be +int+, +unsigned int+, +long+, +unsigned long+, +long long+, +unsigned long long+ and +float+. Though +float+ is only available for the +store+, +load+ and +exchange+ member functions. For any data type +T+ which is 64 bit, the member functions of the atomic class are required to compile however are only guaranteed to execute if the 64 bit atomic extension +cl_khr_int64_base_atomics+ or +cl_khr_int64_extended_atomics+ (depending on which extension provides support for each given member function) is supported by the SYCL device which is executing the SYCL kernel function. If a member function is called with a 64 bit data type and the necessary extension is not supported by the SYCL device which is executing the SYCL kernel function, the SYCL runtime must throw a SYCL feature_not_supported exception. For more detail see Section 5.2. +| `T* fetch_sub(ptrdiff_t operand, memory_order order = default_read_modify_write_order) const` +| Atomically subtracts _operand_ from the value of the object referenced by this +atomic_ref+ and assigns the result to the value of the referenced object. Returns the original value of the referenced object. -==== To: +| `T* operator-=(ptrdiff_t operand) const` +| Equivalent to +fetch_sub(operand)+. -The SYCL specification provides atomic operations based on the {cpp}11 library syntax. The only available ordering, due to constraints of the OpenCL 1.2 memory model, is +memory_order_relaxed+. No default order is supported because a default order would imply sequential consistency. The SYCL atomic library may map directly to the underlying {cpp}11 library in host code, and must interact safely with the host {cpp}11 atomic library when used in host code. The SYCL library must be used in device code to ensure that only the limited subset of functionality is available. SYCL 1.2.1 device compilers should give a compilation error on use of the +std::atomic+ classes and functions in device code. +| `T* operator--(int) const` +| Equivalent to +fetch_sub(1)+. -The template parameter +addressSpace+ is permitted to be +access::address_space::global_space+ or +access::address_space::local_space+. +| `T* operator--() const` +| Equivalent to +fetch_sub(1) + 1+. -The data type +T+ is permitted to be +int+, +unsigned int+, +long+, +unsigned long+, +long long+, +unsigned long long+, +float+ or +double+. For any data type +T+ which is 64 bit, the member functions of the atomic class are required to compile however are only guaranteed to execute if the 64 bit atomic extension +cl_khr_int64_base_atomics+ or +cl_khr_int64_extended_atomics+ (depending on which extension provides support for each given member function) is supported by the SYCL device which is executing the SYCL kernel function. For +float+ and +double+, the member functions of the atomic class may be emulated, and may use a different floating-point environment to those defined by +info::device::single_fp_config+ and +info::device::double_fp_config+ (i.e. floating-point atomics may use different rounding modes and may have different exception behavior). If a member function is called with a 64 bit data type and the necessary extension is not supported by the SYCL device which is executing the SYCL kernel function, the SYCL runtime must throw a SYCL +feature_not_supported+ exception. For more detail see Section 5.2. +|=== -==== From: +==== Sample Header [source,c++] ---- namespace cl { namespace sycl { -enum class memory_order : int { - relaxed +namespace intel { +enum class memory_order : /* unspecified */ { + relaxed, acquire, release, acq_rel, seq_cst }; -template -class atomic { - public: - template - atomic(multi_ptr ptr); - - void store(T operand, memory_order memoryOrder = - memory_order::relaxed); - - T load(memory_order memoryOrder = memory_order::relaxed) const; - - T exchange(T operand, memory_order memoryOrder = - memory_order::relaxed); - - /* Available only when: T != float */ - bool compare_exchange_strong(T &expected, T desired, - memory_order successMemoryOrder = memory_order::relaxed, - memory_order failMemoryOrder = memory_order::relaxed); - - /* Available only when: T != float */ - T fetch_add(T operand, memory_order memoryOrder = - memory_order::relaxed); - - /* Available only when: T != float */ - T fetch_sub(T operand, memory_order memoryOrder = - memory_order::relaxed); - - /* Available only when: T != float */ - T fetch_and(T operand, memory_order memoryOrder = - memory_order::relaxed); - - /* Available only when: T != float */ - T fetch_or(T operand, memory_order memoryOrder = - memory_order::relaxed); - - /* Available only when: T != float */ - T fetch_xor(T operand, memory_order memoryOrder = - memory_order::relaxed); - - /* Available only when: T != float */ - T fetch_min(T operand, memory_order memoryOrder = - memory_order::relaxed); - - /* Available only when: T != float */ - T fetch_max(T operand, memory_order memoryOrder = - memory_order::relaxed); +inline constexpr memory_order memory_order_relaxed = memory_order::relaxed; +inline constexpr memory_order memory_order_acquire = memory_order::acquire; +inline constexpr memory_order memory_order_release = memory_order::release; +inline constexpr memory_order memory_order_acq_rel = memory_order::acq_rel; +inline constexpr memory_order memory_order_seq_cst = memory_order::seq_cst; + +// Exposition only +template +struct memory_order_traits; + +template <> +struct memory_order_traits { + static constexpr memory_order read_order = memory_order::relaxed; + static constexpr memory_order write_order = memory_order::relaxed; }; -} // namespace sycl -} // namespace cl ----- -==== To: +template <> +struct memory_order_traits { + static constexpr memory_order read_order = memory_order::acquire; + static constexpr memory_order write_order = memory_order::release; +}; -[source,c++] ----- -namespace cl { -namespace sycl { -enum class memory_order : int { - relaxed +template <> +struct memory_order_traits { + static constexpr memory_order read_order = memory_order::seq_cst; + static constexpr memory_order write_order = memory_order::seq_cst; }; -namespace intel { -template -class atomic { + +template +class atomic_ref { public: - atomic(multi_ptr ptr); - atomic(const atomic&); - atomic& operator=(const atomic&) = delete; + using value_type = T; + static constexpr size_t required_alignment = /* implementation-defined */; + static constexpr bool is_always_lock_free = /* implementation-defined */; + static constexpr memory_order default_read_order = memory_order_traits::read_order; + static constexpr memory_order default_write_order = memory_order_traits::write_order; + static constexpr memory_order default_read_modify_write_order = DefaultOrder; bool is_lock_free() const; - void store(T operand, memory_order order = - memory_order::relaxed); + explicit atomic_ref(T&); + atomic_ref(const atomic_ref&); + atomic_ref& operator=(const atomic_ref&) = delete; + + void store(T operand, + memory_order order = default_write_order) const; - T operator=(T desired); + T operator=(T desired) const; - T load(memory_order order = memory_order::relaxed) const; + T load(memory_order order = default_read_order) const; operator T() const; - T exchange(T operand, memory_order order = - memory_order::relaxed); + T exchange(T operand, + memory_order order = default_read_modify_write_order) const; bool compare_exchange_weak(T &expected, T desired, - memory_order success = memory_order::relaxed, - memory_order failure = memory_order::relaxed); + memory_order success = default_read_modify_write_order, + memory_order failure = default_read_order) const; bool compare_exchange_weak(T &expected, T desired, - memory_order order = memory_order::relaxed); + memory_order order = default_read_modify_write_order) const; bool compare_exchange_strong(T &expected, T desired, - memory_order success = memory_order::relaxed, - memory_order failure = memory_order::relaxed); + memory_order success = default_read_modify_write_order, + memory_order failure = default_read_order) const; bool compare_exchange_strong(T &expected, T desired, - memory_order order = memory_order::relaxed); + memory_order order = default_read_modify_write_order) const; +}; - T fetch_add(T operand, memory_order order = - memory_order::relaxed); +// Partial specialization for integral types +template +class atomic_ref { - T fetch_sub(T operand, memory_order order = - memory_order::relaxed); + /* All other members from atomic_ref are available */ - T fetch_min(T operand, memory_order order = - memory_order::relaxed); + using difference_type = value_type; - T fetch_max(T operand, memory_order order = - memory_order::relaxed); + Integral fetch_add(Integral operand, + memory_order order = default_read_modify_write_order); - T operator+=(T operand); - T operator-=(T operand); + Integral fetch_sub(Integral operand, + memory_order order = default_read_modify_write_order); - /* Available only when T is Integral */ - T fetch_and(T operand, memory_order order = - memory_order::relaxed); + Integral fetch_and(Integral operand, + memory_order order = default_read_modify_write_order); - T fetch_or(T operand, memory_order order = - memory_order::relaxed); + Integral fetch_or(Integral operand, + memory_order order = default_read_modify_write_order); - T fetch_xor(T operand, memory_order order = - memory_order::relaxed); + Integral fetch_min(Integral operand, + memory_order order = default_read_modify_write_order); - T operator++(int operand); - T operator--(int operand); - T operator++(); - T operator--(); - T operator&= (T operand); - T operator|= (T operand); - T operator^= (T operand); -}; -} // namespace intel -} // namespace sycl -} // namespace cl ----- - -=== Modify Table 4.100 - -==== From: - -|=== -|Constructor|Description + Integral fetch_max(Integral operand, + memory_order order = default_read_modify_write_order); -|+template atomic(multi_ptr ptr)+ -|Permitted data types for +pointerT+ are any valid scalar data type which is the same size in bytes as +T+. Constructs an instance of SYCL +atomic+ which is associated with the pointer +ptr+, converted to a pointer of data type +T+. -|=== + Integral operator++(int) const noexcept; + Integral operator--(int) const noexcept; + Integral operator++() const noexcept; + Integral operator--() const noexcept; + Integral operator+=(Integral) const noexcept; + Integral operator-=(Integral) const noexcept; + Integral operator&=(Integral) const noexcept; + Integral operator|=(Integral) const noexcept; + Integral operator^=(Integral) const noexcept; -==== To: +}; -|=== -|Constructor|Description +// Partial specialization for floating-point types +template +class atomic_ref { -|+atomic(multi_ptr ptr)+ -|Constructs an instance of SYCL +atomic+ which is associated with the pointer +ptr+. -|=== + /* All other members from atomic_ref are available */ -=== Modify Table 4.101 + using difference_type = value_type; -==== Add: + Floating fetch_add(Floating operand, + memory_order order = default_read_modify_write_order); -|=== -|Member function|Description -|+bool is_lock_free() const+ -|Return +true+ if the atomic operations provided by this SYCL +atomic+ are lock-free. + Floating fetch_sub(Floating operand, + memory_order order = default_read_modify_write_order); -|+bool compare_exchange_weak(T &expected, T desired, memory_order order = memory_order::relaxed)+ -|Atomically compares the value at the address of the +multi_ptr+ associated with this SYCL +atomic+ against the value of +expected+. If the values are equal attempts to replaces value at address of the +multi_ptr+ associated with this SYCL +atomic+ with the value of +desired+, otherwise assigns the original value at the address of the +multi_ptr+ associated with this SYCL +atomic+ to +expected+. Returns +true+ if the comparison operation and replacement operation were successful. The memory order of this atomic operation must be +memory_order::relaxed+ for both success and fail. + Floating fetch_min(Floating operand, + memory_order order = default_read_modify_write_order); -|+bool compare_exchange_weak(T &expected, T desired, memory_order order = memory_order::relaxed)+ -|Equivalent to +compare_exchange_weak(expected, desired, order, order)+. + Floating fetch_max(Floating operand, + memory_order order = default_read_modify_write_order); -|+bool compare_exchange_strong(T &expected, T desired, memory_order order = memory_order::relaxed)+ -|Equivalent to +compare_exchange_strong(expected, desired, order, order)+. + Floating operator++(int) const noexcept; + Floating operator--(int) const noexcept; + Floating operator++() const noexcept; + Floating operator--() const noexcept; + Floating operator+=(Floating) const noexcept; + Floating operator-=(Floating) const noexcept; -|+operator T() const+ -|Equivalent to +load()+. +}; -|+T operator=(T desired)+ -|Equivalent to +store(desired)+. Returns +desired+. +// Partial specialization for pointers +template