[SYCL] Remove sycl_ext_intel_usm_address_spaces extension (#20916)

Remove in scope of #16929

Author: againull

sycl/doc/design/CompilerAndRuntimeDesign.md

@@ -983,8 +983,6 @@ space attributes in SYCL mode:
 | Address space attribute | SYCL address_space enumeration |
 |-------------------------|--------------------------------|
 | `__attribute__((opencl_global))` | global_space, constant_space |
-| `__attribute__((opencl_global_host))` | ext_intel_global_host_space |
-| `__attribute__((opencl_global_device))` | ext_intel_global_device_space |
 | `__attribute__((opencl_local))` | local_space |
 | `__attribute__((opencl_private))` | private_space |
 | `__attribute__((opencl_constant))` | N/A

…cl_ext_intel_usm_address_spaces.asciidoc …cl_ext_intel_usm_address_spaces.asciidocsycl/doc/extensions/supported/sycl_ext_intel_usm_address_spaces.asciidoc renamed to sycl/doc/extensions/removed/sycl_ext_intel_usm_address_spaces.asciidoc sycl/doc/extensions/supported/sycl_ext_intel_usm_address_spaces.asciidoc renamed to sycl/doc/extensions/removed/sycl_ext_intel_usm_address_spaces.asciidoc

@@ -54,9 +54,7 @@ enum class address_space : int {
                                        "space is deprecated since SYCL 2020") =
       2,
   local_space = 3,
-  ext_intel_global_device_space = 4,
-  ext_intel_global_host_space = 5,
-  generic_space = 6, // TODO generic_space address space is not supported yet
+  generic_space = 4, // TODO generic_space address space is not supported yet
 };
 
 enum class decorated : int { no = 0, yes = 1, legacy = 2 };
@@ -112,20 +110,11 @@ template <> struct NegateDecorated<access::decorated::no> {
 
 #ifdef __SYCL_DEVICE_ONLY__
 #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global))
-#ifdef __ENABLE_USM_ADDR_SPACE__
-#define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device))
-#define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host))
-#else
-#define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global))
-#define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global))
-#endif // __ENABLE_USM_ADDR_SPACE__
 #define __OPENCL_LOCAL_AS__ __attribute__((opencl_local))
 #define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant))
 #define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private))
 #else
 #define __OPENCL_GLOBAL_AS__
-#define __OPENCL_GLOBAL_DEVICE_AS__
-#define __OPENCL_GLOBAL_HOST_AS__
 #define __OPENCL_LOCAL_AS__
 #define __OPENCL_CONSTANT_AS__
 #define __OPENCL_PRIVATE_AS__
@@ -136,13 +125,6 @@ template <access::target accessTarget> struct TargetToAS {
       access::address_space::global_space;
 };
 
-#ifdef __ENABLE_USM_ADDR_SPACE__
-template <> struct TargetToAS<access::target::device> {
-  constexpr static access::address_space AS =
-      access::address_space::ext_intel_global_device_space;
-};
-#endif // __ENABLE_USM_ADDR_SPACE__
-
 template <> struct TargetToAS<access::target::local> {
   constexpr static access::address_space AS =
       access::address_space::local_space;
@@ -171,18 +153,6 @@ struct DecoratedType<ElementType, access::address_space::global_space> {
   using type = __OPENCL_GLOBAL_AS__ ElementType;
 };
 
-template <typename ElementType>
-struct DecoratedType<ElementType,
-                     access::address_space::ext_intel_global_device_space> {
-  using type = __OPENCL_GLOBAL_DEVICE_AS__ ElementType;
-};
-
-template <typename ElementType>
-struct DecoratedType<ElementType,
-                     access::address_space::ext_intel_global_host_space> {
-  using type = __OPENCL_GLOBAL_HOST_AS__ ElementType;
-};
-
 template <typename ElementType>
 struct DecoratedType<ElementType, access::address_space::constant_space> {
   // Current implementation of address spaces handling leads to possibility
@@ -211,18 +181,6 @@ template <class T> struct deduce_AS_impl {
       access::address_space::generic_space;
 };
 
-#ifdef __ENABLE_USM_ADDR_SPACE__
-template <class T> struct deduce_AS_impl<__OPENCL_GLOBAL_DEVICE_AS__ T> {
-  static constexpr access::address_space value =
-      access::address_space::ext_intel_global_device_space;
-};
-
-template <class T> struct deduce_AS_impl<__OPENCL_GLOBAL_HOST_AS__ T> {
-  static constexpr access::address_space value =
-      access::address_space::ext_intel_global_host_space;
-};
-#endif // __ENABLE_USM_ADDR_SPACE__
-
 template <class T> struct deduce_AS_impl<__OPENCL_GLOBAL_AS__ T> {
   static constexpr access::address_space value =
       access::address_space::global_space;
@@ -259,19 +217,6 @@ template <typename T> struct remove_decoration_impl<__OPENCL_GLOBAL_AS__ T> {
   using type = T;
 };
 
-#ifdef __ENABLE_USM_ADDR_SPACE__
-template <typename T>
-struct remove_decoration_impl<__OPENCL_GLOBAL_DEVICE_AS__ T> {
-  using type = T;
-};
-
-template <typename T>
-struct remove_decoration_impl<__OPENCL_GLOBAL_HOST_AS__ T> {
-  using type = T;
-};
-
-#endif // __ENABLE_USM_ADDR_SPACE__
-
 template <typename T> struct remove_decoration_impl<__OPENCL_PRIVATE_AS__ T> {
   using type = T;
 };
@@ -319,8 +264,6 @@ template <typename T>
 using remove_decoration_t = typename remove_decoration<T>::type;
 
 #undef __OPENCL_GLOBAL_AS__
-#undef __OPENCL_GLOBAL_DEVICE_AS__
-#undef __OPENCL_GLOBAL_HOST_AS__
 #undef __OPENCL_LOCAL_AS__
 #undef __OPENCL_CONSTANT_AS__
 #undef __OPENCL_PRIVATE_AS__

sycl/include/sycl/access/access.hpp

@@ -1701,13 +1701,8 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
 
   template <int Dims = Dimensions>
   operator typename std::enable_if_t<
-      Dims == 0 && AccessMode == access::mode::atomic,
-#ifdef __ENABLE_USM_ADDR_SPACE__
-      atomic<DataT, access::address_space::global_space>
-#else
-      atomic<DataT, AS>
-#endif
-      >() const {
+      Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
+      const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
         getQualifiedPtr() + LinearIndex));

sycl/include/sycl/accessor.hpp

@@ -42,10 +42,8 @@ template <typename T> struct IsValidAtomicType {
 };
 
 template <sycl::access::address_space AS> struct IsValidAtomicAddressSpace {
-  static constexpr bool value =
-      (AS == access::address_space::global_space ||
-       AS == access::address_space::local_space ||
-       AS == access::address_space::ext_intel_global_device_space);
+  static constexpr bool value = (AS == access::address_space::global_space ||
+                                 AS == access::address_space::local_space);
 };
 
 // Type trait to translate a sycl::access::address_space to
@@ -54,11 +52,6 @@ template <access::address_space AS> struct GetSpirvMemoryScope {};
 template <> struct GetSpirvMemoryScope<access::address_space::global_space> {
   static constexpr auto scope = __spv::Scope::Device;
 };
-template <>
-struct GetSpirvMemoryScope<
-    access::address_space::ext_intel_global_device_space> {
-  static constexpr auto scope = __spv::Scope::Device;
-};
 template <> struct GetSpirvMemoryScope<access::address_space::local_space> {
   static constexpr auto scope = __spv::Scope::Workgroup;
 };
@@ -174,7 +167,7 @@ class __SYCL2020_DEPRECATED(
                 "long long, float");
   static_assert(detail::IsValidAtomicAddressSpace<addressSpace>::value,
                 "Invalid SYCL atomic address_space. Valid address spaces are: "
-                "global_space, local_space, ext_intel_global_device_space");
+                "global_space and local_space");
   static constexpr auto SpirvScope =
       detail::GetSpirvMemoryScope<addressSpace>::scope;
 
@@ -201,27 +194,6 @@ class __SYCL2020_DEPRECATED(
                   "T and pointerT must be same size");
   }
 
-#ifdef __ENABLE_USM_ADDR_SPACE__
-  // Create atomic in global_space with one from ext_intel_global_device_space
-  template <access::address_space _Space = addressSpace,
-            typename = typename std::enable_if_t<
-                _Space == addressSpace &&
-                addressSpace == access::address_space::global_space>>
-  atomic(const atomic<T, access::address_space::ext_intel_global_device_space>
-             &RHS) {
-    Ptr = RHS.Ptr;
-  }
-
-  template <access::address_space _Space = addressSpace,
-            typename = typename std::enable_if_t<
-                _Space == addressSpace &&
-                addressSpace == access::address_space::global_space>>
-  atomic(
-      atomic<T, access::address_space::ext_intel_global_device_space> &&RHS) {
-    Ptr = RHS.Ptr;
-  }
-#endif // __ENABLE_USM_ADDR_SPACE__
-
   void store(T Operand, memory_order Order = memory_order::relaxed) {
     __spirv_AtomicStore(Ptr, SpirvScope,
                         detail::getSPIRVMemorySemanticsMask(Order), Operand);

sycl/include/sycl/atomic.hpp

@@ -49,11 +49,9 @@ template <typename T> struct IsValidAtomicRefType {
 };
 
 template <sycl::access::address_space AS> struct IsValidAtomicRefAddressSpace {
-  static constexpr bool value =
-      (AS == access::address_space::global_space ||
-       AS == access::address_space::local_space ||
-       AS == access::address_space::ext_intel_global_device_space ||
-       AS == access::address_space::generic_space);
+  static constexpr bool value = (AS == access::address_space::global_space ||
+                                 AS == access::address_space::local_space ||
+                                 AS == access::address_space::generic_space);
 };
 
 // DefaultOrder parameter is limited to read-modify-write orders
@@ -129,8 +127,7 @@ class atomic_ref_base {
       "and pointer types");
   static_assert(detail::IsValidAtomicRefAddressSpace<AddressSpace>::value,
                 "Invalid atomic address_space.  Valid address spaces are: "
-                "global_space, local_space, ext_intel_global_device_space, "
-                "generic_space");
+                "global_space, local_space, generic_space");
   static_assert(
       detail::IsValidDefaultOrder<DefaultOrder>::value,
       "Invalid default memory_order for atomics.  Valid defaults are: "

sycl/include/sycl/atomic_ref.hpp

@@ -30,17 +30,6 @@ address_space_cast_is_possible(access::address_space Src,
   if (Src == Dst || Src == generic_space || Dst == generic_space)
     return true;
 
-  // global_host/global_device could be casted to/from global
-  auto global_space = access::address_space::global_space;
-  auto global_device = access::address_space::ext_intel_global_device_space;
-  auto global_host = access::address_space::ext_intel_global_host_space;
-
-  if (Src == global_space || Dst == global_space) {
-    auto Other = Src == global_space ? Dst : Src;
-    if (Other == global_device || Other == global_host)
-      return true;
-  }
-
   // No more compatible combinations.
   return false;
 }
@@ -70,10 +59,6 @@ auto dynamic_address_cast(ElementType *Ptr) {
   constexpr auto global_space = access::address_space::global_space;
   constexpr auto local_space = access::address_space::local_space;
   constexpr auto private_space = access::address_space::private_space;
-  constexpr auto global_device =
-      access::address_space::ext_intel_global_device_space;
-  constexpr auto global_host =
-      access::address_space::ext_intel_global_host_space;
 
   constexpr auto SrcAS = deduce_AS<ElementType *>::value;
   using dst_type = typename DecoratedType<
@@ -84,21 +69,6 @@ auto dynamic_address_cast(ElementType *Ptr) {
     return (dst_type) nullptr;
   } else if constexpr (Space == generic_space) {
     return (dst_type)Ptr;
-  } else if constexpr (Space == global_space &&
-                       (SrcAS == global_device || SrcAS == global_host)) {
-    return (dst_type)Ptr;
-  } else if constexpr (SrcAS == global_space &&
-                       (Space == global_device || Space == global_host)) {
-#if defined(__ENABLE_USM_ADDR_SPACE__)
-    static_assert(SupressNotImplementedAssert || Space != Space,
-                  "Not supported yet!");
-    return detail::static_address_cast<Space>(Ptr);
-#else
-    // If __ENABLE_USM_ADDR_SPACE__ isn't defined then both
-    // global_device/global_host are just aliases for global_space.
-    static_assert(std::is_same_v<dst_type, ElementType *>);
-    return (dst_type)Ptr;
-#endif
   } else if constexpr (Space == global_space) {
     return (dst_type)__spirv_GenericCastToPtrExplicit_ToGlobal(
         const_cast<RemoveCvT *>(Ptr), __spv::StorageClass::CrossWorkgroup);
@@ -108,12 +78,6 @@ auto dynamic_address_cast(ElementType *Ptr) {
   } else if constexpr (Space == private_space) {
     return (dst_type)__spirv_GenericCastToPtrExplicit_ToPrivate(
         const_cast<RemoveCvT *>(Ptr), __spv::StorageClass::Function);
-#if !defined(__ENABLE_USM_ADDR_SPACE__)
-  } else if constexpr (SrcAS == generic_space &&
-                       (Space == global_device || Space == global_host)) {
-    return (dst_type)__spirv_GenericCastToPtrExplicit_ToGlobal(
-        const_cast<RemoveCvT *>(Ptr), __spv::StorageClass::CrossWorkgroup);
-#endif
   } else {
     static_assert(SupressNotImplementedAssert || Space != Space,
                   "Not supported yet!");

sycl/include/sycl/detail/address_space_cast.hpp

@@ -263,17 +263,7 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
   using reference =
       sycl::ext::oneapi::experimental::annotated_ref<T, property_list_t>;
 
-#ifdef __ENABLE_USM_ADDR_SPACE__
-  using global_pointer_t = std::conditional_t<
-      detail::IsUsmKindDevice<property_list_t>::value,
-      typename sycl::ext::intel::decorated_device_ptr<T>::pointer,
-      std::conditional_t<
-          detail::IsUsmKindHost<property_list_t>::value,
-          typename sycl::ext::intel::decorated_host_ptr<T>::pointer,
-          typename decorated_global_ptr<T>::pointer>>;
-#else
   using global_pointer_t = typename decorated_global_ptr<T>::pointer;
-#endif // __ENABLE_USM_ADDR_SPACE__
 
   T *m_Ptr;
 

sycl/include/sycl/ext/intel/usm_pointers.hpp

@@ -54,27 +54,6 @@ struct PropertyMetaInfo<usm_kind_key::value_t<Kind>> {
   static constexpr const char *name = "sycl-usm-kind";
   static constexpr sycl::usm::alloc value = Kind;
 };
-
-template <typename PropertyListT, sycl::usm::alloc Kind>
-inline constexpr bool is_usm_kind = []() constexpr {
-  if constexpr (PropertyListT::template has_property<usm_kind_key>())
-    return PropertyListT::template get_property<usm_kind_key>() ==
-           usm_kind<Kind>;
-  else
-    return false;
-}();
-
-template <typename PropertyListT>
-struct IsUsmKindDevice
-    : std::bool_constant<is_usm_kind<PropertyListT, sycl::usm::alloc::device>> {
-};
-template <typename PropertyListT>
-struct IsUsmKindHost
-    : std::bool_constant<is_usm_kind<PropertyListT, sycl::usm::alloc::host>> {};
-template <typename PropertyListT>
-struct IsUsmKindShared
-    : std::bool_constant<is_usm_kind<PropertyListT, sycl::usm::alloc::shared>> {
-};
 } // namespace detail
 
 } // namespace experimental