[SYCL] Fix SYCL-CTS regression after intel#15543

Apparently, we perform some hacks around constant address space that I
wasn't aware of. Make sure new address cast helpers are called
consistently with the previous behavior.

Author: aelovikov-intel

devops/cts_exclude_filter_L0_GPU

@@ -3,5 +3,3 @@ kernel_bundle
 marray
 # fix: https://github.com/KhronosGroup/SYCL-CTS/pull/964
 accessor_legacy
-# CMPLRLLVM-62822
-multi_ptr

devops/cts_exclude_filter_OCL_CPU

@@ -7,5 +7,3 @@ math_builtin_api
 hierarchical
 # fix: https://github.com/KhronosGroup/SYCL-CTS/pull/964
 accessor_legacy
-# CMPLRLLVM-62822
-multi_ptr

sycl/include/sycl/multi_ptr.hpp

@@ -732,12 +732,28 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<void, Space, DecorateAddress> {
   decorated_type *m_Pointer;
 };
 
+namespace detail {
+// See access.hpp's DecoratedType<..., access::address_space::constant_space>.
+//
+// This is only applicable to `access::decorated::legacy` mode because constant
+// AS is deprecated itself and is only accessible in legacy modes.
+template <auto Space>
+#ifdef __SYCL_DEVICE_ONLY__
+inline constexpr auto decoration_space =
+    deduce_AS<typename DecoratedType<void, Space>::type>::value;
+#else
+inline constexpr auto decoration_space = Space;
+#endif
+}
+
 // Legacy specialization of multi_ptr.
 // TODO: Add deprecation warning here when possible.
 template <typename ElementType, access::address_space Space>
 class __SYCL2020_DEPRECATED(
     "decorated::legacy multi_ptr specialization is deprecated since SYCL 2020.")
     multi_ptr<ElementType, Space, access::decorated::legacy> {
+  static constexpr auto DecorationSpace = detail::decoration_space<Space>;
+
 public:
   using value_type = ElementType;
   using element_type =
@@ -777,7 +793,8 @@ class __SYCL2020_DEPRECATED(
 
   multi_ptr(ElementType *pointer)
       : m_Pointer(detail::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
   }
@@ -786,7 +803,8 @@ class __SYCL2020_DEPRECATED(
   template <typename = typename detail::const_if_const_AS<Space, ElementType>>
   multi_ptr(const ElementType *pointer)
       : m_Pointer(detail::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {}
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {}
 #endif
 
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
@@ -814,7 +832,7 @@ class __SYCL2020_DEPRECATED(
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
     m_Pointer = detail::dynamic_address_cast<
-        Space, /* SupressNotImplementedAssert = */ true>(pointer);
+        DecorationSpace, /* SupressNotImplementedAssert = */ true>(pointer);
     return *this;
   }
 
@@ -856,8 +874,8 @@ class __SYCL2020_DEPRECATED(
   multi_ptr(accessor<ElementType, dimensions, Mode, target::device,
                      isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(
-            detail::static_address_cast<Space>(Accessor.get_pointer().get())) {}
+      : multi_ptr(detail::static_address_cast<DecorationSpace>(
+            Accessor.get_pointer().get())) {}
 
   // Only if Space == local_space || generic_space
   template <
@@ -1088,6 +1106,8 @@ template <access::address_space Space>
 class __SYCL2020_DEPRECATED(
     "decorated::legacy multi_ptr specialization is deprecated since SYCL 2020.")
     multi_ptr<void, Space, access::decorated::legacy> {
+  static constexpr auto DecorationSpace = detail::decoration_space<Space>;
+
 public:
   using value_type = void;
   using element_type = void;
@@ -1113,17 +1133,17 @@ class __SYCL2020_DEPRECATED(
                                   !std::is_same_v<RelayPointerT, void *>>>
   multi_ptr(void *pointer)
       : m_Pointer(detail::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {
+                  DecorationSpace, /* 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::dynamic_address_cast<
-                pointer_t, /* SupressNotImplementedAssert = */ true>(pointer)) {
-  }
+      : m_Pointer(detail::dynamic_address_cast<
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {}
 #endif
 #endif
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
@@ -1154,7 +1174,7 @@ class __SYCL2020_DEPRECATED(
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
     m_Pointer = detail::dynamic_address_cast<
-        Space, /* SupressNotImplementedAssert = */ true>(pointer);
+        DecorationSpace, /* SupressNotImplementedAssert = */ true>(pointer);
     return *this;
   }
 #endif
@@ -1249,6 +1269,8 @@ template <access::address_space Space>
 class __SYCL2020_DEPRECATED(
     "decorated::legacy multi_ptr specialization is deprecated since SYCL 2020.")
     multi_ptr<const void, Space, access::decorated::legacy> {
+  static constexpr auto DecorationSpace = detail::decoration_space<Space>;
+
 public:
   using value_type = const void;
   using element_type = const void;
@@ -1263,6 +1285,7 @@ class __SYCL2020_DEPRECATED(
 
   static constexpr access::address_space address_space = Space;
 
+public:
   // Constructors
   multi_ptr() : m_Pointer(nullptr) {}
   multi_ptr(const multi_ptr &) = default;
@@ -1275,15 +1298,17 @@ class __SYCL2020_DEPRECATED(
                                   !std::is_same_v<RelayPointerT, const void *>>>
   multi_ptr(const void *pointer)
       : m_Pointer(detail::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {
+                  DecorationSpace, /* 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::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {}
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {}
 #endif
 #endif
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
@@ -1314,7 +1339,7 @@ class __SYCL2020_DEPRECATED(
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
     m_Pointer = detail::dynamic_address_cast<
-        pointer_t, /* SupressNotImplementedAssert = */ true>(pointer);
+        DecorationSpace, /* SupressNotImplementedAssert = */ true>(pointer);
     return *this;
   }
 #endif
@@ -1442,7 +1467,7 @@ 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::dynamic_address_cast<Space,
+      detail::dynamic_address_cast<detail::decoration_space<Space>,
                                    /* SupressNotImplementedAssert = */ true>(
           pointer));
 }

sycl/test-e2e/Basic/multi_ptr_legacy.hpp

@@ -157,6 +157,7 @@ template <typename T> void testMultPtrArrowOperator() {
     buffer<point<T>, 1> bufferData_2(data_2, numOfItems);
     buffer<point<T>, 1> bufferData_3(data_3, numOfItems);
     buffer<point<T>, 1> bufferData_4(data_4, numOfItems);
+    buffer<bool, 1> result_buf{1};
     queue myQueue;
     myQueue.submit([&](handler &cgh) {
       accessor<point<T>, 1, access::mode::read, access::target::device,
@@ -170,8 +171,12 @@ template <typename T> void testMultPtrArrowOperator() {
                access::placeholder::false_t>
           accessorData_4(bufferData_4, cgh);
 
+      accessor result{result_buf, cgh};
+
       cgh.parallel_for<class testMultPtrArrowOperatorKernel<T>>(
           sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1>) {
+            // Initialize local memory:
+            accessorData_3[0] = T{3};
             auto ptr_1 =
                 make_ptr<const point<T>, access::address_space::global_space,
                          access::decorated::legacy>(
@@ -195,6 +200,12 @@ template <typename T> void testMultPtrArrowOperator() {
             auto x3 = ptr_3->x;
             auto x4 = ptr_4->x;
 
+            result[0] = true;
+            result[0] &= x1 == T{1};
+            result[0] &= x2 == T{2};
+            result[0] &= x3 == T{3};
+            result[0] &= x4 == T{4};
+
             static_assert(std::is_same<decltype(x1), T>::value,
                           "Expected decltype(ptr_1->x) == T");
             static_assert(std::is_same<decltype(x2), T>::value,
@@ -205,6 +216,8 @@ template <typename T> void testMultPtrArrowOperator() {
                           "Expected decltype(ptr_4->x) == T");
           });
     });
+
+    assert(sycl::host_accessor{result_buf}[0]);
   }
 }