[ESIMD] Add non-experimental version of fence() for DG2 and PVC (intel#11923)

The new esimd::fence() accepts 3 template enum parameters and is usable
only on DG2 and PVC.
It implements the subset of functionality supported by
experimental::esimd::lsc_fence(). The functionality that was not
included into the public version is a) not supported by hardware and/or
GPU drivers
b) too specific and is too risky to use now.

---------

Signed-off-by: Klochkov, Vyacheslav N <[email protected]>
Co-authored-by: Victor Mustya <[email protected]>

Author: v-klochkov

sycl/include/sycl/ext/intel/esimd/common.hpp

@@ -384,6 +384,62 @@ enum class cache_hint : uint8_t {
   const_cached = 7
 };
 
+/// The scope that fence() operation should apply to.
+/// Supported platforms: DG2, PVC
+enum class fence_scope : uint8_t {
+  /// Wait until all previous memory transactions from this thread are observed
+  /// within the local thread-group.
+  group = 0,
+
+  /// Wait until all previous memory transactions from this thread are observed
+  /// within the local sub-slice.
+  local = 1,
+
+  /// Wait until all previous memory transactions from this thread are observed
+  /// in the local tile.
+  tile = 2,
+
+  /// Wait until all previous memory transactions from this thread are observed
+  /// in the local GPU.
+  gpu = 3,
+
+  /// Wait until all previous memory transactions from this thread are observed
+  /// across all GPUs in the system.
+  gpus = 4,
+
+  /// Global memory data-port only: wait until all previous memory transactions
+  /// from this thread are observed at the "system" level.
+  system = 5,
+
+  /// Global memory data-port only: for GPUs that do not follow
+  /// PCIe Write ordering for downstream writes targeting device memory,
+  /// this op will commit to device memory all downstream and peer writes that
+  /// have reached the device.
+  system_acquire = 6
+};
+
+/// The cache flush operation to apply to caches after fence() is complete.
+/// Supported platforms: DG2, PVC
+enum class fence_flush_op : uint8_t {
+  none = 0,       /// no operation;
+  evict = 1,      /// R/W: evict dirty lines; R/W and RO: invalidate clean lines
+  invalidate = 2, /// R/W and RO: invalidate all clean lines;
+
+  // enum with the value 3 is reserved;
+
+  clean = 4 /// R/W: dirty lines are written to memory, but retained in
+            /// cache in clean state; RO: no effect.
+};
+
+/// The target memory kind for fence() operation.
+/// Supported platforms: DG2, PVC
+enum class memory_kind : uint8_t {
+  global = 0, /// untyped global memory
+  // enum with the value 1 is reserved;
+  image = 2, /// image (also known as typed global memory)
+  local = 3, /// shared local memory
+};
+
 /// L1, L2 or L3 cache hint levels. L3 is reserved for future use.
 enum class cache_level : uint8_t { L1 = 1, L2 = 2, L3 = 3 };
 

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

@@ -641,6 +641,24 @@ __ESIMD_INTRIN void __esimd_fence(uint8_t cntl)
 }
 #endif // __SYCL_DEVICE_ONLY__
 
+/// Memory fence.
+/// Supported platforms: DG2, PVC
+///
+/// @tparam Kind is the Sfid shaded function.
+/// @tparam FenceOp is the fence operation.
+/// @tparam Scope is the operation scope.
+/// @tparam N is the SIMD size of operation (the number of addresses to access)
+/// @param pred is predicates.
+template <uint8_t Kind, uint8_t FenceOp, uint8_t Scope, int N>
+__ESIMD_INTRIN void __esimd_lsc_fence(__ESIMD_DNS::simd_mask_storage_t<N> pred)
+#ifdef __SYCL_DEVICE_ONLY__
+    ;
+#else  // __SYCL_DEVICE_ONLY__
+{
+  __ESIMD_UNSUPPORTED_ON_HOST;
+}
+#endif // __SYCL_DEVICE_ONLY__
+
 // Predicated (masked) scaled gather from a surface.
 //
 // Template (compile-time constant) parameters:

sycl/include/sycl/ext/intel/esimd/memory.hpp

@@ -5627,6 +5627,26 @@ template <uint8_t cntl> __ESIMD_API void fence() { __esimd_fence(cntl); }
 __SYCL_DEPRECATED("use fence<fence_mask>()")
 __ESIMD_API void fence(fence_mask cntl) { __esimd_fence(cntl); }
 
+/// Memory fence.
+/// Supported platforms: DG2, PVC
+///
+/// @tparam Kind is the memory kind.
+/// @tparam FenceOp is the fence cache flush operation to apply after fence.
+/// @tparam Scope is the fence operation scope.
+template <memory_kind Kind = memory_kind::global,
+          fence_flush_op FenceOp = fence_flush_op::none,
+          fence_scope Scope = fence_scope::group>
+__ESIMD_API void fence() {
+  static_assert(
+      Kind != memory_kind::local ||
+          (FenceOp == fence_flush_op::none && Scope == fence_scope::group),
+      "SLM fence must have 'none' lsc_fence_op and 'group' scope");
+  constexpr int N = 16;
+  simd_mask<N> Mask = 1;
+  __esimd_lsc_fence<static_cast<uint8_t>(Kind), static_cast<uint8_t>(FenceOp),
+                    static_cast<uint8_t>(Scope), N>(Mask.data());
+}
+
 /// Generic work-group barrier.
 /// Performs barrier synchronization for all threads within the same thread
 /// group. The barrier instruction causes the executing thread to wait until

sycl/include/sycl/ext/intel/experimental/esimd/common.hpp

@@ -27,7 +27,8 @@ namespace ext::intel::experimental::esimd {
 
 /// The scope that lsc_fence operation should apply to
 /// Supported platforms: DG2, PVC
-enum class lsc_scope : uint8_t {
+enum class __SYCL_DEPRECATED(
+    "use sycl::ext::intel::esimd::fence_scope") lsc_scope : uint8_t {
   group = 0,  /// flush out to the threadgroup's scope
   local = 1,  /// flush out to the local scope
   tile = 2,   /// tile, flush out to several DSSs
@@ -39,24 +40,26 @@ enum class lsc_scope : uint8_t {
 
 /// The lsc_fence operation to apply to caches
 /// Supported platforms: DG2, PVC
-enum class lsc_fence_op : uint8_t {
-  none = 0,       /// no operation
-  evict = 1,      /// dirty lines evicted and invalidated from L1
-  invalidate = 2, /// invalidate all clean lines
-  discard = 3,    /// direct and clean lines are discarded w/o eviction
-  clean = 4,      /// dirty lines are written to memory, but retained in cache
-                  /// in clean state
-  flushl3 = 5,    /// flush only L3
-};
+enum class __SYCL_DEPRECATED("use sycl::ext::intel::esimd::fence_flush_op")
+    lsc_fence_op : uint8_t {
+      none = 0,       /// no operation
+      evict = 1,      /// dirty lines evicted and invalidated from L1
+      invalidate = 2, /// invalidate all clean lines
+      discard = 3,    /// direct and clean lines are discarded w/o eviction
+      clean = 4,   /// dirty lines are written to memory, but retained in cache
+                   /// in clean state
+      flushl3 = 5, /// flush only L3
+    };
 
 /// The specific LSC shared function to fence with lsc_fence
 /// Supported platforms: DG2, PVC
-enum class lsc_memory_kind : uint8_t {
-  untyped_global = 0,         /// untyped global memory
-  untyped_global_low_pri = 1, /// low-priority untyped global memory
-  typed_global = 2,           /// typed global memory
-  shared_local = 3,           /// shared local memory
-};
+enum class __SYCL_DEPRECATED("use sycl::ext::intel::esimd::memory_kind")
+    lsc_memory_kind : uint8_t {
+      untyped_global = 0,         /// untyped global memory
+      untyped_global_low_pri = 1, /// low-priority untyped global memory
+      typed_global = 2,           /// typed global memory
+      shared_local = 3,           /// shared local memory
+    };
 
 using lsc_data_size = __ESIMD_DNS::lsc_data_size;
 

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

@@ -416,25 +416,6 @@ __esimd_lsc_xatomic_slm_2(
 }
 #endif // __SYCL_DEVICE_ONLY__
 
-/// Memory fence.
-/// Supported platforms: DG2, PVC
-///
-/// @tparam Kind is the Sfid shaded function.
-/// @tparam FenceOp is the fence operation.
-/// @tparam Scope is the operation scope.
-/// @tparam N is the SIMD size of operation (the number of addresses to access)
-/// @param pred is predicates.
-template <__ESIMD_ENS::lsc_memory_kind Kind, __ESIMD_ENS::lsc_fence_op FenceOp,
-          __ESIMD_ENS::lsc_scope Scope, int N>
-__ESIMD_INTRIN void __esimd_lsc_fence(__ESIMD_DNS::simd_mask_storage_t<N> pred)
-#ifdef __SYCL_DEVICE_ONLY__
-    ;
-#else  // __SYCL_DEVICE_ONLY__
-{
-  __ESIMD_UNSUPPORTED_ON_HOST;
-}
-#endif // __SYCL_DEVICE_ONLY__
-
 __ESIMD_INTRIN uint32_t __esimd_slm_alloc(uint32_t size)
 #ifdef __SYCL_DEVICE_ONLY__
     ;

sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp

@@ -3091,12 +3091,17 @@ lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
 template <lsc_memory_kind Kind = lsc_memory_kind::untyped_global,
           lsc_fence_op FenceOp = lsc_fence_op::none,
           lsc_scope Scope = lsc_scope::group, int N = 16>
+__SYCL_DEPRECATED("use sycl::ext::intel::esimd::fence<Kind, FenceOp, Scope>()")
 __ESIMD_API void lsc_fence(__ESIMD_NS::simd_mask<N> pred = 1) {
   static_assert(
       Kind != lsc_memory_kind::shared_local ||
           (FenceOp == lsc_fence_op::none && Scope == lsc_scope::group),
       "SLM fence must have 'none' lsc_fence_op and 'group' scope");
-  __esimd_lsc_fence<Kind, FenceOp, Scope, N>(pred.data());
+  static_assert(Kind != lsc_memory_kind::untyped_global_low_pri,
+                "lsc_memory_kind::untyped_global_low_pri is not supported in HW"
+                " and/or GPU drivers");
+  __esimd_lsc_fence<static_cast<uint8_t>(Kind), static_cast<uint8_t>(FenceOp),
+                    static_cast<uint8_t>(Scope), N>(pred.data());
 }
 
 /// @} sycl_esimd_memory_lsc