[SYCL][RTC] Propagate -Xs options along in-memory pipeline (intel#17329)

Ensures that `-Xs` option arguments are baked into the device image(s),
from which they are picked up by the progam manager.

Signed-off-by: Julian Oppermann <[email protected]>

Author: jopperm

sycl-jit/common/include/Kernel.h

@@ -409,7 +409,14 @@ struct RTCDevImgInfo {
   RTCDevImgInfo &operator=(RTCDevImgInfo &&) = default;
 };
 
-using RTCBundleInfo = DynArray<RTCDevImgInfo>;
+struct RTCBundleInfo {
+  DynArray<RTCDevImgInfo> DevImgInfos;
+  sycl::detail::string CompileOptions;
+
+  RTCBundleInfo() = default;
+  RTCBundleInfo(RTCBundleInfo &&) = default;
+  RTCBundleInfo &operator=(RTCBundleInfo &&) = default;
+};
 
 // LLVM's APIs prefer `char *` for byte buffers.
 using RTCDeviceCodeIR = DynArray<char>;

sycl-jit/jit-compiler/lib/KernelFusion.cpp

@@ -361,7 +361,8 @@ compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
   }
   auto [BundleInfo, Modules] = std::move(*PostLinkResultOrError);
 
-  for (auto [DevImgInfo, Module] : llvm::zip_equal(BundleInfo, Modules)) {
+  for (auto [DevImgInfo, Module] :
+       llvm::zip_equal(BundleInfo.DevImgInfos, Modules)) {
     auto BinaryInfoOrError =
         translation::KernelTranslator::translateDevImgToSPIRV(
             *Module, JITContext::getInstance());
@@ -372,6 +373,8 @@ compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
     DevImgInfo.BinaryInfo = std::move(*BinaryInfoOrError);
   }
 
+  encodeBuildOptions(BundleInfo, UserArgList);
+
   if (llvm::timeTraceProfilerEnabled()) {
     auto Error = llvm::timeTraceProfilerWrite(
         TraceFileName, /*FallbackFileName=*/"trace.json");

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

@@ -661,12 +661,13 @@ jit_compiler::performPostLink(std::unique_ptr<llvm::Module> Module,
 
   // TODO: This allocation assumes that there are no further splits required,
   //       i.e. there are no mixed SYCL/ESIMD modules.
-  RTCBundleInfo BundleInfo{Splitter->remainingSplits()};
+  RTCBundleInfo BundleInfo;
+  BundleInfo.DevImgInfos = DynArray<RTCDevImgInfo>{Splitter->remainingSplits()};
   SmallVector<std::unique_ptr<llvm::Module>> Modules;
 
-  auto *DevImgInfoIt = BundleInfo.begin();
+  auto *DevImgInfoIt = BundleInfo.DevImgInfos.begin();
   while (Splitter->hasMoreSplits()) {
-    assert(DevImgInfoIt != BundleInfo.end());
+    assert(DevImgInfoIt != BundleInfo.DevImgInfos.end());
 
     ModuleDesc MDesc = Splitter->nextSplit();
     RTCDevImgInfo &DevImgInfo = *DevImgInfoIt++;
@@ -796,6 +797,32 @@ jit_compiler::parseUserArgs(View<const char *> UserArgs) {
   return std::move(AL);
 }
 
+void jit_compiler::encodeBuildOptions(RTCBundleInfo &BundleInfo,
+                                      const InputArgList &UserArgList) {
+  std::string CompileOptions;
+  raw_string_ostream COSOS{CompileOptions};
+
+  for (Arg *A : UserArgList.getArgs()) {
+    if (!(A->getOption().matches(OPT_Xs) ||
+          A->getOption().matches(OPT_Xs_separate))) {
+      continue;
+    }
+
+    // Trim first and last quote if they exist, but no others.
+    StringRef AV{A->getValue()};
+    AV = AV.trim();
+    if (AV.front() == AV.back() && (AV.front() == '\'' || AV.front() == '"')) {
+      AV = AV.drop_front().drop_back();
+    }
+
+    COSOS << (CompileOptions.empty() ? "" : " ") << AV;
+  }
+
+  if (!CompileOptions.empty()) {
+    BundleInfo.CompileOptions = CompileOptions;
+  }
+}
+
 void jit_compiler::configureDiagnostics(LLVMContext &Context,
                                         std::string &BuildLog) {
   Context.setDiagnosticHandler(

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

@@ -44,6 +44,9 @@ performPostLink(std::unique_ptr<llvm::Module> Module,
 llvm::Expected<llvm::opt::InputArgList>
 parseUserArgs(View<const char *> UserArgs);
 
+void encodeBuildOptions(RTCBundleInfo &BundleInfo,
+                        const llvm::opt::InputArgList &UserArgList);
+
 void configureDiagnostics(llvm::LLVMContext &Context, std::string &BuildLog);
 
 } // namespace jit_compiler

sycl/source/detail/jit_compiler.cpp

@@ -1155,7 +1155,7 @@ sycl_device_binaries jit_compiler::createDeviceBinaries(
     const std::string &Prefix) {
   auto Collection = std::make_unique<DeviceBinariesCollection>();
 
-  for (const auto &DevImgInfo : BundleInfo) {
+  for (const auto &DevImgInfo : BundleInfo.DevImgInfos) {
     DeviceBinaryContainer Binary;
     for (const auto &Symbol : DevImgInfo.SymbolTable) {
       // Create an offload entry for each kernel. We prepend a unique prefix to
@@ -1182,6 +1182,8 @@ sycl_device_binaries jit_compiler::createDeviceBinaries(
         }
       }
       Binary.addProperty(std::move(PropSet));
+
+      Binary.setCompileOptions(BundleInfo.CompileOptions.c_str());
     }
 
     Collection->addDeviceBinary(std::move(Binary),

sycl/source/detail/jit_compiler.hpp

@@ -27,10 +27,10 @@ class JITContext;
 struct SYCLKernelInfo;
 struct SYCLKernelAttribute;
 struct RTCDevImgInfo;
+struct RTCBundleInfo;
 template <typename T> class DynArray;
 using ArgUsageMask = DynArray<uint8_t>;
 using JITEnvVar = DynArray<char>;
-using RTCBundleInfo = DynArray<RTCDevImgInfo>;
 } // namespace jit_compiler
 
 namespace sycl {

sycl/source/detail/jit_device_binaries.cpp

@@ -89,14 +89,26 @@ void DeviceBinaryContainer::addProperty(PropertySetContainer &&Cont) {
   PropertySets.push_back(std::move(Cont));
 }
 
+void DeviceBinaryContainer::setCompileOptions(std::string_view CompileOpts) {
+  // Forbid calls to this method after the first PI struct has been created.
+  assert(Fused && "Reallocating string would invalidate existing UR structs");
+  if (CompileOpts.empty()) {
+    CompileOptions.reset();
+    return;
+  }
+  CompileOptions.reset(new char[CompileOpts.length() + 1]);
+  std::memcpy(CompileOptions.get(), CompileOpts.data(),
+              CompileOpts.length() + 1);
+}
+
 sycl_device_binary_struct DeviceBinaryContainer::getPIDeviceBinary(
     const unsigned char *BinaryStart, size_t BinarySize, const char *TargetSpec,
     sycl_device_binary_type Format) {
   sycl_device_binary_struct DeviceBinary;
   DeviceBinary.Version = SYCL_DEVICE_BINARY_VERSION;
   DeviceBinary.Kind = SYCL_DEVICE_BINARY_OFFLOAD_KIND_SYCL;
   DeviceBinary.Format = Format;
-  DeviceBinary.CompileOptions = "";
+  DeviceBinary.CompileOptions = CompileOptions ? CompileOptions.get() : "";
   DeviceBinary.LinkOptions = "";
   DeviceBinary.ManifestStart = nullptr;
   DeviceBinary.ManifestEnd = nullptr;

sycl/source/detail/jit_device_binaries.hpp

@@ -13,6 +13,7 @@
 
 #include <cstring>
 #include <memory>
+#include <string_view>
 
 namespace sycl {
 inline namespace _V1 {
@@ -113,6 +114,8 @@ class DeviceBinaryContainer {
 
   void addProperty(PropertySetContainer &&Cont);
 
+  void setCompileOptions(std::string_view CompileOpts);
+
   sycl_device_binary_struct getPIDeviceBinary(const unsigned char *BinaryStart,
                                               size_t BinarySize,
                                               const char *TargetSpec,
@@ -124,6 +127,7 @@ class DeviceBinaryContainer {
   std::vector<_sycl_offload_entry_struct> PIOffloadEntries;
   std::vector<PropertySetContainer> PropertySets;
   std::vector<_sycl_device_binary_property_set_struct> PIPropertySets;
+  std::unique_ptr<char[]> CompileOptions;
 };
 
 /// Representation of sycl_device_binaries_struct for creation of JIT device

sycl/test-e2e/KernelCompiler/sycl_device_flags.cpp

@@ -13,7 +13,8 @@
 // UNSUPPORTED-INTENDED:  IGC shader dump not available on Windows.
 
 // RUN: %{build} -o %t.out
-// RUN: env IGC_DumpToCustomDir=%T.dump IGC_ShaderDumpEnable=1 NEO_CACHE_PERSISTENT=0 %{run} %t.out %T.dump/
+// RUN: env IGC_DumpToCustomDir=%T.dump_0 IGC_ShaderDumpEnable=1 NEO_CACHE_PERSISTENT=0 %{run} %t.out sycl %T.dump_0/
+// RUN: env IGC_DumpToCustomDir=%T.dump_1 IGC_ShaderDumpEnable=1 NEO_CACHE_PERSISTENT=0 %{run} %t.out sycl_jit %T.dump_1/
 
 // clang-format off
 /*
@@ -105,28 +106,32 @@ int test_dump(std::string &dump_dir) {
 
 int main(int argc, char *argv[]) {
 
-  if (argc != 2) {
-    std::cerr << "Usage: " << argv[0] << " <dump_directory>" << std::endl;
+  if (argc != 3) {
+    std::cerr << "Usage: " << argv[0] << " <lang> <dump_directory>"
+              << std::endl;
     return 1;
   }
-  std::string dump_dir = argv[1];
 
   namespace syclex = sycl::ext::oneapi::experimental;
   using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
   using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;
 
+  syclex::source_language lang = std::strcmp(argv[1], "sycl_jit") == 0
+                                     ? syclex::source_language::sycl_jit
+                                     : syclex::source_language::sycl;
+  std::string dump_dir = argv[2];
+
   sycl::queue q;
   sycl::context ctx = q.get_context();
 
-  bool ok =
-      q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
+  bool ok = q.get_device().ext_oneapi_can_compile(lang);
   if (!ok) {
     std::cout << "compiling from SYCL source not supported" << std::endl;
     return 0; // if kernel compilation is not supported, do nothing.
   }
 
-  source_kb kbSrc = syclex::create_kernel_bundle_from_source(
-      ctx, syclex::source_language::sycl, SYCLSource);
+  source_kb kbSrc =
+      syclex::create_kernel_bundle_from_source(ctx, lang, SYCLSource);
 
   // Flags with and without space, inner quotes.
   std::vector<std::string> flags{"-Xs '-doubleGRF'",