From 783d2b93d7924879489943e3a8fa7d0c05f97ea5 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Thu, 12 Sep 2024 01:49:56 +0900 Subject: [PATCH] [SYCL][ClangLinkerWrapper] Support old-style objects and static archives (#15216) This PR finishes up work our intern Jason was working on [here](https://github.com/intel/llvm/pull/15156). Most of the code here is not new, it is old code that was removed [here](https://github.com/intel/llvm/commit/ece73ad61b49eaf9ecb6e2060e5f20e09e26def6). This code is not intended to be permanent or upstreamed. It's intended to be temporary to ease the work to enabling the new offload model by default. Both object files and static archives are supported and tested, but SPIR-V fat objects are not, I don't think any customers are using that anyway. Signed-off-by: Sarnie, Nick Co-authored-by: Li, Jason --- .../ClangLinkerWrapper.cpp | 139 +++++++++++++++++- .../test-e2e/NewOffloadDriver/multisource.cpp | 17 +++ 2 files changed, 154 insertions(+), 2 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index eb37fa583d63a..65e25fa1e034e 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -251,8 +251,8 @@ Expected getInputBitcodeLibrary(StringRef Input) { Image.StringData["arch"] = Arch; Image.Image = std::move(*ImageOrError); - std::unique_ptr Binary = - MemoryBuffer::getMemBufferCopy(OffloadBinary::write(Image)); + std::unique_ptr Binary = MemoryBuffer::getMemBufferCopy( + OffloadBinary::write(Image), Image.Image->getBufferIdentifier()); auto NewBinaryOrErr = OffloadBinary::create(*Binary); if (!NewBinaryOrErr) return NewBinaryOrErr.takeError(); @@ -1358,6 +1358,135 @@ static Expected linkDevice(ArrayRef InputFiles, return *DeviceLinkedFile; } +static bool isStaticArchiveFile(const StringRef Filename) { + if (!llvm::sys::path::has_extension(Filename)) + // Any file with no extension should not be considered an Archive. + return false; + llvm::file_magic Magic; + llvm::identify_magic(Filename, Magic); + // Only archive files are to be considered. + // TODO: .lib check to be added + return (Magic == llvm::file_magic::archive); +} + +static Expected listSection(StringRef Filename, + const ArgList &Args) { + Expected OffloadBundlerPath = findProgram( + "clang-offload-bundler", {getMainExecutable("clang-offload-bundler")}); + if (!OffloadBundlerPath) + return OffloadBundlerPath.takeError(); + BumpPtrAllocator Alloc; + StringSaver Saver(Alloc); + + SmallVector CmdArgs; + CmdArgs.push_back(*OffloadBundlerPath); + bool IsArchive = isStaticArchiveFile(Filename); + CmdArgs.push_back(IsArchive ? "-type=aoo" : "-type=o"); + CmdArgs.push_back(Saver.save("-input=" + Filename)); + CmdArgs.push_back("-list"); + auto Output = createOutputFile("bundled-targets", "list"); + if (!Output) + return Output.takeError(); + SmallVector> Redirects{std::nullopt, *Output, + std::nullopt}; + int ErrCode = llvm::sys::ExecuteAndWait(*OffloadBundlerPath, CmdArgs, + std::nullopt, Redirects); + if (ErrCode != 0) + return createStringError(inconvertibleErrorCode(), + "Failed to list targets"); + return *Output; +} + +// This routine is used to run the clang-offload-bundler tool and unbundle +// device inputs that have been created with an older compiler where the +// device object is bundled into a host object. +static Expected unbundle(StringRef Filename, const ArgList &Args, + llvm::Triple Triple) { + Expected OffloadBundlerPath = findProgram( + "clang-offload-bundler", {getMainExecutable("clang-offload-bundler")}); + if (!OffloadBundlerPath) + return OffloadBundlerPath.takeError(); + + // Create a new file to write the unbundled file to. + auto TempFileOrErr = + createOutputFile(sys::path::filename(ExecutableName), "ir"); + if (!TempFileOrErr) + return TempFileOrErr.takeError(); + + BumpPtrAllocator Alloc; + StringSaver Saver(Alloc); + + SmallVector CmdArgs; + CmdArgs.push_back(*OffloadBundlerPath); + bool IsArchive = isStaticArchiveFile(Filename); + CmdArgs.push_back(IsArchive ? "-type=aoo" : "-type=o"); + auto *Target = Args.MakeArgString(Twine("-targets=sycl-") + Triple.str()); + CmdArgs.push_back(Target); + CmdArgs.push_back(Saver.save("-input=" + Filename)); + CmdArgs.push_back(Saver.save("-output=" + *TempFileOrErr)); + CmdArgs.push_back("-unbundle"); + CmdArgs.push_back("-allow-missing-bundles"); + if (Error Err = executeCommands(*OffloadBundlerPath, CmdArgs)) + return std::move(Err); + return *TempFileOrErr; +} + +Error extractBundledObjects(StringRef Filename, const ArgList &Args, + SmallVector &Binaries) { + auto List = listSection(Filename, Args); + if (!List) + return List.takeError(); + SmallVector TriplesInFile; + llvm::ErrorOr> TripleList = + llvm::MemoryBuffer::getFileOrSTDIN(*List, /*isText=*/true); + if (std::error_code EC = TripleList.getError()) + return createFileError(*List, EC); + (*TripleList) + ->getBuffer() + .split(TriplesInFile, '\n', /*MaxSplit=*/-1, /*KeepEmpty=*/false); + for (StringRef TripleStr : TriplesInFile) { + StringRef SYCLPrefix = "sycl-"; + if (!TripleStr.starts_with(SYCLPrefix)) + continue; + llvm::Triple Triple(TripleStr.substr(SYCLPrefix.size())); + auto UnbundledFile = unbundle(Filename, Args, Triple); + if (!UnbundledFile) + return UnbundledFile.takeError(); + if (*UnbundledFile == Filename) + continue; + + SmallVector ObjectFilePaths; + if (sycl::isStaticArchiveFile(Filename)) { + llvm::ErrorOr> ObjList = + llvm::MemoryBuffer::getFileOrSTDIN(*UnbundledFile, /*isText=*/true); + if (std::error_code EC = ObjList.getError()) + return createFileError(*UnbundledFile, EC); + (*ObjList)->getBuffer().split(ObjectFilePaths, '\n', /*MaxSplit=*/-1, + /*KeepEmpty=*/false); + } else { + ObjectFilePaths.push_back(*UnbundledFile); + } + for (StringRef ObjectFilePath : ObjectFilePaths) { + llvm::file_magic Magic; + llvm::identify_magic(ObjectFilePath, Magic); + if (Magic == file_magic::spirv_object) + return createStringError( + "SPIR-V fat objects must be generated with --offload-new-driver"); + auto Arg = Args.MakeArgString( + "sycl-" + + (Triple.isSPIROrSPIRV() ? Triple.str() + "-" : Triple.str()) + "=" + + ObjectFilePath); + auto Binary = getInputBitcodeLibrary(Arg); + + if (!Binary) + return Binary.takeError(); + + Binaries.push_back(std::move(*Binary)); + } + } + return Error::success(); +} + } // namespace sycl namespace generic { @@ -2634,8 +2763,14 @@ getDeviceInput(const ArgList &Args) { if (identify_magic(Buffer.getBuffer()) == file_magic::elf_shared_object) continue; SmallVector Binaries; + size_t OldSize = Binaries.size(); if (Error Err = extractOffloadBinaries(Buffer, Binaries)) return std::move(Err); + if (Binaries.size() == OldSize) { + if (Error Err = sycl::extractBundledObjects(*Filename, Args, Binaries)) + return std::move(Err); + } + for (auto &OffloadFile : Binaries) { if (identify_magic(Buffer.getBuffer()) == file_magic::archive && !WholeArchive) diff --git a/sycl/test-e2e/NewOffloadDriver/multisource.cpp b/sycl/test-e2e/NewOffloadDriver/multisource.cpp index 83223d4b68e09..0d7bec07654fb 100644 --- a/sycl/test-e2e/NewOffloadDriver/multisource.cpp +++ b/sycl/test-e2e/NewOffloadDriver/multisource.cpp @@ -21,6 +21,23 @@ // RUN: %clangxx -Wno-error=unused-command-line-argument -fsycl -fsycl-targets=%{sycl_triple} --offload-new-driver %t.init.o %t.calc.o %t.main.o -o %t.fat // RUN: %{run} %t.fat +// Multiple sources with kernel code with old-style objects +// Test with `--offload-new-driver` +// RUN: %{build} --no-offload-new-driver -c -o %t.init.o -DINIT_KERNEL +// RUN: %{build} --no-offload-new-driver -c -o %t.calc.o -DCALC_KERNEL +// RUN: %{build} --no-offload-new-driver -c -o %t.main.o -DMAIN_APP +// RUN: %clangxx -Wno-error=unused-command-line-argument -fsycl -fsycl-targets=%{sycl_triple} --offload-new-driver %t.init.o %t.calc.o %t.main.o -o %t.fat +// RUN: %{run} %t.fat + +// Multiple sources with kernel code with old-style objects in a static archive +// Test with `--offload-new-driver` +// RUN: %{build} --no-offload-new-driver -c -o %t.init.o -DINIT_KERNEL +// RUN: %{build} --no-offload-new-driver -c -o %t.calc.o -DCALC_KERNEL +// RUN: %{build} --no-offload-new-driver -c -o %t.main.o -DMAIN_APP +// RUN: llvm-ar r %t.a %t.init.o %t.calc.o +// RUN: %clangxx -Wno-error=unused-command-line-argument -fsycl -fsycl-targets=%{sycl_triple} --offload-new-driver %t.main.o %t.a -o %t.fat +// RUN: %{run} %t.fat + #include #include