Skip to content
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

Jump to bottom

[SYCL] Fix static|dynamic_address_cast to generic #15394

Merged
merged 3 commits into from
Sep 16, 2024
Merged

[SYCL] Fix static|dynamic_address_cast to generic #15394

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
21c3f33
[SYCL] Add a test showing issues with sycl_ext_oneapi_address_cast impl
aelovikov-intel Sep 13, 2024
7d4af87
[SYCL] Fix `static|dynamic_address_cast` to `generic`
aelovikov-intel Sep 13, 2024
f052b05
Disable newly added test on Windows
aelovikov-intel Sep 13, 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
26 changes: 22 additions & 4 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 @@ -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!");
Comment on lines +25 to +27
Copy link
Contributor

Choose a reason for hiding this comment

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

Out of curiosity: is this something that would be difficult to fix? Why not just use remove_decoration_t to create an undecorated raw pointer here, and then use the undecorated pointer for the rest of the function?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I plan to do that in the next PR(s). It has to come with the doc change as well.

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
102 changes: 102 additions & 0 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
@@ -0,0 +1,102 @@
// 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 <sycl/sycl.hpp>

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<access::address_space::global_space,
access::decorated::yes>(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<access::address_space::global_space,
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 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: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA23:![0-9]+]], !alias.scope [[META25:![0-9]+]]
Copy link
Contributor Author

Choose a reason for hiding this comment

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

Use per-commit diff to see how this PR affects code generation here. First commit is test-only without functional changes.

// 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 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: 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,
access::decorated::no>(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 [[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 [[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 [[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 [[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 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: 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 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: 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,
access::decorated::no>(p);
}
} // namespace dynamic_as_cast
Loading