-
Notifications
You must be signed in to change notification settings - Fork 796
[SYCL] PoC implementation of kernel compiler extension with libtooling and sycl-jit #15701
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Changes from 2 commits
6e892f9
246bc6d
46e7127
7e06ae1
ef90b4b
bfe9b43
2c2ac32
f5871cc
c60a528
ce333dc
0ad693d
7a928b3
c2eb448
407124b
8c4597e
4a5c809
863d4b5
b727bb1
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -14,6 +14,7 @@ | |
| #include "fusion/FusionPipeline.h" | ||
| #include "helper/ConfigHelper.h" | ||
| #include "helper/ErrorHandling.h" | ||
| #include "rtc/DeviceCompilation.h" | ||
| #include "translation/KernelTranslation.h" | ||
| #include "translation/SPIRVLLVMTranslation.h" | ||
| #include <llvm/Support/Error.h> | ||
|
|
@@ -235,6 +236,31 @@ extern "C" JITResult fuseKernels(View<SYCLKernelInfo> KernelInformation, | |
| return JITResult{FusedKernelInfo}; | ||
| } | ||
|
|
||
| extern "C" JITResult compileSYCL(const char *SYCLSource, | ||
| View<IncludePair> IncludePairs, | ||
| View<const char *> UserArgs, | ||
| const char *DPCPPRoot) { | ||
|
||
| std::unique_ptr<llvm::Module> Module = | ||
| compileDeviceCode(SYCLSource, IncludePairs, UserArgs, DPCPPRoot); | ||
| if (!Module) { | ||
| return JITResult{"Device code compilation failed"}; | ||
| } | ||
|
|
||
| SYCLKernelInfo Kernel; | ||
| auto Error = translation::KernelTranslator::translateKernel( | ||
| Kernel, *Module, JITContext::getInstance(), BinaryFormat::SPIRV); | ||
|
|
||
| auto *LLVMCtx = &Module->getContext(); | ||
| Module.reset(); | ||
| delete LLVMCtx; | ||
sommerlukas marked this conversation as resolved.
Show resolved
Hide resolved
|
||
|
|
||
| if (Error) { | ||
| return errorToFusionResult(std::move(Error), "SPIR-V translation failed"); | ||
| } | ||
|
|
||
| return JITResult{Kernel}; | ||
| } | ||
|
|
||
| extern "C" void resetJITConfiguration() { ConfigHelper::reset(); } | ||
|
|
||
| extern "C" void addToJITConfiguration(OptionStorage &&Opt) { | ||
|
|
||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,111 @@ | ||
| //==---------------------- DeviceCompilation.cpp ---------------------------==// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #include "DeviceCompilation.h" | ||
|
|
||
| #include <clang/Basic/Version.h> | ||
| #include <clang/CodeGen/CodeGenAction.h> | ||
| #include <clang/Driver/Compilation.h> | ||
| #include <clang/Frontend/CompilerInstance.h> | ||
| #include <clang/Tooling/CompilationDatabase.h> | ||
| #include <clang/Tooling/Tooling.h> | ||
|
|
||
| #include <llvm/IR/Module.h> | ||
|
|
||
| namespace { | ||
| using namespace clang; | ||
| using namespace clang::tooling; | ||
| using namespace clang::driver; | ||
|
|
||
| struct GetLLVMModuleAction : public ToolAction { | ||
| // Code adapted from `FrontendActionFactory::runInvocation`. | ||
| bool runInvocation(std::shared_ptr<CompilerInvocation> Invocation, | ||
| FileManager *Files, | ||
| std::shared_ptr<PCHContainerOperations> PCHContainerOps, | ||
| DiagnosticConsumer *DiagConsumer) override { | ||
| assert(!Module && "Action should only be invoked on a single file"); | ||
|
|
||
| // Create a compiler instance to handle the actual work. | ||
| CompilerInstance Compiler(std::move(PCHContainerOps)); | ||
| Compiler.setInvocation(std::move(Invocation)); | ||
| Compiler.setFileManager(Files); | ||
|
|
||
| // Create the compiler's actual diagnostics engine. | ||
| Compiler.createDiagnostics(DiagConsumer, /*ShouldOwnClient=*/false); | ||
| if (!Compiler.hasDiagnostics()) { | ||
| return false; | ||
| } | ||
|
|
||
| Compiler.createSourceManager(*Files); | ||
|
|
||
| // Ignore `Compiler.getFrontendOpts().ProgramAction` (would be `EmitBC`) and | ||
| // create/execute an `EmitLLVMOnlyAction` (= codegen to LLVM module without | ||
| // emitting anything) instead. | ||
| EmitLLVMOnlyAction ELOA; | ||
| const bool Success = Compiler.ExecuteAction(ELOA); | ||
| Files->clearStatCache(); | ||
| if (!Success) { | ||
| return false; | ||
| } | ||
|
|
||
| // Take the module and its context to extend the objects' lifetime. | ||
| Module = ELOA.takeModule(); | ||
| ELOA.takeLLVMContext(); | ||
|
|
||
| return true; | ||
| } | ||
|
|
||
| std::unique_ptr<llvm::Module> Module; | ||
| }; | ||
|
|
||
| } // anonymous namespace | ||
|
|
||
| std::unique_ptr<llvm::Module> jit_compiler::compileDeviceCode( | ||
| const char *SYCLSource, View<IncludePair> IncludePairs, | ||
| View<const char *> UserArgs, const char *DPCPPRoot) { | ||
|
|
||
| SmallVector<std::string> CommandLine = {"-fsycl-device-only"}; | ||
| // TODO: Allow instrumentation again when device library linking is | ||
| // implemented. | ||
| CommandLine.push_back("-fno-sycl-instrument-device-code"); | ||
| CommandLine.append(UserArgs.begin(), UserArgs.end()); | ||
| clang::tooling::FixedCompilationDatabase DB{"./", CommandLine}; | ||
|
|
||
| constexpr auto SourcePath = "rtc.cpp"; | ||
|
||
| clang::tooling::ClangTool Tool{DB, {SourcePath}}; | ||
|
|
||
| // Set up in-memory filesystem. | ||
| Tool.mapVirtualFile(SourcePath, SYCLSource); | ||
| for (const auto &IP : IncludePairs) { | ||
| Tool.mapVirtualFile(IP.Path, IP.Contents); | ||
| } | ||
|
|
||
| // Reset argument adjusters to drop the `-fsyntax-only` flag which is added by | ||
| // default by this API. | ||
| Tool.clearArgumentsAdjusters(); | ||
| // Then, modify argv[0] and set the resource directory so that the driver | ||
| // picks up the correct SYCL environment. | ||
| Tool.appendArgumentsAdjuster( | ||
| [&DPCPPRoot](const CommandLineArguments &Args, | ||
| StringRef Filename) -> CommandLineArguments { | ||
| (void)Filename; | ||
| CommandLineArguments NewArgs = Args; | ||
| NewArgs[0] = (Twine(DPCPPRoot) + "/bin/clang++").str(); | ||
| NewArgs.push_back((Twine("-resource-dir=") + DPCPPRoot + "/lib/clang/" + | ||
| Twine(CLANG_VERSION_MAJOR)) | ||
| .str()); | ||
|
Comment on lines
+134
to
+136
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Do these paths also apply in a packaged release? There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Yes, a release |
||
| return NewArgs; | ||
| }); | ||
|
|
||
| GetLLVMModuleAction Action; | ||
| if (!Tool.run(&Action)) { | ||
| return std::move(Action.Module); | ||
| } | ||
|
|
||
| return {}; | ||
| } | ||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,30 @@ | ||
| //==---- DeviceCompilation.h - Compile SYCL device code with libtooling ----==// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #ifndef SYCL_JIT_COMPILER_RTC_DEVICE_COMPILATION_H | ||
| #define SYCL_JIT_COMPILER_RTC_DEVICE_COMPILATION_H | ||
|
|
||
| #include "Kernel.h" | ||
| #include "View.h" | ||
|
|
||
| #include <memory> | ||
|
|
||
| namespace llvm { | ||
| class Module; | ||
|
||
| } // namespace llvm | ||
|
|
||
| namespace jit_compiler { | ||
|
|
||
| std::unique_ptr<llvm::Module> compileDeviceCode(const char *SYCLSource, | ||
| View<IncludePair> IncludePairs, | ||
| View<const char *> UserArgs, | ||
| const char *DPCPPRoot); | ||
|
|
||
| } // namespace jit_compiler | ||
|
|
||
| #endif // SYCL_JIT_COMPILER_RTC_DEVICE_COMPILATION_H | ||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -18,6 +18,9 @@ | |
| #include <sycl/detail/ur.hpp> | ||
| #include <sycl/kernel_bundle.hpp> | ||
|
|
||
| #include <dlfcn.h> | ||
| #include <link.h> | ||
|
|
||
| namespace sycl { | ||
| inline namespace _V1 { | ||
| namespace detail { | ||
|
|
@@ -74,6 +77,31 @@ jit_compiler::jit_compiler() { | |
| return false; | ||
| } | ||
|
|
||
| this->CompileSYCLHandle = reinterpret_cast<CompileSYCLFuncT>( | ||
| sycl::detail::ur::getOsLibraryFuncAddress(LibraryPtr, "compileSYCL")); | ||
| if (!this->CompileSYCLHandle) { | ||
| printPerformanceWarning( | ||
| "Cannot resolve JIT library function entry point"); | ||
|
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. This sounds more serious than a mere performance warning :) There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more.
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. The reasoning behind that name is that so far, when something went wrong in sycl-jit, we would simply fall back to non-fused execution of kernels (in case of kernel fusion), which still leads to correct results, but potentially lower performance than expected. But I'm happy to rename this function to something more suitable. |
||
| return false; | ||
| } | ||
|
|
||
| // TODO: Move this query to a more appropriate location (e.g. add | ||
| // `sycl::detail::ur::getOsLibraryPath`), and handle non-POSIX OSs. For now, | ||
| // it should be fine here because the JIT is not built on Windows. | ||
| link_map *Map = nullptr; | ||
| if (dlinfo(LibraryPtr, RTLD_DI_LINKMAP, &Map) == 0) { | ||
| std::string LoadedLibraryPath = Map->l_name; | ||
| std::string JITLibraryPathSuffix = "/lib/" + JITLibraryName; | ||
| auto Pos = LoadedLibraryPath.rfind(JITLibraryPathSuffix); | ||
| if (Pos != std::string::npos) { | ||
| this->DPCPPRoot = LoadedLibraryPath.substr(0, Pos); | ||
| } | ||
| } | ||
| if (this->DPCPPRoot.empty()) { | ||
| printPerformanceWarning("Cannot determine JIT library location"); | ||
| return false; | ||
| } | ||
|
|
||
| return true; | ||
| }; | ||
| Available = checkJITLibrary(); | ||
|
|
@@ -1143,6 +1171,45 @@ std::vector<uint8_t> jit_compiler::encodeReqdWorkGroupSize( | |
| return Encoded; | ||
| } | ||
|
|
||
| std::vector<uint8_t> jit_compiler::compileSYCL( | ||
| const std::string &SYCLSource, | ||
| const std::vector<std::pair<std::string, std::string>> &IncludePairs, | ||
| const std::vector<std::string> &UserArgs, std::string *LogPtr, | ||
| const std::vector<std::string> &RegisteredKernelNames) { | ||
|
|
||
| // TODO: Handle situation. | ||
| assert(RegisteredKernelNames.empty() && | ||
| "Instantiation of kernel templates NYI"); | ||
|
||
|
|
||
| std::vector<::jit_compiler::IncludePair> IncludePairsView; | ||
| IncludePairsView.reserve(IncludePairs.size()); | ||
| std::transform(IncludePairs.begin(), IncludePairs.end(), | ||
| std::back_inserter(IncludePairsView), [](const auto &Pair) { | ||
| return ::jit_compiler::IncludePair{Pair.first.c_str(), | ||
| Pair.second.c_str()}; | ||
| }); | ||
| std::vector<const char *> UserArgsView; | ||
| UserArgsView.reserve(UserArgs.size()); | ||
| std::transform(UserArgs.begin(), UserArgs.end(), | ||
| std::back_inserter(UserArgsView), | ||
| [](const auto &Arg) { return Arg.c_str(); }); | ||
|
|
||
| auto Result = CompileSYCLHandle(SYCLSource.c_str(), IncludePairsView, | ||
| UserArgsView, DPCPPRoot.c_str()); | ||
|
|
||
| if (Result.failed()) { | ||
| throw sycl::exception(sycl::errc::build, Result.getErrorMessage()); | ||
| } | ||
|
|
||
| // TODO: We currently don't have a meaningful build log. | ||
| (void)LogPtr; | ||
|
|
||
| const auto &BI = Result.getKernelInfo().BinaryInfo; | ||
| assert(BI.Format == ::jit_compiler::BinaryFormat::SPIRV); | ||
| std::vector<uint8_t> SPV(BI.BinaryStart, BI.BinaryStart + BI.BinarySize); | ||
| return SPV; | ||
| } | ||
|
|
||
| } // namespace detail | ||
| } // namespace _V1 | ||
| } // namespace sycl | ||
|
|
||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Does LLVM CMake not define a variable for that?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Changed this to use
${LLVM_EXTERNAL_CLANG_SOURCE_DIR}.Unfortunately there doesn't seem to be an equivalent for the build directory, from which an
.incfile that definesCLANG_VERSION_MAJORis included.