diff --git a/sycl/include/sycl/__spirv/spirv_ops.hpp b/sycl/include/sycl/__spirv/spirv_ops.hpp
index d5e6a62d13a6e..29cb4ffa6aa2f 100644
--- a/sycl/include/sycl/__spirv/spirv_ops.hpp
+++ b/sycl/include/sycl/__spirv/spirv_ops.hpp
@@ -540,190 +540,6 @@ __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max)
 #undef __SPIRV_ATOMIC_UNSIGNED
 #undef __SPIRV_ATOMIC_XOR
 
-template <typename dataT>
-extern __attribute__((opencl_global)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToGlobal(void *Ptr) noexcept {
-  return (__attribute__((opencl_global)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToGlobal(
-          Ptr, __spv::StorageClass::CrossWorkgroup);
-}
-
-template <typename dataT>
-extern const __attribute__((opencl_global)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToGlobal(const void *Ptr) noexcept {
-  return (const __attribute__((opencl_global)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToGlobal(
-          Ptr, __spv::StorageClass::CrossWorkgroup);
-}
-
-template <typename dataT>
-extern volatile __attribute__((opencl_global)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToGlobal(volatile void *Ptr) noexcept {
-  return (volatile __attribute__((opencl_global)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToGlobal(
-          Ptr, __spv::StorageClass::CrossWorkgroup);
-}
-
-template <typename dataT>
-extern const volatile __attribute__((opencl_global)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToGlobal(const volatile void *Ptr) noexcept {
-  return (const volatile __attribute__((opencl_global)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToGlobal(
-          Ptr, __spv::StorageClass::CrossWorkgroup);
-}
-
-template <typename dataT>
-extern __attribute__((opencl_local)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToLocal(void *Ptr) noexcept {
-  return (__attribute__((opencl_local)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
-                                               __spv::StorageClass::Workgroup);
-}
-
-template <typename dataT>
-extern const __attribute__((opencl_local)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToLocal(const void *Ptr) noexcept {
-  return (const __attribute__((opencl_local)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
-                                               __spv::StorageClass::Workgroup);
-}
-
-template <typename dataT>
-extern volatile __attribute__((opencl_local)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToLocal(volatile void *Ptr) noexcept {
-  return (volatile __attribute__((opencl_local)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
-                                               __spv::StorageClass::Workgroup);
-}
-
-template <typename dataT>
-extern const volatile __attribute__((opencl_local)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToLocal(const volatile void *Ptr) noexcept {
-  return (const volatile __attribute__((opencl_local)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
-                                               __spv::StorageClass::Workgroup);
-}
-
-template <typename dataT>
-extern __attribute__((opencl_private)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToPrivate(void *Ptr) noexcept {
-  return (__attribute__((opencl_private)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
-                                                 __spv::StorageClass::Function);
-}
-
-template <typename dataT>
-extern const __attribute__((opencl_private)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToPrivate(const void *Ptr) noexcept {
-  return (const __attribute__((opencl_private)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
-                                                 __spv::StorageClass::Function);
-}
-
-template <typename dataT>
-extern volatile __attribute__((opencl_private)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToPrivate(volatile void *Ptr) noexcept {
-  return (volatile __attribute__((opencl_private)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
-                                                 __spv::StorageClass::Function);
-}
-
-template <typename dataT>
-extern const volatile __attribute__((opencl_private)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToPrivate(const volatile void *Ptr) noexcept {
-  return (const volatile __attribute__((opencl_private)) dataT *)
-      __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
-                                                 __spv::StorageClass::Function);
-}
-
-template <typename dataT>
-extern __attribute__((opencl_global)) dataT *
-__SYCL_GenericCastToPtr_ToGlobal(void *Ptr) noexcept {
-  return (__attribute__((opencl_global)) dataT *)
-      __spirv_GenericCastToPtr_ToGlobal(Ptr,
-                                        __spv::StorageClass::CrossWorkgroup);
-}
-
-template <typename dataT>
-extern const __attribute__((opencl_global)) dataT *
-__SYCL_GenericCastToPtr_ToGlobal(const void *Ptr) noexcept {
-  return (const __attribute__((opencl_global)) dataT *)
-      __spirv_GenericCastToPtr_ToGlobal(Ptr,
-                                        __spv::StorageClass::CrossWorkgroup);
-}
-
-template <typename dataT>
-extern volatile __attribute__((opencl_global)) dataT *
-__SYCL_GenericCastToPtr_ToGlobal(volatile void *Ptr) noexcept {
-  return (volatile __attribute__((opencl_global)) dataT *)
-      __spirv_GenericCastToPtr_ToGlobal(Ptr,
-                                        __spv::StorageClass::CrossWorkgroup);
-}
-
-template <typename dataT>
-extern const volatile __attribute__((opencl_global)) dataT *
-__SYCL_GenericCastToPtr_ToGlobal(const volatile void *Ptr) noexcept {
-  return (const volatile __attribute__((opencl_global)) dataT *)
-      __spirv_GenericCastToPtr_ToGlobal(Ptr,
-                                        __spv::StorageClass::CrossWorkgroup);
-}
-
-template <typename dataT>
-extern __attribute__((opencl_local)) dataT *
-__SYCL_GenericCastToPtr_ToLocal(void *Ptr) noexcept {
-  return (__attribute__((opencl_local)) dataT *)
-      __spirv_GenericCastToPtr_ToLocal(Ptr, __spv::StorageClass::Workgroup);
-}
-
-template <typename dataT>
-extern const __attribute__((opencl_local)) dataT *
-__SYCL_GenericCastToPtr_ToLocal(const void *Ptr) noexcept {
-  return (const __attribute__((opencl_local)) dataT *)
-      __spirv_GenericCastToPtr_ToLocal(Ptr, __spv::StorageClass::Workgroup);
-}
-
-template <typename dataT>
-extern volatile __attribute__((opencl_local)) dataT *
-__SYCL_GenericCastToPtr_ToLocal(volatile void *Ptr) noexcept {
-  return (volatile __attribute__((opencl_local)) dataT *)
-      __spirv_GenericCastToPtr_ToLocal(Ptr, __spv::StorageClass::Workgroup);
-}
-
-template <typename dataT>
-extern const volatile __attribute__((opencl_local)) dataT *
-__SYCL_GenericCastToPtr_ToLocal(const volatile void *Ptr) noexcept {
-  return (const volatile __attribute__((opencl_local)) dataT *)
-      __spirv_GenericCastToPtr_ToLocal(Ptr, __spv::StorageClass::Workgroup);
-}
-
-template <typename dataT>
-extern __attribute__((opencl_private)) dataT *
-__SYCL_GenericCastToPtr_ToPrivate(void *Ptr) noexcept {
-  return (__attribute__((opencl_private)) dataT *)
-      __spirv_GenericCastToPtr_ToPrivate(Ptr, __spv::StorageClass::Function);
-}
-
-template <typename dataT>
-extern const __attribute__((opencl_private)) dataT *
-__SYCL_GenericCastToPtr_ToPrivate(const void *Ptr) noexcept {
-  return (const __attribute__((opencl_private)) dataT *)
-      __spirv_GenericCastToPtr_ToPrivate(Ptr, __spv::StorageClass::Function);
-}
-
-template <typename dataT>
-extern volatile __attribute__((opencl_private)) dataT *
-__SYCL_GenericCastToPtr_ToPrivate(volatile void *Ptr) noexcept {
-  return (volatile __attribute__((opencl_private)) dataT *)
-      __spirv_GenericCastToPtr_ToPrivate(Ptr, __spv::StorageClass::Function);
-}
-
-template <typename dataT>
-extern const volatile __attribute__((opencl_private)) dataT *
-__SYCL_GenericCastToPtr_ToPrivate(const volatile void *Ptr) noexcept {
-  return (const volatile __attribute__((opencl_private)) dataT *)
-      __spirv_GenericCastToPtr_ToPrivate(Ptr, __spv::StorageClass::Function);
-}
-
 template <typename dataT>
 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
 __spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;
diff --git a/sycl/include/sycl/access/access.hpp b/sycl/include/sycl/access/access.hpp
index b35bf36ec7fe1..ad95479e9d885 100644
--- a/sycl/include/sycl/access/access.hpp
+++ b/sycl/include/sycl/access/access.hpp
@@ -325,58 +325,154 @@ template <typename T>
 using remove_decoration_t = typename remove_decoration<T>::type;
 
 namespace detail {
-
-// Helper function for selecting appropriate casts between address spaces.
-template <typename ToT, typename FromT> inline ToT cast_AS(FromT from) {
 #ifdef __SYCL_DEVICE_ONLY__
-  constexpr access::address_space ToAS = deduce_AS<ToT>::value;
-  constexpr access::address_space FromAS = deduce_AS<FromT>::value;
-  if constexpr (FromAS == access::address_space::generic_space) {
-#if defined(__NVPTX__) || defined(__AMDGCN__) || defined(__SYCL_NATIVE_CPU__)
-    // TODO: NVPTX and AMDGCN backends do not currently support the
-    //       __spirv_GenericCastToPtrExplicit_* builtins, so to work around this
-    //       we do C-style casting. This may produce warnings when targetting
-    //       these backends.
-    return (ToT)from;
+inline constexpr bool
+address_space_cast_is_possible(access::address_space Src,
+                               access::address_space Dst) {
+  // constant_space is unique and is not interchangeable with any other.
+  auto constant_space = access::address_space::constant_space;
+  if (Src == constant_space || Dst == constant_space)
+    return Src == Dst;
+
+  auto generic_space = access::address_space::generic_space;
+  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;
+}
+
+template <access::address_space Space, typename ElementType>
+auto static_address_cast(ElementType *Ptr) {
+  constexpr auto generic_space = access::address_space::generic_space;
+  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;
+  static_assert(address_space_cast_is_possible(SrcAS, Space));
+
+  using dst_type = typename DecoratedType<
+      std::remove_pointer_t<remove_decoration_t<ElementType *>>, Space>::type *;
+
+  // Note: reinterpret_cast isn't enough for some of the casts between different
+  // address spaces, use C-style cast instead.
+#if !defined(__SPIR__)
+  return (dst_type)Ptr;
 #else
-    using ToElemT = std::remove_pointer_t<remove_decoration_t<ToT>>;
-    if constexpr (ToAS == access::address_space::global_space)
-      return __SYCL_GenericCastToPtrExplicit_ToGlobal<ToElemT>(from);
-    else if constexpr (ToAS == access::address_space::local_space)
-      return __SYCL_GenericCastToPtrExplicit_ToLocal<ToElemT>(from);
-    else if constexpr (ToAS == access::address_space::private_space)
-      return __SYCL_GenericCastToPtrExplicit_ToPrivate<ToElemT>(from);
-#ifdef __ENABLE_USM_ADDR_SPACE__
-    else if constexpr (ToAS == access::address_space::
-                                   ext_intel_global_device_space ||
-                       ToAS ==
-                           access::address_space::ext_intel_global_host_space)
-      // For extended address spaces we do not currently have a SPIR-V
-      // conversion function, so we do a C-style cast. This may produce
-      // warnings.
-      return (ToT)from;
-#endif // __ENABLE_USM_ADDR_SPACE__
-    else
-      return reinterpret_cast<ToT>(from);
-#endif // defined(__NVPTX__) || defined(__AMDGCN__)
-  } else
-#ifdef __ENABLE_USM_ADDR_SPACE__
-      if constexpr (FromAS == access::address_space::global_space &&
-                    (ToAS ==
-                         access::address_space::ext_intel_global_device_space ||
-                     ToAS ==
-                         access::address_space::ext_intel_global_host_space)) {
-    // Casting from global address space to the global device and host address
-    // spaces is allowed.
-    return (ToT)from;
-  } else
-#endif // __ENABLE_USM_ADDR_SPACE__
-#endif // __SYCL_DEVICE_ONLY__
-  {
-    return reinterpret_cast<ToT>(from);
+  if constexpr (SrcAS != generic_space) {
+    return (dst_type)Ptr;
+  } else if constexpr (Space == global_space) {
+    return (dst_type)__spirv_GenericCastToPtr_ToGlobal(
+        Ptr, __spv::StorageClass::CrossWorkgroup);
+  } else if constexpr (Space == local_space) {
+    return (dst_type)__spirv_GenericCastToPtr_ToLocal(
+        Ptr, __spv::StorageClass::Workgroup);
+  } else if constexpr (Space == private_space) {
+    return (dst_type)__spirv_GenericCastToPtr_ToPrivate(
+        Ptr, __spv::StorageClass::Function);
+#if !defined(__ENABLE_USM_ADDR_SPACE__)
+  } else if constexpr (Space == global_device || Space == global_host) {
+    // If __ENABLE_USM_ADDR_SPACE__ isn't defined then both
+    // global_device/global_host are just aliases for global_space.
+    return (dst_type)__spirv_GenericCastToPtr_ToGlobal(
+        Ptr, __spv::StorageClass::CrossWorkgroup);
+#endif
+  } else {
+    return (dst_type)Ptr;
   }
+#endif
 }
 
+// Previous implementation (`castAS`, used in `multi_ptr` ctors among other
+// places), used C-style cast instead of a proper dynamic check for some
+// backends/spaces. `SupressNotImplementedAssert = true` parameter is emulating
+// that previous behavior until the proper support is added for compatibility
+// reasons.
+template <access::address_space Space, bool SupressNotImplementedAssert = false,
+          typename ElementType>
+auto dynamic_address_cast(ElementType *Ptr) {
+  constexpr auto generic_space = access::address_space::generic_space;
+  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<
+      std::remove_pointer_t<remove_decoration_t<ElementType *>>, Space>::type *;
+
+  if constexpr (!address_space_cast_is_possible(SrcAS, Space)) {
+    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 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
+#if defined(__SPIR__)
+  } else if constexpr (Space == global_space) {
+    return (dst_type)__spirv_GenericCastToPtrExplicit_ToGlobal(
+        Ptr, __spv::StorageClass::CrossWorkgroup);
+  } else if constexpr (Space == local_space) {
+    return (dst_type)__spirv_GenericCastToPtrExplicit_ToLocal(
+        Ptr, __spv::StorageClass::Workgroup);
+  } else if constexpr (Space == private_space) {
+    return (dst_type)__spirv_GenericCastToPtrExplicit_ToPrivate(
+        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(
+        Ptr, __spv::StorageClass::CrossWorkgroup);
+#endif
+#endif
+  } else {
+    static_assert(SupressNotImplementedAssert || Space != Space,
+                  "Not supported yet!");
+    return static_address_cast<Space>(Ptr);
+  }
+}
+#else  // __SYCL_DEVICE_ONLY__
+template <access::address_space Space, typename ElementType>
+auto static_address_cast(ElementType *Ptr) {
+  return Ptr;
+}
+template <access::address_space Space, bool SupressNotImplementedAssert = false,
+          typename ElementType>
+auto dynamic_address_cast(ElementType *Ptr) {
+  return Ptr;
+}
+#endif // __SYCL_DEVICE_ONLY__
 } // namespace detail
 
 #undef __OPENCL_GLOBAL_AS__
diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp
index d3113bc1c113e..98ca92b6c988d 100644
--- a/sycl/include/sycl/detail/spirv.hpp
+++ b/sycl/include/sycl/detail/spirv.hpp
@@ -1377,30 +1377,6 @@ __SYCL_GROUP_COLLECTIVE_OVERLOAD(BitwiseAnd, KHR)
 __SYCL_GROUP_COLLECTIVE_OVERLOAD(LogicalAnd, KHR)
 __SYCL_GROUP_COLLECTIVE_OVERLOAD(LogicalOr, KHR)
 
-template <access::address_space Space, typename T>
-auto GenericCastToPtr(T *Ptr) ->
-    typename multi_ptr<T, Space, access::decorated::yes>::pointer {
-  if constexpr (Space == access::address_space::global_space) {
-    return __SYCL_GenericCastToPtr_ToGlobal<T>(Ptr);
-  } else if constexpr (Space == access::address_space::local_space) {
-    return __SYCL_GenericCastToPtr_ToLocal<T>(Ptr);
-  } else if constexpr (Space == access::address_space::private_space) {
-    return __SYCL_GenericCastToPtr_ToPrivate<T>(Ptr);
-  }
-}
-
-template <access::address_space Space, typename T>
-auto GenericCastToPtrExplicit(T *Ptr) ->
-    typename multi_ptr<T, Space, access::decorated::yes>::pointer {
-  if constexpr (Space == access::address_space::global_space) {
-    return __SYCL_GenericCastToPtrExplicit_ToGlobal<T>(Ptr);
-  } else if constexpr (Space == access::address_space::local_space) {
-    return __SYCL_GenericCastToPtrExplicit_ToLocal<T>(Ptr);
-  } else if constexpr (Space == access::address_space::private_space) {
-    return __SYCL_GenericCastToPtrExplicit_ToPrivate<T>(Ptr);
-  }
-}
-
 } // namespace spirv
 } // namespace detail
 } // namespace _V1
diff --git a/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp b/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp
index 8e901d7d2877a..4708cd2224e19 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp
@@ -13,47 +13,25 @@
 
 namespace sycl {
 inline namespace _V1 {
-namespace ext {
-namespace oneapi {
-namespace experimental {
+namespace ext::oneapi ::experimental {
 namespace detail {
 using namespace sycl::detail;
 }
 // Shorthands for address space names
-constexpr inline access::address_space global_space = access::address_space::global_space;
-constexpr inline access::address_space local_space = access::address_space::local_space;
-constexpr inline access::address_space private_space = access::address_space::private_space;
-constexpr inline access::address_space generic_space = access::address_space::generic_space;
+constexpr inline access::address_space global_space =
+    access::address_space::global_space;
+constexpr inline access::address_space local_space =
+    access::address_space::local_space;
+constexpr inline access::address_space private_space =
+    access::address_space::private_space;
+constexpr inline access::address_space generic_space =
+    access::address_space::generic_space;
 
 template <access::address_space Space, typename ElementType>
 multi_ptr<ElementType, Space, access::decorated::no>
 static_address_cast(ElementType *Ptr) {
   using ret_ty = multi_ptr<ElementType, Space, access::decorated::no>;
-#ifdef __SYCL_DEVICE_ONLY__
-  static_assert(std::is_same_v<ElementType, remove_decoration_t<ElementType>>,
-                "The extension expects undecorated raw pointers only!");
-  if constexpr (Space == generic_space) {
-    // Undecorated raw pointer is in generic AS already, no extra casts needed.
-    return ret_ty(Ptr);
-  } else if constexpr (Space == access::address_space::
-                                    ext_intel_global_device_space ||
-                       Space ==
-                           access::address_space::ext_intel_global_host_space) {
-#ifdef __ENABLE_USM_ADDR_SPACE__
-    // No SPIR-V intrinsic for this yet.
-    using raw_type = detail::DecoratedType<ElementType, Space>::type *;
-    auto CastPtr = (raw_type)(Ptr);
-#else
-    auto CastPtr = sycl::detail::spirv::GenericCastToPtr<global_space>(Ptr);
-#endif
-    return ret_ty(CastPtr);
-  } else {
-    auto CastPtr = sycl::detail::spirv::GenericCastToPtr<Space>(Ptr);
-    return ret_ty(CastPtr);
-  }
-#else
-  return ret_ty(Ptr);
-#endif
+  return ret_ty{detail::static_address_cast<Space>(Ptr)};
 }
 
 template <access::address_space Space, access::decorated DecorateAddress,
@@ -63,39 +41,14 @@ multi_ptr<ElementType, Space, DecorateAddress> static_address_cast(
   if constexpr (Space == generic_space)
     return Ptr;
   else
-    return {static_address_cast<Space>(Ptr.get_raw())};
+    return {static_address_cast<Space>(Ptr.get_decorated())};
 }
 
 template <access::address_space Space, typename ElementType>
 multi_ptr<ElementType, Space, access::decorated::no>
 dynamic_address_cast(ElementType *Ptr) {
   using ret_ty = multi_ptr<ElementType, Space, access::decorated::no>;
-#ifdef __SYCL_DEVICE_ONLY__
-  static_assert(std::is_same_v<ElementType, remove_decoration_t<ElementType>>,
-                "The extension expects undecorated raw pointers only!");
-  if constexpr (Space == generic_space) {
-    return ret_ty(Ptr);
-  } else if constexpr (Space == access::address_space::
-                                    ext_intel_global_device_space ||
-                       Space ==
-                           access::address_space::ext_intel_global_host_space) {
-#ifdef __ENABLE_USM_ADDR_SPACE__
-    static_assert(
-        Space != access::address_space::ext_intel_global_device_space &&
-            Space != access::address_space::ext_intel_global_host_space,
-        "Not supported yet!");
-    return ret_ty(nullptr);
-#else
-    auto CastPtr = sycl::detail::spirv::GenericCastToPtr<global_space>(Ptr);
-    return ret_ty(CastPtr);
-#endif
-  } else {
-    auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit<Space>(Ptr);
-    return ret_ty(CastPtr);
-  }
-#else
-  return ret_ty(Ptr);
-#endif
+  return ret_ty{detail::dynamic_address_cast<Space>(Ptr)};
 }
 
 template <access::address_space Space, access::decorated DecorateAddress,
@@ -105,11 +58,9 @@ multi_ptr<ElementType, Space, DecorateAddress> dynamic_address_cast(
   if constexpr (Space == generic_space)
     return Ptr;
   else
-    return {dynamic_address_cast<Space>(Ptr.get_raw())};
+    return {dynamic_address_cast<Space>(Ptr.get_decorated())};
 }
 
-} // namespace experimental
-} // namespace oneapi
-} // namespace ext
+} // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp
index d3945fefd8eba..e98e8848f0e59 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp
@@ -217,11 +217,10 @@ auto get_block_op_ptr(IteratorT iter, [[maybe_unused]] Properties props) {
     if constexpr (AS == access::address_space::global_space) {
       return is_aligned ? reinterpret_cast<block_pointer_type>(iter) : nullptr;
     } else if constexpr (AS == access::address_space::generic_space) {
-      return is_aligned
-                 ? reinterpret_cast<block_pointer_type>(
-                       __SYCL_GenericCastToPtrExplicit_ToGlobal<value_type>(
-                           iter))
-                 : nullptr;
+      return is_aligned ? reinterpret_cast<block_pointer_type>(
+                              detail::dynamic_address_cast<
+                                  access::address_space::global_space>(iter))
+                        : nullptr;
     } else {
       return nullptr;
     }
diff --git a/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp b/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
index 127d62450cb4c..c5be683830183 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
@@ -49,6 +49,8 @@ inline constexpr prefetch_hint_key::value_t<cache_level::L4, nontemporal>
     prefetch_hint_L4_nt;
 
 namespace detail {
+using namespace sycl::detail;
+
 template <> struct IsCompileTimeProperty<prefetch_hint_key> : std::true_type {};
 
 template <cache_level Level, typename Hint>
@@ -71,7 +73,10 @@ inline constexpr bool check_prefetch_acc_mode =
 template <typename T, typename Properties>
 void prefetch_impl(T *ptr, size_t bytes, Properties properties) {
 #ifdef __SYCL_DEVICE_ONLY__
-  auto *ptrGlobalAS = __SYCL_GenericCastToPtrExplicit_ToGlobal<const char>(ptr);
+  auto *ptrGlobalAS =
+      reinterpret_cast<__attribute__((opencl_global)) const char *>(
+          detail::static_address_cast<access::address_space::global_space>(
+              const_cast<const T *>(ptr)));
   const __attribute__((opencl_global)) char *ptrAnnotated = nullptr;
   if constexpr (!properties.template has_property<prefetch_hint_key>()) {
     ptrAnnotated = __builtin_intel_sycl_ptr_annotation(
diff --git a/sycl/include/sycl/group.hpp b/sycl/include/sycl/group.hpp
index e5b16cce825d6..50b5fe6e586ad 100644
--- a/sycl/include/sycl/group.hpp
+++ b/sycl/include/sycl/group.hpp
@@ -449,12 +449,12 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
     using QualSrcT =
         std::conditional_t<std::is_const_v<SrcT>, const uint8_t, uint8_t>;
     auto DestP = multi_ptr<uint8_t, DestS, access::decorated::yes>(
-        detail::cast_AS<typename multi_ptr<uint8_t, DestS,
-                                           access::decorated::yes>::pointer>(
+        reinterpret_cast<typename multi_ptr<uint8_t, DestS,
+                                            access::decorated::yes>::pointer>(
             Dest.get_decorated()));
     auto SrcP = multi_ptr<QualSrcT, SrcS, access::decorated::yes>(
-        detail::cast_AS<typename multi_ptr<QualSrcT, SrcS,
-                                           access::decorated::yes>::pointer>(
+        reinterpret_cast<typename multi_ptr<QualSrcT, SrcS,
+                                            access::decorated::yes>::pointer>(
             Src.get_decorated()));
     return async_work_group_copy(DestP, SrcP, NumElements, Stride);
   }
@@ -478,12 +478,12 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
     using QualSrcVecT =
         std::conditional_t<std::is_const_v<SrcT>, std::add_const_t<VecT>, VecT>;
     auto DestP = multi_ptr<VecT, DestS, access::decorated::yes>(
-        detail::cast_AS<
+        reinterpret_cast<
             typename multi_ptr<VecT, DestS, access::decorated::yes>::pointer>(
             Dest.get_decorated()));
     auto SrcP = multi_ptr<QualSrcVecT, SrcS, access::decorated::yes>(
-        detail::cast_AS<typename multi_ptr<QualSrcVecT, SrcS,
-                                           access::decorated::yes>::pointer>(
+        reinterpret_cast<typename multi_ptr<QualSrcVecT, SrcS,
+                                            access::decorated::yes>::pointer>(
             Src.get_decorated()));
     return async_work_group_copy(DestP, SrcP, NumElements, Stride);
   }
diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp
index e7d9bbcfbe7bd..9bca96ade3879 100644
--- a/sycl/include/sycl/multi_ptr.hpp
+++ b/sycl/include/sycl/multi_ptr.hpp
@@ -127,7 +127,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
   multi_ptr(accessor<ElementType, Dimensions, Mode, target::device,
                      isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(detail::cast_AS<decorated_type *>(
+      : multi_ptr(detail::static_address_cast<Space>(
             Accessor.template get_multi_ptr<DecorateAddress>()
                 .get_decorated())) {}
 
@@ -198,7 +198,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
   multi_ptr(accessor<typename std::remove_const_t<RelayElementType>, Dimensions,
                      Mode, target::device, isPlaceholder, PropertyListT>
                 Accessor)
-      : m_Pointer(detail::cast_AS<decorated_type *>(
+      : m_Pointer(detail::static_address_cast<Space>(
             Accessor.template get_multi_ptr<DecorateAddress>()
                 .get_decorated())) {}
 
@@ -272,7 +272,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
                            OtherSpace != access::address_space::constant_space>>
   multi_ptr &
   operator=(const multi_ptr<value_type, OtherSpace, OtherIsDecorated> &Other) {
-    m_Pointer = detail::cast_AS<decorated_type *>(Other.get_decorated());
+    m_Pointer = detail::static_address_cast<Space>(Other.get_decorated());
     return *this;
   }
   template <
@@ -282,7 +282,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
                            OtherSpace != access::address_space::constant_space>>
   multi_ptr &
   operator=(multi_ptr<value_type, OtherSpace, OtherIsDecorated> &&Other) {
-    m_Pointer = detail::cast_AS<decorated_type *>(std::move(Other.m_Pointer));
+    m_Pointer = detail::static_address_cast<Space>(std::move(Other.m_Pointer));
     return *this;
   }
 
@@ -290,7 +290,10 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
   pointer operator->() const { return get(); }
   reference operator[](difference_type index) const { return m_Pointer[index]; }
 
-  pointer get() const { return detail::cast_AS<pointer>(m_Pointer); }
+  pointer get() const {
+    return detail::static_address_cast<
+        is_decorated ? Space : access::address_space::generic_space>(m_Pointer);
+  }
   decorated_type *get_decorated() const { return m_Pointer; }
   std::add_pointer_t<value_type> get_raw() const {
     return reinterpret_cast<std::add_pointer_t<value_type>>(get());
@@ -312,9 +315,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
   explicit
   operator multi_ptr<value_type, OtherSpace, OtherIsDecorated>() const {
     return multi_ptr<value_type, OtherSpace, OtherIsDecorated>{
-        detail::cast_AS<typename multi_ptr<value_type, OtherSpace,
-                                           access::decorated::yes>::pointer>(
-            get_decorated())};
+        detail::static_address_cast<OtherSpace>(get_decorated())};
   }
 
   template <access::address_space OtherSpace,
@@ -331,9 +332,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
   explicit
   operator multi_ptr<const value_type, OtherSpace, OtherIsDecorated>() const {
     return multi_ptr<const value_type, OtherSpace, OtherIsDecorated>{
-        detail::cast_AS<typename multi_ptr<const value_type, OtherSpace,
-                                           access::decorated::yes>::pointer>(
-            get_decorated())};
+        detail::static_address_cast<OtherSpace>(get_decorated())};
   }
 
   template <access::decorated ConvIsDecorated,
@@ -342,7 +341,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
                 std::is_same_v<RelayElementType, ElementType> &&
                 !std::is_const_v<RelayElementType>>>
   operator multi_ptr<void, Space, ConvIsDecorated>() const {
-    return multi_ptr<void, Space, ConvIsDecorated>{detail::cast_AS<
+    return multi_ptr<void, Space, ConvIsDecorated>{static_cast<
         typename multi_ptr<void, Space, access::decorated::yes>::pointer>(
         get_decorated())};
   }
@@ -353,17 +352,14 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
                 std::is_same_v<RelayElementType, ElementType> &&
                 std::is_const_v<RelayElementType>>>
   operator multi_ptr<const void, Space, ConvIsDecorated>() const {
-    return multi_ptr<const void, Space, ConvIsDecorated>{detail::cast_AS<
+    return multi_ptr<const void, Space, ConvIsDecorated>{static_cast<
         typename multi_ptr<const void, Space, access::decorated::yes>::pointer>(
         get_decorated())};
   }
 
   template <access::decorated ConvIsDecorated>
   operator multi_ptr<const value_type, Space, ConvIsDecorated>() const {
-    return multi_ptr<const value_type, Space, ConvIsDecorated>{
-        detail::cast_AS<typename multi_ptr<const value_type, Space,
-                                           access::decorated::yes>::pointer>(
-            get_decorated())};
+    return multi_ptr<const value_type, Space, ConvIsDecorated>{get_decorated()};
   }
 
   operator multi_ptr<value_type, Space,
@@ -386,11 +382,8 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
            Space == access::address_space::ext_intel_global_host_space)>>
   explicit
   operator multi_ptr<ElementType, GlobalSpace, DecorateAddress>() const {
-    using global_pointer_t =
-        typename multi_ptr<ElementType, GlobalSpace,
-                           access::decorated::yes>::pointer;
     return multi_ptr<ElementType, GlobalSpace, DecorateAddress>(
-        detail::cast_AS<global_pointer_t>(get_decorated()));
+        detail::static_address_cast<GlobalSpace>(get_decorated()));
   }
 
   // Only if Space == global_space
@@ -487,7 +480,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<const void, Space, DecorateAddress> {
   multi_ptr(accessor<ElementType, Dimensions, Mode, target::device,
                      isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(detail::cast_AS<decorated_type *>(
+      : multi_ptr(detail::static_address_cast<Space>(
             Accessor.template get_multi_ptr<DecorateAddress>()
                 .get_decorated())) {}
 
@@ -541,7 +534,10 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<const void, Space, DecorateAddress> {
     return *this;
   }
 
-  pointer get() const { return detail::cast_AS<pointer>(m_Pointer); }
+  pointer get() const {
+    return detail::static_address_cast<
+        is_decorated ? Space : access::address_space::generic_space>(m_Pointer);
+  }
 
   // Conversion to the underlying pointer type
   __SYCL2020_DEPRECATED("Conversion to pointer type is deprecated since SYCL "
@@ -553,8 +549,8 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<const void, Space, DecorateAddress> {
             typename = typename std::enable_if_t<std::is_const_v<ElementType>>>
   explicit operator multi_ptr<ElementType, Space, DecorateAddress>() const {
     return multi_ptr<ElementType, Space, DecorateAddress>{
-        detail::cast_AS<typename multi_ptr<ElementType, Space,
-                                           access::decorated::yes>::pointer>(
+        static_cast<typename multi_ptr<ElementType, Space,
+                                       access::decorated::yes>::pointer>(
             m_Pointer)};
   }
 
@@ -579,11 +575,8 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<const void, Space, DecorateAddress> {
            Space == access::address_space::ext_intel_global_host_space)>>
   explicit
   operator multi_ptr<const void, GlobalSpace, DecorateAddress>() const {
-    using global_pointer_t =
-        typename multi_ptr<const void, GlobalSpace,
-                           access::decorated::yes>::pointer;
-    return multi_ptr<const void, GlobalSpace, DecorateAddress>(
-        detail::cast_AS<global_pointer_t>(m_Pointer));
+    return multi_ptr<const void, GlobalSpace, DecorateAddress>{
+        detail::static_address_cast<GlobalSpace>(m_Pointer)};
   }
 
 private:
@@ -638,7 +631,7 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<void, Space, DecorateAddress> {
   multi_ptr(accessor<ElementType, Dimensions, Mode, target::device,
                      isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(detail::cast_AS<decorated_type *>(
+      : multi_ptr(detail::static_address_cast<Space>(
             Accessor.template get_multi_ptr<DecorateAddress>()
                 .get_decorated())) {}
 
@@ -692,7 +685,10 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<void, Space, DecorateAddress> {
     return *this;
   }
 
-  pointer get() const { return detail::cast_AS<pointer>(m_Pointer); }
+  pointer get() const {
+    return detail::static_address_cast<
+        is_decorated ? Space : access::address_space::generic_space>(m_Pointer);
+  }
 
   // Conversion to the underlying pointer type
   __SYCL2020_DEPRECATED("Conversion to pointer type is deprecated since SYCL "
@@ -703,8 +699,8 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<void, Space, DecorateAddress> {
   template <typename ElementType>
   explicit operator multi_ptr<ElementType, Space, DecorateAddress>() const {
     return multi_ptr<ElementType, Space, DecorateAddress>{
-        detail::cast_AS<typename multi_ptr<ElementType, Space,
-                                           access::decorated::yes>::pointer>(
+        static_cast<typename multi_ptr<ElementType, Space,
+                                       access::decorated::yes>::pointer>(
             m_Pointer)};
   }
 
@@ -728,10 +724,8 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<void, Space, DecorateAddress> {
           (Space == access::address_space::ext_intel_global_device_space ||
            Space == access::address_space::ext_intel_global_host_space)>>
   explicit operator multi_ptr<void, GlobalSpace, DecorateAddress>() const {
-    using global_pointer_t =
-        typename multi_ptr<void, GlobalSpace, access::decorated::yes>::pointer;
     return multi_ptr<void, GlobalSpace, DecorateAddress>(
-        detail::cast_AS<global_pointer_t>(m_Pointer));
+        detail::static_address_cast<GlobalSpace>(m_Pointer));
   }
 
 private:
@@ -782,14 +776,17 @@ class __SYCL2020_DEPRECATED(
 #endif
 
   multi_ptr(ElementType *pointer)
-      : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {
+      : m_Pointer(detail::dynamic_address_cast<
+                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
   }
 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
+  // TODO: This isn't SFINAE, is it? How does this work?
   template <typename = typename detail::const_if_const_AS<Space, ElementType>>
   multi_ptr(const ElementType *pointer)
-      : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {}
+      : m_Pointer(detail::dynamic_address_cast<
+                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {}
 #endif
 
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
@@ -816,7 +813,8 @@ class __SYCL2020_DEPRECATED(
   multi_ptr &operator=(ElementType *pointer) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
-    m_Pointer = detail::cast_AS<pointer_t>(pointer);
+    m_Pointer = detail::dynamic_address_cast<
+        Space, /* SupressNotImplementedAssert = */ true>(pointer);
     return *this;
   }
 
@@ -858,7 +856,8 @@ class __SYCL2020_DEPRECATED(
   multi_ptr(accessor<ElementType, dimensions, Mode, target::device,
                      isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(detail::cast_AS<pointer_t>(Accessor.get_pointer().get())) {}
+      : multi_ptr(
+            detail::static_address_cast<Space>(Accessor.get_pointer().get())) {}
 
   // Only if Space == local_space || generic_space
   template <
@@ -1112,14 +1111,19 @@ class __SYCL2020_DEPRECATED(
       typename RelayPointerT = pointer_t,
       typename = std::enable_if_t<std::is_same_v<RelayPointerT, pointer_t> &&
                                   !std::is_same_v<RelayPointerT, void *>>>
-  multi_ptr(void *pointer) : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {
+  multi_ptr(void *pointer)
+      : m_Pointer(detail::dynamic_address_cast<
+                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
   }
 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
   template <typename = typename detail::const_if_const_AS<Space, void>>
   multi_ptr(const void *pointer)
-      : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {}
+      : m_Pointer(
+            detail::dynamic_address_cast<
+                pointer_t, /* SupressNotImplementedAssert = */ true>(pointer)) {
+  }
 #endif
 #endif
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
@@ -1149,7 +1153,8 @@ class __SYCL2020_DEPRECATED(
   multi_ptr &operator=(void *pointer) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
-    m_Pointer = detail::cast_AS<pointer_t>(pointer);
+    m_Pointer = detail::dynamic_address_cast<
+        Space, /* SupressNotImplementedAssert = */ true>(pointer);
     return *this;
   }
 #endif
@@ -1269,14 +1274,16 @@ class __SYCL2020_DEPRECATED(
       typename = std::enable_if_t<std::is_same_v<RelayPointerT, pointer_t> &&
                                   !std::is_same_v<RelayPointerT, const void *>>>
   multi_ptr(const void *pointer)
-      : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {
+      : m_Pointer(detail::dynamic_address_cast<
+                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
   }
 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
   template <typename = typename detail::const_if_const_AS<Space, void>>
   multi_ptr(const void *pointer)
-      : m_Pointer(detail::cast_AS<pointer_t>(pointer)) {}
+      : m_Pointer(detail::dynamic_address_cast<
+                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {}
 #endif
 #endif
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
@@ -1306,7 +1313,8 @@ class __SYCL2020_DEPRECATED(
   multi_ptr &operator=(const void *pointer) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
-    m_Pointer = detail::cast_AS<pointer_t>(pointer);
+    m_Pointer = detail::dynamic_address_cast<
+        pointer_t, /* SupressNotImplementedAssert = */ true>(pointer);
     return *this;
   }
 #endif
@@ -1434,8 +1442,8 @@ address_space_cast(ElementType *pointer) {
   // space is not compatible with Space.
   // Use LegacyPointerTypes here to also allow constant_space
   return multi_ptr<ElementType, Space, DecorateAddress>(
-      detail::cast_AS<
-          typename detail::LegacyPointerTypes<ElementType, Space>::pointer_t>(
+      detail::dynamic_address_cast<Space,
+                                   /* SupressNotImplementedAssert = */ true>(
           pointer));
 }
 
diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp
index 052596b46c424..3abd919a2a09a 100644
--- a/sycl/include/sycl/nd_item.hpp
+++ b/sycl/include/sycl/nd_item.hpp
@@ -384,12 +384,12 @@ template <int Dimensions = 1> class nd_item {
     using QualSrcT =
         std::conditional_t<std::is_const_v<SrcT>, const uint8_t, uint8_t>;
     auto DestP = multi_ptr<uint8_t, DestS, access::decorated::yes>(
-        detail::cast_AS<typename multi_ptr<uint8_t, DestS,
-                                           access::decorated::yes>::pointer>(
+        reinterpret_cast<typename multi_ptr<uint8_t, DestS,
+                                            access::decorated::yes>::pointer>(
             Dest.get_decorated()));
     auto SrcP = multi_ptr<QualSrcT, SrcS, access::decorated::yes>(
-        detail::cast_AS<typename multi_ptr<QualSrcT, SrcS,
-                                           access::decorated::yes>::pointer>(
+        reinterpret_cast<typename multi_ptr<QualSrcT, SrcS,
+                                            access::decorated::yes>::pointer>(
             Src.get_decorated()));
     return async_work_group_copy(DestP, SrcP, NumElements, Stride);
   }
@@ -413,12 +413,12 @@ template <int Dimensions = 1> class nd_item {
     using QualSrcVecT =
         std::conditional_t<std::is_const_v<SrcT>, std::add_const_t<VecT>, VecT>;
     auto DestP = multi_ptr<VecT, DestS, access::decorated::yes>(
-        detail::cast_AS<
+        reinterpret_cast<
             typename multi_ptr<VecT, DestS, access::decorated::yes>::pointer>(
             Dest.get_decorated()));
     auto SrcP = multi_ptr<QualSrcVecT, SrcS, access::decorated::yes>(
-        detail::cast_AS<typename multi_ptr<QualSrcVecT, SrcS,
-                                           access::decorated::yes>::pointer>(
+        reinterpret_cast<typename multi_ptr<QualSrcVecT, SrcS,
+                                            access::decorated::yes>::pointer>(
             Src.get_decorated()));
     return async_work_group_copy(DestP, SrcP, NumElements, Stride);
   }
diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp
index d9e15d021063f..9f0cf4aa007f7 100644
--- a/sycl/include/sycl/sub_group.hpp
+++ b/sycl/include/sycl/sub_group.hpp
@@ -228,12 +228,14 @@ struct sub_group {
 #if defined(__NVPTX__) || defined(__AMDGCN__)
     return src[get_local_id()[0]];
 #else  // __NVPTX__ || __AMDGCN__
-    auto l = __SYCL_GenericCastToPtrExplicit_ToLocal<T>(src);
-    if (l)
+    if (auto l =
+            detail::dynamic_address_cast<access::address_space::local_space>(
+                src))
       return load(l);
 
-    auto g = __SYCL_GenericCastToPtrExplicit_ToGlobal<T>(src);
-    if (g)
+    if (auto g =
+            detail::dynamic_address_cast<access::address_space::global_space>(
+                src))
       return load(g);
 
     // Sub-group load() is supported for local or global pointers only.
@@ -418,14 +420,16 @@ struct sub_group {
 #if defined(__NVPTX__) || defined(__AMDGCN__)
     dst[get_local_id()[0]] = x;
 #else  // __NVPTX__ || __AMDGCN__
-    auto l = __SYCL_GenericCastToPtrExplicit_ToLocal<T>(dst);
-    if (l) {
+    if (auto l =
+            detail::dynamic_address_cast<access::address_space::local_space>(
+                dst)) {
       store(l, x);
       return;
     }
 
-    auto g = __SYCL_GenericCastToPtrExplicit_ToGlobal<T>(dst);
-    if (g) {
+    if (auto g =
+            detail::dynamic_address_cast<access::address_space::global_space>(
+                dst)) {
       store(g, x);
       return;
     }
diff --git a/sycl/test/check_device_code/extensions/sub_group_as.cpp b/sycl/test/check_device_code/extensions/sub_group_as.cpp
index 2e345b1fb3f66..022b64ad50554 100644
--- a/sycl/test/check_device_code/extensions/sub_group_as.cpp
+++ b/sycl/test/check_device_code/extensions/sub_group_as.cpp
@@ -1,7 +1,4 @@
-// RUN: %clangxx -fsycl-device-only -O3 -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s --check-prefix CHECK-O3
-// RUN: %clangxx -fsycl-device-only -O0 -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s --check-prefix CHECK-O0
-// Test compilation with -O3 when all methods are inlined in kernel function
-// and -O0 when helper methods are preserved.
+// RUN: %clangxx -fsycl-device-only -O3 -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s
 #include <cstdint>
 #include <cstdio>
 #include <cstdlib>
@@ -38,49 +35,30 @@ SYCL_EXTERNAL void test(sycl::accessor<int, 1, sycl::access::mode::read_write,
 }
 
 // clang-format off
-// CHECK-O3:  call spir_func void {{.*}}spirv_ControlBarrierjjj
+// CHECK:  call spir_func void {{.*}}spirv_ControlBarrierjjj
 
 // load() for global address space
-// CHECK-O3: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
-// CHECK-O3: {{.*}}SubgroupLocalInvocationId
-// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
-// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
+// CHECK: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
+// CHECK: {{.*}}SubgroupLocalInvocationId
+// CHECK: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
+// CHECK: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
 
 
 // load() for local address space
-// CHECK-O3: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
-// CHECK-O3: {{.*}}SubgroupLocalInvocationId
-// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
-// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
+// CHECK: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
+// CHECK: {{.*}}SubgroupLocalInvocationId
+// CHECK: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
+// CHECK: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
 
 // load() for private address space
-// CHECK-O3: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
-// CHECK-O3: {{.*}}SubgroupLocalInvocationId
-// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
-// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
+// CHECK: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
+// CHECK: {{.*}}SubgroupLocalInvocationId
+// CHECK: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
+// CHECK: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
 
 // store() for global address space
 // NOTE: Call to __spirv_GenericCastToPtrExplicit_ToLocal is consolidated with an earlier call to it.
-// CHECK-O3: {{.*}}SubgroupLocalInvocationId
-// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
-// CHECK-O3: call spir_func void {{.*}}spirv_SubgroupBlockWriteINTEL{{.*}}(ptr addrspace(1)
-
-// load() accepting raw pointers method
-// CHECK-O0: define{{.*}}spir_func i32 {{.*}}4sycl3_V19sub_group4load{{.*}}addrspace(4) %
-// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
-// CHECK-O0: call spir_func i32 {{.*}}sycl3_V19sub_group4load{{.*}}ptr addrspace(3) %
-// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
-// CHECK-O0: call spir_func i32 {{.*}}sycl3_V19sub_group4load{{.*}}ptr addrspace(1) %
-
-// store() accepting raw pointers method
-// CHECK-O0: define{{.*}}spir_func void {{.*}}4sycl3_V19sub_group5store{{.*}}ptr addrspace(4) %
-// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
-// CHECK-O0: call spir_func void {{.*}}4sycl3_V19sub_group5store{{.*}}, ptr addrspace(3) %
-// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
-// CHECK-O0: call spir_func void {{.*}}4sycl3_V19sub_group5store{{.*}}, ptr addrspace(1) %
-
-// CHECK-O0: define {{.*}}spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4) %
-// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
-// CHECK-O0: define {{.*}}spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4) %
-// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
+// CHECK: {{.*}}SubgroupLocalInvocationId
+// CHECK: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
+// CHECK: call spir_func void {{.*}}spirv_SubgroupBlockWriteINTEL{{.*}}(ptr addrspace(1)
 // clang-format off
diff --git a/sycl/test/extensions/address_cast_negative.cpp b/sycl/test/extensions/address_cast_negative.cpp
index 4324314405be6..f325d283eed9e 100644
--- a/sycl/test/extensions/address_cast_negative.cpp
+++ b/sycl/test/extensions/address_cast_negative.cpp
@@ -5,10 +5,10 @@
 using namespace sycl::ext::oneapi::experimental;
 
 SYCL_EXTERNAL void test(int *p) {
-  // expected-error-re@sycl/ext/oneapi/experimental/address_cast.hpp:* {{{{.*}}Not supported yet!}}
+  // expected-error-re@sycl/access/access.hpp:* {{{{.*}}Not supported yet!}}
   std::ignore = dynamic_address_cast<
       sycl::access::address_space::ext_intel_global_device_space>(p);
-  // expected-error-re@sycl/ext/oneapi/experimental/address_cast.hpp:* {{{{.*}}Not supported yet!}}
+  // expected-error-re@sycl/access/access.hpp:* {{{{.*}}Not supported yet!}}
   std::ignore = dynamic_address_cast<
       sycl::access::address_space::ext_intel_global_host_space>(p);
 }