Skip to content

[SYCL][Doc] Add kernel_function lambda wrapper #17633

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.

Sign up for GitHub

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

Open
wants to merge 6 commits into
base: sycl
Choose a base branch
from
Open

[SYCL][Doc] Add kernel_function lambda wrapper #17633

Changes from 2 commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
f6581d6
[SYCL][Doc] Add kernel_function lambda wrapper
Pennycook Mar 25, 2025
82de663
Fix a few references to old behavior
Pennycook Mar 25, 2025
be618af
Remove implementation details from synopsis
Pennycook Mar 25, 2025
3d580ff
Add && to Function and Args
Pennycook Mar 25, 2025
160fff7
Remove syclx:: from synopsis and definitions
Pennycook Mar 25, 2025
1743542
Clarify meaning of "extend" kernel_function
Pennycook Apr 4, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
77 changes: 67 additions & 10 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -287,16 +287,17 @@ by the property, the implementation must throw a synchronous exception with the

|===

== Embedding Properties into a Kernel
=== Embedding Properties into a Kernel

In other situations it may be useful to embed a kernel's properties directly
into its type, to ensure that a kernel cannot be launched without a property
that it depends upon for correctness.
A kernel's properties are embedded directly into its type, to ensure that a
kernel cannot be launched without a property that it depends upon for
correctness.

To enable this use-case, this extension adds a mechanism for implementations to
extract a property list from a kernel functor, if a kernel functor declares
a member function named `get` accepting a `sycl::ext::oneapi::experimental::properties_tag`
tag type and returning an instance of `sycl::ext::oneapi::experimental::properties`.
To enable this, this extension adds a mechanism for implementations to extract
a property list from a kernel functor, if a kernel functor declares a member
function named `get` accepting a
`sycl::ext::oneapi::experimental::properties_tag` tag type and returning an
instance of `sycl::ext::oneapi::experimental::properties`.

```c++
namespace sycl {
Expand All @@ -323,8 +324,8 @@ attributes to be applied to different call operators within the same
functor. An embedded property list applies to all call operators in
the functor.

The example below shows how the kernel from the previous section could be
rewritten to leverage an embedded property list:
The example below shows how a simple vector addition kernel could be
written to leverage an embedded property list:

```c++
struct KernelFunctor {
Expand Down Expand Up @@ -363,6 +364,62 @@ diagnostic; invalid combinations that can only be detected at run-time should
result in an implementation throwing an `exception` with the `errc::invalid`
error code.

=== Using Properties with Lambda Expressions

When a SYCL kernel is defined via a lambda expression, there is no way to
define a `get` member function and subsequently no way to embed kernel
properties. Instead, developers must wrap the lambda expression in an object.

To simplify this usage pattern, this extension defines a `kernel_function`
that encapsulates a kernel function (which may be a lambda expression) and a
property list.

NOTE: Developers are free to extend `kernel_function` or define their own
wrapper classes (e.g., to attach commonly used property lists).

```c++
namespace sycl::ext::oneapi::experimental {

template <typename Function, typename Properties = empty_properties_t>
struct kernel_function {

kernel_function(Function f, Properties p = syclx::properties{});
Copy link
Contributor

Choose a reason for hiding this comment

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

Suggested change
kernel_function(Function f, Properties p = syclx::properties{});
kernel_function(Function &&f, Properties p = syclx::properties{});

And then std::move(f) inside implementation. That somewhat limits the applicability, but I'd rather start with that and extend later if needed.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Added in 3d580ff.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Okay, adding && to my prototype actually breaks it... Whether I use const && or &&, I get an error about expecting an rvalue. I thought && would accept both, but it seems not to work.

We want to support both uses below. What syntax do we need?

auto lambda = [=]() {};
auto kernel = syclx::kernel_function(lambda); // lambda is an l-value

auto kernel = syclx::kernel_function([=]() {}); // lambda is an r-value

Copy link
Contributor

Choose a reason for hiding this comment

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

https://godbolt.org/z/fK36njGse works, but I don't know if that's the correct/idiomatic way. Otherwise, two overloads work too.


// Available only if Function is invocable with Args
template <typename... Args>
void operator()(Args... args) const {
Copy link
Contributor

@aelovikov-intel aelovikov-intel Mar 25, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

Suggested change
void operator()(Args... args) const {
void operator()(Args&& ...args) const {

I'm not sure about const too.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Added in 3d580ff.

f(std::forward<Args>(args)...);
}

auto get(syclx::properties_tag) {
return Properties{};
}
Copy link
Contributor

@aelovikov-intel aelovikov-intel Mar 25, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

IMO, we should be storing properties object, in case it will contain runtime properties in future.

I think we also need to have two version of this - one static constexpr if Properties is std::empty_v, similarly to how properties' get_property is implemented.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

IMO, we should be storing properties object, in case it will contain runtime properties in future.

The reason I didn't do this is that anything stored in the function object has to be transferred to the device as one of the kernel arguments. I think it would only make sense to allow this if the run-time properties were intended to be consumed within the kernel.

We'd also have to add an explicit specialization for the case where the property list is empty or contains only compile-time properties, to avoid transferring 1 byte unnecessarily. But perhaps this is a quality of implementation thing.

I think we also need to have two version of this - one static constexpr if Properties is std::empty_v, similarly to how properties' get_property is implemented.

Can you say more about this? Why would we only want it to be static constexpr in this case? Could we make it static constexpr if it doesn't contain any run-time properties, instead?

Copy link
Contributor

@aelovikov-intel aelovikov-intel Mar 25, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

I think we also need to have two version of this - one static constexpr if Properties is std::empty_v, similarly to how properties' get_property is implemented.

Can you say more about this? Why would we only want it to be static constexpr in this case? Could we make it static constexpr if it doesn't contain any run-time properties, instead?

std::is_empty is exactly the check for compile-time-only (no run-time), see

llvm/sycl/include/sycl/ext/oneapi/properties/properties.hpp

Lines 256 to 273 in 5c5954e

// Compile-time property.
template <typename property_key_t>
static constexpr auto
get_property() -> std::enable_if_t<std::is_empty_v<prop_t<property_key_t>>,
prop_t<property_key_t>> {
return prop_t<property_key_t>{};
}
// Runtime property.
// Extra operand to make MSVC happy as it complains otherwise:
// https://godbolt.org/z/WGqdqrejj
template <typename property_key_t>
constexpr auto get_property(int = 0) const
-> std::enable_if_t<!std::is_empty_v<prop_t<property_key_t>>,
prop_t<property_key_t>> {
return get_property_impl(detail::property_key_tag<property_key_t>{});
}
};

I'm not sure if we have dedicated traits in the properties extension though. Maybe we can workaround that by using plain English in constraints/requirements?

Copy link
Contributor

@aelovikov-intel aelovikov-intel Mar 25, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

But perhaps this is a quality of implementation thing.

Maybe we could protected inherit from the properties to do that, but I agree about QoI.

This does sound like an issue in the "base" part of the extension though. IMO, having a "getter" for the properties without one for the actual kernel is the root cause of this issue.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I've tried to address this in be618af; I've removed the implementation details from the synopsis, so that the implementation is no longer normative, and tried to describe the constraints we want.

The synopsis no longer shows what members are stored inside the kernel_function object, so implementations have the freedom to do whatever they want to meet the specified behavior.

Copy link
Contributor

Choose a reason for hiding this comment

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

Why are we making accommodations for runtime kernel properties? Aren't kernel properties always compile-time by their very nature? If you think you want a runtime kernel property, then you must really want a launch property (whose value can change each time a kernel is launched). We already have a way to specify launch properties whose values are defined at runtime.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I think we can make either work. I thought compile-time only properties made sense, but several people have asked about run-time properties now.

If we restrict this to compile-time properties only, then we're saying that the only properties that can be attached to a kernel function are properties that affect the actual compilation of the kernel. There'd be no room for properties that were "properties of the kernel" rather than "properties of how the kernel should be launched", and we couldn't change this later.

I'm happy to stick with compile-time only if everybody is confident that we'll never need run-time properties, or that we'll always be able to express run-time properties as a property of the launch in a way that makes sense semantically.

Copy link
Contributor

Choose a reason for hiding this comment

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

We added kernel properties as a replacement for the C++ attribute decorations we have in SYCL 2020. C++ attributes are always a compile-time thing, so we do not need runtime properties to achieve our original goal.

I'm having a hard time imagining a runtime kernel property. It seems like you could either use a launch property (if it affected the way the kernel was launched) or just pass a kernel argument (if it changed the way the kernel runs). Do you have an example of a runtime kernel property that wouldn't fit into either of these categories?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I don't have any run-time kernel properties in mind, no. @aelovikov-intel - Do you have any in mind?

If nobody has any suggestions for run-time kernel properties... Since the current plan is to do this via two different overloads, could we just remove the non-static constexpr one, and revisit this later? We could always add the other overload later.

It would still make sense to have the Constraint that the static constexpr overload is only available with compile-time property lists, because we wouldn't be exposing run-time properties. As long as we have that constraint, there'd be no problem with introducing another overload later.

Copy link
Contributor

Choose a reason for hiding this comment

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

Trunk version has

llvm/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc

Lines 341 to 344 in a123a74

auto get(sycl::ext::oneapi::experimental::properties_tag) {
return sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::work_group_size<8, 8>,
sycl::ext::oneapi::experimental::sub_group_size<8>};
}

If only compile-time properties are accepted, the extension needs to be consistent throughout. Also, if compile-time only, why static getter? Can't we just have an optional type alias? That will always be compile-time only...


private:
const Function f; // exposition only

} // namespace sycl::ext::oneapi::experimental
```

The example below shows how the `KernelFunctor` example from the previous
section can be written using this wrapper:

```c++
namespace syclx = sycl::ext::oneapi::experimental;

...

auto lambda = [=](id<1> i) const {
a[i] = b[i] + c[i];
}
auto props = syclx::properties{syclx::work_group_size<8, 8>, syclx::sub_group_size<8>};
auto kernel = syclx::kernel_function(lambda, props);

...

q.parallel_for(range<2>{16, 16}, kernel).wait();
```

=== Querying Properties in a Compiled Kernel

Any properties embedded into a kernel type via a property list are reflected
Expand Down