diff --git a/sycl/include/sycl/ext/intel/experimental/esimd.hpp b/sycl/include/sycl/ext/intel/experimental/esimd.hpp
index e52fbe0c7641d..6a5b6bd4e9ab6 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd.hpp
@@ -10,16 +10,9 @@
 
 #pragma once
 
-#ifdef __SYCL_DEVICE_ONLY__
-#define SYCL_ESIMD_KERNEL __attribute__((sycl_explicit_simd))
-#define SYCL_ESIMD_FUNCTION __attribute__((sycl_explicit_simd))
-#else
-#define SYCL_ESIMD_KERNEL
-#define SYCL_ESIMD_FUNCTION
-#endif
-
 /// \defgroup sycl_esimd DPC++ Explicit SIMD API
 
+#include <sycl/ext/intel/experimental/esimd/common.hpp>
 #include <sycl/ext/intel/experimental/esimd/math.hpp>
 #include <sycl/ext/intel/experimental/esimd/memory.hpp>
 #include <sycl/ext/intel/experimental/esimd/simd.hpp>
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/common.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/common.hpp
index e5fa31dd70b25..5d3061ca68b45 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/common.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/common.hpp
@@ -14,6 +14,14 @@
 
 #include <cstdint> // for uint* types
 
+#ifdef __SYCL_DEVICE_ONLY__
+#define SYCL_ESIMD_KERNEL __attribute__((sycl_explicit_simd))
+#define SYCL_ESIMD_FUNCTION __attribute__((sycl_explicit_simd))
+#else
+#define SYCL_ESIMD_KERNEL
+#define SYCL_ESIMD_FUNCTION
+#endif
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace ext {
@@ -50,57 +58,162 @@ using uint = unsigned int;
 // functions defined in a header.
 #define ESIMD_INLINE inline __attribute__((always_inline))
 
-// Enums
-// TODO FIXME convert the two enums below to nested enum or class enum to
-// remove enum values from the global namespace
-enum { GENX_NOSAT = 0, GENX_SAT };
-
-enum ChannelMaskType {
-  ESIMD_R_ENABLE = 1,
-  ESIMD_G_ENABLE = 2,
-  ESIMD_GR_ENABLE = 3,
-  ESIMD_B_ENABLE = 4,
-  ESIMD_BR_ENABLE = 5,
-  ESIMD_BG_ENABLE = 6,
-  ESIMD_BGR_ENABLE = 7,
-  ESIMD_A_ENABLE = 8,
-  ESIMD_AR_ENABLE = 9,
-  ESIMD_AG_ENABLE = 10,
-  ESIMD_AGR_ENABLE = 11,
-  ESIMD_AB_ENABLE = 12,
-  ESIMD_ABR_ENABLE = 13,
-  ESIMD_ABG_ENABLE = 14,
-  ESIMD_ABGR_ENABLE = 15
+// Macros for internal use
+#define __ESIMD_NS sycl::ext::intel::experimental::esimd
+#define __ESIMD_QUOTE1(m) #m
+#define __ESIMD_QUOTE(m) __ESIMD_QUOTE1(m)
+#define __ESIMD_NS_QUOTED __ESIMD_QUOTE(__ESIMD_NS)
+#define __ESIMD_DEPRECATED(new_api)                                            \
+  __SYCL_DEPRECATED("use " __ESIMD_NS_QUOTED "::" __ESIMD_QUOTE(new_api))
+// Defines a deprecated enum value. Use of this value will cause a deprecation
+// message printed out by the compiler.
+#define __ESIMD_DEPR_ENUM_V(old, new, t)                                       \
+  old __ESIMD_DEPRECATED(new) = static_cast<t>(new)
+
+/// Gen hardware supports applying saturation to results of some operation.
+/// This enum allows to control this behavior.
+enum class saturation : uint8_t { off, on };
+
+/// Integer type short-cut to saturation::off.
+static inline constexpr uint8_t saturation_off =
+    static_cast<uint8_t>(saturation::off);
+/// Integer type short-cut to saturation::on.
+static inline constexpr uint8_t saturation_on =
+    static_cast<uint8_t>(saturation::on);
+
+enum {
+  __ESIMD_DEPR_ENUM_V(GENX_NOSAT, saturation::off, uint8_t),
+  __ESIMD_DEPR_ENUM_V(GENX_SAT, saturation::on, uint8_t)
 };
 
-#define NumChannels(Mask)                                                      \
-  ((Mask & 1) + ((Mask & 2) >> 1) + ((Mask & 4) >> 2) + ((Mask & 8) >> 3))
-
-#define HasR(Mask) ((Mask & 1) == 1)
-#define HasG(Mask) ((Mask & 2) >> 1 == 1)
-#define HasB(Mask) ((Mask & 4) >> 2 == 1)
-#define HasA(Mask) ((Mask & 8) >> 3 == 1)
-
-enum class EsimdAtomicOpType : uint16_t {
-  ATOMIC_ADD = 0x0,
-  ATOMIC_SUB = 0x1,
-  ATOMIC_INC = 0x2,
-  ATOMIC_DEC = 0x3,
-  ATOMIC_MIN = 0x4,
-  ATOMIC_MAX = 0x5,
-  ATOMIC_XCHG = 0x6,
-  ATOMIC_CMPXCHG = 0x7,
-  ATOMIC_AND = 0x8,
-  ATOMIC_OR = 0x9,
-  ATOMIC_XOR = 0xa,
-  ATOMIC_MINSINT = 0xb,
-  ATOMIC_MAXSINT = 0xc,
-  ATOMIC_FMAX = 0x10,
-  ATOMIC_FMIN = 0x11,
-  ATOMIC_FCMPWR = 0x12,
-  ATOMIC_PREDEC = 0xff
+/// Represents a pixel's channel.
+enum class rgba_channel : uint8_t { R, G, B, A };
+
+namespace detail {
+template <rgba_channel Ch>
+static inline constexpr uint8_t ch = 1 << static_cast<int>(Ch);
+static inline constexpr uint8_t chR = ch<rgba_channel::R>;
+static inline constexpr uint8_t chG = ch<rgba_channel::G>;
+static inline constexpr uint8_t chB = ch<rgba_channel::B>;
+static inline constexpr uint8_t chA = ch<rgba_channel::A>;
+} // namespace detail
+
+/// Represents a pixel's channel mask - all possible combinations of enabled
+/// channels.
+enum class rgba_channel_mask : uint8_t {
+  R = detail::chR,
+  G = detail::chG,
+  GR = detail::chG | detail::chR,
+  B = detail::chB,
+  BR = detail::chB | detail::chR,
+  BG = detail::chB | detail::chG,
+  BGR = detail::chB | detail::chG | detail::chR,
+  A = detail::chA,
+  AR = detail::chA | detail::chR,
+  AG = detail::chA | detail::chG,
+  AGR = detail::chA | detail::chG | detail::chR,
+  AB = detail::chA | detail::chB,
+  ABR = detail::chA | detail::chB | detail::chR,
+  ABG = detail::chA | detail::chB | detail::chG,
+  ABGR = detail::chA | detail::chB | detail::chG | detail::chR,
+  // For backward compatibility ('ChannelMaskType::ESIMD_R_ENABLE' usage style):
+  __ESIMD_DEPR_ENUM_V(ESIMD_R_ENABLE, rgba_channel_mask::R, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_G_ENABLE, rgba_channel_mask::G, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_GR_ENABLE, rgba_channel_mask::GR, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_B_ENABLE, rgba_channel_mask::B, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_BR_ENABLE, rgba_channel_mask::BR, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_BG_ENABLE, rgba_channel_mask::BG, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_BGR_ENABLE, rgba_channel_mask::BGR, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_A_ENABLE, rgba_channel_mask::A, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_AR_ENABLE, rgba_channel_mask::AR, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_AG_ENABLE, rgba_channel_mask::AG, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_AGR_ENABLE, rgba_channel_mask::AGR, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_AB_ENABLE, rgba_channel_mask::AB, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_ABR_ENABLE, rgba_channel_mask::ABR, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_ABG_ENABLE, rgba_channel_mask::ABG, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ESIMD_ABGR_ENABLE, rgba_channel_mask::ABGR, uint8_t)
 };
 
+#define __ESIMD_DEPR_CONST(old, new)                                           \
+  static inline constexpr auto old __ESIMD_DEPRECATED(new) = new
+
+// For backward compatibility ('ESIMD_R_ENABLE' usage style):
+__ESIMD_DEPR_CONST(ESIMD_R_ENABLE, rgba_channel_mask::R);
+__ESIMD_DEPR_CONST(ESIMD_G_ENABLE, rgba_channel_mask::G);
+__ESIMD_DEPR_CONST(ESIMD_GR_ENABLE, rgba_channel_mask::GR);
+__ESIMD_DEPR_CONST(ESIMD_B_ENABLE, rgba_channel_mask::B);
+__ESIMD_DEPR_CONST(ESIMD_BR_ENABLE, rgba_channel_mask::BR);
+__ESIMD_DEPR_CONST(ESIMD_BG_ENABLE, rgba_channel_mask::BG);
+__ESIMD_DEPR_CONST(ESIMD_BGR_ENABLE, rgba_channel_mask::BGR);
+__ESIMD_DEPR_CONST(ESIMD_A_ENABLE, rgba_channel_mask::A);
+__ESIMD_DEPR_CONST(ESIMD_AR_ENABLE, rgba_channel_mask::AR);
+__ESIMD_DEPR_CONST(ESIMD_AG_ENABLE, rgba_channel_mask::AG);
+__ESIMD_DEPR_CONST(ESIMD_AGR_ENABLE, rgba_channel_mask::AGR);
+__ESIMD_DEPR_CONST(ESIMD_AB_ENABLE, rgba_channel_mask::AB);
+__ESIMD_DEPR_CONST(ESIMD_ABR_ENABLE, rgba_channel_mask::ABR);
+__ESIMD_DEPR_CONST(ESIMD_ABG_ENABLE, rgba_channel_mask::ABG);
+__ESIMD_DEPR_CONST(ESIMD_ABGR_ENABLE, rgba_channel_mask::ABGR);
+
+#undef __ESIMD_DEPR_CONST
+
+// For backward compatibility:
+using ChannelMaskType = rgba_channel_mask;
+
+constexpr int is_channel_enabled(rgba_channel_mask M, rgba_channel Ch) {
+  int Pos = static_cast<int>(Ch);
+  return (static_cast<int>(M) & (1 << Pos)) >> Pos;
+}
+
+constexpr int get_num_channels_enabled(rgba_channel_mask M) {
+  return is_channel_enabled(M, rgba_channel::R) +
+         is_channel_enabled(M, rgba_channel::G) +
+         is_channel_enabled(M, rgba_channel::B) +
+         is_channel_enabled(M, rgba_channel::A);
+}
+
+/// Represents an atomic operation.
+enum class atomic_op : uint8_t {
+  add = 0x0,
+  sub = 0x1,
+  inc = 0x2,
+  dec = 0x3,
+  min = 0x4,
+  max = 0x5,
+  xchg = 0x6,
+  cmpxchg = 0x7,
+  bit_and = 0x8,
+  bit_or = 0x9,
+  bit_xor = 0xa,
+  minsint = 0xb,
+  maxsint = 0xc,
+  fmax = 0x10,
+  fmin = 0x11,
+  fcmpwr = 0x12,
+  predec = 0xff,
+  // For backward compatibility:
+  __ESIMD_DEPR_ENUM_V(ATOMIC_ADD, atomic_op::add, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_SUB, atomic_op::sub, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_INC, atomic_op::inc, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_DEC, atomic_op::dec, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_MIN, atomic_op::min, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_MAX, atomic_op::max, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_XCHG, atomic_op::xchg, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_CMPXCHG, atomic_op::cmpxchg, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_AND, atomic_op::bit_and, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_OR, atomic_op::bit_or, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_XOR, atomic_op::bit_xor, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_MINSINT, atomic_op::minsint, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_MAXSINT, atomic_op::maxsint, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_FMAX, atomic_op::fmax, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_FMIN, atomic_op::fmin, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_FCMPWR, atomic_op::fcmpwr, uint8_t),
+  __ESIMD_DEPR_ENUM_V(ATOMIC_PREDEC, atomic_op::predec, uint8_t)
+};
+
+// For backward compatibility:
+using EsimdAtomicOpType = atomic_op;
+
+// TODO Cache hints APIs are being reworked.
 // L1 or L3 cache hint kinds.
 enum class CacheHint : uint8_t {
   None = 0,
@@ -111,11 +224,22 @@ enum class CacheHint : uint8_t {
   ReadInvalidate = 5
 };
 
-enum class EsimdSbarrierType : uint8_t {
-  WAIT = 0,  // split barrier wait
-  SIGNAL = 1 // split barrier signal
+/// Represents a split barrier action.
+enum class split_barrier_action : uint8_t {
+  wait = 0,   // split barrier wait
+  signal = 1, // split barrier signal
+  // For backward compatibility:
+  __ESIMD_DEPR_ENUM_V(WAIT, split_barrier_action::wait, uint8_t),
+  __ESIMD_DEPR_ENUM_V(SIGNAL, split_barrier_action::signal, uint8_t)
 };
 
+// For backward compatibility:
+using EsimdSbarrierType = split_barrier_action;
+
+#undef __ESIMD_DEPR_ENUM_V
+
+// Since EsimdSbarrierType values are deprecated, these macros will generate
+// deprecation message.
 #define ESIMD_SBARRIER_WAIT EsimdSbarrierType::WAIT
 #define ESIMD_SBARRIER_SIGNAL EsimdSbarrierType::SIGNAL
 
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp
index 397e7db84464d..03ec0ab45b171 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp
@@ -121,22 +121,22 @@ __esimd_block_write(SurfIndAliasTy surf_ind, uint32_t offset,
                     __SEIEED::vector_type_t<Ty, N> vals);
 
 // flat_read4 does flat-address gather4
-template <typename Ty, int N, __SEIEE::ChannelMaskType Mask,
+template <typename Ty, int N, __SEIEE::rgba_channel_mask Mask,
           __SEIEE::CacheHint L1H = __SEIEE::CacheHint::None,
           __SEIEE::CacheHint L3H = __SEIEE::CacheHint::None>
-__SEIEED::vector_type_t<Ty, N * NumChannels(Mask)>
+__SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)>
     SYCL_EXTERNAL SYCL_ESIMD_FUNCTION
     __esimd_flat_read4(__SEIEED::vector_type_t<uint64_t, N> addrs,
                        __SEIEED::vector_type_t<uint16_t, N> pred = 1);
 
 // flat_write does flat-address scatter
-template <typename Ty, int N, __SEIEE::ChannelMaskType Mask,
+template <typename Ty, int N, __SEIEE::rgba_channel_mask Mask,
           __SEIEE::CacheHint L1H = __SEIEE::CacheHint::None,
           __SEIEE::CacheHint L3H = __SEIEE::CacheHint::None>
-SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void
-__esimd_flat_write4(__SEIEED::vector_type_t<uint64_t, N> addrs,
-                    __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)> vals,
-                    __SEIEED::vector_type_t<uint16_t, N> pred = 1);
+SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_flat_write4(
+    __SEIEED::vector_type_t<uint64_t, N> addrs,
+    __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)> vals,
+    __SEIEED::vector_type_t<uint16_t, N> pred = 1);
 
 // Low-level surface-based gather. Collects elements located at given offsets in
 // a surface and returns them as a single \ref simd object. Element can be
@@ -224,14 +224,14 @@ __esimd_surf_write(__SEIEED::vector_type_t<uint16_t, N> pred, int16_t scale,
 // correponsing BE intrinsicics parameter order.
 
 // flat_atomic: flat-address atomic
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N,
+template <__SEIEE::atomic_op Op, typename Ty, int N,
           __SEIEE::CacheHint L1H = __SEIEE::CacheHint::None,
           __SEIEE::CacheHint L3H = __SEIEE::CacheHint::None>
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<Ty, N>
 __esimd_flat_atomic0(__SEIEED::vector_type_t<uint64_t, N> addrs,
                      __SEIEED::vector_type_t<uint16_t, N> pred);
 
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N,
+template <__SEIEE::atomic_op Op, typename Ty, int N,
           __SEIEE::CacheHint L1H = __SEIEE::CacheHint::None,
           __SEIEE::CacheHint L3H = __SEIEE::CacheHint::None>
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<Ty, N>
@@ -239,7 +239,7 @@ __esimd_flat_atomic1(__SEIEED::vector_type_t<uint64_t, N> addrs,
                      __SEIEED::vector_type_t<Ty, N> src0,
                      __SEIEED::vector_type_t<uint16_t, N> pred);
 
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N,
+template <__SEIEE::atomic_op Op, typename Ty, int N,
           __SEIEE::CacheHint L1H = __SEIEE::CacheHint::None,
           __SEIEE::CacheHint L3H = __SEIEE::CacheHint::None>
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<Ty, N>
@@ -253,7 +253,7 @@ SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_barrier();
 
 // generic work-group split barrier
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void
-__esimd_sbarrier(__SEIEE::EsimdSbarrierType flag);
+__esimd_sbarrier(__SEIEE::split_barrier_action flag);
 
 // slm_fence sets the SLM read/write order
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_slm_fence(uint8_t cntl);
@@ -282,32 +282,32 @@ SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void
 __esimd_slm_block_write(uint32_t addr, __SEIEED::vector_type_t<Ty, N> vals);
 
 // slm_read4 does SLM gather4
-template <typename Ty, int N, __SEIEE::ChannelMaskType Mask>
-SYCL_EXTERNAL
-    SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)>
+template <typename Ty, int N, __SEIEE::rgba_channel_mask Mask>
+SYCL_EXTERNAL SYCL_ESIMD_FUNCTION
+    __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)>
     __esimd_slm_read4(__SEIEED::vector_type_t<uint32_t, N> addrs,
                       __SEIEED::vector_type_t<uint16_t, N> pred = 1);
 
 // slm_write4 does SLM scatter4
-template <typename Ty, int N, __SEIEE::ChannelMaskType Mask>
-SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void
-__esimd_slm_write4(__SEIEED::vector_type_t<uint32_t, N> addrs,
-                   __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)> vals,
-                   __SEIEED::vector_type_t<uint16_t, N> pred = 1);
+template <typename Ty, int N, __SEIEE::rgba_channel_mask Mask>
+SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_slm_write4(
+    __SEIEED::vector_type_t<uint32_t, N> addrs,
+    __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)> vals,
+    __SEIEED::vector_type_t<uint16_t, N> pred = 1);
 
 // slm_atomic: SLM atomic
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N>
+template <__SEIEE::atomic_op Op, typename Ty, int N>
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<Ty, N>
 __esimd_slm_atomic0(__SEIEED::vector_type_t<uint32_t, N> addrs,
                     __SEIEED::vector_type_t<uint16_t, N> pred);
 
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N>
+template <__SEIEE::atomic_op Op, typename Ty, int N>
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<Ty, N>
 __esimd_slm_atomic1(__SEIEED::vector_type_t<uint32_t, N> addrs,
                     __SEIEED::vector_type_t<Ty, N> src0,
                     __SEIEED::vector_type_t<uint16_t, N> pred);
 
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N>
+template <__SEIEE::atomic_op Op, typename Ty, int N>
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<Ty, N>
 __esimd_slm_atomic2(__SEIEED::vector_type_t<uint32_t, N> addrs,
                     __SEIEED::vector_type_t<Ty, N> src0,
@@ -547,15 +547,15 @@ __esimd_flat_read(__SEIEED::vector_type_t<uint64_t, N> addrs, int ElemsPerAddr,
   return V;
 }
 
-template <typename Ty, int N, __SEIEE::ChannelMaskType Mask,
+template <typename Ty, int N, __SEIEE::rgba_channel_mask Mask,
           __SEIEE::CacheHint L1H, __SEIEE::CacheHint L3H>
-inline __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)>
+inline __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)>
 __esimd_flat_read4(__SEIEED::vector_type_t<uint64_t, N> addrs,
                    __SEIEED::vector_type_t<uint16_t, N> pred) {
-  __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)> V;
+  __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)> V;
   unsigned int Next = 0;
 
-  if constexpr (HasR(Mask)) {
+  if constexpr (__SEIEE::is_channel_enabled(Mask, __SEIEE::rgba_channel::R)) {
     for (int I = 0; I < N; I++, Next++) {
       if (pred[I]) {
         Ty *Addr = reinterpret_cast<Ty *>(addrs[I]);
@@ -564,7 +564,7 @@ __esimd_flat_read4(__SEIEED::vector_type_t<uint64_t, N> addrs,
     }
   }
 
-  if constexpr (HasG(Mask)) {
+  if constexpr (__SEIEE::is_channel_enabled(Mask, __SEIEE::rgba_channel::G)) {
     for (int I = 0; I < N; I++, Next++) {
       if (pred[I]) {
         Ty *Addr = reinterpret_cast<Ty *>(addrs[I] + sizeof(Ty));
@@ -573,7 +573,7 @@ __esimd_flat_read4(__SEIEED::vector_type_t<uint64_t, N> addrs,
     }
   }
 
-  if constexpr (HasB(Mask)) {
+  if constexpr (__SEIEE::is_channel_enabled(Mask, __SEIEE::rgba_channel::B)) {
     for (int I = 0; I < N; I++, Next++) {
       if (pred[I]) {
         Ty *Addr = reinterpret_cast<Ty *>(addrs[I] + sizeof(Ty) + sizeof(Ty));
@@ -582,7 +582,7 @@ __esimd_flat_read4(__SEIEED::vector_type_t<uint64_t, N> addrs,
     }
   }
 
-  if constexpr (HasA(Mask)) {
+  if constexpr (__SEIEE::is_channel_enabled(Mask, __SEIEE::rgba_channel::A)) {
     for (int I = 0; I < N; I++, Next++) {
       if (pred[I]) {
         Ty *Addr = reinterpret_cast<Ty *>(addrs[I] + sizeof(Ty) + sizeof(Ty) +
@@ -621,16 +621,16 @@ inline void __esimd_flat_write(
   }
 }
 
-template <typename Ty, int N, __SEIEE::ChannelMaskType Mask,
+template <typename Ty, int N, __SEIEE::rgba_channel_mask Mask,
           __SEIEE::CacheHint L1H, __SEIEE::CacheHint L3H>
-inline void
-__esimd_flat_write4(__SEIEED::vector_type_t<uint64_t, N> addrs,
-                    __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)> vals,
-                    __SEIEED::vector_type_t<uint16_t, N> pred) {
-  __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)> V;
+inline void __esimd_flat_write4(
+    __SEIEED::vector_type_t<uint64_t, N> addrs,
+    __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)> vals,
+    __SEIEED::vector_type_t<uint16_t, N> pred) {
+  __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)> V;
   unsigned int Next = 0;
 
-  if constexpr (HasR(Mask)) {
+  if constexpr (__SEIEE::is_channel_enabled(Mask, __SEIEE::rgba_channel::R)) {
     for (int I = 0; I < N; I++, Next++) {
       if (pred[I]) {
         Ty *Addr = reinterpret_cast<Ty *>(addrs[I]);
@@ -639,7 +639,7 @@ __esimd_flat_write4(__SEIEED::vector_type_t<uint64_t, N> addrs,
     }
   }
 
-  if constexpr (HasG(Mask)) {
+  if constexpr (__SEIEE::is_channel_enabled(Mask, __SEIEE::rgba_channel::G)) {
     for (int I = 0; I < N; I++, Next++) {
       if (pred[I]) {
         Ty *Addr = reinterpret_cast<Ty *>(addrs[I] + sizeof(Ty));
@@ -648,7 +648,7 @@ __esimd_flat_write4(__SEIEED::vector_type_t<uint64_t, N> addrs,
     }
   }
 
-  if constexpr (HasB(Mask)) {
+  if constexpr (__SEIEE::is_channel_enabled(Mask, __SEIEE::rgba_channel::B)) {
     for (int I = 0; I < N; I++, Next++) {
       if (pred[I]) {
         Ty *Addr = reinterpret_cast<Ty *>(addrs[I] + sizeof(Ty) + sizeof(Ty));
@@ -657,7 +657,7 @@ __esimd_flat_write4(__SEIEED::vector_type_t<uint64_t, N> addrs,
     }
   }
 
-  if constexpr (HasA(Mask)) {
+  if constexpr (__SEIEE::is_channel_enabled(Mask, __SEIEE::rgba_channel::A)) {
     for (int I = 0; I < N; I++, Next++) {
       if (pred[I]) {
         Ty *Addr = reinterpret_cast<Ty *>(addrs[I] + sizeof(Ty) + sizeof(Ty) +
@@ -823,7 +823,7 @@ __esimd_dp4(__SEIEED::vector_type_t<Ty, N> v1,
 /// TODO
 inline void __esimd_barrier() {}
 
-inline void __esimd_sbarrier(__SEIEE::EsimdSbarrierType flag) {}
+inline void __esimd_sbarrier(__SEIEE::split_barrier_action flag) {}
 
 inline void __esimd_slm_fence(uint8_t cntl) {}
 
@@ -854,23 +854,23 @@ inline void __esimd_slm_block_write(uint32_t addr,
                                     __SEIEED::vector_type_t<Ty, N> vals) {}
 
 // slm_read4 does SLM gather4
-template <typename Ty, int N, __SEIEE::ChannelMaskType Mask>
-inline __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)>
+template <typename Ty, int N, __SEIEE::rgba_channel_mask Mask>
+inline __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)>
 __esimd_slm_read4(__SEIEED::vector_type_t<uint32_t, N> addrs,
                   __SEIEED::vector_type_t<uint16_t, N> pred) {
-  __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)> retv;
+  __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)> retv;
   return retv;
 }
 
 // slm_write4 does SLM scatter4
-template <typename Ty, int N, __SEIEE::ChannelMaskType Mask>
-inline void
-__esimd_slm_write4(__SEIEED::vector_type_t<uint32_t, N> addrs,
-                   __SEIEED::vector_type_t<Ty, N * NumChannels(Mask)> vals,
-                   __SEIEED::vector_type_t<uint16_t, N> pred) {}
+template <typename Ty, int N, __SEIEE::rgba_channel_mask Mask>
+inline void __esimd_slm_write4(
+    __SEIEED::vector_type_t<uint32_t, N> addrs,
+    __SEIEED::vector_type_t<Ty, N * get_num_channels_enabled(Mask)> vals,
+    __SEIEED::vector_type_t<uint16_t, N> pred) {}
 
 // slm_atomic: SLM atomic
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N>
+template <__SEIEE::atomic_op Op, typename Ty, int N>
 inline __SEIEED::vector_type_t<Ty, N>
 __esimd_slm_atomic0(__SEIEED::vector_type_t<uint32_t, N> addrs,
                     __SEIEED::vector_type_t<uint16_t, N> pred) {
@@ -878,7 +878,7 @@ __esimd_slm_atomic0(__SEIEED::vector_type_t<uint32_t, N> addrs,
   return retv;
 }
 
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N>
+template <__SEIEE::atomic_op Op, typename Ty, int N>
 inline __SEIEED::vector_type_t<Ty, N>
 __esimd_slm_atomic1(__SEIEED::vector_type_t<uint32_t, N> addrs,
                     __SEIEED::vector_type_t<Ty, N> src0,
@@ -887,7 +887,7 @@ __esimd_slm_atomic1(__SEIEED::vector_type_t<uint32_t, N> addrs,
   return retv;
 }
 
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N>
+template <__SEIEE::atomic_op Op, typename Ty, int N>
 inline __SEIEED::vector_type_t<Ty, N>
 __esimd_slm_atomic2(__SEIEED::vector_type_t<uint32_t, N> addrs,
                     __SEIEED::vector_type_t<Ty, N> src0,
@@ -897,8 +897,8 @@ __esimd_slm_atomic2(__SEIEED::vector_type_t<uint32_t, N> addrs,
   return retv;
 }
 
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N,
-          __SEIEE::CacheHint L1H, __SEIEE::CacheHint L3H>
+template <__SEIEE::atomic_op Op, typename Ty, int N, __SEIEE::CacheHint L1H,
+          __SEIEE::CacheHint L3H>
 inline __SEIEED::vector_type_t<Ty, N>
 __esimd_flat_atomic0(__SEIEED::vector_type_t<uint64_t, N> addrs,
                      __SEIEED::vector_type_t<uint16_t, N> pred) {
@@ -906,8 +906,8 @@ __esimd_flat_atomic0(__SEIEED::vector_type_t<uint64_t, N> addrs,
   return retv;
 }
 
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N,
-          __SEIEE::CacheHint L1H, __SEIEE::CacheHint L3H>
+template <__SEIEE::atomic_op Op, typename Ty, int N, __SEIEE::CacheHint L1H,
+          __SEIEE::CacheHint L3H>
 inline __SEIEED::vector_type_t<Ty, N>
 __esimd_flat_atomic1(__SEIEED::vector_type_t<uint64_t, N> addrs,
                      __SEIEED::vector_type_t<Ty, N> src0,
@@ -916,8 +916,8 @@ __esimd_flat_atomic1(__SEIEED::vector_type_t<uint64_t, N> addrs,
   return retv;
 }
 
-template <__SEIEE::EsimdAtomicOpType Op, typename Ty, int N,
-          __SEIEE::CacheHint L1H, __SEIEE::CacheHint L3H>
+template <__SEIEE::atomic_op Op, typename Ty, int N, __SEIEE::CacheHint L1H,
+          __SEIEE::CacheHint L3H>
 inline __SEIEED::vector_type_t<Ty, N>
 __esimd_flat_atomic2(__SEIEED::vector_type_t<uint64_t, N> addrs,
                      __SEIEED::vector_type_t<Ty, N> src0,
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
index 7383e756b5490..88a9e0d9c38c3 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
@@ -52,9 +52,9 @@ namespace detail {
 
 template <typename T0, typename T1, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE simd<T0, SZ>
-__esimd_abs_common_internal(simd<T1, SZ> src0, int flag = GENX_NOSAT) {
+__esimd_abs_common_internal(simd<T1, SZ> src0, int flag = saturation_off) {
   simd<T1, SZ> Result = __esimd_abs<T1, SZ>(src0.data());
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T0>(Result);
@@ -64,7 +64,7 @@ template <typename T0, typename T1>
 ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
     detail::is_esimd_scalar<T0>::value && detail::is_esimd_scalar<T1>::value,
     typename sycl::detail::remove_const_t<T0>>
-__esimd_abs_common_internal(T1 src0, int flag = GENX_NOSAT) {
+__esimd_abs_common_internal(T1 src0, int flag = saturation_off) {
   typedef typename sycl::detail::remove_const_t<T0> TT0;
   typedef typename sycl::detail::remove_const_t<T1> TT1;
 
@@ -79,7 +79,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
     !std::is_same<typename sycl::detail::remove_const_t<T0>,
                   typename sycl::detail::remove_const_t<T1>>::value,
     simd<T0, SZ>>
-esimd_abs(simd<T1, SZ> src0, int flag = GENX_NOSAT) {
+esimd_abs(simd<T1, SZ> src0, int flag = saturation_off) {
   return detail::__esimd_abs_common_internal<T0, T1, SZ>(src0, flag);
 }
 
@@ -90,13 +90,13 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_esimd_scalar<T0>::value &&
         detail::is_esimd_scalar<T1>::value,
     typename sycl::detail::remove_const_t<T0>>
-esimd_abs(T1 src0, int flag = GENX_NOSAT) {
+esimd_abs(T1 src0, int flag = saturation_off) {
   return detail::__esimd_abs_common_internal<T0, T1>(src0, flag);
 }
 
 template <typename T1, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE simd<T1, SZ> esimd_abs(simd<T1, SZ> src0,
-                                                  int flag = GENX_NOSAT) {
+                                                  int flag = saturation_off) {
   return detail::__esimd_abs_common_internal<T1, T1, SZ>(src0, flag);
 }
 
@@ -104,7 +104,7 @@ template <typename T1>
 ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
     detail::is_esimd_scalar<T1>::value,
     typename sycl::detail::remove_const_t<T1>>
-esimd_abs(T1 src0, int flag = GENX_NOSAT) {
+esimd_abs(T1 src0, int flag = saturation_off) {
   return detail::__esimd_abs_common_internal<T1, T1>(src0, flag);
 }
 
@@ -115,13 +115,13 @@ ESIMD_NODEBUG ESIMD_INLINE
                                            std::is_integral<T1>::value &&
                                            std::is_integral<U>::value,
                                        simd<T0, SZ>>
-    esimd_shl(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
+    esimd_shl(simd<T1, SZ> src0, U src1, int flag = saturation_off) {
   typedef
       typename detail::computation_type<decltype(src0), U>::type ComputationTy;
   typename detail::simd_type<ComputationTy>::type Src0 = src0;
   typename detail::simd_type<ComputationTy>::type Src1 = src1;
 
-  if (flag != GENX_SAT) {
+  if (flag != saturation_on) {
     if constexpr (std::is_unsigned<T0>::value) {
       if constexpr (std::is_unsigned<T1>::value)
         return __esimd_uushl_sat<T0, T1, SZ>(Src0.data(), Src1.data());
@@ -154,7 +154,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_esimd_scalar<T2>::value && std::is_integral<T0>::value &&
         std::is_integral<T1>::value && std::is_integral<T2>::value,
     typename sycl::detail::remove_const_t<T0>>
-esimd_shl(T1 src0, T2 src1, int flag = GENX_NOSAT) {
+esimd_shl(T1 src0, T2 src1, int flag = saturation_off) {
   typedef typename detail::computation_type<T1, T2>::type ComputationTy;
   typename detail::simd_type<ComputationTy>::type Src0 = src0;
   typename detail::simd_type<ComputationTy>::type Src1 = src1;
@@ -169,7 +169,7 @@ ESIMD_NODEBUG ESIMD_INLINE
                                            std::is_integral<T1>::value &&
                                            std::is_integral<U>::value,
                                        simd<T0, SZ>>
-    esimd_shr(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
+    esimd_shr(simd<T1, SZ> src0, U src1, int flag = saturation_off) {
   typedef
       typename detail::computation_type<decltype(src0), U>::type ComputationTy;
   typename detail::simd_type<ComputationTy>::type Src0 = src0;
@@ -177,7 +177,7 @@ ESIMD_NODEBUG ESIMD_INLINE
   typename detail::simd_type<ComputationTy>::type Result =
       Src0.data() >> Src1.data();
 
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T0>(Result);
@@ -189,7 +189,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_esimd_scalar<T2>::value && std::is_integral<T0>::value &&
         std::is_integral<T1>::value && std::is_integral<T2>::value,
     typename sycl::detail::remove_const_t<T0>>
-esimd_shr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
+esimd_shr(T1 src0, T2 src1, int flag = saturation_off) {
   typedef typename detail::computation_type<T1, T2>::type ComputationTy;
   typename detail::simd_type<ComputationTy>::type Src0 = src0;
   typename detail::simd_type<ComputationTy>::type Src1 = src1;
@@ -276,13 +276,13 @@ ESIMD_NODEBUG ESIMD_INLINE
                                            std::is_integral<T1>::value &&
                                            std::is_integral<U>::value,
                                        simd<T0, SZ>>
-    esimd_lsr(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
+    esimd_lsr(simd<T1, SZ> src0, U src1, int flag = saturation_off) {
   typedef typename detail::computation_type<T1, T1>::type IntermedTy;
   typedef typename std::make_unsigned<IntermedTy>::type ComputationTy;
   simd<ComputationTy, SZ> Src0 = src0;
   simd<ComputationTy, SZ> Result = Src0.data() >> src1.data();
 
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T0>(Result);
@@ -294,7 +294,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_esimd_scalar<T2>::value && std::is_integral<T0>::value &&
         std::is_integral<T1>::value && std::is_integral<T2>::value,
     typename sycl::detail::remove_const_t<T0>>
-esimd_lsr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
+esimd_lsr(T1 src0, T2 src1, int flag = saturation_off) {
   typedef typename detail::computation_type<T1, T2>::type ComputationTy;
   typename detail::simd_type<ComputationTy>::type Src0 = src0;
   typename detail::simd_type<ComputationTy>::type Src1 = src1;
@@ -308,7 +308,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         std::is_integral<T0>::value && std::is_integral<T1>::value &&
         std::is_integral<T2>::value,
     decltype(esimd_lsr<T0>(T2(), T1()))>
-esimd_lsr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
+esimd_lsr(T1 src0, T2 src1, int flag = saturation_off) {
   return esimd_lsr<T0>(src1, src0, flag);
 }
 
@@ -319,13 +319,13 @@ ESIMD_NODEBUG ESIMD_INLINE
                                            std::is_integral<T1>::value &&
                                            std::is_integral<U>::value,
                                        simd<T0, SZ>>
-    esimd_asr(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
+    esimd_asr(simd<T1, SZ> src0, U src1, int flag = saturation_off) {
   typedef typename detail::computation_type<T1, T1>::type IntermedTy;
   typedef typename std::make_signed<IntermedTy>::type ComputationTy;
   simd<ComputationTy, SZ> Src0 = src0;
   simd<ComputationTy, SZ> Result = Src0 >> src1;
 
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T0>(Result);
@@ -337,7 +337,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_esimd_scalar<T2>::value && std::is_integral<T0>::value &&
         std::is_integral<T1>::value && std::is_integral<T2>::value,
     typename sycl::detail::remove_const_t<T0>>
-esimd_asr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
+esimd_asr(T1 src0, T2 src1, int flag = saturation_off) {
   typedef typename detail::computation_type<T1, T2>::type ComputationTy;
   typename detail::simd_type<ComputationTy>::type Src0 = src0;
   typename detail::simd_type<ComputationTy>::type Src1 = src1;
@@ -351,7 +351,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         std::is_integral<T0>::value && std::is_integral<T1>::value &&
         std::is_integral<T2>::value,
     decltype(esimd_asr<T0>(T2(), T1()))>
-esimd_asr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
+esimd_asr(T1 src0, T2 src1, int flag = saturation_off) {
   return esimd_asr<T0>(src1, src0, flag);
 }
 
@@ -505,18 +505,18 @@ esimd_div(simd<typename std::remove_const<RT>, 1> &remainder, T0 src0,
 //
 template <typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ>
-esimd_max(simd<T, SZ> src0, simd<T, SZ> src1, int flag = GENX_NOSAT) {
+esimd_max(simd<T, SZ> src0, simd<T, SZ> src1, int flag = saturation_off) {
   if constexpr (std::is_floating_point<T>::value) {
     auto Result = __esimd_fmax<T, SZ>(src0.data(), src1.data());
-    return (flag == GENX_NOSAT) ? Result : __esimd_satf<T, T, SZ>(Result);
+    return (flag == saturation_off) ? Result : __esimd_satf<T, T, SZ>(Result);
   } else if constexpr (std::is_unsigned<T>::value) {
     auto Result = __esimd_umax<T, SZ>(src0.data(), src1.data());
-    return (flag == GENX_NOSAT) ? Result
-                                : __esimd_uutrunc_sat<T, T, SZ>(Result);
+    return (flag == saturation_off) ? Result
+                                    : __esimd_uutrunc_sat<T, T, SZ>(Result);
   } else {
     auto Result = __esimd_smax<T, SZ>(src0.data(), src1.data());
-    return (flag == GENX_NOSAT) ? Result
-                                : __esimd_sstrunc_sat<T, T, SZ>(Result);
+    return (flag == saturation_off) ? Result
+                                    : __esimd_sstrunc_sat<T, T, SZ>(Result);
   }
 }
 
@@ -524,7 +524,7 @@ template <typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::is_esimd_scalar<T>::value,
                                        simd<T, SZ>>
-    esimd_max(simd<T, SZ> src0, T src1, int flag = GENX_NOSAT) {
+    esimd_max(simd<T, SZ> src0, T src1, int flag = saturation_off) {
   simd<T, SZ> Src1 = src1;
   simd<T, SZ> Result = esimd_max<T>(src0, Src1, flag);
   return Result;
@@ -534,7 +534,7 @@ template <typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::is_esimd_scalar<T>::value,
                                        simd<T, SZ>>
-    esimd_max(T src0, simd<T, SZ> src1, int flag = GENX_NOSAT) {
+    esimd_max(T src0, simd<T, SZ> src1, int flag = saturation_off) {
   simd<T, SZ> Src0 = src0;
   simd<T, SZ> Result = esimd_max<T>(Src0, src1, flag);
   return Result;
@@ -543,7 +543,7 @@ ESIMD_NODEBUG ESIMD_INLINE
 template <typename T>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::is_esimd_scalar<T>::value, T>
-    esimd_max(T src0, T src1, int flag = GENX_NOSAT) {
+    esimd_max(T src0, T src1, int flag = saturation_off) {
   simd<T, 1> Src0 = src0;
   simd<T, 1> Src1 = src1;
   simd<T, 1> Result = esimd_max<T>(Src0, Src1, flag);
@@ -552,18 +552,18 @@ ESIMD_NODEBUG ESIMD_INLINE
 
 template <typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ>
-esimd_min(simd<T, SZ> src0, simd<T, SZ> src1, int flag = GENX_NOSAT) {
+esimd_min(simd<T, SZ> src0, simd<T, SZ> src1, int flag = saturation_off) {
   if constexpr (std::is_floating_point<T>::value) {
     auto Result = __esimd_fmin<T, SZ>(src0.data(), src1.data());
-    return (flag == GENX_NOSAT) ? Result : __esimd_satf<T, T, SZ>(Result);
+    return (flag == saturation_off) ? Result : __esimd_satf<T, T, SZ>(Result);
   } else if constexpr (std::is_unsigned<T>::value) {
     auto Result = __esimd_umin<T, SZ>(src0.data(), src1.data());
-    return (flag == GENX_NOSAT) ? Result
-                                : __esimd_uutrunc_sat<T, T, SZ>(Result);
+    return (flag == saturation_off) ? Result
+                                    : __esimd_uutrunc_sat<T, T, SZ>(Result);
   } else {
     auto Result = __esimd_smin<T, SZ>(src0.data(), src1.data());
-    return (flag == GENX_NOSAT) ? Result
-                                : __esimd_sstrunc_sat<T, T, SZ>(Result);
+    return (flag == saturation_off) ? Result
+                                    : __esimd_sstrunc_sat<T, T, SZ>(Result);
   }
 }
 
@@ -571,7 +571,7 @@ template <typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::is_esimd_scalar<T>::value,
                                        simd<T, SZ>>
-    esimd_min(simd<T, SZ> src0, T src1, int flag = GENX_NOSAT) {
+    esimd_min(simd<T, SZ> src0, T src1, int flag = saturation_off) {
   simd<T, SZ> Src1 = src1;
   simd<T, SZ> Result = esimd_min<T>(src0, Src1, flag);
   return Result;
@@ -581,7 +581,7 @@ template <typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::is_esimd_scalar<T>::value,
                                        simd<T, SZ>>
-    esimd_min(T src0, simd<T, SZ> src1, int flag = GENX_NOSAT) {
+    esimd_min(T src0, simd<T, SZ> src1, int flag = saturation_off) {
   simd<T, SZ> Src0 = src0;
   simd<T, SZ> Result = esimd_min<T>(Src0, src1, flag);
   return Result;
@@ -589,7 +589,7 @@ ESIMD_NODEBUG ESIMD_INLINE
 template <typename T>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::is_esimd_scalar<T>::value, T>
-    esimd_min(T src0, T src1, int flag = GENX_NOSAT) {
+    esimd_min(T src0, T src1, int flag = saturation_off) {
   simd<T, 1> Src0 = src0;
   simd<T, 1> Src1 = src1;
   simd<T, 1> Result = esimd_min<T>(Src0, Src1, flag);
@@ -601,12 +601,12 @@ ESIMD_NODEBUG ESIMD_INLINE
     defined(ESIMD_GEN9) || defined(ESIMD_GEN9_5)
 template <typename T0, typename T1, int SZ, typename U>
 ESIMD_NODEBUG ESIMD_INLINE simd<T0, SZ> esimd_dp2(simd<T1, SZ> src0, U src1,
-                                                  int flag = GENX_NOSAT) {
+                                                  int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
   simd<float, SZ> Src0 = src0;
   simd<float, SZ> Src1 = src1;
   simd<float, SZ> Result = __esimd_dp2(Src0.data(), Src1.data());
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T0>(Result);
@@ -614,12 +614,12 @@ ESIMD_NODEBUG ESIMD_INLINE simd<T0, SZ> esimd_dp2(simd<T1, SZ> src0, U src1,
 
 template <typename T0, typename T1, int SZ, typename U>
 ESIMD_NODEBUG ESIMD_INLINE simd<T0, SZ> esimd_dp3(simd<T1, SZ> src0, U src1,
-                                                  int flag = GENX_NOSAT) {
+                                                  int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
   simd<float, SZ> Src0 = src0;
   simd<float, SZ> Src1 = src1;
   simd<float, SZ> Result = __esimd_dp3(Src0.data(), Src1.data());
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T0>(Result);
@@ -627,12 +627,12 @@ ESIMD_NODEBUG ESIMD_INLINE simd<T0, SZ> esimd_dp3(simd<T1, SZ> src0, U src1,
 
 template <typename T0, typename T1, int SZ, typename U>
 ESIMD_NODEBUG ESIMD_INLINE simd<T0, SZ> esimd_dp4(simd<T1, SZ> src0, U src1,
-                                                  int flag = GENX_NOSAT) {
+                                                  int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
   simd<float, SZ> Src0 = src0;
   simd<float, SZ> Src1 = src1;
   simd<float, SZ> Result = __esimd_dp4(Src0.data(), Src1.data());
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T0>(Result);
@@ -640,12 +640,12 @@ ESIMD_NODEBUG ESIMD_INLINE simd<T0, SZ> esimd_dp4(simd<T1, SZ> src0, U src1,
 
 template <typename T0, typename T1, typename U, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE simd<T0, SZ> esimd_dph(simd<T1, SZ> src0, U src1,
-                                                  int flag = GENX_NOSAT) {
+                                                  int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
   simd<float, SZ> Src0 = src0;
   simd<float, SZ> Src1 = src1;
   simd<float, SZ> Result = __esimd_dph(Src0.data(), Src1.data());
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T0>(Result);
@@ -653,7 +653,7 @@ ESIMD_NODEBUG ESIMD_INLINE simd<T0, SZ> esimd_dph(simd<T1, SZ> src0, U src1,
 
 template <typename RT, typename T1, typename T2, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE simd<RT, SZ>
-esimd_line(simd<T1, 4> src0, simd<T2, SZ> src1, int flag = GENX_NOSAT) {
+esimd_line(simd<T1, 4> src0, simd<T2, SZ> src1, int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
   simd<float, 4> Src0 = src0;
@@ -661,7 +661,7 @@ esimd_line(simd<T1, 4> src0, simd<T2, SZ> src1, int flag = GENX_NOSAT) {
   simd<float, SZ> Result = __esimd_line(Src0.data(), Src1.data());
 
   simd<RT, SZ> Result;
-  if (flag == GENX_SAT)
+  if (flag == saturation_on)
     Result = esimd_sat<RT>(Result);
   else
     Result = Result;
@@ -671,7 +671,7 @@ esimd_line(simd<T1, 4> src0, simd<T2, SZ> src1, int flag = GENX_NOSAT) {
 
 template <typename RT, typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE simd<RT, SZ>
-esimd_line(float P, float Q, simd<T, SZ> src1, int flag = GENX_NOSAT) {
+esimd_line(float P, float Q, simd<T, SZ> src1, int flag = saturation_off) {
   simd<float, 4> Src0 = P;
   Src0(3) = Q;
   return esimd_line<RT>(Src0, src1, flag);
@@ -695,7 +695,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_fp_or_dword_type<U>::value &&
         std::is_floating_point<U>::value,
     simd<T0, SZ>>
-esimd_dp2(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
+esimd_dp2(simd<T1, SZ> src0, U src1, int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
   simd<float, SZ> Src1 = src1;
@@ -704,7 +704,7 @@ esimd_dp2(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
   for (int i = 0; i < SZ; i += 4) {
     Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1];
   }
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T1>(Result);
@@ -717,7 +717,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_fp_or_dword_type<U>::value &&
         std::is_floating_point<U>::value,
     simd<T0, SZ>>
-esimd_dp3(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
+esimd_dp3(simd<T1, SZ> src0, U src1, int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
   simd<float, SZ> Src1 = src1;
@@ -727,7 +727,7 @@ esimd_dp3(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
     Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
                              src0[i + 2] * Src1[i + 2];
   }
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T1>(Result);
@@ -740,7 +740,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_fp_or_dword_type<U>::value &&
         std::is_floating_point<U>::value,
     simd<T0, SZ>>
-esimd_dp4(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
+esimd_dp4(simd<T1, SZ> src0, U src1, int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
   simd<T1, SZ> Src1 = src1;
@@ -751,7 +751,7 @@ esimd_dp4(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
                              src0[i + 2] * Src1[i + 2] +
                              src0[i + 3] * Src1[i + 3];
   }
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T1>(Result);
@@ -763,7 +763,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_fp_or_dword_type<U>::value &&
         std::is_floating_point<U>::value,
     simd<T, SZ>>
-esimd_dph(simd<T, SZ> src0, U src1, int flag = GENX_NOSAT) {
+esimd_dph(simd<T, SZ> src0, U src1, int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
   simd<float, SZ> Src1 = src1;
@@ -773,7 +773,7 @@ esimd_dph(simd<T, SZ> src0, U src1, int flag = GENX_NOSAT) {
     Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
                              src0[i + 2] * Src1[i + 2] + 1.0 * Src1[i + 3];
   }
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T>(Result);
@@ -784,7 +784,7 @@ ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::is_fp_or_dword_type<T>::value &&
                                            std::is_floating_point<T>::value,
                                        simd<T, SZ>>
-    esimd_line(simd<T, 4> src0, simd<T, SZ> src1, int flag = GENX_NOSAT) {
+    esimd_line(simd<T, 4> src0, simd<T, SZ> src1, int flag = saturation_off) {
   static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
   simd<T, SZ> Src1 = src1;
@@ -794,7 +794,7 @@ ESIMD_NODEBUG ESIMD_INLINE
     Result.select<4, 1>(i) = src0[0] * src1[i] + src0[3];
   }
 
-  if (flag == GENX_SAT)
+  if (flag == saturation_on)
     Result = esimd_sat<T>(Result);
 
   return Result;
@@ -805,7 +805,7 @@ ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::is_fp_or_dword_type<T>::value &&
                                            std::is_floating_point<T>::value,
                                        simd<T, SZ>>
-    esimd_line(float P, float Q, simd<T, SZ> src1, int flag = GENX_NOSAT) {
+    esimd_line(float P, float Q, simd<T, SZ> src1, int flag = saturation_off) {
   simd<T, 4> Src0 = P;
   Src0(3) = Q;
   return esimd_line<T>(Src0, src1, flag);
@@ -828,7 +828,7 @@ template <typename T> ESIMD_NODEBUG ESIMD_INLINE T esimd_frc(T src0) {
 // esimd_lzd
 template <typename RT, typename T0, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE simd<RT, SZ> esimd_lzd(simd<T0, SZ> src0,
-                                                  int flag = GENX_NOSAT) {
+                                                  int flag = saturation_off) {
   // Saturation parameter ignored
   simd<uint, SZ> Src0 = src0;
   return __esimd_lzd<uint>(Src0);
@@ -838,7 +838,7 @@ template <typename RT, typename T0>
 ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
     detail::is_esimd_scalar<RT>::value && detail::is_esimd_scalar<T0>::value,
     typename sycl::detail::remove_const_t<RT>>
-esimd_lzd(T0 src0, int flag = GENX_NOSAT) {
+esimd_lzd(T0 src0, int flag = saturation_off) {
   simd<T0, 1> Src0 = src0;
   simd<RT, 1> Result = esimd_lzd<RT>(Src0);
   return Result[0];
@@ -850,14 +850,14 @@ esimd_lzd(T0 src0, int flag = GENX_NOSAT) {
 
 template <int SZ, typename U, typename V>
 ESIMD_NODEBUG ESIMD_INLINE simd<float, SZ>
-esimd_lrp(simd<float, SZ> src0, U src1, V src2, int flag = GENX_NOSAT) {
+esimd_lrp(simd<float, SZ> src0, U src1, V src2, int flag = saturation_off) {
   static_assert(SZ >= 4 && (SZ & 0x3) == 0,
                 "vector size must be a multiple of 4");
   simd<float, SZ> Src1 = src1;
   simd<float, SZ> Src2 = src2;
   simd<float, SZ> Result = __esimd_lrp<SZ>(src0, Src1, Src2);
 
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<float>(Result);
@@ -880,13 +880,13 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_fp_or_dword_type<U>::value &&
         std::is_floating_point<U>::value,
     simd<T, SZ>>
-esimd_lrp(simd<T, SZ> src0, U src1, V src2, int flag = GENX_NOSAT) {
+esimd_lrp(simd<T, SZ> src0, U src1, V src2, int flag = saturation_off) {
 
   simd<float, SZ> Src1 = src1;
   simd<float, SZ> Src2 = src2;
   simd<float, SZ> Result;
   Result = Src1 * src0 + Src2 * (1.0f - src0);
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
   return esimd_sat<T>(Result);
 }
@@ -896,7 +896,7 @@ esimd_lrp(simd<T, SZ> src0, U src1, V src2, int flag = GENX_NOSAT) {
 template <int SZ>
 ESIMD_NODEBUG ESIMD_INLINE simd<float, SZ>
 esimd_pln(simd<float, 4> src0, simd<float, SZ> src1, simd<float, SZ> src2,
-          int flag = GENX_NOSAT) {
+          int flag = saturation_off) {
   static_assert(SZ >= 8 && (SZ & 0x7) == 0,
                 "vector size must be a multiple of 8");
 
@@ -915,7 +915,7 @@ esimd_pln(simd<float, 4> src0, simd<float, SZ> src1, simd<float, SZ> src2,
 
   simd<float, SZ> Result = __esimd_pln<SZ>(src0, Src12.read());
 
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<float>(Result);
@@ -1000,9 +1000,9 @@ esimd_bf_extract(T1 src0, T2 src1, T3 src2) {
 //
 // template <int SZ>
 // simd<float, SZ>
-// ESIMD_INLINE esimd_inv(simd<float, SZ> src0, int flag = GENX_NOSAT) {
+// ESIMD_INLINE esimd_inv(simd<float, SZ> src0, int flag = saturation_off) {
 //   simd<float, SZ> Result = __esimd_inv(src0);
-//   if (flag != GENX_SAT)
+//   if (flag != saturation_on)
 //     return Result;
 //   return __esimd_sat<float>(Result);
 // }
@@ -1010,12 +1010,12 @@ esimd_bf_extract(T1 src0, T2 src1, T3 src2) {
 // template <int N1, int N2>
 // ESIMD_NODEBUG ESIMD_INLINE
 // simd<float, N1 * N2>
-// esimd_inv(matrix<float, N1, N2> src0, int flag = GENX_NOSAT) {
+// esimd_inv(matrix<float, N1, N2> src0, int flag = saturation_off) {
 //   simd<float, N1 * N2> Src0 = src0;
 //   return esimd_inv(Src0, flag);
 // }
 //
-// ESIMD_INLINE float esimd_inv(float src0, int flag = GENX_NOSAT) {
+// ESIMD_INLINE float esimd_inv(float src0, int flag = saturation_off) {
 //   simd<float, 1> Src0 = src0;
 //   simd<float, 1> Result = esimd_inv(Src0, flag);
 //   return Result[0];
@@ -1028,15 +1028,15 @@ esimd_bf_extract(T1 src0, T2 src1, T3 src2) {
 #define ESIMD_INTRINSIC_DEF(type, name)                                        \
   template <int SZ>                                                            \
   ESIMD_NODEBUG ESIMD_INLINE simd<type, SZ> esimd_##name(                      \
-      simd<type, SZ> src0, int flag = GENX_NOSAT) {                            \
+      simd<type, SZ> src0, int flag = saturation_off) {                        \
     simd<type, SZ> Result = __esimd_##name<SZ>(src0.data());                   \
-    if (flag != GENX_SAT)                                                      \
+    if (flag != saturation_on)                                                 \
       return Result;                                                           \
     return esimd_sat<type>(Result);                                            \
   }                                                                            \
   template <typename T = void>                                                 \
   ESIMD_NODEBUG ESIMD_INLINE type esimd_##name(type src0,                      \
-                                               int flag = GENX_NOSAT) {        \
+                                               int flag = saturation_off) {    \
     simd<type, 1> Src0 = src0;                                                 \
     simd<type, 1> Result = esimd_##name(Src0, flag);                           \
     return Result[0];                                                          \
@@ -1058,10 +1058,10 @@ ESIMD_INTRINSIC_DEF(double, sqrt_ieee)
 #define ESIMD_INTRINSIC_DEF(ftype, name)                                       \
   template <int SZ, typename U>                                                \
   ESIMD_NODEBUG ESIMD_INLINE simd<ftype, SZ> esimd_##name(                     \
-      simd<ftype, SZ> src0, U src1, int flag = GENX_NOSAT) {                   \
+      simd<ftype, SZ> src0, U src1, int flag = saturation_off) {               \
     simd<ftype, SZ> Src1 = src1;                                               \
     simd<ftype, SZ> Result = __esimd_##name<SZ>(src0.data(), Src1.data());     \
-    if (flag != GENX_SAT)                                                      \
+    if (flag != saturation_on)                                                 \
       return Result;                                                           \
                                                                                \
     return esimd_sat<ftype>(Result);                                           \
@@ -1070,12 +1070,13 @@ ESIMD_INTRINSIC_DEF(double, sqrt_ieee)
   ESIMD_NODEBUG ESIMD_INLINE                                                   \
       typename sycl::detail::enable_if_t<detail::is_esimd_scalar<U>::value,    \
                                          simd<ftype, SZ>>                      \
-          esimd_##name(U src0, simd<ftype, SZ> src1, int flag = GENX_NOSAT) {  \
+          esimd_##name(U src0, simd<ftype, SZ> src1,                           \
+                       int flag = saturation_off) {                            \
     simd<ftype, SZ> Src0 = src0;                                               \
     return esimd_##name(Src0, src1, flag);                                     \
   }                                                                            \
   ESIMD_NODEBUG ESIMD_INLINE ftype esimd_##name(ftype src0, ftype src1,        \
-                                                int flag = GENX_NOSAT) {       \
+                                                int flag = saturation_off) {   \
     simd<ftype, 1> Src0 = src0;                                                \
     simd<ftype, 1> Src1 = src1;                                                \
     simd<ftype, 1> Result = esimd_##name(Src0, Src1, flag);                    \
@@ -1092,7 +1093,7 @@ ESIMD_INTRINSIC_DEF(double, div_ieee)
 // esimd_sincos
 template <int SZ, typename U>
 ESIMD_NODEBUG ESIMD_INLINE simd<float, SZ>
-esimd_sincos(simd<float, SZ> &dstcos, U src0, int flag = GENX_NOSAT) {
+esimd_sincos(simd<float, SZ> &dstcos, U src0, int flag = saturation_off) {
   dstcos = esimd_cos(src0, flag);
   return esimd_sin(src0, flag);
 }
@@ -1105,7 +1106,7 @@ template <typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<std::is_floating_point<T>::value,
                                        simd<T, SZ>>
-    esimd_atan(simd<T, SZ> src0, int flag = GENX_NOSAT) {
+    esimd_atan(simd<T, SZ> src0, int flag = saturation_off) {
   simd<T, SZ> Src0 = esimd_abs(src0);
 
   simd<ushort, SZ> Neg = src0 < T(0.0);
@@ -1127,7 +1128,7 @@ ESIMD_NODEBUG ESIMD_INLINE
   Result.merge(Result - T(ESIMD_HDR_CONST_PI / 2.0), Gt1);
   Result.merge(Result, Neg);
 
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T>(Result);
@@ -1136,7 +1137,7 @@ ESIMD_NODEBUG ESIMD_INLINE
 template <typename T>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<std::is_floating_point<T>::value, T>
-    esimd_atan(T src0, int flag = GENX_NOSAT) {
+    esimd_atan(T src0, int flag = saturation_off) {
   simd<T, 1> Src0 = src0;
   simd<T, 1> Result = esimd_atan(Src0, flag);
   return Result[0];
@@ -1148,7 +1149,7 @@ template <typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<std::is_floating_point<T>::value,
                                        simd<T, SZ>>
-    esimd_acos(simd<T, SZ> src0, int flag = GENX_NOSAT) {
+    esimd_acos(simd<T, SZ> src0, int flag = saturation_off) {
   simd<T, SZ> Src0 = esimd_abs(src0);
 
   simd<ushort, SZ> Neg = src0 < T(0.0);
@@ -1172,7 +1173,7 @@ ESIMD_NODEBUG ESIMD_INLINE
   Result.merge(T(0.0), TooBig);
   Result.merge(T(ESIMD_HDR_CONST_PI) - Result, Neg);
 
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T>(Result);
@@ -1181,7 +1182,7 @@ ESIMD_NODEBUG ESIMD_INLINE
 template <typename T>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<std::is_floating_point<T>::value, T>
-    esimd_acos(T src0, int flag = GENX_NOSAT) {
+    esimd_acos(T src0, int flag = saturation_off) {
   simd<T, 1> Src0 = src0;
   simd<T, 1> Result = esimd_acos(Src0, flag);
   return Result[0];
@@ -1193,7 +1194,7 @@ template <typename T, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<std::is_floating_point<T>::value,
                                        simd<T, SZ>>
-    esimd_asin(simd<T, SZ> src0, int flag = GENX_NOSAT) {
+    esimd_asin(simd<T, SZ> src0, int flag = saturation_off) {
   simd<ushort, SZ> Neg = src0 < T(0.0);
 
   simd<T, SZ> Result =
@@ -1201,7 +1202,7 @@ ESIMD_NODEBUG ESIMD_INLINE
 
   Result.merge(-Result, Neg);
 
-  if (flag != GENX_SAT)
+  if (flag != saturation_on)
     return Result;
 
   return esimd_sat<T>(Result);
@@ -1210,7 +1211,7 @@ ESIMD_NODEBUG ESIMD_INLINE
 template <typename T>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<std::is_floating_point<T>::value, T>
-    esimd_asin(T src0, int flag = GENX_NOSAT) {
+    esimd_asin(T src0, int flag = saturation_off) {
   simd<T, 1> Src0 = src0;
   simd<T, 1> Result = esimd_asin(Src0, flag);
   return Result[0];
@@ -1222,16 +1223,16 @@ ESIMD_NODEBUG ESIMD_INLINE
 
 #define ESIMD_INTRINSIC_DEF(name)                                              \
   template <typename T, int SZ>                                                \
-  ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ> esimd_##name(simd<float, SZ> src0,    \
-                                                      int flag = GENX_NOSAT) { \
+  ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ> esimd_##name(                         \
+      simd<float, SZ> src0, int flag = saturation_off) {                       \
     simd<float, SZ> Result = __esimd_##name<SZ>(src0.data());                  \
-    if (flag != GENX_SAT)                                                      \
+    if (flag != saturation_on)                                                 \
       return Result;                                                           \
     return esimd_sat<T>(Result);                                               \
   }                                                                            \
   template <typename T>                                                        \
   ESIMD_NODEBUG ESIMD_INLINE T esimd_##name(float src0,                        \
-                                            int flag = GENX_NOSAT) {           \
+                                            int flag = saturation_off) {       \
     simd<float, 1> Src0 = src0;                                                \
     simd<T, 1> Result = esimd_##name<T>(Src0, flag);                           \
     return Result[0];                                                          \
@@ -1274,7 +1275,7 @@ ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<std::is_integral<T>::value,
                                        simd<uint, N>>
     esimd_cbit(simd<T, N> src0) {
-  return __esimd_cbit(src0.data());
+  return __esimd_cbit<T, N>(src0.data());
 }
 
 template <typename T>
@@ -1330,7 +1331,7 @@ template <typename T = void> simd<uint, 4> esimd_rdtsc();
 ///
 /// @param src2 the third source operand of dp4a operation.
 ///
-/// @param flag saturation flag, which has default value of GENX_NOSAT.
+/// @param flag saturation flag, which has default value of saturation_off.
 ///
 /// Returns simd vector of the dp4a operation result.
 ///
@@ -1340,14 +1341,14 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
         detail::is_dword_type<T3>::value && detail::is_dword_type<T4>::value,
     simd<T1, N>>
 esimd_dp4a(simd<T2, N> src0, simd<T3, N> src1, simd<T4, N> src2,
-           int flag = GENX_NOSAT) {
+           int flag = saturation_off) {
   simd<T2, N> Src0 = src0;
   simd<T3, N> Src1 = src1;
   simd<T4, N> Src2 = src2;
   simd<T1, N> Result;
 
 #if defined(__SYCL_DEVICE_ONLY__)
-  if (flag == GENX_NOSAT) {
+  if (flag == saturation_off) {
     if constexpr (std::is_unsigned<T1>::value) {
       if constexpr (std::is_unsigned<T2>::value) {
         Result = __esimd_uudp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
@@ -1388,7 +1389,7 @@ esimd_dp4a(simd<T2, N> src0, simd<T3, N> src1, simd<T4, N> src2,
   simd<T2, N> tmp =
       __esimd_dp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(), Src2.data());
 
-  if (flag == GENX_SAT)
+  if (flag == saturation_on)
     Result = esimd_sat<T1>(tmp);
   else
     Result = convert<T1>(tmp);
@@ -1501,7 +1502,7 @@ ESIMD_INLINE simd<float, N> esimd_atan2_fast(simd<float, N> y, simd<float, N> x,
   a1 += (xy / (x2 + y2 * 0.28f + ESIMD_DBL_EPSILON));
 
   atan2.merge(a1, a0, y2 <= x2);
-  if (flags & GENX_SAT)
+  if (flags & saturation_on)
     atan2 = esimd_sat<float>(atan2);
   return atan2;
 }
@@ -1530,7 +1531,7 @@ ESIMD_INLINE simd<float, N> esimd_atan2(simd<float, N> y, simd<float, N> x,
   v_distance = esimd_sqrt(x * x + y * y);
   mask = (esimd_abs<float>(y) < 0.000001f);
   atan2.merge(v_y0, (2 * esimd_atan((v_distance - x) / y)), mask);
-  if (flags & GENX_SAT)
+  if (flags & saturation_on)
     atan2 = esimd_sat<float>(atan2);
 
   return atan2;
@@ -1548,7 +1549,7 @@ template <> ESIMD_INLINE float esimd_atan2(float y, float x, const uint flags) {
   v_distance = esimd_sqrt<float>(x * x + y * y);
   mask = (esimd_abs<float>(y) < 0.000001f);
   atan2.merge(v_y0, (2 * esimd_atan((v_distance - x) / y)), mask);
-  if (flags & GENX_SAT)
+  if (flags & saturation_on)
     atan2 = esimd_sat<float>(atan2);
 
   return atan2[0];
@@ -1564,7 +1565,7 @@ ESIMD_INLINE simd<float, N> esimd_fmod(simd<float, N> y, simd<float, N> x,
 
   v_quot = convert<int>(y / x);
   fmod = y - x * convert<float>(v_quot);
-  if (flags & GENX_SAT)
+  if (flags & saturation_on)
     fmod = esimd_sat<float>(fmod);
 
   return fmod;
@@ -1577,7 +1578,7 @@ template <> ESIMD_INLINE float esimd_fmod(float y, float x, const uint flags) {
 
   v_quot = (int)(y / x);
   fmod = y - x * v_quot;
-  if (flags & GENX_SAT)
+  if (flags & saturation_on)
     fmod = esimd_sat<float>(fmod);
 
   return fmod[0];
@@ -1618,7 +1619,7 @@ ESIMD_INLINE simd<float, N> esimd_sin_emu(simd<float, N> x, const uint flags) {
                                                     (OneP - x2 * 0.0090909f))));
   fTrig *= sign;
 
-  if (flags & GENX_SAT)
+  if (flags & saturation_on)
     fTrig = esimd_sat<float>(fTrig);
 
   return fTrig;
@@ -1656,7 +1657,7 @@ template <typename T> ESIMD_INLINE float esimd_sin_emu(T x0, const uint flags) {
                                                     (OneP - x2 * 0.0090909f))));
   fTrig *= sign;
 
-  if (flags & GENX_SAT)
+  if (flags & saturation_on)
     fTrig = esimd_sat<float>(fTrig);
 
   return fTrig[0];
@@ -1695,7 +1696,7 @@ ESIMD_INLINE simd<float, N> esimd_cos_emu(simd<float, N> x, const uint flags) {
                                                     (OneP - x2 * 0.0090909f))));
   fTrig *= sign;
 
-  if (flags & GENX_SAT)
+  if (flags & saturation_on)
     fTrig = esimd_sat<float>(fTrig);
 
   return fTrig;
@@ -1732,7 +1733,7 @@ template <typename T> ESIMD_INLINE float esimd_cos_emu(T x0, const uint flags) {
                                                     (OneP - x2 * 0.0090909f))));
   fTrig *= sign;
 
-  if (flags & GENX_SAT)
+  if (flags & saturation_on)
     fTrig = esimd_sat<float>(fTrig);
 
   return fTrig[0];
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp
index ba2b1b939a146..3ab8cacb5b96e 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp
@@ -380,12 +380,12 @@ ESIMD_INLINE ESIMD_NODEBUG void scalar_store(AccessorTy acc, uint32_t offset,
 /// Flat-address gather4.
 /// Only allow simd-16 and simd-32.
 /// \ingroup sycl_esimd
-template <typename T, int n, ChannelMaskType Mask,
+template <typename T, int n, rgba_channel_mask Mask,
           CacheHint L1H = CacheHint::None, CacheHint L3H = CacheHint::None>
-ESIMD_INLINE ESIMD_NODEBUG
-    typename sycl::detail::enable_if_t<(n == 16 || n == 32) && (sizeof(T) == 4),
-                                       simd<T, n * NumChannels(Mask)>>
-    gather4(T *p, simd<uint32_t, n> offsets, simd<uint16_t, n> pred = 1) {
+ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t<
+    (n == 16 || n == 32) && (sizeof(T) == 4),
+    simd<T, n * get_num_channels_enabled(Mask)>>
+gather4(T *p, simd<uint32_t, n> offsets, simd<uint16_t, n> pred = 1) {
 
   simd<uint64_t, n> offsets_i = convert<uint64_t>(offsets);
   simd<uint64_t, n> addrs(reinterpret_cast<uint64_t>(p));
@@ -395,12 +395,12 @@ ESIMD_INLINE ESIMD_NODEBUG
 
 /// Flat-address scatter4.
 /// \ingroup sycl_esimd
-template <typename T, int n, ChannelMaskType Mask,
+template <typename T, int n, rgba_channel_mask Mask,
           CacheHint L1H = CacheHint::None, CacheHint L3H = CacheHint::None>
 ESIMD_INLINE ESIMD_NODEBUG
     typename sycl::detail::enable_if_t<(n == 16 || n == 32) && (sizeof(T) == 4),
                                        void>
-    scatter4(T *p, simd<T, n * NumChannels(Mask)> vals,
+    scatter4(T *p, simd<T, n * get_num_channels_enabled(Mask)> vals,
              simd<uint32_t, n> offsets, simd<uint16_t, n> pred = 1) {
   simd<uint64_t, n> offsets_i = convert<uint64_t>(offsets);
   simd<uint64_t, n> addrs(reinterpret_cast<uint64_t>(p));
@@ -412,7 +412,7 @@ ESIMD_INLINE ESIMD_NODEBUG
 namespace detail {
 /// Check the legality of an atomic call in terms of size and type.
 /// \ingroup sycl_esimd
-template <EsimdAtomicOpType Op, typename T, int N, unsigned NumSrc>
+template <atomic_op Op, typename T, int N, unsigned NumSrc>
 constexpr bool check_atomic() {
   if constexpr (!detail::isPowerOf2(N, 32)) {
     static_assert((detail::isPowerOf2(N, 32)),
@@ -421,8 +421,7 @@ constexpr bool check_atomic() {
   }
 
   // No source operand.
-  if constexpr (Op == EsimdAtomicOpType::ATOMIC_INC ||
-                Op == EsimdAtomicOpType::ATOMIC_DEC) {
+  if constexpr (Op == atomic_op::inc || Op == atomic_op::dec) {
     if constexpr (NumSrc != 0) {
       static_assert(NumSrc == 0, "No source operands are expected");
       return false;
@@ -436,29 +435,22 @@ constexpr bool check_atomic() {
   }
 
   // One source integer operand.
-  if constexpr (Op == EsimdAtomicOpType::ATOMIC_ADD ||
-                Op == EsimdAtomicOpType::ATOMIC_SUB ||
-                Op == EsimdAtomicOpType::ATOMIC_MIN ||
-                Op == EsimdAtomicOpType::ATOMIC_MAX ||
-                Op == EsimdAtomicOpType::ATOMIC_XCHG ||
-                Op == EsimdAtomicOpType::ATOMIC_AND ||
-                Op == EsimdAtomicOpType::ATOMIC_OR ||
-                Op == EsimdAtomicOpType::ATOMIC_XOR ||
-                Op == EsimdAtomicOpType::ATOMIC_MINSINT ||
-                Op == EsimdAtomicOpType::ATOMIC_MAXSINT) {
+  if constexpr (Op == atomic_op::add || Op == atomic_op::sub ||
+                Op == atomic_op::min || Op == atomic_op::max ||
+                Op == atomic_op::xchg || Op == atomic_op::bit_and ||
+                Op == atomic_op::bit_or || Op == atomic_op::bit_xor ||
+                Op == atomic_op::minsint || Op == atomic_op::maxsint) {
     if constexpr (NumSrc != 1) {
       static_assert(NumSrc == 1, "One source operand is expected");
       return false;
     }
-    if constexpr ((Op != EsimdAtomicOpType::ATOMIC_MINSINT &&
-                   Op != EsimdAtomicOpType::ATOMIC_MAXSINT) &&
+    if constexpr ((Op != atomic_op::minsint && Op != atomic_op::maxsint) &&
                   !is_type<T, uint16_t, uint32_t, uint64_t>()) {
       static_assert((is_type<T, uint16_t, uint32_t, uint64_t>()),
                     "Type UW, UD or UQ is expected");
       return false;
     }
-    if constexpr ((Op == EsimdAtomicOpType::ATOMIC_MINSINT ||
-                   Op == EsimdAtomicOpType::ATOMIC_MAXSINT) &&
+    if constexpr ((Op == atomic_op::minsint || Op == atomic_op::maxsint) &&
                   !is_type<T, int16_t, int32_t, int64_t>()) {
       static_assert((is_type<T, int16_t, int32_t, int64_t>()),
                     "Type W, D or Q is expected");
@@ -468,8 +460,7 @@ constexpr bool check_atomic() {
   }
 
   // One source float operand.
-  if constexpr (Op == EsimdAtomicOpType::ATOMIC_FMAX ||
-                Op == EsimdAtomicOpType::ATOMIC_FMIN) {
+  if constexpr (Op == atomic_op::fmax || Op == atomic_op::fmin) {
     if constexpr (NumSrc != 1) {
       static_assert(NumSrc == 1, "One source operand is expected");
       return false;
@@ -484,19 +475,18 @@ constexpr bool check_atomic() {
   }
 
   // Two scouce operands.
-  if constexpr (Op == EsimdAtomicOpType::ATOMIC_CMPXCHG ||
-                Op == EsimdAtomicOpType::ATOMIC_FCMPWR) {
+  if constexpr (Op == atomic_op::cmpxchg || Op == atomic_op::fcmpwr) {
     if constexpr (NumSrc != 2) {
       static_assert(NumSrc == 2, "Two source operands are expected");
       return false;
     }
-    if constexpr (Op == EsimdAtomicOpType::ATOMIC_CMPXCHG &&
+    if constexpr (Op == atomic_op::cmpxchg &&
                   !is_type<T, uint16_t, uint32_t, uint64_t>()) {
       static_assert((is_type<T, uint16_t, uint32_t, uint64_t>()),
                     "Type UW, UD or UQ is expected");
       return false;
     }
-    if constexpr (Op == EsimdAtomicOpType::ATOMIC_FCMPWR &&
+    if constexpr (Op == atomic_op::fcmpwr &&
                   !is_type<T, float, cl::sycl::detail::half_impl::StorageT>()) {
       static_assert(
           (is_type<T, float, cl::sycl::detail::half_impl::StorageT>()),
@@ -520,8 +510,8 @@ constexpr bool check_atomic() {
 
 /// Flat-address atomic, zero source operand: inc and dec.
 /// \ingroup sycl_esimd
-template <EsimdAtomicOpType Op, typename T, int n,
-          CacheHint L1H = CacheHint::None, CacheHint L3H = CacheHint::None>
+template <atomic_op Op, typename T, int n, CacheHint L1H = CacheHint::None,
+          CacheHint L3H = CacheHint::None>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::check_atomic<Op, T, n, 0>(),
                                        simd<T, n>>
@@ -534,8 +524,8 @@ ESIMD_NODEBUG ESIMD_INLINE
 
 /// Flat-address atomic, one source operand, add/sub/min/max etc.
 /// \ingroup sycl_esimd
-template <EsimdAtomicOpType Op, typename T, int n,
-          CacheHint L1H = CacheHint::None, CacheHint L3H = CacheHint::None>
+template <atomic_op Op, typename T, int n, CacheHint L1H = CacheHint::None,
+          CacheHint L3H = CacheHint::None>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::check_atomic<Op, T, n, 1>(),
                                        simd<T, n>>
@@ -550,8 +540,8 @@ ESIMD_NODEBUG ESIMD_INLINE
 
 /// Flat-address atomic, two source operands.
 /// \ingroup sycl_esimd
-template <EsimdAtomicOpType Op, typename T, int n,
-          CacheHint L1H = CacheHint::None, CacheHint L3H = CacheHint::None>
+template <atomic_op Op, typename T, int n, CacheHint L1H = CacheHint::None,
+          CacheHint L3H = CacheHint::None>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::check_atomic<Op, T, n, 2>(),
                                        simd<T, n>>
@@ -608,7 +598,7 @@ inline ESIMD_NODEBUG void esimd_barrier() {
 }
 
 /// Generic work-group split barrier
-inline ESIMD_NODEBUG void esimd_sbarrier(EsimdSbarrierType flag) {
+inline ESIMD_NODEBUG void esimd_sbarrier(split_barrier_action flag) {
   __esimd_sbarrier(flag);
 }
 
@@ -641,21 +631,20 @@ ESIMD_INLINE ESIMD_NODEBUG
 /// SLM gather4.
 ///
 /// Only allow simd-8, simd-16 and simd-32.
-template <typename T, int n, ChannelMaskType Mask>
-ESIMD_INLINE ESIMD_NODEBUG
-    typename sycl::detail::enable_if_t<(n == 8 || n == 16 || n == 32) &&
-                                           (sizeof(T) == 4),
-                                       simd<T, n * NumChannels(Mask)>>
-    slm_load4(simd<uint32_t, n> offsets, simd<uint16_t, n> pred = 1) {
+template <typename T, int n, rgba_channel_mask Mask>
+ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t<
+    (n == 8 || n == 16 || n == 32) && (sizeof(T) == 4),
+    simd<T, n * get_num_channels_enabled(Mask)>>
+slm_load4(simd<uint32_t, n> offsets, simd<uint16_t, n> pred = 1) {
   return __esimd_slm_read4<T, n, Mask>(offsets.data(), pred.data());
 }
 
 /// SLM scatter4.
-template <typename T, int n, ChannelMaskType Mask>
-typename sycl::detail::enable_if_t<
+template <typename T, int n, rgba_channel_mask Mask>
+ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t<
     (n == 8 || n == 16 || n == 32) && (sizeof(T) == 4), void>
-slm_store4(simd<T, n * NumChannels(Mask)> vals, simd<uint32_t, n> offsets,
-           simd<uint16_t, n> pred = 1) {
+slm_store4(simd<T, n * get_num_channels_enabled(Mask)> vals,
+           simd<uint32_t, n> offsets, simd<uint16_t, n> pred = 1) {
   __esimd_slm_write4<T, n, Mask>(offsets.data(), vals.data(), pred.data());
 }
 
@@ -694,7 +683,7 @@ ESIMD_INLINE ESIMD_NODEBUG void slm_block_store(uint32_t offset,
 }
 
 /// SLM atomic, zero source operand: inc and dec.
-template <EsimdAtomicOpType Op, typename T, int n>
+template <atomic_op Op, typename T, int n>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::check_atomic<Op, T, n, 0>(),
                                        simd<T, n>>
@@ -703,7 +692,7 @@ ESIMD_NODEBUG ESIMD_INLINE
 }
 
 /// SLM atomic, one source operand, add/sub/min/max etc.
-template <EsimdAtomicOpType Op, typename T, int n>
+template <atomic_op Op, typename T, int n>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::check_atomic<Op, T, n, 1>(),
                                        simd<T, n>>
@@ -714,7 +703,7 @@ ESIMD_NODEBUG ESIMD_INLINE
 }
 
 /// SLM atomic, two source operands.
-template <EsimdAtomicOpType Op, typename T, int n>
+template <atomic_op Op, typename T, int n>
 ESIMD_NODEBUG ESIMD_INLINE
     typename sycl::detail::enable_if_t<detail::check_atomic<Op, T, n, 2>(),
                                        simd<T, n>>
diff --git a/sycl/test/esimd/enums.cpp b/sycl/test/esimd/enums.cpp
new file mode 100644
index 0000000000000..0f686d1038043
--- /dev/null
+++ b/sycl/test/esimd/enums.cpp
@@ -0,0 +1,31 @@
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
+
+// This test checks compilation of various ESIMD enum types. Those which are
+// deprecated must produce deprecation messages.
+
+#include <sycl/ext/intel/experimental/esimd/common.hpp>
+
+using namespace sycl::ext::intel::experimental::esimd;
+
+void foo() SYCL_ESIMD_FUNCTION {
+  // These should produce deprecation messages:
+  int x;
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
+  x = static_cast<int>(ESIMD_SBARRIER_WAIT);
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
+  x = static_cast<int>(EsimdAtomicOpType::ATOMIC_ADD);
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
+  x = static_cast<int>(ChannelMaskType::ESIMD_R_ENABLE);
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
+  x = static_cast<int>(GENX_NOSAT);
+
+  // These should compile cleanly:
+  x = static_cast<int>(split_barrier_action::wait);
+  x = static_cast<int>(atomic_op::add);
+  x = static_cast<int>(rgba_channel_mask::R);
+  x = static_cast<int>(saturation::off);
+}
diff --git a/sycl/test/esimd/flat_atomic.cpp b/sycl/test/esimd/flat_atomic.cpp
index 777723fdb73bd..86a1800ca2158 100644
--- a/sycl/test/esimd/flat_atomic.cpp
+++ b/sycl/test/esimd/flat_atomic.cpp
@@ -1,5 +1,7 @@
-// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
-// expected-no-diagnostics
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
+
+// This test checks compilation of ESIMD atomic APIs. Those which are deprecated
+// must produce deprecation messages.
 
 #include <CL/sycl.hpp>
 #include <sycl/ext/intel/experimental/esimd.hpp>
@@ -9,22 +11,35 @@
 using namespace sycl::ext::intel::experimental::esimd;
 using namespace cl::sycl;
 
-void kernel0(accessor<uint32_t, 1, access::mode::read_write, access::target::global_buffer> &buf) __attribute__((sycl_device)) {
+void kernel0(accessor<uint32_t, 1, access::mode::read_write,
+                      access::target::global_buffer> &buf) SYCL_ESIMD_FUNCTION {
   simd<uint32_t, 32> offsets(0, 1);
 
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
   flat_atomic<EsimdAtomicOpType::ATOMIC_INC, uint32_t, 32>(buf.get_pointer(), offsets, 1);
+  flat_atomic<atomic_op::inc, uint32_t, 32>(buf.get_pointer(), offsets, 1);
 }
 
-void kernel1(accessor<uint32_t, 1, access::mode::read_write, access::target::global_buffer> &buf) __attribute__((sycl_device)) {
+void kernel1(accessor<uint32_t, 1, access::mode::read_write,
+                      access::target::global_buffer> &buf) SYCL_ESIMD_FUNCTION {
   simd<uint32_t, 32> offsets(0, 1);
   simd<uint32_t, 32> v1(0, 1);
 
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
   flat_atomic<EsimdAtomicOpType::ATOMIC_ADD, uint32_t, 32>(buf.get_pointer(), offsets, v1, 1);
+  flat_atomic<atomic_op::add, uint32_t, 32>(buf.get_pointer(), offsets, v1, 1);
 }
 
-void kernel2(accessor<uint32_t, 1, access::mode::read_write, access::target::global_buffer> &buf) __attribute__((sycl_device)) {
+void kernel2(accessor<uint32_t, 1, access::mode::read_write,
+                      access::target::global_buffer> &buf) SYCL_ESIMD_FUNCTION {
   simd<uint32_t, 32> offsets(0, 1);
   simd<uint32_t, 32> v1(0, 1);
 
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
   flat_atomic<EsimdAtomicOpType::ATOMIC_CMPXCHG, uint32_t, 32>(buf.get_pointer(), offsets, v1, v1, 1);
+  flat_atomic<atomic_op::cmpxchg, uint32_t, 32>(buf.get_pointer(), offsets, v1,
+                                                v1, 1);
 }
diff --git a/sycl/test/esimd/gather4_scatter4.cpp b/sycl/test/esimd/gather4_scatter4.cpp
index e097233e3ceff..31fb3ee241a4f 100644
--- a/sycl/test/esimd/gather4_scatter4.cpp
+++ b/sycl/test/esimd/gather4_scatter4.cpp
@@ -1,5 +1,7 @@
-// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
-// expected-no-diagnostics
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
+
+// This test checks compilation of ESIMD slm gather4/scatter4 APIs. Those which
+// are deprecated must produce deprecation messages.
 
 #include <CL/sycl.hpp>
 #include <sycl/ext/intel/experimental/esimd.hpp>
@@ -10,14 +12,27 @@ using namespace sycl::ext::intel::experimental::esimd;
 using namespace cl::sycl;
 
 void kernel(accessor<int, 1, access::mode::read_write,
-                     access::target::global_buffer> &buf)
-    __attribute__((sycl_device)) {
+                     access::target::global_buffer> &buf) SYCL_ESIMD_FUNCTION {
   simd<uint32_t, 32> offsets(0, 1);
   simd<int, 32 * 4> v1(0, 1);
 
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
   auto v0 = gather4<int, 32, ESIMD_ABGR_ENABLE>(buf.get_pointer(), offsets);
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
+  v0 = gather4<int, 32, ChannelMaskType::ESIMD_ABGR_ENABLE>(buf.get_pointer(),
+                                                            offsets);
+  v0 = gather4<int, 32, rgba_channel_mask::ABGR>(buf.get_pointer(), offsets);
 
   v0 = v0 + v1;
 
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
   scatter4<int, 32, ESIMD_ABGR_ENABLE>(buf.get_pointer(), v0, offsets);
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
+  scatter4<int, 32, ChannelMaskType::ESIMD_ABGR_ENABLE>(buf.get_pointer(), v0,
+                                                        offsets);
+  scatter4<int, 32, rgba_channel_mask::ABGR>(buf.get_pointer(), v0, offsets);
 }
diff --git a/sycl/test/esimd/slm_atomic.cpp b/sycl/test/esimd/slm_atomic.cpp
index 6981cb5c19982..fce324854a6e4 100644
--- a/sycl/test/esimd/slm_atomic.cpp
+++ b/sycl/test/esimd/slm_atomic.cpp
@@ -1,5 +1,7 @@
-// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
-// expected-no-diagnostics
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
+
+// This test checks compilation of ESIMD slm atomic APIs. Those which are
+// deprecated must produce deprecation messages.
 
 #include <CL/sycl.hpp>
 #include <sycl/ext/intel/experimental/esimd.hpp>
@@ -9,22 +11,31 @@
 using namespace sycl::ext::intel::experimental::esimd;
 using namespace cl::sycl;
 
-void kernel0() __attribute__((sycl_device)) {
+void kernel0() SYCL_ESIMD_FUNCTION {
   simd<uint32_t, 32> offsets(0, 1);
 
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
   slm_atomic<EsimdAtomicOpType::ATOMIC_INC, uint32_t, 32>(offsets, 1);
+  slm_atomic<atomic_op::inc, uint32_t, 32>(offsets, 1);
 }
 
-void kernel1() __attribute__((sycl_device)) {
+void kernel1() SYCL_ESIMD_FUNCTION {
   simd<uint32_t, 32> offsets(0, 1);
   simd<uint32_t, 32> v1(0, 1);
 
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
   slm_atomic<EsimdAtomicOpType::ATOMIC_ADD, uint32_t, 32>(offsets, v1, 1);
+  slm_atomic<atomic_op::add, uint32_t, 32>(offsets, v1, 1);
 }
 
-void kernel2() __attribute__((sycl_device)) {
+void kernel2() SYCL_ESIMD_FUNCTION {
   simd<uint32_t, 32> offsets(0, 1);
   simd<uint32_t, 32> v1(0, 1);
 
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
   slm_atomic<EsimdAtomicOpType::ATOMIC_CMPXCHG, uint32_t, 32>(offsets, v1, v1, 1);
+  slm_atomic<atomic_op::cmpxchg, uint32_t, 32>(offsets, v1, v1, 1);
 }
diff --git a/sycl/test/esimd/slm_load4.cpp b/sycl/test/esimd/slm_load4.cpp
index fa02e1bca602e..163c06010c89c 100644
--- a/sycl/test/esimd/slm_load4.cpp
+++ b/sycl/test/esimd/slm_load4.cpp
@@ -1,30 +1,28 @@
-// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
-// expected-no-diagnostics
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
+
+// This test checks compilation of ESIMD slm load4/store4 APIs. Those which are
+// deprecated must produce deprecation messages.
 
-#include <CL/sycl.hpp>
 #include <sycl/ext/intel/experimental/esimd.hpp>
-#include <limits>
-#include <utility>
 
 using namespace sycl::ext::intel::experimental::esimd;
 using namespace cl::sycl;
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_call(Func kernelFunc) {
-  kernelFunc();
-}
-
-void caller() {
-  kernel_call<class EsimdKernel>([=]() SYCL_ESIMD_KERNEL {
-    simd<uint32_t, 32> offsets(0, 1);
-    simd<int, 128> v1(0, 1);
+void caller() SYCL_ESIMD_FUNCTION {
+  simd<uint32_t, 32> offsets(0, 1);
+  simd<int, 128> v1(0, 1);
 
-    slm_init(1024);
+  slm_init(1024);
 
-    auto v0 = slm_load4<int, 32, ESIMD_ABGR_ENABLE>(offsets);
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
+  auto v0 = slm_load4<int, 32, ESIMD_ABGR_ENABLE>(offsets);
+  v0 = slm_load4<int, 32, rgba_channel_mask::ABGR>(offsets);
 
-    v0 = v0 + v1;
+  v0 = v0 + v1;
 
-    slm_store4<int, 32, ESIMD_ABGR_ENABLE>(v0, offsets);
-  });
+  // expected-warning@+2 {{deprecated}}
+  // expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
+  slm_store4<int, 32, ESIMD_ABGR_ENABLE>(v0, offsets);
+  slm_store4<int, 32, rgba_channel_mask::ABGR>(v0, offsets);
 }