1//===----------------------------------------------------------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// This tool executes a sequence of steps required to link device code in SYCL
10// device images. SYCL device code linking requires a complex sequence of steps
11// that include linking of llvm bitcode files, linking bitcode library files
12// with the fully linked source bitcode file(s), running several SYCL specific
13// post-link steps on the fully linked bitcode file(s), and finally generating
14// target-specific device code.
15//
16//===----------------------------------------------------------------------===//
17
18#include "clang/Basic/OffloadArch.h"
19#include "clang/Basic/Version.h"
20
21#include "llvm/ADT/STLExtras.h"
22#include "llvm/ADT/StringExtras.h"
23#include "llvm/ADT/StringMap.h"
24#include "llvm/ADT/StringSwitch.h"
25#include "llvm/BinaryFormat/Magic.h"
26#include "llvm/Bitcode/BitcodeReader.h"
27#include "llvm/Bitcode/BitcodeWriter.h"
28#include "llvm/CodeGen/CommandFlags.h"
29#include "llvm/Frontend/Offloading/Utility.h"
30#include "llvm/IR/DiagnosticPrinter.h"
31#include "llvm/IR/LLVMContext.h"
32#include "llvm/IRReader/IRReader.h"
33#include "llvm/LTO/LTO.h"
34#include "llvm/Linker/Linker.h"
35#include "llvm/MC/TargetRegistry.h"
36#include "llvm/Object/Archive.h"
37#include "llvm/Object/Binary.h"
38#include "llvm/Object/IRObjectFile.h"
39#include "llvm/Object/IRSymtab.h"
40#include "llvm/Object/OffloadBinary.h"
41#include "llvm/Option/ArgList.h"
42#include "llvm/Option/OptTable.h"
43#include "llvm/Option/Option.h"
44#include "llvm/Support/CommandLine.h"
45#include "llvm/Support/FileOutputBuffer.h"
46#include "llvm/Support/FileSystem.h"
47#include "llvm/Support/FormatVariadic.h"
48#include "llvm/Support/InitLLVM.h"
49#include "llvm/Support/MemoryBuffer.h"
50#include "llvm/Support/Path.h"
51#include "llvm/Support/Program.h"
52#include "llvm/Support/Signals.h"
53#include "llvm/Support/StringSaver.h"
54#include "llvm/Support/TargetSelect.h"
55#include "llvm/Support/TimeProfiler.h"
56#include "llvm/Support/WithColor.h"
57#include "llvm/Target/TargetMachine.h"
58#include "llvm/Transforms/Utils/SplitModuleByCategory.h"
59
60using namespace llvm;
61using namespace llvm::opt;
62using namespace llvm::object;
63using namespace clang;
64
65/// Print commands with arguments without executing.
66static bool DryRun = false;
67
68/// Print verbose output.
69static bool Verbose = false;
70
71/// Filename of the output being created.
72static StringRef OutputFile;
73
74/// Directory to dump SPIR-V IR if requested by user.
75static SmallString<128> SPIRVDumpDir;
76
77using OffloadingImage = OffloadBinary::OffloadingImage;
78
79static void printVersion(raw_ostream &OS) {
80 OS << clang::getClangToolFullVersion(ToolName: "clang-sycl-linker") << '\n';
81}
82
83/// The value of `argv[0]` when run.
84static const char *Executable;
85
86/// Temporary files to be cleaned up.
87static SmallVector<SmallString<128>> TempFiles;
88
89namespace {
90// Must not overlap with llvm::opt::DriverFlag.
91enum LinkerFlags { LinkerOnlyOption = (1 << 4) };
92
93enum ID {
94 OPT_INVALID = 0, // This is not an option ID.
95#define OPTION(...) LLVM_MAKE_OPT_ID(__VA_ARGS__),
96#include "SYCLLinkOpts.inc"
97 LastOption
98#undef OPTION
99};
100
101#define OPTTABLE_CODE
102#include "SYCLLinkOpts.inc"
103
104class LinkerOptTable : public opt::OptTable {
105public:
106 LinkerOptTable() : opt::OptTable(optionTables()) {}
107};
108} // namespace
109
110static const OptTable &getOptTable() {
111 static const LinkerOptTable *Table = []() {
112 auto Result = std::make_unique<LinkerOptTable>();
113 return Result.release();
114 }();
115 return *Table;
116}
117
118[[noreturn]] static void reportError(Error E) {
119 outs().flush();
120 logAllUnhandledErrors(E: std::move(E), OS&: WithColor::error(OS&: errs(), Prefix: Executable));
121 exit(EXIT_FAILURE);
122}
123
124static std::string getMainExecutable(const char *Name) {
125 void *Ptr = (void *)(intptr_t)&getMainExecutable;
126 auto COWPath = sys::fs::getMainExecutable(argv0: Name, MainExecAddr: Ptr);
127 return sys::path::parent_path(path: COWPath).str();
128}
129
130static Expected<StringRef>
131createTempFile(const ArgList &Args, const Twine &Prefix, StringRef Extension) {
132 SmallString<128> Path;
133 if (Args.hasArg(Ids: OPT_save_temps) || DryRun) {
134 // Generate a unique path name without creating a file
135 sys::fs::createUniquePath(Model: Prefix + "-%%%%%%." + Extension, ResultPath&: Path,
136 /*MakeAbsolute=*/false);
137 } else {
138 if (std::error_code EC =
139 sys::fs::createTemporaryFile(Prefix, Suffix: Extension, ResultPath&: Path))
140 return createFileError(F: Path, EC);
141 }
142
143 TempFiles.emplace_back(Args: std::move(Path));
144 return TempFiles.back();
145}
146
147static Expected<std::string> findProgram(const ArgList &Args, StringRef Name,
148 ArrayRef<StringRef> Paths) {
149 if (DryRun)
150 return Name.str();
151 ErrorOr<std::string> Path = sys::findProgramByName(Name, Paths);
152 if (!Path)
153 Path = sys::findProgramByName(Name);
154 if (!Path)
155 return createStringError(EC: Path.getError(),
156 S: "unable to find '" + Name + "' in path");
157 return *Path;
158}
159
160static void printCommands(ArrayRef<StringRef> CmdArgs) {
161 if (CmdArgs.empty())
162 return;
163
164 llvm::errs() << " \"" << CmdArgs.front() << "\" ";
165 llvm::errs() << llvm::join(Begin: std::next(x: CmdArgs.begin()), End: CmdArgs.end(), Separator: " ")
166 << "\n";
167}
168
169/// Execute the command \p ExecutablePath with the arguments \p Args.
170static Error executeCommands(StringRef ExecutablePath,
171 ArrayRef<StringRef> Args) {
172 if (Verbose || DryRun)
173 printCommands(CmdArgs: Args);
174
175 if (DryRun)
176 return Error::success();
177
178 if (sys::ExecuteAndWait(Program: ExecutablePath, Args))
179 return createStringError(Fmt: "'%s' failed",
180 Vals: sys::path::filename(path: ExecutablePath).str().c_str());
181 return Error::success();
182}
183
184namespace {
185/// A minimal symbol interface used to drive archive member extraction. Only the
186/// flags required by the symbol-resolution fixed-point loop are tracked.
187struct Symbol {
188 enum Flags {
189 None = 0,
190 Undefined = 1 << 0,
191 Weak = 1 << 1,
192 };
193
194 Symbol() : SymFlags(None) {}
195 Symbol(Symbol::Flags F) : SymFlags(F) {}
196 Symbol(const irsymtab::Reader::SymbolRef Sym) : SymFlags(0) {
197 if (Sym.isUndefined())
198 SymFlags |= Undefined;
199 if (Sym.isWeak())
200 SymFlags |= Weak;
201 }
202
203 bool isWeak() const { return SymFlags & Weak; }
204 bool isUndefined() const { return SymFlags & Undefined; }
205
206 uint32_t SymFlags;
207};
208
209/// Description of a single input (positional file or -l library).
210struct InputDesc {
211 enum class Kind { File, Library };
212
213 StringRef Value; // File path, or library name for -l (the value after -l).
214 Kind InputKind = Kind::File;
215 bool WholeArchive = false; // --whole-archive state in effect at this input.
216};
217
218/// An input buffer pending archive-member resolution, together with its parsed
219/// IR symbol table. The symbol table is parsed once and reused across all
220/// fixed-point passes so members are not re-parsed on every pass.
221struct PendingInput {
222 std::unique_ptr<MemoryBuffer> Buffer;
223 bool IsLazy = false;
224 bool FromArchive = false;
225 IRSymtabFile Symtab;
226};
227
228/// Resolved input buffers and their target triple.
229struct ResolvedInputs {
230 SmallVector<std::unique_ptr<MemoryBuffer>> Buffers;
231 llvm::Triple TargetTriple;
232 StringRef TripleSource; // Source of the triple (--triple= or filename)
233};
234} // namespace
235
236static std::optional<std::string> findFile(StringRef Dir, const Twine &Name) {
237 SmallString<128> Path;
238 sys::path::append(path&: Path, a: Dir, b: Name);
239 // Skip directories so a directory whose name matches the requested library
240 // does not stop the search; a later -L path may hold the real archive.
241 if (sys::fs::exists(Path) && !sys::fs::is_directory(Path))
242 return static_cast<std::string>(Path);
243 return std::nullopt;
244}
245
246static std::optional<std::string>
247findFromSearchPaths(StringRef Name, ArrayRef<StringRef> SearchPaths) {
248 for (StringRef Dir : SearchPaths)
249 if (std::optional<std::string> File = findFile(Dir, Name))
250 return File;
251 return std::nullopt;
252}
253
254/// Search for static libraries in the linker's library path given input like
255/// `-lfoo`, `-l:libfoo.a`, or `-l/absolute/path/to/lib.a`.
256static std::optional<std::string>
257searchLibrary(StringRef Input, ArrayRef<StringRef> SearchPaths) {
258 // An absolute path is taken as-is; -L paths are only consulted for relative
259 // names.
260 if (sys::path::is_absolute(path: Input)) {
261 if (sys::fs::exists(Path: Input) && !sys::fs::is_directory(Path: Input))
262 return Input.str();
263 return std::nullopt;
264 }
265
266 if (Input.starts_with(Prefix: ":"))
267 return findFromSearchPaths(Name: Input.drop_front(), SearchPaths);
268 SmallString<128> LibName("lib");
269 LibName += Input;
270 LibName += ".a";
271 return findFromSearchPaths(Name: LibName, SearchPaths);
272}
273
274/// Scan a member's pre-parsed IR symbol table against \p LinkerSymtab and
275/// return true if the member should be extracted: it is non-lazy, or it defines
276/// a symbol that resolves a currently-undefined reference. Mirrors a linker's
277/// archive member selection.
278static bool scanSymbols(const IRSymtabFile &MemberSymtab,
279 StringMap<Symbol> &LinkerSymtab, bool IsLazy) {
280 bool Extracted = !IsLazy;
281 StringMap<Symbol> PendingSymbols;
282 for (unsigned ModIdx = 0; ModIdx != MemberSymtab.Mods.size(); ++ModIdx) {
283 for (const auto &IRSym : MemberSymtab.TheReader.module_symbols(I: ModIdx)) {
284 if (IRSym.isFormatSpecific() || !IRSym.isGlobal())
285 continue;
286
287 bool IsNewSymbol = IsLazy && !LinkerSymtab.count(Key: IRSym.getName());
288 StringMap<Symbol> &Target = IsNewSymbol ? PendingSymbols : LinkerSymtab;
289 Symbol Sym(IRSym);
290 auto [It, Inserted] = Target.try_emplace(Key: IRSym.getName(), Args&: Sym);
291 // A freshly inserted entry has no prior symbol to resolve or upgrade, so
292 // it cannot trigger extraction.
293 if (Inserted)
294 continue;
295
296 Symbol &OldSym = It->second;
297 bool ResolvesReference =
298 !Sym.isUndefined() &&
299 (OldSym.isUndefined() || (OldSym.isWeak() && !Sym.isWeak())) &&
300 !(OldSym.isWeak() && OldSym.isUndefined() && IsLazy);
301 Extracted |= ResolvesReference;
302
303 if (ResolvesReference)
304 OldSym = Sym;
305 }
306 }
307 if (Extracted && IsLazy)
308 for (const auto &[Name, Sym] : PendingSymbols)
309 LinkerSymtab[Name] = Sym;
310 return Extracted;
311}
312
313/// Parse \p Buffer's IR symbol table and append it to \p Inputs. Errors if the
314/// buffer is not LLVM bitcode (the only member type the SYCL linker supports).
315static Error addBitcodeInput(SmallVector<PendingInput> &Inputs,
316 std::unique_ptr<MemoryBuffer> Buffer, bool IsLazy,
317 bool FromArchive) {
318 if (identify_magic(magic: Buffer->getBuffer()) != file_magic::bitcode)
319 return createStringError(S: "unsupported file type: '" +
320 Buffer->getBufferIdentifier() + "'");
321 Expected<IRSymtabFile> SymtabOrErr = readIRSymtab(MBRef: Buffer->getMemBufferRef());
322 if (!SymtabOrErr)
323 return SymtabOrErr.takeError();
324 Inputs.push_back(
325 Elt: {.Buffer: std::move(Buffer), .IsLazy: IsLazy, .FromArchive: FromArchive, .Symtab: std::move(*SymtabOrErr)});
326 return Error::success();
327}
328
329/// Resolve archive members from the given inputs using a symbol-driven
330/// fixed-point algorithm. For each input:
331/// - If it's a Library, search for lib<name>.a or :<name> in SearchPaths
332/// - If it's a File, use the path directly
333/// - Archives are expanded and members are lazily extracted based on symbol
334/// references unless WholeArchive is true
335/// - Non-archive bitcode inputs are always included
336///
337/// Returns the buffers to link, in extraction order, along with the resolved
338/// target triple. All returned buffers have compatible target triples;
339/// incompatible archive members are filtered during resolution.
340static Expected<ResolvedInputs> resolveArchiveMembers(
341 ArrayRef<InputDesc> Order, ArrayRef<StringRef> SearchPaths,
342 ArrayRef<StringRef> ForcedUndefs, StringRef TargetTripleArgValue) {
343 // Collect every candidate member, parsing each one's IR symbol table once.
344 SmallVector<PendingInput> Inputs;
345
346 for (const InputDesc &Desc : Order) {
347 std::optional<std::string> Filename;
348
349 if (Desc.InputKind == InputDesc::Kind::Library) {
350 Filename = searchLibrary(Input: Desc.Value, SearchPaths);
351 if (!Filename)
352 return createStringError(S: "unable to find library -l" + Desc.Value);
353 } else {
354 if (!sys::fs::exists(Path: Desc.Value))
355 return createStringError(S: "input file not found: '" + Desc.Value + "'");
356 if (sys::fs::is_directory(Path: Desc.Value))
357 return createStringError(S: "'" + Desc.Value + "': is a directory");
358 Filename = Desc.Value.str();
359 }
360
361 auto BufferOrErr =
362 errorOrToExpected(EO: MemoryBuffer::getFileOrSTDIN(Filename: *Filename));
363 if (!BufferOrErr)
364 return createFileError(F: *Filename, E: BufferOrErr.takeError());
365
366 MemoryBufferRef Buffer = (*BufferOrErr)->getMemBufferRef();
367 switch (identify_magic(magic: Buffer.getBuffer())) {
368 case file_magic::bitcode:
369 if (Error Err = addBitcodeInput(Inputs, Buffer: std::move(*BufferOrErr),
370 /*IsLazy=*/false, /*FromArchive=*/false))
371 return Err;
372 break;
373 case file_magic::archive: {
374 Expected<std::unique_ptr<object::Archive>> LibFile =
375 object::Archive::create(Source: Buffer);
376 if (!LibFile)
377 return LibFile.takeError();
378 Error Err = Error::success();
379 for (auto Child : (*LibFile)->children(Err)) {
380 auto ChildBufferOrErr = Child.getMemoryBufferRef();
381 if (!ChildBufferOrErr)
382 return ChildBufferOrErr.takeError();
383 // Include archive name in buffer identifier for better diagnostics.
384 std::string BufferIdentifier =
385 (*Filename + "(" + ChildBufferOrErr->getBufferIdentifier() + ")")
386 .str();
387 std::unique_ptr<MemoryBuffer> ChildBuffer =
388 MemoryBuffer::getMemBufferCopy(InputData: ChildBufferOrErr->getBuffer(),
389 BufferName: BufferIdentifier);
390 if (Error E = addBitcodeInput(Inputs, Buffer: std::move(ChildBuffer),
391 IsLazy: !Desc.WholeArchive, /*FromArchive=*/true))
392 return E;
393 }
394 if (Err)
395 return Err;
396 break;
397 }
398 default:
399 return createStringError(S: "unsupported file type: '" + *Filename + "'");
400 }
401 }
402
403 // Resolve the target triple: use --triple= if provided, otherwise infer from
404 // the first non-archive input with a non-empty triple.
405 llvm::Triple TargetTriple(TargetTripleArgValue);
406 StringRef TripleSource = TargetTriple.empty() ? "" : "--triple=";
407
408 if (TargetTriple.empty()) {
409 for (const PendingInput &In : Inputs) {
410 if (!In.FromArchive && In.Symtab.Mods.size() > 0) {
411 StringRef Triple = In.Symtab.TheReader.getTargetTriple();
412 if (!Triple.empty()) {
413 TargetTriple = llvm::Triple(Triple);
414 TripleSource = In.Buffer->getBufferIdentifier();
415 break;
416 }
417 }
418 }
419 }
420
421 // Seed symbol table with forced undefined symbols.
422 StringMap<Symbol> SymTab;
423 for (StringRef Sym : ForcedUndefs)
424 SymTab[Sym] = Symbol(Symbol::Undefined);
425
426 // Fixed-point loop to extract archive members. Each pass may resolve symbols
427 // that unlock further members; iterate until no new member is extracted.
428 SmallVector<std::unique_ptr<MemoryBuffer>> Resolved;
429 bool KeepExtracting = true;
430 while (KeepExtracting) {
431 KeepExtracting = false;
432 for (PendingInput &In : Inputs) {
433 if (!In.Buffer)
434 continue;
435
436 // Filter archive members by target triple before symbol scanning.
437 // Members built for a different target are silently skipped, matching how
438 // a real linker treats device libraries built for other architectures.
439 if (In.FromArchive) {
440 StringRef MemberTriple = In.Symtab.TheReader.getTargetTriple();
441 if (!MemberTriple.empty() &&
442 llvm::Triple(MemberTriple) != TargetTriple) {
443 if (Verbose)
444 errs() << formatv(
445 Fmt: "archive resolution: skipping {0}: triple {1} != {2}\n",
446 Vals: In.Buffer->getBufferIdentifier(), Vals&: MemberTriple,
447 Vals: TargetTriple.str());
448 In.Buffer.reset();
449 In.Symtab = {};
450 continue;
451 }
452 }
453
454 if (!scanSymbols(MemberSymtab: In.Symtab, LinkerSymtab&: SymTab, IsLazy: In.IsLazy))
455 continue;
456 KeepExtracting = true;
457 Resolved.push_back(Elt: std::move(In.Buffer));
458 }
459 }
460
461 return ResolvedInputs{.Buffers: std::move(Resolved), .TargetTriple: std::move(TargetTriple),
462 .TripleSource: TripleSource};
463}
464
465static Expected<ResolvedInputs> getInput(const ArgList &Args) {
466 // Build input descriptors for the archive resolver.
467 SmallVector<InputDesc> InputDescs;
468 bool WholeArchive = false;
469 for (const opt::Arg *Arg : Args.filtered(
470 Ids: OPT_INPUT, Ids: OPT_library, Ids: OPT_whole_archive, Ids: OPT_no_whole_archive)) {
471 if (Arg->getOption().matches(ID: OPT_whole_archive) ||
472 Arg->getOption().matches(ID: OPT_no_whole_archive)) {
473 WholeArchive = Arg->getOption().matches(ID: OPT_whole_archive);
474 continue;
475 }
476
477 InputDesc Desc;
478 Desc.Value = Arg->getValue();
479 Desc.InputKind = Arg->getOption().matches(ID: OPT_library)
480 ? InputDesc::Kind::Library
481 : InputDesc::Kind::File;
482 Desc.WholeArchive = WholeArchive;
483 InputDescs.push_back(Elt: Desc);
484 }
485
486 if (InputDescs.empty())
487 return createStringError(Fmt: "no input files provided");
488
489 // Gather search paths and forced undefined symbols.
490 SmallVector<StringRef> LibraryPaths;
491 for (const opt::Arg *Arg : Args.filtered(Ids: OPT_library_path))
492 LibraryPaths.push_back(Elt: Arg->getValue());
493
494 // getAllArgValues returns a temporary vector; retain it so the StringRefs
495 // remain valid through the resolveArchiveMembers call.
496 std::vector<std::string> ForcedUndefStorage = Args.getAllArgValues(Id: OPT_u);
497 SmallVector<StringRef> ForcedUndefs(ForcedUndefStorage.begin(),
498 ForcedUndefStorage.end());
499
500 // Get target triple from command line if specified.
501 StringRef TargetTripleStr = Args.getLastArgValue(Id: OPT_triple_EQ);
502
503 Expected<ResolvedInputs> ResolvedOrErr = resolveArchiveMembers(
504 Order: InputDescs, SearchPaths: LibraryPaths, ForcedUndefs, TargetTripleArgValue: TargetTripleStr);
505 if (!ResolvedOrErr)
506 return ResolvedOrErr.takeError();
507
508 if (ResolvedOrErr->Buffers.empty())
509 return createStringError(Fmt: "no input files could be resolved");
510
511 if (ResolvedOrErr->TargetTriple.empty())
512 return createStringError(
513 Fmt: "target triple must be specified or inferable from inputs");
514
515 return std::move(*ResolvedOrErr);
516}
517
518namespace {
519struct LinkResult {
520 std::unique_ptr<Module> LinkedModule;
521 SmallString<256> BitcodeFile;
522 llvm::Triple TargetTriple;
523};
524} // namespace
525
526/// Link all resolved input bitcode images into one module. All resolved inputs
527/// are guaranteed to have compatible target triples (incompatible archive
528/// members are filtered during archive resolution). Triple conflicts between
529/// regular (non-archive) inputs are hard errors caught before running
530/// linkInModule.
531static Expected<LinkResult>
532linkInputs(ArrayRef<std::unique_ptr<MemoryBuffer>> Inputs,
533 const llvm::Triple &TargetTriple, StringRef TripleSource,
534 const ArgList &Args, LLVMContext &C) {
535 llvm::TimeTraceScope TimeScope("Link code");
536
537 assert(Inputs.size() && "No inputs to link");
538
539 // Create a new file to write the linked file to.
540 auto BitcodeOutput =
541 createTempFile(Args, Prefix: sys::path::filename(path: OutputFile), Extension: "bc");
542 if (!BitcodeOutput)
543 return BitcodeOutput.takeError();
544
545 if (Verbose) {
546 std::string InputList =
547 llvm::join(R: llvm::map_range(C&: Inputs,
548 F: [](const auto &Buffer) {
549 return Buffer->getBufferIdentifier();
550 }),
551 Separator: ", ");
552 errs() << formatv(Fmt: "link: inputs: {0} output: {1}\n", Vals&: InputList,
553 Vals&: *BitcodeOutput);
554 }
555
556 auto LinkerOutput = std::make_unique<Module>(args: "linker-output", args&: C);
557 Linker L(*LinkerOutput);
558
559 for (const auto &Buffer : Inputs) {
560 auto ModOrErr = parseBitcodeFile(Buffer: Buffer->getMemBufferRef(), Context&: C);
561 if (!ModOrErr)
562 return ModOrErr.takeError();
563
564 const llvm::Triple &T = (*ModOrErr)->getTargetTriple();
565 if (!T.empty() && T != TargetTriple) {
566 // All incompatible archive members should have been filtered during
567 // resolution, so this is a conflict between regular inputs.
568 return createStringError(S: "conflicting target triples: '" +
569 TargetTriple.str() + "' (from " + TripleSource +
570 ") vs '" + T.str() + "' (from " +
571 Buffer->getBufferIdentifier() + ")");
572 }
573
574 if (L.linkInModule(Src: std::move(*ModOrErr)))
575 return createStringError(Fmt: "could not link IR");
576 }
577
578 // Dump linked output for testing.
579 if (Args.hasArg(Ids: OPT_print_linked_module))
580 outs() << *LinkerOutput;
581
582 // Write the final output into 'BitcodeOutput' file.
583 if (!DryRun) {
584 int FD = -1;
585 if (std::error_code EC = sys::fs::openFileForWrite(Name: *BitcodeOutput, ResultFD&: FD))
586 return errorCodeToError(EC);
587 llvm::raw_fd_ostream OS(FD, true);
588 WriteBitcodeToFile(M: *LinkerOutput, Out&: OS);
589 }
590
591 return LinkResult{.LinkedModule: std::move(LinkerOutput), .BitcodeFile: SmallString<256>(*BitcodeOutput),
592 .TargetTriple: std::move(TargetTriple)};
593}
594
595/// Run Code Generation using LLVM backend.
596/// \param File The input LLVM IR bitcode file.
597/// \param TargetTriple The resolved target triple.
598/// \param Args encompasses all arguments required for linking device code and
599/// will be parsed to generate options required to be passed into the backend.
600/// \param OutputFile The output file name.
601/// \param C The LLVM context.
602static Error runCodeGen(StringRef File, const llvm::Triple &TargetTriple,
603 const ArgList &Args, StringRef OutputFile,
604 LLVMContext &C) {
605 llvm::TimeTraceScope TimeScope("Code generation");
606
607 if (Verbose || DryRun)
608 errs() << formatv(Fmt: "LLVM backend: input: {0}, output: {1}\n", Vals&: File,
609 Vals&: OutputFile);
610
611 if (DryRun)
612 return Error::success();
613
614 // Parse input module.
615 SMDiagnostic Err;
616 std::unique_ptr<Module> M = parseIRFile(Filename: File, Err, Context&: C);
617 if (!M)
618 return createStringError(S: Err.getMessage());
619
620 if (Error MatErr = M->materializeAll())
621 return MatErr;
622
623 M->setTargetTriple(TargetTriple);
624
625 // Get a handle to a target backend.
626 std::string Msg;
627 const Target *T = TargetRegistry::lookupTarget(TheTriple: M->getTargetTriple(), Error&: Msg);
628 if (!T)
629 return createStringError(S: Msg + ": " + M->getTargetTriple().str());
630
631 // Allocate target machine.
632 TargetOptions Options;
633 std::optional<Reloc::Model> RM;
634 std::optional<CodeModel::Model> CM;
635 std::unique_ptr<TargetMachine> TM(
636 T->createTargetMachine(TT: M->getTargetTriple(), /*CPU=*/"",
637 /*Features=*/"", Options, RM, CM));
638 if (!TM)
639 return createStringError(Fmt: "could not allocate target machine");
640
641 // Set data layout if needed.
642 if (M->getDataLayout().isDefault())
643 M->setDataLayout(TM->createDataLayout());
644
645 // Open output file for writing.
646 int FD = -1;
647 if (std::error_code EC = sys::fs::openFileForWrite(Name: OutputFile, ResultFD&: FD))
648 return errorCodeToError(EC);
649 auto OS = std::make_unique<llvm::raw_fd_ostream>(args&: FD, args: true);
650
651 legacy::PassManager CodeGenPasses;
652 TargetLibraryInfoImpl TLII(M->getTargetTriple());
653 CodeGenPasses.add(P: new TargetLibraryInfoWrapperPass(TLII));
654 if (TM->addPassesToEmitFile(CodeGenPasses, *OS, nullptr,
655 CodeGenFileType::ObjectFile))
656 return createStringError(Fmt: "failed to execute LLVM backend");
657 CodeGenPasses.run(M&: *M);
658
659 return Error::success();
660}
661
662/// Run AOT compilation for Intel CPU.
663/// Calls opencl-aot tool to generate device code for the Intel OpenCL CPU
664/// Runtime.
665/// \param InputFile The input SPIR-V file.
666/// \param OutputFile The output file name.
667/// \param Args Encompasses all arguments required for linking and wrapping
668/// device code and will be parsed to generate options required to be passed
669/// into the AOT compilation step.
670static Error runAOTCompileIntelCPU(StringRef InputFile, StringRef OutputFile,
671 const ArgList &Args) {
672 SmallVector<StringRef, 8> CmdArgs;
673 Expected<std::string> OpenCLAOTPath =
674 findProgram(Args, Name: "opencl-aot", Paths: {getMainExecutable(Name: "opencl-aot")});
675 if (!OpenCLAOTPath)
676 return OpenCLAOTPath.takeError();
677
678 CmdArgs.push_back(Elt: *OpenCLAOTPath);
679 CmdArgs.push_back(Elt: "--device=cpu");
680 StringRef ExtraArgs = Args.getLastArgValue(Id: OPT_opencl_aot_options_EQ);
681 ExtraArgs.split(A&: CmdArgs, Separator: " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false);
682 CmdArgs.push_back(Elt: "-o");
683 CmdArgs.push_back(Elt: OutputFile);
684 CmdArgs.push_back(Elt: InputFile);
685 if (Error Err = executeCommands(ExecutablePath: *OpenCLAOTPath, Args: CmdArgs))
686 return Err;
687 return Error::success();
688}
689
690/// Run AOT compilation for Intel GPU.
691/// Calls ocloc tool to generate device code for the Intel Graphics Compute
692/// Runtime.
693/// \param InputFile The input SPIR-V file.
694/// \param OutputFile The output file name.
695/// \param Args Encompasses all arguments required for linking and wrapping
696/// device code and will be parsed to generate options required to be passed
697/// into the AOT compilation step.
698static Error runAOTCompileIntelGPU(StringRef InputFile, StringRef OutputFile,
699 const ArgList &Args) {
700 SmallVector<StringRef, 8> CmdArgs;
701 Expected<std::string> OclocPath =
702 findProgram(Args, Name: "ocloc", Paths: {getMainExecutable(Name: "ocloc")});
703 if (!OclocPath)
704 return OclocPath.takeError();
705
706 CmdArgs.push_back(Elt: *OclocPath);
707 // The next line prevents ocloc from modifying the image name
708 CmdArgs.push_back(Elt: "-output_no_suffix");
709 CmdArgs.push_back(Elt: "-spirv_input");
710
711 StringRef Arch(Args.getLastArgValue(Id: OPT_arch_EQ));
712 assert(!Arch.empty() && "Arch must be specified for AOT compilation");
713 CmdArgs.push_back(Elt: "-device");
714 CmdArgs.push_back(Elt: Arch);
715
716 // getAllArgValues returns a temporary vector; retain it so the StringRefs
717 // remain valid through the executeCommands call below.
718 std::vector<std::string> ExtraArgsStorage =
719 Args.getAllArgValues(Id: OPT_ocloc_options_EQ);
720 llvm::append_range(C&: CmdArgs, R&: ExtraArgsStorage);
721
722 CmdArgs.push_back(Elt: "-output");
723 CmdArgs.push_back(Elt: OutputFile);
724 CmdArgs.push_back(Elt: "-file");
725 CmdArgs.push_back(Elt: InputFile);
726 if (Error Err = executeCommands(ExecutablePath: *OclocPath, Args: CmdArgs))
727 return Err;
728 return Error::success();
729}
730
731/// Run AOT compilation for Intel CPU/GPU.
732/// \param InputFile The input SPIR-V file.
733/// \param OutputFile The output file name.
734/// \param Args Encompasses all arguments required for linking and wrapping
735/// device code and will be parsed to generate options required to be passed
736/// into the AOT compilation step.
737static Error runAOTCompile(StringRef InputFile, StringRef OutputFile,
738 const ArgList &Args) {
739 StringRef Arch = Args.getLastArgValue(Id: OPT_arch_EQ);
740 OffloadArch OA = StringToOffloadArch(S: Arch);
741 if (OA.isIntelGPU())
742 return runAOTCompileIntelGPU(InputFile, OutputFile, Args);
743 if (OA.isIntelCPU())
744 return runAOTCompileIntelCPU(InputFile, OutputFile, Args);
745
746 llvm_unreachable("runAOTCompile dispatched on unsupported arch");
747}
748
749static constexpr char AttrSYCLModuleId[] = "sycl-module-id";
750
751namespace {
752/// SYCL device code module split mode.
753enum class IRSplitMode {
754 SPLIT_PER_TU, // one module per translation unit
755 SPLIT_PER_KERNEL, // one module per kernel
756 SPLIT_NONE // no splitting
757};
758} // namespace
759
760/// Parses the value of \p --module-split-mode.
761static std::optional<IRSplitMode> convertStringToSplitMode(StringRef S) {
762 return StringSwitch<std::optional<IRSplitMode>>(S)
763 .Case(S: "translation_unit", Value: IRSplitMode::SPLIT_PER_TU)
764 .Case(S: "kernel", Value: IRSplitMode::SPLIT_PER_KERNEL)
765 .Case(S: "link_unit", Value: IRSplitMode::SPLIT_NONE)
766 .Default(Value: std::nullopt);
767}
768
769static StringRef splitModeToString(IRSplitMode Mode) {
770 switch (Mode) {
771 case IRSplitMode::SPLIT_PER_TU:
772 return "translation_unit";
773 case IRSplitMode::SPLIT_PER_KERNEL:
774 return "kernel";
775 case IRSplitMode::SPLIT_NONE:
776 return "link_unit";
777 }
778 llvm_unreachable("bad split mode");
779}
780
781namespace {
782/// Result of splitting a device module: the bitcode file path and the
783/// serialized symbol table for each device image.
784struct SplitModule {
785 SmallString<256> ModuleFilePath;
786 SmallString<0> Symbols;
787};
788} // namespace
789
790static bool isEntryPoint(const Function &F, bool EmitOnlyKernelsAsEntryPoints) {
791 if (F.isDeclaration())
792 return false;
793 if (F.hasKernelCallingConv())
794 return true;
795 if (EmitOnlyKernelsAsEntryPoints)
796 return false;
797 // sycl_external functions carry the "sycl-module-id" attribute.
798 return F.hasFnAttribute(Kind: AttrSYCLModuleId);
799}
800
801/// Collect entry point names from \p M and serialize them into a symbol table.
802static SmallString<0> collectEntryPoints(const Module &M,
803 bool EmitOnlyKernelsAsEntryPoints) {
804 SmallVector<StringRef> Names;
805 for (const Function &F : M)
806 if (isEntryPoint(F, EmitOnlyKernelsAsEntryPoints))
807 Names.push_back(Elt: F.getName());
808 SmallString<0> SymbolData;
809 llvm::offloading::sycl::writeSymbolTable(Names, Out&: SymbolData);
810 return SymbolData;
811}
812
813namespace {
814/// Functor passed to splitModuleTransitiveFromEntryPoints. For each input
815/// function \p F, returns a numeric group ID (if \p F is an entry point)
816/// determining which device image it lands in, or std::nullopt (for
817/// non-entry-points). SPLIT_PER_KERNEL \p Mode gives each kernel its own ID;
818/// SPLIT_PER_TU \p Mode groups kernels by their "sycl-module-id" attribute
819/// value.
820class EntryPointCategorizer {
821public:
822 EntryPointCategorizer(IRSplitMode Mode, bool EmitOnlyKernelsAsEntryPoints)
823 : Mode(Mode), OnlyKernelsAreEntryPoints(EmitOnlyKernelsAsEntryPoints) {}
824
825 std::optional<int> operator()(const Function &F) {
826 if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints: OnlyKernelsAreEntryPoints))
827 return std::nullopt;
828
829 std::string Key;
830 switch (Mode) {
831 case IRSplitMode::SPLIT_PER_KERNEL:
832 Key = F.getName().str();
833 break;
834 case IRSplitMode::SPLIT_PER_TU:
835 Key = F.getFnAttribute(Kind: AttrSYCLModuleId).getValueAsString().str();
836 break;
837 case IRSplitMode::SPLIT_NONE:
838 llvm_unreachable("categorizer cannot be used for SPLIT_NONE");
839 }
840
841 auto [It, Inserted] =
842 StrToId.try_emplace(Key: std::move(Key), Args: static_cast<int>(StrToId.size()));
843 return It->second;
844 }
845
846private:
847 IRSplitMode Mode;
848 bool OnlyKernelsAreEntryPoints;
849 llvm::StringMap<int> StrToId;
850};
851} // namespace
852
853/// Splits the fully linked device \p M into one bitcode file per device image
854/// according to \p Mode and returns the list of split images with their symbol
855/// tables. The module is split transitively from entry points; each part is
856/// written to a fresh temporary bitcode file.
857static Expected<SmallVector<SplitModule, 0>>
858splitDeviceCode(std::unique_ptr<Module> M, StringRef LinkedBitcodeFile,
859 IRSplitMode Mode, bool EmitOnlyKernelsAsEntryPoints,
860 const ArgList &Args) {
861 assert(Mode != IRSplitMode::SPLIT_NONE && "SPLIT_NONE is unsupported");
862
863 SmallVector<SplitModule, 0> SplitModules;
864 EntryPointCategorizer Categorizer(Mode, EmitOnlyKernelsAsEntryPoints);
865
866 auto SplitCallback = [&](std::unique_ptr<Module> Part) -> Error {
867 Expected<StringRef> BitcodeFileOrErr =
868 createTempFile(Args, Prefix: sys::path::filename(path: OutputFile), Extension: "bc");
869 if (!BitcodeFileOrErr)
870 return BitcodeFileOrErr.takeError();
871
872 if (!DryRun) {
873 int FD = -1;
874 if (std::error_code EC = sys::fs::openFileForWrite(Name: *BitcodeFileOrErr, ResultFD&: FD))
875 return errorCodeToError(EC);
876 raw_fd_ostream OS(FD, /*shouldClose=*/true);
877 WriteBitcodeToFile(M: *Part, Out&: OS);
878 }
879
880 SplitModules.push_back(
881 Elt: {.ModuleFilePath: SmallString<256>(*BitcodeFileOrErr),
882 .Symbols: collectEntryPoints(M: *Part, EmitOnlyKernelsAsEntryPoints)});
883 return Error::success();
884 };
885
886 if (Error Err = splitModuleTransitiveFromEntryPoints(
887 M: std::move(M), EntryPointCategorizer: Categorizer, Callback: SplitCallback))
888 return Err;
889
890 if (Verbose) {
891 errs() << formatv(Fmt: "sycl-module-split: input: {0}, mode: {1}\n",
892 Vals&: LinkedBitcodeFile, Vals: splitModeToString(Mode));
893 for (const SplitModule &SI : SplitModules) {
894 errs() << formatv(Fmt: "{0} [", Vals: SI.ModuleFilePath);
895 llvm::offloading::sycl::forEachSymbol(
896 Symbols: SI.Symbols, Callback: [](StringRef Name) { errs() << Name << " "; });
897 errs() << "]\n";
898 }
899 }
900
901 return SplitModules;
902}
903
904/// Returns true if module splitting can be skipped: either \p Mode is
905/// SPLIT_NONE, or \p M contains no entry points (nothing to split from).
906static bool canSkipModuleSplit(IRSplitMode Mode, const Module &M,
907 bool EmitOnlyKernelsAsEntryPoints) {
908 if (Mode == IRSplitMode::SPLIT_NONE)
909 return true;
910 return llvm::none_of(Range: M.functions(), P: [&](const Function &F) {
911 return isEntryPoint(F, EmitOnlyKernelsAsEntryPoints);
912 });
913}
914
915/// Performs the following steps:
916/// 1. Link all input bitcode files together with library files.
917/// 2. Optionally split the linked module according to the requested
918/// IRSplitMode.
919/// 3. Run SPIR-V code generation on each (split) module.
920/// 4. Optionally run AOT compilation when targeting an Intel HW arch.
921/// 5. Pack the resulting images into a single OffloadBinary written to the
922/// output file.
923static Error runSYCLLink(ArrayRef<std::unique_ptr<MemoryBuffer>> Inputs,
924 const llvm::Triple &TargetTriple,
925 StringRef TripleSource, const ArgList &Args) {
926 llvm::TimeTraceScope TimeScope("SYCL linking");
927
928 LLVMContext C;
929
930 // Link all input bitcode files and library files.
931 Expected<LinkResult> LinkedOrErr =
932 linkInputs(Inputs, TargetTriple, TripleSource, Args, C);
933 if (!LinkedOrErr)
934 return LinkedOrErr.takeError();
935 LinkResult &Result = *LinkedOrErr;
936
937 // Determine the requested module split mode.
938 IRSplitMode SplitMode = IRSplitMode::SPLIT_PER_TU;
939 if (Arg *A = Args.getLastArg(Ids: OPT_module_split_mode_EQ)) {
940 std::optional<IRSplitMode> ModeOrNone =
941 convertStringToSplitMode(S: A->getValue());
942 if (!ModeOrNone)
943 return createStringError(S: formatv(
944 Fmt: "module-split-mode value isn't recognized: {0}", Vals: A->getValue()));
945 SplitMode = *ModeOrNone;
946 }
947
948 // TODO: Expose this as a command-line option and default it to false when
949 // device-image dynamic linking is supported, so that sycl_external functions
950 // can be called across device image boundaries.
951 bool EmitOnlyKernelsAsEntryPoints = true;
952
953 SmallVector<SplitModule, 0> SplitModules;
954 if (canSkipModuleSplit(Mode: SplitMode, M: *Result.LinkedModule,
955 EmitOnlyKernelsAsEntryPoints)) {
956 SplitModules.push_back(Elt: {.ModuleFilePath: SmallString<256>(Result.BitcodeFile),
957 .Symbols: collectEntryPoints(M: *Result.LinkedModule,
958 EmitOnlyKernelsAsEntryPoints)});
959 } else {
960 Expected<SmallVector<SplitModule, 0>> SplitModulesOrErr =
961 splitDeviceCode(M: std::move(Result.LinkedModule), LinkedBitcodeFile: Result.BitcodeFile,
962 Mode: SplitMode, EmitOnlyKernelsAsEntryPoints, Args);
963 if (!SplitModulesOrErr)
964 return SplitModulesOrErr.takeError();
965
966 SplitModules = std::move(*SplitModulesOrErr);
967 }
968
969 bool IsAOTCompileNeeded =
970 StringToOffloadArch(S: Args.getLastArgValue(Id: OPT_arch_EQ)).isIntel();
971
972 StringRef OutputFileNameExt = ".spv";
973
974 // Code generation step.
975 for (size_t I = 0, E = SplitModules.size(); I != E; ++I) {
976 StringRef Stem = OutputFile.rsplit(Separator: '.').first;
977 std::string CodeGenFile = (Stem + "_" + Twine(I) + OutputFileNameExt).str();
978
979 if (Error Err = runCodeGen(File: SplitModules[I].ModuleFilePath,
980 TargetTriple: Result.TargetTriple, Args, OutputFile: CodeGenFile, C))
981 return Err;
982
983 if (!SPIRVDumpDir.empty() && !DryRun) {
984 SmallString<128> DumpFile(SPIRVDumpDir);
985 sys::path::append(path&: DumpFile, a: sys::path::filename(path: CodeGenFile));
986 if (std::error_code EC = sys::fs::copy_file(From: CodeGenFile, To: DumpFile))
987 return createFileError(F: DumpFile, EC);
988 }
989
990 SplitModules[I].ModuleFilePath = CodeGenFile;
991 if (IsAOTCompileNeeded) {
992 std::string AOTFile = (Stem + "_" + Twine(I) + ".out").str();
993 if (Error Err = runAOTCompile(InputFile: CodeGenFile, OutputFile: AOTFile, Args))
994 return Err;
995 SplitModules[I].ModuleFilePath = AOTFile;
996 }
997 }
998
999 // Collect all images to be packed into a single OffloadBinary.
1000 SmallVector<OffloadingImage> Images;
1001 for (SplitModule &SI : SplitModules) {
1002 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> FileOrErr =
1003 DryRun ? llvm::MemoryBuffer::getMemBuffer(InputData: "")
1004 : llvm::MemoryBuffer::getFileOrSTDIN(Filename: SI.ModuleFilePath);
1005 if (!FileOrErr)
1006 return createFileError(F: SI.ModuleFilePath, EC: FileOrErr.getError());
1007
1008 OffloadingImage TheImage{};
1009 TheImage.TheImageKind = IsAOTCompileNeeded ? IMG_Object : IMG_SPIRV;
1010 TheImage.TheOffloadKind = OFK_SYCL;
1011 TheImage.StringData["triple"] =
1012 Args.MakeArgString(Str: Result.TargetTriple.str());
1013 TheImage.StringData["arch"] =
1014 Args.MakeArgString(Str: Args.getLastArgValue(Id: OPT_arch_EQ));
1015 TheImage.StringData["symbols"] = SI.Symbols;
1016 TheImage.Image = std::move(*FileOrErr);
1017 Images.emplace_back(Args: std::move(TheImage));
1018 }
1019
1020 if (Verbose) {
1021 for (const OffloadingImage &Image : Images)
1022 errs() << formatv(
1023 Fmt: "sycl-bundle: image kind: {0}, triple: {1}, arch: {2}\n",
1024 Vals: getImageKindName(Name: Image.TheImageKind),
1025 Vals: Image.StringData.lookup(Key: "triple"), Vals: Image.StringData.lookup(Key: "arch"));
1026 }
1027
1028 llvm::SmallString<0> Buffer = OffloadBinary::write(OffloadingData: Images);
1029 if (Buffer.size() % OffloadBinary::getAlignment() != 0)
1030 return createStringError(Fmt: "offload binary has invalid size alignment");
1031
1032 if (DryRun)
1033 return Error::success();
1034
1035 auto OutputOrErr = FileOutputBuffer::create(FilePath: OutputFile, Size: Buffer.size());
1036 if (!OutputOrErr)
1037 return OutputOrErr.takeError();
1038 llvm::copy(Range&: Buffer, Out: (*OutputOrErr)->getBufferStart());
1039 return (*OutputOrErr)->commit();
1040}
1041
1042int main(int argc, char **argv) {
1043 InitLLVM X(argc, argv);
1044 InitializeAllTargetInfos();
1045 InitializeAllTargets();
1046 InitializeAllTargetMCs();
1047 InitializeAllAsmParsers();
1048 InitializeAllAsmPrinters();
1049
1050 Executable = argv[0];
1051 sys::PrintStackTraceOnErrorSignal(Argv0: argv[0]);
1052
1053 const OptTable &Tbl = getOptTable();
1054 BumpPtrAllocator Alloc;
1055 StringSaver Saver(Alloc);
1056 auto Args = Tbl.parseArgs(Argc: argc, Argv: argv, Unknown: OPT_UNKNOWN, Saver, ErrorFn: [](StringRef Err) {
1057 reportError(E: createStringError(S: Err));
1058 });
1059
1060 if (Args.hasArg(Ids: OPT_help) || Args.hasArg(Ids: OPT_help_hidden)) {
1061 Tbl.printHelp(
1062 OS&: outs(), Usage: "clang-sycl-linker [options] <input bitcode files>",
1063 Title: "A utility that wraps around the SYCL device code linking process.\n"
1064 "This enables LLVM IR linking, post-linking and code generation for "
1065 "SPIR-V JIT and AOT targets.",
1066 ShowHidden: Args.hasArg(Ids: OPT_help_hidden), ShowAllAliases: Args.hasArg(Ids: OPT_help_hidden));
1067 return EXIT_SUCCESS;
1068 }
1069
1070 if (Args.hasArg(Ids: OPT_version)) {
1071 printVersion(OS&: outs());
1072 return EXIT_SUCCESS;
1073 }
1074
1075 Verbose = Args.hasArg(Ids: OPT_verbose);
1076 DryRun = Args.hasArg(Ids: OPT_dry_run);
1077
1078 if (!Args.hasArg(Ids: OPT_o))
1079 reportError(E: createStringError(Fmt: "output file must be specified"));
1080 OutputFile = Args.getLastArgValue(Id: OPT_o);
1081
1082 // Get the input buffers to pass to the linking stage.
1083 auto ResolvedInputsOrErr = getInput(Args);
1084 if (!ResolvedInputsOrErr)
1085 reportError(E: ResolvedInputsOrErr.takeError());
1086
1087 if (auto *A = Args.getLastArg(Ids: OPT_spirv_dump_device_code_EQ)) {
1088 StringRef V = A->getValue();
1089 if (V.empty())
1090 reportError(E: createStringError(
1091 EC: std::make_error_code(e: std::errc::invalid_argument),
1092 S: "--spirv-dump-device-code= requires a non-empty path"));
1093 SPIRVDumpDir = V;
1094 // The directory is shared across all split modules, which use the
1095 // "<output-stem>_<index>.spv" naming scheme. Concurrent invocations
1096 // sharing a dump dir may overwrite each other's files.
1097 if (!DryRun)
1098 if (std::error_code EC = sys::fs::create_directories(path: SPIRVDumpDir))
1099 reportError(E: createStringError(
1100 EC, S: "cannot create SPIR-V dump directory '" + SPIRVDumpDir + "'"));
1101 }
1102
1103 // Run SYCL linking process on the generated inputs.
1104 if (Error Err = runSYCLLink(Inputs: ResolvedInputsOrErr->Buffers,
1105 TargetTriple: ResolvedInputsOrErr->TargetTriple,
1106 TripleSource: ResolvedInputsOrErr->TripleSource, Args))
1107 reportError(E: std::move(Err));
1108
1109 // Remove the temporary files created.
1110 if (!Args.hasArg(Ids: OPT_save_temps) && !DryRun)
1111 for (const auto &TempFile : TempFiles)
1112 if (std::error_code EC = sys::fs::remove(path: TempFile))
1113 reportError(E: createFileError(F: TempFile, EC));
1114
1115 return EXIT_SUCCESS;
1116}
1117