[SYCL] Use static address space cast for atomic_ref ctor in SPIR-V path
#15384
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
…V path From SYCL 2020 specification: > The sycl::atomic_ref class also has a template parameter AddressSpace, > which allows the application to make an assertion about the address > space of the object of type T that it references. The default value > for this parameter is access::address_space::generic_space, which > indicates that the object could be in either the global or local > address spaces. If the application knows the address space, it can set > this template parameter to either access::address_space::global_space > or access::address_space::local_space as an assertion to the > implementation. Specifying the address space via this template > parameter may allow the implementation to perform certain > optimizations. Specifying an address space that does not match the > object’s actual address space results in undefined behavior We use `ext::oneapi::experimental::static_address_cast` to do that. It's not implemented for CUDA/HIP yet, that path continues using `sycl::address_space_cast` that performs runtime checks: > An implementation must return nullptr if the run-time value of pointer > is not compatible with Space, and must issue a compiletime diagnostic > if the deduced address space for pointer is not compatible with Space.
aelovikov-intel
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
aelovikov-intel
commented
Sep 12, 2024
| // 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]+]] |
There was a problem hiding this comment.
One can review individual commits in this PR to see before/after change.
aelovikov-intel
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
aelovikov-intel
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
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.
aelovikov-intel
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
|
@aelovikov-intel I see you marked this as draft. Is this good to review or is it still work in progress? |
Ready now. I had to wait until #15394 is merged and this is updated to include it via |
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
rolandschulz
approved these changes
Sep 17, 2024
| @@ -157,8 +158,16 @@ class atomic_ref_base { | |||
| } | |||
|
|
|||
| #ifdef __SYCL_DEVICE_ONLY__ | |||
| #if defined(__SPIR__) | |||
There was a problem hiding this comment.
Why is feature test macro SYCL_EXT_ONEAPI_ADDRESS_CAST not defined for the __SPIR__ case?
There was a problem hiding this comment.
Extension implementation seems to have missed that completely. I'm making other changes to the extension, will work on that outside this PR.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
From SYCL 2020 specification:
We use
ext::oneapi::experimental::static_address_castto do that. It'snot implemented for CUDA/HIP yet, that path continues using
sycl::address_space_castthat performs runtime checks: