[SYCL] Pass driver options to JIT compiler and linker (#5476)

This patch implements passing compiler and link options to sycl::compile() and sycl::link() respectively.

Co-authored-by: Romanov Vlad <[email protected]>
Co-authored-by: Steffen Larsen <[email protected]>

Author: 3 people

sycl/source/detail/program_manager/program_manager.cpp

@@ -365,37 +365,52 @@ RT::PiProgram ProgramManager::createPIProgram(const RTDeviceBinaryImage &Img,
 
   return Res;
 }
-static void applyOptionsFromImage(std::string &CompileOpts,
-                                  std::string &LinkOpts,
-                                  const RTDeviceBinaryImage &Img) {
+
+static void appendLinkOptionsFromImage(std::string &LinkOpts,
+                                       const RTDeviceBinaryImage &Img) {
+  static const char *LinkOptsEnv = SYCLConfig<SYCL_PROGRAM_LINK_OPTIONS>::get();
+  // Update only if link options are not overwritten by environment variable
+  if (!LinkOptsEnv) {
+    const char *TemporaryStr = Img.getLinkOptions();
+    if (TemporaryStr != nullptr) {
+      if (!LinkOpts.empty())
+        LinkOpts += " ";
+      LinkOpts += std::string(TemporaryStr);
+    }
+  }
+}
+
+static void appendCompileOptionsFromImage(std::string &CompileOpts,
+                                          const RTDeviceBinaryImage &Img) {
   // Build options are overridden if environment variables are present.
   // Environment variables are not changed during program lifecycle so it
   // is reasonable to use static here to read them only once.
   static const char *CompileOptsEnv =
       SYCLConfig<SYCL_PROGRAM_COMPILE_OPTIONS>::get();
-  static const char *LinkOptsEnv = SYCLConfig<SYCL_PROGRAM_LINK_OPTIONS>::get();
+  pi_device_binary_property isEsimdImage = Img.getProperty("isEsimdImage");
   // Update only if compile options are not overwritten by environment
   // variable
   if (!CompileOptsEnv) {
     if (!CompileOpts.empty())
       CompileOpts += " ";
-    CompileOpts += Img.getCompileOptions();
+    const char *TemporaryStr = Img.getCompileOptions();
+    if (TemporaryStr != nullptr)
+      CompileOpts += std::string(TemporaryStr);
   }
-
   // The -vc-codegen option is always preserved for ESIMD kernels, regardless
   // of the contents SYCL_PROGRAM_COMPILE_OPTIONS environment variable.
-  pi_device_binary_property isEsimdImage = Img.getProperty("isEsimdImage");
   if (isEsimdImage && pi::DeviceBinaryProperty(isEsimdImage).asUint32()) {
     if (!CompileOpts.empty())
       CompileOpts += " ";
     CompileOpts += "-vc-codegen";
   }
+}
 
-  // Update only if link options are not overwritten by environment variable
-  if (!LinkOptsEnv)
-    if (!LinkOpts.empty())
-      LinkOpts += " ";
-  LinkOpts += Img.getLinkOptions();
+static void applyOptionsFromImage(std::string &CompileOpts,
+                                  std::string &LinkOpts,
+                                  const RTDeviceBinaryImage &Img) {
+  appendCompileOptionsFromImage(CompileOpts, Img);
+  appendLinkOptionsFromImage(LinkOpts, Img);
 }
 
 static void applyOptionsFromEnvironment(std::string &CompileOpts,
@@ -1000,9 +1015,12 @@ ProgramManager::ProgramPtr ProgramManager::build(
 
   const detail::plugin &Plugin = Context->getPlugin();
   if (LinkPrograms.empty() && !ForceLink) {
+    const std::string &Options = LinkOptions.empty()
+                                     ? CompileOptions
+                                     : (CompileOptions + " " + LinkOptions);
     RT::PiResult Error = Plugin.call_nocheck<PiApiKind::piProgramBuild>(
-        Program.get(), /*num devices =*/1, &Device, CompileOptions.c_str(),
-        nullptr, nullptr);
+        Program.get(), /*num devices =*/1, &Device, Options.c_str(), nullptr,
+        nullptr);
     if (Error != PI_SUCCESS)
       throw compile_program_error(getProgramBuildLog(Program.get(), Context),
                                   Error);
@@ -1674,10 +1692,12 @@ ProgramManager::compile(const device_image_plain &DeviceImage,
   // TODO: Set spec constatns here.
 
   // TODO: Handle zero sized Device list.
+  std::string CompileOptions;
+  appendCompileOptionsFromImage(CompileOptions,
+                                *(InputImpl->get_bin_image_ref()));
   RT::PiResult Error = Plugin.call_nocheck<PiApiKind::piProgramCompile>(
       ObjectImpl->get_program_ref(), /*num devices=*/Devs.size(),
-      PIDevices.data(),
-      /*options=*/nullptr,
+      PIDevices.data(), CompileOptions.c_str(),
       /*num_input_headers=*/0, /*input_headers=*/nullptr,
       /*header_include_names=*/nullptr,
       /*pfn_notify=*/nullptr, /*user_data*/ nullptr);
@@ -1706,15 +1726,21 @@ ProgramManager::link(const std::vector<device_image_plain> &DeviceImages,
   for (const device &Dev : Devs)
     PIDevices.push_back(getSyclObjImpl(Dev)->getHandleRef());
 
+  std::string LinkOptionsStr;
+  for (const device_image_plain &DeviceImage : DeviceImages) {
+    const std::shared_ptr<device_image_impl> &InputImpl =
+        getSyclObjImpl(DeviceImage);
+    appendLinkOptionsFromImage(LinkOptionsStr,
+                               *(InputImpl->get_bin_image_ref()));
+  }
   const context &Context = getSyclObjImpl(DeviceImages[0])->get_context();
   const ContextImplPtr ContextImpl = getSyclObjImpl(Context);
-
   const detail::plugin &Plugin = ContextImpl->getPlugin();
 
   RT::PiProgram LinkedProg = nullptr;
   RT::PiResult Error = Plugin.call_nocheck<PiApiKind::piProgramLink>(
       ContextImpl->getHandleRef(), PIDevices.size(), PIDevices.data(),
-      /*options=*/nullptr, PIPrograms.size(), PIPrograms.data(),
+      /*options=*/LinkOptionsStr.c_str(), PIPrograms.size(), PIPrograms.data(),
       /*pfn_notify=*/nullptr,
       /*user_data=*/nullptr, &LinkedProg);
 

sycl/unittests/misc/KernelBuildOptions.cpp

@@ -243,7 +243,7 @@ TEST(KernelBuildOptions, KernelBundleBasic) {
   sycl::kernel_bundle KernelBundle =
       sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev});
   auto ExecBundle = sycl::build(KernelBundle);
-  EXPECT_EQ(BuildOpts, "-compile-img -vc-codegen");
+  EXPECT_EQ(BuildOpts, "-compile-img -vc-codegen -link-img");
 
   auto ObjBundle = sycl::compile(KernelBundle, KernelBundle.get_devices());
   // TODO: uncomment when image options are passed to BE

sycl/unittests/program_manager/CMakeLists.txt

@@ -5,5 +5,6 @@ add_sycl_unittest(ProgramManagerTests OBJECT
   EliminatedArgMask.cpp
   itt_annotations.cpp
   SubDevices.cpp
+  passing_link_and_compile_options.cpp
 )