[SYCL] Fix is_device_copyable with range rounding #4478
Merged
Conversation
This file contains hidden or 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
This patch precommits the test to expose the issue.
Suppose we have the following code:
```
template <> struct is_device_copyable<DeviceCopyable>
: std::true_type {};
...
DeviceCopyable DC(0);
Q.submit([=](sycl::handler& cgh){
const sycl::range<2> range(1026, 1026);
cgh.parallel_for(range,[=](sycl::item<2> item) {
(void)DC;
});
});
```
This code doesn't compile because range rounding optimization
wraps kernel lambda function, so we have something like this:
```
|- WrapperLambda
|-KernelLambda
||-DC
|-Other Wrapper captures
```
According to the implementation of is_device_copyable,
we check whether all the fields of WrapperLambda are device
copyable. KernelLambda is not device copyable since there is
no corresponding template specialization of is_device_copyable.
It's not possible to provide a template specialization for is_device_copyable with kernel lambda since lambda types don't have the name. To fix this issue, we create a functor class RoundedRangeKernel and provide a is_device_copyable trait for this class, which simply forwards check to nested kernel lambda.
|
Test that exposes the problem in original code is in the first commit in this PR. |
v-klochkov
reviewed
Sep 8, 2021
There was a problem hiding this comment.
Looks good overall. I have 2 questions.
@erichkeane - would you also take a quick look at this fix?
v-klochkov
reviewed
Sep 8, 2021
v-klochkov
reviewed
Sep 9, 2021
v-klochkov
approved these changes
Sep 9, 2021
|
@turinevgeny - please review. This PR requires your approval. |
turinevgeny
approved these changes
Sep 10, 2021
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.
Suppose we have the following code:
This code doesn't compile because range rounding optimization
wraps kernel lambda function, so we have something like this:
According to the implementation of is_device_copyable,
we check whether all the fields of WrapperLambda are device
copyable. KernelLambda is not device copyable since there is
no corresponding template specialization of is_device_copyable.
It's not possible to provide a template specialization for
is_device_copyable with kernel lambda since lambda types don't
have the name. To fix this issue, we create a functor class
RoundedRangeKernel and provide an is_device_copyable trait for this
class, which simply forwards check to nested kernel lambda.