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

Implements a new `compile` variant for source-based `kernel_bundle` and
the corresponding linking functionality.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

llvm/include/llvm/SYCLPostLink/ModuleSplitter.h

@@ -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,

llvm/lib/SYCLPostLink/ModuleSplitter.cpp

@@ -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;
+
   EntryPointsGroupScope Scope =
       selectDeviceCodeGroupScope(MD.getModule(), Mode, IROutputOnly);
 

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

@@ -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
@@ -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()) {

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc

@@ -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
@@ -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>& source,
+    const std::vector<device>& devs, PropertyListT props={})
+
+template<typename PropertyListT = empty_properties_t>                 (2)
+kernel_bundle<bundle_state::executable> build(
+  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`.
+
+* 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
+  `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
@@ -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|

sycl/include/sycl/kernel_bundle.hpp

@@ -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
 
 /////////////////////////
@@ -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
 /////////////////////////