diff --git a/llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h b/llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h index bfa641f01c0e2..ef379b815b8de 100644 --- a/llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h +++ b/llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h @@ -17,6 +17,8 @@ #include "llvm/ADT/SmallVector.h" #include "llvm/Support/Error.h" +#include + namespace llvm { namespace sycl { @@ -54,8 +56,8 @@ bool lowerESIMDConstructs(llvm::module_split::ModuleDesc &MD, /// \p Modified value indicates whether the Module has been modified. /// \p SplitOccurred value indicates whether split has occurred before or during /// function's invocation. -Expected> -handleESIMD(llvm::module_split::ModuleDesc MDesc, +Expected, 2>> +handleESIMD(std::unique_ptr MDesc, const ESIMDProcessingOptions &Options, bool &Modified, bool &SplitOccurred); diff --git a/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h b/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h index 144f99b753e7e..aa6afea1f05c1 100644 --- a/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h +++ b/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h @@ -141,7 +141,7 @@ class ModuleDesc { std::string Name = ""; Properties Props; - ModuleDesc(std::unique_ptr &&M, StringRef Name = "TOP-LEVEL") + ModuleDesc(std::unique_ptr M, StringRef Name = "TOP-LEVEL") : M(std::move(M)), IsTopLevel(true), Name(Name) { // DeviceLib module doesn't include any entry point,it can be constructed // using ctor without any entry point related parameter. @@ -153,13 +153,13 @@ class ModuleDesc { } } - ModuleDesc(std::unique_ptr &&M, EntryPointGroup &&EntryPoints, + ModuleDesc(std::unique_ptr M, EntryPointGroup &&EntryPoints, const Properties &Props) : M(std::move(M)), EntryPoints(std::move(EntryPoints)), Props(Props) { Name = this->EntryPoints.GroupId; } - ModuleDesc(std::unique_ptr &&M, const std::vector &Names, + ModuleDesc(std::unique_ptr M, const std::vector &Names, StringRef Name = "NoName") : M(std::move(M)), Name(Name) { rebuildEntryPoints(Names); @@ -225,7 +225,7 @@ class ModuleDesc { bool isSpecConstantDefault() const; void setSpecConstantDefault(bool Value); - ModuleDesc clone() const; + std::unique_ptr clone() const; std::string makeSymbolTable() const; @@ -252,7 +252,7 @@ class ModuleDesc { // from input module that should be included in a split module. class ModuleSplitterBase { protected: - ModuleDesc Input; + std::unique_ptr Input; EntryPointGroupVec Groups; bool AllowDeviceImageDependencies; @@ -264,14 +264,15 @@ class ModuleSplitterBase { return Res; } - Module &getInputModule() { return Input.getModule(); } + Module &getInputModule() { return Input->getModule(); } std::unique_ptr releaseInputModule() { - return Input.releaseModulePtr(); + return Input->releaseModulePtr(); } public: - ModuleSplitterBase(ModuleDesc &&MD, EntryPointGroupVec &&GroupVec, + ModuleSplitterBase(std::unique_ptr MD, + EntryPointGroupVec &&GroupVec, bool AllowDeviceImageDependencies) : Input(std::move(MD)), Groups(std::move(GroupVec)), AllowDeviceImageDependencies(AllowDeviceImageDependencies) { @@ -288,7 +289,7 @@ class ModuleSplitterBase { // Gets next subsequence of entry points in an input module and provides split // submodule containing these entry points and their dependencies. - virtual ModuleDesc nextSplit() = 0; + virtual std::unique_ptr nextSplit() = 0; // Returns a number of remaining modules, which can be split out using this // splitter. The value is reduced by 1 each time nextSplit is called. @@ -298,13 +299,13 @@ class ModuleSplitterBase { bool hasMoreSplits() const { return remainingSplits() > 0; } }; -SmallVector splitByESIMD(ModuleDesc &&MD, - bool EmitOnlyKernelsAsEntryPoints, - bool AllowDeviceImageDependencies); +SmallVector, 2> +splitByESIMD(std::unique_ptr MD, bool EmitOnlyKernelsAsEntryPoints, + bool AllowDeviceImageDependencies); std::unique_ptr -getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, - bool EmitOnlyKernelsAsEntryPoints, +getDeviceCodeSplitter(std::unique_ptr MD, IRSplitMode Mode, + bool IROutputOnly, bool EmitOnlyKernelsAsEntryPoints, bool AllowDeviceImageDependencies); #if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP) diff --git a/llvm/include/llvm/SYCLPostLink/SpecializationConstants.h b/llvm/include/llvm/SYCLPostLink/SpecializationConstants.h index afdb7a4be1c93..e01adfc153f5c 100644 --- a/llvm/include/llvm/SYCLPostLink/SpecializationConstants.h +++ b/llvm/include/llvm/SYCLPostLink/SpecializationConstants.h @@ -16,6 +16,7 @@ #include "llvm/SYCLLowerIR/SpecConstants.h" #include "llvm/SYCLPostLink/ModuleSplitter.h" +#include #include namespace llvm { @@ -31,9 +32,10 @@ namespace sycl { /// \returns Boolean value indicating whether the lowering has changed the input /// modules. bool handleSpecializationConstants( - llvm::SmallVectorImpl &MDs, + llvm::SmallVectorImpl> &MDs, std::optional Mode, - llvm::SmallVectorImpl &NewModuleDescs, + llvm::SmallVectorImpl> + &NewModuleDescs, bool GenerateModuleDescWithDefaultSpecConsts); } // namespace sycl diff --git a/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp b/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp index f0df7fa2b3a59..a1c1a0832d945 100644 --- a/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp +++ b/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp @@ -63,20 +63,22 @@ buildESIMDLoweringPipeline(const sycl::ESIMDProcessingOptions &Options) { return MPM; } -Expected linkModules(ModuleDesc MD1, ModuleDesc MD2) { +Expected> +linkModules(std::unique_ptr MD1, std::unique_ptr MD2) { std::vector Names; - MD1.saveEntryPointNames(Names); - MD2.saveEntryPointNames(Names); + MD1->saveEntryPointNames(Names); + MD2->saveEntryPointNames(Names); bool LinkError = - llvm::Linker::linkModules(MD1.getModule(), MD2.releaseModulePtr()); + llvm::Linker::linkModules(MD1->getModule(), MD2->releaseModulePtr()); if (LinkError) return createStringError( - formatv("link failed. Module names: {0}, {1}", MD1.Name, MD2.Name)); + formatv("link failed. Module names: {0}, {1}", MD1->Name, MD2->Name)); - ModuleDesc Res(MD1.releaseModulePtr(), std::move(Names)); - Res.assignMergedProperties(MD1, MD2); - Res.Name = (Twine("linked[") + MD1.Name + "," + MD2.Name + "]").str(); + auto Res = + std::make_unique(MD1->releaseModulePtr(), std::move(Names)); + Res->assignMergedProperties(*MD1, *MD2); + Res->Name = (Twine("linked[") + MD1->Name + "," + MD2->Name + "]").str(); return std::move(Res); } @@ -110,11 +112,11 @@ bool sycl::lowerESIMDConstructs(ModuleDesc &MD, return !Res.areAllPreserved(); } -Expected> -llvm::sycl::handleESIMD(ModuleDesc MDesc, +Expected, 2>> +llvm::sycl::handleESIMD(std::unique_ptr MDesc, const sycl::ESIMDProcessingOptions &Options, bool &Modified, bool &SplitOccurred) { - SmallVector Result = + SmallVector, 2> Result = splitByESIMD(std::move(MDesc), Options.EmitOnlyKernelsAsEntryPoints, Options.AllowDeviceImageDependencies); @@ -123,32 +125,32 @@ llvm::sycl::handleESIMD(ModuleDesc MDesc, SplitOccurred |= Result.size() > 1; - for (ModuleDesc &MD : Result) - if (Options.LowerESIMD && MD.isESIMD()) - Modified |= lowerESIMDConstructs(MD, Options); + for (std::unique_ptr &MD : Result) + if (Options.LowerESIMD && MD->isESIMD()) + Modified |= lowerESIMDConstructs(*MD, Options); if (Options.SplitESIMD || Result.size() == 1) return std::move(Result); // SYCL/ESIMD splitting is not requested, link back into single module. - int ESIMDInd = Result[0].isESIMD() ? 0 : 1; + int ESIMDInd = Result[0]->isESIMD() ? 0 : 1; int SYCLInd = 1 - ESIMDInd; - assert(Result[SYCLInd].isSYCL() && - "Result[SYCLInd].isSYCL() expected to be true."); + assert(Result[SYCLInd]->isSYCL() && + "Result[SYCLInd]->isSYCL() expected to be true."); // Make sure that no link conflicts occur. - Result[ESIMDInd].renameDuplicatesOf(Result[SYCLInd].getModule(), ".esimd"); + Result[ESIMDInd]->renameDuplicatesOf(Result[SYCLInd]->getModule(), ".esimd"); auto LinkedOrErr = linkModules(std::move(Result[0]), std::move(Result[1])); if (!LinkedOrErr) return LinkedOrErr.takeError(); - ModuleDesc &Linked = *LinkedOrErr; - Linked.restoreLinkageOfDirectInvokeSimdTargets(); + std::unique_ptr &Linked = *LinkedOrErr; + Linked->restoreLinkageOfDirectInvokeSimdTargets(); std::vector Names; - Linked.saveEntryPointNames(Names); + Linked->saveEntryPointNames(Names); // Cleanup may remove some entry points, need to save/rebuild. - Linked.cleanup(Options.AllowDeviceImageDependencies); - Linked.rebuildEntryPoints(Names); + Linked->cleanup(Options.AllowDeviceImageDependencies); + Linked->rebuildEntryPoints(Names); Result.clear(); Result.emplace_back(std::move(Linked)); Modified = true; diff --git a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp b/llvm/lib/SYCLPostLink/ModuleSplitter.cpp index 30647bd826040..33b4dfb26d740 100644 --- a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp +++ b/llvm/lib/SYCLPostLink/ModuleSplitter.cpp @@ -365,9 +365,9 @@ void processSubModuleNamedMetadata(Module *M) { } } -ModuleDesc extractSubModule(const ModuleDesc &MD, - const SetVector GVs, - EntryPointGroup &&ModuleEntryPoints) { +std::unique_ptr +extractSubModule(const ModuleDesc &MD, const SetVector GVs, + EntryPointGroup &&ModuleEntryPoints) { const Module &M = MD.getModule(); // For each group of entry points collect all dependencies. ValueToValueMapTy VMap; @@ -382,13 +382,14 @@ ModuleDesc extractSubModule(const ModuleDesc &MD, NewEPs.insert(cast(VMap[F])); }); ModuleEntryPoints.Functions = std::move(NewEPs); - return ModuleDesc{std::move(SubM), std::move(ModuleEntryPoints), MD.Props}; + return std::make_unique(std::move(SubM), + std::move(ModuleEntryPoints), MD.Props); } // The function produces a copy of input LLVM IR module M with only those // functions and globals that can be called from entry points that are specified // in ModuleEntryPoints vector, in addition to the entry point functions. -ModuleDesc extractCallGraph( +std::unique_ptr extractCallGraph( const ModuleDesc &MD, EntryPointGroup &&ModuleEntryPoints, const DependencyGraph &CG, bool AllowDeviceImageDependencies, const std::function &IncludeFunctionPredicate = @@ -397,14 +398,14 @@ ModuleDesc extractCallGraph( collectFunctionsAndGlobalVariablesToExtract( GVs, MD.getModule(), ModuleEntryPoints, CG, IncludeFunctionPredicate); - ModuleDesc SplitM = + std::unique_ptr SplitM = extractSubModule(MD, std::move(GVs), std::move(ModuleEntryPoints)); // TODO: cleanup pass is now called for each output module at the end of // sycl-post-link. This call is redundant. However, we subsequently run // GenXSPIRVWriterAdaptor pass that relies on this cleanup. This cleanup call // can be removed once that pass no longer depends on this cleanup. - SplitM.cleanup(AllowDeviceImageDependencies); - checkForCallsToUndefinedFunctions(SplitM.getModule(), + SplitM->cleanup(AllowDeviceImageDependencies); + checkForCallsToUndefinedFunctions(SplitM->getModule(), AllowDeviceImageDependencies); return SplitM; @@ -413,7 +414,7 @@ ModuleDesc extractCallGraph( // The function is similar to 'extractCallGraph', but it produces a copy of // input LLVM IR module M with _all_ ESIMD functions and kernels included, // regardless of whether or not they are listed in ModuleEntryPoints. -ModuleDesc extractESIMDSubModule( +std::unique_ptr extractESIMDSubModule( const ModuleDesc &MD, EntryPointGroup &&ModuleEntryPoints, const DependencyGraph &CG, bool AllowDeviceImageDependencies, const std::function &IncludeFunctionPredicate = @@ -426,13 +427,13 @@ ModuleDesc extractESIMDSubModule( collectFunctionsAndGlobalVariablesToExtract( GVs, MD.getModule(), ModuleEntryPoints, CG, IncludeFunctionPredicate); - ModuleDesc SplitM = + std::unique_ptr SplitM = extractSubModule(MD, std::move(GVs), std::move(ModuleEntryPoints)); // TODO: cleanup pass is now called for each output module at the end of // sycl-post-link. This call is redundant. However, we subsequently run // GenXSPIRVWriterAdaptor pass that relies on this cleanup. This cleanup call // can be removed once that pass no longer depends on this cleanup. - SplitM.cleanup(AllowDeviceImageDependencies); + SplitM->cleanup(AllowDeviceImageDependencies); return SplitM; } @@ -441,29 +442,30 @@ class ModuleCopier : public ModuleSplitterBase { public: using ModuleSplitterBase::ModuleSplitterBase; // to inherit base constructors - ModuleDesc nextSplit() override { - ModuleDesc Desc{releaseInputModule(), nextGroup(), Input.Props}; + std::unique_ptr nextSplit() override { + auto Desc = std::make_unique(releaseInputModule(), nextGroup(), + Input->Props); // Do some basic optimization like unused symbol removal // even if there was no split. // TODO: cleanup pass is now called for each output module at the end of // sycl-post-link. This call is redundant. However, we subsequently run // GenXSPIRVWriterAdaptor pass that relies on this cleanup. This cleanup // call can be removed once that pass no longer depends on this cleanup. - Desc.cleanup(AllowDeviceImageDependencies); + Desc->cleanup(AllowDeviceImageDependencies); return Desc; } }; class ModuleSplitter : public ModuleSplitterBase { public: - ModuleSplitter(ModuleDesc &&MD, EntryPointGroupVec &&GroupVec, + ModuleSplitter(std::unique_ptr MD, EntryPointGroupVec &&GroupVec, bool AllowDeviceImageDependencies) : ModuleSplitterBase(std::move(MD), std::move(GroupVec), AllowDeviceImageDependencies), - CG(Input.getModule(), AllowDeviceImageDependencies) {} + CG(Input->getModule(), AllowDeviceImageDependencies) {} - ModuleDesc nextSplit() override { - return extractCallGraph(Input, nextGroup(), CG, + std::unique_ptr nextSplit() override { + return extractCallGraph(*Input, nextGroup(), CG, AllowDeviceImageDependencies); } @@ -760,10 +762,10 @@ void ModuleDesc::setSpecConstantDefault(bool Value) { Props.IsSpecConstantDefault = Value; } -ModuleDesc ModuleDesc::clone() const { +std::unique_ptr ModuleDesc::clone() const { std::unique_ptr NewModule = CloneModule(getModule()); - ModuleDesc NewMD(std::move(NewModule)); - NewMD.EntryPoints.Props = EntryPoints.Props; + auto NewMD = std::make_unique(std::move(NewModule)); + NewMD->EntryPoints.Props = EntryPoints.Props; return NewMD; } @@ -1042,18 +1044,18 @@ std::string computeFuncCategoryForSplitting(const Function &F, } // namespace std::unique_ptr -getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, - bool EmitOnlyKernelsAsEntryPoints, +getDeviceCodeSplitter(std::unique_ptr MD, IRSplitMode Mode, + bool IROutputOnly, bool EmitOnlyKernelsAsEntryPoints, bool AllowDeviceImageDependencies) { EntryPointsGroupScope Scope = - selectDeviceCodeGroupScope(MD.getModule(), Mode, IROutputOnly); + selectDeviceCodeGroupScope(MD->getModule(), Mode, IROutputOnly); // std::map is used here to ensure stable ordering of entry point groups, // which is based on their contents, this greatly helps LIT tests std::map EntryPointsMap; // Only process module entry points: - for (auto &F : MD.getModule().functions()) { + for (auto &F : MD->getModule().functions()) { if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints)) continue; @@ -1069,7 +1071,7 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, } else { Groups.reserve(EntryPointsMap.size()); // Start with properties of a source module - EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props; + EntryPointGroup::Properties MDProps = MD->getEntryPointGroup().Props; for (auto &[Key, EntryPoints] : EntryPointsMap) { bool HasVirtualFunctions = false; for (auto *F : EntryPoints) { @@ -1115,21 +1117,21 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, // invoke_simd, for example), the modules has to be linked back together to // avoid undefined behavior at later stages. That is done at higher level, // outside of this function. -SmallVector splitByESIMD(ModuleDesc &&MD, - bool EmitOnlyKernelsAsEntryPoints, - bool AllowDeviceImageDependencies) { +SmallVector, 2> +splitByESIMD(std::unique_ptr MD, bool EmitOnlyKernelsAsEntryPoints, + bool AllowDeviceImageDependencies) { - SmallVector Result; + SmallVector, 2> Result; EntryPointGroupVec EntryPointGroups{}; EntryPointSet SYCLEntryPoints, ESIMDEntryPoints; bool hasESIMDFunctions = false; // Only process module entry points: - for (Function &F : MD.getModule().functions()) { + for (Function &F : MD->getModule().functions()) { if (isESIMDFunction(F)) hasESIMDFunctions = true; if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints) || - !MD.isEntryPointCandidate(F)) + !MD->isEntryPointCandidate(F)) continue; if (isESIMDFunction(F)) ESIMDEntryPoints.insert(&F); @@ -1142,30 +1144,31 @@ SmallVector splitByESIMD(ModuleDesc &&MD, // can lower the ESIMD functions. if (!ESIMDEntryPoints.empty() || hasESIMDFunctions) { EntryPointGroups.emplace_back(ESIMD_SCOPE_NAME, std::move(ESIMDEntryPoints), - MD.getEntryPointGroup().Props); + MD->getEntryPointGroup().Props); EntryPointGroup &G = EntryPointGroups.back(); G.Props.HasESIMD = SyclEsimdSplitStatus::ESIMD_ONLY; } if (!SYCLEntryPoints.empty() || EntryPointGroups.empty()) { EntryPointGroups.emplace_back(SYCL_SCOPE_NAME, std::move(SYCLEntryPoints), - MD.getEntryPointGroup().Props); + MD->getEntryPointGroup().Props); EntryPointGroup &G = EntryPointGroups.back(); G.Props.HasESIMD = SyclEsimdSplitStatus::SYCL_ONLY; } if (EntryPointGroups.size() == 1) { - Result.emplace_back(MD.releaseModulePtr(), std::move(EntryPointGroups[0]), - MD.Props); + auto MD2 = std::make_unique( + MD->releaseModulePtr(), std::move(EntryPointGroups[0]), MD->Props); + Result.emplace_back(std::move(MD2)); return Result; } - DependencyGraph CG(MD.getModule(), AllowDeviceImageDependencies); + DependencyGraph CG(MD->getModule(), AllowDeviceImageDependencies); for (auto &Group : EntryPointGroups) { if (Group.isEsimd()) { // For ESIMD module, we use full call graph of all entry points and all // ESIMD functions. - Result.emplace_back(extractESIMDSubModule(MD, std::move(Group), CG, + Result.emplace_back(extractESIMDSubModule(*MD, std::move(Group), CG, AllowDeviceImageDependencies)); } else { // For non-ESIMD module we only use non-ESIMD functions. Additional filter @@ -1175,7 +1178,7 @@ SmallVector splitByESIMD(ModuleDesc &&MD, // were processed and therefore it is fine to return an "incomplete" // module here. Result.emplace_back(extractCallGraph( - MD, std::move(Group), CG, AllowDeviceImageDependencies, + *MD, std::move(Group), CG, AllowDeviceImageDependencies, [=](const Function *F) -> bool { return !isESIMDFunction(*F); })); } } @@ -1314,7 +1317,7 @@ bool runPreSplitProcessingPipeline(Module &M) { Expected> splitSYCLModule(std::unique_ptr M, ModuleSplitterSettings Settings) { - ModuleDesc MD = std::move(M); // makeModuleDesc() ? + auto MD = std::make_unique(std::move(M)); // FIXME: false arguments are temporary for now. auto Splitter = getDeviceCodeSplitter(std::move(MD), Settings.Mode, /*IROutputOnly=*/false, @@ -1324,12 +1327,12 @@ splitSYCLModule(std::unique_ptr M, ModuleSplitterSettings Settings) { size_t ID = 0; std::vector OutputImages; while (Splitter->hasMoreSplits()) { - ModuleDesc MD2 = Splitter->nextSplit(); - MD2.fixupLinkageOfDirectInvokeSimdTargets(); + std::unique_ptr MD2 = Splitter->nextSplit(); + MD2->fixupLinkageOfDirectInvokeSimdTargets(); std::string OutIRFileName = (Settings.OutputPrefix + "_" + Twine(ID)).str(); auto SplittedImageOrErr = - saveModuleDesc(MD2, OutIRFileName, Settings.OutputAssembly); + saveModuleDesc(*MD2, OutIRFileName, Settings.OutputAssembly); if (!SplittedImageOrErr) return SplittedImageOrErr.takeError(); diff --git a/llvm/lib/SYCLPostLink/SpecializationConstants.cpp b/llvm/lib/SYCLPostLink/SpecializationConstants.cpp index d009be542c95a..8a36db0296204 100644 --- a/llvm/lib/SYCLPostLink/SpecializationConstants.cpp +++ b/llvm/lib/SYCLPostLink/SpecializationConstants.cpp @@ -22,11 +22,11 @@ #include using namespace llvm; +using namespace llvm::module_split; namespace { -bool lowerSpecConstants(module_split::ModuleDesc &MD, - SpecConstantsPass::HandlingMode Mode) { +bool lowerSpecConstants(ModuleDesc &MD, SpecConstantsPass::HandlingMode Mode) { ModulePassManager RunSpecConst; ModuleAnalysisManager MAM; SpecConstantsPass SCP(Mode); @@ -45,14 +45,12 @@ bool lowerSpecConstants(module_split::ModuleDesc &MD, /// Specialization constants are replaced by corresponding default values. /// If the Module in \p MD doesn't contain specialization constants then /// std::nullopt is returned. -std::optional -cloneModuleWithSpecConstsReplacedByDefaultValues( - const module_split::ModuleDesc &MD) { - std::optional NewMD; +std::optional> +cloneModuleWithSpecConstsReplacedByDefaultValues(const ModuleDesc &MD) { if (!checkModuleContainsSpecConsts(MD.getModule())) - return NewMD; + return std::nullopt; - NewMD = MD.clone(); + std::unique_ptr NewMD = MD.clone(); NewMD->setSpecConstantDefault(true); ModulePassManager MPM; @@ -68,25 +66,25 @@ cloneModuleWithSpecConstsReplacedByDefaultValues( "SpecConstsMet should be true since the presence of SpecConsts " "has been checked before the run of the pass"); NewMD->rebuildEntryPoints(); - return NewMD; + return std::move(NewMD); } } // namespace bool llvm::sycl::handleSpecializationConstants( - SmallVectorImpl &MDs, + SmallVectorImpl> &MDs, std::optional Mode, - SmallVectorImpl &NewModuleDescs, + SmallVectorImpl> &NewModuleDescs, bool GenerateModuleDescWithDefaultSpecConsts) { bool Modified = false; - for (module_split::ModuleDesc &MD : MDs) { + for (std::unique_ptr &MD : MDs) { if (GenerateModuleDescWithDefaultSpecConsts) - if (std::optional NewMD = - cloneModuleWithSpecConstsReplacedByDefaultValues(MD)) + if (std::optional> NewMD = + cloneModuleWithSpecConstsReplacedByDefaultValues(*MD)) NewModuleDescs.push_back(std::move(*NewMD)); if (Mode) - Modified |= lowerSpecConstants(MD, *Mode); + Modified |= lowerSpecConstants(*MD, *Mode); } return Modified; diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 0b51b42f5680d..1d7d6686e1ad7 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -540,8 +540,9 @@ processInputModule(std::unique_ptr M, const StringRef OutputPrefix) { std::unique_ptr Splitter = module_split::getDeviceCodeSplitter( - module_split::ModuleDesc{std::move(M)}, SplitMode, IROutputOnly, - EmitOnlyKernelsAsEntryPoints, AllowDeviceImageDependencies); + std::make_unique(std::move(M)), SplitMode, + IROutputOnly, EmitOnlyKernelsAsEntryPoints, + AllowDeviceImageDependencies); bool SplitOccurred = Splitter->remainingSplits() > 1; Modified |= SplitOccurred; @@ -562,10 +563,10 @@ processInputModule(std::unique_ptr M, const StringRef OutputPrefix) { // same time, because it leads to a huge RAM consumption by the tool on bigger // inputs. while (Splitter->hasMoreSplits()) { - module_split::ModuleDesc MDesc = Splitter->nextSplit(); - DUMP_ENTRY_POINTS(MDesc.entries(), MDesc.Name.c_str(), 1); + std::unique_ptr MDesc = Splitter->nextSplit(); + DUMP_ENTRY_POINTS(MDesc->entries(), MDesc->Name.c_str(), 1); - MDesc.fixupLinkageOfDirectInvokeSimdTargets(); + MDesc->fixupLinkageOfDirectInvokeSimdTargets(); ESIMDProcessingOptions Options = {SplitMode, EmitOnlyKernelsAsEntryPoints, @@ -577,9 +578,11 @@ processInputModule(std::unique_ptr M, const StringRef OutputPrefix) { auto ModulesOrErr = handleESIMD(std::move(MDesc), Options, Modified, SplitOccurred); CHECK_AND_EXIT(ModulesOrErr.takeError()); - SmallVector &MMs = *ModulesOrErr; + SmallVector, 2> &MMs = + *ModulesOrErr; assert(MMs.size() && "at least one module is expected after ESIMD split"); - SmallVector MMsWithDefaultSpecConsts; + SmallVector, 2> + MMsWithDefaultSpecConsts; Modified |= handleSpecializationConstants(MMs, SCMode, MMsWithDefaultSpecConsts, GenerateDeviceImageWithDefaultSpecConsts); @@ -589,8 +592,8 @@ processInputModule(std::unique_ptr M, const StringRef OutputPrefix) { error("some modules had to be split, '-" + IROutputOnly.ArgStr + "' can't be used"); } - MMs.front().cleanup(AllowDeviceImageDependencies); - saveModuleIR(MMs.front().getModule(), OutputFiles[0].Filename); + MMs.front()->cleanup(AllowDeviceImageDependencies); + saveModuleIR(MMs.front()->getModule(), OutputFiles[0].Filename); return Tables; } // Empty IR file name directs saveModule to generate one and save IR to @@ -603,18 +606,19 @@ processInputModule(std::unique_ptr M, const StringRef OutputPrefix) { errs() << "sycl-post-link NOTE: no modifications to the input LLVM IR " "have been made\n"; } - for (module_split::ModuleDesc &IrMD : MMs) { - IsBF16DeviceLibUsed |= isSYCLDeviceLibBF16Used(IrMD.getModule()); - saveModule(Tables, IrMD, ID, OutputPrefix, OutIRFileName); + for (const std::unique_ptr &IrMD : MMs) { + IsBF16DeviceLibUsed |= isSYCLDeviceLibBF16Used(IrMD->getModule()); + saveModule(Tables, *IrMD, ID, OutputPrefix, OutIRFileName); } ++ID; if (!MMsWithDefaultSpecConsts.empty()) { for (size_t i = 0; i != MMsWithDefaultSpecConsts.size(); ++i) { - module_split::ModuleDesc &IrMD = MMsWithDefaultSpecConsts[i]; - IsBF16DeviceLibUsed |= isSYCLDeviceLibBF16Used(IrMD.getModule()); - saveModule(Tables, IrMD, ID, OutputPrefix, OutIRFileName); + const std::unique_ptr &IrMD = + MMsWithDefaultSpecConsts[i]; + IsBF16DeviceLibUsed |= isSYCLDeviceLibBF16Used(IrMD->getModule()); + saveModule(Tables, *IrMD, ID, OutputPrefix, OutIRFileName); } ++ID; diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index 1bc0b92d84db0..82a38d6ce870a 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -52,6 +52,7 @@ #include #include +#include #include using namespace clang; @@ -795,7 +796,7 @@ jit_compiler::performPostLink(ModuleUPtr Module, } std::unique_ptr Splitter = getDeviceCodeSplitter( - ModuleDesc{std::move(Module)}, SplitMode, + std::make_unique(std::move(Module)), SplitMode, /*IROutputOnly=*/false, EmitOnlyKernelsAsEntryPoints, AllowDeviceImageDependencies); assert(Splitter->hasMoreSplits()); @@ -814,30 +815,30 @@ jit_compiler::performPostLink(ModuleUPtr Module, bool IsBF16DeviceLibUsed = false; while (Splitter->hasMoreSplits()) { - ModuleDesc MDesc = Splitter->nextSplit(); + std::unique_ptr MDesc = Splitter->nextSplit(); // TODO: Call `MDesc.fixupLinkageOfDirectInvokeSimdTargets()` when // `invoke_simd` is supported. - SmallVector ESIMDSplits = + SmallVector, 2> ESIMDSplits = splitByESIMD(std::move(MDesc), EmitOnlyKernelsAsEntryPoints, AllowDeviceImageDependencies); for (auto &ES : ESIMDSplits) { MDesc = std::move(ES); - if (MDesc.isESIMD()) { + if (MDesc->isESIMD()) { // `sycl-post-link` has a `-lower-esimd` option, but there's no clang // driver option to influence it. Rather, the driver sets it // unconditionally in the multi-file output mode, which we are mimicking // here. - lowerEsimdConstructs(MDesc, PerformOpts); + lowerEsimdConstructs(*MDesc, PerformOpts); } - MDesc.saveSplitInformationAsMetadata(); + MDesc->saveSplitInformationAsMetadata(); RTCDevImgInfo &DevImgInfo = DevImgInfoVec.emplace_back(); - DevImgInfo.SymbolTable = FrozenSymbolTable{MDesc.entries().size()}; - transform(MDesc.entries(), DevImgInfo.SymbolTable.begin(), + DevImgInfo.SymbolTable = FrozenSymbolTable{MDesc->entries().size()}; + transform(MDesc->entries(), DevImgInfo.SymbolTable.begin(), [](Function *F) { return F->getName(); }); // TODO: Determine what is requested. @@ -848,7 +849,7 @@ jit_compiler::performPostLink(ModuleUPtr Module, /*EmitImportedSymbols=*/true, /*DeviceGlobals=*/true}; PropertySetRegistry Properties = - computeModuleProperties(MDesc.getModule(), MDesc.entries(), PropReq, + computeModuleProperties(MDesc->getModule(), MDesc->entries(), PropReq, AllowDeviceImageDependencies); // When the split mode is none, the required work group size will be added @@ -866,8 +867,8 @@ jit_compiler::performPostLink(ModuleUPtr Module, encodeProperties(Properties, DevImgInfo); - IsBF16DeviceLibUsed |= isSYCLDeviceLibBF16Used(MDesc.getModule()); - Modules.push_back(MDesc.releaseModulePtr()); + IsBF16DeviceLibUsed |= isSYCLDeviceLibBF16Used(MDesc->getModule()); + Modules.push_back(MDesc->releaseModulePtr()); } }