From fbb0fd101c1c48a1456edf3f855b5d701b5b3a59 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Mon, 16 Sep 2024 10:15:38 -0700 Subject: [PATCH 1/2] [SYCL] Align `sycl_ext_oneapi_address_cast` impl with the spec Reflects spec changes from https://github.com/intel/llvm/pull/12689 --- .../sycl_ext_oneapi_address_cast.asciidoc | 8 +- .../ext/oneapi/experimental/address_cast.hpp | 49 +++++++--- .../extensions/address_cast.cpp | 92 ++++++++++--------- 3 files changed, 89 insertions(+), 60 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_address_cast.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_address_cast.asciidoc index 5ad284907debe..deb597b7de02f 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_address_cast.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_address_cast.asciidoc @@ -94,10 +94,10 @@ implementation supports. namespace sycl::ext::oneapi::experimental { // Shorthands for address space names -constexpr inline address_space global_space = sycl::access::address_space::global_space; -constexpr inline address_space local_space = sycl::access::address_space::local_space; -constexpr inline address_space private_space = sycl::access::address_space::private_space; -constexpr inline address_space generic_space = sycl::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 diff --git a/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp b/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp index a3b74a24f2fe6..4814d01b9c1ef 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp @@ -16,49 +16,74 @@ inline namespace _V1 { namespace ext { namespace oneapi { namespace experimental { +// 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; -template -multi_ptr +template +multi_ptr static_address_cast(ElementType *Ptr) { + using ret_ty = multi_ptr; #ifdef __SYCL_DEVICE_ONLY__ // TODO: Remove this restriction. static_assert(std::is_same_v>, "The extension expect undecorated raw pointers only!"); - if constexpr (Space == access::address_space::generic_space) { + if constexpr (Space == generic_space) { // Undecorated raw pointer is in generic AS already, no extra casts needed. // Note for future, for `OpPtrCastToGeneric`, `Pointer` must point to one of // `Storage Classes` that doesn't include `Generic`, so this will have to // remain a special case even if the restriction above is lifted. - return multi_ptr(Ptr); + return ret_ty(Ptr); } else { auto CastPtr = sycl::detail::spirv::GenericCastToPtr(Ptr); - return multi_ptr(CastPtr); + return ret_ty(CastPtr); } #else - return multi_ptr(Ptr); + return ret_ty(Ptr); #endif } template -multi_ptr +multi_ptr static_address_cast( + multi_ptr Ptr) { + if constexpr (Space == generic_space) + return Ptr; + else + return {static_address_cast(Ptr.get_raw())}; +} + +template +multi_ptr dynamic_address_cast(ElementType *Ptr) { + using ret_ty = multi_ptr; #ifdef __SYCL_DEVICE_ONLY__ // TODO: Remove this restriction. static_assert(std::is_same_v>, "The extension expect undecorated raw pointers only!"); - if constexpr (Space == access::address_space::generic_space) { - return multi_ptr(Ptr); + if constexpr (Space == generic_space) { + return ret_ty(Ptr); } else { auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit(Ptr); - return multi_ptr(CastPtr); + return ret_ty(CastPtr); } #else - return multi_ptr(Ptr); + return ret_ty(Ptr); #endif } +template +multi_ptr dynamic_address_cast( + multi_ptr Ptr) { + if constexpr (Space == generic_space) + return Ptr; + else + return {dynamic_address_cast(Ptr.get_raw())}; +} + } // namespace experimental } // namespace oneapi } // namespace ext diff --git a/sycl/test/check_device_code/extensions/address_cast.cpp b/sycl/test/check_device_code/extensions/address_cast.cpp index 51aa87c211eb9..e8013a3b6cebe 100644 --- a/sycl/test/check_device_code/extensions/address_cast.cpp +++ b/sycl/test/check_device_code/extensions/address_cast.cpp @@ -11,92 +11,96 @@ using namespace sycl; using namespace sycl::ext::oneapi::experimental; +// FIXME: should be removed when https://github.com/intel/llvm/pull/15389 is merged in. +template +using decorated_generic_ptr = + multi_ptr; + namespace static_as_cast { -// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast19to_global_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !srcloc [[META6:![0-9]+]] !sycl_fixed_targets [[META7:![0-9]+]] { +// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast19to_global_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE( +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !srcloc [[META6:![0-9]+]] !sycl_fixed_targets [[META7:![0-9]+]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR3:[0-9]+]] -// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8:![0-9]+]], !alias.scope [[META13:![0-9]+]] +// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8:![0-9]+]] +// CHECK-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) +// CHECK-NEXT: [[CALL_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[TMP1]], i32 noundef 5) #[[ATTR5:[0-9]+]] +// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12:![0-9]+]], !alias.scope [[META14:![0-9]+]] // CHECK-NEXT: ret void // -SYCL_EXTERNAL auto to_global_decorated(int *p) { - return static_address_cast(p); +SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr p) { + return static_address_cast(p); } // CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast23to_global_not_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META16:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] !srcloc [[META19:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR3]] -// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA17:![0-9]+]], !alias.scope [[META19:![0-9]+]] +// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR5]] +// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20:![0-9]+]], !alias.scope [[META22:![0-9]+]] // CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_global_not_decorated(int *p) { - return static_address_cast(p); + return static_address_cast(p); } -// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast20to_generic_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] !srcloc [[META22:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast20to_generic_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE( +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR3:[0-9]+]] !srcloc [[META25:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA23:![0-9]+]], !alias.scope [[META25:![0-9]+]] +// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]] +// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META26:![0-9]+]] // CHECK-NEXT: ret void // -SYCL_EXTERNAL auto to_generic_decorated(int *p) { - return static_address_cast(p); +SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr p) { + return static_address_cast(p); } // CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast24to_generic_not_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META28:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4:[0-9]+]] !srcloc [[META29:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA29:![0-9]+]], !alias.scope [[META31:![0-9]+]] +// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30:![0-9]+]], !alias.scope [[META32:![0-9]+]] // CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_generic_not_decorated(int *p) { - return static_address_cast(p); + return static_address_cast(p); } } // namespace static_as_cast namespace dynamic_as_cast { -// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast19to_global_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META34:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast19to_global_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE( +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META35:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR3]] -// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META35:![0-9]+]] +// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]] +// CHECK-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) +// CHECK-NEXT: [[CALL_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[TMP1]], i32 noundef 5) #[[ATTR5]] +// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12]], !alias.scope [[META36:![0-9]+]] // CHECK-NEXT: ret void // -SYCL_EXTERNAL auto to_global_decorated(int *p) { - return dynamic_address_cast(p); +SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr p) { + return dynamic_address_cast(p); } // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast23to_global_not_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META38:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META41:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR3]] -// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA17]], !alias.scope [[META39:![0-9]+]] +// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR5]] +// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META42:![0-9]+]] // CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_global_not_decorated(int *p) { - return dynamic_address_cast(p); + return dynamic_address_cast(p); } -// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast20to_generic_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META42:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast20to_generic_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE( +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR3]] !srcloc [[META45:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA23]], !alias.scope [[META43:![0-9]+]] +// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]] +// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META46:![0-9]+]] // CHECK-NEXT: ret void // -SYCL_EXTERNAL auto to_generic_decorated(int *p) { - return dynamic_address_cast(p); +SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr p) { + return dynamic_address_cast(p); } // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast24to_generic_not_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META46:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META49:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA29]], !alias.scope [[META47:![0-9]+]] +// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META50:![0-9]+]] // CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_generic_not_decorated(int *p) { - return dynamic_address_cast(p); + return dynamic_address_cast(p); } } // namespace dynamic_as_cast From 1244fc28f6f65d7c339cf526435cb94f8acc3f45 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Mon, 16 Sep 2024 11:28:30 -0700 Subject: [PATCH 2/2] Update e2e tests as well --- .../AddressCast/dynamic_address_cast.cpp | 36 +++++++++---------- .../AddressCast/static_address_cast.cpp | 12 +++---- 2 files changed, 24 insertions(+), 24 deletions(-) diff --git a/sycl/test-e2e/AddressCast/dynamic_address_cast.cpp b/sycl/test-e2e/AddressCast/dynamic_address_cast.cpp index 797c0e6c022f7..441fe486564b3 100644 --- a/sycl/test-e2e/AddressCast/dynamic_address_cast.cpp +++ b/sycl/test-e2e/AddressCast/dynamic_address_cast.cpp @@ -42,16 +42,16 @@ int main() { { auto GlobalPointer = sycl::ext::oneapi::experimental::dynamic_address_cast< - sycl::access::address_space::global_space, - sycl::access::decorated::no>(RawGlobalPointer); + sycl::access::address_space::global_space>( + RawGlobalPointer); auto LocalPointer = sycl::ext::oneapi::experimental::dynamic_address_cast< - sycl::access::address_space::local_space, - sycl::access::decorated::no>(RawGlobalPointer); + sycl::access::address_space::local_space>( + RawGlobalPointer); auto PrivatePointer = sycl::ext::oneapi::experimental::dynamic_address_cast< - sycl::access::address_space::private_space, - sycl::access::decorated::no>(RawGlobalPointer); + sycl::access::address_space::private_space>( + RawGlobalPointer); Success &= reinterpret_cast(RawGlobalPointer) == reinterpret_cast(GlobalPointer.get_raw()); Success &= LocalPointer.get_raw() == nullptr; @@ -62,16 +62,16 @@ int main() { { auto GlobalPointer = sycl::ext::oneapi::experimental::dynamic_address_cast< - sycl::access::address_space::global_space, - sycl::access::decorated::no>(RawLocalPointer); + sycl::access::address_space::global_space>( + RawLocalPointer); auto LocalPointer = sycl::ext::oneapi::experimental::dynamic_address_cast< - sycl::access::address_space::local_space, - sycl::access::decorated::no>(RawLocalPointer); + sycl::access::address_space::local_space>( + RawLocalPointer); auto PrivatePointer = sycl::ext::oneapi::experimental::dynamic_address_cast< - sycl::access::address_space::private_space, - sycl::access::decorated::no>(RawLocalPointer); + sycl::access::address_space::private_space>( + RawLocalPointer); Success &= GlobalPointer.get_raw() == nullptr; Success &= reinterpret_cast(RawLocalPointer) == reinterpret_cast(LocalPointer.get_raw()); @@ -83,16 +83,16 @@ int main() { { auto GlobalPointer = sycl::ext::oneapi::experimental::dynamic_address_cast< - sycl::access::address_space::global_space, - sycl::access::decorated::no>(RawPrivatePointer); + sycl::access::address_space::global_space>( + RawPrivatePointer); auto LocalPointer = sycl::ext::oneapi::experimental::dynamic_address_cast< - sycl::access::address_space::local_space, - sycl::access::decorated::no>(RawPrivatePointer); + sycl::access::address_space::local_space>( + RawPrivatePointer); auto PrivatePointer = sycl::ext::oneapi::experimental::dynamic_address_cast< - sycl::access::address_space::private_space, - sycl::access::decorated::no>(RawPrivatePointer); + sycl::access::address_space::private_space>( + RawPrivatePointer); Success &= GlobalPointer.get_raw() == nullptr; Success &= LocalPointer.get_raw() == nullptr; Success &= reinterpret_cast(RawPrivatePointer) == diff --git a/sycl/test-e2e/AddressCast/static_address_cast.cpp b/sycl/test-e2e/AddressCast/static_address_cast.cpp index 15a997ba81e53..a30068cb3fafb 100644 --- a/sycl/test-e2e/AddressCast/static_address_cast.cpp +++ b/sycl/test-e2e/AddressCast/static_address_cast.cpp @@ -39,16 +39,16 @@ int main() { int *RawGlobalPointer = &GlobalAccessor[Index]; auto GlobalPointer = sycl::ext::oneapi::experimental::static_address_cast< - sycl::access::address_space::global_space, - sycl::access::decorated::no>(RawGlobalPointer); + sycl::access::address_space::global_space>( + RawGlobalPointer); Success &= reinterpret_cast(RawGlobalPointer) == reinterpret_cast(GlobalPointer.get_raw()); int *RawLocalPointer = &LocalAccessor[0]; auto LocalPointer = sycl::ext::oneapi::experimental::static_address_cast< - sycl::access::address_space::local_space, - sycl::access::decorated::no>(RawLocalPointer); + sycl::access::address_space::local_space>( + RawLocalPointer); Success &= reinterpret_cast(RawLocalPointer) == reinterpret_cast(LocalPointer.get_raw()); @@ -56,8 +56,8 @@ int main() { int *RawPrivatePointer = &PrivateVariable; auto PrivatePointer = sycl::ext::oneapi::experimental::static_address_cast< - sycl::access::address_space::private_space, - sycl::access::decorated::no>(RawPrivatePointer); + sycl::access::address_space::private_space>( + RawPrivatePointer); Success &= reinterpret_cast(RawPrivatePointer) == reinterpret_cast(PrivatePointer.get_raw());