Skip to content

Commit

Permalink
[SYCL] Fix static|dynamic_address_cast to generic
Browse files Browse the repository at this point in the history 
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.
  • Loading branch information
@aelovikov-intel
aelovikov-intel committed Sep 13, 2024
1 parent 21c3f33 commit 7d4af87
Show file tree
Hide file tree
Showing 2 changed files with 38 additions and 16 deletions.
26 changes: 22 additions & 4 deletions sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,8 +22,19 @@ template <access::address_space Space, access::decorated DecorateAddress,
multi_ptr<ElementType, Space, DecorateAddress>
static_address_cast(ElementType *Ptr) {
#ifdef __SYCL_DEVICE_ONLY__
auto CastPtr = sycl::detail::spirv::GenericCastToPtr<Space>(Ptr);
return multi_ptr<ElementType, Space, DecorateAddress>(CastPtr);
// TODO: Remove this restriction.
static_assert(std::is_same_v<ElementType, remove_decoration_t<ElementType>>,
"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<ElementType, Space, DecorateAddress>(Ptr);
} else {
auto CastPtr = sycl::detail::spirv::GenericCastToPtr<Space>(Ptr);
return multi_ptr<ElementType, Space, DecorateAddress>(CastPtr);
}
#else
return multi_ptr<ElementType, Space, DecorateAddress>(Ptr);
#endif
Expand All @@ -34,8 +45,15 @@ template <access::address_space Space, access::decorated DecorateAddress,
multi_ptr<ElementType, Space, DecorateAddress>
dynamic_address_cast(ElementType *Ptr) {
#ifdef __SYCL_DEVICE_ONLY__
auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit<Space>(Ptr);
return multi_ptr<ElementType, Space, DecorateAddress>(CastPtr);
// TODO: Remove this restriction.
static_assert(std::is_same_v<ElementType, remove_decoration_t<ElementType>>,
"The extension expect undecorated raw pointers only!");
if constexpr (Space == access::address_space::generic_space) {
return multi_ptr<ElementType, Space, DecorateAddress>(Ptr);
} else {
auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit<Space>(Ptr);
return multi_ptr<ElementType, Space, DecorateAddress>(CastPtr);
}
#else
return multi_ptr<ElementType, Space, DecorateAddress>(Ptr);
#endif
Expand Down
28 changes: 16 additions & 12 deletions sycl/test/check_device_code/extensions/address_cast.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<access::address_space::generic_space,
access::decorated::yes>(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<access::address_space::generic_space,
Expand All @@ -51,40 +53,42 @@ SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {

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-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-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: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META35:![0-9]+]]
// CHECK-NEXT: ret void
//
SYCL_EXTERNAL auto to_global_decorated(int *p) {
return dynamic_address_cast<access::address_space::global_space,
access::decorated::yes>(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) {
return dynamic_address_cast<access::address_space::global_space,
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<access::address_space::generic_space,
access::decorated::yes>(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<access::address_space::generic_space,
Expand Down

0 comments on commit 7d4af87

Please sign in to comment.