Skip to content

[SYCL][Docs] Implement compile and link for source-based kernel bundles #17442

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 14 commits into
base: sycl
Choose a base branch
from
Open

[SYCL][Docs] Implement compile and link for source-based kernel bundles #17442

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
14 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
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
5 changes: 5 additions & 0 deletions llvm/include/llvm/SYCLPostLink/ModuleSplitter.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -300,6 +300,11 @@ std::unique_ptr<ModuleSplitterBase>
getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
bool EmitOnlyKernelsAsEntryPoints);

std::unique_ptr<ModuleSplitterBase>
getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
bool EmitOnlyKernelsAsEntryPoints,
bool OverwriteAllowDeviceImageDependencies);

#ifndef NDEBUG
void dumpEntryPoints(const EntryPointSet &C, const char *Msg = "", int Tab = 0);
void dumpEntryPoints(const Module &M, bool OnlyKernelsAreEntryPoints = false,
Expand Down
11 changes: 11 additions & 0 deletions llvm/lib/SYCLPostLink/ModuleSplitter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1153,8 +1153,19 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const {
std::unique_ptr<ModuleSplitterBase>
getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
bool EmitOnlyKernelsAsEntryPoints) {
return getDeviceCodeSplitter(std::move(MD), Mode, IROutputOnly,
EmitOnlyKernelsAsEntryPoints,
AllowDeviceImageDependencies);
}

std::unique_ptr<ModuleSplitterBase>
getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
bool EmitOnlyKernelsAsEntryPoints,
bool OverwriteAllowDeviceImageDependencies) {
FunctionsCategorizer Categorizer;

AllowDeviceImageDependencies = OverwriteAllowDeviceImageDependencies;
Copy link
Contributor

Choose a reason for hiding this comment

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

Please, describe a use-case that you want to cover.

This change looks controversial since the interface becomes sophisticated.

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, I am not married to this solution, but it is a result of what I think is a bad design around the module splitter. That is, the module splitter isn't an independent tool, yet it defines its own option "allow-device-image-dependencies". As an alternative solution, I can move this option to the tools using it, i.e. sycl-post-link, sycl-module-split, clang-linker-wrapper and the SYCL kernel compiler. Thoughts?

If we agree, I can open a separate PR doing that. I figure it can be done neatly in separation.

Copy link
Contributor

Choose a reason for hiding this comment

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

cl::opt<bool> AllowDeviceImageDependencies can be removed and replaced by a member in ModuleSplitterBase class.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

My thought exactly! Something like #18060?


EntryPointsGroupScope Scope =
selectDeviceCodeGroupScope(MD.getModule(), Mode, IROutputOnly);

Expand Down
11 changes: 9 additions & 2 deletions sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -648,10 +648,16 @@ jit_compiler::performPostLink(std::unique_ptr<llvm::Module> Module,

const auto SplitMode = getDeviceCodeSplitMode(UserArgList);

const bool AllowDeviceImageDependencies = UserArgList.hasFlag(
options::OPT_fsycl_allow_device_image_dependencies,
options::OPT_fno_sycl_allow_device_image_dependencies, false);

// TODO: EmitOnlyKernelsAsEntryPoints is controlled by
// `shouldEmitOnlyKernelsAsEntryPoints` in
// `clang/lib/Driver/ToolChains/Clang.cpp`.
const bool EmitOnlyKernelsAsEntryPoints = true;
// If we allow device image dependencies, we should definitely not only emit
// kernels as entry points.
const bool EmitOnlyKernelsAsEntryPoints = !AllowDeviceImageDependencies;

// TODO: The optlevel passed to `sycl-post-link` is determined by
// `getSYCLPostLinkOptimizationLevel` in
Expand Down Expand Up @@ -684,7 +690,8 @@ jit_compiler::performPostLink(std::unique_ptr<llvm::Module> Module,

std::unique_ptr<ModuleSplitterBase> Splitter = getDeviceCodeSplitter(
ModuleDesc{std::move(Module)}, SplitMode,
/*IROutputOnly=*/false, EmitOnlyKernelsAsEntryPoints);
/*IROutputOnly=*/false, EmitOnlyKernelsAsEntryPoints,
AllowDeviceImageDependencies);
assert(Splitter->hasMoreSplits());

if (auto Err = Splitter->verifyNoCrossModuleDeviceGlobalUsage()) {
Expand Down
82 changes: 78 additions & 4 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc
View file
Open in desktop
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.


* 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
View file
Open in desktop
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