From 21c3f33ea9b045ae149633fe4ad871e36bcafe9f Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Fri, 13 Sep 2024 11:09:42 -0700 Subject: [PATCH 1/3] [SYCL] Add a test showing issues with sycl_ext_oneapi_address_cast impl --- .../extensions/address_cast.cpp | 93 +++++++++++++++++++ 1 file changed, 93 insertions(+) create mode 100644 sycl/test/check_device_code/extensions/address_cast.cpp diff --git a/sycl/test/check_device_code/extensions/address_cast.cpp b/sycl/test/check_device_code/extensions/address_cast.cpp new file mode 100644 index 0000000000000..f73c7233e1d10 --- /dev/null +++ b/sycl/test/check_device_code/extensions/address_cast.cpp @@ -0,0 +1,93 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// RUN: %clangxx -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code -o - %s | FileCheck %s + +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental; + +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-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: ret void +// +SYCL_EXTERNAL auto to_global_decorated(int *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-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: ret void +// +SYCL_EXTERNAL auto to_global_not_decorated(int *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 readnone sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] !srcloc [[META22:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: unreachable +// +SYCL_EXTERNAL auto to_generic_decorated(int *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 readnone sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META23:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: unreachable +// +SYCL_EXTERNAL auto to_generic_not_decorated(int *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 [[META24:![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 [[META25:![0-9]+]] +// CHECK-NEXT: ret void +// +SYCL_EXTERNAL auto to_global_decorated(int *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 [[META28:![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 [[META29:![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_decoratedEPi( +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture readnone sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META32:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: unreachable +// +SYCL_EXTERNAL auto to_generic_decorated(int *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 readnone sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META33:![0-9]+]] !sycl_fixed_targets [[META7]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: unreachable +// +SYCL_EXTERNAL auto to_generic_not_decorated(int *p) { + return dynamic_address_cast(p); +} +} // namespace dynamic_as_cast From 7d4af8793824d5cabf4ec8c716610d17851a4a9c Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Fri, 13 Sep 2024 11:11:18 -0700 Subject: [PATCH 2/3] [SYCL] Fix `static|dynamic_address_cast` to `generic` SPIRV operations are defined such that `OpGenericCastToPtr` and `OpGenericCastToPtrExplicit` cannot be used when target `Storage Class` is `Generic`, yet we were generating such code. This PR fixes that. --- .../ext/oneapi/experimental/address_cast.hpp | 26 ++++++++++++++--- .../extensions/address_cast.cpp | 28 +++++++++++-------- 2 files changed, 38 insertions(+), 16 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp b/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp index e51219cfee462..a3b74a24f2fe6 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp @@ -22,8 +22,19 @@ template static_address_cast(ElementType *Ptr) { #ifdef __SYCL_DEVICE_ONLY__ - auto CastPtr = sycl::detail::spirv::GenericCastToPtr(Ptr); - return multi_ptr(CastPtr); + // 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) { + // 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); + } else { + auto CastPtr = sycl::detail::spirv::GenericCastToPtr(Ptr); + return multi_ptr(CastPtr); + } #else return multi_ptr(Ptr); #endif @@ -34,8 +45,15 @@ template dynamic_address_cast(ElementType *Ptr) { #ifdef __SYCL_DEVICE_ONLY__ - auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit(Ptr); - return multi_ptr(CastPtr); + // 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); + } else { + auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit(Ptr); + return multi_ptr(CastPtr); + } #else return multi_ptr(Ptr); #endif diff --git a/sycl/test/check_device_code/extensions/address_cast.cpp b/sycl/test/check_device_code/extensions/address_cast.cpp index f73c7233e1d10..7abbce22d5cfc 100644 --- a/sycl/test/check_device_code/extensions/address_cast.cpp +++ b/sycl/test/check_device_code/extensions/address_cast.cpp @@ -30,18 +30,20 @@ SYCL_EXTERNAL auto to_global_not_decorated(int *p) { access::decorated::no>(p); } // CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast20to_generic_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture readnone sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] !srcloc [[META22:![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 [[META22:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: unreachable +// 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: ret void // SYCL_EXTERNAL auto to_generic_decorated(int *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 readnone sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META23:![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 #[[ATTR2]] !srcloc [[META28:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: unreachable +// 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: ret void // SYCL_EXTERNAL auto to_generic_not_decorated(int *p) { return static_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 [[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.0") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META38:![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 [[META29:![0-9]+]] +// 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: ret void // SYCL_EXTERNAL auto to_global_not_decorated(int *p) { @@ -73,18 +75,20 @@ SYCL_EXTERNAL auto to_global_not_decorated(int *p) { access::decorated::no>(p); } // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast20to_generic_decoratedEPi( -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture readnone sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META32:![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 [[META42:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: unreachable +// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA23]], !alias.scope [[META43:![0-9]+]] +// CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_generic_decorated(int *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 readnone sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META33:![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 #[[ATTR2]] !srcloc [[META46:![0-9]+]] !sycl_fixed_targets [[META7]] { // CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: unreachable +// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA29]], !alias.scope [[META47:![0-9]+]] +// CHECK-NEXT: ret void // SYCL_EXTERNAL auto to_generic_not_decorated(int *p) { return dynamic_address_cast Date: Fri, 13 Sep 2024 13:10:41 -0700 Subject: [PATCH 3/3] Disable newly added test on Windows --- sycl/test/check_device_code/extensions/address_cast.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/sycl/test/check_device_code/extensions/address_cast.cpp b/sycl/test/check_device_code/extensions/address_cast.cpp index 7abbce22d5cfc..51aa87c211eb9 100644 --- a/sycl/test/check_device_code/extensions/address_cast.cpp +++ b/sycl/test/check_device_code/extensions/address_cast.cpp @@ -1,6 +1,11 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 // RUN: %clangxx -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code -o - %s | FileCheck %s +// Linux/Windows have minor differences in the generated IR (e.g. TBAA +// metadata). Having linux-only checks eases the maintenance without sacrifising +// coverage of what's important for this test. +// REQUIRES: linux + #include using namespace sycl;