-
Notifications
You must be signed in to change notification settings - Fork 795
[SYCL][ESIMD] Add support for local accessors to lsc API #10340
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.
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
Conversation
sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp
Outdated
Show resolved
Hide resolved
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Mostly LGTM, just a couple of questions.
!std::is_same_v<Toffset, uint64_t>, | ||
__ESIMD_NS::simd<T, N * NElts>> | ||
__ESIMD_API std::enable_if_t< | ||
!std::is_pointer_v<AccessorTy> && |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Just a note here:
!std::is_pointer_v is a bad check (not because of your PR). It allows using write-only accessors lsc_gather() operations. Same/vice-versa is for read-only accessors usable in write operations in other places in memory.hpp here.
I created internal tracker for this issue.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
The much better check is used for copy_to() and copy_from(), e.g. copy_from():
detail::is_sycl_accessor_with<
AccessorT, accessor_mode_cap::can_read,
sycl::access::target::device>::value)
No description provided.