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);
86 GPUSan = (DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
87 options::OPT_fno_gpu_sanitize,
true) &&
91void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) {
92 assert(!Path.empty());
94 const StringRef Suffix(
".bc");
95 const StringRef Suffix2(
".amdgcn.bc");
98 for (llvm::vfs::directory_iterator LI = D.getVFS().dir_begin(Path, EC), LE;
99 !EC && LI != LE; LI = LI.increment(EC)) {
100 StringRef FilePath = LI->path();
101 StringRef
FileName = llvm::sys::path::filename(FilePath);
107 BaseName =
FileName.drop_back(Suffix2.size());
108 else if (
FileName.ends_with(Suffix))
109 BaseName =
FileName.drop_back(Suffix.size());
111 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
112 if (BaseName ==
"ocml") {
114 }
else if (BaseName ==
"ockl") {
116 }
else if (BaseName ==
"opencl") {
118 }
else if (BaseName ==
"asanrtl") {
120 }
else if (BaseName ==
"oclc_finite_only_off") {
121 FiniteOnly.Off = FilePath;
122 }
else if (BaseName ==
"oclc_finite_only_on") {
123 FiniteOnly.On = FilePath;
124 }
else if (BaseName ==
"oclc_unsafe_math_on") {
125 UnsafeMath.On = FilePath;
126 }
else if (BaseName ==
"oclc_unsafe_math_off") {
127 UnsafeMath.Off = FilePath;
128 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
129 WavefrontSize64.On = FilePath;
130 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
131 WavefrontSize64.Off = FilePath;
132 }
else if (BaseName.starts_with(ABIVersionPrefix)) {
133 unsigned ABIVersionNumber;
134 if (BaseName.drop_front(ABIVersionPrefix.size())
135 .getAsInteger(0, ABIVersionNumber))
137 ABIVersionMap[ABIVersionNumber] = FilePath.str();
141 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
142 if (!BaseName.starts_with(DeviceLibPrefix))
145 StringRef IsaVersionNumber =
146 BaseName.drop_front(DeviceLibPrefix.size());
148 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
150 LibDeviceMap.insert({GfxName.toStringRef(Tmp), FilePath.str()});
157bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
158 SmallVector<StringRef, 4> VersionParts;
159 V.split(VersionParts,
'\n');
160 unsigned Major = ~0U;
161 unsigned Minor = ~0U;
162 for (
auto Part : VersionParts) {
163 auto Splits = Part.rtrim().split(
'=');
164 if (Splits.first ==
"HIP_VERSION_MAJOR") {
165 if (Splits.second.getAsInteger(0, Major))
167 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
168 if (Splits.second.getAsInteger(0, Minor))
170 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
171 VersionPatch = Splits.second.str();
173 if (Major == ~0U || Minor == ~0U)
175 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
177 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
183const SmallVectorImpl<RocmInstallationDetector::Candidate> &
184RocmInstallationDetector::getInstallationPathCandidates() {
187 if (!ROCmSearchDirs.empty())
188 return ROCmSearchDirs;
190 auto DoPrintROCmSearchDirs = [&]() {
191 if (PrintROCmSearchDirs)
192 for (
auto Cand : ROCmSearchDirs) {
193 llvm::errs() <<
"ROCm installation search path: " << Cand.Path <<
'\n';
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;
213 StringRef InstallDir = D.Dir;
218 auto DeduceROCmPath = [](StringRef ClangPath) {
220 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
221 StringRef ParentName = llvm::sys::path::filename(ParentDir);
224 if (ParentName ==
"bin") {
225 ParentDir = llvm::sys::path::parent_path(ParentDir);
226 ParentName = llvm::sys::path::filename(ParentDir);
231 if (ParentName ==
"llvm" || ParentName.starts_with(
"aomp")) {
232 ParentDir = llvm::sys::path::parent_path(ParentDir);
233 ParentName = llvm::sys::path::filename(ParentDir);
237 if (ParentName ==
"lib")
238 ParentDir = llvm::sys::path::parent_path(ParentDir);
241 return Candidate(ParentDir.str(),
true);
246 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
250 llvm::SmallString<256> RealClangPath;
251 llvm::sys::fs::real_path(D.getClangProgramPath(), RealClangPath);
252 auto ParentPath = llvm::sys::path::parent_path(RealClangPath);
253 if (ParentPath != InstallDir)
254 ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));
257 auto ClangRoot = llvm::sys::path::parent_path(InstallDir);
258 auto RealClangRoot = llvm::sys::path::parent_path(ParentPath);
259 ROCmSearchDirs.emplace_back(ClangRoot.str(),
true);
260 if (RealClangRoot != ClangRoot)
261 ROCmSearchDirs.emplace_back(RealClangRoot.str(),
true);
262 ROCmSearchDirs.emplace_back(D.ResourceDir,
265 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/rocm",
270 std::string LatestROCm;
271 llvm::VersionTuple LatestVer;
273 auto GetROCmVersion = [](StringRef DirName) {
274 llvm::VersionTuple
V;
275 std::string VerStr = DirName.drop_front(strlen(
"rocm-")).str();
278 llvm::replace(VerStr,
'-',
'.');
282 for (llvm::vfs::directory_iterator
283 File = D.getVFS().dir_begin(D.SysRoot +
"/opt", EC),
285 File != FileEnd && !EC;
File.increment(EC)) {
286 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
289 if (LatestROCm.empty()) {
291 LatestVer = GetROCmVersion(LatestROCm);
294 auto Ver = GetROCmVersion(
FileName);
295 if (LatestVer < Ver) {
300 if (!LatestROCm.empty())
301 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/" + LatestROCm,
304 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr/local",
306 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr",
309 DoPrintROCmSearchDirs();
310 return ROCmSearchDirs;
314 const Driver &D,
const llvm::Triple &HostTriple,
315 const llvm::opt::ArgList &Args,
bool DetectHIPRuntime)
317 Verbose = Args.hasArg(options::OPT_v);
318 RocmPathArg = Args.getLastArgValue(options::OPT_rocm_path_EQ);
319 PrintROCmSearchDirs = Args.hasArg(options::OPT_print_rocm_search_dirs);
320 RocmDeviceLibPathArg =
321 Args.getAllArgValues(options::OPT_rocm_device_lib_path_EQ);
322 HIPPathArg = Args.getLastArgValue(options::OPT_hip_path_EQ);
323 HIPStdParPathArg = Args.getLastArgValue(options::OPT_hipstdpar_path_EQ);
324 HasHIPStdParLibrary =
325 !HIPStdParPathArg.empty() && D.getVFS().exists(HIPStdParPathArg +
326 "/hipstdpar_lib.hpp");
327 HIPRocThrustPathArg =
328 Args.getLastArgValue(options::OPT_hipstdpar_thrust_path_EQ);
329 HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&
330 D.getVFS().exists(HIPRocThrustPathArg +
"/thrust");
331 HIPRocPrimPathArg = Args.getLastArgValue(options::OPT_hipstdpar_prim_path_EQ);
332 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
333 D.getVFS().exists(HIPRocPrimPathArg +
"/rocprim");
335 if (
auto *A = Args.getLastArg(options::OPT_hip_version_EQ)) {
336 HIPVersionArg = A->getValue();
337 unsigned Major = ~0U;
338 unsigned Minor = ~0U;
339 SmallVector<StringRef, 3> Parts;
340 HIPVersionArg.split(Parts,
'.');
342 Parts[0].getAsInteger(0, Major);
343 if (Parts.size() > 1)
344 Parts[1].getAsInteger(0, Minor);
345 if (Parts.size() > 2)
346 VersionPatch = Parts[2].str();
347 if (VersionPatch.empty())
349 if (Major != ~0U && Minor == ~0U)
351 if (Major == ~0U || Minor == ~0U)
352 D.Diag(diag::err_drv_invalid_value)
353 << A->getAsString(Args) << HIPVersionArg;
355 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
357 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
359 VersionPatch = DefaultVersionPatch;
361 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
362 DetectedVersion = (Twine(DefaultVersionMajor) +
"." +
363 Twine(DefaultVersionMinor) +
"." + VersionPatch)
367 if (DetectHIPRuntime)
372 assert(LibDevicePath.empty());
374 if (!RocmDeviceLibPathArg.empty())
375 LibDevicePath = RocmDeviceLibPathArg.back();
376 else if (std::optional<std::string> LibPathEnv =
377 llvm::sys::Process::GetEnv(
"HIP_DEVICE_LIB_PATH"))
378 LibDevicePath = std::move(*LibPathEnv);
380 auto &FS = D.getVFS();
381 if (!LibDevicePath.empty()) {
385 if (!FS.exists(LibDevicePath))
388 scanLibDevicePath(LibDevicePath);
389 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
394 auto CheckDeviceLib = [&](StringRef Path,
bool StrictChecking) {
395 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
396 if (CheckLibDevice && !FS.exists(Path))
399 scanLibDevicePath(Path);
401 if (!NoBuiltinLibs) {
403 if (!allGenericLibsValid())
408 if (LibDeviceMap.empty())
415 LibDevicePath = D.ResourceDir;
416 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
417 "amdgcn",
"bitcode");
418 HasDeviceLibrary = CheckDeviceLib(LibDevicePath,
true);
419 if (HasDeviceLibrary)
424 auto &ROCmDirs = getInstallationPathCandidates();
425 for (
const auto &Candidate : ROCmDirs) {
426 LibDevicePath = Candidate.Path;
427 llvm::sys::path::append(LibDevicePath,
"amdgcn",
"bitcode");
428 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);
429 if (HasDeviceLibrary)
436 if (!HIPPathArg.empty())
437 HIPSearchDirs.emplace_back(HIPPathArg.str());
438 else if (std::optional<std::string> HIPPathEnv =
439 llvm::sys::Process::GetEnv(
"HIP_PATH")) {
440 if (!HIPPathEnv->empty())
441 HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));
443 if (HIPSearchDirs.empty())
444 HIPSearchDirs.append(getInstallationPathCandidates());
445 auto &FS = D.getVFS();
447 for (
const auto &Candidate : HIPSearchDirs) {
448 InstallPath = Candidate.Path;
449 if (InstallPath.empty() || !FS.exists(InstallPath))
452 BinPath = InstallPath;
453 llvm::sys::path::append(BinPath,
"bin");
454 IncludePath = InstallPath;
455 llvm::sys::path::append(IncludePath,
"include");
456 LibPath = InstallPath;
457 llvm::sys::path::append(LibPath,
"lib");
458 SharePath = InstallPath;
459 llvm::sys::path::append(SharePath,
"share");
462 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
463 llvm::sys::path::append(ParentSharePath,
"share");
466 const Twine &
c =
"",
const Twine &d =
"") {
468 llvm::sys::path::append(newpath, a,
b,
c, d);
472 std::vector<SmallString<0>> VersionFilePaths = {
473 Append(SharePath,
"hip",
"version"),
474 InstallPath != D.SysRoot +
"/usr/local"
475 ?
Append(ParentSharePath,
"hip",
"version")
477 Append(BinPath,
".hipVersion")};
479 for (
const auto &VersionFilePath : VersionFilePaths) {
480 if (VersionFilePath.empty())
482 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
483 FS.getBufferForFile(VersionFilePath);
486 if (HIPVersionArg.empty() && VersionFile)
487 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
490 HasHIPRuntime =
true;
495 if (!Candidate.StrictChecking) {
496 HasHIPRuntime =
true;
500 HasHIPRuntime =
false;
505 OS <<
"Found HIP installation: " << InstallPath <<
", version "
506 << DetectedVersion <<
'\n';
510 ArgStringList &CC1Args)
const {
511 bool UsesRuntimeWrapper = VersionMajorMinor > llvm::VersionTuple(3, 5) &&
512 !DriverArgs.hasArg(options::OPT_nohipwrapperinc);
513 bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
515 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
530 if (UsesRuntimeWrapper)
531 llvm::sys::path::append(P,
"include",
"cuda_wrappers");
532 CC1Args.push_back(
"-internal-isystem");
533 CC1Args.push_back(DriverArgs.MakeArgString(P));
536 const auto HandleHipStdPar = [=, &DriverArgs, &CC1Args]() {
538 auto &FS = D.getVFS();
541 if (!HIPStdParPathArg.empty() ||
542 !FS.exists(Inc +
"/thrust/system/hip/hipstdpar/hipstdpar_lib.hpp")) {
543 D.Diag(diag::err_drv_no_hipstdpar_lib);
546 if (!HasRocThrustLibrary && !FS.exists(Inc +
"/thrust")) {
547 D.Diag(diag::err_drv_no_hipstdpar_thrust_lib);
550 if (!HasRocPrimLibrary && !FS.exists(Inc +
"/rocprim")) {
551 D.Diag(diag::err_drv_no_hipstdpar_prim_lib);
554 const char *ThrustPath;
555 if (HasRocThrustLibrary)
556 ThrustPath = DriverArgs.MakeArgString(HIPRocThrustPathArg);
558 ThrustPath = DriverArgs.MakeArgString(Inc +
"/thrust");
560 const char *HIPStdParPath;
562 HIPStdParPath = DriverArgs.MakeArgString(HIPStdParPathArg);
564 HIPStdParPath = DriverArgs.MakeArgString(StringRef(ThrustPath) +
565 "/system/hip/hipstdpar");
567 const char *PrimPath;
568 if (HasRocPrimLibrary)
569 PrimPath = DriverArgs.MakeArgString(HIPRocPrimPathArg);
571 PrimPath = DriverArgs.MakeArgString(
getIncludePath() +
"/rocprim");
573 CC1Args.append({
"-idirafter", ThrustPath,
"-idirafter", PrimPath,
574 "-idirafter", HIPStdParPath,
"-include",
575 "hipstdpar_lib.hpp"});
578 if (!DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
587 D.Diag(diag::err_drv_no_hip_runtime);
591 CC1Args.push_back(
"-idirafter");
593 if (UsesRuntimeWrapper)
594 CC1Args.append({
"-include",
"__clang_hip_runtime_wrapper.h"});
603 const char *LinkingOutput)
const {
605 ArgStringList CmdArgs;
606 if (!Args.hasArg(options::OPT_r)) {
607 CmdArgs.push_back(
"--no-undefined");
608 CmdArgs.push_back(
"-shared");
611 if (
C.getDriver().isUsingLTO()) {
612 const bool ThinLTO = (
C.getDriver().getLTOMode() ==
LTOK_Thin);
614 }
else if (Args.hasArg(options::OPT_mcpu_EQ)) {
615 CmdArgs.push_back(Args.MakeArgString(
616 "-plugin-opt=mcpu=" +
618 Args.getLastArgValue(options::OPT_mcpu_EQ))));
622 Args.AddAllArgs(CmdArgs, options::OPT_L);
626 std::vector<StringRef> Features;
629 if (!Features.empty()) {
631 Args.MakeArgString(
"-plugin-opt=-mattr=" + llvm::join(Features,
",")));
636 if (Args.hasArg(options::OPT_stdlib))
637 CmdArgs.append({
"-lc",
"-lm"});
638 if (Args.hasArg(options::OPT_startfiles)) {
639 std::optional<std::string> IncludePath =
getToolChain().getStdlibPath();
641 IncludePath =
"/lib";
643 llvm::sys::path::append(P,
"crt1.o");
644 CmdArgs.push_back(Args.MakeArgString(P));
647 CmdArgs.push_back(
"-o");
649 C.addCommand(std::make_unique<Command>(
651 CmdArgs, Inputs, Output));
655 const llvm::Triple &Triple,
656 const llvm::opt::ArgList &Args,
657 std::vector<StringRef> &Features) {
661 if (Args.hasArg(options::OPT_mcpu_EQ))
662 TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
663 else if (Args.hasArg(options::OPT_march_EQ))
664 TargetID = Args.getLastArgValue(options::OPT_march_EQ);
665 if (!TargetID.empty()) {
666 llvm::StringMap<bool> FeatureMap;
667 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
668 if (OptionalGpuArch) {
669 StringRef GpuArch = *OptionalGpuArch;
675 auto Pos = FeatureMap.find(
Feature);
676 if (Pos == FeatureMap.end())
678 Features.push_back(Args.MakeArgStringRef(
679 (Twine(Pos->second ?
"+" :
"-") +
Feature).str()));
684 if (Args.hasFlag(options::OPT_mwavefrontsize64,
685 options::OPT_mno_wavefrontsize64,
false))
686 Features.push_back(
"+wavefrontsize64");
688 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
689 options::OPT_mno_amdgpu_precise_memory_op,
false))
690 Features.push_back(
"+precise-memory");
693 options::OPT_m_amdgpu_Features_Group);
701 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
717 DerivedArgList *DAL =
723 DAL =
new DerivedArgList(Args.getBaseArgs());
729 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
730 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
731 DAL->eraseArg(options::OPT_mcpu_EQ);
735 << llvm::Triple::getArchTypeName(
getArch())
736 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
738 auto &GPUs = *GPUsOrErr;
739 if (GPUs.size() > 1) {
741 << llvm::Triple::getArchTypeName(
getArch())
742 << llvm::join(GPUs,
", ") <<
"-mcpu";
744 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
745 Args.MakeArgString(GPUs.front()));
751 if (Args.getLastArgValue(options::OPT_x) !=
"cl")
755 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
756 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
758 : options::OPT_m32));
762 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
764 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
772 llvm::AMDGPU::GPUKind Kind) {
775 if (Kind == llvm::AMDGPU::GK_NONE)
778 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
782 const bool BothDenormAndFMAFast =
783 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
784 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
785 return !BothDenormAndFMAFast;
789 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
790 const llvm::fltSemantics *FPType)
const {
792 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
793 return llvm::DenormalMode::getIEEE();
798 auto Kind = llvm::AMDGPU::parseArchAMDGCN(
Arch);
799 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
800 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
801 options::OPT_fno_gpu_flush_denormals_to_zero,
803 return llvm::DenormalMode::getPreserveSign();
805 return llvm::DenormalMode::getIEEE();
808 const StringRef GpuArch =
getGPUArch(DriverArgs);
809 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
813 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
818 return DAZ ? llvm::DenormalMode::getPreserveSign() :
819 llvm::DenormalMode::getIEEE();
823 llvm::AMDGPU::GPUKind Kind) {
824 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
825 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
827 return !HasWave32 || DriverArgs.hasFlag(
828 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
836 if (Triple.getEnvironment() != llvm::Triple::LLVM)
841 const llvm::opt::ArgList &DriverArgs,
842 llvm::opt::ArgStringList &CC1Args,
846 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
847 options::OPT_fvisibility_ms_compat)) {
848 CC1Args.push_back(
"-fvisibility=hidden");
849 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
863 !DriverArgs.hasArg(options::OPT_disable_llvm_optzns))
864 CC1Args.push_back(
"-disable-llvm-optzns");
873 CC1Args.push_back(
"-Werror=atomic-alignment");
877 ArgStringList &CC1Args)
const {
878 if (DriverArgs.hasArg(options::OPT_nostdinc) ||
879 DriverArgs.hasArg(options::OPT_nostdlibinc))
889 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
894 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
895 if (TargetID.empty())
896 return {std::nullopt, std::nullopt, std::nullopt};
898 llvm::StringMap<bool> FeatureMap;
900 if (!OptionalGpuArch)
901 return {TargetID.str(), std::nullopt, std::nullopt};
903 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
907 const llvm::opt::ArgList &DriverArgs)
const {
909 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
911 << *PTID.OptionalTargetID;
919 if (Arg *A = Args.getLastArg(options::OPT_offload_arch_tool_EQ))
920 Program = A->getValue();
926 return StdoutOrErr.takeError();
929 for (StringRef
Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
931 GPUArchs.push_back(
Arch.str());
933 if (GPUArchs.empty())
934 return llvm::createStringError(std::error_code(),
935 "No AMD GPU detected in the system");
937 return std::move(GPUArchs);
941 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
944 DeviceOffloadingKind);
949 DriverArgs.hasArg(options::OPT_nostdlib))
952 if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
964 if (TT.getEnvironment() == llvm::Triple::LLVM)
968 const StringRef GpuArch =
getGPUArch(DriverArgs);
969 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
970 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
984 DriverArgs, LibDeviceFile, GpuArch, DeviceOffloadingKind,
987 for (
auto [BCFile, Internalize] : BCLibs) {
989 CC1Args.push_back(
"-mlink-builtin-bitcode");
991 CC1Args.push_back(
"-mlink-bitcode-file");
992 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
997 StringRef GPUArch, StringRef LibDeviceFile,
1000 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
1003 if (LibDeviceFile.empty()) {
1004 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
1011 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString() << 0;
1013 D.Diag(diag::err_drv_no_rocm_device_lib)
1014 << 2 << ABIVer.
toString() << 1 <<
"6.3";
1022 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
1024 const bool NeedsASanRT)
const {
1027 CommonBitcodeLibsPreferences Pref{D, DriverArgs, GPUArch,
1028 DeviceOffloadingKind, NeedsASanRT};
1031 bool Internalize =
true) {
1032 if (!BCLib.
Path.empty()) {
1034 BCLibs.emplace_back(BCLib);
1037 auto AddSanBCLibs = [&]() {
1046 else if (Pref.GPUSan && Pref.IsOpenMP)
1051 AddBCLib(LibDeviceFile);
1053 if (!ABIVerPath.empty())
1054 AddBCLib(ABIVerPath);
1061 const llvm::opt::ArgList &DriverArgs,
const std::string &GPUArch,
1063 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1064 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1074 DriverArgs, LibDeviceFile, GPUArch, DeviceOffloadingKind,
1079 const ToolChain &TC,
const llvm::opt::ArgList &DriverArgs,
1080 StringRef TargetID,
const llvm::opt::Arg *A)
const {
1082 bool IsExplicitDevice =
1083 A->getBaseArg().getOption().matches(options::OPT_Xarch_device);
1086 llvm::StringRef Processor =
1088 auto ProcKind = TC.
getTriple().isAMDGCN()
1089 ? llvm::AMDGPU::parseArchAMDGCN(Processor)
1090 : llvm::AMDGPU::parseArchR600(Processor);
1091 auto Features = TC.
getTriple().isAMDGCN()
1092 ? llvm::AMDGPU::getArchAttrAMDGCN(ProcKind)
1093 : llvm::AMDGPU::getArchAttrR600(ProcKind);
1094 if (Features & llvm::AMDGPU::FEATURE_XNACK_ALWAYS)
1098 llvm::StringMap<bool> FeatureMap;
1100 assert(OptionalGpuArch &&
"Invalid Target ID");
1101 (void)OptionalGpuArch;
1102 auto Loc = FeatureMap.find(
"xnack");
1103 if (Loc == FeatureMap.end() || !Loc->second) {
1104 if (IsExplicitDevice) {
1106 clang::diag::err_drv_unsupported_option_for_offload_arch_req_feature)
1107 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
1110 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1111 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
static StringRef getTriple(const Command &Job)
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 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)
StringRef getOCKLPath() const
StringRef getFiniteOnlyPath(bool Enabled) const
bool hasDeviceLibrary() const
Check whether we detected a valid ROCm device library.
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()