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] Use static address space cast for atomic_ref ctor in SPIR-V path #15384

Merged
merged 11 commits into from
Sep 18, 2024
Merged

[SYCL] Use static address space cast for atomic_ref ctor in SPIR-V path #15384

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
0e5d540
[SYCL] Add a test for generated device code for `atomic_ref`
aelovikov-intel Sep 12, 2024
a22d4bd
[SYCL] Use `static` address space cast for `atomic_ref` ctor in SPIR-…
aelovikov-intel Sep 12, 2024
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
b30b1c4
Merge remote-tracking branch 'my/atomic-ref-static-as-cast' into HEAD
aelovikov-intel Sep 13, 2024
d40110d
clang-format
aelovikov-intel Sep 13, 2024
f052b05
Disable newly added test on Windows
aelovikov-intel Sep 13, 2024
78cde19
Merge remote-tracking branch 'my/fix-casts' into HEAD
aelovikov-intel Sep 13, 2024
fd7848b
Merge remote-tracking branch 'origin/sycl' into HEAD
aelovikov-intel Sep 16, 2024
d62ab6d
Merge remote-tracking branch 'origin/sycl' into HEAD
aelovikov-intel Sep 17, 2024
56e8f2f
Update the fix to the recent extension change
aelovikov-intel Sep 17, 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
11 changes: 10 additions & 1 deletion sycl/include/sycl/atomic_ref.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,8 @@

#include <sycl/access/access.hpp> // for address_space
#include <sycl/bit_cast.hpp> // for bit_cast
#include <sycl/memory_enums.hpp> // for getStdMemoryOrder, memory_order
#include <sycl/ext/oneapi/experimental/address_cast.hpp>
#include <sycl/memory_enums.hpp> // for getStdMemoryOrder, memory_order

#ifdef __SYCL_DEVICE_ONLY__
#include <sycl/detail/spirv.hpp>
Expand Down Expand Up @@ -157,8 +158,16 @@ class atomic_ref_base {
}

#ifdef __SYCL_DEVICE_ONLY__
#if defined(__SPIR__)
Copy link
Contributor

Choose a reason for hiding this comment

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

Why is feature test macro SYCL_EXT_ONEAPI_ADDRESS_CAST not defined for the __SPIR__ case?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Extension implementation seems to have missed that completely. I'm making other changes to the extension, will work on that outside this PR.

explicit atomic_ref_base(T &ref)
: ptr(ext::oneapi::experimental::static_address_cast<AddressSpace>(
&ref)) {}
#else
// CUDA/HIP don't support `ext::oneapi::experimental::static_address_cast`
// yet.
explicit atomic_ref_base(T &ref)
: ptr(address_space_cast<AddressSpace, access::decorated::no>(&ref)) {}
#endif
#else
// FIXME: This reinterpret_cast is UB, but happens to work for now
explicit atomic_ref_base(T &ref)
Expand Down
20 changes: 20 additions & 0 deletions sycl/test/check_device_code/atomic_ref.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,20 @@
// 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

#include <sycl/sycl.hpp>

// CHECK-LABEL: define dso_local spir_func noundef i32 @_Z17atomic_ref_globalRi(
// CHECK-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) [[I:%.*]]) 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_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[I]], i32 noundef 5) #[[ATTR3:[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.

One can review individual commits in this PR to see before/after change.

// CHECK-NEXT: [[CALL3_I_I:%.*]] = tail call spir_func noundef i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(ptr addrspace(1) noundef [[CALL_I_I_I_I_I_I]], i32 noundef 1, i32 noundef 898) #[[ATTR4:[0-9]+]]
// CHECK-NEXT: ret i32 [[CALL3_I_I]]
//
SYCL_EXTERNAL auto atomic_ref_global(int &i) {
// Verify that we use _Z33__spirv_GenericCastToPtr_ToGlobalPvi that doesn't
// perform dynamic address space validation.
sycl::atomic_ref<int, sycl::memory_order::acq_rel, sycl::memory_scope::device,
sycl::access::address_space::global_space>
a(i);
return a.load();
}
Loading