From 8761d403787f32348377d88e5def1aa37dccae8e Mon Sep 17 00:00:00 2001 From: Maksim Sabianin Date: Mon, 15 Jul 2024 17:36:31 +0200 Subject: [PATCH] [SYCL][ClangLinkerWrapper] Add SYCL Module split library usage in clang-linker-wrapper (#13806) This patch introduces -sycl-module-split-mode command line option in clang-linker-wrapper tool and makes the tool invoke either library or sycl-post-link tool depending on the presense of -sycl-module-split-mode option. Also the patch removes funcionality of reading/writing of Table files in clang-linker-wrapper. Instead of them the tool starts using in-memory structures. --- clang/test/Driver/linker-wrapper-sycl.cpp | 21 +- .../tools/clang-linker-wrapper/CMakeLists.txt | 1 + .../ClangLinkerWrapper.cpp | 288 ++++++++---------- .../clang-linker-wrapper/LinkerWrapperOpts.td | 5 + .../include/llvm/SYCLLowerIR/ModuleSplitter.h | 8 + llvm/lib/SYCLLowerIR/ModuleSplitter.cpp | 69 +++++ 6 files changed, 222 insertions(+), 170 deletions(-) diff --git a/clang/test/Driver/linker-wrapper-sycl.cpp b/clang/test/Driver/linker-wrapper-sycl.cpp index f8f2d14586500..954a616b5953f 100644 --- a/clang/test/Driver/linker-wrapper-sycl.cpp +++ b/clang/test/Driver/linker-wrapper-sycl.cpp @@ -15,7 +15,7 @@ // // RUN: clang-linker-wrapper -sycl-device-libraries=%t.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS %s // CHK-CMDS: "{{.*}}spirv-to-ir-wrapper" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global -// CHK-CMDS-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings +// CHK-CMDS-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-CMDS-NEXT: "{{.*}}llvm-link" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-CMDS-NEXT: "{{.*}}sycl-post-link"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc // CHK-CMDS-NEXT: "{{.*}}llvm-spirv"{{.*}} LLVM_SPIRV_OPTIONS -o {{.*}} @@ -23,6 +23,19 @@ // CHK-CMDS-NEXT: "{{.*}}llc" -filetype=obj -o [[LLCOUT:.*]] [[WRAPPEROUT]].bc // CHK-CMDS-NEXT: "{{.*}}/ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT]] HOST_LIB_PATH HOST_STAT_LIB {{.*}}.o +// Check sycl-module-split-mode command line option. +// This option uses split library instead of sycl-post-link tool. +// RUN: clang-linker-wrapper -sycl-module-split-mode=auto -sycl-device-libraries=%t.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--triple=spir64" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-SPLIT-CMDS %s +// 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: "{{.*}}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 +// CHK-SPLIT-CMDS-NEXT: "{{.*}}llc" -filetype=obj -o [[LLCOUT:.*]] [[WRAPPEROUT]].bc +// CHK-SPLIT-CMDS-NEXT: "{{.*}}/ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT]] HOST_LIB_PATH HOST_STAT_LIB {{.*}}.o + /// check for PIC for device wrap compilation when using -shared // RUN: clang-linker-wrapper -sycl-device-libraries=%t.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--triple=spir64" "--linker-path=/usr/bin/ld" -shared "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-SHARED %s // CHK-SHARED: "{{.*}}llc"{{.*}} -relocation-model=pic @@ -42,7 +55,7 @@ // // RUN: clang-linker-wrapper -sycl-device-libraries=%t1.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS-AOT-GEN %s // CHK-CMDS-AOT-GEN: "{{.*}}spirv-to-ir-wrapper" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global -// CHK-CMDS-AOT-GEN-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings +// CHK-CMDS-AOT-GEN-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-CMDS-AOT-GEN-NEXT: "{{.*}}llvm-link" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-CMDS-AOT-GEN-NEXT: "{{.*}}sycl-post-link"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc // CHK-CMDS-AOT-GEN-NEXT: "{{.*}}llvm-spirv"{{.*}} LLVM_SPIRV_OPTIONS -o {{.*}} @@ -66,7 +79,7 @@ // // RUN: clang-linker-wrapper -sycl-device-libraries=%t2.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t2.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS-AOT-CPU %s // CHK-CMDS-AOT-CPU: "{{.*}}spirv-to-ir-wrapper" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global -// CHK-CMDS-AOT-CPU-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings +// CHK-CMDS-AOT-CPU-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-CMDS-AOT-CPU-NEXT: "{{.*}}llvm-link" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-CMDS-AOT-CPU-NEXT: "{{.*}}sycl-post-link"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc // CHK-CMDS-AOT-CPU-NEXT: "{{.*}}llvm-spirv"{{.*}} LLVM_SPIRV_OPTIONS -o {{.*}} @@ -90,7 +103,7 @@ // // RUN: clang-linker-wrapper -sycl-device-libraries=%t3.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t3.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS-AOT-NV %s // CHK-CMDS-AOT-NV: "{{.*}}spirv-to-ir-wrapper" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global -// CHK-CMDS-AOT-NV-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings +// CHK-CMDS-AOT-NV-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-CMDS-AOT-NV-NEXT: "{{.*}}llvm-link" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-CMDS-AOT-NV-NEXT: "{{.*}}sycl-post-link"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc // CHK-CMDS-AOT-NV-NEXT: "{{.*}}clang"{{.*}} -o [[CLANGOUT:.*]] --target=nvptx64-nvidia-cuda -march={{.*}} diff --git a/clang/tools/clang-linker-wrapper/CMakeLists.txt b/clang/tools/clang-linker-wrapper/CMakeLists.txt index 5556869affaa6..2692160fb429f 100644 --- a/clang/tools/clang-linker-wrapper/CMakeLists.txt +++ b/clang/tools/clang-linker-wrapper/CMakeLists.txt @@ -16,6 +16,7 @@ set(LLVM_LINK_COMPONENTS CodeGen LTO FrontendOffloading + SYCLLowerIR ) set(LLVM_TARGET_DEFINITIONS LinkerWrapperOpts.td) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 57273b022dbae..cddc66ff3074e 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -38,6 +38,7 @@ #include "llvm/Option/ArgList.h" #include "llvm/Option/OptTable.h" #include "llvm/Option/Option.h" +#include "llvm/SYCLLowerIR/ModuleSplitter.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Errc.h" #include "llvm/Support/FileOutputBuffer.h" @@ -99,6 +100,8 @@ static codegen::RegisterCodeGenFlags CodeGenFlags; /// Global flag to indicate that the LTO pipeline threw an error. static std::atomic LTOError; +static std::optional SYCLModuleSplitMode; + using OffloadingImage = OffloadBinary::OffloadingImage; namespace llvm { @@ -602,8 +605,8 @@ 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 runSYCLPostLink(ArrayRef InputFiles, - const ArgList &Args) { +static Expected> +runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args) { Expected SYCLPostLinkPath = findProgram("sycl-post-link", {getMainExecutable("sycl-post-link")}); if (!SYCLPostLinkPath) @@ -630,79 +633,85 @@ static Expected runSYCLPostLink(ArrayRef InputFiles, CmdArgs.push_back(File); if (Error Err = executeCommands(*SYCLPostLinkPath, CmdArgs)) return std::move(Err); - return *TempFileOrErr; -} -// This table is used to manage the output table populated by sycl-post-link. -struct Table { - struct SYCLTableEntry { - std::string IRFile; - std::string PropFile; - std::string SymFile; - }; + if (DryRun) { + // In DryRun we need a dummy entry in order to continue the whole pipeline. + auto ImageFileOrErr = createOutputFile( + sys::path::filename(ExecutableName) + ".sycl.split.image", "bc"); + if (!ImageFileOrErr) + return ImageFileOrErr.takeError(); - SmallVector Entries; + std::vector Modules = {module_split::SplitModule( + *ImageFileOrErr, util::PropertySetRegistry(), "")}; + return Modules; + } - SmallVector getListOfIRFiles(void) { - SmallVector Files; - for (auto &Entry : Entries) { - Files.push_back(Entry.IRFile); - } - return Files; + return llvm::module_split::parseSplitModulesFromFile(*TempFileOrErr); +} + +/// Invokes SYCL Split 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. +/// \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; } - Expected writeSYCLTableToFile(void) { - // Create a new file. - auto TempFileOrErr = - createOutputFile(sys::path::filename(ExecutableName), "table"); - if (!TempFileOrErr) - return TempFileOrErr.takeError(); - std::error_code EC; - raw_fd_ostream TableFile(*TempFileOrErr, EC, sys::fs::OF_None); - if (EC) - reportError(errorCodeToError(EC)); - TableFile << "[Code|Properties|Symbols]\n"; - for (auto &Entry : Entries) { - TableFile << Entry.IRFile << "|"; - TableFile << Entry.PropFile << "|"; - TableFile << Entry.SymFile << "\n"; - } - return *TempFileOrErr; + llvm::module_split::ModuleSplitterSettings Settings; + Settings.Mode = Mode; + Settings.OutputPrefix = ""; + + for (StringRef InputFile : InputFiles) { + 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 &NewSplitModules = *SplitModulesOrErr; + SplitModules.insert(SplitModules.end(), NewSplitModules.begin(), + NewSplitModules.end()); } - Error populateSYCLTable(StringRef EntriesFile) { - llvm::ErrorOr> MBOrErr = - llvm::MemoryBuffer::getFileOrSTDIN(EntriesFile); - if (std::error_code EC = MBOrErr.getError()) - return createFileError(EntriesFile, EC); - int LineNumber = -1; - for (line_iterator LI(**MBOrErr); !LI.is_at_eof(); ++LI) { - // Skip first line - StringRef Line = *LI; - if (LineNumber == -1) { - if (Line != "[Code|Properties|Symbols]") - return createStringError(inconvertibleErrorCode(), - "Invalid SYCL Table file."); - LineNumber++; - continue; - } - if (Line.empty()) - return createStringError(inconvertibleErrorCode(), - "Invalid SYCL Table file."); - auto [FirstWord, Rem1] = Line.split("|"); - SYCLTableEntry Entry; - Entry.IRFile = FirstWord.str(); - if (Rem1.empty()) - return createStringError(inconvertibleErrorCode(), - "Invalid SYCL Table file."); - auto [SecondWord, ThirdWord] = Rem1.split("|"); - Entry.PropFile = SecondWord.str(); - Entry.SymFile = ThirdWord.str(); - Entries.push_back(Entry); + 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; } - return Error::success(); + + errs() << formatv("sycl-module-split: input: {0}, output: {1}\n", + InputFilesStr, SplitOutputFilesStr); } -}; + + return SplitModules; +} /// Add any llvm-spirv option that relies on a specific Triple in addition /// to user supplied options. @@ -924,80 +933,13 @@ static Expected runAOTCompile(StringRef InputFile, "Unsupported SYCL Triple and Arch"); } -Expected readTextFile(StringRef File) { - auto MBOrErr = MemoryBuffer::getFile(File, /*IsText*/ true, - /*RequiresNullTerminator */ true); - if (!MBOrErr) - return createFileError(File, MBOrErr.getError()); - - auto &MB = *MBOrErr; - return std::string(MB->getBufferStart(), MB->getBufferEnd()); -} - -Expected> -readPropertyRegistryFromFile(StringRef File) { - auto MBOrErr = MemoryBuffer::getFile(File, /*IsText*/ true); - if (!MBOrErr) - return createFileError(File, MBOrErr.getError()); - - auto &MB = *MBOrErr; - return util::PropertySetRegistry::read(&*MB); -} - -// The table format is the following: -// [Code|Properties|Symbols] -// a_0.bin|a_0.prop|a_0.sym -// . -// a_n.bin|a_n.prop|a_n.sym -// -// .bin extension might be a bc, spv or other native extension. -Expected> -readSYCLImagesFromTable(StringRef TableFile, const ArgList &Args) { - auto TableOrErr = util::SimpleTable::read(TableFile); - if (!TableOrErr) - return TableOrErr.takeError(); - - std::unique_ptr Table = std::move(*TableOrErr); - int CodeIndex = Table->getColumnId("Code"); - int PropertiesIndex = Table->getColumnId("Properties"); - int SymbolsIndex = Table->getColumnId("Symbols"); - if (CodeIndex == -1 || PropertiesIndex == -1 || SymbolsIndex == -1) - return createStringError( - inconvertibleErrorCode(), - "expected columns in the table: Code, Properties and Symbols"); - - SmallVector Images; - for (const util::SimpleTable::Row &row : Table->rows()) { - auto ImagePath = row.getCell("Code"); - auto ImageOrErr = MemoryBuffer::getFile(ImagePath); - if (!ImageOrErr) - return createFileError(ImagePath, ImageOrErr.getError()); - - auto PropertiesOrErr = - readPropertyRegistryFromFile(row.getCell("Properties")); - if (!PropertiesOrErr) - return PropertiesOrErr.takeError(); - - auto SymbolsOrErr = readTextFile(row.getCell("Symbols")); - if (!SymbolsOrErr) - return SymbolsOrErr.takeError(); - - offloading::SYCLImage Image; - Image.Image = std::move(*ImageOrErr); - Image.PropertyRegistry = std::move(**PropertiesOrErr); - Image.Entries = std::move(*SymbolsOrErr); - Images.push_back(std::move(Image)); - } - - return std::move(Images); -} - /// Reads device images from the given \p InputFile and wraps them /// in one LLVM IR Module as a constant data. /// /// \returns A path to the LLVM Module that contains wrapped images. -Expected wrapSYCLBinariesFromFile(StringRef InputFile, - const ArgList &Args) { +Expected +wrapSYCLBinariesFromFile(std::vector &SplitModules, + const ArgList &Args) { auto OutputFileOrErr = createOutputFile( sys::path::filename(ExecutableName) + ".sycl.image.wrapper", "bc"); if (!OutputFileOrErr) @@ -1005,30 +947,38 @@ Expected wrapSYCLBinariesFromFile(StringRef InputFile, StringRef OutputFilePath = *OutputFileOrErr; if (Verbose || DryRun) { - errs() << formatv(" offload-wrapper: input: {0}, output: {1}\n", InputFile, + std::string InputFiles; + for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { + InputFiles += SplitModules[I].ModuleFilePath; + if (I + 1 < E) + InputFiles += ','; + } + + errs() << formatv(" offload-wrapper: input: {0}, output: {1}\n", InputFiles, OutputFilePath); if (DryRun) return OutputFilePath; } - auto ImagesOrErr = readSYCLImagesFromTable(InputFile, Args); - if (!ImagesOrErr) - return ImagesOrErr.takeError(); - - auto &Images = *ImagesOrErr; StringRef Target = Args.getLastArgValue(OPT_triple_EQ); if (Target.empty()) return createStringError( inconvertibleErrorCode(), "can't wrap SYCL image. -triple argument is missed."); + SmallVector Images; // SYCL runtime currently works for spir64 target triple and not for // spir64-unknown-unknown. // TODO: Fix SYCL runtime to accept both triple llvm::Triple T(Target); StringRef A(T.getArchName()); - for (offloading::SYCLImage &Image : Images) - Image.Target = A; + for (auto &SI : SplitModules) { + auto MBOrDesc = MemoryBuffer::getFile(SI.ModuleFilePath); + if (!MBOrDesc) + return createFileError(SI.ModuleFilePath, MBOrDesc.getError()); + + Images.emplace_back(std::move(*MBOrDesc), SI.Properties, SI.Symbols, A); + } LLVMContext C; Module M("offload.wrapper.object", C); @@ -1060,7 +1010,7 @@ Expected wrapSYCLBinariesFromFile(StringRef InputFile, return E; if (Args.hasArg(OPT_print_wrapped_module)) - errs() << M; + errs() << "Wrapped Module\n" << M; // TODO: Once "llc tool->runCompile" migration is finished we need to remove // this scope and use community flow. @@ -1106,9 +1056,10 @@ static Expected runCompile(StringRef &InputFile, } // Run wrapping library and llc -static Expected runWrapperAndCompile(StringRef &InputFile, - const ArgList &Args) { - auto OutputFile = sycl::wrapSYCLBinariesFromFile(InputFile, Args); +static Expected +runWrapperAndCompile(std::vector &SplitModules, + const ArgList &Args) { + auto OutputFile = sycl::wrapSYCLBinariesFromFile(SplitModules, Args); if (!OutputFile) return OutputFile.takeError(); // call to llc @@ -2099,30 +2050,25 @@ Expected> linkAndWrapDeviceFiles( return TmpOutputOrErr.takeError(); SmallVector InputFilesSYCL; InputFilesSYCL.emplace_back(*TmpOutputOrErr); - auto SYCLPostLinkFile = sycl::runSYCLPostLink(InputFilesSYCL, LinkerArgs); - if (!SYCLPostLinkFile) - return SYCLPostLinkFile.takeError(); - sycl::Table LiveSYCLTable; - if (Error Err = LiveSYCLTable.populateSYCLTable(*SYCLPostLinkFile)) - return std::move(Err); - auto PostLinkedFiles = LiveSYCLTable.getListOfIRFiles(); - if (DryRun) - PostLinkedFiles.push_back("dummy"); - for (unsigned I = 0; I < PostLinkedFiles.size(); ++I) { - SmallVector Files; - Files.emplace_back(PostLinkedFiles[I]); + auto SplitModulesOrErr = + SYCLModuleSplitMode + ? sycl::runSYCLSplitLibrary(InputFilesSYCL, LinkerArgs, + *SYCLModuleSplitMode) + : sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs); + if (!SplitModulesOrErr) + return SplitModulesOrErr.takeError(); + + auto &SplitModules = *SplitModulesOrErr; + for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { + SmallVector Files = {SplitModules[I].ModuleFilePath}; auto LinkedFileFinalOrErr = linkDevice(Files, LinkerArgs, true /* IsSYCLKind */); if (!LinkedFileFinalOrErr) return LinkedFileFinalOrErr.takeError(); - if (!DryRun) - LiveSYCLTable.Entries[I].IRFile = *LinkedFileFinalOrErr; + SplitModules[I].ModuleFilePath = *LinkedFileFinalOrErr; } - auto WrapperInput = LiveSYCLTable.writeSYCLTableToFile(); - if (!WrapperInput) - return WrapperInput.takeError(); // TODO(NOM7): Remove this call and use community flow for bundle/wrap - auto OutputFile = sycl::runWrapperAndCompile(*WrapperInput, LinkerArgs); + auto OutputFile = sycl::runWrapperAndCompile(SplitModules, LinkerArgs); if (!OutputFile) return OutputFile.takeError(); @@ -2616,6 +2562,16 @@ int main(int Argc, char **Argv) { timeTraceProfilerInitialize(Granularity, Argv[0]); } + if (Args.hasArg(OPT_sycl_module_split_mode_EQ)) { + StringRef StrMode = Args.getLastArgValue(OPT_sycl_module_split_mode_EQ); + SYCLModuleSplitMode = module_split::convertStringToSplitMode(StrMode); + if (!SYCLModuleSplitMode) + reportError(createStringError( + inconvertibleErrorCode(), + formatv("sycl-module-split-mode value isn't recognized: {0}", + StrMode))); + } + { llvm::TimeTraceScope TimeScope("Execute linker wrapper"); diff --git a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td index ac5b2a3b3dda1..fe086788c2fa3 100644 --- a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td +++ b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td @@ -169,6 +169,11 @@ def sycl_post_link_options_EQ : Joined<["--", "-"], "sycl-post-link-options=">, Flags<[WrapperOnlyOption]>, HelpText<"Options that will control sycl-post-link step">; +def sycl_module_split_mode_EQ : + Joined<["--", "-"], "sycl-module-split-mode=">, + Flags<[WrapperOnlyOption]>, + HelpText<"Option that turns on split library with the given split mode">; + // Special option to pass in llvm-spirv options def llvm_spirv_options_EQ : Joined<["--", "-"], "llvm-spirv-options=">, Flags<[WrapperOnlyOption]>, diff --git a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h index f4214bb82cecf..5f7e5ba78ed73 100644 --- a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h +++ b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h @@ -22,6 +22,7 @@ #include "llvm/Support/PropertySetIO.h" #include +#include #include #include @@ -45,6 +46,10 @@ enum IRSplitMode { SPLIT_NONE // no splitting }; +// \returns IRSplitMode value if \p S is recognized. Otherwise, std::nullopt is +// returned. +std::optional convertStringToSplitMode(StringRef S); + // A vector that contains all entry point functions in a split module. using EntryPointSet = SetVector; @@ -302,6 +307,9 @@ struct ModuleSplitterSettings { StringRef OutputPrefix; }; +/// Parses the output table file from sycl-post-link tool. +Expected> parseSplitModulesFromFile(StringRef File); + /// Splits the given module \p M according to the given \p Settings. Expected> splitSYCLModule(std::unique_ptr M, ModuleSplitterSettings Settings); diff --git a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp index ade77543cf60d..824f5de06f53e 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/LineIterator.h" #include "llvm/Transforms/IPO.h" #include "llvm/Transforms/IPO/GlobalDCE.h" #include "llvm/Transforms/IPO/Internalize.h" @@ -443,6 +444,19 @@ class ModuleSplitter : public ModuleSplitterBase { namespace llvm { namespace module_split { +std::optional convertStringToSplitMode(StringRef S) { + static const StringMap Values = {{"kernel", SPLIT_PER_KERNEL}, + {"source", SPLIT_PER_TU}, + {"auto", SPLIT_AUTO}, + {"none", SPLIT_NONE}}; + + auto It = Values.find(S); + if (It == Values.end()) + return std::nullopt; + + return It->second; +} + bool isESIMDFunction(const Function &F) { return F.getMetadata(ESIMD_MARKER_MD) != nullptr; } @@ -1215,6 +1229,61 @@ static Expected saveModuleDesc(ModuleDesc &MD, std::string Prefix, return SM; } +Expected> parseSplitModulesFromFile(StringRef File) { + auto EntriesMBOrErr = llvm::MemoryBuffer::getFile(File); + + if (!EntriesMBOrErr) + return createFileError(File, EntriesMBOrErr.getError()); + + line_iterator LI(**EntriesMBOrErr); + if (LI.is_at_eof() || *LI != "[Code|Properties|Symbols]") + return createStringError(inconvertibleErrorCode(), + "invalid SYCL Table file."); + + ++LI; + std::vector Modules; + while (!LI.is_at_eof()) { + StringRef Line = *LI; + if (Line.empty()) + return createStringError(inconvertibleErrorCode(), + "invalid SYCL table row."); + + SmallVector Parts; + Line.split(Parts, "|"); + if (Parts.size() != 3) + return createStringError(inconvertibleErrorCode(), + "invalid SYCL Table row."); + + auto [IRFilePath, PropertyFilePath, SymbolsFilePath] = + std::tie(Parts[0], Parts[1], Parts[2]); + if (PropertyFilePath.empty() || SymbolsFilePath.empty()) + return createStringError(inconvertibleErrorCode(), + "invalid SYCL Table row."); + + auto MBOrErr = MemoryBuffer::getFile(PropertyFilePath); + if (!MBOrErr) + return createFileError(PropertyFilePath, MBOrErr.getError()); + + auto &MB = **MBOrErr; + auto PropSetOrErr = llvm::util::PropertySetRegistry::read(&MB); + if (!PropSetOrErr) + return PropSetOrErr.takeError(); + + llvm::util::PropertySetRegistry Properties = std::move(**PropSetOrErr); + MBOrErr = MemoryBuffer::getFile(SymbolsFilePath); + if (!MBOrErr) + return createFileError(SymbolsFilePath, MBOrErr.getError()); + + auto &MB2 = *MBOrErr; + std::string Symbols = + std::string(MB2->getBufferStart(), MB2->getBufferEnd()); + Modules.emplace_back(IRFilePath, std::move(Properties), std::move(Symbols)); + ++LI; + } + + return Modules; +} + Expected> splitSYCLModule(std::unique_ptr M, ModuleSplitterSettings Settings) { ModuleDesc MD = std::move(M); // makeModuleDesc() ?