From a726365eecf2e3a871d5884c8dc5f49016dc61b3 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 18 Nov 2024 09:17:25 +0000 Subject: [PATCH 01/12] WIP: Symbols and properties available in jit_compiler (runtime) Signed-off-by: Julian Oppermann --- sycl-jit/common/include/Kernel.h | 50 ++++++++ sycl-jit/jit-compiler/CMakeLists.txt | 2 + sycl-jit/jit-compiler/include/KernelFusion.h | 28 ++++- sycl-jit/jit-compiler/lib/KernelFusion.cpp | 47 ++++--- .../lib/rtc/DeviceCompilation.cpp | 115 ++++++++++++++++- .../jit-compiler/lib/rtc/DeviceCompilation.h | 4 + .../jit-compiler/lib/rtc/PostLinkActions.cpp | 116 ++++++++++++++++++ .../jit-compiler/lib/rtc/PostLinkActions.h | 50 ++++++++ .../lib/translation/KernelTranslation.cpp | 14 +++ .../lib/translation/KernelTranslation.h | 3 + sycl/source/detail/jit_compiler.cpp | 2 +- 11 files changed, 413 insertions(+), 18 deletions(-) create mode 100644 sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp create mode 100644 sycl-jit/jit-compiler/lib/rtc/PostLinkActions.h diff --git a/sycl-jit/common/include/Kernel.h b/sycl-jit/common/include/Kernel.h index bbb9a802995ff..794126a462e20 100644 --- a/sycl-jit/common/include/Kernel.h +++ b/sycl-jit/common/include/Kernel.h @@ -17,6 +17,7 @@ #include #include #include +#include #include namespace jit_compiler { @@ -350,11 +351,60 @@ struct SYCLKernelInfo { : Name{KernelName}, Args{NumArgs}, Attributes{}, NDR{}, BinaryInfo{} {} }; +// RTC-related datastructures +// TODO: Consider moving into separate header. + struct InMemoryFile { const char *Path; const char *Contents; }; +using RTCBundleBinaryInfo = SYCLKernelBinaryInfo; +using FrozenSymbolTable = DynArray; + +// Note: `FrozenPropertyValue` and `FrozenPropertySet` constructors take +// `std::string_view` arguments instead of `const char *` because they will be +// created from `llvm::SmallString`s, which don't contain the trailing '\0' +// byte. Hence obtaining a C-string would cause an additional copy. + +struct FrozenPropertyValue { + sycl::detail::string Name; + bool IsUIntValue; + uint32_t UIntValue; + DynArray Bytes; + + FrozenPropertyValue() = default; + FrozenPropertyValue(FrozenPropertyValue &&) = default; + FrozenPropertyValue &operator=(FrozenPropertyValue &&) = default; + + FrozenPropertyValue(std::string_view Name, uint32_t Value) + : Name{Name}, IsUIntValue{true}, UIntValue{Value}, Bytes{0} {} + FrozenPropertyValue(std::string_view Name, const uint8_t *Ptr, size_t Size) + : Name{Name}, IsUIntValue{false}, Bytes{Size} { + std::memcpy(Bytes.begin(), Ptr, Size); + } +}; + +struct FrozenPropertySet { + sycl::detail::string Name; + DynArray Values; + + FrozenPropertySet() = default; + FrozenPropertySet(FrozenPropertySet &&) = default; + FrozenPropertySet &operator=(FrozenPropertySet &&) = default; + + FrozenPropertySet(std::string_view Name, size_t Size) + : Name{Name}, Values{Size} {} +}; + +using FrozenPropertyRegistry = DynArray; + +struct RTCBundleInfo { + RTCBundleBinaryInfo BinaryInfo; + FrozenSymbolTable SymbolTable; + FrozenPropertyRegistry Properties; +}; + } // namespace jit_compiler #endif // SYCL_FUSION_COMMON_KERNEL_H diff --git a/sycl-jit/jit-compiler/CMakeLists.txt b/sycl-jit/jit-compiler/CMakeLists.txt index 98385e8ee4f2a..a704b4c91a503 100644 --- a/sycl-jit/jit-compiler/CMakeLists.txt +++ b/sycl-jit/jit-compiler/CMakeLists.txt @@ -8,6 +8,7 @@ add_llvm_library(sycl-jit lib/fusion/JITContext.cpp lib/fusion/ModuleHelper.cpp lib/rtc/DeviceCompilation.cpp + lib/rtc/PostLinkActions.cpp lib/helper/ConfigHelper.cpp SHARED @@ -31,6 +32,7 @@ add_llvm_library(sycl-jit Target TargetParser MC + SYCLLowerIR ${LLVM_TARGETS_TO_BUILD} LINK_LIBS diff --git a/sycl-jit/jit-compiler/include/KernelFusion.h b/sycl-jit/jit-compiler/include/KernelFusion.h index d3575f33189aa..76636beb937d9 100644 --- a/sycl-jit/jit-compiler/include/KernelFusion.h +++ b/sycl-jit/jit-compiler/include/KernelFusion.h @@ -56,6 +56,32 @@ class JITResult { sycl::detail::string ErrorMessage; }; +class RTCResult { +public: + explicit RTCResult(const char *ErrorMessage) + : Failed{true}, BundleInfo{}, ErrorMessage{ErrorMessage} {} + + explicit RTCResult(RTCBundleInfo &&BundleInfo) + : Failed{false}, BundleInfo{std::move(BundleInfo)}, ErrorMessage{} {} + + bool failed() const { return Failed; } + + const char *getErrorMessage() const { + assert(failed() && "No error message present"); + return ErrorMessage.c_str(); + } + + const RTCBundleInfo &getBundleInfo() const { + assert(!failed() && "No bundle info"); + return BundleInfo; + } + +private: + bool Failed; + RTCBundleInfo BundleInfo; + sycl::detail::string ErrorMessage; +}; + extern "C" { #ifdef __clang__ @@ -77,7 +103,7 @@ KF_EXPORT_SYMBOL JITResult materializeSpecConstants( const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo, View SpecConstBlob); -KF_EXPORT_SYMBOL JITResult compileSYCL(InMemoryFile SourceFile, +KF_EXPORT_SYMBOL RTCResult compileSYCL(InMemoryFile SourceFile, View IncludeFiles, View UserArgs); diff --git a/sycl-jit/jit-compiler/lib/KernelFusion.cpp b/sycl-jit/jit-compiler/lib/KernelFusion.cpp index 9f8bb09ec9f26..d5bae664dc886 100644 --- a/sycl-jit/jit-compiler/lib/KernelFusion.cpp +++ b/sycl-jit/jit-compiler/lib/KernelFusion.cpp @@ -25,8 +25,8 @@ using namespace jit_compiler; using FusedFunction = helper::FusionHelper::FusedFunction; using FusedFunctionList = std::vector; -static JITResult errorToFusionResult(llvm::Error &&Err, - const std::string &Msg) { +template +static ResultType wrapError(llvm::Error &&Err, const std::string &Msg) { std::stringstream ErrMsg; ErrMsg << Msg << "\nDetailed information:\n"; llvm::handleAllErrors(std::move(Err), @@ -35,7 +35,16 @@ static JITResult errorToFusionResult(llvm::Error &&Err, // compiled without exception support. ErrMsg << "\t" << StrErr.getMessage() << "\n"; }); - return JITResult{ErrMsg.str().c_str()}; + return ResultType{ErrMsg.str().c_str()}; +} + +static JITResult errorToFusionResult(llvm::Error &&Err, + const std::string &Msg) { + return wrapError(std::move(Err), Msg); +} + +static RTCResult errorToRTCResult(llvm::Error &&Err, const std::string &Msg) { + return wrapError(std::move(Err), Msg); } static std::vector @@ -234,20 +243,20 @@ fuseKernels(View KernelInformation, const char *FusedKernelName, return JITResult{FusedKernelInfo}; } -extern "C" KF_EXPORT_SYMBOL JITResult +extern "C" KF_EXPORT_SYMBOL RTCResult compileSYCL(InMemoryFile SourceFile, View IncludeFiles, View UserArgs) { auto UserArgListOrErr = parseUserArgs(UserArgs); if (!UserArgListOrErr) { - return errorToFusionResult(UserArgListOrErr.takeError(), - "Parsing of user arguments failed"); + return errorToRTCResult(UserArgListOrErr.takeError(), + "Parsing of user arguments failed"); } llvm::opt::InputArgList UserArgList = std::move(*UserArgListOrErr); auto ModuleOrErr = compileDeviceCode(SourceFile, IncludeFiles, UserArgList); if (!ModuleOrErr) { - return errorToFusionResult(ModuleOrErr.takeError(), - "Device compilation failed"); + return errorToRTCResult(ModuleOrErr.takeError(), + "Device compilation failed"); } std::unique_ptr Context; @@ -255,16 +264,26 @@ compileSYCL(InMemoryFile SourceFile, View IncludeFiles, Context.reset(&Module->getContext()); if (auto Error = linkDeviceLibraries(*Module, UserArgList)) { - return errorToFusionResult(std::move(Error), "Device linking failed"); + return errorToRTCResult(std::move(Error), "Device linking failed"); } - SYCLKernelInfo Kernel; - if (auto Error = translation::KernelTranslator::translateKernel( - Kernel, *Module, JITContext::getInstance(), BinaryFormat::SPIRV)) { - return errorToFusionResult(std::move(Error), "SPIR-V translation failed"); + auto BundleInfoOrError = performPostLink(*Module, UserArgList); + if (!BundleInfoOrError) { + return errorToRTCResult(BundleInfoOrError.takeError(), + "Post-link phase failed"); + } + auto BundleInfo = std::move(*BundleInfoOrError); + + auto BinaryInfoOrError = + translation::KernelTranslator::translateBundleToSPIRV( + *Module, JITContext::getInstance()); + if (!BinaryInfoOrError) { + return errorToRTCResult(BinaryInfoOrError.takeError(), + "SPIR-V translation failed"); } + BundleInfo.BinaryInfo = std::move(*BinaryInfoOrError); - return JITResult{Kernel}; + return RTCResult{std::move(BundleInfo)}; } 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 1bdfe7d63b641..5931458990d8c 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -8,6 +8,8 @@ #include "DeviceCompilation.h" +#include "PostLinkActions.h" + #include #include #include @@ -20,8 +22,10 @@ #include #include - -#include +#include +#include +#include +#include using namespace clang; using namespace clang::tooling; @@ -29,6 +33,11 @@ using namespace clang::driver; using namespace clang::driver::options; using namespace llvm; using namespace llvm::opt; +using namespace llvm::sycl; +using namespace llvm::module_split; +using namespace llvm::util; +using namespace jit_compiler; +using namespace jit_compiler::post_link; #ifdef _GNU_SOURCE #include @@ -356,6 +365,96 @@ Error jit_compiler::linkDeviceLibraries(llvm::Module &Module, return Error::success(); } +Expected jit_compiler::performPostLink( + llvm::Module &Module, [[maybe_unused]] const InputArgList &UserArgList) { + // This is a simplified version of `processInputModule` in + // `llvm/tools/sycl-post-link.cpp`. Assertions/TODOs point to functionality + // left out of the algorithm for now. + + // After linking device bitcode "llvm.used" holds references to the kernels + // that are defined in the device image. But after splitting device image into + // separate kernels we may end up with having references to kernel declaration + // originating from "llvm.used" in the IR that is passed to llvm-spirv tool, + // and these declarations cause an assertion in llvm-spirv. To workaround this + // issue remove "llvm.used" from the input module before performing any other + // actions. + removeSYCLKernelsConstRefArray(Module); + + // There may be device_global variables kept alive in "llvm.compiler.used" + // to keep the optimizer from wrongfully removing them. llvm.compiler.used + // symbols are usually removed at backend lowering, but this is handled here + // for SPIR-V since SYCL compilation uses llvm-spirv, not the SPIR-V backend. + removeDeviceGlobalFromCompilerUsed(Module); + + assert(!isModuleUsingAsan(Module)); + // Otherwise: Need to instrument each image scope device globals if the module + // has been instrumented by sanitizer pass. + + // Transform Joint Matrix builtin calls to align them with SPIR-V friendly + // LLVM IR specification. + runModulePass(Module); + + // TODO: Implement actual device code splitting. We're just using the splitter + // to obtain additional information about the module for now. + // TODO: EmitOnlyKernelsAsEntryPoints is controlled by + // `shouldEmitOnlyKernelsAsEntryPoints` in + // `clang/lib/Driver/ToolChains/Clang.cpp`. + std::unique_ptr Splitter = getDeviceCodeSplitter( + ModuleDesc{std::unique_ptr{&Module}}, SPLIT_NONE, + /*IROutputOnly=*/false, + /*EmitOnlyKernelsAsEntryPoints=*/true); + bool SplitOccurred = Splitter->remainingSplits() > 1; + assert(!SplitOccurred); + + // TODO: Call `verifyNoCrossModuleDeviceGlobalUsage` if device globals shall + // be processed. + + assert(Splitter->hasMoreSplits()); + ModuleDesc MDesc = Splitter->nextSplit(); + assert(&Module == &MDesc.getModule()); + MDesc.saveSplitInformationAsMetadata(); + + RTCBundleInfo BundleInfo; + BundleInfo.SymbolTable = + decltype(BundleInfo.SymbolTable){MDesc.entries().size()}; + transform(MDesc.entries(), BundleInfo.SymbolTable.begin(), + [](Function *F) { return F->getName(); }); + + // TODO: Determine what is requested. + GlobalBinImageProps PropReq{ + /*EmitKernelParamInfo=*/true, /*EmitProgramMetadata=*/true, + /*EmitExportedSymbols=*/true, /*EmitImportedSymbols=*/true, + /*DeviceGlobals=*/false}; + PropertySetRegistry Properties = + computeModuleProperties(MDesc.getModule(), MDesc.entries(), PropReq); + // TODO: Manually add `compile_target` property as in + // `saveModuleProperties`? + const auto &PropertySets = Properties.getPropSets(); + + BundleInfo.Properties = decltype(BundleInfo.Properties){PropertySets.size()}; + for (auto &&[KV, FrozenPropSet] : zip(PropertySets, BundleInfo.Properties)) { + const auto &PropertySetName = KV.first; + const auto &PropertySet = KV.second; + FrozenPropertySet FPS{PropertySetName.str(), PropertySet.size()}; + for (auto &&[KV2, FrozenProp] : zip(PropertySet, FPS.Values)) { + const auto &PropertyName = KV2.first; + const auto &PropertyValue = KV2.second; + FrozenProp = PropertyValue.getType() == PropertyValue::Type::UINT32 + ? FrozenPropertyValue{PropertyName.str(), + PropertyValue.asUint32()} + : FrozenPropertyValue{ + PropertyName.str(), PropertyValue.asRawByteArray(), + PropertyValue.getRawByteArraySize()}; + } + FrozenPropSet = std::move(FPS); + }; + + // Regain ownership of the module. + MDesc.releaseModulePtr().release(); + + return BundleInfo; +} + Expected jit_compiler::parseUserArgs(View UserArgs) { unsigned MissingArgIndex, MissingArgCount; @@ -410,5 +509,17 @@ jit_compiler::parseUserArgs(View UserArgs) { } } + if (auto DCSMode = AL.getLastArgValue(OPT_fsycl_device_code_split_EQ, "none"); + DCSMode != "none" && DCSMode != "auto") { + return createStringError("Device code splitting is not yet supported"); + } + + if (AL.hasArg(OPT_fsycl_device_code_split_esimd, + OPT_fno_sycl_device_code_split_esimd)) { + // TODO: There are more ESIMD-related options. + return createStringError( + "Runtime compilation of ESIMD kernels is not yet supported"); + } + return Expected{std::move(AL)}; } diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h index ec890a8213827..8aa47939e3b1d 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h @@ -27,6 +27,10 @@ compileDeviceCode(InMemoryFile SourceFile, View IncludeFiles, llvm::Error linkDeviceLibraries(llvm::Module &Module, const llvm::opt::InputArgList &UserArgList); +llvm::Expected +performPostLink(llvm::Module &Module, + const llvm::opt::InputArgList &UserArgList); + llvm::Expected parseUserArgs(View UserArgs); diff --git a/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp b/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp new file mode 100644 index 0000000000000..82fac2f824445 --- /dev/null +++ b/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp @@ -0,0 +1,116 @@ +//==------------------------ PostLinkActions.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 "PostLinkActions.h" + +#include +#include +#include + +using namespace llvm; + +bool jit_compiler::post_link::removeSYCLKernelsConstRefArray(Module &M) { + GlobalVariable *GV = M.getGlobalVariable("llvm.used"); + + if (!GV) { + return false; + } + assert(GV->user_empty() && "Unexpected llvm.used users"); + Constant *Initializer = GV->getInitializer(); + GV->setInitializer(nullptr); + GV->eraseFromParent(); + + // Destroy the initializer and all operands of it. + SmallVector IOperands; + for (auto It = Initializer->op_begin(); It != Initializer->op_end(); It++) + IOperands.push_back(cast(*It)); + assert(llvm::isSafeToDestroyConstant(Initializer) && + "Cannot remove initializer of llvm.used global"); + Initializer->destroyConstant(); + for (auto It = IOperands.begin(); It != IOperands.end(); It++) { + auto Op = (*It)->stripPointerCasts(); + auto *F = dyn_cast(Op); + if (llvm::isSafeToDestroyConstant(*It)) { + (*It)->destroyConstant(); + } else if (F && F->getCallingConv() == CallingConv::SPIR_KERNEL && + !F->use_empty()) { + // The element in "llvm.used" array has other users. That is Ok for + // specialization constants, but is wrong for kernels. + llvm::report_fatal_error("Unexpected usage of SYCL kernel"); + } + + // Remove unused kernel declarations to avoid LLVM IR check fails. + if (F && F->isDeclaration() && F->use_empty()) + F->eraseFromParent(); + } + return true; +} + +// Removes all device_global variables from the llvm.compiler.used global +// variable. A device_global with internal linkage will be in llvm.compiler.used +// to avoid the compiler wrongfully removing it during optimizations. However, +// as an effect the device_global variables will also be distributed across +// binaries, even if llvm.compiler.used has served its purpose. To avoid +// polluting other binaries with unused device_global variables, we remove them +// from llvm.compiler.used and erase them if they have no further uses. +bool jit_compiler::post_link::removeDeviceGlobalFromCompilerUsed(Module &M) { + GlobalVariable *GV = M.getGlobalVariable("llvm.compiler.used"); + if (!GV) + return false; + + // Erase the old llvm.compiler.used. A new one will be created at the end if + // there are other values in it (other than device_global). + assert(GV->user_empty() && "Unexpected llvm.compiler.used users"); + Constant *Initializer = GV->getInitializer(); + const auto *VAT = cast(GV->getValueType()); + GV->setInitializer(nullptr); + GV->eraseFromParent(); + + // Destroy the initializer. Keep the operands so we keep the ones we need. + SmallVector IOperands; + for (auto It = Initializer->op_begin(); It != Initializer->op_end(); It++) + IOperands.push_back(cast(*It)); + assert(llvm::isSafeToDestroyConstant(Initializer) && + "Cannot remove initializer of llvm.compiler.used global"); + Initializer->destroyConstant(); + + // Iterate through all operands. If they are device_global then we drop them + // and erase them if they have no uses afterwards. All other values are kept. + SmallVector NewOperands; + for (auto It = IOperands.begin(); It != IOperands.end(); It++) { + Constant *Op = *It; + auto *DG = dyn_cast(Op->stripPointerCasts()); + + // If it is not a device_global we keep it. + if (!DG || !isDeviceGlobalVariable(*DG)) { + NewOperands.push_back(Op); + continue; + } + + // Destroy the device_global operand. + if (llvm::isSafeToDestroyConstant(Op)) + Op->destroyConstant(); + + // Remove device_global if it no longer has any uses. + if (!DG->isConstantUsed()) + DG->eraseFromParent(); + } + + // If we have any operands left from the original llvm.compiler.used we create + // a new one with the new size. + if (!NewOperands.empty()) { + ArrayType *ATy = ArrayType::get(VAT->getElementType(), NewOperands.size()); + GlobalVariable *NGV = + new GlobalVariable(M, ATy, false, GlobalValue::AppendingLinkage, + ConstantArray::get(ATy, NewOperands), ""); + NGV->setName("llvm.compiler.used"); + NGV->setSection("llvm.metadata"); + } + + return true; +} diff --git a/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.h b/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.h new file mode 100644 index 0000000000000..878e4244b2996 --- /dev/null +++ b/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.h @@ -0,0 +1,50 @@ +//==------ PostLinkActions.h - Fork of sycl-post-link actions for RTC ------==// +// +// 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_POST_LINK_ACTIONS_H +#define SYCL_JIT_COMPILER_RTC_POST_LINK_ACTIONS_H + +#include +#include +#include + +namespace jit_compiler::post_link { + +using namespace llvm; + +template bool runModulePass(Module &M) { + ModulePassManager MPM; + ModuleAnalysisManager MAM; + // Register required analysis + MAM.registerPass([&] { return PassInstrumentationAnalysis(); }); + MPM.addPass(PassClass{}); + PreservedAnalyses Res = MPM.run(M, MAM); + return !Res.areAllPreserved(); +} + +// Removes the global variable "llvm.used" and returns true on success. +// "llvm.used" is a global constant array containing references to kernels +// available in the module and callable from host code. The elements of +// the array are ConstantExpr bitcast to i8*. +// The variable must be removed as it is a) has done the job to the moment +// of this function call and b) the references to the kernels callable from +// host must not have users. +bool removeSYCLKernelsConstRefArray(Module &M); + +// Removes all device_global variables from the llvm.compiler.used global +// variable. A device_global with internal linkage will be in llvm.compiler.used +// to avoid the compiler wrongfully removing it during optimizations. However, +// as an effect the device_global variables will also be distributed across +// binaries, even if llvm.compiler.used has served its purpose. To avoid +// polluting other binaries with unused device_global variables, we remove them +// from llvm.compiler.used and erase them if they have no further uses. +bool removeDeviceGlobalFromCompilerUsed(llvm::Module &M); + +} // namespace jit_compiler::post_link + +#endif // SYCL_JIT_COMPILER_RTC_POST_LINK_ACTIONS_H diff --git a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp index 40eec5d241ea0..61ca038dea4cf 100644 --- a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp +++ b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp @@ -222,6 +222,20 @@ llvm::Error KernelTranslator::translateKernel(SYCLKernelInfo &Kernel, return Error::success(); } +llvm::Expected +KernelTranslator::translateBundleToSPIRV(llvm::Module &Mod, + JITContext &JITCtx) { + llvm::Expected BinaryOrError = translateToSPIRV(Mod, JITCtx); + if (auto Error = BinaryOrError.takeError()) { + return Error; + } + KernelBinary *Binary = *BinaryOrError; + RTCBundleBinaryInfo BBI{BinaryFormat::SPIRV, + Mod.getDataLayout().getPointerSizeInBits(), + Binary->address(), Binary->size()}; + return BBI; +} + llvm::Expected KernelTranslator::translateToSPIRV(llvm::Module &Mod, JITContext &JITCtx) { return SPIRVLLVMTranslator::translateLLVMtoSPIRV(Mod, JITCtx); diff --git a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.h b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.h index 809c8fab2e42f..e71aa2b0d19f3 100644 --- a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.h +++ b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.h @@ -27,6 +27,9 @@ class KernelTranslator { static llvm::Error translateKernel(SYCLKernelInfo &Kernel, llvm::Module &Mod, JITContext &JITCtx, BinaryFormat Format); + static llvm::Expected + translateBundleToSPIRV(llvm::Module &Mod, JITContext &JITCtx); + private: /// /// Pair of address and size to represent a binary blob. diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 19f1915943f05..5ce1376eda66e 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1207,7 +1207,7 @@ std::vector jit_compiler::compileSYCL( // TODO: We currently don't have a meaningful build log. (void)LogPtr; - const auto &BI = Result.getKernelInfo().BinaryInfo; + const auto &BI = Result.getBundleInfo().BinaryInfo; assert(BI.Format == ::jit_compiler::BinaryFormat::SPIRV); std::vector SPV(BI.BinaryStart, BI.BinaryStart + BI.BinarySize); return SPV; From b1eeac360c76817b6dbc7ee13d7ea88765184da4 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 18 Nov 2024 14:50:04 +0000 Subject: [PATCH 02/12] WIP, UR program is created via program manager Signed-off-by: Julian Oppermann --- sycl/source/detail/jit_compiler.cpp | 50 ++++++++- sycl/source/detail/jit_compiler.hpp | 8 +- sycl/source/detail/jit_device_binaries.cpp | 2 +- sycl/source/detail/jit_device_binaries.hpp | 2 +- sycl/source/detail/kernel_bundle_impl.hpp | 104 +++++++++--------- .../kernel_compiler/kernel_compiler_sycl.cpp | 2 +- .../kernel_compiler/kernel_compiler_sycl.hpp | 4 +- 7 files changed, 111 insertions(+), 61 deletions(-) diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 5ce1376eda66e..2bc8e4b7cf5dd 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1119,6 +1119,49 @@ sycl_device_binaries jit_compiler::createPIDeviceBinary( return JITDeviceBinaries.back().getPIDeviceStruct(); } +const RTDeviceBinaryImage &jit_compiler::createDeviceBinaryImage( + const ::jit_compiler::RTCBundleInfo &BundleInfo) { + DeviceBinaryContainer Binary; + for (const auto &Symbol : BundleInfo.SymbolTable) { + // Create an offload entry each kernel. + // It seems to be OK to set zero for most of the information here, at least + // that is the case for compiled SPIR-V binaries. + OffloadEntryContainer Entry{Symbol.c_str(), nullptr, 0, 0, 0}; + Binary.addOffloadEntry(std::move(Entry)); + } + + for (const auto &FPS : BundleInfo.Properties) { + PropertySetContainer PropSet{FPS.Name.c_str()}; + for (const auto &FPV : FPS.Values) { + if (FPV.IsUIntValue) { + PropSet.addProperty(PropertyContainer{FPV.Name.c_str(), FPV.UIntValue}); + } else { + PropSet.addProperty(PropertyContainer{ + FPV.Name.c_str(), FPV.Bytes.begin(), FPV.Bytes.size(), + sycl_property_type::SYCL_PROPERTY_TYPE_BYTE_ARRAY}); + } + } + Binary.addProperty(std::move(PropSet)); + } + + DeviceBinariesCollection Collection; + Collection.addDeviceBinary(std::move(Binary), + BundleInfo.BinaryInfo.BinaryStart, + BundleInfo.BinaryInfo.BinarySize, + (BundleInfo.BinaryInfo.AddressBits == 64) + ? __SYCL_DEVICE_BINARY_TARGET_SPIRV64 + : __SYCL_DEVICE_BINARY_TARGET_SPIRV32, + SYCL_DEVICE_BINARY_TYPE_SPIRV); + JITDeviceBinaries.push_back(std::move(Collection)); + // TODO: If we want to handle multiple device binary images, we should instead + // return `sycl_device_binaries`, to be passed to + // `program_manager::addImages`. The program manager then creates and + // owns the `RTDeviceBinaryImage` instances. + RTCDeviceBinaryImages.emplace_back( + &JITDeviceBinaries.back().getPIDeviceStruct()->DeviceBinaries[0]); + return RTCDeviceBinaryImages.back(); +} + std::vector jit_compiler::encodeArgUsageMask( const ::jit_compiler::ArgUsageMask &Mask) const { // This must match the decoding logic in program_manager.cpp. @@ -1167,7 +1210,7 @@ std::vector jit_compiler::encodeReqdWorkGroupSize( return Encoded; } -std::vector jit_compiler::compileSYCL( +const RTDeviceBinaryImage &jit_compiler::compileSYCL( const std::string &Id, const std::string &SYCLSource, const std::vector> &IncludePairs, const std::vector &UserArgs, std::string *LogPtr, @@ -1207,10 +1250,7 @@ std::vector jit_compiler::compileSYCL( // TODO: We currently don't have a meaningful build log. (void)LogPtr; - const auto &BI = Result.getBundleInfo().BinaryInfo; - assert(BI.Format == ::jit_compiler::BinaryFormat::SPIRV); - std::vector SPV(BI.BinaryStart, BI.BinaryStart + BI.BinarySize); - return SPV; + return createDeviceBinaryImage(Result.getBundleInfo()); } } // namespace detail diff --git a/sycl/source/detail/jit_compiler.hpp b/sycl/source/detail/jit_compiler.hpp index 1908defa42e77..d84c6bc7b4f08 100644 --- a/sycl/source/detail/jit_compiler.hpp +++ b/sycl/source/detail/jit_compiler.hpp @@ -44,7 +44,7 @@ class jit_compiler { const std::string &KernelName, const std::vector &SpecConstBlob); - std::vector compileSYCL( + const RTDeviceBinaryImage &compileSYCL( const std::string &Id, const std::string &SYCLSource, const std::vector> &IncludePairs, const std::vector &UserArgs, std::string *LogPtr, @@ -69,6 +69,9 @@ class jit_compiler { createPIDeviceBinary(const ::jit_compiler::SYCLKernelInfo &FusedKernelInfo, ::jit_compiler::BinaryFormat Format); + const RTDeviceBinaryImage & + createDeviceBinaryImage(const ::jit_compiler::RTCBundleInfo &BundleInfo); + std::vector encodeArgUsageMask(const ::jit_compiler::ArgUsageMask &Mask) const; @@ -81,6 +84,9 @@ class jit_compiler { // Manages the lifetime of the UR structs for device binaries. std::vector JITDeviceBinaries; + // Manages the lifetime of the runtime wrappers for device binary images. + std::vector RTCDeviceBinaryImages; + #if SYCL_EXT_JIT_ENABLE // Handles to the entry points of the lazily loaded JIT library. using FuseKernelsFuncT = decltype(::jit_compiler::fuseKernels) *; diff --git a/sycl/source/detail/jit_device_binaries.cpp b/sycl/source/detail/jit_device_binaries.cpp index b5f3b67c08ecd..f90be2c27ec3a 100644 --- a/sycl/source/detail/jit_device_binaries.cpp +++ b/sycl/source/detail/jit_device_binaries.cpp @@ -27,7 +27,7 @@ _sycl_offload_entry_struct OffloadEntryContainer::getPIOffloadEntry() { EntryFlags, EntryReserved}; } -PropertyContainer::PropertyContainer(const std::string &Name, void *Data, +PropertyContainer::PropertyContainer(const std::string &Name, const void *Data, size_t Size, uint32_t Type) : PropName{new char[Name.length() + 1]}, Value{new unsigned char[Size]}, ValueSize{Size}, PropType{Type} { diff --git a/sycl/source/detail/jit_device_binaries.hpp b/sycl/source/detail/jit_device_binaries.hpp index 49445e07920e7..bca83839f39e8 100644 --- a/sycl/source/detail/jit_device_binaries.hpp +++ b/sycl/source/detail/jit_device_binaries.hpp @@ -50,7 +50,7 @@ class OffloadEntryContainer { class PropertyContainer { public: - PropertyContainer(const std::string &Name, void *Data, size_t Size, + PropertyContainer(const std::string &Name, const void *Data, size_t Size, uint32_t Type); // Set a UR_PROPERTY_TYPE_UINT32 property PropertyContainer(const std::string &Name, uint32_t Data); diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 7ce5971711291..ecea8e079b32f 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -468,57 +468,59 @@ class kernel_bundle_impl { } if (!FetchedFromCache) { - const auto spirv = [&]() -> std::vector { - if (Language == syclex::source_language::opencl) { - // if successful, the log is empty. if failed, throws an error with - // the compilation log. - std::vector IPVersionVec(Devices.size()); - std::transform(DeviceVec.begin(), DeviceVec.end(), - IPVersionVec.begin(), [&](ur_device_handle_t d) { - uint32_t ipVersion = 0; - Adapter->call( - d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t), - &ipVersion, nullptr); - return ipVersion; - }); - return syclex::detail::OpenCLC_to_SPIRV(*SourceStrPtr, IPVersionVec, - BuildOptions, LogPtr); - } - if (Language == syclex::source_language::spirv) { - const auto &SourceBytes = - std::get>(this->Source); - std::vector Result(SourceBytes.size()); - std::transform(SourceBytes.cbegin(), SourceBytes.cend(), - Result.begin(), - [](std::byte B) { return static_cast(B); }); - return Result; - } - if (Language == syclex::source_language::sycl) { - return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs, - BuildOptions, LogPtr, - RegisteredKernelNames); - } - if (Language == syclex::source_language::sycl_jit) { - const auto &SourceStr = std::get(this->Source); - return syclex::detail::SYCL_JIT_to_SPIRV(SourceStr, IncludePairs, - BuildOptions, LogPtr, - RegisteredKernelNames); - } - throw sycl::exception( - make_error_code(errc::invalid), - "SYCL C++, OpenCL C and SPIR-V are the only supported " - "languages at this time"); - }(); - - Adapter->call( - ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr, - &UrProgram); - // program created by urProgramCreateWithIL is implicitly retained. - if (UrProgram == nullptr) - throw sycl::exception( - sycl::make_error_code(errc::invalid), - "urProgramCreateWithIL resulted in a null program handle."); - + if (Language == syclex::source_language::sycl_jit) { + const auto &SourceStr = std::get(this->Source); + const auto &Img = syclex::detail::SYCL_JIT_to_SPIRV( + SourceStr, IncludePairs, BuildOptions, LogPtr, + RegisteredKernelNames); + UrProgram = ProgramManager::getInstance().createURProgram(Img, MContext, + MDevices); + } else { + const auto spirv = [&]() -> std::vector { + if (Language == syclex::source_language::opencl) { + // if successful, the log is empty. if failed, throws an error with + // the compilation log. + std::vector IPVersionVec(Devices.size()); + std::transform(DeviceVec.begin(), DeviceVec.end(), + IPVersionVec.begin(), [&](ur_device_handle_t d) { + uint32_t ipVersion = 0; + Adapter->call( + d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t), + &ipVersion, nullptr); + return ipVersion; + }); + return syclex::detail::OpenCLC_to_SPIRV(*SourceStrPtr, IPVersionVec, + BuildOptions, LogPtr); + } + if (Language == syclex::source_language::spirv) { + const auto &SourceBytes = + std::get>(this->Source); + std::vector Result(SourceBytes.size()); + std::transform(SourceBytes.cbegin(), SourceBytes.cend(), + Result.begin(), + [](std::byte B) { return static_cast(B); }); + return Result; + } + if (Language == syclex::source_language::sycl) { + return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs, + BuildOptions, LogPtr, + RegisteredKernelNames); + } + throw sycl::exception( + make_error_code(errc::invalid), + "SYCL C++, OpenCL C and SPIR-V are the only supported " + "languages at this time"); + }(); + + Adapter->call( + ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr, + &UrProgram); + // program created by urProgramCreateWithIL is implicitly retained. + if (UrProgram == nullptr) + throw sycl::exception( + sycl::make_error_code(errc::invalid), + "urProgramCreateWithIL resulted in a null program handle."); + } } // if(!FetchedFromCache) std::string XsFlags = extractXsFlags(BuildOptions); diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 89a8a548a6d8a..10e1755370e00 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -342,7 +342,7 @@ bool SYCL_JIT_Compilation_Available() { #endif } -spirv_vec_t SYCL_JIT_to_SPIRV( +const sycl::detail::RTDeviceBinaryImage &SYCL_JIT_to_SPIRV( [[maybe_unused]] const std::string &SYCLSource, [[maybe_unused]] include_pairs_t IncludePairs, [[maybe_unused]] const std::vector &UserArgs, diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index 948f199c1b3cc..8fce2a2045c32 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -12,6 +12,8 @@ #include // __SYCL_EXPORT #include +#include + #include // std::accumulate #include #include @@ -33,7 +35,7 @@ bool SYCL_Compilation_Available(); std::string userArgsAsString(const std::vector &UserArguments); -spirv_vec_t +const sycl::detail::RTDeviceBinaryImage & SYCL_JIT_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs, const std::vector &UserArgs, std::string *LogPtr, const std::vector &RegisteredKernelNames); From db562cbc8a247a11f2c535cb80bf04d0561ba6a3 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Wed, 20 Nov 2024 13:20:46 +0000 Subject: [PATCH 03/12] Return `sycl_device_binaries` from JIT library Signed-off-by: Julian Oppermann --- sycl/source/detail/jit_compiler.cpp | 12 +- sycl/source/detail/jit_compiler.hpp | 7 +- sycl/source/detail/kernel_bundle_impl.hpp | 107 +++++++++--------- .../kernel_compiler/kernel_compiler_sycl.cpp | 2 +- .../kernel_compiler/kernel_compiler_sycl.hpp | 2 +- 5 files changed, 61 insertions(+), 69 deletions(-) diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 2bc8e4b7cf5dd..742769f76af81 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1119,7 +1119,7 @@ sycl_device_binaries jit_compiler::createPIDeviceBinary( return JITDeviceBinaries.back().getPIDeviceStruct(); } -const RTDeviceBinaryImage &jit_compiler::createDeviceBinaryImage( +sycl_device_binaries jit_compiler::createDeviceBinaryImage( const ::jit_compiler::RTCBundleInfo &BundleInfo) { DeviceBinaryContainer Binary; for (const auto &Symbol : BundleInfo.SymbolTable) { @@ -1153,13 +1153,7 @@ const RTDeviceBinaryImage &jit_compiler::createDeviceBinaryImage( : __SYCL_DEVICE_BINARY_TARGET_SPIRV32, SYCL_DEVICE_BINARY_TYPE_SPIRV); JITDeviceBinaries.push_back(std::move(Collection)); - // TODO: If we want to handle multiple device binary images, we should instead - // return `sycl_device_binaries`, to be passed to - // `program_manager::addImages`. The program manager then creates and - // owns the `RTDeviceBinaryImage` instances. - RTCDeviceBinaryImages.emplace_back( - &JITDeviceBinaries.back().getPIDeviceStruct()->DeviceBinaries[0]); - return RTCDeviceBinaryImages.back(); + return JITDeviceBinaries.back().getPIDeviceStruct(); } std::vector jit_compiler::encodeArgUsageMask( @@ -1210,7 +1204,7 @@ std::vector jit_compiler::encodeReqdWorkGroupSize( return Encoded; } -const RTDeviceBinaryImage &jit_compiler::compileSYCL( +sycl_device_binaries jit_compiler::compileSYCL( const std::string &Id, const std::string &SYCLSource, const std::vector> &IncludePairs, const std::vector &UserArgs, std::string *LogPtr, diff --git a/sycl/source/detail/jit_compiler.hpp b/sycl/source/detail/jit_compiler.hpp index d84c6bc7b4f08..982246b48da5f 100644 --- a/sycl/source/detail/jit_compiler.hpp +++ b/sycl/source/detail/jit_compiler.hpp @@ -44,7 +44,7 @@ class jit_compiler { const std::string &KernelName, const std::vector &SpecConstBlob); - const RTDeviceBinaryImage &compileSYCL( + sycl_device_binaries compileSYCL( const std::string &Id, const std::string &SYCLSource, const std::vector> &IncludePairs, const std::vector &UserArgs, std::string *LogPtr, @@ -69,7 +69,7 @@ class jit_compiler { createPIDeviceBinary(const ::jit_compiler::SYCLKernelInfo &FusedKernelInfo, ::jit_compiler::BinaryFormat Format); - const RTDeviceBinaryImage & + sycl_device_binaries createDeviceBinaryImage(const ::jit_compiler::RTCBundleInfo &BundleInfo); std::vector @@ -84,9 +84,6 @@ class jit_compiler { // Manages the lifetime of the UR structs for device binaries. std::vector JITDeviceBinaries; - // Manages the lifetime of the runtime wrappers for device binary images. - std::vector RTCDeviceBinaryImages; - #if SYCL_EXT_JIT_ENABLE // Handles to the entry points of the lazily loaded JIT library. using FuseKernelsFuncT = decltype(::jit_compiler::fuseKernels) *; diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index ecea8e079b32f..58e605e85c458 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -468,59 +468,60 @@ class kernel_bundle_impl { } if (!FetchedFromCache) { - if (Language == syclex::source_language::sycl_jit) { - const auto &SourceStr = std::get(this->Source); - const auto &Img = syclex::detail::SYCL_JIT_to_SPIRV( - SourceStr, IncludePairs, BuildOptions, LogPtr, - RegisteredKernelNames); - UrProgram = ProgramManager::getInstance().createURProgram(Img, MContext, - MDevices); - } else { - const auto spirv = [&]() -> std::vector { - if (Language == syclex::source_language::opencl) { - // if successful, the log is empty. if failed, throws an error with - // the compilation log. - std::vector IPVersionVec(Devices.size()); - std::transform(DeviceVec.begin(), DeviceVec.end(), - IPVersionVec.begin(), [&](ur_device_handle_t d) { - uint32_t ipVersion = 0; - Adapter->call( - d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t), - &ipVersion, nullptr); - return ipVersion; - }); - return syclex::detail::OpenCLC_to_SPIRV(*SourceStrPtr, IPVersionVec, - BuildOptions, LogPtr); - } - if (Language == syclex::source_language::spirv) { - const auto &SourceBytes = - std::get>(this->Source); - std::vector Result(SourceBytes.size()); - std::transform(SourceBytes.cbegin(), SourceBytes.cend(), - Result.begin(), - [](std::byte B) { return static_cast(B); }); - return Result; - } - if (Language == syclex::source_language::sycl) { - return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs, - BuildOptions, LogPtr, - RegisteredKernelNames); - } - throw sycl::exception( - make_error_code(errc::invalid), - "SYCL C++, OpenCL C and SPIR-V are the only supported " - "languages at this time"); - }(); - - Adapter->call( - ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr, - &UrProgram); - // program created by urProgramCreateWithIL is implicitly retained. - if (UrProgram == nullptr) - throw sycl::exception( - sycl::make_error_code(errc::invalid), - "urProgramCreateWithIL resulted in a null program handle."); - } + const auto spirv = [&]() -> std::vector { + if (Language == syclex::source_language::opencl) { + // if successful, the log is empty. if failed, throws an error with + // the compilation log. + std::vector IPVersionVec(Devices.size()); + std::transform(DeviceVec.begin(), DeviceVec.end(), + IPVersionVec.begin(), [&](ur_device_handle_t d) { + uint32_t ipVersion = 0; + Adapter->call( + d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t), + &ipVersion, nullptr); + return ipVersion; + }); + return syclex::detail::OpenCLC_to_SPIRV(*SourceStrPtr, IPVersionVec, + BuildOptions, LogPtr); + } + if (Language == syclex::source_language::spirv) { + const auto &SourceBytes = + std::get>(this->Source); + std::vector Result(SourceBytes.size()); + std::transform(SourceBytes.cbegin(), SourceBytes.cend(), + Result.begin(), + [](std::byte B) { return static_cast(B); }); + return Result; + } + if (Language == syclex::source_language::sycl) { + return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs, + BuildOptions, LogPtr, + RegisteredKernelNames); + } + if (Language == syclex::source_language::sycl_jit) { + auto *Binaries = syclex::detail::SYCL_JIT_to_SPIRV( + *SourceStrPtr, IncludePairs, BuildOptions, LogPtr, + RegisteredKernelNames); + assert(Binaries->NumDeviceBinaries == 1 && + "Device code splitting is not yet supported"); + return std::vector(Binaries->DeviceBinaries->BinaryStart, + Binaries->DeviceBinaries->BinaryEnd); + } + throw sycl::exception( + make_error_code(errc::invalid), + "SYCL C++, OpenCL C and SPIR-V are the only supported " + "languages at this time"); + }(); + + Adapter->call( + ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr, + &UrProgram); + // program created by urProgramCreateWithIL is implicitly retained. + if (UrProgram == nullptr) + throw sycl::exception( + sycl::make_error_code(errc::invalid), + "urProgramCreateWithIL resulted in a null program handle."); + } // if(!FetchedFromCache) std::string XsFlags = extractXsFlags(BuildOptions); diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 10e1755370e00..6362bf355cfc5 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -342,7 +342,7 @@ bool SYCL_JIT_Compilation_Available() { #endif } -const sycl::detail::RTDeviceBinaryImage &SYCL_JIT_to_SPIRV( +sycl_device_binaries SYCL_JIT_to_SPIRV( [[maybe_unused]] const std::string &SYCLSource, [[maybe_unused]] include_pairs_t IncludePairs, [[maybe_unused]] const std::vector &UserArgs, diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index 8fce2a2045c32..a9ac1410fa52b 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -35,7 +35,7 @@ bool SYCL_Compilation_Available(); std::string userArgsAsString(const std::vector &UserArguments); -const sycl::detail::RTDeviceBinaryImage & +sycl_device_binaries SYCL_JIT_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs, const std::vector &UserArgs, std::string *LogPtr, const std::vector &RegisteredKernelNames); From 0a65e1b5e5530532fcae37f6047d0e5d05eb0eee Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Wed, 20 Nov 2024 13:21:23 +0000 Subject: [PATCH 04/12] Add missing forward declaration Signed-off-by: Julian Oppermann --- sycl/source/detail/jit_compiler.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/source/detail/jit_compiler.hpp b/sycl/source/detail/jit_compiler.hpp index 982246b48da5f..b673e4d37b8fa 100644 --- a/sycl/source/detail/jit_compiler.hpp +++ b/sycl/source/detail/jit_compiler.hpp @@ -23,6 +23,7 @@ enum class BinaryFormat : uint32_t; class JITContext; struct SYCLKernelInfo; struct SYCLKernelAttribute; +struct RTCBundleInfo; template class DynArray; using ArgUsageMask = DynArray; using JITEnvVar = DynArray; From be9bebe8d0a17b96e280b34ef6dbef88e2204486 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Wed, 20 Nov 2024 13:40:07 +0000 Subject: [PATCH 05/12] Cleanup Signed-off-by: Julian Oppermann --- sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp | 6 +++--- sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp | 2 +- sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp | 2 ++ 3 files changed, 6 insertions(+), 4 deletions(-) diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index 5931458990d8c..98a838d7b9cf9 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -435,8 +435,9 @@ Expected jit_compiler::performPostLink( for (auto &&[KV, FrozenPropSet] : zip(PropertySets, BundleInfo.Properties)) { const auto &PropertySetName = KV.first; const auto &PropertySet = KV.second; - FrozenPropertySet FPS{PropertySetName.str(), PropertySet.size()}; - for (auto &&[KV2, FrozenProp] : zip(PropertySet, FPS.Values)) { + FrozenPropSet = + FrozenPropertySet{PropertySetName.str(), PropertySet.size()}; + for (auto &&[KV2, FrozenProp] : zip(PropertySet, FrozenPropSet.Values)) { const auto &PropertyName = KV2.first; const auto &PropertyValue = KV2.second; FrozenProp = PropertyValue.getType() == PropertyValue::Type::UINT32 @@ -446,7 +447,6 @@ Expected jit_compiler::performPostLink( PropertyName.str(), PropertyValue.asRawByteArray(), PropertyValue.getRawByteArraySize()}; } - FrozenPropSet = std::move(FPS); }; // Regain ownership of the module. diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index a9ac1410fa52b..eee0f8a78f602 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -12,7 +12,7 @@ #include // __SYCL_EXPORT #include -#include +#include // sycl_device_binaries #include // std::accumulate #include diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index 01f25f813b826..8977caf7d82d4 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -179,6 +179,8 @@ int test_unsupported_options() { CheckUnsupported({"-Xsycl-target-frontend", "-fsanitize=address"}); CheckUnsupported({"-Xsycl-target-frontend=spir64", "-fsanitize=address"}); CheckUnsupported({"-Xarch_device", "-fsanitize=address"}); + CheckUnsupported({"-fsycl-device-code-split=kernel"}); + CheckUnsupported({"-fsycl-device-code-split-esimd"}); return 0; } From 58d035df7c6623a6fafb565e278c6f33d51351ef Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Thu, 21 Nov 2024 11:33:55 +0100 Subject: [PATCH 06/12] Update sycl/source/detail/jit_compiler.cpp Co-authored-by: Lukas Sommer --- sycl/source/detail/jit_compiler.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 742769f76af81..76ea2774ca83c 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1123,7 +1123,7 @@ sycl_device_binaries jit_compiler::createDeviceBinaryImage( const ::jit_compiler::RTCBundleInfo &BundleInfo) { DeviceBinaryContainer Binary; for (const auto &Symbol : BundleInfo.SymbolTable) { - // Create an offload entry each kernel. + // Create an offload entry for each kernel. // It seems to be OK to set zero for most of the information here, at least // that is the case for compiled SPIR-V binaries. OffloadEntryContainer Entry{Symbol.c_str(), nullptr, 0, 0, 0}; From e93076cd6102c62f1d029e089177fb00f123fc3a Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Thu, 21 Nov 2024 10:50:53 +0000 Subject: [PATCH 07/12] Drop uncessary `decltype`. Signed-off-by: Julian Oppermann --- sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index 98a838d7b9cf9..ca308dd50ffdd 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -415,8 +415,7 @@ Expected jit_compiler::performPostLink( MDesc.saveSplitInformationAsMetadata(); RTCBundleInfo BundleInfo; - BundleInfo.SymbolTable = - decltype(BundleInfo.SymbolTable){MDesc.entries().size()}; + BundleInfo.SymbolTable = FrozenSymbolTable{MDesc.entries().size()}; transform(MDesc.entries(), BundleInfo.SymbolTable.begin(), [](Function *F) { return F->getName(); }); @@ -431,7 +430,7 @@ Expected jit_compiler::performPostLink( // `saveModuleProperties`? const auto &PropertySets = Properties.getPropSets(); - BundleInfo.Properties = decltype(BundleInfo.Properties){PropertySets.size()}; + BundleInfo.Properties = FrozenPropertyRegistry{PropertySets.size()}; for (auto &&[KV, FrozenPropSet] : zip(PropertySets, BundleInfo.Properties)) { const auto &PropertySetName = KV.first; const auto &PropertySet = KV.second; From 00fe093dd61d76fe965f9d79e3f000129647bf7b Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Thu, 21 Nov 2024 10:55:38 +0000 Subject: [PATCH 08/12] Use templated error wrapper directly Signed-off-by: Julian Oppermann --- sycl-jit/jit-compiler/lib/KernelFusion.cpp | 47 +++++++++------------- 1 file changed, 19 insertions(+), 28 deletions(-) diff --git a/sycl-jit/jit-compiler/lib/KernelFusion.cpp b/sycl-jit/jit-compiler/lib/KernelFusion.cpp index d5bae664dc886..fae9a3c29dcf4 100644 --- a/sycl-jit/jit-compiler/lib/KernelFusion.cpp +++ b/sycl-jit/jit-compiler/lib/KernelFusion.cpp @@ -26,7 +26,7 @@ using FusedFunction = helper::FusionHelper::FusedFunction; using FusedFunctionList = std::vector; template -static ResultType wrapError(llvm::Error &&Err, const std::string &Msg) { +static ResultType errorTo(llvm::Error &&Err, const std::string &Msg) { std::stringstream ErrMsg; ErrMsg << Msg << "\nDetailed information:\n"; llvm::handleAllErrors(std::move(Err), @@ -38,15 +38,6 @@ static ResultType wrapError(llvm::Error &&Err, const std::string &Msg) { return ResultType{ErrMsg.str().c_str()}; } -static JITResult errorToFusionResult(llvm::Error &&Err, - const std::string &Msg) { - return wrapError(std::move(Err), Msg); -} - -static RTCResult errorToRTCResult(llvm::Error &&Err, const std::string &Msg) { - return wrapError(std::move(Err), Msg); -} - static std::vector gatherNDRanges(llvm::ArrayRef KernelInformation) { std::vector NDRanges; @@ -104,7 +95,7 @@ extern "C" KF_EXPORT_SYMBOL JITResult materializeSpecConstants( translation::KernelTranslator::loadKernels(*JITCtx.getLLVMContext(), ModuleInfo.kernels()); if (auto Error = ModOrError.takeError()) { - return errorToFusionResult(std::move(Error), "Failed to load kernels"); + return errorTo(std::move(Error), "Failed to load kernels"); } std::unique_ptr NewMod = std::move(*ModOrError); if (!fusion::FusionPipeline::runMaterializerPasses( @@ -116,8 +107,8 @@ extern "C" KF_EXPORT_SYMBOL JITResult materializeSpecConstants( SYCLKernelInfo &MaterializerKernelInfo = *ModuleInfo.getKernelFor(KernelName); if (auto Error = translation::KernelTranslator::translateKernel( MaterializerKernelInfo, *NewMod, JITCtx, TargetFormat)) { - return errorToFusionResult(std::move(Error), - "Translation to output format failed"); + return errorTo(std::move(Error), + "Translation to output format failed"); } return JITResult{MaterializerKernelInfo}; @@ -142,7 +133,7 @@ fuseKernels(View KernelInformation, const char *FusedKernelName, llvm::Expected FusedNDR = jit_compiler::FusedNDRange::get(NDRanges); if (llvm::Error Err = FusedNDR.takeError()) { - return errorToFusionResult(std::move(Err), "Illegal ND-range combination"); + return errorTo(std::move(Err), "Illegal ND-range combination"); } if (!isTargetFormatSupported(TargetFormat)) { @@ -189,7 +180,7 @@ fuseKernels(View KernelInformation, const char *FusedKernelName, translation::KernelTranslator::loadKernels(*JITCtx.getLLVMContext(), ModuleInfo.kernels()); if (auto Error = ModOrError.takeError()) { - return errorToFusionResult(std::move(Error), "SPIR-V translation failed"); + return errorTo(std::move(Error), "SPIR-V translation failed"); } std::unique_ptr LLVMMod = std::move(*ModOrError); @@ -206,8 +197,8 @@ fuseKernels(View KernelInformation, const char *FusedKernelName, llvm::Expected> NewModOrError = helper::FusionHelper::addFusedKernel(LLVMMod.get(), FusedKernelList); if (auto Error = NewModOrError.takeError()) { - return errorToFusionResult(std::move(Error), - "Insertion of fused kernel stub failed"); + return errorTo(std::move(Error), + "Insertion of fused kernel stub failed"); } std::unique_ptr NewMod = std::move(*NewModOrError); @@ -230,8 +221,8 @@ fuseKernels(View KernelInformation, const char *FusedKernelName, if (auto Error = translation::KernelTranslator::translateKernel( FusedKernelInfo, *NewMod, JITCtx, TargetFormat)) { - return errorToFusionResult(std::move(Error), - "Translation to output format failed"); + return errorTo(std::move(Error), + "Translation to output format failed"); } FusedKernelInfo.NDR = FusedNDR->getNDR(); @@ -248,15 +239,15 @@ compileSYCL(InMemoryFile SourceFile, View IncludeFiles, View UserArgs) { auto UserArgListOrErr = parseUserArgs(UserArgs); if (!UserArgListOrErr) { - return errorToRTCResult(UserArgListOrErr.takeError(), - "Parsing of user arguments failed"); + return errorTo(UserArgListOrErr.takeError(), + "Parsing of user arguments failed"); } llvm::opt::InputArgList UserArgList = std::move(*UserArgListOrErr); auto ModuleOrErr = compileDeviceCode(SourceFile, IncludeFiles, UserArgList); if (!ModuleOrErr) { - return errorToRTCResult(ModuleOrErr.takeError(), - "Device compilation failed"); + return errorTo(ModuleOrErr.takeError(), + "Device compilation failed"); } std::unique_ptr Context; @@ -264,13 +255,13 @@ compileSYCL(InMemoryFile SourceFile, View IncludeFiles, Context.reset(&Module->getContext()); if (auto Error = linkDeviceLibraries(*Module, UserArgList)) { - return errorToRTCResult(std::move(Error), "Device linking failed"); + return errorTo(std::move(Error), "Device linking failed"); } auto BundleInfoOrError = performPostLink(*Module, UserArgList); if (!BundleInfoOrError) { - return errorToRTCResult(BundleInfoOrError.takeError(), - "Post-link phase failed"); + return errorTo(BundleInfoOrError.takeError(), + "Post-link phase failed"); } auto BundleInfo = std::move(*BundleInfoOrError); @@ -278,8 +269,8 @@ compileSYCL(InMemoryFile SourceFile, View IncludeFiles, translation::KernelTranslator::translateBundleToSPIRV( *Module, JITContext::getInstance()); if (!BinaryInfoOrError) { - return errorToRTCResult(BinaryInfoOrError.takeError(), - "SPIR-V translation failed"); + return errorTo(BinaryInfoOrError.takeError(), + "SPIR-V translation failed"); } BundleInfo.BinaryInfo = std::move(*BinaryInfoOrError); From 8b2c1968f2a71ec1dde78ba7e62b7ab250f2fa80 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Thu, 21 Nov 2024 15:55:38 +0000 Subject: [PATCH 09/12] Help old GCC figure out the construction of the `Expected` return object. Signed-off-by: Julian Oppermann --- sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index ca308dd50ffdd..8bbcdf0c0febd 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -451,7 +451,7 @@ Expected jit_compiler::performPostLink( // Regain ownership of the module. MDesc.releaseModulePtr().release(); - return BundleInfo; + return std::move(BundleInfo); } Expected @@ -520,5 +520,5 @@ jit_compiler::parseUserArgs(View UserArgs) { "Runtime compilation of ESIMD kernels is not yet supported"); } - return Expected{std::move(AL)}; + return std::move(AL); } From c5cc2f36f165e427b7abbfc81176c54a3145bfdb Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 25 Nov 2024 09:54:42 +0000 Subject: [PATCH 10/12] Comments for literal arguments Signed-off-by: Julian Oppermann --- sycl/source/detail/jit_compiler.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 76ea2774ca83c..8e466b97d75cf 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1126,7 +1126,8 @@ sycl_device_binaries jit_compiler::createDeviceBinaryImage( // Create an offload entry for each kernel. // It seems to be OK to set zero for most of the information here, at least // that is the case for compiled SPIR-V binaries. - OffloadEntryContainer Entry{Symbol.c_str(), nullptr, 0, 0, 0}; + OffloadEntryContainer Entry{Symbol.c_str(), /*Addr=*/nullptr, /*Size=*/0, + /*Flags=*/0, /*Reserved=*/0}; Binary.addOffloadEntry(std::move(Entry)); } From ec89eeb9c48cb4bb8a7e999e53cfe07075b8f20b Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 25 Nov 2024 10:04:53 +0000 Subject: [PATCH 11/12] Drop duplicate comment. Signed-off-by: Julian Oppermann --- sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp b/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp index 82fac2f824445..7a35325cddaf5 100644 --- a/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp @@ -51,13 +51,6 @@ bool jit_compiler::post_link::removeSYCLKernelsConstRefArray(Module &M) { return true; } -// Removes all device_global variables from the llvm.compiler.used global -// variable. A device_global with internal linkage will be in llvm.compiler.used -// to avoid the compiler wrongfully removing it during optimizations. However, -// as an effect the device_global variables will also be distributed across -// binaries, even if llvm.compiler.used has served its purpose. To avoid -// polluting other binaries with unused device_global variables, we remove them -// from llvm.compiler.used and erase them if they have no further uses. bool jit_compiler::post_link::removeDeviceGlobalFromCompilerUsed(Module &M) { GlobalVariable *GV = M.getGlobalVariable("llvm.compiler.used"); if (!GV) From ffdb19747732d19f187973bd54765e5fb902d2bd Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 25 Nov 2024 11:07:45 +0000 Subject: [PATCH 12/12] Drop copied transformation from sycl-post-link Signed-off-by: Julian Oppermann --- sycl-jit/jit-compiler/CMakeLists.txt | 1 - .../lib/rtc/DeviceCompilation.cpp | 33 +++--- .../jit-compiler/lib/rtc/PostLinkActions.cpp | 109 ------------------ .../jit-compiler/lib/rtc/PostLinkActions.h | 50 -------- 4 files changed, 16 insertions(+), 177 deletions(-) delete mode 100644 sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp delete mode 100644 sycl-jit/jit-compiler/lib/rtc/PostLinkActions.h diff --git a/sycl-jit/jit-compiler/CMakeLists.txt b/sycl-jit/jit-compiler/CMakeLists.txt index a704b4c91a503..82d2356691c3f 100644 --- a/sycl-jit/jit-compiler/CMakeLists.txt +++ b/sycl-jit/jit-compiler/CMakeLists.txt @@ -8,7 +8,6 @@ add_llvm_library(sycl-jit lib/fusion/JITContext.cpp lib/fusion/ModuleHelper.cpp lib/rtc/DeviceCompilation.cpp - lib/rtc/PostLinkActions.cpp lib/helper/ConfigHelper.cpp SHARED diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index 8bbcdf0c0febd..a89635b79ed46 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -8,8 +8,6 @@ #include "DeviceCompilation.h" -#include "PostLinkActions.h" - #include #include #include @@ -20,6 +18,8 @@ #include #include +#include +#include #include #include #include @@ -37,7 +37,6 @@ using namespace llvm::sycl; using namespace llvm::module_split; using namespace llvm::util; using namespace jit_compiler; -using namespace jit_compiler::post_link; #ifdef _GNU_SOURCE #include @@ -365,26 +364,26 @@ Error jit_compiler::linkDeviceLibraries(llvm::Module &Module, return Error::success(); } +template static bool runModulePass(llvm::Module &M) { + ModulePassManager MPM; + ModuleAnalysisManager MAM; + // Register required analysis + MAM.registerPass([&] { return PassInstrumentationAnalysis(); }); + MPM.addPass(PassClass{}); + PreservedAnalyses Res = MPM.run(M, MAM); + return !Res.areAllPreserved(); +} + Expected jit_compiler::performPostLink( llvm::Module &Module, [[maybe_unused]] const InputArgList &UserArgList) { // This is a simplified version of `processInputModule` in // `llvm/tools/sycl-post-link.cpp`. Assertions/TODOs point to functionality // left out of the algorithm for now. - // After linking device bitcode "llvm.used" holds references to the kernels - // that are defined in the device image. But after splitting device image into - // separate kernels we may end up with having references to kernel declaration - // originating from "llvm.used" in the IR that is passed to llvm-spirv tool, - // and these declarations cause an assertion in llvm-spirv. To workaround this - // issue remove "llvm.used" from the input module before performing any other - // actions. - removeSYCLKernelsConstRefArray(Module); - - // There may be device_global variables kept alive in "llvm.compiler.used" - // to keep the optimizer from wrongfully removing them. llvm.compiler.used - // symbols are usually removed at backend lowering, but this is handled here - // for SPIR-V since SYCL compilation uses llvm-spirv, not the SPIR-V backend. - removeDeviceGlobalFromCompilerUsed(Module); + assert(!Module.getGlobalVariable("llvm.used") && + !Module.getGlobalVariable("llvm.compiler.used")); + // Otherwise: Port over the `removeSYCLKernelsConstRefArray` and + // `removeDeviceGlobalFromCompilerUsed` methods. assert(!isModuleUsingAsan(Module)); // Otherwise: Need to instrument each image scope device globals if the module diff --git a/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp b/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp deleted file mode 100644 index 7a35325cddaf5..0000000000000 --- a/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.cpp +++ /dev/null @@ -1,109 +0,0 @@ -//==------------------------ PostLinkActions.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 "PostLinkActions.h" - -#include -#include -#include - -using namespace llvm; - -bool jit_compiler::post_link::removeSYCLKernelsConstRefArray(Module &M) { - GlobalVariable *GV = M.getGlobalVariable("llvm.used"); - - if (!GV) { - return false; - } - assert(GV->user_empty() && "Unexpected llvm.used users"); - Constant *Initializer = GV->getInitializer(); - GV->setInitializer(nullptr); - GV->eraseFromParent(); - - // Destroy the initializer and all operands of it. - SmallVector IOperands; - for (auto It = Initializer->op_begin(); It != Initializer->op_end(); It++) - IOperands.push_back(cast(*It)); - assert(llvm::isSafeToDestroyConstant(Initializer) && - "Cannot remove initializer of llvm.used global"); - Initializer->destroyConstant(); - for (auto It = IOperands.begin(); It != IOperands.end(); It++) { - auto Op = (*It)->stripPointerCasts(); - auto *F = dyn_cast(Op); - if (llvm::isSafeToDestroyConstant(*It)) { - (*It)->destroyConstant(); - } else if (F && F->getCallingConv() == CallingConv::SPIR_KERNEL && - !F->use_empty()) { - // The element in "llvm.used" array has other users. That is Ok for - // specialization constants, but is wrong for kernels. - llvm::report_fatal_error("Unexpected usage of SYCL kernel"); - } - - // Remove unused kernel declarations to avoid LLVM IR check fails. - if (F && F->isDeclaration() && F->use_empty()) - F->eraseFromParent(); - } - return true; -} - -bool jit_compiler::post_link::removeDeviceGlobalFromCompilerUsed(Module &M) { - GlobalVariable *GV = M.getGlobalVariable("llvm.compiler.used"); - if (!GV) - return false; - - // Erase the old llvm.compiler.used. A new one will be created at the end if - // there are other values in it (other than device_global). - assert(GV->user_empty() && "Unexpected llvm.compiler.used users"); - Constant *Initializer = GV->getInitializer(); - const auto *VAT = cast(GV->getValueType()); - GV->setInitializer(nullptr); - GV->eraseFromParent(); - - // Destroy the initializer. Keep the operands so we keep the ones we need. - SmallVector IOperands; - for (auto It = Initializer->op_begin(); It != Initializer->op_end(); It++) - IOperands.push_back(cast(*It)); - assert(llvm::isSafeToDestroyConstant(Initializer) && - "Cannot remove initializer of llvm.compiler.used global"); - Initializer->destroyConstant(); - - // Iterate through all operands. If they are device_global then we drop them - // and erase them if they have no uses afterwards. All other values are kept. - SmallVector NewOperands; - for (auto It = IOperands.begin(); It != IOperands.end(); It++) { - Constant *Op = *It; - auto *DG = dyn_cast(Op->stripPointerCasts()); - - // If it is not a device_global we keep it. - if (!DG || !isDeviceGlobalVariable(*DG)) { - NewOperands.push_back(Op); - continue; - } - - // Destroy the device_global operand. - if (llvm::isSafeToDestroyConstant(Op)) - Op->destroyConstant(); - - // Remove device_global if it no longer has any uses. - if (!DG->isConstantUsed()) - DG->eraseFromParent(); - } - - // If we have any operands left from the original llvm.compiler.used we create - // a new one with the new size. - if (!NewOperands.empty()) { - ArrayType *ATy = ArrayType::get(VAT->getElementType(), NewOperands.size()); - GlobalVariable *NGV = - new GlobalVariable(M, ATy, false, GlobalValue::AppendingLinkage, - ConstantArray::get(ATy, NewOperands), ""); - NGV->setName("llvm.compiler.used"); - NGV->setSection("llvm.metadata"); - } - - return true; -} diff --git a/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.h b/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.h deleted file mode 100644 index 878e4244b2996..0000000000000 --- a/sycl-jit/jit-compiler/lib/rtc/PostLinkActions.h +++ /dev/null @@ -1,50 +0,0 @@ -//==------ PostLinkActions.h - Fork of sycl-post-link actions for RTC ------==// -// -// 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_POST_LINK_ACTIONS_H -#define SYCL_JIT_COMPILER_RTC_POST_LINK_ACTIONS_H - -#include -#include -#include - -namespace jit_compiler::post_link { - -using namespace llvm; - -template bool runModulePass(Module &M) { - ModulePassManager MPM; - ModuleAnalysisManager MAM; - // Register required analysis - MAM.registerPass([&] { return PassInstrumentationAnalysis(); }); - MPM.addPass(PassClass{}); - PreservedAnalyses Res = MPM.run(M, MAM); - return !Res.areAllPreserved(); -} - -// Removes the global variable "llvm.used" and returns true on success. -// "llvm.used" is a global constant array containing references to kernels -// available in the module and callable from host code. The elements of -// the array are ConstantExpr bitcast to i8*. -// The variable must be removed as it is a) has done the job to the moment -// of this function call and b) the references to the kernels callable from -// host must not have users. -bool removeSYCLKernelsConstRefArray(Module &M); - -// Removes all device_global variables from the llvm.compiler.used global -// variable. A device_global with internal linkage will be in llvm.compiler.used -// to avoid the compiler wrongfully removing it during optimizations. However, -// as an effect the device_global variables will also be distributed across -// binaries, even if llvm.compiler.used has served its purpose. To avoid -// polluting other binaries with unused device_global variables, we remove them -// from llvm.compiler.used and erase them if they have no further uses. -bool removeDeviceGlobalFromCompilerUsed(llvm::Module &M); - -} // namespace jit_compiler::post_link - -#endif // SYCL_JIT_COMPILER_RTC_POST_LINK_ACTIONS_H