Skip to content
Merged
Show file tree
Hide file tree
Changes from 17 commits
Commits
Show all changes
35 commits
Select commit Hold shift + click to select a range
70be072
[SYCL] Implement compile and link for source-based kernel bundles
steffenlarsen Mar 12, 2025
419929a
Add trailing newline to tests
steffenlarsen Apr 7, 2025
147cb3d
Address binding capture warning
steffenlarsen Apr 7, 2025
7fc5202
Skip creating image when no devices support it
steffenlarsen Apr 7, 2025
4ef1ce4
Move unsupported skip to earlier
steffenlarsen Apr 7, 2025
14b44c4
Fix test mixup
steffenlarsen Apr 7, 2025
3592e0d
More tests and a correction
steffenlarsen Apr 8, 2025
18a0310
Actually include the tests
steffenlarsen Apr 8, 2025
5472cc9
Fix typos
steffenlarsen Apr 8, 2025
3d5ef96
Another typo
steffenlarsen Apr 8, 2025
a7d51ce
Rename arg to match
steffenlarsen Apr 8, 2025
6d7a894
Avoid exporting RTC image symbols
steffenlarsen Apr 8, 2025
6bb6050
Fix formatting
steffenlarsen Apr 8, 2025
d947479
Check RTC kernel conflict in link and filter unique images
steffenlarsen Apr 9, 2025
8df3327
Merge remote-tracking branch 'intel/sycl' into steffen/kernel_bundle_…
steffenlarsen Apr 22, 2025
3683732
Merge remote-tracking branch 'intel/sycl' into steffen/kernel_bundle_…
steffenlarsen Apr 22, 2025
b03367f
Remove leftover ctor
steffenlarsen Apr 22, 2025
8fa8545
Merge branch 'sycl' into steffen/kernel_bundle_compile_link_2
steffenlarsen Apr 23, 2025
2861344
Clarify build_options
steffenlarsen Apr 23, 2025
449437f
Split ext_oneapi_can_compile in two
steffenlarsen Apr 23, 2025
c33808d
Change exception condition for compile
steffenlarsen Apr 23, 2025
70c692a
Remove the compile source language restriction
steffenlarsen Apr 28, 2025
2f42578
Move language support requirement to build/compile
steffenlarsen Apr 28, 2025
2989462
Address comments
steffenlarsen Apr 28, 2025
cdec73e
Address comments
steffenlarsen Apr 29, 2025
aa9f502
Rename more functions
steffenlarsen Apr 29, 2025
f900a6a
Change to using string_view for options
steffenlarsen Apr 29, 2025
3cc2029
Fix formatting
steffenlarsen Apr 29, 2025
02fca8f
Merge remote-tracking branch 'intel/sycl' into steffen/kernel_bundle_…
steffenlarsen Apr 29, 2025
b5324f0
Fix build error
steffenlarsen Apr 29, 2025
d6dc083
Remove unused code
steffenlarsen Apr 29, 2025
81943ea
Fix windows build
steffenlarsen Apr 29, 2025
62f06fc
Switch to unordered_map
steffenlarsen Apr 29, 2025
5a189dc
Update sycl/include/sycl/kernel_bundle.hpp
steffenlarsen May 1, 2025
daeacc3
Update sycl/include/sycl/kernel_bundle.hpp
steffenlarsen May 1, 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
Original file line number Diff line number Diff line change
Expand Up @@ -271,8 +271,8 @@ kernel_bundle<bundle_state::executable> build(

_Constraints:_ Available only when `PropertyListT` is an instance of
`sycl::ext::oneapi::experimental::properties` which contains no properties
other than those listed below in the section "New properties for the `build`
function".
other than those listed below in the section "New properties for the `build` and
`compile` functions".

_Effects (1):_ The source code from `sourceBundle` is translated into one or more
device images of state `bundle_state::executable`, and a new kernel bundle is
Expand Down Expand Up @@ -317,6 +317,80 @@ source code used to create the kernel bundle being printed to the terminal.
In situations where this is undesirable, developers must ensure that the
exception is caught and handled appropriately.
_{endnote}_]

a|
[frame=all,grid=none]
!====
a!
[source]
----
namespace sycl::ext::oneapi::experimental {

template<typename PropertyListT = empty_properties_t> (1)
kernel_bundle<bundle_state::object> compile(
const kernel_bundle<bundle_state::ext_oneapi_source>& sourceBundle,
const std::vector<device>& devs, PropertyListT props={})

template<typename PropertyListT = empty_properties_t> (2)
kernel_bundle<bundle_state::object> compile(
const kernel_bundle<bundle_state::ext_oneapi_source>& sourceBundle,
PropertyListT props = {})

} // namespace sycl::ext::oneapi::experimental
----
!====


_Constraints:_ Available only when `PropertyListT` is an instance of
`sycl::ext::oneapi::experimental::properties` which contains no properties
other than those listed below in the section "New properties for the `build` and
`compile` functions".

_Effects (1):_ The source code from `sourceBundle` is translated into one or
more device images of state `bundle_state::object`, and a new kernel bundle is
created to contain these device images.
The new bundle represents all of the kernels in `sourceBundle` that are
compatible with at least one of the devices in `devs`.
Any remaining kernels (those that are not compatible with any of the devices in
`devs`) are not represented in the new kernel bundle.

The new bundle has the same associated context as `sourceBundle`, and the new
bundle's set of associated devices is `devs` (with duplicate devices removed).

_Effects (2)_: Equivalent to
`compile(sourceBundle, sourceBundle.get_devices(), props)`.

_Returns:_ The newly created kernel bundle, which has `object` state.

_Throws:_

* An `exception` with the `errc::invalid` error code if `source` was not created
with `source_language::sycl` or was the result of `sycl::join` taking one or
more `kernel_bundle` objects not created with `source_language::sycl`.
Copy link
Contributor

Choose a reason for hiding this comment

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

There is no similar restriction for build. Is there a reason we need this restriction for compile? Can we not compile other languages like OpenCL or SPIR-V?

It seems weird to list the sycl language specifically here. Should we instead split the query ext_oneapi_can_compile into two queries (for compile and build)? Or, do you think the limitation with sycl is only temporary, and that's why you didn't add a query?

The whole business about join seems like a can of worms. As you allude, the application could joint two source bundles from different languages! I think it's better to avoid the whole can of worms by adding a constraint to join, preventing it from joining bundles of source state.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

There is no similar restriction for build. Is there a reason we need this restriction for compile? Can we not compile other languages like OpenCL or SPIR-V?

I believe we could relax this restriction to allow compiling OpenCL, SPIR-V, etc, but by allowing these source languages in object state complicates the logic for linking a little, as we wouldn't know their imported/exported symbols. There are ways we could solve that, but since my focus was on SYCL RTC kernel bundles in this initial variant, I elected to disallow them for now.

It seems weird to list the sycl language specifically here. Should we instead split the query ext_oneapi_can_compile into two queries (for compile and build)? Or, do you think the limitation with sycl is only temporary, and that's why you didn't add a query?

I did not consider splitting ext_oneapi_can_compile. I think it would make sense (unless we know we want to expand the above and don't think there will be cases where a target can build but not compile,) but since it would be ext_oneapi_can_compile for compile we have to repurpose the one we already have and probably change the uses of it around existing code. Would it make sense to do this in a follow-up for the sake of change granularity?

The whole business about join seems like a can of worms. As you allude, the application could joint two source bundles from different languages! I think it's better to avoid the whole can of worms by adding a constraint to join, preventing it from joining bundles of source state.

I'm still not sure it's necessarily that big a can of worms. I agree, the wording here is a little awkward because of there not currently being any restrictions on join and I think it might make sense to at least specify that joining kernel bundles in source state must be of the same state. Logically that makes sense in my head, as I think the source language of a source state bundle is intrinsic to it. I would be happy to sketch up some wording for a section describing the behavior of join with source state bundles. 😄

Copy link
Contributor

Choose a reason for hiding this comment

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

I agree that it makes sense to support only SYCL for compilation at this point. In that case, I think we need to have a query that tells the application which source languages can be compiled.

Let's repurpose ext_oneapi_can_compile and introduce ext_oneapi_can_build. I realize this may break some existing code. We'll have to reach out to our existing users and let them know.

I'm still not convinced about join. A kernel bundle in source state just contains a string (or a byte array) that is the source code for the kernel. The join operation can't just concatenate them. In the case of SYCL, this would result in a source code string that includes <sycl/sycl.hpp> twice. I don't think any customer is requesting join support for source bundles. In the spirit of providing an MVP, I think we should just disallow join for source bundles. This is easy to do by adding a constraint.

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 agree that it makes sense to support only SYCL for compilation at this point. In that case, I think we need to have a query that tells the application which source languages can be compiled.

Let's repurpose ext_oneapi_can_compile and introduce ext_oneapi_can_build. I realize this may break some existing code. We'll have to reach out to our existing users and let them know.

Sure!

I'm still not convinced about join. A kernel bundle in source state just contains a string (or a byte array) that is the source code for the kernel. The join operation can't just concatenate them. In the case of SYCL, this would result in a source code string that includes <sycl/sycl.hpp> twice. I don't think any customer is requesting join support for source bundles. In the spirit of providing an MVP, I think we should just disallow join for source bundles. This is easy to do by adding a constraint.

I suppose my mental model for join doesn't append the sources, just like two kernel_bundle being joined don't combine their underlying binaries. I suppose those are not exposed to the user either, so maybe your way of looking at it is the more natural one. Another way to look at it, in core SYCL 2020 you can get the same kernel_bundle through regular constructors as you could using join... Maybe with some interop exceptions.

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 added ext_oneapi_can_build and repurposed ext_oneapi_can_compile. For the mention of sycl::join in the exception condition of compile, I've changed the wording to ignore join for now. Since join support was added in a previous patch, I'll make the note about join in a separate patch.

Copy link
Contributor

Choose a reason for hiding this comment

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

Deferring the restriction on join to another PR is fine.

The latest wording still calls out the "sycl" language specifically:

An exception with the errc::invalid error code if source was not created with source_language::sycl.

I thought we added ext_oneapi_can_compile in order to generalize this restriction.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Good catch! That line has been removed.


* An `exception` with the `errc::invalid` error code if any of the devices in
`devs` is not contained by the context associated with `sourceBundle`.

* An `exception` with the `errc::invalid` error code if any of the devices in
`devs` does not support compilation of kernels in the source language of
`sourceBundle`.

* An `exception` with the `errc::invalid` error code if `props` contains an
`build_options` property that specifies an invalid option.

* An `exception` with the `errc::build` error code if the compilation operation
fails. In this case, the exception `what` string provides a full build log,
including descriptions of any errors, warning messages, and other
diagnostics.
This string is intended for human consumption, and the format may not be
stable across implementations of this extension.

[_Note:_ An uncaught `errc::build` exception may result in some or all of the
source code used to create the kernel bundle being printed to the terminal.
In situations where this is undesirable, developers must ensure that the
exception is caught and handled appropriately.
_{endnote}_]

|====

=== New properties for the `create_kernel_bundle_from_source` function
Expand Down Expand Up @@ -384,10 +458,10 @@ _Throws (3):_
entry with `name` in this property.
|====

=== New properties for the `build` function
=== New properties for the `build` and `compile` functions

This extension adds the following properties, which can be used in conjunction
with the `build` function that is defined above:
with the `build` and `compile` function that is defined above:

|====
a|
Expand Down
63 changes: 63 additions & 0 deletions sycl/include/sycl/kernel_bundle.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -1183,6 +1183,36 @@ build_from_source(kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
}
return build_from_source(SourceKB, Devices, Options, nullptr, KernelNames);
}

__SYCL_EXPORT kernel_bundle<bundle_state::object> compile_from_source(
kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
const std::vector<device> &Devices,
const std::vector<sycl::detail::string_view> &CompileOptions,
sycl::detail::string *LogPtr,
const std::vector<sycl::detail::string_view> &RegisteredKernelNames);

inline kernel_bundle<bundle_state::object>
compile_from_source(kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
const std::vector<device> &Devices,
const std::vector<std::string> &CompileOptions,
std::string *LogPtr,
const std::vector<std::string> &RegisteredKernelNames) {
std::vector<sycl::detail::string_view> Options;
for (const std::string &opt : CompileOptions)
Options.push_back(sycl::detail::string_view{opt});

std::vector<sycl::detail::string_view> KernelNames;
for (const std::string &name : RegisteredKernelNames)
KernelNames.push_back(sycl::detail::string_view{name});

sycl::detail::string Log;
auto result = compile_from_source(SourceKB, Devices, Options,
LogPtr ? &Log : nullptr, KernelNames);
if (LogPtr)
*LogPtr = Log.c_str();
return result;
}

} // namespace detail

/////////////////////////
Expand Down Expand Up @@ -1220,6 +1250,39 @@ kernel_bundle<bundle_state::ext_oneapi_source> create_kernel_bundle_from_source(
}
#endif

/////////////////////////
// syclex::compile(source_kb) => obj_kb
/////////////////////////

template <typename PropertyListT = empty_properties_t,
typename = std::enable_if_t<detail::all_are_properties_of_v<
detail::build_source_bundle_props, PropertyListT>>>
kernel_bundle<bundle_state::object>
compile(kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
const std::vector<device> &Devices, PropertyListT props = {}) {
std::vector<std::string> CompileOptionsVec;
std::string *LogPtr = nullptr;
std::vector<std::string> RegisteredKernelNamesVec;
if constexpr (props.template has_property<build_options>())
CompileOptionsVec = props.template get_property<build_options>().opts;
if constexpr (props.template has_property<save_log>())
LogPtr = props.template get_property<save_log>().log;
if constexpr (props.template has_property<registered_names>())
RegisteredKernelNamesVec =
props.template get_property<registered_names>().names;
return detail::compile_from_source(SourceKB, Devices, CompileOptionsVec,
LogPtr, RegisteredKernelNamesVec);
}

template <typename PropertyListT = empty_properties_t,
typename = std::enable_if_t<detail::all_are_properties_of_v<
detail::build_source_bundle_props, PropertyListT>>>
kernel_bundle<bundle_state::object>
compile(kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
PropertyListT props = {}) {
return compile<PropertyListT>(SourceKB, SourceKB.get_devices(), props);
}

/////////////////////////
// syclex::build(source_kb) => exe_kb
/////////////////////////
Expand Down
Loading
Loading