[SYCL][RTC] Add support for errors and build log in JIT-based RTC (#16132)

Adds support for correct error message routing and the `save_log`
property to the SYCL-RTC implementation based on the SYCL-JIT library.

If the compilation fails, any diagnostic messages encountered during
runtime compilation are stored in the `what` of the exception thrown.

If the user passes the `save_log` property defined in the
`kernel_compiler` extension, all diagnostics encountered during
compilation (errors, warnings, notes and remarks) are stored in the
user-provided string.

---------

Signed-off-by: Lukas Sommer <[email protected]>

Author: sommerlukas

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

@@ -58,18 +58,15 @@ class JITResult {
 
 class RTCResult {
 public:
-  explicit RTCResult(const char *ErrorMessage)
-      : Failed{true}, BundleInfo{}, ErrorMessage{ErrorMessage} {}
+  explicit RTCResult(const char *BuildLog)
+      : Failed{true}, BundleInfo{}, BuildLog{BuildLog} {}
 
-  explicit RTCResult(RTCBundleInfo &&BundleInfo)
-      : Failed{false}, BundleInfo{std::move(BundleInfo)}, ErrorMessage{} {}
+  RTCResult(RTCBundleInfo &&BundleInfo, const char *BuildLog)
+      : Failed{false}, BundleInfo{std::move(BundleInfo)}, BuildLog{BuildLog} {}
 
   bool failed() const { return Failed; }
 
-  const char *getErrorMessage() const {
-    assert(failed() && "No error message present");
-    return ErrorMessage.c_str();
-  }
+  const char *getBuildLog() const { return BuildLog.c_str(); }
 
   const RTCBundleInfo &getBundleInfo() const {
     assert(!failed() && "No bundle info");
@@ -79,7 +76,7 @@ class RTCResult {
 private:
   bool Failed;
   RTCBundleInfo BundleInfo;
-  sycl::detail::string ErrorMessage;
+  sycl::detail::string BuildLog;
 };
 
 extern "C" {

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

@@ -244,7 +244,10 @@ compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
   }
   llvm::opt::InputArgList UserArgList = std::move(*UserArgListOrErr);
 
-  auto ModuleOrErr = compileDeviceCode(SourceFile, IncludeFiles, UserArgList);
+  std::string BuildLog;
+
+  auto ModuleOrErr =
+      compileDeviceCode(SourceFile, IncludeFiles, UserArgList, BuildLog);
   if (!ModuleOrErr) {
     return errorTo<RTCResult>(ModuleOrErr.takeError(),
                               "Device compilation failed");
@@ -254,7 +257,7 @@ compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
   std::unique_ptr<llvm::Module> Module = std::move(*ModuleOrErr);
   Context.reset(&Module->getContext());
 
-  if (auto Error = linkDeviceLibraries(*Module, UserArgList)) {
+  if (auto Error = linkDeviceLibraries(*Module, UserArgList, BuildLog)) {
     return errorTo<RTCResult>(std::move(Error), "Device linking failed");
   }
 
@@ -274,7 +277,7 @@ compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
   }
   BundleInfo.BinaryInfo = std::move(*BinaryInfoOrError);
 
-  return RTCResult{std::move(BundleInfo)};
+  return RTCResult{std::move(BundleInfo), BuildLog.c_str()};
 }
 
 extern "C" KF_EXPORT_SYMBOL void resetJITConfiguration() {

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

@@ -13,11 +13,15 @@
 #include <clang/CodeGen/CodeGenAction.h>
 #include <clang/Driver/Compilation.h>
 #include <clang/Driver/Options.h>
+#include <clang/Frontend/ChainedDiagnosticConsumer.h>
 #include <clang/Frontend/CompilerInstance.h>
 #include <clang/Frontend/TextDiagnosticBuffer.h>
+#include <clang/Frontend/TextDiagnosticPrinter.h>
 #include <clang/Tooling/CompilationDatabase.h>
 #include <clang/Tooling/Tooling.h>
 
+#include <llvm/IR/DiagnosticInfo.h>
+#include <llvm/IR/DiagnosticPrinter.h>
 #include <llvm/IR/PassInstrumentation.h>
 #include <llvm/IR/PassManager.h>
 #include <llvm/IRReader/IRReader.h>
@@ -27,6 +31,10 @@
 #include <llvm/SYCLLowerIR/SYCLJointMatrixTransform.h>
 #include <llvm/Support/PropertySetIO.h>
 
+#include <algorithm>
+#include <array>
+#include <sstream>
+
 using namespace clang;
 using namespace clang::tooling;
 using namespace clang::driver;
@@ -132,6 +140,9 @@ struct GetLLVMModuleAction : public ToolAction {
     CompilerInstance Compiler(std::move(PCHContainerOps));
     Compiler.setInvocation(std::move(Invocation));
     Compiler.setFileManager(Files);
+    // Suppress summary with number of warnings and errors being printed to
+    // stdout.
+    Compiler.setVerboseOutputStream(std::make_unique<llvm::raw_null_ostream>());
 
     // Create the compiler's actual diagnostics engine.
     Compiler.createDiagnostics(DiagConsumer, /*ShouldOwnClient=*/false);
@@ -161,12 +172,59 @@ struct GetLLVMModuleAction : public ToolAction {
   std::unique_ptr<llvm::Module> Module;
 };
 
+class ClangDiagnosticWrapper {
+
+  llvm::raw_string_ostream LogStream;
+
+  std::unique_ptr<clang::TextDiagnosticPrinter> LogPrinter;
+
+public:
+  ClangDiagnosticWrapper(std::string &LogString, DiagnosticOptions *DiagOpts)
+      : LogStream(LogString),
+        LogPrinter(
+            std::make_unique<TextDiagnosticPrinter>(LogStream, DiagOpts)) {}
+
+  clang::TextDiagnosticPrinter *consumer() { return LogPrinter.get(); }
+
+  llvm::raw_ostream &stream() { return LogStream; }
+};
+
+class LLVMDiagnosticWrapper : public llvm::DiagnosticHandler {
+  llvm::raw_string_ostream LogStream;
+
+  DiagnosticPrinterRawOStream LogPrinter;
+
+public:
+  LLVMDiagnosticWrapper(std::string &BuildLog)
+      : LogStream(BuildLog), LogPrinter(LogStream) {}
+
+  bool handleDiagnostics(const DiagnosticInfo &DI) override {
+    auto Prefix = [](DiagnosticSeverity Severity) -> llvm::StringLiteral {
+      switch (Severity) {
+      case llvm::DiagnosticSeverity::DS_Error:
+        return "ERROR";
+      case llvm::DiagnosticSeverity::DS_Warning:
+        return "WARNING";
+      case llvm::DiagnosticSeverity::DS_Note:
+        return "NOTE:";
+      case llvm::DiagnosticSeverity::DS_Remark:
+        return "REMARK:";
+      default:
+        llvm_unreachable("Unhandled case");
+      }
+    }(DI.getSeverity());
+    LogPrinter << Prefix;
+    DI.print(LogPrinter);
+    LogPrinter << "\n";
+    return true;
+  }
+};
+
 } // anonymous namespace
 
-Expected<std::unique_ptr<llvm::Module>>
-jit_compiler::compileDeviceCode(InMemoryFile SourceFile,
-                                View<InMemoryFile> IncludeFiles,
-                                const InputArgList &UserArgList) {
+Expected<std::unique_ptr<llvm::Module>> jit_compiler::compileDeviceCode(
+    InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
+    const InputArgList &UserArgList, std::string &BuildLog) {
   const std::string &DPCPPRoot = getDPCPPRoot();
   if (DPCPPRoot == InvalidDPCPPRoot) {
     return createStringError("Could not locate DPCPP root directory");
@@ -197,6 +255,12 @@ jit_compiler::compileDeviceCode(InMemoryFile SourceFile,
   FixedCompilationDatabase DB{".", CommandLine};
   ClangTool Tool{DB, {SourceFile.Path}};
 
+  IntrusiveRefCntPtr<DiagnosticOptions> DiagOpts{new DiagnosticOptions};
+  ClangDiagnosticWrapper Wrapper(BuildLog, DiagOpts.get());
+  Tool.setDiagnosticConsumer(Wrapper.consumer());
+  // Suppress message "Error while processing" being printed to stdout.
+  Tool.setPrintErrorMessage(false);
+
   // Set up in-memory filesystem.
   Tool.mapVirtualFile(SourceFile.Path, SourceFile.Contents);
   for (const auto &IF : IncludeFiles) {
@@ -222,15 +286,15 @@ jit_compiler::compileDeviceCode(InMemoryFile SourceFile,
     return std::move(Action.Module);
   }
 
-  // TODO: Capture compiler errors from the ClangTool.
-  return createStringError("Unable to obtain LLVM module");
+  return createStringError(BuildLog);
 }
 
 // This function is a simplified copy of the device library selection process in
 // `clang::driver::tools::SYCL::getDeviceLibraries`, assuming a SPIR-V target
 // (no AoT, no third-party GPUs, no native CPU). Keep in sync!
-static SmallVector<std::string, 8>
-getDeviceLibraries(const ArgList &Args, DiagnosticsEngine &Diags) {
+static bool getDeviceLibraries(const ArgList &Args,
+                               SmallVectorImpl<std::string> &LibraryList,
+                               DiagnosticsEngine &Diags) {
   struct DeviceLibOptInfo {
     StringRef DeviceLibName;
     StringRef DeviceLibOption;
@@ -247,6 +311,8 @@ getDeviceLibraries(const ArgList &Args, DiagnosticsEngine &Diags) {
   // libraries cannot be affected via -fno-sycl-device-lib.
   bool ExcludeDeviceLibs = false;
 
+  bool FoundUnknownLib = false;
+
   if (Arg *A = Args.getLastArg(OPT_fsycl_device_lib_EQ,
                                OPT_fno_sycl_device_lib_EQ)) {
     if (A->getValues().size() == 0) {
@@ -268,6 +334,7 @@ getDeviceLibraries(const ArgList &Args, DiagnosticsEngine &Diags) {
         if (LinkInfoIter == DeviceLibLinkInfo.end() || Val == "internal") {
           Diags.Report(diag::err_drv_unsupported_option_argument)
               << A->getSpelling() << Val;
+          FoundUnknownLib = true;
         }
         DeviceLibLinkInfo[Val] = !ExcludeDeviceLibs;
       }
@@ -292,7 +359,6 @@ getDeviceLibraries(const ArgList &Args, DiagnosticsEngine &Diags) {
       {"libsycl-itt-compiler-wrappers", "internal"},
       {"libsycl-itt-stubs", "internal"}};
 
-  SmallVector<std::string, 8> LibraryList;
   StringRef LibSuffix = ".bc";
   auto AddLibraries = [&](const SYCLDeviceLibsList &LibsList) {
     for (const DeviceLibOptInfo &Lib : LibsList) {
@@ -312,37 +378,33 @@ getDeviceLibraries(const ArgList &Args, DiagnosticsEngine &Diags) {
     AddLibraries(SYCLDeviceAnnotationLibs);
   }
 
-  return LibraryList;
+  return FoundUnknownLib;
 }
 
 Error jit_compiler::linkDeviceLibraries(llvm::Module &Module,
-                                        const InputArgList &UserArgList) {
+                                        const InputArgList &UserArgList,
+                                        std::string &BuildLog) {
   const std::string &DPCPPRoot = getDPCPPRoot();
   if (DPCPPRoot == InvalidDPCPPRoot) {
     return createStringError("Could not locate DPCPP root directory");
   }
 
-  // TODO: Seems a bit excessive to set up this machinery for one warning and
-  //       one error. Rethink when implementing the build log/error reporting as
-  //       mandated by the extension.
   IntrusiveRefCntPtr<DiagnosticIDs> DiagID{new DiagnosticIDs};
   IntrusiveRefCntPtr<DiagnosticOptions> DiagOpts{new DiagnosticOptions};
-  TextDiagnosticBuffer *DiagBuffer = new TextDiagnosticBuffer;
-  DiagnosticsEngine Diags(DiagID, DiagOpts, DiagBuffer);
-
-  auto LibNames = getDeviceLibraries(UserArgList, Diags);
-  if (std::distance(DiagBuffer->err_begin(), DiagBuffer->err_end()) > 0) {
-    std::string DiagMsg;
-    raw_string_ostream SOS{DiagMsg};
-    interleave(
-        DiagBuffer->err_begin(), DiagBuffer->err_end(),
-        [&](const auto &D) { SOS << D.second; }, [&]() { SOS << '\n'; });
+  ClangDiagnosticWrapper Wrapper(BuildLog, DiagOpts.get());
+  DiagnosticsEngine Diags(DiagID, DiagOpts, Wrapper.consumer(),
+                          /* ShouldOwnClient=*/false);
+
+  SmallVector<std::string> LibNames;
+  bool FoundUnknownLib = getDeviceLibraries(UserArgList, LibNames, Diags);
+  if (FoundUnknownLib) {
     return createStringError("Could not determine list of device libraries: %s",
-                             DiagMsg.c_str());
+                             BuildLog.c_str());
   }
-  // TODO: Add warnings to build log.
 
   LLVMContext &Context = Module.getContext();
+  Context.setDiagnosticHandler(
+      std::make_unique<LLVMDiagnosticWrapper>(BuildLog));
   for (const std::string &LibName : LibNames) {
     std::string LibPath = DPCPPRoot + "/lib/" + LibName;
 
@@ -356,10 +418,8 @@ Error jit_compiler::linkDeviceLibraries(llvm::Module &Module,
     }
 
     if (Linker::linkModules(Module, std::move(Lib), Linker::LinkOnlyNeeded)) {
-      // TODO: Obtain detailed error message from the context's diagnostics
-      //       handler.
-      return createStringError("Unable to link device library: %s",
-                               LibPath.c_str());
+      return createStringError("Unable to link device library %s: %s",
+                               LibPath.c_str(), BuildLog.c_str());
     }
   }
 

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

@@ -17,15 +17,18 @@
 #include <llvm/Support/Error.h>
 
 #include <memory>
+#include <string>
 
 namespace jit_compiler {
 
 llvm::Expected<std::unique_ptr<llvm::Module>>
 compileDeviceCode(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
-                  const llvm::opt::InputArgList &UserArgList);
+                  const llvm::opt::InputArgList &UserArgList,
+                  std::string &BuildLog);
 
 llvm::Error linkDeviceLibraries(llvm::Module &Module,
-                                const llvm::opt::InputArgList &UserArgList);
+                                const llvm::opt::InputArgList &UserArgList,
+                                std::string &BuildLog);
 
 llvm::Expected<RTCBundleInfo>
 performPostLink(llvm::Module &Module,

sycl/source/detail/jit_compiler.cpp

@@ -1242,12 +1242,13 @@ sycl_device_binaries jit_compiler::compileSYCL(
 
   auto Result = CompileSYCLHandle(SourceFile, IncludeFilesView, UserArgsView);
 
-  if (Result.failed()) {
-    throw sycl::exception(sycl::errc::build, Result.getErrorMessage());
+  if (LogPtr) {
+    LogPtr->append(Result.getBuildLog());
   }
 
-  // TODO: We currently don't have a meaningful build log.
-  (void)LogPtr;
+  if (Result.failed()) {
+    throw sycl::exception(sycl::errc::build, Result.getBuildLog());
+  }
 
   return createDeviceBinaryImage(Result.getBundleInfo());
 }