diff --git a/devops/cts_exclude_filter_L0_GPU b/devops/cts_exclude_filter_L0_GPU index ddc388477b6e6..dfd8b4623bae1 100644 --- a/devops/cts_exclude_filter_L0_GPU +++ b/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 diff --git a/devops/cts_exclude_filter_OCL_CPU b/devops/cts_exclude_filter_OCL_CPU index 18e6f333094f5..24f4a5c9eb41b 100644 --- a/devops/cts_exclude_filter_OCL_CPU +++ b/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 diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp index 9bca96ade3879..e84e1094d2ecb 100644 --- a/sycl/include/sycl/multi_ptr.hpp +++ b/sycl/include/sycl/multi_ptr.hpp @@ -732,12 +732,28 @@ class __SYCL_TYPE(multi_ptr) multi_ptr { 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 +#ifdef __SYCL_DEVICE_ONLY__ +inline constexpr auto decoration_space = + deduce_AS::type>::value; +#else +inline constexpr auto decoration_space = Space; +#endif +} // namespace detail + // Legacy specialization of multi_ptr. // TODO: Add deprecation warning here when possible. template class __SYCL2020_DEPRECATED( "decorated::legacy multi_ptr specialization is deprecated since SYCL 2020.") multi_ptr { + static constexpr auto DecorationSpace = detail::decoration_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 > 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 Accessor) - : multi_ptr( - detail::static_address_cast(Accessor.get_pointer().get())) {} + : multi_ptr(detail::static_address_cast( + Accessor.get_pointer().get())) {} // Only if Space == local_space || generic_space template < @@ -1088,6 +1106,8 @@ template class __SYCL2020_DEPRECATED( "decorated::legacy multi_ptr specialization is deprecated since SYCL 2020.") multi_ptr { + static constexpr auto DecorationSpace = detail::decoration_space; + public: using value_type = void; using element_type = void; @@ -1113,17 +1133,17 @@ class __SYCL2020_DEPRECATED( !std::is_same_v>> 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 > 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 class __SYCL2020_DEPRECATED( "decorated::legacy multi_ptr specialization is deprecated since SYCL 2020.") multi_ptr { + static constexpr auto DecorationSpace = detail::decoration_space; + public: using value_type = const void; using element_type = const void; @@ -1275,7 +1297,8 @@ class __SYCL2020_DEPRECATED( !std::is_same_v>> 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. } @@ -1283,7 +1306,8 @@ class __SYCL2020_DEPRECATED( template > 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 +1338,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 +1466,7 @@ address_space_cast(ElementType *pointer) { // space is not compatible with Space. // Use LegacyPointerTypes here to also allow constant_space return multi_ptr( - detail::dynamic_address_cast, /* SupressNotImplementedAssert = */ true>( pointer)); } diff --git a/sycl/test-e2e/Basic/multi_ptr_legacy.hpp b/sycl/test-e2e/Basic/multi_ptr_legacy.hpp index 7e854c98036a8..6860b21ee0683 100644 --- a/sycl/test-e2e/Basic/multi_ptr_legacy.hpp +++ b/sycl/test-e2e/Basic/multi_ptr_legacy.hpp @@ -157,6 +157,7 @@ template void testMultPtrArrowOperator() { buffer, 1> bufferData_2(data_2, numOfItems); buffer, 1> bufferData_3(data_3, numOfItems); buffer, 1> bufferData_4(data_4, numOfItems); + buffer result_buf{1}; queue myQueue; myQueue.submit([&](handler &cgh) { accessor, 1, access::mode::read, access::target::device, @@ -170,8 +171,12 @@ template void testMultPtrArrowOperator() { access::placeholder::false_t> accessorData_4(bufferData_4, cgh); + accessor result{result_buf, cgh}; + cgh.parallel_for>( sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1>) { + // Initialize local memory: + accessorData_3[0] = T{3}; auto ptr_1 = make_ptr, access::address_space::global_space, access::decorated::legacy>( @@ -195,6 +200,12 @@ template 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::value, "Expected decltype(ptr_1->x) == T"); static_assert(std::is_same::value, @@ -205,6 +216,8 @@ template void testMultPtrArrowOperator() { "Expected decltype(ptr_4->x) == T"); }); }); + + assert(sycl::host_accessor{result_buf}[0]); } }