-
Notifications
You must be signed in to change notification settings - Fork 752
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] sycl_khr_free_function_commands extension #16770
base: sycl
Are you sure you want to change the base?
Conversation
Co-authored-by: Steffen Larsen <[email protected]>
Co-authored-by: Steffen Larsen <[email protected]>
template <typename KernelType, typename... Reductions> | ||
void launch_reduce(handler &h, range<1> r, const KernelType &k, | ||
Reductions &&...reductions) { | ||
h.parallel_for(r, std::forward<Reductions>(reductions)..., k); |
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.
Is there a reason we don't pass the KernelType
as a name here? Same applies to a couple other cases down the line.
h.parallel_for(r, std::forward<Reductions>(reductions)..., k); | |
h.parallel_for<KernelType>(r, std::forward<Reductions>(reductions)..., k); |
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.
It's a temporary fix because whenever we pass KernelType
as a name, there are compilation errors due to KernelType
being an invalid name. The compiler does not know how to handle it and treats the whole kernel as unnamed.
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 took a look at this. I believe there is a distinction to be made between the KernelName
passed to parallel_for
and the actual type of the lambda used to construct the kernel i.e KernelType
. Passing the latter in place of the former is very odd and certainly not the right way to do things. I don't think I can commit Stephen's suggestion because of this as its an issue with the test itself. Instead, KernelName
should be passed explicitly as a template argument to the function at the call site. It can be anything as long as its unique.
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.
@lbushi25 - The extension does not have kernel names as explicit parameters. I somewhat agree that the extension doesn't mandate using the type as the name, from what I can tell, but what would be wrong with passing a lambda type as a kernel name w.r.t. the requirements of the specification?
Note: Based on what I've been told, the failing kernel name is wrapping the lambda name in another type as part of the multi-kernel implementation of reductions. Even if we don't plan on using KernelType
as the name here, if it is valid to use lambda types as kernel names then this is still a bug in the reduction implementation.
Co-authored-by: Steffen Larsen <[email protected]>
This PR is implementation of
sycl_khr_free_function_commands
KHR extension.KhronosGroup/SYCL-Docs#644