Skip to content

Commit e8ebe3f

Browse files
aelovikov-intelaelovikov-intel
committed
[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`.
1 parent 628487a commit e8ebe3f

File tree

3 files changed

+63
-16
lines changed

3 files changed

+63
-16
lines changed

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

Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,9 @@ inline namespace _V1 {
1616
namespace ext {
1717
namespace oneapi {
1818
namespace experimental {
19+
namespace detail {
20+
using namespace sycl::detail;
21+
}
1922
// Shorthands for address space names
2023
constexpr inline access::address_space global_space = access::address_space::global_space;
2124
constexpr inline access::address_space local_space = access::address_space::local_space;
@@ -32,6 +35,18 @@ static_address_cast(ElementType *Ptr) {
3235
if constexpr (Space == generic_space) {
3336
// Undecorated raw pointer is in generic AS already, no extra casts needed.
3437
return ret_ty(Ptr);
38+
} else if constexpr (Space == access::address_space::
39+
ext_intel_global_device_space ||
40+
Space ==
41+
access::address_space::ext_intel_global_host_space) {
42+
#ifdef __ENABLE_USM_ADDR_SPACE__
43+
// No SPIR-V intrinsic for this yet.
44+
using raw_type = detail::DecoratedType<ElementType, Space>::type *;
45+
auto CastPtr = (raw_type)(Ptr);
46+
#else
47+
auto CastPtr = sycl::detail::spirv::GenericCastToPtr<global_space>(Ptr);
48+
#endif
49+
return ret_ty(CastPtr);
3550
} else {
3651
auto CastPtr = sycl::detail::spirv::GenericCastToPtr<Space>(Ptr);
3752
return ret_ty(CastPtr);
@@ -60,6 +75,20 @@ dynamic_address_cast(ElementType *Ptr) {
6075
"The extension expects undecorated raw pointers only!");
6176
if constexpr (Space == generic_space) {
6277
return ret_ty(Ptr);
78+
} else if constexpr (Space == access::address_space::
79+
ext_intel_global_device_space ||
80+
Space ==
81+
access::address_space::ext_intel_global_host_space) {
82+
#ifdef __ENABLE_USM_ADDR_SPACE__
83+
static_assert(
84+
Space != access::address_space::ext_intel_global_device_space &&
85+
Space != access::address_space::ext_intel_global_host_space,
86+
"Not supported yet!");
87+
return ret_ty(nullptr);
88+
#else
89+
auto CastPtr = sycl::detail::spirv::GenericCastToPtr<global_space>(Ptr);
90+
return ret_ty(CastPtr);
91+
#endif
6392
} else {
6493
auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit<Space>(Ptr);
6594
return ret_ty(CastPtr);

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

Lines changed: 20 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@ namespace static_as_cast {
2323
// CHECK-NEXT: [[ENTRY:.*:]]
2424
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8:![0-9]+]]
2525
// CHECK-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)
26-
// 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]+]]
26+
// 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]+]]
2727
// 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]+]]
2828
// CHECK-NEXT: ret void
2929
//
@@ -33,7 +33,7 @@ SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr<int> p) {
3333
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast23to_global_not_decoratedEPi(
3434
// 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]] {
3535
// CHECK-NEXT: [[ENTRY:.*:]]
36-
// 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]]
36+
// 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]]
3737
// 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]+]]
3838
// CHECK-NEXT: ret void
3939
//
@@ -61,18 +61,22 @@ SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {
6161
}
6262

6363
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast16to_global_deviceEPi(
64-
// 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]] {
64+
// 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]] {
6565
// CHECK-NEXT: [[ENTRY:.*:]]
66-
// CHECK-NEXT: unreachable
66+
// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(5)
67+
// CHECK-NEXT: store ptr addrspace(5) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA36:![0-9]+]], !alias.scope [[META38:![0-9]+]]
68+
// CHECK-NEXT: ret void
6769
//
6870
SYCL_EXTERNAL auto to_global_device(int *p) {
6971
return static_address_cast<access::address_space::ext_intel_global_device_space>(p);
7072
}
7173

7274
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast14to_global_hostEPi(
73-
// 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]] {
75+
// 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]] {
7476
// CHECK-NEXT: [[ENTRY:.*:]]
75-
// CHECK-NEXT: unreachable
77+
// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(6)
78+
// CHECK-NEXT: store ptr addrspace(6) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA42:![0-9]+]], !alias.scope [[META44:![0-9]+]]
79+
// CHECK-NEXT: ret void
7680
//
7781
SYCL_EXTERNAL auto to_global_host(int *p) {
7882
return static_address_cast<access::address_space::ext_intel_global_host_space>(p);
@@ -81,41 +85,41 @@ SYCL_EXTERNAL auto to_global_host(int *p) {
8185

8286
namespace dynamic_as_cast {
8387
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast19to_global_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
84-
// 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]] {
88+
// 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]] {
8589
// CHECK-NEXT: [[ENTRY:.*:]]
8690
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
8791
// CHECK-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)
88-
// 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]]
89-
// 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]+]]
92+
// 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]]
93+
// 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]+]]
9094
// CHECK-NEXT: ret void
9195
//
9296
SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr<int> p) {
9397
return dynamic_address_cast<access::address_space::global_space>(p);
9498
}
9599
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast23to_global_not_decoratedEPi(
96-
// 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]] {
100+
// 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]] {
97101
// CHECK-NEXT: [[ENTRY:.*:]]
98-
// 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]]
99-
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META44:![0-9]+]]
102+
// 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]]
103+
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META54:![0-9]+]]
100104
// CHECK-NEXT: ret void
101105
//
102106
SYCL_EXTERNAL auto to_global_not_decorated(int *p) {
103107
return dynamic_address_cast<access::address_space::global_space>(p);
104108
}
105109
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast20to_generic_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
106-
// 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]] {
110+
// 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]] {
107111
// CHECK-NEXT: [[ENTRY:.*:]]
108112
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
109-
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META48:![0-9]+]]
113+
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META58:![0-9]+]]
110114
// CHECK-NEXT: ret void
111115
//
112116
SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
113117
return dynamic_address_cast<access::address_space::generic_space>(p);
114118
}
115119
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast24to_generic_not_decoratedEPi(
116-
// 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]] {
120+
// 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]] {
117121
// CHECK-NEXT: [[ENTRY:.*:]]
118-
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META52:![0-9]+]]
122+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META62:![0-9]+]]
119123
// CHECK-NEXT: ret void
120124
//
121125
SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {

‎sycl/test/extensions/address_cast_negative.cpp‎

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
// RUN: %clangxx -D__ENABLE_USM_ADDR_SPACE__ -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=warning,note %s
2+
3+
#include <sycl/sycl.hpp>
4+
5+
using namespace sycl::ext::oneapi::experimental;
6+
7+
SYCL_EXTERNAL void test(int *p) {
8+
// expected-error-re@sycl/ext/oneapi/experimental/address_cast.hpp:* {{{{.*}}Not supported yet!}}
9+
std::ignore = dynamic_address_cast<
10+
sycl::access::address_space::ext_intel_global_device_space>(p);
11+
// expected-error-re@sycl/ext/oneapi/experimental/address_cast.hpp:* {{{{.*}}Not supported yet!}}
12+
std::ignore = dynamic_address_cast<
13+
sycl::access::address_space::ext_intel_global_host_space>(p);
14+
}

0 commit comments

Comments
 (0)