[SYCL][Docs] Add kernel enqueue functions for kernel and properties

[steffenlarsen]

The current version of the sycl_ext_oneapi_kernel_properties extension does not have overloads for single_task and parallel_for that takes a sycl::kernel argument. This was likely omitted as all the properties added by the extension had direct effects on how the compiler would handle the kernel code, which cannot be done when the kernel argument is a sycl::kernel object. However, as more kernel property extension are written, some of them affect how the runtime handles the launch of kernels and could as such be applied to such kernel enqueue functions.

This commit adds member function overloads to sycl::handler for single_task and parallel_for that takes a sycl::kernel argument in the sycl_ext_oneapi_kernel_properties extension. Likewise, these new overloads are added to the implementation of sycl::handler.

[steffenlarsen]

Tag @AD2605 as this was discussed in #14113 (comment)

[gmlueck]

I would rather not expose this as a new API. I've been thinking instead that we should extend the meaning of is_property_key_of to allow arbitrary classification of properties. For example, we could define two types like this:

struct kernel_properties;
struct kernel_launch_properties;

We would then document each kernel property indicating which classification it has: kernel_property vs. kernel_launch_property. We could then define constraints for APIs that take a property list using something like is_property_key_of<Prop, kernel_launch_properties>. Application code could use this trait for its own constraints.

There is still some name bikeshedding to do. For example, are these classification types directly in sycl::ext::oneapi::experimental or are they in a sub-namespace? I think there should also be some sort of trait to tell if all properties in a list have a particular classification (e.g. is_property_list_key_of<PropertyListT, kernel_launch_properties>. This name could also use some bikeshedding.

I think there are a couple options for this PR. You could push the ideas I outline above forward into something more concrete, and then change this PR to use those new APIs.

Or, you could simply remove has_compile_time_kernel_effect from this PR and and use less precise language about the constraints on the single_task. etc. functions below. For example, the constraint can just say that the properties must be launch properties without defining a specific trait. This would allow us to make the language more precise later once we refine the property classification traits. In this case, the implementation could still define traits in the detail namespace, but they would be undocumented.

[steffenlarsen]

I went with the latter to buy us time so we can settle on a good solution here.

[gmlueck]

I also wonder if we should invest effort into extending these forms of single_task and parallel_for. The recently implemented sycl_ext_oneapi_enqueue_functions has a cleaner separation between launch properties vs. kernel properties.

I think our longer term strategy may be to drop the forms of single_task and parallel_for in sycl_ext_oneapi_kernel_properties and use the ones in sycl_ext_oneapi_enqueue_functions instead.

@Pennycook what were your thoughts on these two extension regarding ways of specifying kernel properties?

[Pennycook]

I think our longer term strategy may be to drop the forms of single_task and parallel_for in sycl_ext_oneapi_kernel_properties and use the ones in sycl_ext_oneapi_enqueue_functions instead.

I agree with this.

sycl_ext_oneapi_kernel_properties was essentially our first attempt at defining a kernel launch interface that accepted properties. The sycl_ext_oneapi_enqueue_functions design addresses a lot of feedback we received from users and other implementers, and it's much more aligned with where we expect future versions of SYCL to go.

My preference would be that we move to sycl_ext_oneapi_enqueue_functions as soon as possible, and deprecate these property overloads. If there's a short-term need to expose these overloads with minimal effort, I'm not opposed, but we should figure out an implementation plan for the new extension.

[gmlueck]

If there's a short-term need to expose these overloads with minimal effort, I'm not opposed, but we should figure out an implementation plan for the new extension.

I think sycl_ext_oneapi_enqueue_functions is already implemented.

@steffenlarsen I'm not sure who is asking for this change. Can we ask them to use sycl_ext_oneapi_enqueue_functions instead of making the changes in this PR?

[steffenlarsen]

Discussed #14707 (comment) offline and we agreed that this is the way we want to go about it. The implementation of the free functions will need additional work to allow this, which I will repurpose these changes for.