diff --git a/sycl-jit/jit-compiler/include/KernelFusion.h b/sycl-jit/jit-compiler/include/KernelFusion.h index 76636beb937d9..e79124f016c68 100644 --- a/sycl-jit/jit-compiler/include/KernelFusion.h +++ b/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" { diff --git a/sycl-jit/jit-compiler/lib/KernelFusion.cpp b/sycl-jit/jit-compiler/lib/KernelFusion.cpp index fae9a3c29dcf4..7b706ba88a16b 100644 --- a/sycl-jit/jit-compiler/lib/KernelFusion.cpp +++ b/sycl-jit/jit-compiler/lib/KernelFusion.cpp @@ -244,7 +244,10 @@ compileSYCL(InMemoryFile SourceFile, View 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(ModuleOrErr.takeError(), "Device compilation failed"); @@ -254,7 +257,7 @@ compileSYCL(InMemoryFile SourceFile, View IncludeFiles, std::unique_ptr Module = std::move(*ModuleOrErr); Context.reset(&Module->getContext()); - if (auto Error = linkDeviceLibraries(*Module, UserArgList)) { + if (auto Error = linkDeviceLibraries(*Module, UserArgList, BuildLog)) { return errorTo(std::move(Error), "Device linking failed"); } @@ -274,7 +277,7 @@ compileSYCL(InMemoryFile SourceFile, View 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() { diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index 1d75136160e99..b41e07082701b 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -13,11 +13,15 @@ #include #include #include +#include #include #include +#include #include #include +#include +#include #include #include #include @@ -27,6 +31,10 @@ #include #include +#include +#include +#include + 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()); // Create the compiler's actual diagnostics engine. Compiler.createDiagnostics(DiagConsumer, /*ShouldOwnClient=*/false); @@ -161,12 +172,59 @@ struct GetLLVMModuleAction : public ToolAction { std::unique_ptr Module; }; +class ClangDiagnosticWrapper { + + llvm::raw_string_ostream LogStream; + + std::unique_ptr LogPrinter; + +public: + ClangDiagnosticWrapper(std::string &LogString, DiagnosticOptions *DiagOpts) + : LogStream(LogString), + LogPrinter( + std::make_unique(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> -jit_compiler::compileDeviceCode(InMemoryFile SourceFile, - View IncludeFiles, - const InputArgList &UserArgList) { +Expected> jit_compiler::compileDeviceCode( + InMemoryFile SourceFile, View 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 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 -getDeviceLibraries(const ArgList &Args, DiagnosticsEngine &Diags) { +static bool getDeviceLibraries(const ArgList &Args, + SmallVectorImpl &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 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 DiagID{new DiagnosticIDs}; IntrusiveRefCntPtr 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 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(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()); } } diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h index 8aa47939e3b1d..c378bb695a8f5 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h @@ -17,15 +17,18 @@ #include #include +#include namespace jit_compiler { llvm::Expected> compileDeviceCode(InMemoryFile SourceFile, View 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 performPostLink(llvm::Module &Module, diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index c366996aadc9b..ce498356e6191 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/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()); } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index 563f75d313e95..3a3e402ede786 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -10,8 +10,8 @@ // UNSUPPORTED: accelerator // RUN: %{build} -o %t.out -// RUN: %{run} %t.out -// RUN: %{l0_leak_check} %{run} %t.out +// RUN: %{run} %t.out 1 +// RUN: %{l0_leak_check} %{run} %t.out 1 // -- Test again, with caching. @@ -77,6 +77,34 @@ void ff_templated(T *ptr, T *unused) { } )==="; +auto constexpr BadSource = R"===( +#include + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_cp(int *ptr) { + + sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + + sycl::id<1> GId = Item.get_global_id() + no semi colon !! + ptr[GId.get(0)] = GId.get(0) + 41; +} +)==="; + +auto constexpr WarningSource = R"===( +#include + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_cp(int *ptr) { + + // intentionally using deprecated routine, as opposed to this_work_item::get_nd_item<1>() + // to provoke a warning. + sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>(); + + sycl::id<1> GId = Item.get_global_id(); + ptr[GId.get(0)] = GId.get(0) + 41; +} +)==="; + void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) { constexpr int Range = 10; int *usmPtr = sycl::malloc_shared(Range, Queue); @@ -211,10 +239,66 @@ int test_unsupported_options() { return 0; } -int main() { +int test_error() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl_jit); + if (!ok) { + return 0; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl_jit, BadSource); + try { + exe_kb kbExe = syclex::build(kbSrc); + assert(false && "we should not be here"); + } catch (sycl::exception &e) { + // yas! + assert(e.code() == sycl::errc::build); + assert(std::string(e.what()).find( + "error: expected ';' at end of declaration") != + std::string::npos); + } + return 0; +} + +int test_warning() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl_jit); + if (!ok) { + return 0; + } + std::string build_log; + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl_jit, WarningSource); + exe_kb kbExe = + syclex::build(kbSrc, syclex::properties{syclex::save_log{&build_log}}); + bool found_warning = + (build_log.find("warning: 'this_nd_item<1>' is deprecated") != + std::string::npos); + assert(found_warning); + return 0; +} +int main(int argc, char **) { #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER - return test_build_and_run() || test_unsupported_options(); + int optional_tests = (argc > 1) ? test_warning() : 0; + return test_build_and_run() || test_unsupported_options() || test_error() || + optional_tests; #else static_assert(false, "Kernel Compiler feature test macro undefined"); #endif