[ESIMD][NFC][DOC] Add 'restriction' section to atomic_update() doc (#13202)

Signed-off-by: Klochkov, Vyacheslav N <[email protected]>

Author: v-klochkov

sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd.md

@@ -46,7 +46,7 @@
 
 ## Introduction
 
-The main motivation for introducing the "Explicit SIMD" SYCL extension 
+The main motivation for introducing the "Explicit SIMD" SYCL extension
 (or simply "ESIMD") is enabling efficient low-level programming for Intel graphics
 architectures. It provides APIs close to the Intel GPU ISA
 and allows writing explicitly vectorized device code.
@@ -356,7 +356,7 @@ See more details on the API documentation [page TODO](https://intel.github.io/ll
 ### Memory access APIs
 
 Explicit SIMD memory access interface is quite different from the standard SYCL
-memory access interface. It supports main SYCL's device memory representations:
+memory access interface. It supports main SYCL device memory representations:
 - USM pointers
 - SYCL accessors
   - 1D global accessors
@@ -651,7 +651,7 @@ ESIMD supports the following non-standard math functions implemented in hardware
 - Fraction - `frc`,  extracts the fractional parts of the input vector elements.
 - Count leading zeroes - `lzd`.
 - Linear interpolation - `lrp`. Basically computes `src1 * src0 + src2 * (1.0f - src0)`
-- Plane equation - `plane`. Solves a component-wise plane equation 
+- Plane equation - `plane`. Solves a component-wise plane equation
   `w = p*u + q*v + r` where `u`, `v`, `w` are vectors and `p`, `q`, `r` are scalars.
 
 
@@ -865,7 +865,7 @@ There are other useful miscellaneous APIs provided by ESIMD.
   types with saturation.
 - Conversion - `convert`. Converts between vectors with different element data
   types.
-- Reverse bits - `bf_reverse`. 
+- Reverse bits - `bf_reverse`.
 - Insert bit field - `bf_insert`.
 - Extract bit field - `bf_extract`.
 - Convert mask to integer and back - `pack_mask`, `unpack_mask`.
@@ -978,7 +978,7 @@ More examples of the unwrap/merge process:
       B6 b;
       char x;
       char y;
-    
+
       C6 foo() { return *this; }
     };
     ```
@@ -989,7 +989,7 @@ More examples of the unwrap/merge process:
       ```
       %struct.C6 = type { %struct.B6, i8, i8 }
       %struct.B6 = type { i32 addrspace(4)*, i32 }
-      ``` 
+      ```
 
 Note that `__regcall` does not guarantee passing through registers in the final
 generated code. For example, compiler will use a threshold for argument or
@@ -1162,8 +1162,7 @@ inside ESIMD kernels and functions. Most of missing SYCL features listed below
 must be supported eventually:
 - 2D and 3D target::device accessor and local_accessor;
 - Constant accessors;
-- `sycl::accessor::get_pointer()` and `sycl::accessor::operator[]` are supported only with `-fsycl-esimd-force-stateless-mem`. Otherwise, All memory accesses through an accessor are
-done via explicit APIs; e.g. `sycl::ext::intel::esimd::block_store(acc, offset)`
+- `sycl::accessor::get_pointer()` and `sycl::accessor::operator[]` are not supported with with `-fno-sycl-esimd-force-stateless-mem` compilation switch.
 - Accessors with non-zero offsets to accessed buffer;
 - Accessors with access/memory range specified;
 - `sycl::image`, `sycl::sampler` and `sycl::stream` classes.

sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd_functions.md

@@ -122,7 +122,7 @@ The optional [compile-time properties](#compile-time-properties) list `props` ma
 
 `N` - the valid values may depend on usage of cache-hints or passing of the `pred` argument:
 
-| `Function` | `Condition` | Requirement for `N` | Required/supported Intel GPU |
+| `Function` | `Condition` | Requirement for `N` | Required Intel GPU |
 |-|-|-|-|
 | `(usm-bl-*)` | (no cache-hints) and (`pred` is not passed) | `N` is any positive number | Any Intel GPU |
 | `(usm-bl-*)` | (cache-hints) or (`pred` is passed) | `N` must be from [Table1 below](#table1---valid-values-of-n-if-cache-hints-used-or-pred-parameter-is-passed) | DG2 or PVC |
@@ -195,7 +195,7 @@ The optional [compile-time properties](#compile-time-properties) list `props` ma
 
 `N` - the valid values may depend on usage of cache-hints or passing of the `pred` argument:
 
-| `Function` | `Condition` | Requirement for `N` | Required/supported Intel GPU |
+| `Function` | `Condition` | Requirement for `N` | Required Intel GPU |
 |-|-|-|-|
 | `(usm-bs-*)` | (no cache-hints) and (`pred` is not passed) | `N` is any positive number | Any Intel GPU |
 | `(usm-bs-*)` | (cache-hints) or (`pred` is passed) | `N` must be from [Table2 below](#table1---valid-values-of-n-if-cache-hints-used-or-pred-parameter-is-passed) | DG2 or PVC |
@@ -338,7 +338,7 @@ template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename Prop
 `(slm-ga-*)`: Loads ("gathers") elements of the type `T` from shared local memory locations addressed by `byte_offsets`.  
 The parameter `byte_offset` is a vector of any integral type elements for `(usm-ga-*)`, 32-bit integer elements for `(lacc-ga-*)` and `(slm-ga-*)`, any integral type integer elements for `(acc-ga-*)` in [stateless](#statelessstateful-memory-mode) mode(default),
 and up-to-32-bit integer elements for `(acc-ga-*)` in [stateful](#statelessstateful-memory-mode) mode.  
-The optional parameter `pred` provides a `simd_mask`. If some element in `pred` is zero, then the load of the corresponding memory location is skipped and the element of the result is copied from `pass_thru` (if it is passed) or it is undefined (if `pass_thru` is omitted).  
+The optional parameter `mask` provides a `simd_mask`. If some element in `mask` is zero, then the load of the corresponding memory location is skipped and the element of the result is copied from `pass_thru` (if it is passed) or it is undefined (if `pass_thru` is omitted).  
 The optional [compile-time properties](#compile-time-properties) list `props` may specify `alignment` and/or `cache-hints`. The cache-hints are ignored for `(lacc-*)` and `(slm-*)` functions.  
 The template parameter `N` can be any positive number.  
 The optional template parameter `VS` must be one of `{1, 2, 3, 4, 8, 16, 32, 64}` values. It specifies how many conseсutive elements are loaded per each element in `byte_offsets`.   
@@ -354,7 +354,7 @@ simd<float, 8> vec8 = gather<float, 8, 2>(ptr, offsets);
 ```
 
 ### Restrictions
-| `Function` | `Condition` | Required/supported Intel GPU |
+| `Function` | `Condition` | Required Intel GPU |
 |-|-|-|
 | `(usm-ga-1,4,7)`,`(acc-ga-1,4,7)` | true (`pass_thru` arg is passed) | DG2 or PVC |
 | `(usm-ga-2,3,8,9)`, `(acc-ga-2,3,8,9)` | !(cache-hints) and (`VS` == 1) and (`N` == 1,2,4,8,16,32) | Any Intel GPU |
@@ -439,7 +439,7 @@ template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename Prop
 `(slm-sc-*)`: Stores ("scatters") the vector `vals` to shared local memory locations addressed by `byte_offsets`.  
 The parameter `byte_offset` is a vector of any integral type elements for `(usm-sc-*)`, 32-bit integer elements for `(lacc-sc-*)` and `(slm-sc-*)`, any integral type integer elements for `(acc-sc-*)` in [stateless](#statelessstateful-memory-mode) mode(default),
 and up-to-32-bit integer elements for `(acc-sc-*)` in [stateful](#statelessstateful-memory-mode) mode.  
-The optional parameter `pred` provides a `simd_mask`. If some element in `pred` is zero, then the store to the corresponding memory location is skipped.  
+The optional parameter `mask` provides a `simd_mask`. If some element in `mask` is zero, then the store to the corresponding memory location is skipped.  
 The optional [compile-time properties](#compile-time-properties) list `props` may specify `alignment` and/or `cache-hints`. The cache-hints are ignored for `(lacc-sc-*)` and `(slm-sc-*)` functions.
 The template parameter `N` can be any positive number.  
 The optional template parameter `VS` must be one of `{1, 2, 3, 4, 8, 16, 32, 64}` values. It specifies how many conseсutive elements are written per each element in `byte_offsets`.   
@@ -457,7 +457,7 @@ scatter<float, 8, 2>(ptr, offsets4);
 ```
 
 ### Restrictions
-| `Function` | `Condition` | Required/supported Intel GPU |
+| `Function` | `Condition` | Required Intel GPU |
 |-|-|-|
 | `(usm-sc-*)`, `(acc-sc-*)` | !(cache-hints) and (`VS` == 1) and (`N` == 1,2,4,8,16,32) | Any Intel GPU |
 | `(usm-sc-*)`, `(acc-sc-*)` | (cache-hints) or (`VS` > 1) or (`N` != 1,2,4,8,16,32) | DG2 or PVC |
@@ -502,7 +502,7 @@ template <atomic_op Op, typename T, int N, typename OffsetObjT, typename Accesso
 // Atomic update the memory locations referenced by local-accessor (SLM) - zero operands (dec, load, etc.).
 template <atomic_op Op, typename T, int N, typename AccessorT>
 /*lacc-au0-1*/ simd<T, N> atomic_update(AccessorT lacc, simd<uint32_t, N> byte_offset,
-                                        simd_mask<1> pred = 1);
+                                        simd_mask<1> mask = 1);
 
 // Atomic update the shared local memory (SLM) - zero operands (dec, load, etc.).
 template <atomic_op Op, typename T, int N>
@@ -543,7 +543,7 @@ template <atomic_op Op, typename T, int N, typename OffsetObjT, typename Accesso
 // Atomic update the memory locations referenced by local-accessor (SLM) - one operand (add, max, etc.).
 template <atomic_op Op, typename T, int N, typename AccessorT>
 /*lacc-au1-1*/ simd<T, N> atomic_update(AccessorT lacc, simd<uint32_t, N> byte_offset,
-                                        simd<T, N> src0, simd_mask<1> pred = 1);
+                                        simd<T, N> src0, simd_mask<1> mask = 1);
 
 // Atomic update the shared local memory (SLM) - one operand (add, max etc.).
 template <atomic_op Op, typename T, int N>
@@ -583,7 +583,7 @@ template <atomic_op Op, typename T, int N, typename OffsetObjT, typename Accesso
 // Atomic update the memory locations referenced by local-accessor (SLM) - two operands: cmpxchg, fcmpxchg.
 template <atomic_op Op, typename T, int N, typename AccessorT>
 /*lacc-au2-1*/ simd<T, N> atomic_update(AccessorT lacc, simd<uint32_t, N> byte_offset,
-                                        simd<T, N> src0, simd<T, N> src1, simd_mask<1> pred = 1);
+                                        simd<T, N> src0, simd<T, N> src1, simd_mask<1> mask = 1);
 
 // Atomic update the shared local memory (SLM) - two operands: cmpxchg, fcmpxchg.
 template <atomic_op Op, typename T, int N>
@@ -597,8 +597,25 @@ template <atomic_op Op, typename T, int N>
 `(slm-*)`: Atomically updates the shared memory locations addressed by `byte_offset`.  
 The parameter `byte_offset` is a vector of any integral type elements for `(usm-*)`, 32-bit integer elements for `(lacc-*)` and `(slm-*)`, any integral type integer elements for `(acc-*)` in [stateless](#statelessstateful-memory-mode) mode(default),
 and up-to-32-bit integer elements for `(acc-*)` in [stateful](#statelessstateful-memory-mode) mode.  
-The optional parameter `pred` provides a `simd_mask`. If some element in `pred` is zero, then the corresponding memory location is not updated.  
-`(usm-*)`, `(acc-*)`: The optional [compile-time properties](#compile-time-properties) list `props` may specify `cache-hints`.
+The optional parameter `mask` provides a `simd_mask`. If some element in `mask` is zero, then the corresponding memory location is not updated.  
+`(usm-*)`, `(acc-*)`: The optional [compile-time properties](#compile-time-properties) list `props` may specify `cache-hints`.  
+The template parameter `Op` specifies the atomic operation applied to the memory.  
+The template parameter `T` specifies the type of the elements used in the atomic_update operation. Only 2,4,8-byte types are supported.  
+The template parameter `N` is the number of elements being atomically updated.
+
+### Restrictions
+| `Function` | `Condition` | Required Intel GPU |
+|-|-|-|
+| `(usm-au0-*)`, `(acc-au0-*)` | !(cache-hints) and (`N` == 1,2,4,8,16,32) and (sizeof(T) >= 4) | Any Intel GPU |
+| `(usm-au0-*)`, `(acc-au0-*)` | (cache-hints) or (`N` != 1,2,4,8,16,32) or (sizeof(T) == 2) | DG2 or PVC |
+| `(usm-au1-*)`, `(acc-au1-*)`, `(usm-au2-*)`, `(acc-au2-*)`  | !(cache-hints) and (`N` == 1,2,4,8,16,32) and (sizeof(T) >= 4) and (`Op` is integral operation) | Any Intel GPU |
+| `(usm-au1-*)`, `(acc-au1-*)`, `(usm-au2-*)`, `(acc-au2-*)` | (cache-hints) or (`N` != 1,2,4,8,16,32) or (sizeof(T) == 2) or (`Op` is FP operation) | DG2 or PVC |
+|-|-|-|
+| `(slm-au0-*)`, `(lacc-au0-*)` | (`N` == 1,2,4,8,16,32) and (sizeof(T) == 4) | Any Intel GPU |
+| `(slm-au0-*)`, `(lacc-au0-*)` | (`N` != 1,2,4,8,16,32) or (sizeof(T) == 2) or (sizeof(T) == 8)| DG2 or PVC |
+| `(slm-au1-*)`, `(lacc-au1-*)`, `(slm-au2-*)`, `(lacc-au2-*)` | (`N` == 1,2,4,8,16,32) and (sizeof(T) == 4) and (`Op` is integral operation) | Any Intel GPU |
+| `(slm-au1-*)`, `(lacc-au1-*)`, `(slm-au2-*)`, `(lacc-au2-*)` | (`N` != 1,2,4,8,16,32) or (sizeof(T) == 2) or (sizeof(T) == 8) or (`Op` is FP operation)| DG2 or PVC |
+
 
 ## prefetch(...)
 ```C++