11#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 "llvm/TargetParser/TargetParser.h"
27#include <system_error>
35RocmInstallationDetector::CommonBitcodeLibsPreferences::
36 CommonBitcodeLibsPreferences(
const Driver &D,
37 const llvm::opt::ArgList &DriverArgs,
40 const bool NeedsASanRT)
43 const auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
44 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
48 const bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
50 !HasWave32 || DriverArgs.hasFlag(options::OPT_mwavefrontsize64,
51 options::OPT_mno_wavefrontsize64,
false);
58 const bool DefaultDAZ =
59 (Kind == llvm::AMDGPU::GK_NONE)
61 : !((ArchAttr &
llvm::
AMDGPU::FEATURE_FAST_FMA_F32) &&
62 (ArchAttr &
llvm::
AMDGPU::FEATURE_FAST_DENORMAL_F32));
65 DAZ = IsKnownOffloading
66 ? DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
67 options::OPT_fno_gpu_flush_denormals_to_zero,
69 : DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) || DefaultDAZ;
71 FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only) ||
72 DriverArgs.hasFlag(options::OPT_ffinite_math_only,
73 options::OPT_fno_finite_math_only,
false);
76 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations) ||
77 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
78 options::OPT_fno_unsafe_math_optimizations,
false);
80 FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math) ||
81 DriverArgs.hasFlag(options::OPT_ffast_math,
82 options::OPT_fno_fast_math,
false);
84 const bool DefaultSqrt = IsKnownOffloading ?
true :
false;
86 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt) ||
88 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
89 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, DefaultSqrt);
92 GPUSan = (DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
93 options::OPT_fno_gpu_sanitize,
true) &&
97void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) {
98 assert(!Path.empty());
100 const StringRef Suffix(
".bc");
101 const StringRef Suffix2(
".amdgcn.bc");
104 for (llvm::vfs::directory_iterator LI = D.getVFS().dir_begin(Path, EC), LE;
105 !EC && LI != LE; LI = LI.increment(EC)) {
106 StringRef FilePath = LI->path();
107 StringRef
FileName = llvm::sys::path::filename(FilePath);
113 BaseName =
FileName.drop_back(Suffix2.size());
114 else if (
FileName.ends_with(Suffix))
115 BaseName =
FileName.drop_back(Suffix.size());
117 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
118 if (BaseName ==
"ocml") {
120 }
else if (BaseName ==
"ockl") {
122 }
else if (BaseName ==
"opencl") {
124 }
else if (BaseName ==
"asanrtl") {
126 }
else if (BaseName ==
"oclc_finite_only_off") {
127 FiniteOnly.Off = FilePath;
128 }
else if (BaseName ==
"oclc_finite_only_on") {
129 FiniteOnly.On = FilePath;
130 }
else if (BaseName ==
"oclc_daz_opt_on") {
131 DenormalsAreZero.On = FilePath;
132 }
else if (BaseName ==
"oclc_daz_opt_off") {
133 DenormalsAreZero.Off = FilePath;
134 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_on") {
135 CorrectlyRoundedSqrt.On = FilePath;
136 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_off") {
137 CorrectlyRoundedSqrt.Off = FilePath;
138 }
else if (BaseName ==
"oclc_unsafe_math_on") {
139 UnsafeMath.On = FilePath;
140 }
else if (BaseName ==
"oclc_unsafe_math_off") {
141 UnsafeMath.Off = FilePath;
142 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
143 WavefrontSize64.On = FilePath;
144 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
145 WavefrontSize64.Off = FilePath;
146 }
else if (BaseName.starts_with(ABIVersionPrefix)) {
147 unsigned ABIVersionNumber;
148 if (BaseName.drop_front(ABIVersionPrefix.size())
149 .getAsInteger(0, ABIVersionNumber))
151 ABIVersionMap[ABIVersionNumber] = FilePath.str();
155 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
156 if (!BaseName.starts_with(DeviceLibPrefix))
159 StringRef IsaVersionNumber =
160 BaseName.drop_front(DeviceLibPrefix.size());
162 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
165 std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
172bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
173 SmallVector<StringRef, 4> VersionParts;
174 V.split(VersionParts,
'\n');
175 unsigned Major = ~0U;
176 unsigned Minor = ~0U;
177 for (
auto Part : VersionParts) {
178 auto Splits = Part.rtrim().split(
'=');
179 if (Splits.first ==
"HIP_VERSION_MAJOR") {
180 if (Splits.second.getAsInteger(0, Major))
182 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
183 if (Splits.second.getAsInteger(0, Minor))
185 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
186 VersionPatch = Splits.second.str();
188 if (Major == ~0U || Minor == ~0U)
190 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
192 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
198const SmallVectorImpl<RocmInstallationDetector::Candidate> &
199RocmInstallationDetector::getInstallationPathCandidates() {
202 if (!ROCmSearchDirs.empty())
203 return ROCmSearchDirs;
205 auto DoPrintROCmSearchDirs = [&]() {
206 if (PrintROCmSearchDirs)
207 for (
auto Cand : ROCmSearchDirs) {
208 llvm::errs() <<
"ROCm installation search path: " << Cand.Path <<
'\n';
214 if (!RocmPathArg.empty()) {
215 ROCmSearchDirs.emplace_back(RocmPathArg.str());
216 DoPrintROCmSearchDirs();
217 return ROCmSearchDirs;
218 }
else if (std::optional<std::string> RocmPathEnv =
219 llvm::sys::Process::GetEnv(
"ROCM_PATH")) {
220 if (!RocmPathEnv->empty()) {
221 ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));
222 DoPrintROCmSearchDirs();
223 return ROCmSearchDirs;
228 StringRef InstallDir = D.Dir;
233 auto DeduceROCmPath = [](StringRef ClangPath) {
235 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
236 StringRef ParentName = llvm::sys::path::filename(ParentDir);
239 if (ParentName ==
"bin") {
240 ParentDir = llvm::sys::path::parent_path(ParentDir);
241 ParentName = llvm::sys::path::filename(ParentDir);
246 if (ParentName ==
"llvm" || ParentName.starts_with(
"aomp")) {
247 ParentDir = llvm::sys::path::parent_path(ParentDir);
248 ParentName = llvm::sys::path::filename(ParentDir);
252 if (ParentName ==
"lib")
253 ParentDir = llvm::sys::path::parent_path(ParentDir);
256 return Candidate(ParentDir.str(),
true);
261 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
265 llvm::SmallString<256> RealClangPath;
266 llvm::sys::fs::real_path(D.getClangProgramPath(), RealClangPath);
267 auto ParentPath = llvm::sys::path::parent_path(RealClangPath);
268 if (ParentPath != InstallDir)
269 ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));
272 auto ClangRoot = llvm::sys::path::parent_path(InstallDir);
273 auto RealClangRoot = llvm::sys::path::parent_path(ParentPath);
274 ROCmSearchDirs.emplace_back(ClangRoot.str(),
true);
275 if (RealClangRoot != ClangRoot)
276 ROCmSearchDirs.emplace_back(RealClangRoot.str(),
true);
277 ROCmSearchDirs.emplace_back(D.ResourceDir,
280 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/rocm",
285 std::string LatestROCm;
286 llvm::VersionTuple LatestVer;
288 auto GetROCmVersion = [](StringRef DirName) {
289 llvm::VersionTuple
V;
290 std::string VerStr = DirName.drop_front(strlen(
"rocm-")).str();
293 llvm::replace(VerStr,
'-',
'.');
297 for (llvm::vfs::directory_iterator
298 File = D.getVFS().dir_begin(D.SysRoot +
"/opt", EC),
300 File != FileEnd && !EC;
File.increment(EC)) {
301 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
304 if (LatestROCm.empty()) {
306 LatestVer = GetROCmVersion(LatestROCm);
309 auto Ver = GetROCmVersion(
FileName);
310 if (LatestVer < Ver) {
315 if (!LatestROCm.empty())
316 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/" + LatestROCm,
319 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr/local",
321 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr",
324 DoPrintROCmSearchDirs();
325 return ROCmSearchDirs;
329 const Driver &D,
const llvm::Triple &HostTriple,
330 const llvm::opt::ArgList &Args,
bool DetectHIPRuntime,
bool DetectDeviceLib)
332 Verbose = Args.hasArg(options::OPT_v);
333 RocmPathArg = Args.getLastArgValue(options::OPT_rocm_path_EQ);
334 PrintROCmSearchDirs = Args.hasArg(options::OPT_print_rocm_search_dirs);
335 RocmDeviceLibPathArg =
336 Args.getAllArgValues(options::OPT_rocm_device_lib_path_EQ);
337 HIPPathArg = Args.getLastArgValue(options::OPT_hip_path_EQ);
338 HIPStdParPathArg = Args.getLastArgValue(options::OPT_hipstdpar_path_EQ);
339 HasHIPStdParLibrary =
340 !HIPStdParPathArg.empty() && D.getVFS().exists(HIPStdParPathArg +
341 "/hipstdpar_lib.hpp");
342 HIPRocThrustPathArg =
343 Args.getLastArgValue(options::OPT_hipstdpar_thrust_path_EQ);
344 HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&
345 D.getVFS().exists(HIPRocThrustPathArg +
"/thrust");
346 HIPRocPrimPathArg = Args.getLastArgValue(options::OPT_hipstdpar_prim_path_EQ);
347 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
348 D.getVFS().exists(HIPRocPrimPathArg +
"/rocprim");
350 if (
auto *A = Args.getLastArg(options::OPT_hip_version_EQ)) {
351 HIPVersionArg = A->getValue();
352 unsigned Major = ~0U;
353 unsigned Minor = ~0U;
354 SmallVector<StringRef, 3> Parts;
355 HIPVersionArg.split(Parts,
'.');
357 Parts[0].getAsInteger(0, Major);
358 if (Parts.size() > 1)
359 Parts[1].getAsInteger(0, Minor);
360 if (Parts.size() > 2)
361 VersionPatch = Parts[2].str();
362 if (VersionPatch.empty())
364 if (Major != ~0U && Minor == ~0U)
366 if (Major == ~0U || Minor == ~0U)
367 D.Diag(diag::err_drv_invalid_value)
368 << A->getAsString(Args) << HIPVersionArg;
370 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
372 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
374 VersionPatch = DefaultVersionPatch;
376 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
377 DetectedVersion = (Twine(DefaultVersionMajor) +
"." +
378 Twine(DefaultVersionMinor) +
"." + VersionPatch)
382 if (DetectHIPRuntime)
389 assert(LibDevicePath.empty());
391 if (!RocmDeviceLibPathArg.empty())
392 LibDevicePath = RocmDeviceLibPathArg[RocmDeviceLibPathArg.size() - 1];
393 else if (std::optional<std::string> LibPathEnv =
394 llvm::sys::Process::GetEnv(
"HIP_DEVICE_LIB_PATH"))
395 LibDevicePath = std::move(*LibPathEnv);
397 auto &FS = D.getVFS();
398 if (!LibDevicePath.empty()) {
402 if (!FS.exists(LibDevicePath))
405 scanLibDevicePath(LibDevicePath);
406 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
411 auto CheckDeviceLib = [&](StringRef Path,
bool StrictChecking) {
412 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
413 if (CheckLibDevice && !FS.exists(Path))
416 scanLibDevicePath(Path);
418 if (!NoBuiltinLibs) {
420 if (!allGenericLibsValid())
425 if (LibDeviceMap.empty())
432 LibDevicePath = D.ResourceDir;
433 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
434 "amdgcn",
"bitcode");
435 HasDeviceLibrary = CheckDeviceLib(LibDevicePath,
true);
436 if (HasDeviceLibrary)
441 auto &ROCmDirs = getInstallationPathCandidates();
442 for (
const auto &Candidate : ROCmDirs) {
443 LibDevicePath = Candidate.Path;
444 llvm::sys::path::append(LibDevicePath,
"amdgcn",
"bitcode");
445 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);
446 if (HasDeviceLibrary)
453 if (!HIPPathArg.empty())
454 HIPSearchDirs.emplace_back(HIPPathArg.str());
455 else if (std::optional<std::string> HIPPathEnv =
456 llvm::sys::Process::GetEnv(
"HIP_PATH")) {
457 if (!HIPPathEnv->empty())
458 HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));
460 if (HIPSearchDirs.empty())
461 HIPSearchDirs.append(getInstallationPathCandidates());
462 auto &FS = D.getVFS();
464 for (
const auto &Candidate : HIPSearchDirs) {
465 InstallPath = Candidate.Path;
466 if (InstallPath.empty() || !FS.exists(InstallPath))
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");
479 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
480 llvm::sys::path::append(ParentSharePath,
"share");
483 const Twine &
c =
"",
const Twine &d =
"") {
485 llvm::sys::path::append(newpath, a,
b,
c, d);
489 std::vector<SmallString<0>> VersionFilePaths = {
490 Append(SharePath,
"hip",
"version"),
491 InstallPath != D.SysRoot +
"/usr/local"
492 ?
Append(ParentSharePath,
"hip",
"version")
494 Append(BinPath,
".hipVersion")};
496 for (
const auto &VersionFilePath : VersionFilePaths) {
497 if (VersionFilePath.empty())
499 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
500 FS.getBufferForFile(VersionFilePath);
503 if (HIPVersionArg.empty() && VersionFile)
504 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
507 HasHIPRuntime =
true;
512 if (!Candidate.StrictChecking) {
513 HasHIPRuntime =
true;
517 HasHIPRuntime =
false;
522 OS <<
"Found HIP installation: " << InstallPath <<
", version "
523 << DetectedVersion <<
'\n';
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);
532 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
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));
553 const auto HandleHipStdPar = [=, &DriverArgs, &CC1Args]() {
555 auto &FS = D.getVFS();
558 if (!HIPStdParPathArg.empty() ||
559 !FS.exists(Inc +
"/thrust/system/hip/hipstdpar/hipstdpar_lib.hpp")) {
560 D.Diag(diag::err_drv_no_hipstdpar_lib);
563 if (!HasRocThrustLibrary && !FS.exists(Inc +
"/thrust")) {
564 D.Diag(diag::err_drv_no_hipstdpar_thrust_lib);
567 if (!HasRocPrimLibrary && !FS.exists(Inc +
"/rocprim")) {
568 D.Diag(diag::err_drv_no_hipstdpar_prim_lib);
571 const char *ThrustPath;
572 if (HasRocThrustLibrary)
573 ThrustPath = DriverArgs.MakeArgString(HIPRocThrustPathArg);
575 ThrustPath = DriverArgs.MakeArgString(Inc +
"/thrust");
577 const char *HIPStdParPath;
579 HIPStdParPath = DriverArgs.MakeArgString(HIPStdParPathArg);
581 HIPStdParPath = DriverArgs.MakeArgString(StringRef(ThrustPath) +
582 "/system/hip/hipstdpar");
584 const char *PrimPath;
585 if (HasRocPrimLibrary)
586 PrimPath = DriverArgs.MakeArgString(HIPRocPrimPathArg);
588 PrimPath = DriverArgs.MakeArgString(
getIncludePath() +
"/rocprim");
590 CC1Args.append({
"-idirafter", ThrustPath,
"-idirafter", PrimPath,
591 "-idirafter", HIPStdParPath,
"-include",
592 "hipstdpar_lib.hpp"});
595 if (!DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
604 D.Diag(diag::err_drv_no_hip_runtime);
608 CC1Args.push_back(
"-idirafter");
610 if (UsesRuntimeWrapper)
611 CC1Args.append({
"-include",
"__clang_hip_runtime_wrapper.h"});
620 const char *LinkingOutput)
const {
622 ArgStringList CmdArgs;
623 if (!Args.hasArg(options::OPT_r)) {
624 CmdArgs.push_back(
"--no-undefined");
625 CmdArgs.push_back(
"-shared");
628 if (
C.getDriver().isUsingLTO()) {
629 const bool ThinLTO = (
C.getDriver().getLTOMode() ==
LTOK_Thin);
631 }
else if (Args.hasArg(options::OPT_mcpu_EQ)) {
632 CmdArgs.push_back(Args.MakeArgString(
633 "-plugin-opt=mcpu=" +
635 Args.getLastArgValue(options::OPT_mcpu_EQ))));
639 Args.AddAllArgs(CmdArgs, options::OPT_L);
643 std::vector<StringRef> Features;
646 if (!Features.empty()) {
648 Args.MakeArgString(
"-plugin-opt=-mattr=" + llvm::join(Features,
",")));
651 if (Args.hasArg(options::OPT_stdlib))
652 CmdArgs.append({
"-lc",
"-lm"});
653 if (Args.hasArg(options::OPT_startfiles)) {
654 std::optional<std::string> IncludePath =
getToolChain().getStdlibPath();
656 IncludePath =
"/lib";
658 llvm::sys::path::append(P,
"crt1.o");
659 CmdArgs.push_back(Args.MakeArgString(P));
662 CmdArgs.push_back(
"-o");
664 C.addCommand(std::make_unique<Command>(
666 CmdArgs, Inputs, Output));
670 const llvm::Triple &Triple,
671 const llvm::opt::ArgList &Args,
672 std::vector<StringRef> &Features) {
676 if (Args.hasArg(options::OPT_mcpu_EQ))
677 TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
678 else if (Args.hasArg(options::OPT_march_EQ))
679 TargetID = Args.getLastArgValue(options::OPT_march_EQ);
680 if (!TargetID.empty()) {
681 llvm::StringMap<bool> FeatureMap;
682 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
683 if (OptionalGpuArch) {
684 StringRef GpuArch = *OptionalGpuArch;
690 auto Pos = FeatureMap.find(
Feature);
691 if (Pos == FeatureMap.end())
693 Features.push_back(Args.MakeArgStringRef(
694 (Twine(Pos->second ?
"+" :
"-") +
Feature).str()));
699 if (Args.hasFlag(options::OPT_mwavefrontsize64,
700 options::OPT_mno_wavefrontsize64,
false))
701 Features.push_back(
"+wavefrontsize64");
703 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
704 options::OPT_mno_amdgpu_precise_memory_op,
false))
705 Features.push_back(
"+precise-memory");
708 options::OPT_m_amdgpu_Features_Group);
716 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
732 DerivedArgList *DAL =
738 DAL =
new DerivedArgList(Args.getBaseArgs());
744 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
745 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
746 DAL->eraseArg(options::OPT_mcpu_EQ);
750 << llvm::Triple::getArchTypeName(
getArch())
751 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
753 auto &GPUs = *GPUsOrErr;
754 if (GPUs.size() > 1) {
756 << llvm::Triple::getArchTypeName(
getArch())
757 << llvm::join(GPUs,
", ") <<
"-mcpu";
759 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
760 Args.MakeArgString(GPUs.front()));
766 if (Args.getLastArgValue(options::OPT_x) !=
"cl")
770 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
771 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
773 : options::OPT_m32));
777 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
779 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
787 llvm::AMDGPU::GPUKind Kind) {
790 if (Kind == llvm::AMDGPU::GK_NONE)
793 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
797 const bool BothDenormAndFMAFast =
798 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
799 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
800 return !BothDenormAndFMAFast;
804 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
805 const llvm::fltSemantics *FPType)
const {
807 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
808 return llvm::DenormalMode::getIEEE();
813 auto Kind = llvm::AMDGPU::parseArchAMDGCN(
Arch);
814 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
815 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
816 options::OPT_fno_gpu_flush_denormals_to_zero,
818 return llvm::DenormalMode::getPreserveSign();
820 return llvm::DenormalMode::getIEEE();
823 const StringRef GpuArch =
getGPUArch(DriverArgs);
824 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
828 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
833 return DAZ ? llvm::DenormalMode::getPreserveSign() :
834 llvm::DenormalMode::getIEEE();
838 llvm::AMDGPU::GPUKind Kind) {
839 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
840 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
842 return !HasWave32 || DriverArgs.hasFlag(
843 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
855 const llvm::opt::ArgList &DriverArgs,
856 llvm::opt::ArgStringList &CC1Args,
860 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
861 options::OPT_fvisibility_ms_compat)) {
862 CC1Args.push_back(
"-fvisibility=hidden");
863 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
877 !DriverArgs.hasArg(options::OPT_disable_llvm_optzns))
878 CC1Args.push_back(
"-disable-llvm-optzns");
887 CC1Args.push_back(
"-Werror=atomic-alignment");
891 ArgStringList &CC1Args)
const {
892 if (DriverArgs.hasArg(options::OPT_nostdinc) ||
893 DriverArgs.hasArg(options::OPT_nostdlibinc))
903 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
908 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
909 if (TargetID.empty())
910 return {std::nullopt, std::nullopt, std::nullopt};
912 llvm::StringMap<bool> FeatureMap;
914 if (!OptionalGpuArch)
915 return {TargetID.str(), std::nullopt, std::nullopt};
917 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
921 const llvm::opt::ArgList &DriverArgs)
const {
923 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
925 << *PTID.OptionalTargetID;
933 if (Arg *A = Args.getLastArg(options::OPT_offload_arch_tool_EQ))
934 Program = A->getValue();
940 return StdoutOrErr.takeError();
943 for (StringRef
Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
945 GPUArchs.push_back(
Arch.str());
947 if (GPUArchs.empty())
948 return llvm::createStringError(std::error_code(),
949 "No AMD GPU detected in the system");
951 return std::move(GPUArchs);
955 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
958 DeviceOffloadingKind);
963 DriverArgs.hasArg(options::OPT_nostdlib))
966 if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
971 const StringRef GpuArch =
getGPUArch(DriverArgs);
972 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
973 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
987 DriverArgs, LibDeviceFile, GpuArch, DeviceOffloadingKind,
990 for (
auto [BCFile, Internalize] : BCLibs) {
992 CC1Args.push_back(
"-mlink-builtin-bitcode");
994 CC1Args.push_back(
"-mlink-bitcode-file");
995 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
1000 StringRef GPUArch, StringRef LibDeviceFile,
1003 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
1006 if (LibDeviceFile.empty()) {
1007 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
1014 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString() << 0;
1016 D.Diag(diag::err_drv_no_rocm_device_lib)
1017 << 2 << ABIVer.
toString() << 1 <<
"6.3";
1025 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
1027 const bool NeedsASanRT)
const {
1030 CommonBitcodeLibsPreferences Pref{D, DriverArgs, GPUArch,
1031 DeviceOffloadingKind, NeedsASanRT};
1034 bool Internalize =
true) {
1036 BCLibs.emplace_back(BCLib);
1038 auto AddSanBCLibs = [&]() {
1047 else if (Pref.GPUSan && Pref.IsOpenMP)
1054 AddBCLib(LibDeviceFile);
1056 if (!ABIVerPath.empty())
1057 AddBCLib(ABIVerPath);
1064 const llvm::opt::ArgList &DriverArgs,
const std::string &GPUArch,
1066 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1067 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1077 DriverArgs, LibDeviceFile, GPUArch, DeviceOffloadingKind,
1082 const ToolChain &TC,
const llvm::opt::ArgList &DriverArgs,
1083 StringRef TargetID,
const llvm::opt::Arg *A)
const {
1085 bool IsExplicitDevice =
1086 A->getBaseArg().getOption().matches(options::OPT_Xarch_device);
1089 llvm::StringRef Processor =
1091 auto ProcKind = TC.
getTriple().isAMDGCN()
1092 ? llvm::AMDGPU::parseArchAMDGCN(Processor)
1093 : llvm::AMDGPU::parseArchR600(Processor);
1094 auto Features = TC.
getTriple().isAMDGCN()
1095 ? llvm::AMDGPU::getArchAttrAMDGCN(ProcKind)
1096 : llvm::AMDGPU::getArchAttrR600(ProcKind);
1097 if (Features & llvm::AMDGPU::FEATURE_XNACK_ALWAYS)
1101 llvm::StringMap<bool> FeatureMap;
1103 assert(OptionalGpuArch &&
"Invalid Target ID");
1104 (void)OptionalGpuArch;
1105 auto Loc = FeatureMap.find(
"xnack");
1106 if (Loc == FeatureMap.end() || !Loc->second) {
1107 if (IsExplicitDevice) {
1109 clang::diag::err_drv_unsupported_option_for_offload_arch_req_feature)
1110 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
1113 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1114 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
static void Append(char *Start, char *End, char *&Buffer, unsigned &BufferSize, unsigned &BufferCapacity)
__device__ __2f16 float c
const char * getOffloadingArch() const
OffloadKind getOffloadingDeviceKind() const
Compilation - A set of tasks to perform for a single driver invocation.
Driver - Encapsulate logic for constructing compilation processes from a set of gcc-driver-like comma...
DiagnosticsEngine & getDiags() const
llvm::Expected< std::unique_ptr< llvm::MemoryBuffer > > executeProgram(llvm::ArrayRef< llvm::StringRef > Args) const
DiagnosticBuilder Diag(unsigned DiagID) const
const llvm::opt::OptTable & getOpts() const
StringRef getCorrectlyRoundedSqrtPath(bool Enabled) const
StringRef getIncludePath() const
Get the detected path to Rocm's bin directory.
StringRef getUnsafeMathPath(bool Enabled) const
StringRef getOCMLPath() const
StringRef getAsanRTLPath() const
Returns empty string of Asan runtime library is not available.
RocmInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args, bool DetectHIPRuntime=true, bool DetectDeviceLib=false)
StringRef getOCKLPath() const
StringRef getFiniteOnlyPath(bool Enabled) const
bool hasDeviceLibrary() const
Check whether we detected a valid ROCm device library.
StringRef getDenormalsAreZeroPath(bool Enabled) const
bool checkCommonBitcodeLibs(StringRef GPUArch, StringRef LibDeviceFile, DeviceLibABIVersion ABIVer) const
Check file paths of default bitcode libraries common to AMDGPU based toolchains.
bool hasHIPStdParLibrary() const
Check whether we detected a valid HIP STDPAR Acceleration library.
StringRef getABIVersionPath(DeviceLibABIVersion ABIVer) const
llvm::SmallVector< ToolChain::BitCodeLibraryInfo, 12 > getCommonBitcodeLibs(const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, StringRef GPUArch, const Action::OffloadKind DeviceOffloadingKind, const bool NeedsASanRT) const
Get file paths of default bitcode libraries common to AMDGPU based toolchains.
bool hasHIPRuntime() const
Check whether we detected a valid HIP runtime.
void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const
void detectDeviceLibrary()
StringRef getWavefrontSize64Path(bool Enabled) const
void print(raw_ostream &OS) const
Print information about the detected ROCm installation.
SmallVector< InputInfo, 4 > InputInfoList
The JSON file list parser is used to communicate input to InstallAPI.
if(T->getSizeExpr()) TRY_TO(TraverseStmt(const_cast< Expr * >(T -> getSizeExpr())))
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.
llvm::StringRef getProcessorFromTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch)
Get processor name from target ID.
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.
Diagnostic wrappers for TextAPI types for error reporting.
ABI version of device library.
unsigned getAsCodeObjectVersion() const
static DeviceLibABIVersion fromCodeObjectVersion(unsigned CodeObjectVersion)
bool requiresLibrary()
Whether ABI version bc file is requested.
static constexpr ResponseFileSupport AtFileCurCP()