[SYCL][RTC] Rework handling of build_options (#17405)

Reject user-supplied arguments a bit more coarsely based on their kind
or relation to an unsupported feature.

---------

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

Author: jopperm

sycl-jit/jit-compiler/include/KernelFusion.h

@@ -86,32 +86,37 @@ class RTCHashResult {
   sycl::detail::string HashOrLog;
 };
 
+enum class RTCErrorCode { SUCCESS, BUILD, INVALID };
+
 class RTCResult {
 public:
-  explicit RTCResult(const char *BuildLog)
-      : Failed{true}, BundleInfo{}, BuildLog{BuildLog} {}
+  explicit RTCResult(const char *BuildLog,
+                     RTCErrorCode ErrorCode = RTCErrorCode::BUILD)
+      : ErrorCode{ErrorCode}, BundleInfo{}, BuildLog{BuildLog} {
+    assert(ErrorCode != RTCErrorCode::SUCCESS);
+  }
 
   RTCResult(RTCBundleInfo &&BundleInfo, RTCDeviceCodeIR &&DeviceCodeIR,
             const char *BuildLog)
-      : Failed{false}, BundleInfo{std::move(BundleInfo)},
+      : ErrorCode{RTCErrorCode::SUCCESS}, BundleInfo{std::move(BundleInfo)},
         DeviceCodeIR(std::move(DeviceCodeIR)), BuildLog{BuildLog} {}
 
-  bool failed() const { return Failed; }
+  RTCErrorCode getErrorCode() const { return ErrorCode; }
 
   const char *getBuildLog() const { return BuildLog.c_str(); }
 
   const RTCBundleInfo &getBundleInfo() const {
-    assert(!failed() && "No bundle info");
+    assert(ErrorCode == RTCErrorCode::SUCCESS && "No bundle info");
     return BundleInfo;
   }
 
   const RTCDeviceCodeIR &getDeviceCodeIR() const {
-    assert(!failed() && "No device code IR");
+    assert(ErrorCode == RTCErrorCode::SUCCESS && "No device code IR");
     return DeviceCodeIR;
   }
 
 private:
-  bool Failed;
+  RTCErrorCode ErrorCode;
   RTCBundleInfo BundleInfo;
   RTCDeviceCodeIR DeviceCodeIR;
   sycl::detail::string BuildLog;

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

@@ -45,11 +45,12 @@ static std::string formatError(llvm::Error &&Err, const std::string &Msg) {
   return ErrMsg.str();
 }
 
-template <typename ResultType>
-static ResultType errorTo(llvm::Error &&Err, const std::string &Msg) {
+template <typename ResultType, typename... ExtraArgTypes>
+static ResultType errorTo(llvm::Error &&Err, const std::string &Msg,
+                          ExtraArgTypes... ExtraArgs) {
   // Cannot throw an exception here if LLVM itself is compiled without exception
   // support.
-  return ResultType{formatError(std::move(Err), Msg).c_str()};
+  return ResultType{formatError(std::move(Err), Msg).c_str(), ExtraArgs...};
 }
 
 static std::vector<jit_compiler::NDRange>
@@ -288,7 +289,8 @@ compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
   auto UserArgListOrErr = parseUserArgs(UserArgs);
   if (!UserArgListOrErr) {
     return errorTo<RTCResult>(UserArgListOrErr.takeError(),
-                              "Parsing of user arguments failed");
+                              "Parsing of user arguments failed",
+                              RTCErrorCode::INVALID);
   }
   llvm::opt::InputArgList UserArgList = std::move(*UserArgListOrErr);
 

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

@@ -321,20 +321,15 @@ static void adjustArgs(const InputArgList &UserArgList,
   DAL.AddJoinedArg(
       nullptr, OptTable.getOption(OPT_resource_dir_EQ),
       (DPCPPRoot + "/lib/clang/" + Twine(CLANG_VERSION_MAJOR)).str());
-  for (auto *Arg : UserArgList) {
-    DAL.append(Arg);
-  }
-  // Remove args that will trigger an unused command line argument warning for
-  // the FrontendAction invocation, but are handled later (e.g. during device
-  // linking).
-  DAL.eraseArg(OPT_fsycl_device_lib_EQ);
-  DAL.eraseArg(OPT_fno_sycl_device_lib_EQ);
-  DAL.eraseArg(OPT_ftime_trace_EQ);
-  DAL.eraseArg(OPT_ftime_trace_granularity_EQ);
-  DAL.eraseArg(OPT_ftime_trace_verbose);
+  // User args may contain options not intended for the frontend, but we can't
+  // claim them here to tell the driver they're used later. Hence, suppress the
+  // unused argument warning.
+  DAL.AddFlagArg(nullptr, OptTable.getOption(OPT_Qunused_arguments));
 
   ArgStringList ASL;
   for_each(DAL, [&DAL, &ASL](Arg *A) { A->render(DAL, ASL); });
+  for_each(UserArgList,
+           [&UserArgList, &ASL](Arg *A) { A->render(UserArgList, ASL); });
   transform(ASL, std::back_inserter(CommandLine),
             [](const char *AS) { return std::string{AS}; });
 }
@@ -811,51 +806,21 @@ jit_compiler::parseUserArgs(View<const char *> UserArgs) {
         UserArgsRef[MissingArgIndex], MissingArgIndex);
   }
 
-  // Check for unsupported options.
-  // TODO: There are probably more, e.g. requesting non-SPIR-V targets.
-  {
-    // -fsanitize=address
-    bool IsDeviceAsanEnabled = false;
-    if (Arg *A = AL.getLastArg(OPT_fsanitize_EQ, OPT_fno_sanitize_EQ)) {
-      if (A->getOption().matches(OPT_fsanitize_EQ) &&
-          A->getValues().size() == 1) {
-        std::string SanitizeVal = A->getValue();
-        IsDeviceAsanEnabled = SanitizeVal == "address";
-      }
-    } else {
-      // User can pass -fsanitize=address to device compiler via
-      // -Xsycl-target-frontend.
-      auto SyclFEArg = AL.getAllArgValues(OPT_Xsycl_frontend);
-      IsDeviceAsanEnabled = (std::count(SyclFEArg.begin(), SyclFEArg.end(),
-                                        "-fsanitize=address") > 0);
-      if (!IsDeviceAsanEnabled) {
-        auto SyclFEArgEq = AL.getAllArgValues(OPT_Xsycl_frontend_EQ);
-        IsDeviceAsanEnabled =
-            (std::count(SyclFEArgEq.begin(), SyclFEArgEq.end(),
-                        "-fsanitize=address") > 0);
-      }
-
-      // User can also enable asan for SYCL device via -Xarch_device option.
-      if (!IsDeviceAsanEnabled) {
-        auto DeviceArchVals = AL.getAllArgValues(OPT_Xarch_device);
-        for (auto DArchVal : DeviceArchVals) {
-          if (DArchVal.find("-fsanitize=address") != std::string::npos) {
-            IsDeviceAsanEnabled = true;
-            break;
-          }
-        }
-      }
-    }
-
-    if (IsDeviceAsanEnabled) {
-      return createStringError(
-          "Device ASAN is not supported for runtime compilation");
-    }
-  }
-
-  if (!AL.hasFlag(OPT_fsycl_device_code_split_esimd,
-                  OPT_fno_sycl_device_code_split_esimd, true)) {
-    return createStringError("ESIMD device code split cannot be deactivated");
+  // Check for options that are unsupported because they would interfere with
+  // the in-memory pipeline.
+  Arg *UnsupportedArg =
+      AL.getLastArg(OPT_Action_Group,     // Actions like -c or -S
+                    OPT_Link_Group,       // Linker flags
+                    OPT_o,                // Output file
+                    OPT_fsycl_targets_EQ, // AoT compilation
+                    OPT_fsycl_link_EQ,    // SYCL linker
+                    OPT_fno_sycl_device_code_split_esimd, // invoke_simd
+                    OPT_fsanitize_EQ                      // Sanitizer
+      );
+  if (UnsupportedArg) {
+    return createStringError(
+        "Option '%s' is not supported for SYCL runtime compilation",
+        UnsupportedArg->getAsString(AL).c_str());
   }
 
   return std::move(AL);
@@ -866,20 +831,14 @@ void jit_compiler::encodeBuildOptions(RTCBundleInfo &BundleInfo,
   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;
+  for (Arg *A : UserArgList.filtered(OPT_Xs, OPT_Xs_separate)) {
+    if (!CompileOptions.empty()) {
+      COSOS << ' ';
     }
-
-    // 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();
+    if (A->getOption().matches(OPT_Xs)) {
+      COSOS << '-';
     }
-
-    COSOS << (CompileOptions.empty() ? "" : " ") << AV;
+    COSOS << A->getValue();
   }
 
   if (!CompileOptions.empty()) {

sycl/source/detail/jit_compiler.cpp

@@ -1306,9 +1306,16 @@ std::pair<sycl_device_binaries, std::string> jit_compiler::compileSYCL(
   auto Result = CompileSYCLHandle(SourceFile, IncludeFilesView, UserArgsView,
                                   CachedIR, /*SaveIR=*/!CacheKey.empty());
 
-  appendToLog(Result.getBuildLog());
-  if (Result.failed()) {
-    throw sycl::exception(sycl::errc::build, Result.getBuildLog());
+  const char *BuildLog = Result.getBuildLog();
+  appendToLog(BuildLog);
+  switch (Result.getErrorCode()) {
+    using RTCErrC = ::jit_compiler::RTCErrorCode;
+  case RTCErrC::BUILD:
+    throw sycl::exception(sycl::errc::build, BuildLog);
+  case RTCErrC::INVALID:
+    throw sycl::exception(sycl::errc::invalid, BuildLog);
+  default: // RTCErrC::SUCCESS
+    break;
   }
 
   const auto &IR = Result.getDeviceCodeIR();

sycl/test-e2e/KernelCompiler/sycl.cpp

@@ -454,23 +454,17 @@ int test_unsupported_options(sycl::queue q) {
       ctx, syclex::source_language::sycl, "");
   std::vector<sycl::device> devs = kbSrc.get_devices();
 
-  auto CheckUnsupported = [&](const std::vector<std::string> &flags) {
-    try {
-      syclex::build(kbSrc, devs,
-                    syclex::properties{syclex::build_options{flags}});
-      assert(false && "unsupported option not detected");
-    } catch (sycl::exception &e) {
-      assert(e.code() == sycl::errc::build);
-      assert(std::string(e.what()).find("Parsing of user arguments failed") !=
-             std::string::npos);
-    }
-  };
-
-  CheckUnsupported({"-fsanitize=address"});
-  CheckUnsupported({"-Xsycl-target-frontend", "-fsanitize=address"});
-  CheckUnsupported({"-Xsycl-target-frontend=spir64", "-fsanitize=address"});
-  CheckUnsupported({"-Xarch_device", "-fsanitize=address"});
-  CheckUnsupported({"-fno-sycl-device-code-split-esimd"});
+  try {
+    // Don't attempt to test exhaustively here...
+    syclex::build(kbSrc, devs,
+                  syclex::properties{
+                      syclex::build_options{"-fsycl-targets=intel_gpu_pvc"}});
+    assert(false && "unsupported option not detected");
+  } catch (sycl::exception &e) {
+    assert(e.code() == sycl::errc::invalid);
+    assert(std::string(e.what()).find("Parsing of user arguments failed") !=
+           std::string::npos);
+  }
 
   return 0;
 }

sycl/test-e2e/KernelCompiler/sycl_device_flags.cpp

@@ -129,9 +129,8 @@ int main(int argc, char *argv[]) {
   source_kb kbSrc = syclex::create_kernel_bundle_from_source(
       ctx, syclex::source_language::sycl, SYCLSource);
 
-  // Flags with and without space, inner quotes.
-  std::vector<std::string> flags{"-Xs '-doubleGRF'",
-                                 "-Xs'-Xfinalizer \"-printregusage\"'"};
+  std::vector<std::string> flags{"-Xs", "-doubleGRF",
+                                 "-XsXfinalizer \"-printregusage\""};
   exe_kb kbExe =
       syclex::build(kbSrc, syclex::properties{syclex::build_options{flags}});