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.

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

Merged
merged 9 commits into from
Oct 23, 2023

Conversation

fineg74
Copy link
Contributor

@fineg74 fineg74 commented Sep 11, 2023

No description provided.

@fineg74 fineg74 requested a review from a team as a code owner September 11, 2023 23:42
@v-klochkov
Copy link
Contributor

@fineg74 - thank you for the fix. It look very good in general!
I have only few comments. Also, this fix needs conflicts resolution.

@fineg74 fineg74 temporarily deployed to WindowsCILock October 12, 2023 20:42 — with GitHub Actions Inactive
@v-klochkov
Copy link
Contributor

Looks good to me.
Because this patch touches several key memory functions, it would be good to have another look at it by @sarnex or @turinevgeny

@fineg74 fineg74 temporarily deployed to WindowsCILock October 12, 2023 21:04 — with GitHub Actions Inactive
@@ -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>
Copy link
Contributor

@sarnex sarnex Oct 13, 2023

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.

Copy link
Contributor Author

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>
Copy link
Contributor

Choose a reason for hiding this comment

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

Same Q as above

Copy link
Contributor

@v-klochkov v-klochkov Oct 13, 2023

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.

Copy link
Contributor

@v-klochkov v-klochkov Oct 13, 2023

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.

Copy link
Contributor Author

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

sycl/test/esimd/block_load_store.cpp Show resolved Hide resolved
@v-klochkov
Copy link
Contributor

v-klochkov commented Oct 13, 2023

@fineg74 - thank you for the fix. It look very good in general! I have only few comments. Also, this fix needs conflicts resolution.

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.
Such cases also may have some alignment issues. Perhaps, that can be fixed later in other PRs.

@fineg74
Copy link
Contributor Author

fineg74 commented Oct 14, 2023

@fineg74 - thank you for the fix. It look very good in general! I have only few comments. Also, this fix needs conflicts resolution.

For the record: The new checks makes 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. Such cases also may have some alignment issues. Perhaps, that can be fixed later in other PRs.

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.

@v-klochkov
Copy link
Contributor

@fineg74 - thank you for the fix. It look very good in general! I have only few comments. Also, this fix needs conflicts resolution.

For the record: The new checks makes 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. Such cases also may have some alignment issues. Perhaps, that can be fixed later in other PRs.

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
@fineg74 fineg74 temporarily deployed to WindowsCILock October 17, 2023 16:21 — with GitHub Actions Inactive
@fineg74 fineg74 temporarily deployed to WindowsCILock October 17, 2023 16:53 — with GitHub Actions Inactive
# Conflicts:
#	sycl/include/sycl/ext/intel/esimd/memory.hpp
# Conflicts:
#	sycl/include/sycl/ext/intel/esimd/memory.hpp
@fineg74 fineg74 temporarily deployed to WindowsCILock October 20, 2023 04:22 — with GitHub Actions Inactive
@fineg74 fineg74 temporarily deployed to WindowsCILock October 20, 2023 04:44 — with GitHub Actions Inactive
@againull againull closed this Oct 20, 2023
@againull againull reopened this Oct 20, 2023
@againull againull temporarily deployed to WindowsCILock October 20, 2023 22:28 — with GitHub Actions Inactive
@againull againull temporarily deployed to WindowsCILock October 20, 2023 22:56 — with GitHub Actions Inactive
@againull
Copy link
Contributor

Failed task was restarted and passed.

@againull againull merged commit 803a77f into intel:sycl Oct 23, 2023
21 of 22 checks passed
@fineg74 fineg74 deleted the accessorCheck1 branch December 2, 2023 01:59
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants