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