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"
11 #include "clang/Basic/TargetID.h"
12 #include "clang/Config/config.h"
15 #include "clang/Driver/InputInfo.h"
16 #include "clang/Driver/Options.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 
29 using namespace clang::driver;
30 using namespace clang::driver::tools;
31 using namespace clang::driver::toolchains;
32 using namespace clang;
33 using 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.
42 RocmInstallationDetector::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 
77 void 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.
154 bool 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.
181 RocmInstallationDetector::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 
521 void RocmInstallationDetector::print(raw_ostream &OS) const {
522  if (hasHIPRuntime())
523  OS << "Found HIP installation: " << InstallPath << ", version "
524  << DetectedVersion << '\n';
525 }
526 
527 void 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;
579  if (hasHIPStdParLibrary())
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 
616 void amdgpu::Linker::ConstructJob(Compilation &C, const JobAction &JA,
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  addGPULibraries(getToolChain(), Args, CmdArgs);
652 
653  CmdArgs.push_back("-o");
654  CmdArgs.push_back(Output.getFilename());
655  C.addCommand(std::make_unique<Command>(
656  JA, *this, ResponseFileSupport::AtFileCurCP(), Args.MakeArgString(Linker),
657  CmdArgs, Inputs, Output));
658 }
659 
661  const llvm::Triple &Triple,
662  const llvm::opt::ArgList &Args,
663  std::vector<StringRef> &Features) {
664  // Add target ID features to -target-feature options. No diagnostics should
665  // be emitted here since invalid target ID is diagnosed at other places.
666  StringRef TargetID;
667  if (Args.hasArg(options::OPT_mcpu_EQ))
668  TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
669  else if (Args.hasArg(options::OPT_march_EQ))
670  TargetID = Args.getLastArgValue(options::OPT_march_EQ);
671  if (!TargetID.empty()) {
672  llvm::StringMap<bool> FeatureMap;
673  auto OptionalGpuArch = parseTargetID(Triple, TargetID, &FeatureMap);
674  if (OptionalGpuArch) {
675  StringRef GpuArch = *OptionalGpuArch;
676  // Iterate through all possible target ID features for the given GPU.
677  // If it is mapped to true, add +feature.
678  // If it is mapped to false, add -feature.
679  // If it is not in the map (default), do not add it
680  for (auto &&Feature : getAllPossibleTargetIDFeatures(Triple, GpuArch)) {
681  auto Pos = FeatureMap.find(Feature);
682  if (Pos == FeatureMap.end())
683  continue;
684  Features.push_back(Args.MakeArgStringRef(
685  (Twine(Pos->second ? "+" : "-") + Feature).str()));
686  }
687  }
688  }
689 
690  if (Args.hasFlag(options::OPT_mwavefrontsize64,
691  options::OPT_mno_wavefrontsize64, false))
692  Features.push_back("+wavefrontsize64");
693 
694  if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
695  options::OPT_mno_amdgpu_precise_memory_op, false))
696  Features.push_back("+precise-memory");
697 
698  handleTargetFeaturesGroup(D, Triple, Args, Features,
699  options::OPT_m_amdgpu_Features_Group);
700 }
701 
702 /// AMDGPU Toolchain
703 AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
704  const ArgList &Args)
705  : Generic_ELF(D, Triple, Args),
706  OptionsDefault(
707  {{options::OPT_O, "3"}, {options::OPT_cl_std_EQ, "CL1.2"}}) {
708  // Check code object version options. Emit warnings for legacy options
709  // and errors for the last invalid code object version options.
710  // It is done here to avoid repeated warning or error messages for
711  // each tool invocation.
713 }
714 
716  return new tools::amdgpu::Linker(*this);
717 }
718 
719 DerivedArgList *
720 AMDGPUToolChain::TranslateArgs(const DerivedArgList &Args, StringRef BoundArch,
721  Action::OffloadKind DeviceOffloadKind) const {
722 
723  DerivedArgList *DAL =
724  Generic_ELF::TranslateArgs(Args, BoundArch, DeviceOffloadKind);
725 
726  const OptTable &Opts = getDriver().getOpts();
727 
728  if (!DAL)
729  DAL = new DerivedArgList(Args.getBaseArgs());
730 
731  for (Arg *A : Args)
732  DAL->append(A);
733 
734  // Replace -mcpu=native with detected GPU.
735  Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
736  if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) == "native") {
737  DAL->eraseArg(options::OPT_mcpu_EQ);
738  auto GPUsOrErr = getSystemGPUArchs(Args);
739  if (!GPUsOrErr) {
740  getDriver().Diag(diag::err_drv_undetermined_gpu_arch)
741  << llvm::Triple::getArchTypeName(getArch())
742  << llvm::toString(GPUsOrErr.takeError()) << "-mcpu";
743  } else {
744  auto &GPUs = *GPUsOrErr;
745  if (GPUs.size() > 1) {
746  getDriver().Diag(diag::warn_drv_multi_gpu_arch)
747  << llvm::Triple::getArchTypeName(getArch())
748  << llvm::join(GPUs, ", ") << "-mcpu";
749  }
750  DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_mcpu_EQ),
751  Args.MakeArgString(GPUs.front()));
752  }
753  }
754 
755  checkTargetID(*DAL);
756 
757  if (Args.getLastArgValue(options::OPT_x) != "cl")
758  return DAL;
759 
760  // Phase 1 (.cl -> .bc)
761  if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
762  DAL->AddFlagArg(nullptr, Opts.getOption(getTriple().isArch64Bit()
763  ? options::OPT_m64
764  : options::OPT_m32));
765 
766  // Have to check OPT_O4, OPT_O0 & OPT_Ofast separately
767  // as they defined that way in Options.td
768  if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
769  options::OPT_Ofast))
770  DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_O),
771  getOptionDefault(options::OPT_O));
772  }
773 
774  return DAL;
775 }
776 
778  llvm::AMDGPU::GPUKind Kind) {
779 
780  // Assume nothing without a specific target.
781  if (Kind == llvm::AMDGPU::GK_NONE)
782  return false;
783 
784  const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
785 
786  // Default to enabling f32 denormals by default on subtargets where fma is
787  // fast with denormals
788  const bool BothDenormAndFMAFast =
789  (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
790  (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
791  return !BothDenormAndFMAFast;
792 }
793 
795  const llvm::opt::ArgList &DriverArgs, const JobAction &JA,
796  const llvm::fltSemantics *FPType) const {
797  // Denormals should always be enabled for f16 and f64.
798  if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
799  return llvm::DenormalMode::getIEEE();
800 
804  auto Kind = llvm::AMDGPU::parseArchAMDGCN(Arch);
805  if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
806  DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
807  options::OPT_fno_gpu_flush_denormals_to_zero,
809  return llvm::DenormalMode::getPreserveSign();
810 
811  return llvm::DenormalMode::getIEEE();
812  }
813 
814  const StringRef GpuArch = getGPUArch(DriverArgs);
815  auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
816 
817  // TODO: There are way too many flags that change this. Do we need to check
818  // them all?
819  bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
821 
822  // Outputs are flushed to zero (FTZ), preserving sign. Denormal inputs are
823  // also implicit treated as zero (DAZ).
824  return DAZ ? llvm::DenormalMode::getPreserveSign() :
825  llvm::DenormalMode::getIEEE();
826 }
827 
828 bool AMDGPUToolChain::isWave64(const llvm::opt::ArgList &DriverArgs,
829  llvm::AMDGPU::GPUKind Kind) {
830  const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
831  bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
832 
833  return !HasWave32 || DriverArgs.hasFlag(
834  options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64, false);
835 }
836 
837 
838 /// ROCM Toolchain
839 ROCMToolChain::ROCMToolChain(const Driver &D, const llvm::Triple &Triple,
840  const ArgList &Args)
841  : AMDGPUToolChain(D, Triple, Args) {
842  RocmInstallation->detectDeviceLibrary();
843 }
844 
846  const llvm::opt::ArgList &DriverArgs,
847  llvm::opt::ArgStringList &CC1Args,
848  Action::OffloadKind DeviceOffloadingKind) const {
849  // Default to "hidden" visibility, as object level linking will not be
850  // supported for the foreseeable future.
851  if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
852  options::OPT_fvisibility_ms_compat)) {
853  CC1Args.push_back("-fvisibility=hidden");
854  CC1Args.push_back("-fapply-global-visibility-to-externs");
855  }
856 }
857 
858 void AMDGPUToolChain::addClangWarningOptions(ArgStringList &CC1Args) const {
859  // AMDGPU does not support atomic lib call. Treat atomic alignment
860  // warnings as errors.
861  CC1Args.push_back("-Werror=atomic-alignment");
862 }
863 
864 StringRef
865 AMDGPUToolChain::getGPUArch(const llvm::opt::ArgList &DriverArgs) const {
867  getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
868 }
869 
871 AMDGPUToolChain::getParsedTargetID(const llvm::opt::ArgList &DriverArgs) const {
872  StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
873  if (TargetID.empty())
874  return {std::nullopt, std::nullopt, std::nullopt};
875 
876  llvm::StringMap<bool> FeatureMap;
877  auto OptionalGpuArch = parseTargetID(getTriple(), TargetID, &FeatureMap);
878  if (!OptionalGpuArch)
879  return {TargetID.str(), std::nullopt, std::nullopt};
880 
881  return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
882 }
883 
885  const llvm::opt::ArgList &DriverArgs) const {
886  auto PTID = getParsedTargetID(DriverArgs);
887  if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
888  getDriver().Diag(clang::diag::err_drv_bad_target_id)
889  << *PTID.OptionalTargetID;
890  }
891 }
892 
894 AMDGPUToolChain::getSystemGPUArchs(const ArgList &Args) const {
895  // Detect AMD GPUs availible on the system.
896  std::string Program;
897  if (Arg *A = Args.getLastArg(options::OPT_amdgpu_arch_tool_EQ))
898  Program = A->getValue();
899  else
900  Program = GetProgramPath("amdgpu-arch");
901 
902  auto StdoutOrErr = executeToolChainProgram(Program);
903  if (!StdoutOrErr)
904  return StdoutOrErr.takeError();
905 
907  for (StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(), "\n"))
908  if (!Arch.empty())
909  GPUArchs.push_back(Arch.str());
910 
911  if (GPUArchs.empty())
912  return llvm::createStringError(std::error_code(),
913  "No AMD GPU detected in the system");
914 
915  return std::move(GPUArchs);
916 }
917 
919  const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
920  Action::OffloadKind DeviceOffloadingKind) const {
921  AMDGPUToolChain::addClangTargetOptions(DriverArgs, CC1Args,
922  DeviceOffloadingKind);
923 
924  // For the OpenCL case where there is no offload target, accept -nostdlib to
925  // disable bitcode linking.
926  if (DeviceOffloadingKind == Action::OFK_None &&
927  DriverArgs.hasArg(options::OPT_nostdlib))
928  return;
929 
930  if (DriverArgs.hasArg(options::OPT_nogpulib))
931  return;
932 
933  // Get the device name and canonicalize it
934  const StringRef GpuArch = getGPUArch(DriverArgs);
935  auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
936  const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
937  StringRef LibDeviceFile = RocmInstallation->getLibDeviceFile(CanonArch);
939  getAMDGPUCodeObjectVersion(getDriver(), DriverArgs));
940  if (!RocmInstallation->checkCommonBitcodeLibs(CanonArch, LibDeviceFile,
941  ABIVer))
942  return;
943 
944  bool Wave64 = isWave64(DriverArgs, Kind);
945 
946  // TODO: There are way too many flags that change this. Do we need to check
947  // them all?
948  bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
950  bool FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only);
951 
952  bool UnsafeMathOpt =
953  DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations);
954  bool FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math);
955  bool CorrectSqrt =
956  DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt);
957 
958  // Add the OpenCL specific bitcode library.
960  BCLibs.push_back(RocmInstallation->getOpenCLPath().str());
961 
962  // Add the generic set of libraries.
963  BCLibs.append(RocmInstallation->getCommonBitcodeLibs(
964  DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
965  FastRelaxedMath, CorrectSqrt, ABIVer, false));
966 
967  if (getSanitizerArgs(DriverArgs).needsAsanRt()) {
968  CC1Args.push_back("-mlink-bitcode-file");
969  CC1Args.push_back(
970  DriverArgs.MakeArgString(RocmInstallation->getAsanRTLPath()));
971  }
972  for (StringRef BCFile : BCLibs) {
973  CC1Args.push_back("-mlink-builtin-bitcode");
974  CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
975  }
976 }
977 
979  StringRef GPUArch, StringRef LibDeviceFile,
980  DeviceLibABIVersion ABIVer) const {
981  if (!hasDeviceLibrary()) {
982  D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
983  return false;
984  }
985  if (LibDeviceFile.empty()) {
986  D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
987  return false;
988  }
989  if (ABIVer.requiresLibrary() && getABIVersionPath(ABIVer).empty()) {
990  D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.toString();
991  return false;
992  }
993  return true;
994 }
995 
998  const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, bool Wave64,
999  bool DAZ, bool FiniteOnly, bool UnsafeMathOpt, bool FastRelaxedMath,
1000  bool CorrectSqrt, DeviceLibABIVersion ABIVer, bool isOpenMP = false) const {
1002 
1003  auto AddBCLib = [&](StringRef BCFile) { BCLibs.push_back(BCFile.str()); };
1004 
1005  AddBCLib(getOCMLPath());
1006  if (!isOpenMP)
1007  AddBCLib(getOCKLPath());
1008  AddBCLib(getDenormalsAreZeroPath(DAZ));
1009  AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath));
1010  AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath));
1011  AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt));
1012  AddBCLib(getWavefrontSize64Path(Wave64));
1013  AddBCLib(LibDeviceFile);
1014  auto ABIVerPath = getABIVersionPath(ABIVer);
1015  if (!ABIVerPath.empty())
1016  AddBCLib(ABIVerPath);
1017 
1018  return BCLibs;
1019 }
1020 
1022  const llvm::opt::ArgList &DriverArgs, const std::string &GPUArch,
1023  const Action::OffloadKind DeviceOffloadingKind, bool isOpenMP) const {
1024  auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1025  const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1026 
1027  StringRef LibDeviceFile = RocmInstallation->getLibDeviceFile(CanonArch);
1029  getAMDGPUCodeObjectVersion(getDriver(), DriverArgs));
1030  if (!RocmInstallation->checkCommonBitcodeLibs(CanonArch, LibDeviceFile,
1031  ABIVer))
1032  return {};
1033 
1034  // If --hip-device-lib is not set, add the default bitcode libraries.
1035  // TODO: There are way too many flags that change this. Do we need to check
1036  // them all?
1037  bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
1038  options::OPT_fno_gpu_flush_denormals_to_zero,
1040  bool FiniteOnly = DriverArgs.hasFlag(
1041  options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only, false);
1042  bool UnsafeMathOpt =
1043  DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
1044  options::OPT_fno_unsafe_math_optimizations, false);
1045  bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,
1046  options::OPT_fno_fast_math, false);
1047  bool CorrectSqrt = false;
1048  if (DeviceOffloadingKind == Action::OFK_SYCL) {
1049  // When using SYCL, sqrt is only correctly rounded if the flag is specified
1050  CorrectSqrt = DriverArgs.hasArg(options::OPT_fsycl_fp32_prec_sqrt);
1051  } else
1052  CorrectSqrt = DriverArgs.hasFlag(
1053  options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
1054  options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, true);
1055  bool Wave64 = isWave64(DriverArgs, Kind);
1056 
1057  return RocmInstallation->getCommonBitcodeLibs(
1058  DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
1059  FastRelaxedMath, CorrectSqrt, ABIVer, isOpenMP);
1060 }
#define V(N, I)
Definition: ASTContext.h:3346
StringRef P
const Decl * D
IndirectLocalPath & Path
enum clang::sema::@1659::IndirectLocalPathEntry::EntryKind Kind
static void Append(char *Start, char *End, char *&Buffer, unsigned &BufferSize, unsigned &BufferCapacity)
__device__ __2f16 b
__device__ __2f16 float c
OffloadKind getOffloadingDeviceKind() const
Definition: Action.h:222
const char * getOffloadingArch() const
Definition: Action.h:223
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:182
llvm::vfs::FileSystem & getVFS() const
Definition: Driver.h:405
DiagnosticBuilder Diag(unsigned DiagID) const
Definition: Driver.h:146
const llvm::opt::OptTable & getOpts() const
Definition: Driver.h:401
std::string ResourceDir
The path to the compiler resource directory.
Definition: Driver.h:166
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:978
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:997
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
const Driver & getDriver() const
Definition: ToolChain.h:269
llvm::Triple::ArchType getArch() const
Definition: ToolChain.h:285
const llvm::Triple & getTriple() const
Definition: ToolChain.h:271
std::string GetProgramPath(const char *Name) const
Definition: ToolChain.cpp:981
SanitizerArgs getSanitizerArgs(const llvm::opt::ArgList &JobArgs) const
Definition: ToolChain.cpp:334
llvm::Expected< std::unique_ptr< llvm::MemoryBuffer > > executeToolChainProgram(StringRef Executable) const
Executes the given Executable and returns the stdout.
Definition: ToolChain.cpp:111
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:794
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:720
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:777
StringRef getGPUArch(const llvm::opt::ArgList &DriverArgs) const
Get GPU arch from -mcpu without checking.
Definition: AMDGPU.cpp:865
virtual void checkTargetID(const llvm::opt::ArgList &DriverArgs) const
Check and diagnose invalid target ID specified by -mcpu.
Definition: AMDGPU.cpp:884
Tool * buildLinker() const override
Definition: AMDGPU.cpp:715
static bool isWave64(const llvm::opt::ArgList &DriverArgs, llvm::AMDGPU::GPUKind Kind)
Definition: AMDGPU.cpp:828
void addClangWarningOptions(llvm::opt::ArgStringList &CC1Args) const override
Common warning options shared by AMDGPU HIP, OpenCL and OpenMP toolchains.
Definition: AMDGPU.cpp:858
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:871
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:845
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:894
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: Gnu.cpp:3568
LazyDetector< RocmInstallationDetector > RocmInstallation
Definition: Gnu.h:296
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:918
llvm::SmallVector< std::string, 12 > getCommonDeviceLibNames(const llvm::opt::ArgList &DriverArgs, const std::string &GPUArch, const Action::OffloadKind DeviceOffloadingKind, bool isOpenMP=false) const
Definition: AMDGPU.cpp:1021
ROCMToolChain(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args)
ROCM Toolchain.
Definition: AMDGPU.cpp:839
void getAMDGPUTargetFeatures(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args, std::vector< StringRef > &Features)
Definition: AMDGPU.cpp:660
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 addGPULibraries(const ToolChain &TC, const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs)
Definition: CommonArgs.cpp:521
void addLinkerCompressDebugSectionsOption(const ToolChain &TC, const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs)
Definition: CommonArgs.cpp:503
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)
bool LE(InterpState &S, CodePtr OpPC)
Definition: Interp.h:1103
std::string toString(const til::SExpr *E)
The JSON file list parser is used to communicate input to InstallAPI.
@ OpenCL
Definition: LangStandard.h:66
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:93
The struct type returned by getParsedTargetID.
Definition: AMDGPU.h:110