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