[SYCL][Doc] Add free function kernel extension spec

[gmlueck]

Add a proposed extension spec that allows a kernel to be defined as a plain function, where the function arguments are the kernel arguments. This is useful for the sycl_ext_oneapi_kernel_compiler because it allows a kernel to be defined in isolation from its host code.

[rolandschulz]

Why? I would prefer if the item/nd_item is the first argument followed by the other function parameters.
This would also remove the need for the range/nd_range/single_task_kernel property. Because this could be inferred from the first argument.

[Pennycook]

Wouldn't we still need to decorate the function somehow, though? I don't think we want to turn every function which accepts a sycl::nd_item into a kernel -- it might just be a device function that's called from a kernel.

[rolandschulz]

I agree we can't use the first argument to determine whether it is a kernel. Only the type of kernel (nd/range/single_task). In the same TU we can turn every function which gets passed to get_kernel into a kernel. Cross TU we do need a decoration. But then the user only needs to provide the decoration cross TU (like SYCL_EXTERNAL) and the FE doesn't need to understand multiple decorations.

If we support kernel function pointers cross TU, we should also support callable types cross TU. And those would need the same decoration.

[bader]

Can we avoid UB by having just one macro with no parameters (e.g. SYCL_EXT_ONEAPI_KERNEL)?
The spec doesn't say why we need 3 different macros, which must match kernel invocation function to avoid UB.

[gmlueck]

I was thinking that by having three different macros we can provide better error messages. For example, we can emit a compile-time error if a kernel decorated with SYCL_EXT_ONEAPI_RANGE_KERNEL calls the function sycl::ext::oneapi::this_kernel::get_nd_item. We can also throw an exception if a kernel decorated with SYCL_EXT_ONEAPI_RANGE_KERNEL is enqueued as an nd-range kernel.

Note that we can diagnose these errors even if the spec says the behavior is UB. We always have the option to diagnose UB with a helpful error.

[bader]

For example, we can emit a compile-time error if a kernel decorated with SYCL_EXT_ONEAPI_RANGE_KERNEL calls the function sycl::ext::oneapi::this_kernel::get_nd_item.

Is this a situation where we need to emit an error? I would expect any SYCL kernel to work for any combination of kernel invocation API and kernel identifiers free functions. nd_item is a wrapper around kernel identifiers, which are available to any kernel regardless of parallel_for overload. Current SYCL API limitations on lambda call operator parameter type are artificial. They are useful for code optimization: compiler has to allocate and initialize id/item/nd_item lambda argument. In this case, it's created by user via free function, so we don't need to optimize it.

[gmlueck]

I do think an error would be helpful. The specification for the free function queries says their behavior is undefined if they are called from a kernel that is launched in an incompatible way:

It is the user’s responsibility to ensure that these functions are called in a manner that is compatible with the kernel’s launch parameters, as detailed in the definition of each function. Calling these functions from an incompatible kernel results in undefined behavior.

(https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_queries.asciidoc#accessing-instances-of-special-sycl-classes)

If we don't say it is UB, then we would need to clarify the spec to explain what is returned in the mismatched cases. For example, what is the work-group size for a simple "range" parallel-for? This seems worse than saying the behavior is UB and diagnosing an error when possible.

[bader]

If we don't say it is UB, then we would need to clarify the spec to explain what is returned in the mismatched cases.

There are no mismatches cases because there is no SYCL kernel argument we have to "match" (or to be compatible with). This argument exists only for call operator of functional object, but this extension defines alternative API, which we don't need to "match".

For example, what is the work-group size for a simple "range" parallel-for?

I expect it to be defined by the implementation. AFAIK, all supported backends provide access to work-group size on device side.
On the host side, when not set explicitly OpenCL back-end chooses the work-group size automatically. Other backends has API to query recommended work-group size, which can be set explicitly.

[gmlueck]

There are no mismatches cases because there is no SYCL kernel argument we have to "match"

By "mismatched" I meant a mismatch between the way the kernel is launched and the queries the kernel makes. For example, launching the kernel as a simple "range" kernel and calling this_kernel::get_nd_item.

I expect it to be defined by the implementation. AFAIK, all supported backends provide access to work-group size on device side.

I really feel this would open a can of worms that would make our life difficult in the future. For example, what if the application launches the kernel as a simple "range" and then uses this_kernel::get_nd_item to get a group object that it uses to call one of the group functions defined in section 4.17.2 or 4.17.3? I really don't think we want to say this is supported. If we were required to support this, I think it would limit our implementation choices in the future.

I also think this presents a confusing story to our users. @Pennycook has been clarifying the SYCL spec to say that simple "range" kernels have no concept of work-group. Therefore, it does not make sense to allow "range" kernels to call this_kernel::get_nd_item.

[bader]

Probably, I don't see much value in "range" kernels. According to my understanding, this mode limits access to low-level HW features for two reasons:

1. Simpler interface.
2. In some cases, it might be more efficient to create item object instead of nd_item.

Free functions remove the second item.
Regarding the first one, this extension makes it even more difficult to use. This proposal requires users to make use of 3 different entities consistent: kernel function declaration, kernel function invocation and index space identifiers. Current spec requires only two and I've seen cases of inconsistent usage.

From my POV, "range" kernel benefits do not justify adding more UB and usage (and implementation) complexity of this proposal.

All low-level APIs we support today have work-group concept, so all limitations enforced by "range" (or even "single") kernels are artificial. We can always create valid meaningful nd_item regardless of kernel invocation API. We add compiler and runtime library support to diagnose inconsistent API uses for very low ROI.

[rdeodhar]

Regarding "decomposition". Shortly after I added support for array kernel parameters a couple of years ago the clang front-end team decided to decompose all non-POD kernel arguments. That includes all structs and arrays passed by value. SYCL accessors are just an instance of a struct type. So any decision made for accessors is really a decision about all non-POD types.

[gmlueck]

Regarding "decomposition". Shortly after I added support for array kernel parameters a couple of years ago the clang front-end team decided to decompose all non-POD kernel arguments. That includes all structs and arrays passed by value. SYCL accessors are just an instance of a struct type. So any decision made for accessors is really a decision about all non-POD types.

Are you sure? I just tried an example with an argument defined like this:

struct pair {
  int i;
  float f;
  int *p;
};

And it seems like it was not decomposed.

In fact, I remember @elizabethandrews telling me that she changed the front-end to stop decomposing struct arguments when a member is a pointer. At this point, I thought we only decomposed a kernel argument if it is a "special" type (e.g. accessor or stream) or if it is a struct with a member that is a "special" type.

[rdeodhar]

My understanding of decomposition was outdated. Yes, it looks like only certain SYCL types are decomposed.

[rdeodhar]

What was the thinking behind the macros SYCL_EXT_ONEAPI_RANGE_KERNEL, SYCL_EXT_ONEAPI_ND_RANGE_KERNEL, or SYCL_EXT_ONEAPI_SINGLE_TASK_KERNEL and the kernel function traits? The macros have somehow to be implemented to declare the traits?

[rdeodhar]

Ok, using properties looks better than the macros.

[gmlueck]

Attention reviewers: I do not plan any more changes to this PR and will ask for it to be merged soon. If you have more comments, please make them soon.

[rdeodhar]

The macro SYCL_EXT_ONEAPI_FUNCTION_PROPERTY is defined in properties.hpp as adding some properties to the sycl-device-has attribute. Is that how the properties range_kernel, etc, are to be found in clang front-end processing? Through the sycl-device-has attribute??

[gmlueck]

The macro SYCL_EXT_ONEAPI_FUNCTION_PROPERTY is defined in properties.hpp as adding some properties to the sycl-device-has attribute. Is that how the properties range_kernel, etc, are to be found in clang front-end processing? Through the sycl-device-has attribute??

No, the SYCL_EXT_ONEAPI_FUNCTION_PROPERTY macro is not specific to sycl::device_has. This macro is intended to be a general way to attach properties to a function. The example here illustrates the macro for sycl::device_has, but that is not the only use.

[rdeodhar]

I find it odd that the macro SYCL_EXT_ONEAPI_FUNCTION_PROPERTY is defined in a header file with no indication that it is to be treated as an example. And then this PR specifically mentions use of that macro to define the new kernel properties range_kernel, nd_range_kernel and single_task_kernel, when simply using that macro does not achieve the desired effect.

[gmlueck]

I find it odd that the macro SYCL_EXT_ONEAPI_FUNCTION_PROPERTY is defined in a header file with no indication that it is to be treated as an example. And then this PR specifically mentions use of that macro to define the new kernel properties range_kernel, nd_range_kernel and single_task_kernel, when simply using that macro does not achieve the desired effect.

Sorry, I forgot to include a hyperlink in my previous response. I meant to say:

The example here illustrates the macro for sycl::device_has, but that is not the only use.

(The word "here" was supposed to be a link to the example I was talking about.)

You can see that the macro is described to be a general way to attach properties to device functions:

This extension provides a mechanism exposing similar behavior, allowing for kernel properties to be associated with a function via the SYCL_EXT_ONEAPI_FUNCTION_PROPERTY macro. Each instance of the SYCL_EXT_ONEAPI_FUNCTION_PROPERTY macro accepts one argument, corresponding to a single property value.

I'm not sure why the header file hard-codes the macro to work only for sycl::device_has. This is either a bug, or someone was intentionally delaying the extra work for generalizing the macro for a later time. I presume that at the point when that header was implemented, the only property that could be used with that macro was sycl::device_has, but that will no longer be the case once the extension in this PR is implemented.

[AerialMantis]

This extension looks great, I've left some comments, mostly clarifications.

Though, when reading through this I thought about a potentially different approach to declaring free function kernels (perhaps this has already been considered), whereby a free function kernel is identified instead by the first parameter to the function. We could introduce a new class template such as sycl::this_kernel which would be templated on the invocation type (range_kernel, nd_range_kernel or single_task_kernel) and compile-time properties.

class single_task_kernel;

template <size_t Dims>
class range_kernel;

template <size_t Dims>
class nd_range_kernel;

template <typename KernelInvocation , typename... Properties>
class this_kernel;

So a free function kernel would be declared as follows:

void iota(syclex::this_kernel<syclex::range_kernel<1>, syclex::work_group_size<32>> this,
               float start, float *ptr) {
  // Get the ID of this kernel iteration.
  size_t id = this.get_id();

  ptr[id] = start + static_cast<float>(id);
}

I believe this could be implemented by a device compiler by having a pass which detects the sycl::this_kernel parameter and transforms the function into a SYCL kernel function as is does for parallel_for etc, and it could be implemented in a library by detecting the first parameter of the function to verify that it's a SYCL kernel function before calling it.

This approach could provide some benefits:

• Removes the need to use a macro for declaring free function kernels.
• Adds the invocation type to the type system, avoiding undefined behaviour.
• Allows a free function kernel to be overloaded by the invocation type, which would improve generic programming.
• Allows compile-time properties on free function kernels to be applied generically.
• Avoids being dependent on the free function queries extension to retrieve the kernel id object.

[gmlueck]

Hi @AerialMantis, Thanks for the review! I've either addressed or responded to your line-by-line comments below.

Regarding your larger comment:

I believe this could be implemented by a device compiler by having a pass which detects the sycl::this_kernel parameter and transforms the function into a SYCL kernel function as is does for parallel_for etc,

This is an interesting idea, but I think the front-end team will not like this direction, and I think it would be hard to get upstreamed. The front-end does not want to get in the business of recognizing specific type names, and I think it would have to do this in order to identify the kernel via the syclex::this_kernel parameter. Note that the front-end does not currently recognize the parallel_for by name. Instead, the header wraps the lambda with a function that is decorated with a special attribute, and the front-end recognizes the attribute. The mechanism is similar for SYCL_EXT_ONEAPI_FUNCTION_PROPERTY. That macro adds a property to the function, which the header translates into a C++ attribute, and the front-end recognizes the attribute.

I agree that your approach has some benefits, but I don't agree with all of them:

• Removes the need to use a macro for declaring free function kernels.

Yes, this is nice.

• Adds the invocation type to the type system, avoiding undefined behaviour.

I think your approach is not type-safe, so there would still be undefined behavior if it is misused. The fundamental problem is that the kernel_id and kernel objects do not convey the type of the kernel or the types of the kernel's arguments. Therefore, the compiler cannot catch errors when the kernel is launched (if the application launches it the wrong way or passes invalid argument types).

This is partly by design. One of our customers likes the fact that any kernel can be represented by a generic kernel object. This allows them to use generic code to manage a library of kernels along with some metadata that provides the information they need to invoke each one.

FWIW, @Pennycook's proposed free-function sycl::parallel_for would be a type-safe way to invoke a free-function kernel. See the first issue at the bottom of the spec. This will be type-safe even with the current proposal that uses SYCL_EXT_ONEAPI_FUNCTION_PROPERTY.

• Allows a free function kernel to be overloaded by the invocation type, which would improve generic programming.
• Allows compile-time properties on free function kernels to be applied generically.

I agree that your proposal is more convenient for generic programming, but I think you can also do it with the current proposal. Free function kernels can be templated, so you could use different template instantiations that correspond to different invocation types or different properties.

• Avoids being dependent on the free function queries extension to retrieve the kernel id object.

I actually don't see this as an advantage. We plan to implement that sycl_ext_oneapi_free_function_queries extension anyways, so I think it's an advantage to using it rather than inventing a new thing.

[AerialMantis]

This is an interesting idea, but I think the front-end team will not like this direction, and I think it would be hard to get upstreamed. The front-end does not want to get in the business of recognizing specific type names, and I think it would have to do this in order to identify the kernel via the syclex::this_kernel parameter. Note that the front-end does not currently recognize the parallel_for by name. Instead, the header wraps the lambda with a function that is decorated with a special attribute, and the front-end recognizes the attribute. The mechanism is similar for SYCL_EXT_ONEAPI_FUNCTION_PROPERTY. That macro adds a property to the function, which the header translates into a C++ attribute, and the front-end recognizes the attribute.

So I don't think it would be necessary for the front-end to identify the sycl::this_kernel by name. The way I would see this working is that the sycl::this_kernel class would have an attribute attached to it, similar to the underlying function called by functions like parallel_for. Then we could have a function pass which would look at non-member and static member functions (and specialized templates) for a parameter with that attribute. Then from there the function would be transformed into a SYCL kernel much like the function called by parallel_for, etc. For the compile-time properties attached to the sycl::this_kernel a similar approach could be taken by which after finding a parameter with the above attribute would then look for compile-time types in the template declaration with further attributes which represent the compile-time properties.

I agree this would be a more involved implementation than the macro approach as the macro can expand the attributes directly to the function declaration, whereas with sycl::this_kernel a function pass would need to inspect the functions.

I think your approach is not type-safe, so there would still be undefined behavior if it is misused. The fundamental problem is that the kernel_id and kernel objects do not convey the type of the kernel or the types of the kernel's arguments. Therefore, the compiler cannot catch errors when the kernel is launched (if the application launches it the wrong way or passes invalid argument types).

Yes, that is true, while there would be compile-time information about the function, this wouldn't translate to the kernel_id type so any information would still need to be exposed at runtime.

This is partly by design. One of our customers likes the fact that any kernel can be represented by a generic kernel object. This allows them to use generic code to manage a library of kernels along with some metadata that provides the information they need to invoke each one.

I agree there is a strong benefit to this, so I wouldn't want to propose that we start adding type information to the kernel_id or kernel classes.

FWIW, @Pennycook's proposed free-function sycl::parallel_for would be a type-safe way to invoke a free-function kernel. See the first issue at the bottom of the spec. This will be type-safe even with the current proposal that uses SYCL_EXT_ONEAPI_FUNCTION_PROPERTY.

Ah okay that's interesting, I'll need to revisit that extension.

I agree that your proposal is more convenient for generic programming, but I think you can also do it with the current proposal. Free function kernels can be templated, so you could use different template instantiations that correspond to different invocation types or different properties.

I think you're right, you could have a template function with some template parameters which describe the invocation method, and then explicitly instantiate them with a different macro. and therefore compile-time property for each.

Though thinking about that, this raises another question, when we template a free function kernel is the macro applied to the template declaration (I think this is how it's described in the example), or on the specializations (explicit instantiation or explicit specialization), or both? I wonder how it would work if it was applied to the template declaration because then the template arguments wouldn't be known until provided when retrieving the kernel bundle (something like get_kernel_bundle<iota<float>>), would this need to trigger deferred template instantiation in the front-end to obtain the specialized function definition in order to transform that to a SYCL kernel?

I actually don't see this as an advantage. We plan to implement that sycl_ext_oneapi_free_function_queries extension anyways, so I think it's an advantage to using it rather than inventing a new thing.

Yeah I was hesitant to state this as a benefit, as we may want to go in the direction of the free function queries, but I'm also wary that the free function queries may be different to implement in a library only or host implementation of SYCL.

[gmlueck]

Regarding the comment thread above with @AerialMantis, we had an offline meeting which also included @Pennycook about this. We decided that it did not make sense to use sycl::this_kernel as an general way to attach properties to device functions. We'd prefer to use the macro SYCL_EXT_ONEAPI_FUNCTION_PROPERTY as the only way to attach properties to functions. However, we did discuss at length whether it should be possible to declare a free function kernel where the initial parameter is the iteration index.

We did not come to a conclusion on this, so we agreed to capture this as an open issue, which I did in ecefa00. We do not want this open issue to block initial implementation of this extension.

[gmlueck]

@intel/llvm-gatekeepers I think this PR is ready to merge. I had a meeting with @AerialMantis, and we agreed to capture his remaining comments as an open issue, which I did.

I also met with @bader. His comments are more of a fundamental question about whether SYCL should make a distinction between the three kernels types ("range", "nd-range", and "single task"), or whether we should simplify the language and eliminate both "range" and "single task" kernels. This seems outside of the scope of this PR, and I think he does not mean to hold up this PR with his comment (but let me know, @bader, if I misunderstood).

[AerialMantis]

@gmlueck thanks for the changes, this looks good, I'm happy for this to be merged now and for the remaining questions as they are described in the open issues section, to be discussed further at a later time.

[steffenlarsen]

Merging this. If there are any open or follow-up issues, please open an issue.