[SYCL][ESIMD] Implement unified memory API - atomic_update(acc,...) zero operands (#11884)

This implements the compile-time properties version of atomic_update
with accessors and zero operands.

I moved the zero operand block to be before the one operand block, so
hopefully it's easier to review. I didn't port the local accessor ones
yet, I'll do that as part of the SLM change.

I also moved the one and two operand intrinsics to supported even though
they aren't used by and supported code yet, it's just preparing for the
next change.

I also fixed some typos in the usm version doc.

---------

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp

@@ -502,6 +502,116 @@ __esimd_lsc_xatomic_stateless_2(
 }
 #endif // __SYCL_DEVICE_ONLY__
 
+/// Accessor-based atomic.
+/// Supported platforms: DG2, PVC
+///
+/// @tparam Ty is element type.
+/// @tparam InternalOp is operation type.
+/// @tparam L1H is L1 cache hint.
+/// @tparam L2H is L2 cache hint.
+/// @tparam AddressScale is the address scale.
+/// @tparam ImmOffset is the immediate offset added to each address.
+/// @tparam DS is the data size.
+/// @tparam VS is the number of elements per address.
+/// @tparam Transposed indicates if the data is transposed during the transfer.
+/// @tparam N is the SIMD size of operation (the number of addresses to access)
+/// @tparam SurfIndAliasTy is the \ref sycl::accessor type.
+/// @param pred is predicates.
+/// @param offsets is the zero-based offsets.
+/// @param surf_ind is the surface index.
+template <typename Ty, int InternalOp, __ESIMD_NS::cache_hint L1H,
+          __ESIMD_NS::cache_hint L2H, uint16_t AddressScale, int ImmOffset,
+          __ESIMD_DNS::lsc_data_size DS, __ESIMD_DNS::lsc_vector_size VS,
+          __ESIMD_DNS::lsc_data_order Transposed, int N,
+          typename SurfIndAliasTy>
+__ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_DNS::to_int<VS>()>
+__esimd_lsc_xatomic_bti_0(__ESIMD_DNS::simd_mask_storage_t<N> pred,
+                          __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
+                          SurfIndAliasTy surf_ind)
+#ifdef __SYCL_DEVICE_ONLY__
+    ;
+#else  // __SYCL_DEVICE_ONLY__
+{
+  __ESIMD_UNSUPPORTED_ON_HOST;
+}
+#endif // __SYCL_DEVICE_ONLY__
+
+/// Accessor-based atomic.
+/// Supported platforms: DG2, PVC
+///
+/// @tparam Ty is element type.
+/// @tparam InternalOp is operation type.
+/// @tparam L1H is L1 cache hint.
+/// @tparam L2H is L2 cache hint.
+/// @tparam AddressScale is the address scale.
+/// @tparam ImmOffset is the immediate offset added to each address.
+/// @tparam DS is the data size.
+/// @tparam VS is the number of elements per address.
+/// @tparam Transposed indicates if the data is transposed during the transfer.
+/// @tparam N is the SIMD size of operation (the number of addresses to access)
+/// @tparam SurfIndAliasTy is the \ref sycl::accessor type.
+/// @param pred is predicates.
+/// @param offsets is the zero-based offsets.
+/// @param src0 is the first atomic operand.
+/// @param surf_ind is the surface index.
+template <typename Ty, int InternalOp, __ESIMD_NS::cache_hint L1H,
+          __ESIMD_NS::cache_hint L2H, uint16_t AddressScale, int ImmOffset,
+          __ESIMD_DNS::lsc_data_size DS, __ESIMD_DNS::lsc_vector_size VS,
+          __ESIMD_DNS::lsc_data_order _Transposed, int N,
+          typename SurfIndAliasTy>
+__ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_DNS::to_int<VS>()>
+__esimd_lsc_xatomic_bti_1(
+    __ESIMD_DNS::simd_mask_storage_t<N> pred,
+    __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
+    __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_DNS::to_int<VS>()> src0,
+    SurfIndAliasTy surf_ind)
+#ifdef __SYCL_DEVICE_ONLY__
+    ;
+#else  // __SYCL_DEVICE_ONLY__
+{
+  __ESIMD_UNSUPPORTED_ON_HOST;
+}
+#endif // __SYCL_DEVICE_ONLY__
+
+/// Accessor-based atomic.
+/// Supported platforms: DG2, PVC
+///
+/// @tparam Ty is element type.
+/// @tparam InternalOp is operation type.
+/// @tparam L1H is L1 cache hint.
+/// @tparam L2H is L2 cache hint.
+/// @tparam AddressScale is the address scale.
+/// @tparam ImmOffset is the immediate offset added to each address.
+/// @tparam DS is the data size.
+/// @tparam VS is the number of elements per address.
+/// @tparam Transposed indicates if the data is transposed during the transfer.
+/// @tparam N is the SIMD size of operation (the number of addresses to access)
+/// @tparam SurfIndAliasTy is the \ref sycl::accessor type.
+/// @param pred is predicates.
+/// @param offsets is the zero-based offsets.
+/// @param src0 is the first atomic operand.
+/// @param src1 is the second atomic operand.
+/// @param surf_ind is the surface index.
+template <typename Ty, int InternalOp, __ESIMD_NS::cache_hint L1H,
+          __ESIMD_NS::cache_hint L2H, uint16_t AddressScale, int ImmOffset,
+          __ESIMD_DNS::lsc_data_size DS, __ESIMD_DNS::lsc_vector_size VS,
+          __ESIMD_DNS::lsc_data_order Transposed, int N,
+          typename SurfIndAliasTy>
+__ESIMD_INTRIN __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_DNS::to_int<VS>()>
+__esimd_lsc_xatomic_bti_2(
+    __ESIMD_DNS::simd_mask_storage_t<N> pred,
+    __ESIMD_DNS::vector_type_t<uint32_t, N> offsets,
+    __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_DNS::to_int<VS>()> src0,
+    __ESIMD_DNS::vector_type_t<Ty, N * __ESIMD_DNS::to_int<VS>()> src1,
+    SurfIndAliasTy surf_ind)
+#ifdef __SYCL_DEVICE_ONLY__
+    ;
+#else  // __SYCL_DEVICE_ONLY__
+{
+  __ESIMD_UNSUPPORTED_ON_HOST;
+}
+#endif // __SYCL_DEVICE_ONLY__
+
 __ESIMD_INTRIN void __esimd_slm_init(uint32_t size)
 #ifdef __SYCL_DEVICE_ONLY__
     ;