From b8cef90bec3f5dbad4afee9d78a8ae811cedcedf Mon Sep 17 00:00:00 2001 From: "Sabianin, Maksim" Date: Wed, 11 Sep 2024 09:32:19 -0700 Subject: [PATCH 1/3] [SYCL] Refactor SYCL Post Link Library Post Link compilation flow consists of Module spliting, specializaition constants processing and ESIMD processing. To correspond to the all functionalities this patch adds renames in order to be less confusing. The following changes are added: * ModuleSplitProcessingSettings structure is renamed to ModuleProcessingSettings. * Added function convertProcessingSettingsToString to add additional piece of information in the testing. * SplitModule structure is renamed to ProcessedModule. * runSYCLSplitLibrary is renamed to runSYCLPostLinkLibrary. * DryRun mode in runSYCLPostLinkLibrary is combined with Verbose mode in order to remove code duplication. * splitSYCLModule is renamed to SYCLPostLinkProcess. --- clang/test/Driver/linker-wrapper-sycl.cpp | 2 +- .../ClangLinkerWrapper.cpp | 124 +++++++++--------- .../include/llvm/SYCLLowerIR/ModuleSplitter.h | 36 +++-- llvm/lib/SYCLLowerIR/ModuleSplitter.cpp | 58 +++++--- .../sycl-module-split/sycl-module-split.cpp | 8 +- 5 files changed, 132 insertions(+), 96 deletions(-) diff --git a/clang/test/Driver/linker-wrapper-sycl.cpp b/clang/test/Driver/linker-wrapper-sycl.cpp index 18cf3c91ad5b7..ef3cfd91513e6 100644 --- a/clang/test/Driver/linker-wrapper-sycl.cpp +++ b/clang/test/Driver/linker-wrapper-sycl.cpp @@ -29,7 +29,7 @@ // CHK-SPLIT-CMDS: "{{.*}}spirv-to-ir-wrapper" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global // CHK-SPLIT-CMDS-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-SPLIT-CMDS-NEXT: "{{.*}}llvm-link" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings -// CHK-SPLIT-CMDS-NEXT: sycl-module-split: input: [[SECONDLLVMLINKOUT]].bc, output: [[SYCLMODULESPLITOUT:.*]].bc +// CHK-SPLIT-CMDS-NEXT: sycl-post-link-library: input: [[SECONDLLVMLINKOUT]].bc, output: [[SYCLMODULESPLITOUT:.*]].bc, split_mode: auto, output_assembly: false, output_prefix: // CHK-SPLIT-CMDS-NEXT: "{{.*}}llvm-spirv"{{.*}} LLVM_SPIRV_OPTIONS -o [[SPIRVOUT:.*]].spv [[SYCLMODULESPLITOUT]].bc // LLVM-SPIRV is not called in dry-run // CHK-SPLIT-CMDS-NEXT: offload-wrapper: input: [[SPIRVOUT]].spv, output: [[WRAPPEROUT:.*]].bc diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index eb37fa583d63a..d6ccb3ee77fb2 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -707,7 +707,7 @@ getTripleBasedSYCLPostLinkOpts(const ArgList &Args, /// 'Args' encompasses all arguments required for linking and wrapping device /// code and will be parsed to generate options required to be passed into the /// sycl-post-link tool. -static Expected> +static Expected> runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args) { Expected SYCLPostLinkPath = findProgram("sycl-post-link", {getMainExecutable("sycl-post-link")}); @@ -743,76 +743,76 @@ runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args) { if (!ImageFileOrErr) return ImageFileOrErr.takeError(); - std::vector Modules = {module_split::SplitModule( + std::vector Modules = {module_split::ProcessedModule( *ImageFileOrErr, util::PropertySetRegistry(), "")}; return Modules; } - return llvm::module_split::parseSplitModulesFromFile(*TempFileOrErr); + return llvm::module_split::parseProcessedModulesFromFile(*TempFileOrErr); } -/// Invokes SYCL Split library for SYCL offloading. +/// Invokes SYCL Post Link library for SYCL offloading. /// /// \param InputFiles the list of input LLVM IR files. /// \param Args Encompasses all arguments for linking and wrapping device code. -/// It will be parsed to generate options required to be passed to SYCL split -/// library. +/// It will be parsed to generate options required to be passed to SYCL Post +/// Link library. /// \param Mode The splitting mode. /// \returns The vector of split modules. -static Expected> -runSYCLSplitLibrary(ArrayRef InputFiles, const ArgList &Args, - module_split::IRSplitMode Mode) { - std::vector SplitModules; - if (DryRun) { - auto OutputFileOrErr = createOutputFile( - sys::path::filename(ExecutableName) + ".sycl.split.image", "bc"); - if (!OutputFileOrErr) - return OutputFileOrErr.takeError(); - - StringRef OutputFilePath = *OutputFileOrErr; - auto InputFilesStr = llvm::join(InputFiles.begin(), InputFiles.end(), ","); - errs() << formatv("sycl-module-split: input: {0}, output: {1}\n", - InputFilesStr, OutputFilePath); - SplitModules.emplace_back(OutputFilePath, util::PropertySetRegistry(), ""); - return SplitModules; - } - - llvm::module_split::ModuleSplitterSettings Settings; +static Expected> +runSYCLPostLinkLibrary(ArrayRef InputFiles, const ArgList &Args, + module_split::IRSplitMode Mode) { + std::vector OutputModules; + llvm::module_split::ModuleProcessingSettings Settings; Settings.Mode = Mode; Settings.OutputPrefix = ""; for (StringRef InputFile : InputFiles) { + if (DryRun) + break; + SMDiagnostic Err; LLVMContext C; std::unique_ptr M = parseIRFile(InputFile, Err, C); if (!M) return createStringError(inconvertibleErrorCode(), Err.getMessage()); - auto SplitModulesOrErr = - module_split::splitSYCLModule(std::move(M), Settings); - if (!SplitModulesOrErr) - return SplitModulesOrErr.takeError(); + auto ModulesOrErr = + module_split::SYCLPostLinkProcess(std::move(M), Settings); + if (!ModulesOrErr) + return ModulesOrErr.takeError(); - auto &NewSplitModules = *SplitModulesOrErr; - SplitModules.insert(SplitModules.end(), NewSplitModules.begin(), - NewSplitModules.end()); + auto &NewModules = *ModulesOrErr; + OutputModules.insert(OutputModules.end(), NewModules.begin(), + NewModules.end()); } - if (Verbose) { + if (Verbose || DryRun) { + if (DryRun) { + auto OutputFileOrErr = createOutputFile( + sys::path::filename(ExecutableName) + ".sycl.split.image", "bc"); + if (!OutputFileOrErr) + return OutputFileOrErr.takeError(); + + OutputModules.emplace_back(*OutputFileOrErr, util::PropertySetRegistry(), + ""); + } + auto InputFilesStr = llvm::join(InputFiles.begin(), InputFiles.end(), ","); std::string SplitOutputFilesStr; - for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { + for (size_t I = 0, E = OutputModules.size(); I != E; ++I) { if (I > 0) SplitOutputFilesStr += ','; - SplitOutputFilesStr += SplitModules[I].ModuleFilePath; + SplitOutputFilesStr += OutputModules[I].ModuleFilePath; } - errs() << formatv("sycl-module-split: input: {0}, output: {1}\n", - InputFilesStr, SplitOutputFilesStr); + errs() << formatv("sycl-post-link-library: input: {0}, output: {1}, {2}\n", + InputFilesStr, SplitOutputFilesStr, + convertProcessingSettingsToString(Settings)); } - return SplitModules; + return OutputModules; } /// Add any llvm-spirv option that relies on a specific Triple in addition @@ -1063,7 +1063,7 @@ static Expected runAOTCompile(StringRef InputFile, /// /// \returns A path to the LLVM Module that contains wrapped images. Expected -wrapSYCLBinariesFromFile(std::vector &SplitModules, +wrapSYCLBinariesFromFile(std::vector &Modules, const ArgList &Args, bool IsEmbeddedIR) { auto OutputFileOrErr = createOutputFile( sys::path::filename(ExecutableName) + ".sycl.image.wrapper", "bc"); @@ -1073,8 +1073,8 @@ wrapSYCLBinariesFromFile(std::vector &SplitModules, StringRef OutputFilePath = *OutputFileOrErr; if (Verbose || DryRun) { std::string InputFiles; - for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { - InputFiles += SplitModules[I].ModuleFilePath; + for (size_t I = 0, E = Modules.size(); I != E; ++I) { + InputFiles += Modules[I].ModuleFilePath; if (I + 1 < E) InputFiles += ','; } @@ -1102,13 +1102,14 @@ wrapSYCLBinariesFromFile(std::vector &SplitModules, if (RegularTarget == "spirv64") RegularTarget = "spir64"; - for (auto &SI : SplitModules) { - auto MBOrDesc = MemoryBuffer::getFile(SI.ModuleFilePath); + for (auto &M : Modules) { + auto MBOrDesc = MemoryBuffer::getFile(M.ModuleFilePath); if (!MBOrDesc) - return createFileError(SI.ModuleFilePath, MBOrDesc.getError()); + return createFileError(M.ModuleFilePath, MBOrDesc.getError()); StringRef ImageTarget = IsEmbeddedIR ? StringRef(EmbeddedIRTarget) : StringRef(RegularTarget); - Images.emplace_back(std::move(*MBOrDesc), SI.Properties, SI.Symbols, ImageTarget); + Images.emplace_back(std::move(*MBOrDesc), M.Properties, M.Symbols, + ImageTarget); } LLVMContext C; @@ -1183,9 +1184,9 @@ static Expected runCompile(StringRef &InputFile, // Run wrapping library and llc static Expected -runWrapperAndCompile(std::vector &SplitModules, +runWrapperAndCompile(std::vector &Modules, const ArgList &Args, bool IsEmbeddedIR = false) { - auto OutputFile = sycl::wrapSYCLBinariesFromFile(SplitModules, Args, IsEmbeddedIR); + auto OutputFile = sycl::wrapSYCLBinariesFromFile(Modules, Args, IsEmbeddedIR); if (!OutputFile) return OutputFile.takeError(); // call to llc @@ -2252,15 +2253,15 @@ Expected> linkAndWrapDeviceFiles( return TmpOutputOrErr.takeError(); SmallVector InputFilesSYCL; InputFilesSYCL.emplace_back(*TmpOutputOrErr); - auto SplitModulesOrErr = + auto ProcessedModulesOrErr = SYCLModuleSplitMode - ? sycl::runSYCLSplitLibrary(InputFilesSYCL, LinkerArgs, - *SYCLModuleSplitMode) + ? sycl::runSYCLPostLinkLibrary(InputFilesSYCL, LinkerArgs, + *SYCLModuleSplitMode) : sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs); - if (!SplitModulesOrErr) - return SplitModulesOrErr.takeError(); + if (!ProcessedModulesOrErr) + return ProcessedModulesOrErr.takeError(); - auto &SplitModules = *SplitModulesOrErr; + auto &ProcessedModules = *ProcessedModulesOrErr; const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); if ((Triple.isNVPTX() || Triple.isAMDGCN()) && LinkerArgs.hasArg(OPT_sycl_embed_ir)) { @@ -2269,14 +2270,14 @@ Expected> linkAndWrapDeviceFiles( // of sycl-post-link (filetable referencing LLVM Bitcode + symbols) // through the offload wrapper and link the resulting object to the // application. - auto OutputFile = - sycl::runWrapperAndCompile(SplitModules, LinkerArgs, /* IsEmbeddedIR */ true); + auto OutputFile = sycl::runWrapperAndCompile( + ProcessedModules, LinkerArgs, /* IsEmbeddedIR */ true); if (!OutputFile) return OutputFile.takeError(); WrappedOutput.push_back(*OutputFile); } - for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { - SmallVector Files = {SplitModules[I].ModuleFilePath}; + for (size_t I = 0, E = ProcessedModules.size(); I != E; ++I) { + SmallVector Files = {ProcessedModules[I].ModuleFilePath}; StringRef Arch = LinkerArgs.getLastArgValue(OPT_arch_EQ); if (Arch.empty()) Arch = "native"; @@ -2298,21 +2299,22 @@ Expected> linkAndWrapDeviceFiles( nvptx::fatbinary(BundlerInputFiles, LinkerArgs); if (!BundledFileOrErr) return BundledFileOrErr.takeError(); - SplitModules[I].ModuleFilePath = *BundledFileOrErr; + ProcessedModules[I].ModuleFilePath = *BundledFileOrErr; } else if (Triple.isAMDGCN()) { BundlerInputFiles.emplace_back(*ClangOutputOrErr, Arch); auto BundledFileOrErr = amdgcn::fatbinary(BundlerInputFiles, LinkerArgs); if (!BundledFileOrErr) return BundledFileOrErr.takeError(); - SplitModules[I].ModuleFilePath = *BundledFileOrErr; + ProcessedModules[I].ModuleFilePath = *BundledFileOrErr; } else { - SplitModules[I].ModuleFilePath = *ClangOutputOrErr; + ProcessedModules[I].ModuleFilePath = *ClangOutputOrErr; } } // TODO(NOM7): Remove this call and use community flow for bundle/wrap - auto OutputFile = sycl::runWrapperAndCompile(SplitModules, LinkerArgs); + auto OutputFile = + sycl::runWrapperAndCompile(ProcessedModules, LinkerArgs); if (!OutputFile) return OutputFile.takeError(); diff --git a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h index 0da3706ad3626..8ce2a7d7f50f5 100644 --- a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h +++ b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h @@ -50,6 +50,8 @@ enum IRSplitMode { // returned. std::optional convertStringToSplitMode(StringRef S); +StringRef convertSplitModeToString(IRSplitMode SM); + // A vector that contains all entry point functions in a split module. using EntryPointSet = SetVector; @@ -289,35 +291,41 @@ void dumpEntryPoints(const Module &M, bool OnlyKernelsAreEntryPoints = false, const char *Msg = "", int Tab = 0); #endif // NDEBUG -struct SplitModule { +struct ProcessedModule { std::string ModuleFilePath; util::PropertySetRegistry Properties; std::string Symbols; - SplitModule() = default; - SplitModule(const SplitModule &) = default; - SplitModule &operator=(const SplitModule &) = default; - SplitModule(SplitModule &&) = default; - SplitModule &operator=(SplitModule &&) = default; + ProcessedModule() = default; + ProcessedModule(const ProcessedModule &) = default; + ProcessedModule &operator=(const ProcessedModule &) = default; + ProcessedModule(ProcessedModule &&) = default; + ProcessedModule &operator=(ProcessedModule &&) = default; - SplitModule(std::string_view File, util::PropertySetRegistry Properties, - std::string Symbols) + ProcessedModule(std::string_view File, util::PropertySetRegistry Properties, + std::string Symbols) : ModuleFilePath(File), Properties(std::move(Properties)), Symbols(std::move(Symbols)) {} }; -struct ModuleSplitterSettings { +struct ModuleProcessingSettings { IRSplitMode Mode; bool OutputAssembly = false; // Bitcode or LLVM IR. StringRef OutputPrefix; }; -/// Parses the output table file from sycl-post-link tool. -Expected> parseSplitModulesFromFile(StringRef File); +SmallString<64> +convertProcessingSettingsToString(const ModuleProcessingSettings &S); -/// Splits the given module \p M according to the given \p Settings. -Expected> -splitSYCLModule(std::unique_ptr M, ModuleSplitterSettings Settings); +/// Parses the output table file from sycl-post-link tool. +Expected> +parseProcessedModulesFromFile(StringRef File); + +/// Performs post-link processing of the given module \p M according to the +/// given \p Settings. +Expected> +SYCLPostLinkProcess(std::unique_ptr M, + ModuleProcessingSettings Settings); bool isESIMDFunction(const Function &F); bool canBeImportedFunction(const Function &F); diff --git a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp index b9eda1376663f..8a88b75a63dd1 100644 --- a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp +++ b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp @@ -29,6 +29,7 @@ #include "llvm/Support/CommandLine.h" #include "llvm/Support/Error.h" #include "llvm/Support/FileSystem.h" +#include "llvm/Support/FormatVariadic.h" #include "llvm/Support/LineIterator.h" #include "llvm/Transforms/IPO.h" #include "llvm/Transforms/IPO/GlobalDCE.h" @@ -458,6 +459,20 @@ std::optional convertStringToSplitMode(StringRef S) { return It->second; } +StringRef convertSplitModeToString(IRSplitMode SM) { + static const DenseMap Values = { + {SPLIT_PER_KERNEL, "kernel"}, + {SPLIT_PER_TU, "source"}, + {SPLIT_AUTO, "auto"}, + {SPLIT_NONE, "none"}}; + + auto It = Values.find(SM); + if (It == Values.end()) + llvm_unreachable("SplitMode value is unhandled!"); + + return It->second; +} + bool isESIMDFunction(const Function &F) { return F.getMetadata(ESIMD_MARKER_MD) != nullptr; } @@ -1285,20 +1300,22 @@ static Error saveModuleIRInFile(Module &M, StringRef FilePath, return Error::success(); } -static Expected saveModuleDesc(ModuleDesc &MD, std::string Prefix, - bool OutputAssembly) { - SplitModule SM; +static Expected +saveModuleDesc(ModuleDesc &MD, std::string Prefix, bool OutputAssembly) { + ProcessedModule PM; Prefix += OutputAssembly ? ".ll" : ".bc"; Error E = saveModuleIRInFile(MD.getModule(), Prefix, OutputAssembly); if (E) return E; - SM.ModuleFilePath = Prefix; - SM.Symbols = MD.makeSymbolTable(); - return SM; + PM.ModuleFilePath = Prefix; + PM.Symbols = MD.makeSymbolTable(); + // TODO: add properties generation. + return PM; } -Expected> parseSplitModulesFromFile(StringRef File) { +Expected> +parseProcessedModulesFromFile(StringRef File) { auto EntriesMBOrErr = llvm::MemoryBuffer::getFile(File); if (!EntriesMBOrErr) @@ -1310,7 +1327,7 @@ Expected> parseSplitModulesFromFile(StringRef File) { "invalid SYCL Table file."); ++LI; - std::vector Modules; + std::vector Modules; while (!LI.is_at_eof()) { StringRef Line = *LI; if (Line.empty()) @@ -1353,27 +1370,36 @@ Expected> parseSplitModulesFromFile(StringRef File) { return Modules; } -Expected> -splitSYCLModule(std::unique_ptr M, ModuleSplitterSettings Settings) { - ModuleDesc MD = std::move(M); // makeModuleDesc() ? +SmallString<64> +convertProcessingSettingsToString(const ModuleProcessingSettings &S) { + return formatv("split_mode: {0}, output_assembly: {1}, output_prefix: {2}", + convertSplitModeToString(S.Mode), S.OutputAssembly, + S.OutputPrefix) + .sstr<64>(); +} + +Expected> +SYCLPostLinkProcess(std::unique_ptr M, + ModuleProcessingSettings Settings) { + ModuleDesc MD = std::move(M); // FIXME: false arguments are temporary for now. auto Splitter = getDeviceCodeSplitter(std::move(MD), Settings.Mode, /*IROutputOnly=*/false, /*EmitOnlyKernelsAsEntryPoints=*/false); size_t ID = 0; - std::vector OutputImages; + std::vector OutputImages; while (Splitter->hasMoreSplits()) { ModuleDesc MD2 = Splitter->nextSplit(); MD2.fixupLinkageOfDirectInvokeSimdTargets(); std::string OutIRFileName = (Settings.OutputPrefix + "_" + Twine(ID)).str(); - auto SplittedImageOrErr = + auto ImageOrErr = saveModuleDesc(MD2, OutIRFileName, Settings.OutputAssembly); - if (!SplittedImageOrErr) - return SplittedImageOrErr.takeError(); + if (!ImageOrErr) + return ImageOrErr.takeError(); - OutputImages.emplace_back(std::move(*SplittedImageOrErr)); + OutputImages.emplace_back(std::move(*ImageOrErr)); ++ID; } diff --git a/llvm/tools/sycl-module-split/sycl-module-split.cpp b/llvm/tools/sycl-module-split/sycl-module-split.cpp index 89d8b9e10b2b7..4aaf71702ac6f 100644 --- a/llvm/tools/sycl-module-split/sycl-module-split.cpp +++ b/llvm/tools/sycl-module-split/sycl-module-split.cpp @@ -75,7 +75,7 @@ void writePropertiesToFile(const PropertySetRegistry &Properties, Properties.write(OS); } -void dumpModulesAsTable(const std::vector &SplitModules, +void dumpModulesAsTable(const std::vector &Modules, StringRef Path) { std::vector Columns = {"Code", "Properties", "Symbols"}; auto TableOrErr = SimpleTable::create(Columns); @@ -85,7 +85,7 @@ void dumpModulesAsTable(const std::vector &SplitModules, } std::unique_ptr Table = std::move(*TableOrErr); - for (const auto &[I, SM] : enumerate(SplitModules)) { + for (const auto &[I, SM] : enumerate(Modules)) { std::string SymbolsFile = (Twine(Path) + "_" + Twine(I) + ".sym").str(); std::string PropertiesFile = (Twine(Path) + "_" + Twine(I) + ".prop").str(); writePropertiesToFile(SM.Properties, PropertiesFile); @@ -116,11 +116,11 @@ int main(int argc, char *argv[]) { return 1; } - ModuleSplitterSettings Settings; + ModuleProcessingSettings Settings; Settings.Mode = SplitMode; Settings.OutputAssembly = OutputAssembly; Settings.OutputPrefix = OutputFilenamePrefix; - auto SplitModulesOrErr = splitSYCLModule(std::move(M), Settings); + auto SplitModulesOrErr = SYCLPostLinkProcess(std::move(M), Settings); if (!SplitModulesOrErr) { Err.print(argv[0], errs()); return 1; From 2addab06b04f7590daf070c98c24e26f44c0c148 Mon Sep 17 00:00:00 2001 From: "Sabianin, Maksim" Date: Thu, 12 Sep 2024 08:00:09 -0700 Subject: [PATCH 2/3] fix code duplication --- .../ClangLinkerWrapper.cpp | 60 ++++++++++--------- 1 file changed, 33 insertions(+), 27 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index d6ccb3ee77fb2..7bd6612257ee2 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -751,6 +751,26 @@ runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args) { return llvm::module_split::parseProcessedModulesFromFile(*TempFileOrErr); } +/// Prints the message for DryRun and Verbose modes. The message contains of +/// input, output and settings. +void logSYCLLibraryInvocation( + ArrayRef InputFiles, + const std::vector &Modules, + const module_split::ModuleProcessingSettings &Settings) { + auto InputFilesStr = llvm::join(InputFiles.begin(), InputFiles.end(), ","); + SmallString<128> SplitOutputFilesStr; + for (size_t I = 0, E = Modules.size(); I != E; ++I) { + if (I > 0) + SplitOutputFilesStr += ','; + + SplitOutputFilesStr += Modules[I].ModuleFilePath; + } + + errs() << formatv("sycl-post-link-library: input: {0}, output: {1}, {2}\n", + InputFilesStr, SplitOutputFilesStr, + convertProcessingSettingsToString(Settings)); +} + /// Invokes SYCL Post Link library for SYCL offloading. /// /// \param InputFiles the list of input LLVM IR files. @@ -766,11 +786,19 @@ runSYCLPostLinkLibrary(ArrayRef InputFiles, const ArgList &Args, llvm::module_split::ModuleProcessingSettings Settings; Settings.Mode = Mode; Settings.OutputPrefix = ""; + if (DryRun) { + auto OutputFileOrErr = createOutputFile( + sys::path::filename(ExecutableName) + ".sycl.split.image", "bc"); + if (!OutputFileOrErr) + return OutputFileOrErr.takeError(); - for (StringRef InputFile : InputFiles) { - if (DryRun) - break; + OutputModules.emplace_back(*OutputFileOrErr, util::PropertySetRegistry(), + ""); + logSYCLLibraryInvocation(InputFiles, OutputModules, Settings); + return OutputModules; + } + for (StringRef InputFile : InputFiles) { SMDiagnostic Err; LLVMContext C; std::unique_ptr M = parseIRFile(InputFile, Err, C); @@ -787,30 +815,8 @@ runSYCLPostLinkLibrary(ArrayRef InputFiles, const ArgList &Args, NewModules.end()); } - if (Verbose || DryRun) { - if (DryRun) { - auto OutputFileOrErr = createOutputFile( - sys::path::filename(ExecutableName) + ".sycl.split.image", "bc"); - if (!OutputFileOrErr) - return OutputFileOrErr.takeError(); - - OutputModules.emplace_back(*OutputFileOrErr, util::PropertySetRegistry(), - ""); - } - - auto InputFilesStr = llvm::join(InputFiles.begin(), InputFiles.end(), ","); - std::string SplitOutputFilesStr; - for (size_t I = 0, E = OutputModules.size(); I != E; ++I) { - if (I > 0) - SplitOutputFilesStr += ','; - - SplitOutputFilesStr += OutputModules[I].ModuleFilePath; - } - - errs() << formatv("sycl-post-link-library: input: {0}, output: {1}, {2}\n", - InputFilesStr, SplitOutputFilesStr, - convertProcessingSettingsToString(Settings)); - } + if (Verbose) + logSYCLLibraryInvocation(InputFiles, OutputModules, Settings); return OutputModules; } From 2efb76ee86394639678ff1986fb0301fe35a907a Mon Sep 17 00:00:00 2001 From: "Sabianin, Maksim" Date: Wed, 2 Oct 2024 06:30:28 -0700 Subject: [PATCH 3/3] address PR feedback --- .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 16 ++++++++-------- llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h | 6 +++--- llvm/lib/SYCLLowerIR/ModuleSplitter.cpp | 2 +- 3 files changed, 12 insertions(+), 12 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 16ff2c8a8cc41..372200ac909f4 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -773,16 +773,16 @@ void logSYCLLibraryInvocation( convertProcessingSettingsToString(Settings)); } -/// Invokes SYCL Post Link library for SYCL offloading. +/// Invokes SYCL processing library for SYCL offload finalization. /// /// \param InputFiles the list of input LLVM IR files. /// \param Args Encompasses all arguments for linking and wrapping device code. -/// It will be parsed to generate options required to be passed to SYCL Post -/// Link library. +/// It will be parsed to generate options required to be passed to SYCL +/// library. /// \param Mode The splitting mode. -/// \returns The vector of split modules. +/// \returns The vector of processed modules. static Expected> -runSYCLPostLinkLibrary(ArrayRef InputFiles, const ArgList &Args, +runSYCLOffloadFinalize(ArrayRef InputFiles, const ArgList &Args, module_split::IRSplitMode Mode) { std::vector OutputModules; llvm::module_split::ModuleProcessingSettings Settings; @@ -808,7 +808,7 @@ runSYCLPostLinkLibrary(ArrayRef InputFiles, const ArgList &Args, return createStringError(inconvertibleErrorCode(), Err.getMessage()); auto ModulesOrErr = - module_split::SYCLPostLinkProcess(std::move(M), Settings); + module_split::SYCLOffloadFinalize(std::move(M), Settings); if (!ModulesOrErr) return ModulesOrErr.takeError(); @@ -2403,8 +2403,8 @@ Expected> linkAndWrapDeviceFiles( auto ProcessedModulesOrErr = UseSYCLPostLinkTool ? sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs) - : sycl::runSYCLSplitLibrary(InputFilesSYCL, LinkerArgs, - *SYCLModuleSplitMode); + : sycl::runSYCLOffloadFinalize(InputFilesSYCL, LinkerArgs, + *SYCLModuleSplitMode); if (!ProcessedModulesOrErr) return ProcessedModulesOrErr.takeError(); diff --git a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h index 8ce2a7d7f50f5..cac964a34e683 100644 --- a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h +++ b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h @@ -321,10 +321,10 @@ convertProcessingSettingsToString(const ModuleProcessingSettings &S); Expected> parseProcessedModulesFromFile(StringRef File); -/// Performs post-link processing of the given module \p M according to the -/// given \p Settings. +/// Performs the offload finale processing of the given module \p M according +/// to the given \p Settings. Expected> -SYCLPostLinkProcess(std::unique_ptr M, +SYCLOffloadFinalize(std::unique_ptr M, ModuleProcessingSettings Settings); bool isESIMDFunction(const Function &F); diff --git a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp index e8aae3214648a..3106ebceed91b 100644 --- a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp +++ b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp @@ -1381,7 +1381,7 @@ convertProcessingSettingsToString(const ModuleProcessingSettings &S) { } Expected> -SYCLPostLinkProcess(std::unique_ptr M, +SYCLOffloadFinalize(std::unique_ptr M, ModuleProcessingSettings Settings) { ModuleDesc MD = std::move(M); // FIXME: false arguments are temporary for now.