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 const char *InstallDir = D.getInstalledDir();
213
214 // Check both a normal Unix prefix position of the clang binary, as well as
215 // the Windows-esque layout the ROCm packages use with the host architecture
216 // subdirectory of bin.
217 auto DeduceROCmPath = [](StringRef ClangPath) {
218 // Strip off directory (usually bin)
219 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
220 StringRef ParentName = llvm::sys::path::filename(ParentDir);
221
222 // Some builds use bin/{host arch}, so go up again.
223 if (ParentName == "bin") {
224 ParentDir = llvm::sys::path::parent_path(ParentDir);
225 ParentName = llvm::sys::path::filename(ParentDir);
226 }
227
228 // Detect ROCm packages built with SPACK.
229 // clang is installed at
230 // <rocm_root>/llvm-amdgpu-<rocm_release_string>-<hash>/bin directory.
231 // We only consider the parent directory of llvm-amdgpu package as ROCm
232 // installation candidate for SPACK.
233 if (ParentName.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 AddLinkerInputs(getToolChain(), Inputs, Args, CmdArgs, JA);
629 if (C.getDriver().isUsingLTO())
630 addLTOOptions(getToolChain(), Args, CmdArgs, Output, Inputs[0],
631 C.getDriver().getLTOMode() == LTOK_Thin);
632 else if (Args.hasArg(options::OPT_mcpu_EQ))
633 CmdArgs.push_back(Args.MakeArgString(
634 "-plugin-opt=mcpu=" + Args.getLastArgValue(options::OPT_mcpu_EQ)));
635 CmdArgs.push_back("-o");
636 CmdArgs.push_back(Output.getFilename());
637 C.addCommand(std::make_unique<Command>(
638 JA, *this, ResponseFileSupport::AtFileCurCP(), Args.MakeArgString(Linker),
639 CmdArgs, Inputs, Output));
640}
641
643 const llvm::Triple &Triple,
644 const llvm::opt::ArgList &Args,
645 std::vector<StringRef> &Features) {
646 // Add target ID features to -target-feature options. No diagnostics should
647 // be emitted here since invalid target ID is diagnosed at other places.
648 StringRef TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
649 if (!TargetID.empty()) {
650 llvm::StringMap<bool> FeatureMap;
651 auto OptionalGpuArch = parseTargetID(Triple, TargetID, &FeatureMap);
652 if (OptionalGpuArch) {
653 StringRef GpuArch = *OptionalGpuArch;
654 // Iterate through all possible target ID features for the given GPU.
655 // If it is mapped to true, add +feature.
656 // If it is mapped to false, add -feature.
657 // If it is not in the map (default), do not add it
658 for (auto &&Feature : getAllPossibleTargetIDFeatures(Triple, GpuArch)) {
659 auto Pos = FeatureMap.find(Feature);
660 if (Pos == FeatureMap.end())
661 continue;
662 Features.push_back(Args.MakeArgStringRef(
663 (Twine(Pos->second ? "+" : "-") + Feature).str()));
664 }
665 }
666 }
667
668 if (Args.hasFlag(options::OPT_mwavefrontsize64,
669 options::OPT_mno_wavefrontsize64, false))
670 Features.push_back("+wavefrontsize64");
671
672 handleTargetFeaturesGroup(D, Triple, Args, Features,
673 options::OPT_m_amdgpu_Features_Group);
674}
675
676/// AMDGPU Toolchain
677AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
678 const ArgList &Args)
679 : Generic_ELF(D, Triple, Args),
680 OptionsDefault(
681 {{options::OPT_O, "3"}, {options::OPT_cl_std_EQ, "CL1.2"}}) {
682 // Check code object version options. Emit warnings for legacy options
683 // and errors for the last invalid code object version options.
684 // It is done here to avoid repeated warning or error messages for
685 // each tool invocation.
687}
688
690 return new tools::amdgpu::Linker(*this);
691}
692
693DerivedArgList *
694AMDGPUToolChain::TranslateArgs(const DerivedArgList &Args, StringRef BoundArch,
695 Action::OffloadKind DeviceOffloadKind) const {
696
697 DerivedArgList *DAL =
698 Generic_ELF::TranslateArgs(Args, BoundArch, DeviceOffloadKind);
699
700 const OptTable &Opts = getDriver().getOpts();
701
702 if (!DAL)
703 DAL = new DerivedArgList(Args.getBaseArgs());
704
705 for (Arg *A : Args)
706 DAL->append(A);
707
708 // Replace -mcpu=native with detected GPU.
709 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
710 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) == "native") {
711 DAL->eraseArg(options::OPT_mcpu_EQ);
712 auto GPUsOrErr = getSystemGPUArchs(Args);
713 if (!GPUsOrErr) {
714 getDriver().Diag(diag::err_drv_undetermined_gpu_arch)
715 << llvm::Triple::getArchTypeName(getArch())
716 << llvm::toString(GPUsOrErr.takeError()) << "-mcpu";
717 } else {
718 auto &GPUs = *GPUsOrErr;
719 if (GPUs.size() > 1) {
720 getDriver().Diag(diag::warn_drv_multi_gpu_arch)
721 << llvm::Triple::getArchTypeName(getArch())
722 << llvm::join(GPUs, ", ") << "-mcpu";
723 }
724 DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_mcpu_EQ),
725 Args.MakeArgString(GPUs.front()));
726 }
727 }
728
729 checkTargetID(*DAL);
730
731 if (!Args.getLastArgValue(options::OPT_x).equals("cl"))
732 return DAL;
733
734 // Phase 1 (.cl -> .bc)
735 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
736 DAL->AddFlagArg(nullptr, Opts.getOption(getTriple().isArch64Bit()
737 ? options::OPT_m64
738 : options::OPT_m32));
739
740 // Have to check OPT_O4, OPT_O0 & OPT_Ofast separately
741 // as they defined that way in Options.td
742 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
743 options::OPT_Ofast))
744 DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_O),
745 getOptionDefault(options::OPT_O));
746 }
747
748 return DAL;
749}
750
752 llvm::AMDGPU::GPUKind Kind) {
753
754 // Assume nothing without a specific target.
755 if (Kind == llvm::AMDGPU::GK_NONE)
756 return false;
757
758 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
759
760 // Default to enabling f32 denormals by default on subtargets where fma is
761 // fast with denormals
762 const bool BothDenormAndFMAFast =
763 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
764 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
765 return !BothDenormAndFMAFast;
766}
767
769 const llvm::opt::ArgList &DriverArgs, const JobAction &JA,
770 const llvm::fltSemantics *FPType) const {
771 // Denormals should always be enabled for f16 and f64.
772 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
773 return llvm::DenormalMode::getIEEE();
774
778 auto Kind = llvm::AMDGPU::parseArchAMDGCN(Arch);
779 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
780 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
781 options::OPT_fno_gpu_flush_denormals_to_zero,
783 return llvm::DenormalMode::getPreserveSign();
784
785 return llvm::DenormalMode::getIEEE();
786 }
787
788 const StringRef GpuArch = getGPUArch(DriverArgs);
789 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
790
791 // TODO: There are way too many flags that change this. Do we need to check
792 // them all?
793 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
795
796 // Outputs are flushed to zero (FTZ), preserving sign. Denormal inputs are
797 // also implicit treated as zero (DAZ).
798 return DAZ ? llvm::DenormalMode::getPreserveSign() :
799 llvm::DenormalMode::getIEEE();
800}
801
802bool AMDGPUToolChain::isWave64(const llvm::opt::ArgList &DriverArgs,
803 llvm::AMDGPU::GPUKind Kind) {
804 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
805 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
806
807 return !HasWave32 || DriverArgs.hasFlag(
808 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64, false);
809}
810
811
812/// ROCM Toolchain
813ROCMToolChain::ROCMToolChain(const Driver &D, const llvm::Triple &Triple,
814 const ArgList &Args)
815 : AMDGPUToolChain(D, Triple, Args) {
816 RocmInstallation->detectDeviceLibrary();
817}
818
820 const llvm::opt::ArgList &DriverArgs,
821 llvm::opt::ArgStringList &CC1Args,
822 Action::OffloadKind DeviceOffloadingKind) const {
823 // Default to "hidden" visibility, as object level linking will not be
824 // supported for the foreseeable future.
825 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
826 options::OPT_fvisibility_ms_compat)) {
827 CC1Args.push_back("-fvisibility=hidden");
828 CC1Args.push_back("-fapply-global-visibility-to-externs");
829 }
830}
831
832void AMDGPUToolChain::addClangWarningOptions(ArgStringList &CC1Args) const {
833 // AMDGPU does not support atomic lib call. Treat atomic alignment
834 // warnings as errors.
835 CC1Args.push_back("-Werror=atomic-alignment");
836}
837
838StringRef
839AMDGPUToolChain::getGPUArch(const llvm::opt::ArgList &DriverArgs) const {
841 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
842}
843
845AMDGPUToolChain::getParsedTargetID(const llvm::opt::ArgList &DriverArgs) const {
846 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
847 if (TargetID.empty())
848 return {std::nullopt, std::nullopt, std::nullopt};
849
850 llvm::StringMap<bool> FeatureMap;
851 auto OptionalGpuArch = parseTargetID(getTriple(), TargetID, &FeatureMap);
852 if (!OptionalGpuArch)
853 return {TargetID.str(), std::nullopt, std::nullopt};
854
855 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
856}
857
859 const llvm::opt::ArgList &DriverArgs) const {
860 auto PTID = getParsedTargetID(DriverArgs);
861 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
862 getDriver().Diag(clang::diag::err_drv_bad_target_id)
863 << *PTID.OptionalTargetID;
864 }
865}
866
868AMDGPUToolChain::getSystemGPUArchs(const ArgList &Args) const {
869 // Detect AMD GPUs availible on the system.
870 std::string Program;
871 if (Arg *A = Args.getLastArg(options::OPT_amdgpu_arch_tool_EQ))
872 Program = A->getValue();
873 else
874 Program = GetProgramPath("amdgpu-arch");
875
876 auto StdoutOrErr = executeToolChainProgram(Program);
877 if (!StdoutOrErr)
878 return StdoutOrErr.takeError();
879
881 for (StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(), "\n"))
882 if (!Arch.empty())
883 GPUArchs.push_back(Arch.str());
884
885 if (GPUArchs.empty())
886 return llvm::createStringError(std::error_code(),
887 "No AMD GPU detected in the system");
888
889 return std::move(GPUArchs);
890}
891
893 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
894 Action::OffloadKind DeviceOffloadingKind) const {
895 AMDGPUToolChain::addClangTargetOptions(DriverArgs, CC1Args,
896 DeviceOffloadingKind);
897
898 // For the OpenCL case where there is no offload target, accept -nostdlib to
899 // disable bitcode linking.
900 if (DeviceOffloadingKind == Action::OFK_None &&
901 DriverArgs.hasArg(options::OPT_nostdlib))
902 return;
903
904 if (DriverArgs.hasArg(options::OPT_nogpulib))
905 return;
906
907 // Get the device name and canonicalize it
908 const StringRef GpuArch = getGPUArch(DriverArgs);
909 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
910 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
911 StringRef LibDeviceFile = RocmInstallation->getLibDeviceFile(CanonArch);
914 if (!RocmInstallation->checkCommonBitcodeLibs(CanonArch, LibDeviceFile,
915 ABIVer))
916 return;
917
918 bool Wave64 = isWave64(DriverArgs, Kind);
919
920 // TODO: There are way too many flags that change this. Do we need to check
921 // them all?
922 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
924 bool FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only);
925
926 bool UnsafeMathOpt =
927 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations);
928 bool FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math);
929 bool CorrectSqrt =
930 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt);
931
932 // Add the OpenCL specific bitcode library.
934 BCLibs.push_back(RocmInstallation->getOpenCLPath().str());
935
936 // Add the generic set of libraries.
937 BCLibs.append(RocmInstallation->getCommonBitcodeLibs(
938 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
939 FastRelaxedMath, CorrectSqrt, ABIVer, false));
940
941 for (StringRef BCFile : BCLibs) {
942 CC1Args.push_back("-mlink-builtin-bitcode");
943 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
944 }
945}
946
948 StringRef GPUArch, StringRef LibDeviceFile,
949 DeviceLibABIVersion ABIVer) const {
950 if (!hasDeviceLibrary()) {
951 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
952 return false;
953 }
954 if (LibDeviceFile.empty()) {
955 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
956 return false;
957 }
958 if (ABIVer.requiresLibrary() && getABIVersionPath(ABIVer).empty()) {
959 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.toString();
960 return false;
961 }
962 return true;
963}
964
967 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, bool Wave64,
968 bool DAZ, bool FiniteOnly, bool UnsafeMathOpt, bool FastRelaxedMath,
969 bool CorrectSqrt, DeviceLibABIVersion ABIVer, bool isOpenMP = false) const {
971
972 auto AddBCLib = [&](StringRef BCFile) { BCLibs.push_back(BCFile.str()); };
973
974 AddBCLib(getOCMLPath());
975 if (!isOpenMP)
976 AddBCLib(getOCKLPath());
977 AddBCLib(getDenormalsAreZeroPath(DAZ));
978 AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath));
979 AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath));
980 AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt));
981 AddBCLib(getWavefrontSize64Path(Wave64));
982 AddBCLib(LibDeviceFile);
983 auto ABIVerPath = getABIVersionPath(ABIVer);
984 if (!ABIVerPath.empty())
985 AddBCLib(ABIVerPath);
986
987 return BCLibs;
988}
989
991ROCMToolChain::getCommonDeviceLibNames(const llvm::opt::ArgList &DriverArgs,
992 const std::string &GPUArch,
993 bool isOpenMP) const {
994 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
995 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
996
997 StringRef LibDeviceFile = RocmInstallation->getLibDeviceFile(CanonArch);
1000 if (!RocmInstallation->checkCommonBitcodeLibs(CanonArch, LibDeviceFile,
1001 ABIVer))
1002 return {};
1003
1004 // If --hip-device-lib is not set, add the default bitcode libraries.
1005 // TODO: There are way too many flags that change this. Do we need to check
1006 // them all?
1007 bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
1008 options::OPT_fno_gpu_flush_denormals_to_zero,
1010 bool FiniteOnly = DriverArgs.hasFlag(
1011 options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only, false);
1012 bool UnsafeMathOpt =
1013 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
1014 options::OPT_fno_unsafe_math_optimizations, false);
1015 bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,
1016 options::OPT_fno_fast_math, false);
1017 bool CorrectSqrt = DriverArgs.hasFlag(
1018 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
1019 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, true);
1020 bool Wave64 = isWave64(DriverArgs, Kind);
1021
1022 return RocmInstallation->getCommonBitcodeLibs(
1023 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
1024 FastRelaxedMath, CorrectSqrt, ABIVer, isOpenMP);
1025}
#define V(N, I)
Definition: ASTContext.h:3255
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:212
OffloadKind getOffloadingDeviceKind() const
Definition: Action.h:211
Compilation - A set of tasks to perform for a single driver invocation.
Definition: Compilation.h:45
Driver - Encapsulate logic for constructing compilation processes from a set of gcc-driver-like comma...
Definition: Driver.h:77
std::string SysRoot
sysroot, if present
Definition: Driver.h:183
const char * getClangProgramPath() const
Get the path to the main clang executable.
Definition: Driver.h:421
DiagnosticBuilder Diag(unsigned DiagID) const
Definition: Driver.h:144
const llvm::opt::OptTable & getOpts() const
Definition: Driver.h:397
std::string ResourceDir
The path to the compiler resource directory.
Definition: Driver.h:167
const char * getInstalledDir() const
Get the path to where the clang executable was installed.
Definition: Driver.h:426
llvm::vfs::FileSystem & getVFS() const
Definition: Driver.h:401
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:947
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:966
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:261
const Driver & getDriver() const
Definition: ToolChain.h:245
virtual llvm::opt::DerivedArgList * TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch, Action::OffloadKind DeviceOffloadKind) const
TranslateArgs - Create a new derived argument list for any argument translations this ToolChain may w...
Definition: ToolChain.h:348
const llvm::Triple & getTriple() const
Definition: ToolChain.h:247
std::string GetProgramPath(const char *Name) const
Definition: ToolChain.cpp:829
llvm::Expected< std::unique_ptr< llvm::MemoryBuffer > > executeToolChainProgram(StringRef Executable) const
Executes the given Executable and returns the stdout.
Definition: ToolChain.cpp:98
Tool - Information on a specific compilation tool.
Definition: Tool.h:32
llvm::DenormalMode getDefaultDenormalModeForType(const llvm::opt::ArgList &DriverArgs, const JobAction &JA, const llvm::fltSemantics *FPType=nullptr) const override
Returns the output denormal handling type in the default floating point environment for the given FPT...
Definition: AMDGPU.cpp:768
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:694
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:751
StringRef getGPUArch(const llvm::opt::ArgList &DriverArgs) const
Get GPU arch from -mcpu without checking.
Definition: AMDGPU.cpp:839
virtual void checkTargetID(const llvm::opt::ArgList &DriverArgs) const
Check and diagnose invalid target ID specified by -mcpu.
Definition: AMDGPU.cpp:858
Tool * buildLinker() const override
Definition: AMDGPU.cpp:689
static bool isWave64(const llvm::opt::ArgList &DriverArgs, llvm::AMDGPU::GPUKind Kind)
Definition: AMDGPU.cpp:802
void addClangWarningOptions(llvm::opt::ArgStringList &CC1Args) const override
Common warning options shared by AMDGPU HIP, OpenCL and OpenMP toolchains.
Definition: AMDGPU.cpp:832
AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args)
AMDGPU Toolchain.
Definition: AMDGPU.cpp:677
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:845
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:819
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:868
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:991
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:892
ROCMToolChain(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args)
ROCM Toolchain.
Definition: AMDGPU.cpp:813
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:642
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:443
void addLTOOptions(const ToolChain &ToolChain, const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs, const InputInfo &Output, const InputInfo &Input, bool IsThinLTO)
void AddLinkerInputs(const ToolChain &TC, const InputInfoList &Inputs, const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs, const JobAction &JA)
unsigned getAMDGPUCodeObjectVersion(const Driver &D, const llvm::opt::ArgList &Args)
std::optional< llvm::StringRef > parseTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch, llvm::StringMap< bool > *FeatureMap)
Parse a target ID to get processor and feature map.
Definition: TargetID.cpp:105
llvm::StringRef getProcessorFromTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch)
Get processor name from target ID.
Definition: TargetID.cpp:54
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