Skip to content

Commit db75b03

Browse files
aelovikov-intelaelovikov-intel
authored
[SYCL] Fix static|dynamic_address_cast to generic (#15394)
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.
1 parent a285d1d commit db75b03

File tree

2 files changed

+124
-4
lines changed

2 files changed

+124
-4
lines changed

‎sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp

Lines changed: 22 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -22,8 +22,19 @@ template <access::address_space Space, access::decorated DecorateAddress,
2222
multi_ptr<ElementType, Space, DecorateAddress>
2323
static_address_cast(ElementType *Ptr) {
2424
#ifdef __SYCL_DEVICE_ONLY__
25-
auto CastPtr = sycl::detail::spirv::GenericCastToPtr<Space>(Ptr);
26-
return multi_ptr<ElementType, Space, DecorateAddress>(CastPtr);
25+
// TODO: Remove this restriction.
26+
static_assert(std::is_same_v<ElementType, remove_decoration_t<ElementType>>,
27+
"The extension expect undecorated raw pointers only!");
28+
if constexpr (Space == access::address_space::generic_space) {
29+
// Undecorated raw pointer is in generic AS already, no extra casts needed.
30+
// Note for future, for `OpPtrCastToGeneric`, `Pointer` must point to one of
31+
// `Storage Classes` that doesn't include `Generic`, so this will have to
32+
// remain a special case even if the restriction above is lifted.
33+
return multi_ptr<ElementType, Space, DecorateAddress>(Ptr);
34+
} else {
35+
auto CastPtr = sycl::detail::spirv::GenericCastToPtr<Space>(Ptr);
36+
return multi_ptr<ElementType, Space, DecorateAddress>(CastPtr);
37+
}
2738
#else
2839
return multi_ptr<ElementType, Space, DecorateAddress>(Ptr);
2940
#endif
@@ -34,8 +45,15 @@ template <access::address_space Space, access::decorated DecorateAddress,
3445
multi_ptr<ElementType, Space, DecorateAddress>
3546
dynamic_address_cast(ElementType *Ptr) {
3647
#ifdef __SYCL_DEVICE_ONLY__
37-
auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit<Space>(Ptr);
38-
return multi_ptr<ElementType, Space, DecorateAddress>(CastPtr);
48+
// TODO: Remove this restriction.
49+
static_assert(std::is_same_v<ElementType, remove_decoration_t<ElementType>>,
50+
"The extension expect undecorated raw pointers only!");
51+
if constexpr (Space == access::address_space::generic_space) {
52+
return multi_ptr<ElementType, Space, DecorateAddress>(Ptr);
53+
} else {
54+
auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit<Space>(Ptr);
55+
return multi_ptr<ElementType, Space, DecorateAddress>(CastPtr);
56+
}
3957
#else
4058
return multi_ptr<ElementType, Space, DecorateAddress>(Ptr);
4159
#endif

‎sycl/test/check_device_code/extensions/address_cast.cpp

Lines changed: 102 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,102 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
2+
// RUN: %clangxx -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code -o - %s | FileCheck %s
3+
4+
// Linux/Windows have minor differences in the generated IR (e.g. TBAA
5+
// metadata). Having linux-only checks eases the maintenance without sacrifising
6+
// coverage of what's important for this test.
7+
// REQUIRES: linux
8+
9+
#include <sycl/sycl.hpp>
10+
11+
using namespace sycl;
12+
using namespace sycl::ext::oneapi::experimental;
13+
14+
namespace static_as_cast {
15+
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast19to_global_decoratedEPi(
16+
// 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]+]] {
17+
// CHECK-NEXT: [[ENTRY:.*:]]
18+
// 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]+]]
19+
// 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]+]]
20+
// CHECK-NEXT: ret void
21+
//
22+
SYCL_EXTERNAL auto to_global_decorated(int *p) {
23+
return static_address_cast<access::address_space::global_space,
24+
access::decorated::yes>(p);
25+
}
26+
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast23to_global_not_decoratedEPi(
27+
// 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]] {
28+
// CHECK-NEXT: [[ENTRY:.*:]]
29+
// 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]]
30+
// 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]+]]
31+
// CHECK-NEXT: ret void
32+
//
33+
SYCL_EXTERNAL auto to_global_not_decorated(int *p) {
34+
return static_address_cast<access::address_space::global_space,
35+
access::decorated::no>(p);
36+
}
37+
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast20to_generic_decoratedEPi(
38+
// 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]] {
39+
// CHECK-NEXT: [[ENTRY:.*:]]
40+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA23:![0-9]+]], !alias.scope [[META25:![0-9]+]]
41+
// CHECK-NEXT: ret void
42+
//
43+
SYCL_EXTERNAL auto to_generic_decorated(int *p) {
44+
return static_address_cast<access::address_space::generic_space,
45+
access::decorated::yes>(p);
46+
}
47+
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast24to_generic_not_decoratedEPi(
48+
// 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]] {
49+
// CHECK-NEXT: [[ENTRY:.*:]]
50+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA29:![0-9]+]], !alias.scope [[META31:![0-9]+]]
51+
// CHECK-NEXT: ret void
52+
//
53+
SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {
54+
return static_address_cast<access::address_space::generic_space,
55+
access::decorated::no>(p);
56+
}
57+
} // namespace static_as_cast
58+
59+
namespace dynamic_as_cast {
60+
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast19to_global_decoratedEPi(
61+
// 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]] {
62+
// CHECK-NEXT: [[ENTRY:.*:]]
63+
// 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]]
64+
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META35:![0-9]+]]
65+
// CHECK-NEXT: ret void
66+
//
67+
SYCL_EXTERNAL auto to_global_decorated(int *p) {
68+
return dynamic_address_cast<access::address_space::global_space,
69+
access::decorated::yes>(p);
70+
}
71+
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast23to_global_not_decoratedEPi(
72+
// 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]] {
73+
// CHECK-NEXT: [[ENTRY:.*:]]
74+
// 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]]
75+
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA17]], !alias.scope [[META39:![0-9]+]]
76+
// CHECK-NEXT: ret void
77+
//
78+
SYCL_EXTERNAL auto to_global_not_decorated(int *p) {
79+
return dynamic_address_cast<access::address_space::global_space,
80+
access::decorated::no>(p);
81+
}
82+
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast20to_generic_decoratedEPi(
83+
// 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]] {
84+
// CHECK-NEXT: [[ENTRY:.*:]]
85+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA23]], !alias.scope [[META43:![0-9]+]]
86+
// CHECK-NEXT: ret void
87+
//
88+
SYCL_EXTERNAL auto to_generic_decorated(int *p) {
89+
return dynamic_address_cast<access::address_space::generic_space,
90+
access::decorated::yes>(p);
91+
}
92+
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast24to_generic_not_decoratedEPi(
93+
// 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]] {
94+
// CHECK-NEXT: [[ENTRY:.*:]]
95+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA29]], !alias.scope [[META47:![0-9]+]]
96+
// CHECK-NEXT: ret void
97+
//
98+
SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {
99+
return dynamic_address_cast<access::address_space::generic_space,
100+
access::decorated::no>(p);
101+
}
102+
} // namespace dynamic_as_cast

0 commit comments

Comments
 (0)