clang 18.0.0git
AMDGPU.cpp
Go to the documentation of this file.
1//===--- AMDGPU.cpp - AMDGPU ToolChain Implementations ----------*- C++ -*-===//
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#include "AMDGPU.h"
10#include "CommonArgs.h"
12#include "clang/Config/config.h"
17#include "llvm/ADT/StringExtras.h"
18#include "llvm/Option/ArgList.h"
19#include "llvm/Support/Error.h"
20#include "llvm/Support/LineIterator.h"
21#include "llvm/Support/Path.h"
22#include "llvm/Support/Process.h"
23#include "llvm/Support/VirtualFileSystem.h"
24#include "llvm/TargetParser/Host.h"
25#include <optional>
26#include <system_error>
27
28using namespace clang::driver;
29using namespace clang::driver::tools;
30using namespace clang::driver::toolchains;
31using namespace clang;
32using namespace llvm::opt;
33
34// Look for sub-directory starts with PackageName under ROCm candidate path.
35// If there is one and only one matching sub-directory found, append the
36// sub-directory to Path. If there is no matching sub-directory or there are
37// more than one matching sub-directories, diagnose them. Returns the full
38// path of the package if there is only one matching sub-directory, otherwise
39// returns an empty string.
41RocmInstallationDetector::findSPACKPackage(const Candidate &Cand,
42 StringRef PackageName) {
43 if (!Cand.isSPACK())
44 return {};
45 std::error_code EC;
46 std::string Prefix = Twine(PackageName + "-" + Cand.SPACKReleaseStr).str();
48 for (llvm::vfs::directory_iterator File = D.getVFS().dir_begin(Cand.Path, EC),
49 FileEnd;
50 File != FileEnd && !EC; File.increment(EC)) {
51 llvm::StringRef FileName = llvm::sys::path::filename(File->path());
52 if (FileName.startswith(Prefix)) {
53 SubDirs.push_back(FileName);
54 if (SubDirs.size() > 1)
55 break;
56 }
57 }
58 if (SubDirs.size() == 1) {
59 auto PackagePath = Cand.Path;
60 llvm::sys::path::append(PackagePath, SubDirs[0]);
61 return PackagePath;
62 }
63 if (SubDirs.size() == 0 && Verbose) {
64 llvm::errs() << "SPACK package " << Prefix << " not found at " << Cand.Path
65 << '\n';
66 return {};
67 }
68
69 if (SubDirs.size() > 1 && Verbose) {
70 llvm::errs() << "Cannot use SPACK package " << Prefix << " at " << Cand.Path
71 << " due to multiple installations for the same version\n";
72 }
73 return {};
74}
75
76void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) {
77 assert(!Path.empty());
78
79 const StringRef Suffix(".bc");
80 const StringRef Suffix2(".amdgcn.bc");
81
82 std::error_code EC;
83 for (llvm::vfs::directory_iterator LI = D.getVFS().dir_begin(Path, EC), LE;
84 !EC && LI != LE; LI = LI.increment(EC)) {
85 StringRef FilePath = LI->path();
86 StringRef FileName = llvm::sys::path::filename(FilePath);
87 if (!FileName.endswith(Suffix))
88 continue;
89
90 StringRef BaseName;
91 if (FileName.endswith(Suffix2))
92 BaseName = FileName.drop_back(Suffix2.size());
93 else if (FileName.endswith(Suffix))
94 BaseName = FileName.drop_back(Suffix.size());
95
96 const StringRef ABIVersionPrefix = "oclc_abi_version_";
97 if (BaseName == "ocml") {
98 OCML = FilePath;
99 } else if (BaseName == "ockl") {
100 OCKL = FilePath;
101 } else if (BaseName == "opencl") {
102 OpenCL = FilePath;
103 } else if (BaseName == "hip") {
104 HIP = FilePath;
105 } else if (BaseName == "asanrtl") {
106 AsanRTL = FilePath;
107 } else if (BaseName == "oclc_finite_only_off") {
108 FiniteOnly.Off = FilePath;
109 } else if (BaseName == "oclc_finite_only_on") {
110 FiniteOnly.On = FilePath;
111 } else if (BaseName == "oclc_daz_opt_on") {
112 DenormalsAreZero.On = FilePath;
113 } else if (BaseName == "oclc_daz_opt_off") {
114 DenormalsAreZero.Off = FilePath;
115 } else if (BaseName == "oclc_correctly_rounded_sqrt_on") {
116 CorrectlyRoundedSqrt.On = FilePath;
117 } else if (BaseName == "oclc_correctly_rounded_sqrt_off") {
118 CorrectlyRoundedSqrt.Off = FilePath;
119 } else if (BaseName == "oclc_unsafe_math_on") {
120 UnsafeMath.On = FilePath;
121 } else if (BaseName == "oclc_unsafe_math_off") {
122 UnsafeMath.Off = FilePath;
123 } else if (BaseName == "oclc_wavefrontsize64_on") {
124 WavefrontSize64.On = FilePath;
125 } else if (BaseName == "oclc_wavefrontsize64_off") {
126 WavefrontSize64.Off = FilePath;
127 } else if (BaseName.startswith(ABIVersionPrefix)) {
128 unsigned ABIVersionNumber;
129 if (BaseName.drop_front(ABIVersionPrefix.size())
130 .getAsInteger(/*Redex=*/0, ABIVersionNumber))
131 continue;
132 ABIVersionMap[ABIVersionNumber] = FilePath.str();
133 } else {
134 // Process all bitcode filenames that look like
135 // ocl_isa_version_XXX.amdgcn.bc
136 const StringRef DeviceLibPrefix = "oclc_isa_version_";
137 if (!BaseName.startswith(DeviceLibPrefix))
138 continue;
139
140 StringRef IsaVersionNumber =
141 BaseName.drop_front(DeviceLibPrefix.size());
142
143 llvm::Twine GfxName = Twine("gfx") + IsaVersionNumber;
144 SmallString<8> Tmp;
145 LibDeviceMap.insert(
146 std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
147 }
148 }
149}
150
151// Parse and extract version numbers from `.hipVersion`. Return `true` if
152// the parsing fails.
153bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef V) {
154 SmallVector<StringRef, 4> VersionParts;
155 V.split(VersionParts, '\n');
156 unsigned Major = ~0U;
157 unsigned Minor = ~0U;
158 for (auto Part : VersionParts) {
159 auto Splits = Part.rtrim().split('=');
160 if (Splits.first == "HIP_VERSION_MAJOR") {
161 if (Splits.second.getAsInteger(0, Major))
162 return true;
163 } else if (Splits.first == "HIP_VERSION_MINOR") {
164 if (Splits.second.getAsInteger(0, Minor))
165 return true;
166 } else if (Splits.first == "HIP_VERSION_PATCH")
167 VersionPatch = Splits.second.str();
168 }
169 if (Major == ~0U || Minor == ~0U)
170 return true;
171 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
172 DetectedVersion =
173 (Twine(Major) + "." + Twine(Minor) + "." + VersionPatch).str();
174 return false;
175}
176
177/// \returns a list of candidate directories for ROCm installation, which is
178/// cached and populated only once.
180RocmInstallationDetector::getInstallationPathCandidates() {
181
182 // Return the cached candidate list if it has already been populated.
183 if (!ROCmSearchDirs.empty())
184 return ROCmSearchDirs;
185
186 auto DoPrintROCmSearchDirs = [&]() {
187 if (PrintROCmSearchDirs)
188 for (auto Cand : ROCmSearchDirs) {
189 llvm::errs() << "ROCm installation search path";
190 if (Cand.isSPACK())
191 llvm::errs() << " (Spack " << Cand.SPACKReleaseStr << ")";
192 llvm::errs() << ": " << Cand.Path << '\n';
193 }
194 };
195
196 // For candidate specified by --rocm-path we do not do strict check, i.e.,
197 // checking existence of HIP version file and device library files.
198 if (!RocmPathArg.empty()) {
199 ROCmSearchDirs.emplace_back(RocmPathArg.str());
200 DoPrintROCmSearchDirs();
201 return ROCmSearchDirs;
202 } else if (std::optional<std::string> RocmPathEnv =
203 llvm::sys::Process::GetEnv("ROCM_PATH")) {
204 if (!RocmPathEnv->empty()) {
205 ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));
206 DoPrintROCmSearchDirs();
207 return ROCmSearchDirs;
208 }
209 }
210
211 // Try to find relative to the compiler binary.
212 const char *InstallDir = D.getInstalledDir();
213
214 // Check both a normal Unix prefix position of the clang binary, as well as
215 // the Windows-esque layout the ROCm packages use with the host architecture
216 // subdirectory of bin.
217 auto DeduceROCmPath = [](StringRef ClangPath) {
218 // Strip off directory (usually bin)
219 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
220 StringRef ParentName = llvm::sys::path::filename(ParentDir);
221
222 // Some builds use bin/{host arch}, so go up again.
223 if (ParentName == "bin") {
224 ParentDir = llvm::sys::path::parent_path(ParentDir);
225 ParentName = llvm::sys::path::filename(ParentDir);
226 }
227
228 // Detect ROCm packages built with SPACK.
229 // clang is installed at
230 // <rocm_root>/llvm-amdgpu-<rocm_release_string>-<hash>/bin directory.
231 // We only consider the parent directory of llvm-amdgpu package as ROCm
232 // installation candidate for SPACK.
233 if (ParentName.startswith("llvm-amdgpu-")) {
234 auto SPACKPostfix =
235 ParentName.drop_front(strlen("llvm-amdgpu-")).split('-');
236 auto SPACKReleaseStr = SPACKPostfix.first;
237 if (!SPACKReleaseStr.empty()) {
238 ParentDir = llvm::sys::path::parent_path(ParentDir);
239 return Candidate(ParentDir.str(), /*StrictChecking=*/true,
240 SPACKReleaseStr);
241 }
242 }
243
244 // Some versions of the rocm llvm package install to /opt/rocm/llvm/bin
245 // Some versions of the aomp package install to /opt/rocm/aomp/bin
246 if (ParentName == "llvm" || ParentName.startswith("aomp"))
247 ParentDir = llvm::sys::path::parent_path(ParentDir);
248
249 return Candidate(ParentDir.str(), /*StrictChecking=*/true);
250 };
251
252 // Deduce ROCm path by the path used to invoke clang. Do not resolve symbolic
253 // link of clang itself.
254 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
255
256 // Deduce ROCm path by the real path of the invoked clang, resolving symbolic
257 // link of clang itself.
258 llvm::SmallString<256> RealClangPath;
259 llvm::sys::fs::real_path(D.getClangProgramPath(), RealClangPath);
260 auto ParentPath = llvm::sys::path::parent_path(RealClangPath);
261 if (ParentPath != InstallDir)
262 ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));
263
264 // Device library may be installed in clang or resource directory.
265 auto ClangRoot = llvm::sys::path::parent_path(InstallDir);
266 auto RealClangRoot = llvm::sys::path::parent_path(ParentPath);
267 ROCmSearchDirs.emplace_back(ClangRoot.str(), /*StrictChecking=*/true);
268 if (RealClangRoot != ClangRoot)
269 ROCmSearchDirs.emplace_back(RealClangRoot.str(), /*StrictChecking=*/true);
270 ROCmSearchDirs.emplace_back(D.ResourceDir,
271 /*StrictChecking=*/true);
272
273 ROCmSearchDirs.emplace_back(D.SysRoot + "/opt/rocm",
274 /*StrictChecking=*/true);
275
276 // Find the latest /opt/rocm-{release} directory.
277 std::error_code EC;
278 std::string LatestROCm;
279 llvm::VersionTuple LatestVer;
280 // Get ROCm version from ROCm directory name.
281 auto GetROCmVersion = [](StringRef DirName) {
282 llvm::VersionTuple V;
283 std::string VerStr = DirName.drop_front(strlen("rocm-")).str();
284 // The ROCm directory name follows the format of
285 // rocm-{major}.{minor}.{subMinor}[-{build}]
286 std::replace(VerStr.begin(), VerStr.end(), '-', '.');
287 V.tryParse(VerStr);
288 return V;
289 };
290 for (llvm::vfs::directory_iterator
291 File = D.getVFS().dir_begin(D.SysRoot + "/opt", EC),
292 FileEnd;
293 File != FileEnd && !EC; File.increment(EC)) {
294 llvm::StringRef FileName = llvm::sys::path::filename(File->path());
295 if (!FileName.startswith("rocm-"))
296 continue;
297 if (LatestROCm.empty()) {
298 LatestROCm = FileName.str();
299 LatestVer = GetROCmVersion(LatestROCm);
300 continue;
301 }
302 auto Ver = GetROCmVersion(FileName);
303 if (LatestVer < Ver) {
304 LatestROCm = FileName.str();
305 LatestVer = Ver;
306 }
307 }
308 if (!LatestROCm.empty())
309 ROCmSearchDirs.emplace_back(D.SysRoot + "/opt/" + LatestROCm,
310 /*StrictChecking=*/true);
311
312 ROCmSearchDirs.emplace_back(D.SysRoot + "/usr/local",
313 /*StrictChecking=*/true);
314 ROCmSearchDirs.emplace_back(D.SysRoot + "/usr",
315 /*StrictChecking=*/true);
316
317 DoPrintROCmSearchDirs();
318 return ROCmSearchDirs;
319}
320
322 const Driver &D, const llvm::Triple &HostTriple,
323 const llvm::opt::ArgList &Args, bool DetectHIPRuntime, bool DetectDeviceLib)
324 : D(D) {
325 Verbose = Args.hasArg(options::OPT_v);
326 RocmPathArg = Args.getLastArgValue(clang::driver::options::OPT_rocm_path_EQ);
327 PrintROCmSearchDirs =
328 Args.hasArg(clang::driver::options::OPT_print_rocm_search_dirs);
329 RocmDeviceLibPathArg =
330 Args.getAllArgValues(clang::driver::options::OPT_rocm_device_lib_path_EQ);
331 HIPPathArg = Args.getLastArgValue(clang::driver::options::OPT_hip_path_EQ);
332 if (auto *A = Args.getLastArg(clang::driver::options::OPT_hip_version_EQ)) {
333 HIPVersionArg = A->getValue();
334 unsigned Major = ~0U;
335 unsigned Minor = ~0U;
337 HIPVersionArg.split(Parts, '.');
338 if (Parts.size())
339 Parts[0].getAsInteger(0, Major);
340 if (Parts.size() > 1)
341 Parts[1].getAsInteger(0, Minor);
342 if (Parts.size() > 2)
343 VersionPatch = Parts[2].str();
344 if (VersionPatch.empty())
345 VersionPatch = "0";
346 if (Major != ~0U && Minor == ~0U)
347 Minor = 0;
348 if (Major == ~0U || Minor == ~0U)
349 D.Diag(diag::err_drv_invalid_value)
350 << A->getAsString(Args) << HIPVersionArg;
351
352 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
353 DetectedVersion =
354 (Twine(Major) + "." + Twine(Minor) + "." + VersionPatch).str();
355 } else {
356 VersionPatch = DefaultVersionPatch;
357 VersionMajorMinor =
358 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
359 DetectedVersion = (Twine(DefaultVersionMajor) + "." +
360 Twine(DefaultVersionMinor) + "." + VersionPatch)
361 .str();
362 }
363
364 if (DetectHIPRuntime)
366 if (DetectDeviceLib)
368}
369
371 assert(LibDevicePath.empty());
372
373 if (!RocmDeviceLibPathArg.empty())
374 LibDevicePath = RocmDeviceLibPathArg[RocmDeviceLibPathArg.size() - 1];
375 else if (std::optional<std::string> LibPathEnv =
376 llvm::sys::Process::GetEnv("HIP_DEVICE_LIB_PATH"))
377 LibDevicePath = std::move(*LibPathEnv);
378
379 auto &FS = D.getVFS();
380 if (!LibDevicePath.empty()) {
381 // Maintain compatability with HIP flag/envvar pointing directly at the
382 // bitcode library directory. This points directly at the library path instead
383 // of the rocm root installation.
384 if (!FS.exists(LibDevicePath))
385 return;
386
387 scanLibDevicePath(LibDevicePath);
388 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
389 return;
390 }
391
392 // Check device library exists at the given path.
393 auto CheckDeviceLib = [&](StringRef Path, bool StrictChecking) {
394 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
395 if (CheckLibDevice && !FS.exists(Path))
396 return false;
397
398 scanLibDevicePath(Path);
399
400 if (!NoBuiltinLibs) {
401 // Check that the required non-target libraries are all available.
402 if (!allGenericLibsValid())
403 return false;
404
405 // Check that we have found at least one libdevice that we can link in
406 // if -nobuiltinlib hasn't been specified.
407 if (LibDeviceMap.empty())
408 return false;
409 }
410 return true;
411 };
412
413 // Find device libraries in <LLVM_DIR>/lib/clang/<ver>/lib/amdgcn/bitcode
414 LibDevicePath = D.ResourceDir;
415 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
416 "amdgcn", "bitcode");
417 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, true);
418 if (HasDeviceLibrary)
419 return;
420
421 // Find device libraries in a legacy ROCm directory structure
422 // ${ROCM_ROOT}/amdgcn/bitcode/*
423 auto &ROCmDirs = getInstallationPathCandidates();
424 for (const auto &Candidate : ROCmDirs) {
425 LibDevicePath = Candidate.Path;
426 llvm::sys::path::append(LibDevicePath, "amdgcn", "bitcode");
427 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);
428 if (HasDeviceLibrary)
429 return;
430 }
431}
432
434 SmallVector<Candidate, 4> HIPSearchDirs;
435 if (!HIPPathArg.empty())
436 HIPSearchDirs.emplace_back(HIPPathArg.str());
437 else if (std::optional<std::string> HIPPathEnv =
438 llvm::sys::Process::GetEnv("HIP_PATH")) {
439 if (!HIPPathEnv->empty())
440 HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));
441 }
442 if (HIPSearchDirs.empty())
443 HIPSearchDirs.append(getInstallationPathCandidates());
444 auto &FS = D.getVFS();
445
446 for (const auto &Candidate : HIPSearchDirs) {
447 InstallPath = Candidate.Path;
448 if (InstallPath.empty() || !FS.exists(InstallPath))
449 continue;
450 // HIP runtime built by SPACK is installed to
451 // <rocm_root>/hip-<rocm_release_string>-<hash> directory.
452 auto SPACKPath = findSPACKPackage(Candidate, "hip");
453 InstallPath = SPACKPath.empty() ? InstallPath : SPACKPath;
454
455 BinPath = InstallPath;
456 llvm::sys::path::append(BinPath, "bin");
457 IncludePath = InstallPath;
458 llvm::sys::path::append(IncludePath, "include");
459 LibPath = InstallPath;
460 llvm::sys::path::append(LibPath, "lib");
461 SharePath = InstallPath;
462 llvm::sys::path::append(SharePath, "share");
463
464 // Get parent of InstallPath and append "share"
465 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
466 llvm::sys::path::append(ParentSharePath, "share");
467
468 auto Append = [](SmallString<0> &path, const Twine &a, const Twine &b = "",
469 const Twine &c = "", const Twine &d = "") {
470 SmallString<0> newpath = path;
471 llvm::sys::path::append(newpath, a, b, c, d);
472 return newpath;
473 };
474 // If HIP version file can be found and parsed, use HIP version from there.
475 for (const auto &VersionFilePath :
476 {Append(SharePath, "hip", "version"),
477 Append(ParentSharePath, "hip", "version"),
478 Append(BinPath, ".hipVersion")}) {
479 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
480 FS.getBufferForFile(VersionFilePath);
481 if (!VersionFile)
482 continue;
483 if (HIPVersionArg.empty() && VersionFile)
484 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
485 continue;
486
487 HasHIPRuntime = true;
488 return;
489 }
490 // Otherwise, if -rocm-path is specified (no strict checking), use the
491 // default HIP version or specified by --hip-version.
492 if (!Candidate.StrictChecking) {
493 HasHIPRuntime = true;
494 return;
495 }
496 }
497 HasHIPRuntime = false;
498}
499
500void RocmInstallationDetector::print(raw_ostream &OS) const {
501 if (hasHIPRuntime())
502 OS << "Found HIP installation: " << InstallPath << ", version "
503 << DetectedVersion << '\n';
504}
505
506void RocmInstallationDetector::AddHIPIncludeArgs(const ArgList &DriverArgs,
507 ArgStringList &CC1Args) const {
508 bool UsesRuntimeWrapper = VersionMajorMinor > llvm::VersionTuple(3, 5) &&
509 !DriverArgs.hasArg(options::OPT_nohipwrapperinc);
510
511 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
512 // HIP header includes standard library wrapper headers under clang
513 // cuda_wrappers directory. Since these wrapper headers include_next
514 // standard C++ headers, whereas libc++ headers include_next other clang
515 // headers. The include paths have to follow this order:
516 // - wrapper include path
517 // - standard C++ include path
518 // - other clang include path
519 // Since standard C++ and other clang include paths are added in other
520 // places after this function, here we only need to make sure wrapper
521 // include path is added.
522 //
523 // ROCm 3.5 does not fully support the wrapper headers. Therefore it needs
524 // a workaround.
526 if (UsesRuntimeWrapper)
527 llvm::sys::path::append(P, "include", "cuda_wrappers");
528 CC1Args.push_back("-internal-isystem");
529 CC1Args.push_back(DriverArgs.MakeArgString(P));
530 }
531
532 if (DriverArgs.hasArg(options::OPT_nogpuinc))
533 return;
534
535 if (!hasHIPRuntime()) {
536 D.Diag(diag::err_drv_no_hip_runtime);
537 return;
538 }
539
540 CC1Args.push_back("-idirafter");
541 CC1Args.push_back(DriverArgs.MakeArgString(getIncludePath()));
542 if (UsesRuntimeWrapper)
543 CC1Args.append({"-include", "__clang_hip_runtime_wrapper.h"});
544}
545
547 const InputInfo &Output,
548 const InputInfoList &Inputs,
549 const ArgList &Args,
550 const char *LinkingOutput) const {
551
552 std::string Linker = getToolChain().GetProgramPath(getShortName());
553 ArgStringList CmdArgs;
554 CmdArgs.push_back("--no-undefined");
555 CmdArgs.push_back("-shared");
556
557 addLinkerCompressDebugSectionsOption(getToolChain(), Args, CmdArgs);
558 Args.AddAllArgs(CmdArgs, options::OPT_L);
559 AddLinkerInputs(getToolChain(), Inputs, Args, CmdArgs, JA);
560 if (C.getDriver().isUsingLTO())
561 addLTOOptions(getToolChain(), Args, CmdArgs, Output, Inputs[0],
562 C.getDriver().getLTOMode() == LTOK_Thin);
563 else if (Args.hasArg(options::OPT_mcpu_EQ))
564 CmdArgs.push_back(Args.MakeArgString(
565 "-plugin-opt=mcpu=" + Args.getLastArgValue(options::OPT_mcpu_EQ)));
566 CmdArgs.push_back("-o");
567 CmdArgs.push_back(Output.getFilename());
568 C.addCommand(std::make_unique<Command>(
569 JA, *this, ResponseFileSupport::AtFileCurCP(), Args.MakeArgString(Linker),
570 CmdArgs, Inputs, Output));
571}
572
574 const llvm::Triple &Triple,
575 const llvm::opt::ArgList &Args,
576 std::vector<StringRef> &Features) {
577 // Add target ID features to -target-feature options. No diagnostics should
578 // be emitted here since invalid target ID is diagnosed at other places.
579 StringRef TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
580 if (!TargetID.empty()) {
581 llvm::StringMap<bool> FeatureMap;
582 auto OptionalGpuArch = parseTargetID(Triple, TargetID, &FeatureMap);
583 if (OptionalGpuArch) {
584 StringRef GpuArch = *OptionalGpuArch;
585 // Iterate through all possible target ID features for the given GPU.
586 // If it is mapped to true, add +feature.
587 // If it is mapped to false, add -feature.
588 // If it is not in the map (default), do not add it
589 for (auto &&Feature : getAllPossibleTargetIDFeatures(Triple, GpuArch)) {
590 auto Pos = FeatureMap.find(Feature);
591 if (Pos == FeatureMap.end())
592 continue;
593 Features.push_back(Args.MakeArgStringRef(
594 (Twine(Pos->second ? "+" : "-") + Feature).str()));
595 }
596 }
597 }
598
599 if (Args.hasFlag(options::OPT_mwavefrontsize64,
600 options::OPT_mno_wavefrontsize64, false))
601 Features.push_back("+wavefrontsize64");
602
603 handleTargetFeaturesGroup(D, Triple, Args, Features,
604 options::OPT_m_amdgpu_Features_Group);
605}
606
607/// AMDGPU Toolchain
608AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
609 const ArgList &Args)
610 : Generic_ELF(D, Triple, Args),
611 OptionsDefault(
612 {{options::OPT_O, "3"}, {options::OPT_cl_std_EQ, "CL1.2"}}) {
613 // Check code object version options. Emit warnings for legacy options
614 // and errors for the last invalid code object version options.
615 // It is done here to avoid repeated warning or error messages for
616 // each tool invocation.
618}
619
621 return new tools::amdgpu::Linker(*this);
622}
623
624DerivedArgList *
625AMDGPUToolChain::TranslateArgs(const DerivedArgList &Args, StringRef BoundArch,
626 Action::OffloadKind DeviceOffloadKind) const {
627
628 DerivedArgList *DAL =
629 Generic_ELF::TranslateArgs(Args, BoundArch, DeviceOffloadKind);
630
631 const OptTable &Opts = getDriver().getOpts();
632
633 if (!DAL)
634 DAL = new DerivedArgList(Args.getBaseArgs());
635
636 for (Arg *A : Args)
637 DAL->append(A);
638
639 // Replace -mcpu=native with detected GPU.
640 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
641 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) == "native") {
642 DAL->eraseArg(options::OPT_mcpu_EQ);
643 auto GPUsOrErr = getSystemGPUArchs(Args);
644 if (!GPUsOrErr) {
645 getDriver().Diag(diag::err_drv_undetermined_gpu_arch)
646 << llvm::Triple::getArchTypeName(getArch())
647 << llvm::toString(GPUsOrErr.takeError()) << "-mcpu";
648 } else {
649 auto &GPUs = *GPUsOrErr;
650 if (GPUs.size() > 1) {
651 getDriver().Diag(diag::warn_drv_multi_gpu_arch)
652 << llvm::Triple::getArchTypeName(getArch())
653 << llvm::join(GPUs, ", ") << "-mcpu";
654 }
655 DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_mcpu_EQ),
656 Args.MakeArgString(GPUs.front()));
657 }
658 }
659
660 checkTargetID(*DAL);
661
662 if (!Args.getLastArgValue(options::OPT_x).equals("cl"))
663 return DAL;
664
665 // Phase 1 (.cl -> .bc)
666 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
667 DAL->AddFlagArg(nullptr, Opts.getOption(getTriple().isArch64Bit()
668 ? options::OPT_m64
669 : options::OPT_m32));
670
671 // Have to check OPT_O4, OPT_O0 & OPT_Ofast separately
672 // as they defined that way in Options.td
673 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
674 options::OPT_Ofast))
675 DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_O),
676 getOptionDefault(options::OPT_O));
677 }
678
679 return DAL;
680}
681
683 llvm::AMDGPU::GPUKind Kind) {
684
685 // Assume nothing without a specific target.
686 if (Kind == llvm::AMDGPU::GK_NONE)
687 return false;
688
689 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
690
691 // Default to enabling f32 denormals by default on subtargets where fma is
692 // fast with denormals
693 const bool BothDenormAndFMAFast =
694 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
695 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
696 return !BothDenormAndFMAFast;
697}
698
700 const llvm::opt::ArgList &DriverArgs, const JobAction &JA,
701 const llvm::fltSemantics *FPType) const {
702 // Denormals should always be enabled for f16 and f64.
703 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
704 return llvm::DenormalMode::getIEEE();
705
709 auto Kind = llvm::AMDGPU::parseArchAMDGCN(Arch);
710 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
711 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
712 options::OPT_fno_gpu_flush_denormals_to_zero,
714 return llvm::DenormalMode::getPreserveSign();
715
716 return llvm::DenormalMode::getIEEE();
717 }
718
719 const StringRef GpuArch = getGPUArch(DriverArgs);
720 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
721
722 // TODO: There are way too many flags that change this. Do we need to check
723 // them all?
724 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
726
727 // Outputs are flushed to zero (FTZ), preserving sign. Denormal inputs are
728 // also implicit treated as zero (DAZ).
729 return DAZ ? llvm::DenormalMode::getPreserveSign() :
730 llvm::DenormalMode::getIEEE();
731}
732
733bool AMDGPUToolChain::isWave64(const llvm::opt::ArgList &DriverArgs,
734 llvm::AMDGPU::GPUKind Kind) {
735 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
736 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
737
738 return !HasWave32 || DriverArgs.hasFlag(
739 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64, false);
740}
741
742
743/// ROCM Toolchain
744ROCMToolChain::ROCMToolChain(const Driver &D, const llvm::Triple &Triple,
745 const ArgList &Args)
746 : AMDGPUToolChain(D, Triple, Args) {
747 RocmInstallation->detectDeviceLibrary();
748}
749
751 const llvm::opt::ArgList &DriverArgs,
752 llvm::opt::ArgStringList &CC1Args,
753 Action::OffloadKind DeviceOffloadingKind) const {
754 // Default to "hidden" visibility, as object level linking will not be
755 // supported for the foreseeable future.
756 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
757 options::OPT_fvisibility_ms_compat)) {
758 CC1Args.push_back("-fvisibility=hidden");
759 CC1Args.push_back("-fapply-global-visibility-to-externs");
760 }
761}
762
763StringRef
764AMDGPUToolChain::getGPUArch(const llvm::opt::ArgList &DriverArgs) const {
766 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
767}
768
770AMDGPUToolChain::getParsedTargetID(const llvm::opt::ArgList &DriverArgs) const {
771 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
772 if (TargetID.empty())
773 return {std::nullopt, std::nullopt, std::nullopt};
774
775 llvm::StringMap<bool> FeatureMap;
776 auto OptionalGpuArch = parseTargetID(getTriple(), TargetID, &FeatureMap);
777 if (!OptionalGpuArch)
778 return {TargetID.str(), std::nullopt, std::nullopt};
779
780 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
781}
782
784 const llvm::opt::ArgList &DriverArgs) const {
785 auto PTID = getParsedTargetID(DriverArgs);
786 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
787 getDriver().Diag(clang::diag::err_drv_bad_target_id)
788 << *PTID.OptionalTargetID;
789 }
790}
791
793AMDGPUToolChain::getSystemGPUArchs(const ArgList &Args) const {
794 // Detect AMD GPUs availible on the system.
795 std::string Program;
796 if (Arg *A = Args.getLastArg(options::OPT_amdgpu_arch_tool_EQ))
797 Program = A->getValue();
798 else
799 Program = GetProgramPath("amdgpu-arch");
800
801 auto StdoutOrErr = executeToolChainProgram(Program);
802 if (!StdoutOrErr)
803 return StdoutOrErr.takeError();
804
806 for (StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(), "\n"))
807 if (!Arch.empty())
808 GPUArchs.push_back(Arch.str());
809
810 if (GPUArchs.empty())
811 return llvm::createStringError(std::error_code(),
812 "No AMD GPU detected in the system");
813
814 return std::move(GPUArchs);
815}
816
818 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
819 Action::OffloadKind DeviceOffloadingKind) const {
820 AMDGPUToolChain::addClangTargetOptions(DriverArgs, CC1Args,
821 DeviceOffloadingKind);
822
823 // For the OpenCL case where there is no offload target, accept -nostdlib to
824 // disable bitcode linking.
825 if (DeviceOffloadingKind == Action::OFK_None &&
826 DriverArgs.hasArg(options::OPT_nostdlib))
827 return;
828
829 if (DriverArgs.hasArg(options::OPT_nogpulib))
830 return;
831
832 // Get the device name and canonicalize it
833 const StringRef GpuArch = getGPUArch(DriverArgs);
834 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
835 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
836 StringRef LibDeviceFile = RocmInstallation->getLibDeviceFile(CanonArch);
839 if (!RocmInstallation->checkCommonBitcodeLibs(CanonArch, LibDeviceFile,
840 ABIVer))
841 return;
842
843 bool Wave64 = isWave64(DriverArgs, Kind);
844
845 // TODO: There are way too many flags that change this. Do we need to check
846 // them all?
847 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
849 bool FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only);
850
851 bool UnsafeMathOpt =
852 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations);
853 bool FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math);
854 bool CorrectSqrt =
855 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt);
856
857 // Add the OpenCL specific bitcode library.
859 BCLibs.push_back(RocmInstallation->getOpenCLPath().str());
860
861 // Add the generic set of libraries.
862 BCLibs.append(RocmInstallation->getCommonBitcodeLibs(
863 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
864 FastRelaxedMath, CorrectSqrt, ABIVer, false));
865
866 for (StringRef BCFile : BCLibs) {
867 CC1Args.push_back("-mlink-builtin-bitcode");
868 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
869 }
870}
871
873 StringRef GPUArch, StringRef LibDeviceFile,
874 DeviceLibABIVersion ABIVer) const {
875 if (!hasDeviceLibrary()) {
876 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
877 return false;
878 }
879 if (LibDeviceFile.empty()) {
880 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
881 return false;
882 }
883 if (ABIVer.requiresLibrary() && getABIVersionPath(ABIVer).empty()) {
884 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.toString();
885 return false;
886 }
887 return true;
888}
889
892 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, bool Wave64,
893 bool DAZ, bool FiniteOnly, bool UnsafeMathOpt, bool FastRelaxedMath,
894 bool CorrectSqrt, DeviceLibABIVersion ABIVer, bool isOpenMP = false) const {
896
897 auto AddBCLib = [&](StringRef BCFile) { BCLibs.push_back(BCFile.str()); };
898
899 AddBCLib(getOCMLPath());
900 AddBCLib(getOCKLPath());
901 AddBCLib(getDenormalsAreZeroPath(DAZ));
902 AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath));
903 AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath));
904 AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt));
905 AddBCLib(getWavefrontSize64Path(Wave64));
906 AddBCLib(LibDeviceFile);
907 auto ABIVerPath = getABIVersionPath(ABIVer);
908 if (!ABIVerPath.empty())
909 AddBCLib(ABIVerPath);
910
911 return BCLibs;
912}
913
915ROCMToolChain::getCommonDeviceLibNames(const llvm::opt::ArgList &DriverArgs,
916 const std::string &GPUArch,
917 bool isOpenMP) const {
918 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
919 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
920
921 StringRef LibDeviceFile = RocmInstallation->getLibDeviceFile(CanonArch);
924 if (!RocmInstallation->checkCommonBitcodeLibs(CanonArch, LibDeviceFile,
925 ABIVer))
926 return {};
927
928 // If --hip-device-lib is not set, add the default bitcode libraries.
929 // TODO: There are way too many flags that change this. Do we need to check
930 // them all?
931 bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
932 options::OPT_fno_gpu_flush_denormals_to_zero,
934 bool FiniteOnly = DriverArgs.hasFlag(
935 options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only, false);
936 bool UnsafeMathOpt =
937 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
938 options::OPT_fno_unsafe_math_optimizations, false);
939 bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,
940 options::OPT_fno_fast_math, false);
941 bool CorrectSqrt = DriverArgs.hasFlag(
942 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
943 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, true);
944 bool Wave64 = isWave64(DriverArgs, Kind);
945
946 return RocmInstallation->getCommonBitcodeLibs(
947 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
948 FastRelaxedMath, CorrectSqrt, ABIVer, isOpenMP);
949}
#define V(N, I)
Definition: ASTContext.h:3233
StringRef P
static void Append(char *Start, char *End, char *&Buffer, unsigned &BufferSize, unsigned &BufferCapacity)
__device__ __2f16 b
__device__ __2f16 float c
const char * getOffloadingArch() const
Definition: Action.h:211
OffloadKind getOffloadingDeviceKind() const
Definition: Action.h:210
Compilation - A set of tasks to perform for a single driver invocation.
Definition: Compilation.h:45
Driver - Encapsulate logic for constructing compilation processes from a set of gcc-driver-like comma...
Definition: Driver.h:77
std::string SysRoot
sysroot, if present
Definition: Driver.h:183
const char * getClangProgramPath() const
Get the path to the main clang executable.
Definition: Driver.h:412
DiagnosticBuilder Diag(unsigned DiagID) const
Definition: Driver.h:144
const llvm::opt::OptTable & getOpts() const
Definition: Driver.h:388
std::string ResourceDir
The path to the compiler resource directory.
Definition: Driver.h:167
const char * getInstalledDir() const
Get the path to where the clang executable was installed.
Definition: Driver.h:417
llvm::vfs::FileSystem & getVFS() const
Definition: Driver.h:392
InputInfo - Wrapper for information about an input source.
Definition: InputInfo.h:22
const char * getFilename() const
Definition: InputInfo.h:83
StringRef getIncludePath() const
Get the detected path to Rocm's bin directory.
Definition: ROCm.h:196
RocmInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args, bool DetectHIPRuntime=true, bool DetectDeviceLib=false)
Definition: AMDGPU.cpp:321
bool checkCommonBitcodeLibs(StringRef GPUArch, StringRef LibDeviceFile, DeviceLibABIVersion ABIVer) const
Check file paths of default bitcode libraries common to AMDGPU based toolchains.
Definition: AMDGPU.cpp:872
bool hasHIPRuntime() const
Check whether we detected a valid HIP runtime.
Definition: ROCm.h:178
llvm::SmallVector< std::string, 12 > getCommonBitcodeLibs(const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, bool Wave64, bool DAZ, bool FiniteOnly, bool UnsafeMathOpt, bool FastRelaxedMath, bool CorrectSqrt, DeviceLibABIVersion ABIVer, bool isOpenMP) const
Get file paths of default bitcode libraries common to AMDGPU based toolchains.
Definition: AMDGPU.cpp:891
void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const
Definition: AMDGPU.cpp:506
void print(raw_ostream &OS) const
Print information about the detected ROCm installation.
Definition: AMDGPU.cpp:500
llvm::Triple::ArchType getArch() const
Definition: ToolChain.h:261
const Driver & getDriver() const
Definition: ToolChain.h:245
virtual llvm::opt::DerivedArgList * TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch, Action::OffloadKind DeviceOffloadKind) const
TranslateArgs - Create a new derived argument list for any argument translations this ToolChain may w...
Definition: ToolChain.h:348
const llvm::Triple & getTriple() const
Definition: ToolChain.h:247
std::string GetProgramPath(const char *Name) const
Definition: ToolChain.cpp:828
llvm::Expected< std::unique_ptr< llvm::MemoryBuffer > > executeToolChainProgram(StringRef Executable) const
Executes the given Executable and returns the stdout.
Definition: ToolChain.cpp:98
Tool - Information on a specific compilation tool.
Definition: Tool.h:32
llvm::DenormalMode getDefaultDenormalModeForType(const llvm::opt::ArgList &DriverArgs, const JobAction &JA, const llvm::fltSemantics *FPType=nullptr) const override
Returns the output denormal handling type in the default floating point environment for the given FPT...
Definition: AMDGPU.cpp:699
llvm::opt::DerivedArgList * TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch, Action::OffloadKind DeviceOffloadKind) const override
TranslateArgs - Create a new derived argument list for any argument translations this ToolChain may w...
Definition: AMDGPU.cpp:625
static bool getDefaultDenormsAreZeroForTarget(llvm::AMDGPU::GPUKind GPUKind)
Return whether denormals should be flushed, and treated as 0 by default for the subtarget.
Definition: AMDGPU.cpp:682
StringRef getGPUArch(const llvm::opt::ArgList &DriverArgs) const
Get GPU arch from -mcpu without checking.
Definition: AMDGPU.cpp:764
virtual void checkTargetID(const llvm::opt::ArgList &DriverArgs) const
Check and diagnose invalid target ID specified by -mcpu.
Definition: AMDGPU.cpp:783
Tool * buildLinker() const override
Definition: AMDGPU.cpp:620
static bool isWave64(const llvm::opt::ArgList &DriverArgs, llvm::AMDGPU::GPUKind Kind)
Definition: AMDGPU.cpp:733
AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args)
AMDGPU Toolchain.
Definition: AMDGPU.cpp:608
ParsedTargetIDType getParsedTargetID(const llvm::opt::ArgList &DriverArgs) const
Get target ID, GPU arch, and target ID features if the target ID is specified and valid.
Definition: AMDGPU.cpp:770
StringRef getOptionDefault(options::ID OptID) const
Definition: AMDGPU.h:54
void addClangTargetOptions(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args, Action::OffloadKind DeviceOffloadKind) const override
Add options that need to be passed to cc1 for this target.
Definition: AMDGPU.cpp:750
virtual Expected< SmallVector< std::string > > getSystemGPUArchs(const llvm::opt::ArgList &Args) const override
Uses amdgpu-arch tool to get arch of the system GPU.
Definition: AMDGPU.cpp:793
LazyDetector< RocmInstallationDetector > RocmInstallation
Definition: Gnu.h:291
llvm::SmallVector< std::string, 12 > getCommonDeviceLibNames(const llvm::opt::ArgList &DriverArgs, const std::string &GPUArch, bool isOpenMP=false) const
Definition: AMDGPU.cpp:915
void addClangTargetOptions(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args, Action::OffloadKind DeviceOffloadKind) const override
Add options that need to be passed to cc1 for this target.
Definition: AMDGPU.cpp:817
ROCMToolChain(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args)
ROCM Toolchain.
Definition: AMDGPU.cpp:744
void ConstructJob(Compilation &C, const JobAction &JA, const InputInfo &Output, const InputInfoList &Inputs, const llvm::opt::ArgList &TCArgs, const char *LinkingOutput) const override
ConstructJob - Construct jobs to perform the action JA, writing to Output and with Inputs,...
Definition: AMDGPU.cpp:546
void getAMDGPUTargetFeatures(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args, std::vector< StringRef > &Features)
Definition: AMDGPU.cpp:573
void handleTargetFeaturesGroup(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args, std::vector< StringRef > &Features, llvm::opt::OptSpecifier Group)
Iterate Args and convert -mxxx to +xxx and -mno-xxx to -xxx and append it to Features.
void checkAMDGPUCodeObjectVersion(const Driver &D, const llvm::opt::ArgList &Args)
void addLinkerCompressDebugSectionsOption(const ToolChain &TC, const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs)
Definition: CommonArgs.cpp:306
void addLTOOptions(const ToolChain &ToolChain, const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs, const InputInfo &Output, const InputInfo &Input, bool IsThinLTO)
void AddLinkerInputs(const ToolChain &TC, const InputInfoList &Inputs, const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs, const JobAction &JA)
unsigned getAMDGPUCodeObjectVersion(const Driver &D, const llvm::opt::ArgList &Args)
std::optional< llvm::StringRef > parseTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch, llvm::StringMap< bool > *FeatureMap)
Parse a target ID to get processor and feature map.
Definition: TargetID.cpp:105
llvm::StringRef getProcessorFromTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch)
Get processor name from target ID.
Definition: TargetID.cpp:54
@ C
Languages that the frontend can parse and compile.
llvm::SmallVector< llvm::StringRef, 4 > getAllPossibleTargetIDFeatures(const llvm::Triple &T, llvm::StringRef Processor)
Get all feature strings that can be used in target ID for Processor.
Definition: TargetID.cpp:38
ABI version of device library.
Definition: ROCm.h:26
static DeviceLibABIVersion fromCodeObjectVersion(unsigned CodeObjectVersion)
Definition: ROCm.h:29
bool requiresLibrary()
Whether ABI version bc file is requested.
Definition: ROCm.h:38
static constexpr ResponseFileSupport AtFileCurCP()
Definition: Job.h:92
The struct type returned by getParsedTargetID.
Definition: AMDGPU.h:110