[SYCL][ESIMD][ABI-Break] Remove deprecated APIs (#9948)

Author: turinevgeny

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

@@ -153,10 +153,8 @@ enum class atomic_op : uint8_t {
   dec = 0x3,
   /// Minimum: <code>*addr = min(*addr, src0)</code>.
   umin = 0x4,
-  min __SYCL_DEPRECATED("use umin") = umin,
   /// Maximum: <code>*addr = max(*addr, src0)</code>.
   umax = 0x5,
-  max __SYCL_DEPRECATED("use smax") = umax,
   /// Exchange. <code>*addr == src0;</code>
   xchg = 0x6,
   /// Compare and exchange. <code>if (*addr == src0) *sddr = src1;</code>
@@ -169,10 +167,8 @@ enum class atomic_op : uint8_t {
   bit_xor = 0xa,
   /// Minimum (signed integer): <code>*addr = min(*addr, src0)</code>.
   smin = 0xb,
-  minsint __SYCL_DEPRECATED("use smin") = smin,
   /// Maximum (signed integer): <code>*addr = max(*addr, src0)</code>.
   smax = 0xc,
-  maxsint __SYCL_DEPRECATED("use smax") = 0xc,
   /// Minimum (floating point): <code>*addr = min(*addr, src0)</code>.
   fmax __SYCL_DEPRECATED("fmax" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x10,
   /// Maximum (floating point): <code>*addr = max(*addr, src0)</code>.

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

@@ -377,12 +377,6 @@ class simd_obj_impl {
   /// with l-value contexts in inline assembly.
   raw_vector_type &data_ref() { return M_data; }
 
-  /// Commit the current stored underlying raw vector to memory.
-  /// This is required when using inline assembly with private global variables.
-  __SYCL_DEPRECATED(
-      "commit is deprecated and will be removed in a future release")
-  void commit() {}
-
   /// @return Newly constructed (from the underlying data) object of the Derived
   /// type.
   Derived read() const { return Derived{data()}; }

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

@@ -740,17 +740,6 @@ gather_rgba(const T *p, Toffset offset, simd_mask<N> mask = 1) {
   return gather_rgba<RGBAMask, T, N>(p, simd<Toffset, N>(offset), mask);
 }
 
-template <typename T, int N, rgba_channel_mask RGBAMask>
-__SYCL_DEPRECATED("use gather_rgba<rgba_channel_mask>()")
-__ESIMD_API std::enable_if_t<
-    (N == 8 || N == 16 || N == 32) && sizeof(T) == 4,
-    simd<T, N * get_num_channels_enabled(
-                    RGBAMask)>> gather_rgba(const T *p,
-                                            simd<uint32_t, N> offsets,
-                                            simd_mask<N> mask = 1) {
-  return gather_rgba<RGBAMask>(p, offsets, mask);
-}
-
 namespace detail {
 template <rgba_channel_mask M> static void validate_rgba_write_channel_mask() {
   using CM = rgba_channel_mask;

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

@@ -21,28 +21,17 @@ namespace ext::intel::esimd::xmx {
 /// it is deducible from the element types of input matrices.
 enum class dpas_argument_type {
   Invalid = 0,
-  u1 = 1, // unsigned 1 bit
-  U1 __SYCL_DEPRECATED("use u1") = u1,
-  s1 = 2, // signed 1 bit
-  S1 __SYCL_DEPRECATED("use s1") = s1,
-  u2 = 3, // unsigned 2 bits
-  U2 __SYCL_DEPRECATED("use u2") = u2,
-  s2 = 4, // signed 2 bits
-  S2 __SYCL_DEPRECATED("use s2") = s2,
-  u4 = 5, // unsigned 4 bits
-  U4 __SYCL_DEPRECATED("use u4") = u4,
-  s4 = 6, // signed 4 bits
-  S4 __SYCL_DEPRECATED("use s4") = s4,
-  u8 = 7, // unsigned 8 bits
-  U8 __SYCL_DEPRECATED("use u8") = u8,
-  s8 = 8, // signed 8 bits
-  S8 __SYCL_DEPRECATED("use s8") = s8,
-  bf16 = 9, // bfloat 16
-  BF16 __SYCL_DEPRECATED("use bf16") = bf16,
+  u1 = 1,    // unsigned 1 bit
+  s1 = 2,    // signed 1 bit
+  u2 = 3,    // unsigned 2 bits
+  s2 = 4,    // signed 2 bits
+  u4 = 5,    // unsigned 4 bits
+  s4 = 6,    // signed 4 bits
+  u8 = 7,    // unsigned 8 bits
+  s8 = 8,    // signed 8 bits
+  bf16 = 9,  // bfloat 16
   fp16 = 10, // half float
-  FP16 __SYCL_DEPRECATED("use fp16") = fp16,
   tf32 = 12, // tensorfloat 32
-  TF32 __SYCL_DEPRECATED("use tf32") = tf32
 };
 
 } // namespace ext::intel::esimd::xmx

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

@@ -24,10 +24,6 @@ namespace ext::intel::experimental::esimd {
 /// @addtogroup sycl_esimd_core
 /// @{
 
-using argument_type
-    __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas_argument_type") =
-        __ESIMD_NS::xmx::dpas_argument_type;
-
 /// The scope that lsc_fence operation should apply to
 /// Supported platforms: DG2, PVC
 enum class lsc_scope : uint8_t {

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

@@ -1664,143 +1664,6 @@ __ESIMD_NS::simd<T, N> dp4(__ESIMD_NS::simd<T, N> v1,
 /// APIs below are used to implement dot product accumulate systolic functions
 /// @ingroup sycl_esimd
 
-/// @addtogroup sycl_esimd_systolic_array_api
-/// @{
-/// DPAS
-/// @param src0 is the source operand that represents accumulator for the dpas
-/// function
-/// @param src1 is the first source perand with data precision type specified
-/// by src1_precision.
-/// @param src2 is the second source operand with data precision type specified
-/// by src2_precision.
-/// @param sat enables/disables the saturation (off by default). Possible
-/// values: saturation_on/saturation_off.
-/// @return the vector value of DPAS computation result.
-template <argument_type src1_precision, argument_type src2_precision,
-          typename T, int systolic_depth, int repeat_count, typename T0,
-          typename T1, typename T2, int N, int N1, int N2,
-          typename Sat = __ESIMD_NS::saturation_off_tag>
-__SYCL_DEPRECATED("use sycl::ext::intel::esimd::native::dpas()")
-__ESIMD_API __ESIMD_NS::simd<T, N> dpas(
-    __ESIMD_NS::simd<T0, N> src0, __ESIMD_NS::simd<T1, N1> src1,
-    __ESIMD_NS::simd<T2, N2> src2,
-    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
-  auto result =
-      __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T0, T1, T2,
-                            src1_precision, src2_precision>(src0, src1, src2);
-  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
-    return result;
-  else
-    return __ESIMD_NS::saturate<T>(result);
-}
-
-/// DPAS
-/// @param src0 is the source operand that represents accumulator for the dpas
-/// function, which must have the same type as return value
-/// @param src1 is the first source perand with data precision type specified
-/// by src1_precision.
-/// @param src2 is the second source operand with data precision type specified
-/// by src2_precision.
-/// @param sat enables/disables the saturation (off by default). Possible
-/// values: saturation_on/saturation_off.
-/// @return the vector value of DPAS computation result.
-template <argument_type src1_precision, argument_type src2_precision,
-          int systolic_depth, int repeat_count, typename T, typename T1,
-          typename T2, int N, int N1, int N2,
-          typename Sat = __ESIMD_NS::saturation_off_tag>
-__SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
-__ESIMD_API __ESIMD_NS::simd<T, N> dpas(
-    __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
-    __ESIMD_NS::simd<T2, N2> src2,
-    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
-  return dpas<src1_precision, src2_precision, T, systolic_depth, repeat_count>(
-      src0, src1, src2, sat);
-}
-
-/// DPAS
-/// @param src1 is the first source perand with data precision type specified
-/// by src1_precision.
-/// @param src2 is the second source operand with data precision type specified
-/// by src2_precision.
-/// @param sat enables/disables the saturation (off by default). Possible
-/// values: saturation_on/saturation_off.
-/// @return the vector value of DPAS computation result.
-template <argument_type src1_precision, argument_type src2_precision,
-          int systolic_depth, int repeat_count, typename T, typename T1,
-          typename T2, int N, int N1, int N2,
-          typename Sat = __ESIMD_NS::saturation_off_tag>
-__SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
-__ESIMD_API __ESIMD_NS::simd<T, N> dpas(
-    __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
-    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
-
-  __ESIMD_NS::simd<T, N> result =
-      __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T1, T2,
-                            src1_precision, src2_precision>(src1, src2);
-
-  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
-    return result;
-  else
-    return __ESIMD_NS::saturate<T>(result);
-}
-
-/// DPASW
-/// @param src0 is the source operand that represents accumulator for the dpas
-/// function, which must have the same type as return value.
-/// @param src1 is the first source perand with data precision type specified
-/// by src1_precision.
-/// @param src2 is the second source operand with data precision type specified
-/// by src2_precision.
-/// @param sat enables/disables the saturation (off by default). Possible
-/// values: saturation_on/saturation_off.
-/// @return the vector value of DPAS computation result.
-template <argument_type src1_precision, argument_type src2_precision,
-          int systolic_depth, int repeat_count, typename T, typename T1,
-          typename T2, int N, int N1, int N2,
-          typename Sat = __ESIMD_NS::saturation_off_tag>
-__SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
-__ESIMD_API __ESIMD_NS::simd<T, N> dpasw(
-    __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
-    __ESIMD_NS::simd<T2, N2> src2,
-    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
-
-  __ESIMD_NS::simd<T, N> result =
-      __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
-                             src1_precision, src2_precision>(src0, src1, src2);
-  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
-    return result;
-  else
-    return __ESIMD_NS::saturate<T>(result);
-}
-
-/// DPASW2
-/// @param src1 is the first source perand with data precision type specified
-/// by src1_precision.
-/// @param src2 is the second source operand with data precision type specified
-/// by src2_precision.
-/// @param sat enables/disables the saturation (off by default). Possible
-/// values: saturation_on/saturation_off.
-/// @return the vector value of DPAS computation result.
-template <argument_type src1_precision, argument_type src2_precision,
-          int systolic_depth, int repeat_count, typename T, typename T1,
-          typename T2, int N, int N1, int N2,
-          typename Sat = __ESIMD_NS::saturation_off_tag>
-__SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
-__ESIMD_API __ESIMD_NS::simd<T, N> dpasw2(
-    __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
-    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
-
-  __ESIMD_NS::simd<T, N> result =
-      __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
-                             src1_precision, src2_precision>(src1, src2);
-
-  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
-    return result;
-  else
-    return __ESIMD_NS::saturate<T>(result);
-}
-/// @} sycl_esimd_systolic_array_api
-
 /// @addtogroup sycl_esimd_logical
 /// @{
 

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

@@ -92,19 +92,6 @@ raw_sends(__ESIMD_NS::simd<T1, n1> msgDst, __ESIMD_NS::simd<T2, n2> msgSrc0,
       msgDesc, msgSrc0.data(), msgSrc1.data(), msgDst.data());
 }
 
-template <typename T1, int n1, typename T2, int n2, typename T3, int n3,
-          int N = 16>
-__SYCL_DEPRECATED("raw_sends_load is deprecated. Use raw_sends")
-__ESIMD_API __ESIMD_NS::simd<T1, n1> raw_sends_load(
-    __ESIMD_NS::simd<T1, n1> msgDst, __ESIMD_NS::simd<T2, n2> msgSrc0,
-    __ESIMD_NS::simd<T3, n3> msgSrc1, uint32_t exDesc, uint32_t msgDesc,
-    uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1,
-    uint8_t numDst, uint8_t isEOT = 0, uint8_t isSendc = 0,
-    __ESIMD_NS::simd_mask<N> mask = 1) {
-  return raw_sends(msgDst, msgSrc0, msgSrc1, exDesc, msgDesc, execSize, sfid,
-                   numSrc0, numSrc1, numDst, isEOT, isSendc);
-}
-
 /// Raw send.
 ///
 /// @param msgDst is the old value of the destination operand.
@@ -146,17 +133,6 @@ raw_send(__ESIMD_NS::simd<T1, n1> msgDst, __ESIMD_NS::simd<T2, n2> msgSrc0,
       msgSrc0.data(), msgDst.data());
 }
 
-template <typename T1, int n1, typename T2, int n2, int N = 16>
-__SYCL_DEPRECATED("raw_send_load is deprecated. Use raw_send")
-__ESIMD_API __ESIMD_NS::simd<T1, n1> raw_send_load(
-    __ESIMD_NS::simd<T1, n1> msgDst, __ESIMD_NS::simd<T2, n2> msgSrc0,
-    uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid,
-    uint8_t numSrc0, uint8_t numDst, uint8_t isEOT = 0, uint8_t isSendc = 0,
-    __ESIMD_NS::simd_mask<N> mask = 1) {
-  return raw_send(msgDst, msgSrc0, exDesc, msgDesc, execSize, sfid, numSrc0,
-                  numDst, isEOT, isSendc, mask);
-}
-
 /// Raw sends. "s" suffix designates "split" variant - i.e. two sources.
 ///
 /// @param msgSrc0 is the first source operand of send message.
@@ -197,19 +173,6 @@ raw_sends(__ESIMD_NS::simd<T1, n1> msgSrc0, __ESIMD_NS::simd<T2, n2> msgSrc1,
       msgSrc0.data(), msgSrc1.data());
 }
 
-template <typename T1, int n1, typename T2, int n2, int N = 16>
-__SYCL_DEPRECATED("raw_sends_store is deprecated. Use raw_sends")
-__ESIMD_API
-    void raw_sends_store(__ESIMD_NS::simd<T1, n1> msgSrc0,
-                         __ESIMD_NS::simd<T2, n2> msgSrc1, uint32_t exDesc,
-                         uint32_t msgDesc, uint8_t execSize, uint8_t sfid,
-                         uint8_t numSrc0, uint8_t numSrc1, uint8_t isEOT = 0,
-                         uint8_t isSendc = 0,
-                         __ESIMD_NS::simd_mask<N> mask = 1) {
-  raw_sends(msgSrc0, msgSrc1, exDesc, msgDesc, execSize, sfid, numSrc0, numSrc1,
-            isEOT, isSendc, mask);
-}
-
 /// Raw send. Generates a \c send or \c sendc instruction for the message
 /// gateway.
 ///
@@ -242,17 +205,6 @@ raw_send(__ESIMD_NS::simd<T1, n1> msgSrc0, uint32_t exDesc, uint32_t msgDesc,
                                             msgSrc0.data());
 }
 
-template <typename T1, int n1, int N = 16>
-__SYCL_DEPRECATED("raw_send_store is deprecated. Use raw_send")
-__ESIMD_API
-    void raw_send_store(__ESIMD_NS::simd<T1, n1> msgSrc0, uint32_t exDesc,
-                        uint32_t msgDesc, uint8_t execSize, uint8_t sfid,
-                        uint8_t numSrc0, uint8_t isEOT = 0, uint8_t isSendc = 0,
-                        __ESIMD_NS::simd_mask<N> mask = 1) {
-  raw_send(msgSrc0, exDesc, msgDesc, execSize, sfid, numSrc0, isEOT, isSendc,
-           mask);
-}
-
 /// @} sycl_esimd_raw_send
 
 #endif // !__ESIMD_FORCE_STATELESS_MEM
@@ -2309,22 +2261,6 @@ lsc_load_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight,
   }
 }
 
-template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1,
-          bool Transposed = false, bool Transformed = false,
-          cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
-          int N = detail::get_lsc_block_2d_data_size<
-              T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>
-__SYCL_DEPRECATED("use lsc_load_2d()")
-__ESIMD_API __ESIMD_NS::simd<T, N> lsc_load2d(const T *Ptr,
-                                              unsigned SurfaceWidth,
-                                              unsigned SurfaceHeight,
-                                              unsigned SurfacePitch, int X,
-                                              int Y) {
-  return lsc_load_2d<T, BlockWidth, BlockHeight, NBlocks, Transposed,
-                     Transformed, L1H, L3H>(Ptr, SurfaceWidth, SurfaceHeight,
-                                            SurfacePitch, X, Y);
-}
-
 /// 2D USM pointer block prefetch.
 /// Supported platforms: PVC
 /// VISA instruction: lsc_load_block2d.ugm
@@ -2368,17 +2304,6 @@ __ESIMD_API void lsc_prefetch_2d(const T *Ptr, unsigned SurfaceWidth,
       pred.data(), surf_addr, SurfaceWidth, SurfaceHeight, SurfacePitch, X, Y);
 }
 
-template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1,
-          cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
-          int N = detail::get_lsc_block_2d_data_size<
-              T, NBlocks, BlockHeight, BlockWidth, false, false>()>
-__SYCL_DEPRECATED("use lsc_prefetch_2d()")
-__ESIMD_API void lsc_prefetch2d(const T *Ptr, unsigned SurfaceWidth,
-                                unsigned SurfaceHeight, unsigned SurfacePitch,
-                                int X, int Y) {
-  lsc_prefetch_2d<T, BlockWidth, BlockHeight, NBlocks, L1H, L3H>(
-      Ptr, SurfaceWidth, SurfaceHeight, SurfacePitch, X, Y);
-}
 /// 2D USM pointer block store.
 /// Supported platforms: PVC
 /// VISA instruction: lsc_store_block2d.ugm
@@ -2439,18 +2364,6 @@ __ESIMD_API void lsc_store_2d(T *Ptr, unsigned SurfaceWidth,
       Raw.data());
 }
 
-template <typename T, int BlockWidth, int BlockHeight = 1,
-          cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none,
-          int N = detail::get_lsc_block_2d_data_size<
-              T, 1u, BlockHeight, BlockWidth, false, false>()>
-__SYCL_DEPRECATED("use lsc_store_2d()")
-__ESIMD_API void lsc_store2d(T *Ptr, unsigned SurfaceWidth,
-                             unsigned SurfaceHeight, unsigned SurfacePitch,
-                             int X, int Y, __ESIMD_NS::simd<T, N> Vals) {
-  lsc_store_2d<T, BlockWidth, BlockHeight, L1H, L3H>(
-      Ptr, SurfaceWidth, SurfaceHeight, SurfacePitch, X, Y, Vals);
-}
-
 /// <summary>
 ///  Container class to hold parameters for \c load2d/store2d \c functions
 /// </summary>