-
Notifications
You must be signed in to change notification settings - Fork 730
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
[SYCL][ESIMD] Add more stringent compile time checks for accessor versions of block_load/block_store, gather/scatter API #11145
Conversation
…_load, block_store, gather, scatter API
@fineg74 - thank you for the fix. It look very good in general! |
# Conflicts: # sycl/include/sycl/ext/intel/esimd/memory.hpp
Looks good to me. |
@@ -3049,8 +3052,7 @@ void simd_obj_impl<T, N, T1, SFINAE>::copy_from( | |||
|
|||
template <typename T, int N, class T1, class SFINAE> | |||
template <typename AccessorT, typename Flags, int ChunkSize, typename> | |||
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, | |||
sycl::access::target::device, void> | |||
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, void> |
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.
I think this allows local accessors now. We have
template <typename T, accessor_mode_cap_val_t Capability, typename RetT>
using EnableIfAccessor =
std::enable_if_t<detail::is_accessor_with_v<T, Capability>, RetT>;
and is_accessor_with_v
is
template <typename T, accessor_mode_cap_val_t Capability>
inline constexpr bool is_accessor_with_v =
is_device_accessor_with_v<T, Capability> ||
is_local_accessor_with_v<T, Capability>;
So won't local accessor be allowed because of the ||
? Or did I misunderstand? I assume we don't want to allow local accessors here.
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.
Fixed
@@ -3211,8 +3213,7 @@ void simd_obj_impl<T, N, T1, SFINAE>::copy_to( | |||
|
|||
template <typename T, int N, class T1, class SFINAE> | |||
template <typename AccessorT, typename Flags, int ChunkSize, typename> | |||
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, | |||
sycl::access::target::device, void> | |||
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, void> |
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.
Same Q as above
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.
Good catch.
This case was supposed to be supported in the series of PRs implementing passing local_accessor to accessor-based memory functions, but it is not.
We need variants of functions copy_from() & copy_to() accepting local accessor, but such variants need to accept uint32_t even with __ESIMD_FORCE_STATELESS_MEM.
This function here should not accept local accessor because it accepts uint64_t in stateless mode.
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.
copy_to() copy_from() need a separate overload accepting local accessor.
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.
Fixed. Introduced a separate version of copy_to/copy_from for local_accessors
For the record: The new checks make the code much safer. Before this PR memory functions could write to memory referenced by read-only accessor, which could cause any problems such as not updated buffers after kernel run, etc. The new checks do not yet verify the accessor element type. So, it is still possible write QWORDS to BYTE buffer ACCESSOR. |
This probably needs more discussion, as I am not sure if there are problems with different data types apart from alignment that would require enforcement of buffer type and accessor type being the same or probably the same size. Alignment is a separate story as technically BYTE buffer could be QWORD aligned by user which means we would need to introduce an additional parameter for the API to indicate buffer alignment. |
Yes, this definitely needs a special discussion. Accessor type-element checks are surely not for this PR. |
…n for local accessors to prevent passing 64 bit offsets to local accessors API under some circumstances
# Conflicts: # sycl/include/sycl/ext/intel/esimd/memory.hpp
# Conflicts: # sycl/include/sycl/ext/intel/esimd/memory.hpp
Failed task was restarted and passed. |
No description provided.