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 // getAllArgValues returns a temporary vector; retain it so the StringRefs
729 // remain valid through the executeCommands call below.
730 std::vector<std::string> ExtraArgsStorage =
731 Args.getAllArgValues(Id: OPT_ocloc_options_EQ);
732 llvm::append_range(C&: CmdArgs, R&: ExtraArgsStorage);
733
734 CmdArgs.push_back(Elt: "-output");
735 CmdArgs.push_back(Elt: OutputFile);
736 CmdArgs.push_back(Elt: "-file");
737 CmdArgs.push_back(Elt: InputFile);
738 if (Error Err = executeCommands(ExecutablePath: *OclocPath, Args: CmdArgs))
739 return Err;
740 return Error::success();
741}
742
743/// Run AOT compilation for Intel CPU/GPU.
744/// \param InputFile The input SPIR-V file.
745/// \param OutputFile The output file name.
746/// \param Args Encompasses all arguments required for linking and wrapping
747/// device code and will be parsed to generate options required to be passed
748/// into the AOT compilation step.
749static Error runAOTCompile(StringRef InputFile, StringRef OutputFile,
750 const ArgList &Args) {
751 StringRef Arch = Args.getLastArgValue(Id: OPT_arch_EQ);
752 OffloadArch OA = StringToOffloadArch(S: Arch);
753 if (OA.isIntelGPU())
754 return runAOTCompileIntelGPU(InputFile, OutputFile, Args);
755 if (OA.isIntelCPU())
756 return runAOTCompileIntelCPU(InputFile, OutputFile, Args);
757
758 llvm_unreachable("runAOTCompile dispatched on unsupported arch");
759}
760
761static constexpr char AttrSYCLModuleId[] = "sycl-module-id";
762
763namespace {
764/// SYCL device code module split mode.
765enum class IRSplitMode {
766 SPLIT_PER_TU, // one module per translation unit
767 SPLIT_PER_KERNEL, // one module per kernel
768 SPLIT_NONE // no splitting
769};
770} // namespace
771
772/// Parses the value of \p --module-split-mode.
773static std::optional<IRSplitMode> convertStringToSplitMode(StringRef S) {
774 return StringSwitch<std::optional<IRSplitMode>>(S)
775 .Case(S: "translation_unit", Value: IRSplitMode::SPLIT_PER_TU)
776 .Case(S: "kernel", Value: IRSplitMode::SPLIT_PER_KERNEL)
777 .Case(S: "link_unit", Value: IRSplitMode::SPLIT_NONE)
778 .Default(Value: std::nullopt);
779}
780
781static StringRef splitModeToString(IRSplitMode Mode) {
782 switch (Mode) {
783 case IRSplitMode::SPLIT_PER_TU:
784 return "translation_unit";
785 case IRSplitMode::SPLIT_PER_KERNEL:
786 return "kernel";
787 case IRSplitMode::SPLIT_NONE:
788 return "link_unit";
789 }
790 llvm_unreachable("bad split mode");
791}
792
793namespace {
794/// Result of splitting a device module: the bitcode file path and the
795/// serialized symbol table for each device image.
796struct SplitModule {
797 SmallString<256> ModuleFilePath;
798 SmallString<0> Symbols;
799};
800} // namespace
801
802static bool isEntryPoint(const Function &F, bool EmitOnlyKernelsAsEntryPoints) {
803 if (F.isDeclaration())
804 return false;
805 if (F.hasKernelCallingConv())
806 return true;
807 if (EmitOnlyKernelsAsEntryPoints)
808 return false;
809 // sycl_external functions carry the "sycl-module-id" attribute.
810 return F.hasFnAttribute(Kind: AttrSYCLModuleId);
811}
812
813/// Collect entry point names from \p M and serialize them into a symbol table.
814static SmallString<0> collectEntryPoints(const Module &M,
815 bool EmitOnlyKernelsAsEntryPoints) {
816 SmallVector<StringRef> Names;
817 for (const Function &F : M)
818 if (isEntryPoint(F, EmitOnlyKernelsAsEntryPoints))
819 Names.push_back(Elt: F.getName());
820 SmallString<0> SymbolData;
821 llvm::offloading::sycl::writeSymbolTable(Names, Out&: SymbolData);
822 return SymbolData;
823}
824
825namespace {
826/// Functor passed to splitModuleTransitiveFromEntryPoints. For each input
827/// function \p F, returns a numeric group ID (if \p F is an entry point)
828/// determining which device image it lands in, or std::nullopt (for
829/// non-entry-points). SPLIT_PER_KERNEL \p Mode gives each kernel its own ID;
830/// SPLIT_PER_TU \p Mode groups kernels by their "sycl-module-id" attribute
831/// value.
832class EntryPointCategorizer {
833public:
834 EntryPointCategorizer(IRSplitMode Mode, bool EmitOnlyKernelsAsEntryPoints)
835 : Mode(Mode), OnlyKernelsAreEntryPoints(EmitOnlyKernelsAsEntryPoints) {}
836
837 std::optional<int> operator()(const Function &F) {
838 if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints: OnlyKernelsAreEntryPoints))
839 return std::nullopt;
840
841 std::string Key;
842 switch (Mode) {
843 case IRSplitMode::SPLIT_PER_KERNEL:
844 Key = F.getName().str();
845 break;
846 case IRSplitMode::SPLIT_PER_TU:
847 Key = F.getFnAttribute(Kind: AttrSYCLModuleId).getValueAsString().str();
848 break;
849 case IRSplitMode::SPLIT_NONE:
850 llvm_unreachable("categorizer cannot be used for SPLIT_NONE");
851 }
852
853 auto [It, Inserted] =
854 StrToId.try_emplace(Key: std::move(Key), Args: static_cast<int>(StrToId.size()));
855 return It->second;
856 }
857
858private:
859 IRSplitMode Mode;
860 bool OnlyKernelsAreEntryPoints;
861 llvm::StringMap<int> StrToId;
862};
863} // namespace
864
865/// Splits the fully linked device \p M into one bitcode file per device image
866/// according to \p Mode and returns the list of split images with their symbol
867/// tables. The module is split transitively from entry points; each part is
868/// written to a fresh temporary bitcode file.
869static Expected<SmallVector<SplitModule, 0>>
870splitDeviceCode(std::unique_ptr<Module> M, StringRef LinkedBitcodeFile,
871 IRSplitMode Mode, bool EmitOnlyKernelsAsEntryPoints,
872 const ArgList &Args) {
873 assert(Mode != IRSplitMode::SPLIT_NONE && "SPLIT_NONE is unsupported");
874
875 SmallVector<SplitModule, 0> SplitModules;
876 EntryPointCategorizer Categorizer(Mode, EmitOnlyKernelsAsEntryPoints);
877
878 auto SplitCallback = [&](std::unique_ptr<Module> Part) -> Error {
879 Expected<StringRef> BitcodeFileOrErr =
880 createTempFile(Args, Prefix: sys::path::filename(path: OutputFile), Extension: "bc");
881 if (!BitcodeFileOrErr)
882 return BitcodeFileOrErr.takeError();
883
884 if (!DryRun) {
885 int FD = -1;
886 if (std::error_code EC = sys::fs::openFileForWrite(Name: *BitcodeFileOrErr, ResultFD&: FD))
887 return errorCodeToError(EC);
888 raw_fd_ostream OS(FD, /*shouldClose=*/true);
889 WriteBitcodeToFile(M: *Part, Out&: OS);
890 }
891
892 SplitModules.push_back(
893 Elt: {.ModuleFilePath: SmallString<256>(*BitcodeFileOrErr),
894 .Symbols: collectEntryPoints(M: *Part, EmitOnlyKernelsAsEntryPoints)});
895 return Error::success();
896 };
897
898 if (Error Err = splitModuleTransitiveFromEntryPoints(
899 M: std::move(M), EntryPointCategorizer: Categorizer, Callback: SplitCallback))
900 return Err;
901
902 if (Verbose) {
903 errs() << formatv(Fmt: "sycl-module-split: input: {0}, mode: {1}\n",
904 Vals&: LinkedBitcodeFile, Vals: splitModeToString(Mode));
905 for (const SplitModule &SI : SplitModules) {
906 errs() << formatv(Fmt: "{0} [", Vals: SI.ModuleFilePath);
907 llvm::offloading::sycl::forEachSymbol(
908 Symbols: SI.Symbols, Callback: [](StringRef Name) { errs() << Name << " "; });
909 errs() << "]\n";
910 }
911 }
912
913 return SplitModules;
914}
915
916/// Returns true if module splitting can be skipped: either \p Mode is
917/// SPLIT_NONE, or \p M contains no entry points (nothing to split from).
918static bool canSkipModuleSplit(IRSplitMode Mode, const Module &M,
919 bool EmitOnlyKernelsAsEntryPoints) {
920 if (Mode == IRSplitMode::SPLIT_NONE)
921 return true;
922 return llvm::none_of(Range: M.functions(), P: [&](const Function &F) {
923 return isEntryPoint(F, EmitOnlyKernelsAsEntryPoints);
924 });
925}
926
927/// Performs the following steps:
928/// 1. Link all input bitcode files together with library files.
929/// 2. Optionally split the linked module according to the requested
930/// IRSplitMode.
931/// 3. Run SPIR-V code generation on each (split) module.
932/// 4. Optionally run AOT compilation when targeting an Intel HW arch.
933/// 5. Pack the resulting images into a single OffloadBinary written to the
934/// output file.
935static Error runSYCLLink(ArrayRef<std::unique_ptr<MemoryBuffer>> Inputs,
936 const llvm::Triple &TargetTriple,
937 StringRef TripleSource, const ArgList &Args) {
938 llvm::TimeTraceScope TimeScope("SYCL linking");
939
940 LLVMContext C;
941
942 // Link all input bitcode files and library files.
943 Expected<LinkResult> LinkedOrErr =
944 linkInputs(Inputs, TargetTriple, TripleSource, Args, C);
945 if (!LinkedOrErr)
946 return LinkedOrErr.takeError();
947 LinkResult &Result = *LinkedOrErr;
948
949 // Determine the requested module split mode.
950 IRSplitMode SplitMode = IRSplitMode::SPLIT_PER_TU;
951 if (Arg *A = Args.getLastArg(Ids: OPT_module_split_mode_EQ)) {
952 std::optional<IRSplitMode> ModeOrNone =
953 convertStringToSplitMode(S: A->getValue());
954 if (!ModeOrNone)
955 return createStringError(S: formatv(
956 Fmt: "module-split-mode value isn't recognized: {0}", Vals: A->getValue()));
957 SplitMode = *ModeOrNone;
958 }
959
960 // TODO: Expose this as a command-line option and default it to false when
961 // device-image dynamic linking is supported, so that sycl_external functions
962 // can be called across device image boundaries.
963 bool EmitOnlyKernelsAsEntryPoints = true;
964
965 SmallVector<SplitModule, 0> SplitModules;
966 if (canSkipModuleSplit(Mode: SplitMode, M: *Result.LinkedModule,
967 EmitOnlyKernelsAsEntryPoints)) {
968 SplitModules.push_back(Elt: {.ModuleFilePath: SmallString<256>(Result.BitcodeFile),
969 .Symbols: collectEntryPoints(M: *Result.LinkedModule,
970 EmitOnlyKernelsAsEntryPoints)});
971 } else {
972 Expected<SmallVector<SplitModule, 0>> SplitModulesOrErr =
973 splitDeviceCode(M: std::move(Result.LinkedModule), LinkedBitcodeFile: Result.BitcodeFile,
974 Mode: SplitMode, EmitOnlyKernelsAsEntryPoints, Args);
975 if (!SplitModulesOrErr)
976 return SplitModulesOrErr.takeError();
977
978 SplitModules = std::move(*SplitModulesOrErr);
979 }
980
981 bool IsAOTCompileNeeded =
982 StringToOffloadArch(S: Args.getLastArgValue(Id: OPT_arch_EQ)).isIntel();
983
984 StringRef OutputFileNameExt = ".spv";
985
986 // Code generation step.
987 for (size_t I = 0, E = SplitModules.size(); I != E; ++I) {
988 StringRef Stem = OutputFile.rsplit(Separator: '.').first;
989 std::string CodeGenFile = (Stem + "_" + Twine(I) + OutputFileNameExt).str();
990
991 if (Error Err = runCodeGen(File: SplitModules[I].ModuleFilePath,
992 TargetTriple: Result.TargetTriple, Args, OutputFile: CodeGenFile, C))
993 return Err;
994
995 if (!SPIRVDumpDir.empty() && !DryRun) {
996 SmallString<128> DumpFile(SPIRVDumpDir);
997 sys::path::append(path&: DumpFile, a: sys::path::filename(path: CodeGenFile));
998 if (std::error_code EC = sys::fs::copy_file(From: CodeGenFile, To: DumpFile))
999 return createFileError(F: DumpFile, EC);
1000 }
1001
1002 SplitModules[I].ModuleFilePath = CodeGenFile;
1003 if (IsAOTCompileNeeded) {
1004 std::string AOTFile = (Stem + "_" + Twine(I) + ".out").str();
1005 if (Error Err = runAOTCompile(InputFile: CodeGenFile, OutputFile: AOTFile, Args))
1006 return Err;
1007 SplitModules[I].ModuleFilePath = AOTFile;
1008 }
1009 }
1010
1011 // Collect all images to be packed into a single OffloadBinary.
1012 SmallVector<OffloadingImage> Images;
1013 for (SplitModule &SI : SplitModules) {
1014 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> FileOrErr =
1015 DryRun ? llvm::MemoryBuffer::getMemBuffer(InputData: "")
1016 : llvm::MemoryBuffer::getFileOrSTDIN(Filename: SI.ModuleFilePath);
1017 if (!FileOrErr)
1018 return createFileError(F: SI.ModuleFilePath, EC: FileOrErr.getError());
1019
1020 OffloadingImage TheImage{};
1021 TheImage.TheImageKind = IsAOTCompileNeeded ? IMG_Object : IMG_SPIRV;
1022 TheImage.TheOffloadKind = OFK_SYCL;
1023 TheImage.StringData["triple"] =
1024 Args.MakeArgString(Str: Result.TargetTriple.str());
1025 TheImage.StringData["arch"] =
1026 Args.MakeArgString(Str: Args.getLastArgValue(Id: OPT_arch_EQ));
1027 TheImage.StringData["symbols"] = SI.Symbols;
1028 TheImage.Image = std::move(*FileOrErr);
1029 Images.emplace_back(Args: std::move(TheImage));
1030 }
1031
1032 if (Verbose) {
1033 for (const OffloadingImage &Image : Images)
1034 errs() << formatv(
1035 Fmt: "sycl-bundle: image kind: {0}, triple: {1}, arch: {2}\n",
1036 Vals: getImageKindName(Name: Image.TheImageKind),
1037 Vals: Image.StringData.lookup(Key: "triple"), Vals: Image.StringData.lookup(Key: "arch"));
1038 }
1039
1040 llvm::SmallString<0> Buffer = OffloadBinary::write(OffloadingData: Images);
1041 if (Buffer.size() % OffloadBinary::getAlignment() != 0)
1042 return createStringError(Fmt: "offload binary has invalid size alignment");
1043
1044 if (DryRun)
1045 return Error::success();
1046
1047 auto OutputOrErr = FileOutputBuffer::create(FilePath: OutputFile, Size: Buffer.size());
1048 if (!OutputOrErr)
1049 return OutputOrErr.takeError();
1050 llvm::copy(Range&: Buffer, Out: (*OutputOrErr)->getBufferStart());
1051 return (*OutputOrErr)->commit();
1052}
1053
1054int main(int argc, char **argv) {
1055 InitLLVM X(argc, argv);
1056 InitializeAllTargetInfos();
1057 InitializeAllTargets();
1058 InitializeAllTargetMCs();
1059 InitializeAllAsmParsers();
1060 InitializeAllAsmPrinters();
1061
1062 Executable = argv[0];
1063 sys::PrintStackTraceOnErrorSignal(Argv0: argv[0]);
1064
1065 const OptTable &Tbl = getOptTable();
1066 BumpPtrAllocator Alloc;
1067 StringSaver Saver(Alloc);
1068 auto Args = Tbl.parseArgs(Argc: argc, Argv: argv, Unknown: OPT_UNKNOWN, Saver, ErrorFn: [](StringRef Err) {
1069 reportError(E: createStringError(S: Err));
1070 });
1071
1072 if (Args.hasArg(Ids: OPT_help) || Args.hasArg(Ids: OPT_help_hidden)) {
1073 Tbl.printHelp(
1074 OS&: outs(), Usage: "clang-sycl-linker [options] <input bitcode files>",
1075 Title: "A utility that wraps around the SYCL device code linking process.\n"
1076 "This enables LLVM IR linking, post-linking and code generation for "
1077 "SPIR-V JIT and AOT targets.",
1078 ShowHidden: Args.hasArg(Ids: OPT_help_hidden), ShowAllAliases: Args.hasArg(Ids: OPT_help_hidden));
1079 return EXIT_SUCCESS;
1080 }
1081
1082 if (Args.hasArg(Ids: OPT_version)) {
1083 printVersion(OS&: outs());
1084 return EXIT_SUCCESS;
1085 }
1086
1087 Verbose = Args.hasArg(Ids: OPT_verbose);
1088 DryRun = Args.hasArg(Ids: OPT_dry_run);
1089
1090 if (!Args.hasArg(Ids: OPT_o))
1091 reportError(E: createStringError(Fmt: "output file must be specified"));
1092 OutputFile = Args.getLastArgValue(Id: OPT_o);
1093
1094 // Get the input buffers to pass to the linking stage.
1095 auto ResolvedInputsOrErr = getInput(Args);
1096 if (!ResolvedInputsOrErr)
1097 reportError(E: ResolvedInputsOrErr.takeError());
1098
1099 if (auto *A = Args.getLastArg(Ids: OPT_spirv_dump_device_code_EQ)) {
1100 StringRef V = A->getValue();
1101 if (V.empty())
1102 reportError(E: createStringError(
1103 EC: std::make_error_code(e: std::errc::invalid_argument),
1104 S: "--spirv-dump-device-code= requires a non-empty path"));
1105 SPIRVDumpDir = V;
1106 // The directory is shared across all split modules, which use the
1107 // "<output-stem>_<index>.spv" naming scheme. Concurrent invocations
1108 // sharing a dump dir may overwrite each other's files.
1109 if (!DryRun)
1110 if (std::error_code EC = sys::fs::create_directories(path: SPIRVDumpDir))
1111 reportError(E: createStringError(
1112 EC, S: "cannot create SPIR-V dump directory '" + SPIRVDumpDir + "'"));
1113 }
1114
1115 // Run SYCL linking process on the generated inputs.
1116 if (Error Err = runSYCLLink(Inputs: ResolvedInputsOrErr->Buffers,
1117 TargetTriple: ResolvedInputsOrErr->TargetTriple,
1118 TripleSource: ResolvedInputsOrErr->TripleSource, Args))
1119 reportError(E: std::move(Err));
1120
1121 // Remove the temporary files created.
1122 if (!Args.hasArg(Ids: OPT_save_temps) && !DryRun)
1123 for (const auto &TempFile : TempFiles)
1124 if (std::error_code EC = sys::fs::remove(path: TempFile))
1125 reportError(E: createFileError(F: TempFile, EC));
1126
1127 return EXIT_SUCCESS;
1128}
1129