clang 19.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.starts_with(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.ends_with(Suffix))
88 continue;
89
90 StringRef BaseName;
91 if (FileName.ends_with(Suffix2))
92 BaseName = FileName.drop_back(Suffix2.size());
93 else if (FileName.ends_with(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.starts_with(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.starts_with(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 StringRef InstallDir = D.Dir;
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.starts_with("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.starts_with("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.starts_with("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 HIPStdParPathArg =
333 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_path_EQ);
334 HasHIPStdParLibrary =
335 !HIPStdParPathArg.empty() && D.getVFS().exists(HIPStdParPathArg +
336 "/hipstdpar_lib.hpp");
337 HIPRocThrustPathArg =
338 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_thrust_path_EQ);
339 HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&
340 D.getVFS().exists(HIPRocThrustPathArg + "/thrust");
341 HIPRocPrimPathArg =
342 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_prim_path_EQ);
343 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
344 D.getVFS().exists(HIPRocPrimPathArg + "/rocprim");
345
346 if (auto *A = Args.getLastArg(clang::driver::options::OPT_hip_version_EQ)) {
347 HIPVersionArg = A->getValue();
348 unsigned Major = ~0U;
349 unsigned Minor = ~0U;
351 HIPVersionArg.split(Parts, '.');
352 if (Parts.size())
353 Parts[0].getAsInteger(0, Major);
354 if (Parts.size() > 1)
355 Parts[1].getAsInteger(0, Minor);
356 if (Parts.size() > 2)
357 VersionPatch = Parts[2].str();
358 if (VersionPatch.empty())
359 VersionPatch = "0";
360 if (Major != ~0U && Minor == ~0U)
361 Minor = 0;
362 if (Major == ~0U || Minor == ~0U)
363 D.Diag(diag::err_drv_invalid_value)
364 << A->getAsString(Args) << HIPVersionArg;
365
366 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
367 DetectedVersion =
368 (Twine(Major) + "." + Twine(Minor) + "." + VersionPatch).str();
369 } else {
370 VersionPatch = DefaultVersionPatch;
371 VersionMajorMinor =
372 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
373 DetectedVersion = (Twine(DefaultVersionMajor) + "." +
374 Twine(DefaultVersionMinor) + "." + VersionPatch)
375 .str();
376 }
377
378 if (DetectHIPRuntime)
380 if (DetectDeviceLib)
382}
383
385 assert(LibDevicePath.empty());
386
387 if (!RocmDeviceLibPathArg.empty())
388 LibDevicePath = RocmDeviceLibPathArg[RocmDeviceLibPathArg.size() - 1];
389 else if (std::optional<std::string> LibPathEnv =
390 llvm::sys::Process::GetEnv("HIP_DEVICE_LIB_PATH"))
391 LibDevicePath = std::move(*LibPathEnv);
392
393 auto &FS = D.getVFS();
394 if (!LibDevicePath.empty()) {
395 // Maintain compatability with HIP flag/envvar pointing directly at the
396 // bitcode library directory. This points directly at the library path instead
397 // of the rocm root installation.
398 if (!FS.exists(LibDevicePath))
399 return;
400
401 scanLibDevicePath(LibDevicePath);
402 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
403 return;
404 }
405
406 // Check device library exists at the given path.
407 auto CheckDeviceLib = [&](StringRef Path, bool StrictChecking) {
408 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
409 if (CheckLibDevice && !FS.exists(Path))
410 return false;
411
412 scanLibDevicePath(Path);
413
414 if (!NoBuiltinLibs) {
415 // Check that the required non-target libraries are all available.
416 if (!allGenericLibsValid())
417 return false;
418
419 // Check that we have found at least one libdevice that we can link in
420 // if -nobuiltinlib hasn't been specified.
421 if (LibDeviceMap.empty())
422 return false;
423 }
424 return true;
425 };
426
427 // Find device libraries in <LLVM_DIR>/lib/clang/<ver>/lib/amdgcn/bitcode
428 LibDevicePath = D.ResourceDir;
429 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
430 "amdgcn", "bitcode");
431 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, true);
432 if (HasDeviceLibrary)
433 return;
434
435 // Find device libraries in a legacy ROCm directory structure
436 // ${ROCM_ROOT}/amdgcn/bitcode/*
437 auto &ROCmDirs = getInstallationPathCandidates();
438 for (const auto &Candidate : ROCmDirs) {
439 LibDevicePath = Candidate.Path;
440 llvm::sys::path::append(LibDevicePath, "amdgcn", "bitcode");
441 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);
442 if (HasDeviceLibrary)
443 return;
444 }
445}
446
448 SmallVector<Candidate, 4> HIPSearchDirs;
449 if (!HIPPathArg.empty())
450 HIPSearchDirs.emplace_back(HIPPathArg.str());
451 else if (std::optional<std::string> HIPPathEnv =
452 llvm::sys::Process::GetEnv("HIP_PATH")) {
453 if (!HIPPathEnv->empty())
454 HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));
455 }
456 if (HIPSearchDirs.empty())
457 HIPSearchDirs.append(getInstallationPathCandidates());
458 auto &FS = D.getVFS();
459
460 for (const auto &Candidate : HIPSearchDirs) {
461 InstallPath = Candidate.Path;
462 if (InstallPath.empty() || !FS.exists(InstallPath))
463 continue;
464 // HIP runtime built by SPACK is installed to
465 // <rocm_root>/hip-<rocm_release_string>-<hash> directory.
466 auto SPACKPath = findSPACKPackage(Candidate, "hip");
467 InstallPath = SPACKPath.empty() ? InstallPath : SPACKPath;
468
469 BinPath = InstallPath;
470 llvm::sys::path::append(BinPath, "bin");
471 IncludePath = InstallPath;
472 llvm::sys::path::append(IncludePath, "include");
473 LibPath = InstallPath;
474 llvm::sys::path::append(LibPath, "lib");
475 SharePath = InstallPath;
476 llvm::sys::path::append(SharePath, "share");
477
478 // Get parent of InstallPath and append "share"
479 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
480 llvm::sys::path::append(ParentSharePath, "share");
481
482 auto Append = [](SmallString<0> &path, const Twine &a, const Twine &b = "",
483 const Twine &c = "", const Twine &d = "") {
484 SmallString<0> newpath = path;
485 llvm::sys::path::append(newpath, a, b, c, d);
486 return newpath;
487 };
488 // If HIP version file can be found and parsed, use HIP version from there.
489 std::vector<SmallString<0>> VersionFilePaths = {
490 Append(SharePath, "hip", "version"),
491 InstallPath != D.SysRoot + "/usr/local"
492 ? Append(ParentSharePath, "hip", "version")
493 : SmallString<0>(),
494 Append(BinPath, ".hipVersion")};
495
496 for (const auto &VersionFilePath : VersionFilePaths) {
497 if (VersionFilePath.empty())
498 continue;
499 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
500 FS.getBufferForFile(VersionFilePath);
501 if (!VersionFile)
502 continue;
503 if (HIPVersionArg.empty() && VersionFile)
504 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
505 continue;
506
507 HasHIPRuntime = true;
508 return;
509 }
510 // Otherwise, if -rocm-path is specified (no strict checking), use the
511 // default HIP version or specified by --hip-version.
512 if (!Candidate.StrictChecking) {
513 HasHIPRuntime = true;
514 return;
515 }
516 }
517 HasHIPRuntime = false;
518}
519
520void RocmInstallationDetector::print(raw_ostream &OS) const {
521 if (hasHIPRuntime())
522 OS << "Found HIP installation: " << InstallPath << ", version "
523 << DetectedVersion << '\n';
524}
525
526void RocmInstallationDetector::AddHIPIncludeArgs(const ArgList &DriverArgs,
527 ArgStringList &CC1Args) const {
528 bool UsesRuntimeWrapper = VersionMajorMinor > llvm::VersionTuple(3, 5) &&
529 !DriverArgs.hasArg(options::OPT_nohipwrapperinc);
530 bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
531
532 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
533 // HIP header includes standard library wrapper headers under clang
534 // cuda_wrappers directory. Since these wrapper headers include_next
535 // standard C++ headers, whereas libc++ headers include_next other clang
536 // headers. The include paths have to follow this order:
537 // - wrapper include path
538 // - standard C++ include path
539 // - other clang include path
540 // Since standard C++ and other clang include paths are added in other
541 // places after this function, here we only need to make sure wrapper
542 // include path is added.
543 //
544 // ROCm 3.5 does not fully support the wrapper headers. Therefore it needs
545 // a workaround.
547 if (UsesRuntimeWrapper)
548 llvm::sys::path::append(P, "include", "cuda_wrappers");
549 CC1Args.push_back("-internal-isystem");
550 CC1Args.push_back(DriverArgs.MakeArgString(P));
551 }
552
553 const auto HandleHipStdPar = [=, &DriverArgs, &CC1Args]() {
554 StringRef Inc = getIncludePath();
555 auto &FS = D.getVFS();
556
557 if (!hasHIPStdParLibrary())
558 if (!HIPStdParPathArg.empty() ||
559 !FS.exists(Inc + "/thrust/system/hip/hipstdpar/hipstdpar_lib.hpp")) {
560 D.Diag(diag::err_drv_no_hipstdpar_lib);
561 return;
562 }
563 if (!HasRocThrustLibrary && !FS.exists(Inc + "/thrust")) {
564 D.Diag(diag::err_drv_no_hipstdpar_thrust_lib);
565 return;
566 }
567 if (!HasRocPrimLibrary && !FS.exists(Inc + "/rocprim")) {
568 D.Diag(diag::err_drv_no_hipstdpar_prim_lib);
569 return;
570 }
571 const char *ThrustPath;
572 if (HasRocThrustLibrary)
573 ThrustPath = DriverArgs.MakeArgString(HIPRocThrustPathArg);
574 else
575 ThrustPath = DriverArgs.MakeArgString(Inc + "/thrust");
576
577 const char *HIPStdParPath;
579 HIPStdParPath = DriverArgs.MakeArgString(HIPStdParPathArg);
580 else
581 HIPStdParPath = DriverArgs.MakeArgString(StringRef(ThrustPath) +
582 "/system/hip/hipstdpar");
583
584 const char *PrimPath;
585 if (HasRocPrimLibrary)
586 PrimPath = DriverArgs.MakeArgString(HIPRocPrimPathArg);
587 else
588 PrimPath = DriverArgs.MakeArgString(getIncludePath() + "/rocprim");
589
590 CC1Args.append({"-idirafter", ThrustPath, "-idirafter", PrimPath,
591 "-idirafter", HIPStdParPath, "-include",
592 "hipstdpar_lib.hpp"});
593 };
594
595 if (DriverArgs.hasArg(options::OPT_nogpuinc)) {
596 if (HasHipStdPar)
597 HandleHipStdPar();
598
599 return;
600 }
601
602 if (!hasHIPRuntime()) {
603 D.Diag(diag::err_drv_no_hip_runtime);
604 return;
605 }
606
607 CC1Args.push_back("-idirafter");
608 CC1Args.push_back(DriverArgs.MakeArgString(getIncludePath()));
609 if (UsesRuntimeWrapper)
610 CC1Args.append({"-include", "__clang_hip_runtime_wrapper.h"});
611 if (HasHipStdPar)
612 HandleHipStdPar();
613}
614
616 const InputInfo &Output,
617 const InputInfoList &Inputs,
618 const ArgList &Args,
619 const char *LinkingOutput) const {
620
621 std::string Linker = getToolChain().GetProgramPath(getShortName());
622 ArgStringList CmdArgs;
623 CmdArgs.push_back("--no-undefined");
624 CmdArgs.push_back("-shared");
625
626 addLinkerCompressDebugSectionsOption(getToolChain(), Args, CmdArgs);
627 Args.AddAllArgs(CmdArgs, options::OPT_L);
628 getToolChain().AddFilePathLibArgs(Args, CmdArgs);
629 AddLinkerInputs(getToolChain(), Inputs, Args, CmdArgs, JA);
630 if (C.getDriver().isUsingLTO())
631 addLTOOptions(getToolChain(), Args, CmdArgs, Output, Inputs[0],
632 C.getDriver().getLTOMode() == LTOK_Thin);
633 else if (Args.hasArg(options::OPT_mcpu_EQ))
634 CmdArgs.push_back(Args.MakeArgString(
635 "-plugin-opt=mcpu=" + Args.getLastArgValue(options::OPT_mcpu_EQ)));
636 CmdArgs.push_back("-o");
637 CmdArgs.push_back(Output.getFilename());
638 C.addCommand(std::make_unique<Command>(
639 JA, *this, ResponseFileSupport::AtFileCurCP(), Args.MakeArgString(Linker),
640 CmdArgs, Inputs, Output));
641}
642
644 const llvm::Triple &Triple,
645 const llvm::opt::ArgList &Args,
646 std::vector<StringRef> &Features) {
647 // Add target ID features to -target-feature options. No diagnostics should
648 // be emitted here since invalid target ID is diagnosed at other places.
649 StringRef TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
650 if (!TargetID.empty()) {
651 llvm::StringMap<bool> FeatureMap;
652 auto OptionalGpuArch = parseTargetID(Triple, TargetID, &FeatureMap);
653 if (OptionalGpuArch) {
654 StringRef GpuArch = *OptionalGpuArch;
655 // Iterate through all possible target ID features for the given GPU.
656 // If it is mapped to true, add +feature.
657 // If it is mapped to false, add -feature.
658 // If it is not in the map (default), do not add it
659 for (auto &&Feature : getAllPossibleTargetIDFeatures(Triple, GpuArch)) {
660 auto Pos = FeatureMap.find(Feature);
661 if (Pos == FeatureMap.end())
662 continue;
663 Features.push_back(Args.MakeArgStringRef(
664 (Twine(Pos->second ? "+" : "-") + Feature).str()));
665 }
666 }
667 }
668
669 if (Args.hasFlag(options::OPT_mwavefrontsize64,
670 options::OPT_mno_wavefrontsize64, false))
671 Features.push_back("+wavefrontsize64");
672
673 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
674 options::OPT_mno_amdgpu_precise_memory_op, false))
675 Features.push_back("+precise-memory");
676
677 handleTargetFeaturesGroup(D, Triple, Args, Features,
678 options::OPT_m_amdgpu_Features_Group);
679}
680
681/// AMDGPU Toolchain
682AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
683 const ArgList &Args)
684 : Generic_ELF(D, Triple, Args),
685 OptionsDefault(
686 {{options::OPT_O, "3"}, {options::OPT_cl_std_EQ, "CL1.2"}}) {
687 // Check code object version options. Emit warnings for legacy options
688 // and errors for the last invalid code object version options.
689 // It is done here to avoid repeated warning or error messages for
690 // each tool invocation.
692}
693
695 return new tools::amdgpu::Linker(*this);
696}
697
698DerivedArgList *
699AMDGPUToolChain::TranslateArgs(const DerivedArgList &Args, StringRef BoundArch,
700 Action::OffloadKind DeviceOffloadKind) const {
701
702 DerivedArgList *DAL =
703 Generic_ELF::TranslateArgs(Args, BoundArch, DeviceOffloadKind);
704
705 const OptTable &Opts = getDriver().getOpts();
706
707 if (!DAL)
708 DAL = new DerivedArgList(Args.getBaseArgs());
709
710 for (Arg *A : Args)
711 DAL->append(A);
712
713 // Replace -mcpu=native with detected GPU.
714 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
715 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) == "native") {
716 DAL->eraseArg(options::OPT_mcpu_EQ);
717 auto GPUsOrErr = getSystemGPUArchs(Args);
718 if (!GPUsOrErr) {
719 getDriver().Diag(diag::err_drv_undetermined_gpu_arch)
720 << llvm::Triple::getArchTypeName(getArch())
721 << llvm::toString(GPUsOrErr.takeError()) << "-mcpu";
722 } else {
723 auto &GPUs = *GPUsOrErr;
724 if (GPUs.size() > 1) {
725 getDriver().Diag(diag::warn_drv_multi_gpu_arch)
726 << llvm::Triple::getArchTypeName(getArch())
727 << llvm::join(GPUs, ", ") << "-mcpu";
728 }
729 DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_mcpu_EQ),
730 Args.MakeArgString(GPUs.front()));
731 }
732 }
733
734 checkTargetID(*DAL);
735
736 if (!Args.getLastArgValue(options::OPT_x).equals("cl"))
737 return DAL;
738
739 // Phase 1 (.cl -> .bc)
740 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
741 DAL->AddFlagArg(nullptr, Opts.getOption(getTriple().isArch64Bit()
742 ? options::OPT_m64
743 : options::OPT_m32));
744
745 // Have to check OPT_O4, OPT_O0 & OPT_Ofast separately
746 // as they defined that way in Options.td
747 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
748 options::OPT_Ofast))
749 DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_O),
750 getOptionDefault(options::OPT_O));
751 }
752
753 return DAL;
754}
755
757 llvm::AMDGPU::GPUKind Kind) {
758
759 // Assume nothing without a specific target.
760 if (Kind == llvm::AMDGPU::GK_NONE)
761 return false;
762
763 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
764
765 // Default to enabling f32 denormals by default on subtargets where fma is
766 // fast with denormals
767 const bool BothDenormAndFMAFast =
768 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
769 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
770 return !BothDenormAndFMAFast;
771}
772
774 const llvm::opt::ArgList &DriverArgs, const JobAction &JA,
775 const llvm::fltSemantics *FPType) const {
776 // Denormals should always be enabled for f16 and f64.
777 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
778 return llvm::DenormalMode::getIEEE();
779
783 auto Kind = llvm::AMDGPU::parseArchAMDGCN(Arch);
784 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
785 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
786 options::OPT_fno_gpu_flush_denormals_to_zero,
788 return llvm::DenormalMode::getPreserveSign();
789
790 return llvm::DenormalMode::getIEEE();
791 }
792
793 const StringRef GpuArch = getGPUArch(DriverArgs);
794 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
795
796 // TODO: There are way too many flags that change this. Do we need to check
797 // them all?
798 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
800
801 // Outputs are flushed to zero (FTZ), preserving sign. Denormal inputs are
802 // also implicit treated as zero (DAZ).
803 return DAZ ? llvm::DenormalMode::getPreserveSign() :
804 llvm::DenormalMode::getIEEE();
805}
806
807bool AMDGPUToolChain::isWave64(const llvm::opt::ArgList &DriverArgs,
808 llvm::AMDGPU::GPUKind Kind) {
809 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
810 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
811
812 return !HasWave32 || DriverArgs.hasFlag(
813 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64, false);
814}
815
816
817/// ROCM Toolchain
818ROCMToolChain::ROCMToolChain(const Driver &D, const llvm::Triple &Triple,
819 const ArgList &Args)
820 : AMDGPUToolChain(D, Triple, Args) {
821 RocmInstallation->detectDeviceLibrary();
822}
823
825 const llvm::opt::ArgList &DriverArgs,
826 llvm::opt::ArgStringList &CC1Args,
827 Action::OffloadKind DeviceOffloadingKind) const {
828 // Default to "hidden" visibility, as object level linking will not be
829 // supported for the foreseeable future.
830 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
831 options::OPT_fvisibility_ms_compat)) {
832 CC1Args.push_back("-fvisibility=hidden");
833 CC1Args.push_back("-fapply-global-visibility-to-externs");
834 }
835}
836
837void AMDGPUToolChain::addClangWarningOptions(ArgStringList &CC1Args) const {
838 // AMDGPU does not support atomic lib call. Treat atomic alignment
839 // warnings as errors.
840 CC1Args.push_back("-Werror=atomic-alignment");
841}
842
843StringRef
844AMDGPUToolChain::getGPUArch(const llvm::opt::ArgList &DriverArgs) const {
846 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
847}
848
850AMDGPUToolChain::getParsedTargetID(const llvm::opt::ArgList &DriverArgs) const {
851 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
852 if (TargetID.empty())
853 return {std::nullopt, std::nullopt, std::nullopt};
854
855 llvm::StringMap<bool> FeatureMap;
856 auto OptionalGpuArch = parseTargetID(getTriple(), TargetID, &FeatureMap);
857 if (!OptionalGpuArch)
858 return {TargetID.str(), std::nullopt, std::nullopt};
859
860 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
861}
862
864 const llvm::opt::ArgList &DriverArgs) const {
865 auto PTID = getParsedTargetID(DriverArgs);
866 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
867 getDriver().Diag(clang::diag::err_drv_bad_target_id)
868 << *PTID.OptionalTargetID;
869 }
870}
871
873AMDGPUToolChain::getSystemGPUArchs(const ArgList &Args) const {
874 // Detect AMD GPUs availible on the system.
875 std::string Program;
876 if (Arg *A = Args.getLastArg(options::OPT_amdgpu_arch_tool_EQ))
877 Program = A->getValue();
878 else
879 Program = GetProgramPath("amdgpu-arch");
880
881 auto StdoutOrErr = executeToolChainProgram(Program);
882 if (!StdoutOrErr)
883 return StdoutOrErr.takeError();
884
886 for (StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(), "\n"))
887 if (!Arch.empty())
888 GPUArchs.push_back(Arch.str());
889
890 if (GPUArchs.empty())
891 return llvm::createStringError(std::error_code(),
892 "No AMD GPU detected in the system");
893
894 return std::move(GPUArchs);
895}
896
898 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
899 Action::OffloadKind DeviceOffloadingKind) const {
900 AMDGPUToolChain::addClangTargetOptions(DriverArgs, CC1Args,
901 DeviceOffloadingKind);
902
903 // For the OpenCL case where there is no offload target, accept -nostdlib to
904 // disable bitcode linking.
905 if (DeviceOffloadingKind == Action::OFK_None &&
906 DriverArgs.hasArg(options::OPT_nostdlib))
907 return;
908
909 if (DriverArgs.hasArg(options::OPT_nogpulib))
910 return;
911
912 // Get the device name and canonicalize it
913 const StringRef GpuArch = getGPUArch(DriverArgs);
914 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
915 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
916 StringRef LibDeviceFile = RocmInstallation->getLibDeviceFile(CanonArch);
919 if (!RocmInstallation->checkCommonBitcodeLibs(CanonArch, LibDeviceFile,
920 ABIVer))
921 return;
922
923 bool Wave64 = isWave64(DriverArgs, Kind);
924
925 // TODO: There are way too many flags that change this. Do we need to check
926 // them all?
927 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
929 bool FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only);
930
931 bool UnsafeMathOpt =
932 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations);
933 bool FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math);
934 bool CorrectSqrt =
935 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt);
936
937 // Add the OpenCL specific bitcode library.
939 BCLibs.push_back(RocmInstallation->getOpenCLPath().str());
940
941 // Add the generic set of libraries.
942 BCLibs.append(RocmInstallation->getCommonBitcodeLibs(
943 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
944 FastRelaxedMath, CorrectSqrt, ABIVer, false));
945
946 for (StringRef BCFile : BCLibs) {
947 CC1Args.push_back("-mlink-builtin-bitcode");
948 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
949 }
950}
951
953 StringRef GPUArch, StringRef LibDeviceFile,
954 DeviceLibABIVersion ABIVer) const {
955 if (!hasDeviceLibrary()) {
956 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
957 return false;
958 }
959 if (LibDeviceFile.empty()) {
960 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
961 return false;
962 }
963 if (ABIVer.requiresLibrary() && getABIVersionPath(ABIVer).empty()) {
964 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.toString();
965 return false;
966 }
967 return true;
968}
969
972 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, bool Wave64,
973 bool DAZ, bool FiniteOnly, bool UnsafeMathOpt, bool FastRelaxedMath,
974 bool CorrectSqrt, DeviceLibABIVersion ABIVer, bool isOpenMP = false) const {
976
977 auto AddBCLib = [&](StringRef BCFile) { BCLibs.push_back(BCFile.str()); };
978
979 AddBCLib(getOCMLPath());
980 if (!isOpenMP)
981 AddBCLib(getOCKLPath());
982 AddBCLib(getDenormalsAreZeroPath(DAZ));
983 AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath));
984 AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath));
985 AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt));
986 AddBCLib(getWavefrontSize64Path(Wave64));
987 AddBCLib(LibDeviceFile);
988 auto ABIVerPath = getABIVersionPath(ABIVer);
989 if (!ABIVerPath.empty())
990 AddBCLib(ABIVerPath);
991
992 return BCLibs;
993}
994
996ROCMToolChain::getCommonDeviceLibNames(const llvm::opt::ArgList &DriverArgs,
997 const std::string &GPUArch,
998 bool isOpenMP) const {
999 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1000 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1001
1002 StringRef LibDeviceFile = RocmInstallation->getLibDeviceFile(CanonArch);
1004 getAMDGPUCodeObjectVersion(getDriver(), DriverArgs));
1005 if (!RocmInstallation->checkCommonBitcodeLibs(CanonArch, LibDeviceFile,
1006 ABIVer))
1007 return {};
1008
1009 // If --hip-device-lib is not set, add the default bitcode libraries.
1010 // TODO: There are way too many flags that change this. Do we need to check
1011 // them all?
1012 bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
1013 options::OPT_fno_gpu_flush_denormals_to_zero,
1015 bool FiniteOnly = DriverArgs.hasFlag(
1016 options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only, false);
1017 bool UnsafeMathOpt =
1018 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
1019 options::OPT_fno_unsafe_math_optimizations, false);
1020 bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,
1021 options::OPT_fno_fast_math, false);
1022 bool CorrectSqrt = DriverArgs.hasFlag(
1023 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
1024 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, true);
1025 bool Wave64 = isWave64(DriverArgs, Kind);
1026
1027 return RocmInstallation->getCommonBitcodeLibs(
1028 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
1029 FastRelaxedMath, CorrectSqrt, ABIVer, isOpenMP);
1030}
#define V(N, I)
Definition: ASTContext.h:3273
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:180
const char * getClangProgramPath() const
Get the path to the main clang executable.
Definition: Driver.h:423
DiagnosticBuilder Diag(unsigned DiagID) const
Definition: Driver.h:144
const llvm::opt::OptTable & getOpts() const
Definition: Driver.h:399
std::string ResourceDir
The path to the compiler resource directory.
Definition: Driver.h:164
llvm::vfs::FileSystem & getVFS() const
Definition: Driver.h:403
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: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:952
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:971
void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const
Definition: AMDGPU.cpp:526
void print(raw_ostream &OS) const
Print information about the detected ROCm installation.
Definition: AMDGPU.cpp:520
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:864
llvm::Expected< std::unique_ptr< llvm::MemoryBuffer > > executeToolChainProgram(StringRef Executable) const
Executes the given Executable and returns the stdout.
Definition: ToolChain.cpp:107
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:773
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:699
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:756
StringRef getGPUArch(const llvm::opt::ArgList &DriverArgs) const
Get GPU arch from -mcpu without checking.
Definition: AMDGPU.cpp:844
virtual void checkTargetID(const llvm::opt::ArgList &DriverArgs) const
Check and diagnose invalid target ID specified by -mcpu.
Definition: AMDGPU.cpp:863
Tool * buildLinker() const override
Definition: AMDGPU.cpp:694
static bool isWave64(const llvm::opt::ArgList &DriverArgs, llvm::AMDGPU::GPUKind Kind)
Definition: AMDGPU.cpp:807
void addClangWarningOptions(llvm::opt::ArgStringList &CC1Args) const override
Common warning options shared by AMDGPU HIP, OpenCL and OpenMP toolchains.
Definition: AMDGPU.cpp:837
AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args)
AMDGPU Toolchain.
Definition: AMDGPU.cpp:682
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:850
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:824
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:873
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:996
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:897
ROCMToolChain(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args)
ROCM Toolchain.
Definition: AMDGPU.cpp:818
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:615
void getAMDGPUTargetFeatures(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args, std::vector< StringRef > &Features)
Definition: AMDGPU.cpp:643
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:444
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:105
llvm::StringRef getProcessorFromTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch)
Get processor name from target ID.
Definition: TargetID.cpp:54
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