[SYCL][ESIMD] Remove deprecated API (#5402)

* [SYCL][ESIMD] Remove deprecated API

Signed-off-by: Sergey Dmitriev <[email protected]>

Author: sndmitriev

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

@@ -62,10 +62,6 @@
 #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)
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -92,11 +88,6 @@ static inline constexpr uint8_t saturation_off =
 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)
-};
-
 /// Represents a pixel's channel.
 enum class rgba_channel : uint8_t { R, G, B, A };
 
@@ -127,49 +118,8 @@ enum class rgba_channel_mask : uint8_t {
   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;
@@ -201,58 +151,14 @@ enum class atomic_op : uint8_t {
   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,
-  Uncached = 1,
-  Cached = 2,
-  WriteBack = 3,
-  WriteThrough = 4,
-  Streaming = 5,
-  ReadInvalidate = 6
 };
 
 /// 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;
-
-// Since EsimdSbarrierType values are deprecated, these macros will generate
-// deprecation message.
-#define ESIMD_SBARRIER_WAIT EsimdSbarrierType::WAIT
-#define ESIMD_SBARRIER_SIGNAL EsimdSbarrierType::SIGNAL
-
 /// Surface index type. Surface is an internal representation of a memory block
 /// addressable by GPU in "stateful" memory model, and each surface is
 /// identified by its "binding table index" - surface index.

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

@@ -150,16 +150,6 @@ struct element_type_traits<T, std::enable_if_t<is_vectorizable_v<T>>> {
   static inline constexpr bool is_floating_point = std::is_floating_point_v<T>;
 };
 
-#ifdef __SYCL_DEVICE_ONLY__
-template <> struct element_type_traits<_Float16, void> {
-  using RawT = _Float16;
-  using EnclosingCppT = _Float16;
-  __SYCL_DEPRECATED("use sycl::half as element type")
-  static inline constexpr bool use_native_cpp_ops = true;
-  static inline constexpr bool is_floating_point = true;
-};
-#endif
-
 // --- Type conversions
 
 // Low-level conversion functions to and from a wrapper element type.

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

@@ -21,11 +21,6 @@ namespace experimental {
 namespace esimd {
 namespace detail {
 
-#define __ESIMD_MASK_DEPRECATION_MSG                                           \
-  "Use of 'simd'/'simd_view<simd,...>' class to represent predicate or mask "  \
-  "is deprecated. Use "                                                        \
-  "'simd_mask'/'simd_view<simd_mask,...>' instead."
-
 template <typename T, int N>
 class simd_mask_impl
     : public detail::simd_obj_impl<
@@ -51,10 +46,6 @@ class simd_mask_impl
   // TODO this should be made inaccessible from user code.
   simd_mask_impl(const raw_vector_type &Val) : base_type(Val) {}
 
-  /// Initializer list constructor.
-  __SYCL_DEPRECATED("use constructor from array, e.g: simd_mask<3> x({0,1,1});")
-  simd_mask_impl(std::initializer_list<T> Ilist) : base_type(Ilist) {}
-
   /// Construct from an array. To allow e.g. simd_mask<N> m({1,0,0,1,...}).
   template <int N1, class = std::enable_if_t<N1 == N>>
   simd_mask_impl(const raw_element_type (&&Arr)[N1]) {
@@ -64,22 +55,6 @@ class simd_mask_impl
   /// Implicit conversion from simd.
   simd_mask_impl(const simd<T, N> &Val) : base_type(Val.data()) {}
 
-  /// Implicit conversion from simd_view<simd,...>.
-  template <
-      // viewed simd class parameters
-      int N1, class T1,
-      // view region
-      class RegionT2,
-      // view element type
-      class T2 = typename __SEIEE::shape_type<RegionT2>::element_type,
-      // view size in elements
-      int N2 = __SEIEE::shape_type<RegionT2>::length,
-      // enable only if view length and element type match this object
-      class = std::enable_if_t<N == N2 && std::is_same_v<T, T2>>>
-  __SYCL_DEPRECATED(__ESIMD_MASK_DEPRECATION_MSG)
-  simd_mask_impl(const simd_view<simd<T1, N1>, RegionT2> &Val)
-      : base_type(Val.read().data()) {}
-
 private:
   static inline constexpr bool mask_size_ok_for_mem_io() {
     constexpr unsigned Sz = sizeof(element_type) * N;

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

@@ -155,22 +155,6 @@ class simd_obj_impl {
     set(Val);
   }
 
-  /// This constructor is deprecated for two reasons:
-  /// 1) it adds confusion between
-  ///   simd s1(1,2); //calls next constructor
-  ///   simd s2{1,2}; //calls this constructor (uniform initialization syntax)
-  /// 2) no compile-time control over the size of the initializer; e.g. the
-  ///    following will compile:
-  ///   simd<int, 2> x = {1, 2, 3, 4};
-  __SYCL_DEPRECATED("use constructor from array, e.g: simd<int,3> x({1,2,3});")
-  simd_obj_impl(std::initializer_list<RawTy> Ilist) noexcept {
-    __esimd_dbg_print(simd_obj_impl(std::initializer_list<RawTy> Ilist));
-    int i = 0;
-    for (auto It = Ilist.begin(); It != Ilist.end() && i < N; ++It) {
-      M_data[i++] = *It;
-    }
-  }
-
   /// Initialize a simd_obj_impl object with an initial value and step.
   simd_obj_impl(Ty Val, Ty Step) noexcept {
     __esimd_dbg_print(simd_obj_impl(Ty Val, Ty Step));
@@ -294,12 +278,6 @@ class simd_obj_impl {
     return RetTy{cast_this_to_derived(), TopRegionTy{0}};
   }
 
-  template <typename EltTy>
-  __SYCL_DEPRECATED("use simd_obj_impl::bit_cast_view.")
-  auto format() & {
-    return bit_cast_view<EltTy>();
-  }
-
   /// View as a 2-dimensional simd_view.
   template <typename EltTy, int Height, int Width>
   auto bit_cast_view() &[[clang::lifetimebound]] {
@@ -308,12 +286,6 @@ class simd_obj_impl {
     return RetTy{cast_this_to_derived(), TopRegionTy{0, 0}};
   }
 
-  template <typename EltTy, int Height, int Width>
-  __SYCL_DEPRECATED("use simd_obj_impl::bit_cast_view.")
-  auto format() & {
-    return bit_cast_view<EltTy, Height, Width>();
-  }
-
   /// 1D region select, apply a region on top of this LValue object.
   ///
   /// \tparam Size is the number of elements to be selected.
@@ -347,22 +319,12 @@ class simd_obj_impl {
   /// Read single element, return value only (not reference).
   Ty operator[](int i) const { return bitcast_to_wrapper_type<Ty>(data()[i]); }
 
-  /// Read single element, return value only (not reference).
-  __SYCL_DEPRECATED("use operator[] form.")
-  Ty operator()(int i) const { return bitcast_to_wrapper_type<Ty>(data()[i]); }
-
   /// Return writable view of a single element.
   simd_view<Derived, region1d_scalar_t<Ty>> operator[](int i)
       [[clang::lifetimebound]] {
     return select<1, 1>(i);
   }
 
-  /// Return writable view of a single element.
-  __SYCL_DEPRECATED("use operator[] form.")
-  simd_view<Derived, region1d_scalar_t<Ty>> operator()(int i) {
-    return select<1, 1>(i);
-  }
-
   // TODO ESIMD_EXPERIMENTAL
   /// Read multiple elements by their indices in vector
   template <int Size>
@@ -398,17 +360,7 @@ class simd_obj_impl {
   /// \tparam Rep is number of times region has to be replicated.
   /// \return replicated simd_obj_impl instance.
   template <int Rep> resize_a_simd_type_t<Derived, Rep * N> replicate() const {
-    return replicate<Rep, N>(0);
-  }
-
-  /// \tparam Rep is number of times region has to be replicated.
-  /// \tparam W is width of src region to replicate.
-  /// \param Offset is offset in number of elements in src region.
-  /// \return replicated simd_obj_impl instance.
-  template <int Rep, int W>
-  __SYCL_DEPRECATED("use simd_obj_impl::replicate_w")
-  resize_a_simd_type_t<Derived, Rep * W> replicate(uint16_t Offset) const {
-    return replicate_w<Rep, W>(Offset);
+    return replicate_w<Rep, N>(0);
   }
 
   /// \tparam Rep is number of times region has to be replicated.
@@ -420,17 +372,6 @@ class simd_obj_impl {
     return replicate_vs_w_hs<Rep, 0, W, 1>(Offset);
   }
 
-  /// \tparam Rep is number of times region has to be replicated.
-  /// \tparam VS vertical stride of src region to replicate.
-  /// \tparam W is width of src region to replicate.
-  /// \param Offset is offset in number of elements in src region.
-  /// \return replicated simd_obj_impl instance.
-  template <int Rep, int VS, int W>
-  __SYCL_DEPRECATED("use simd_obj_impl::replicate_vs_w")
-  resize_a_simd_type_t<Derived, Rep * W> replicate(uint16_t Offset) const {
-    return replicate_vs_w<Rep, VS, W>(Offset);
-  }
-
   /// \tparam Rep is number of times region has to be replicated.
   /// \tparam VS vertical stride of src region to replicate.
   /// \tparam W width of src region to replicate.
@@ -441,18 +382,6 @@ class simd_obj_impl {
     return replicate_vs_w_hs<Rep, VS, W, 1>(Offset);
   }
 
-  /// \tparam Rep is number of times region has to be replicated.
-  /// \tparam VS vertical stride of src region to replicate.
-  /// \tparam W is width of src region to replicate.
-  /// \tparam HS horizontal stride of src region to replicate.
-  /// \param Offset is offset in number of elements in src region.
-  /// \return replicated simd_obj_impl instance.
-  template <int Rep, int VS, int W, int HS>
-  __SYCL_DEPRECATED("use simd_obj_impl::replicate_vs_w_hs")
-  resize_a_simd_type_t<Derived, Rep * W> replicate(uint16_t Offset) const {
-    return replicate_vs_w_hs<Rep, VS, W, HS>(Offset);
-  }
-
   /// \tparam Rep is number of times region has to be replicated.
   /// \tparam VS vertical stride of src region to replicate.
   /// \tparam W is width of src region to replicate.