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