diff --git a/clang/test/Driver/linker-wrapper-sycl.cpp b/clang/test/Driver/linker-wrapper-sycl.cpp index 11faa498dd07d..4a93f02fe806b 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 9173ff2ac48bd..372200ac909f4 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -709,7 +709,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")}); @@ -745,44 +745,61 @@ 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. +/// 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 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 split +/// 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. -static Expected> -runSYCLSplitLibrary(ArrayRef InputFiles, const ArgList &Args, - module_split::IRSplitMode Mode) { - std::vector SplitModules; +/// \returns The vector of processed modules. +static Expected> +runSYCLOffloadFinalize(ArrayRef InputFiles, const ArgList &Args, + module_split::IRSplitMode Mode) { + std::vector OutputModules; + 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(); - 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; + OutputModules.emplace_back(*OutputFileOrErr, util::PropertySetRegistry(), + ""); + logSYCLLibraryInvocation(InputFiles, OutputModules, Settings); + return OutputModules; } - llvm::module_split::ModuleSplitterSettings Settings; - Settings.Mode = Mode; - Settings.OutputPrefix = ""; - for (StringRef InputFile : InputFiles) { SMDiagnostic Err; LLVMContext C; @@ -790,31 +807,20 @@ runSYCLSplitLibrary(ArrayRef InputFiles, const ArgList &Args, 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::SYCLOffloadFinalize(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) { - auto InputFilesStr = llvm::join(InputFiles.begin(), InputFiles.end(), ","); - std::string SplitOutputFilesStr; - for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { - if (I > 0) - SplitOutputFilesStr += ','; - - SplitOutputFilesStr += SplitModules[I].ModuleFilePath; - } - - errs() << formatv("sycl-module-split: input: {0}, output: {1}\n", - InputFilesStr, SplitOutputFilesStr); - } + if (Verbose) + logSYCLLibraryInvocation(InputFiles, OutputModules, Settings); - return SplitModules; + return OutputModules; } /// Add any llvm-spirv option that relies on a specific Triple in addition @@ -1065,7 +1071,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"); @@ -1075,8 +1081,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 += ','; } @@ -1104,13 +1110,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; @@ -1193,9 +1200,9 @@ static Expected runCompile(StringRef &InputFile, // Run wrapping library and clang 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 clang @@ -2393,15 +2400,15 @@ Expected> linkAndWrapDeviceFiles( return TmpOutputOrErr.takeError(); SmallVector InputFilesSYCL; InputFilesSYCL.emplace_back(*TmpOutputOrErr); - auto SplitModulesOrErr = + auto ProcessedModulesOrErr = UseSYCLPostLinkTool ? sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs) - : sycl::runSYCLSplitLibrary(InputFilesSYCL, LinkerArgs, - *SYCLModuleSplitMode); - if (!SplitModulesOrErr) - return SplitModulesOrErr.takeError(); + : sycl::runSYCLOffloadFinalize(InputFilesSYCL, LinkerArgs, + *SYCLModuleSplitMode); + 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)) { @@ -2410,14 +2417,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"; @@ -2439,21 +2446,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..cac964a34e683 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 the offload finale processing of the given module \p M according +/// to the given \p Settings. +Expected> +SYCLOffloadFinalize(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 6068ce58f414f..3106ebceed91b 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; } @@ -1287,20 +1302,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) @@ -1312,7 +1329,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()) @@ -1355,27 +1372,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> +SYCLOffloadFinalize(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;