From e8ebe3f4daa5269788dd5a547aa97a618ec2af49 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Tue, 24 Sep 2024 12:44:11 -0700 Subject: [PATCH] [SYCL] Support `*global_[device|host]_space` in `static_address_cast` When these address spaces are used with regular `sycl::detail::spirv::GenericCastToPtr` they are turned into `unreachable`. We don't have a SPIR-V intrinsic yet (or maybe we shouldn't even have it, and will continue to rely on standard LLVM IR's `addrspacecast`), so use C-style cast and rely on the translator/backend to generate proper operation, similarly to `sycl::detail::cast_AS`. --- .../ext/oneapi/experimental/address_cast.hpp | 29 +++++++++++++++ .../extensions/address_cast.cpp | 36 ++++++++++--------- .../test/extensions/address_cast_negative.cpp | 14 ++++++++ 3 files changed, 63 insertions(+), 16 deletions(-) create mode 100644 sycl/test/extensions/address_cast_negative.cpp diff --git a/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp b/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp index 39542b58b284d..8e901d7d2877a 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp @@ -16,6 +16,9 @@ inline namespace _V1 { namespace ext { namespace oneapi { namespace 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; @@ -32,6 +35,18 @@ static_address_cast(ElementType *Ptr) { 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::type *; + auto CastPtr = (raw_type)(Ptr); +#else + auto CastPtr = sycl::detail::spirv::GenericCastToPtr(Ptr); +#endif + return ret_ty(CastPtr); } else { auto CastPtr = sycl::detail::spirv::GenericCastToPtr(Ptr); return ret_ty(CastPtr); @@ -60,6 +75,20 @@ dynamic_address_cast(ElementType *Ptr) { "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(Ptr); + return ret_ty(CastPtr); +#endif } else { auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit(Ptr); return ret_ty(CastPtr); diff --git a/sycl/test/check_device_code/extensions/address_cast.cpp b/sycl/test/check_device_code/extensions/address_cast.cpp index 0f02caadfd2db..36251cecbbc3b 100644 --- a/sycl/test/check_device_code/extensions/address_cast.cpp +++ b/sycl/test/check_device_code/extensions/address_cast.cpp @@ -23,7 +23,7 @@ namespace static_as_cast { // CHECK-NEXT: [[ENTRY:.*:]] // 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) #[[ATTR6:[0-9]+]] +// 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 // @@ -33,7 +33,7 @@ SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr 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.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) #[[ATTR6]] +// 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 // @@ -61,18 +61,22 @@ SYCL_EXTERNAL auto to_generic_not_decorated(int *p) { } // CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast16to_global_deviceEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture readnone sret(%"class.sycl::_V1::multi_ptr.3") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR5:[0-9]+]] !srcloc [[META35:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.3") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META35:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: unreachable +// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(5) +// CHECK-NEXT: store ptr addrspace(5) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA36:![0-9]+]], !alias.scope [[META38:![0-9]+]] +// CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_global_device(int *p) { return static_address_cast(p); } // CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast14to_global_hostEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture readnone sret(%"class.sycl::_V1::multi_ptr.4") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR5]] !srcloc [[META36:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.4") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META41:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: unreachable +// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(6) +// CHECK-NEXT: store ptr addrspace(6) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA42:![0-9]+]], !alias.scope [[META44:![0-9]+]] +// CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_global_host(int *p) { return static_address_cast(p); @@ -81,41 +85,41 @@ SYCL_EXTERNAL auto to_global_host(int *p) { namespace dynamic_as_cast { // 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 [[META37:![0-9]+]] !sycl_fixed_targets [[META7]] { +// 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 [[META47:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] // 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) #[[ATTR6]] -// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12]], !alias.scope [[META38:![0-9]+]] +// 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 [[META48:![0-9]+]] // CHECK-NEXT: ret void // 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.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META43:![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 [[META53:![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) #[[ATTR6]] -// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META44:![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 [[META54:![0-9]+]] // CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_global_not_decorated(int *p) { return dynamic_address_cast(p); } // 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 [[META47:![0-9]+]] !sycl_fixed_targets [[META7]] { +// 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 [[META57:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] // 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 [[META48:![0-9]+]] +// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META58:![0-9]+]] // CHECK-NEXT: ret void // 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 #[[ATTR4]] !srcloc [[META51:![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 [[META61:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META52:![0-9]+]] +// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META62:![0-9]+]] // CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_generic_not_decorated(int *p) { diff --git a/sycl/test/extensions/address_cast_negative.cpp b/sycl/test/extensions/address_cast_negative.cpp new file mode 100644 index 0000000000000..4324314405be6 --- /dev/null +++ b/sycl/test/extensions/address_cast_negative.cpp @@ -0,0 +1,14 @@ +// RUN: %clangxx -D__ENABLE_USM_ADDR_SPACE__ -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=warning,note %s + +#include + +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!}} + 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!}} + std::ignore = dynamic_address_cast< + sycl::access::address_space::ext_intel_global_host_space>(p); +}