Skip to content

[SYCL] Fix launching a kernel with a default constructed local accessor #13382

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

Merged
merged 3 commits into from
Apr 17, 2024

Conversation

jzc
Copy link
Contributor

@jzc jzc commented Apr 12, 2024

According to 4.7.6.11.1. Interface for local accessors, passing a default constructed local accessor to a kernel function is allowed, as long as no accesses to the data elements are made.

When default constructing a local_accessor, it will default construct its local_accessor_base

local_accessor_base()
: detail::LocalAccessorBaseHost{/*Size*/ sycl::range<3>{0, 0, 0},
/*Dims*/ 0, /*ElemSize*/ sizeof(DataT)} {}

leading its LocalAccessorImplHost to be constructed with MDims == 0. When setting the kernel arguments corresponding to the local accessor, this caused piKernelSetArg to be called with arg_size == 0, which then causes an error. This PR fixes this by checking if MDims would be 0 and then setting an appropriate value.

@jzc jzc requested a review from a team as a code owner April 12, 2024 19:02
@jzc jzc requested a review from againull April 12, 2024 19:02
@jzc jzc temporarily deployed to WindowsCILock April 12, 2024 19:15 — with GitHub Actions Inactive
@jzc jzc temporarily deployed to WindowsCILock April 12, 2024 20:18 — with GitHub Actions Inactive
@jzc jzc temporarily deployed to WindowsCILock April 12, 2024 20:39 — with GitHub Actions Inactive
@@ -0,0 +1,18 @@
// RUN: %{build} -o %t.out -O0
// RUN: %{run} %t.out
Copy link
Contributor Author

Choose a reason for hiding this comment

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

-O0 is necessary; on higher levels of optimization, an error does not occur because of dead argument elimination of the local_accessor.

Copy link
Contributor

Choose a reason for hiding this comment

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

This comment should be moved to the test itself

@jzc jzc temporarily deployed to WindowsCILock April 16, 2024 18:58 — with GitHub Actions Inactive
@jzc jzc temporarily deployed to WindowsCILock April 16, 2024 20:56 — with GitHub Actions Inactive
@againull againull merged commit 80ecd7f into intel:sycl Apr 17, 2024
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.

3 participants