-
Notifications
You must be signed in to change notification settings - Fork 733
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
[SYCL] PoC implementation of kernel compiler extension with libtooling and sycl-jit #15701
base: sycl
Are you sure you want to change the base?
Changes from 5 commits
6e892f9
246bc6d
46e7127
7e06ae1
ef90b4b
bfe9b43
2c2ac32
f5871cc
c60a528
ce333dc
0ad693d
7a928b3
c2eb448
407124b
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; | ||
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. How expensive is it to set up and destroy the 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. If I'm not mistaken, certain things like constants and metadata are stored within At least that is the behavior we discovered a few years ago when we were debugging exceptionally huge memory footprint of 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. It's certainly possible to pass an existing context into the ToolAction, but that also raises questions of thread safety. For the performance implications, a) yes, looks like setting up a context does involve a non-trivial amount of work, and b) still seems true that types, constants and metadata are allocated in the context and not freed when the module is destroyed. I'd propose to keep the simple implementation for now, and will look out for the context setup overhead once we start benchmarking the RTC. |
||
|
||
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"; | ||
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. rather than hard-code "rtc.cpp" in, can this 'fantasy name' for the file be an argument to the API? It might show up in debug information, so it might be useful to users to be able to disambiguate, rather than having every dynamic device compiled kernel originate with the same fictional file name. 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. Good idea, will do 👍 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. Done. I'm passing in a semi-random ID, same as the file-based implementation. There's no property yet in the extension to specify a file name or prefix, correct? 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. What would be the reason that a user might want to change this name? Why is it better to have a semi-random ID as a default rather than a fixed string? I can imagine that the name might show up in error / log messages, for example, when there is a syntax error in the source string. If that's the only case the name is visible, it seems like having a fixed string like I'm not opposed to adding a property which allows the user to set this name, but I think it should be an optional property because I think many people will not care what the name is. 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, that's the only case. I agree that the ID doesn't add much value here because we don't materialise anything on the actual filesystem. I'll keep the plumbing to pass the filename down to the JIT, but will set it |
||
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
+136
to
+138
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 |
---|---|---|
|
@@ -20,7 +20,12 @@ enum class bundle_state : char { | |
|
||
namespace ext::oneapi::experimental { | ||
|
||
enum class source_language : int { opencl = 0, spirv = 1, sycl = 2 /* cuda */ }; | ||
enum class source_language : int { | ||
opencl = 0, | ||
spirv = 1, | ||
sycl = 2 /* cuda */, | ||
sycljit = 99 | ||
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 is part of a public interface. Can we document it somewhere? Also, I don't think you need to make a jump, as long as we don't change it after it's merged. Side note, I personally prefer 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 is only temporary until we complete functionality in follow-up PRs, so not intended to ever be exposed users. 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. I renamed the enum value and added a comment that this is temporary. |
||
}; | ||
|
||
// opencl versions | ||
struct cl_version { | ||
|
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.
|
||
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"); | ||
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. Should we throw instead? In cases where asserts are disabled what would happen if execution continues from here? |
||
|
||
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.
DPCPPRoot is not like the other arguments. Doesn't it seem like the routine should be able to figure that out itself, rather than being provided?
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.
Makes sense; it's detected from inside the JIT library now.