Skip to content

[SYCL] Support *global_[device|host]_space in static_address_cast #15498

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Sep 27, 2024
Merged

[SYCL] Support *global_[device|host]_space in static_address_cast #15498

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
628487a
[SYCL] Add a test showing miscompilation
aelovikov-intel Sep 24, 2024
e8ebe3f
[SYCL] Support `*global_[device|host]_space` in `static_address_cast`
aelovikov-intel Sep 24, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
29 changes: 29 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,9 @@ inline namespace _V1 {
namespace ext {
namespace oneapi {
namespace experimental {
namespace detail {
using namespace sycl::detail;
}
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think it's a bad habit of ours to have any other detail namespace than sycl::detail. Any details can be in sycl::detail, even if it's extension related. That said, not this PR's fault and I have added to those namespaces in the past. Just a bit of venting. 😉

// 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;
Expand All @@ -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<ElementType, Space>::type *;
auto CastPtr = (raw_type)(Ptr);
#else
auto CastPtr = sycl::detail::spirv::GenericCastToPtr<global_space>(Ptr);
#endif
return ret_ty(CastPtr);
} else {
auto CastPtr = sycl::detail::spirv::GenericCastToPtr<Space>(Ptr);
return ret_ty(CastPtr);
Expand Down Expand Up @@ -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<global_space>(Ptr);
return ret_ty(CastPtr);
#endif
} else {
auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit<Space>(Ptr);
return ret_ty(CastPtr);
Expand Down
40 changes: 31 additions & 9 deletions sycl/test/check_device_code/extensions/address_cast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
// 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
// RUN: %clangxx -D__ENABLE_USM_ADDR_SPACE__ -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
Expand Down Expand Up @@ -59,45 +59,67 @@ SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {
return static_address_cast<access::address_space::generic_space>(p);
}

// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast16to_global_deviceEPi(
// 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: [[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<access::address_space::ext_intel_global_device_space>(p);
}

// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast14to_global_hostEPi(
// 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: [[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<access::address_space::ext_intel_global_host_space>(p);
}
} // namespace static_as_cast

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 [[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") 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) #[[ATTR5]]
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12]], !alias.scope [[META36:![0-9]+]]
// 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<int> p) {
return dynamic_address_cast<access::address_space::global_space>(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 [[META41:![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) #[[ATTR5]]
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META42:![0-9]+]]
// 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<access::address_space::global_space>(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 [[META45:![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 [[META46:![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<int> p) {
return dynamic_address_cast<access::address_space::generic_space>(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 [[META49:![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 [[META50:![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) {
Expand Down
14 changes: 14 additions & 0 deletions sycl/test/extensions/address_cast_negative.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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 <sycl/sycl.hpp>

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);
}
Loading