[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

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

@@ -621,7 +621,9 @@ jit_compiler::performPostLink(std::unique_ptr<llvm::Module> Module,
   // TODO: EmitOnlyKernelsAsEntryPoints is controlled by
   //       `shouldEmitOnlyKernelsAsEntryPoints` in
   //       `clang/lib/Driver/ToolChains/Clang.cpp`.
-  const bool EmitOnlyKernelsAsEntryPoints = true;
+  const bool EmitOnlyKernelsAsEntryPoints = !UserArgList.hasFlag(
+      options::OPT_fsycl_allow_device_image_dependencies,
+      options::OPT_fno_sycl_allow_device_image_dependencies, false);
 
   // TODO: The optlevel passed to `sycl-post-link` is determined by
   //       `getSYCLPostLinkOptimizationLevel` in

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
@@ -316,6 +316,81 @@ 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 `source_language::sycl_jit` or was the result
+  of `sycl::join` taking one or more `kernel_bundle` objects not created with
+  `source_language::sycl` or `source_language::sycl_jit`.
+
+* 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
@@ -383,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

@@ -1112,6 +1112,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
 
 /////////////////////////
@@ -1155,6 +1185,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
 /////////////////////////