12#include "clang/Config/config.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"
27#include <system_error>
42RocmInstallationDetector::findSPACKPackage(
const Candidate &Cand,
43 StringRef PackageName) {
47 std::string Prefix = Twine(PackageName +
"-" + Cand.SPACKReleaseStr).str();
49 for (llvm::vfs::directory_iterator
File = D.
getVFS().dir_begin(Cand.Path, EC),
51 File != FileEnd && !EC;
File.increment(EC)) {
52 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
55 if (SubDirs.size() > 1)
59 if (SubDirs.size() == 1) {
60 auto PackagePath = Cand.Path;
61 llvm::sys::path::append(PackagePath, SubDirs[0]);
64 if (SubDirs.size() == 0 && Verbose) {
65 llvm::errs() <<
"SPACK package " << Prefix <<
" not found at " << Cand.Path
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";
77void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef
Path) {
78 assert(!
Path.empty());
80 const StringRef Suffix(
".bc");
81 const StringRef Suffix2(
".amdgcn.bc");
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);
93 BaseName =
FileName.drop_back(Suffix2.size());
95 BaseName =
FileName.drop_back(Suffix.size());
97 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
98 if (BaseName ==
"ocml") {
100 }
else if (BaseName ==
"ockl") {
102 }
else if (BaseName ==
"opencl") {
104 }
else if (BaseName ==
"hip") {
106 }
else if (BaseName ==
"asanrtl") {
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(0, ABIVersionNumber))
133 ABIVersionMap[ABIVersionNumber] = FilePath.str();
137 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
138 if (!BaseName.starts_with(DeviceLibPrefix))
141 StringRef IsaVersionNumber =
142 BaseName.drop_front(DeviceLibPrefix.size());
144 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
147 std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
154bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
156 V.split(VersionParts,
'\n');
157 unsigned Major = ~0
U;
158 unsigned Minor = ~0
U;
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))
164 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
165 if (Splits.second.getAsInteger(0, Minor))
167 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
168 VersionPatch = Splits.second.str();
170 if (Major == ~0
U || Minor == ~0
U)
172 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
174 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
181RocmInstallationDetector::getInstallationPathCandidates() {
184 if (!ROCmSearchDirs.empty())
185 return ROCmSearchDirs;
187 auto DoPrintROCmSearchDirs = [&]() {
188 if (PrintROCmSearchDirs)
189 for (
auto Cand : ROCmSearchDirs) {
190 llvm::errs() <<
"ROCm installation search path";
192 llvm::errs() <<
" (Spack " << Cand.SPACKReleaseStr <<
")";
193 llvm::errs() <<
": " << 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);
234 if (ParentName.starts_with(
"llvm-amdgpu-")) {
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(),
true,
247 if (ParentName ==
"llvm" || ParentName.starts_with(
"aomp"))
248 ParentDir = llvm::sys::path::parent_path(ParentDir);
250 return Candidate(ParentDir.str(),
true);
255 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
261 auto ParentPath = llvm::sys::path::parent_path(RealClangPath);
262 if (ParentPath != InstallDir)
263 ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));
266 auto ClangRoot = llvm::sys::path::parent_path(InstallDir);
267 auto RealClangRoot = llvm::sys::path::parent_path(ParentPath);
268 ROCmSearchDirs.emplace_back(ClangRoot.str(),
true);
269 if (RealClangRoot != ClangRoot)
270 ROCmSearchDirs.emplace_back(RealClangRoot.str(),
true);
274 ROCmSearchDirs.emplace_back(D.
SysRoot +
"/opt/rocm",
279 std::string LatestROCm;
280 llvm::VersionTuple LatestVer;
282 auto GetROCmVersion = [](StringRef DirName) {
283 llvm::VersionTuple
V;
284 std::string VerStr = DirName.drop_front(strlen(
"rocm-")).str();
287 std::replace(VerStr.begin(), VerStr.end(),
'-',
'.');
291 for (llvm::vfs::directory_iterator
294 File != FileEnd && !EC;
File.increment(EC)) {
295 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
298 if (LatestROCm.empty()) {
300 LatestVer = GetROCmVersion(LatestROCm);
303 auto Ver = GetROCmVersion(
FileName);
304 if (LatestVer < Ver) {
309 if (!LatestROCm.empty())
310 ROCmSearchDirs.emplace_back(D.
SysRoot +
"/opt/" + LatestROCm,
313 ROCmSearchDirs.emplace_back(D.
SysRoot +
"/usr/local",
315 ROCmSearchDirs.emplace_back(D.
SysRoot +
"/usr",
318 DoPrintROCmSearchDirs();
319 return ROCmSearchDirs;
323 const Driver &
D,
const llvm::Triple &HostTriple,
324 const llvm::opt::ArgList &Args,
bool DetectHIPRuntime,
bool DetectDeviceLib)
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);
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");
343 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_prim_path_EQ);
344 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
345 D.getVFS().exists(HIPRocPrimPathArg +
"/rocprim");
347 if (
auto *A = Args.getLastArg(clang::driver::options::OPT_hip_version_EQ)) {
348 HIPVersionArg = A->getValue();
349 unsigned Major = ~0
U;
350 unsigned Minor = ~0
U;
352 HIPVersionArg.split(Parts,
'.');
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())
361 if (Major != ~0
U && Minor == ~0
U)
363 if (Major == ~0
U || Minor == ~0
U)
364 D.Diag(diag::err_drv_invalid_value)
365 << A->getAsString(Args) << HIPVersionArg;
367 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
369 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
371 VersionPatch = DefaultVersionPatch;
373 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
374 DetectedVersion = (Twine(DefaultVersionMajor) +
"." +
375 Twine(DefaultVersionMinor) +
"." + VersionPatch)
379 if (DetectHIPRuntime)
386 assert(LibDevicePath.empty());
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);
395 if (!LibDevicePath.empty()) {
399 if (!FS.exists(LibDevicePath))
402 scanLibDevicePath(LibDevicePath);
403 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
408 auto CheckDeviceLib = [&](StringRef
Path,
bool StrictChecking) {
409 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
410 if (CheckLibDevice && !FS.exists(
Path))
413 scanLibDevicePath(
Path);
415 if (!NoBuiltinLibs) {
417 if (!allGenericLibsValid())
422 if (LibDeviceMap.empty())
430 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
431 "amdgcn",
"bitcode");
432 HasDeviceLibrary = CheckDeviceLib(LibDevicePath,
true);
433 if (HasDeviceLibrary)
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)
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));
457 if (HIPSearchDirs.empty())
458 HIPSearchDirs.append(getInstallationPathCandidates());
461 for (
const auto &Candidate : HIPSearchDirs) {
462 InstallPath = Candidate.Path;
463 if (InstallPath.empty() || !FS.exists(InstallPath))
467 auto SPACKPath = findSPACKPackage(Candidate,
"hip");
468 InstallPath = SPACKPath.empty() ? InstallPath : SPACKPath;
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");
480 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
481 llvm::sys::path::append(ParentSharePath,
"share");
484 const Twine &
c =
"",
const Twine &d =
"") {
486 llvm::sys::path::append(newpath, a,
b,
c, d);
490 std::vector<SmallString<0>> VersionFilePaths = {
491 Append(SharePath,
"hip",
"version"),
492 InstallPath != D.
SysRoot +
"/usr/local"
493 ?
Append(ParentSharePath,
"hip",
"version")
495 Append(BinPath,
".hipVersion")};
497 for (
const auto &VersionFilePath : VersionFilePaths) {
498 if (VersionFilePath.empty())
500 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
501 FS.getBufferForFile(VersionFilePath);
504 if (HIPVersionArg.empty() && VersionFile)
505 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
508 HasHIPRuntime =
true;
513 if (!Candidate.StrictChecking) {
514 HasHIPRuntime =
true;
518 HasHIPRuntime =
false;
523 OS <<
"Found HIP installation: " << InstallPath <<
", version "
524 << DetectedVersion <<
'\n';
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);
533 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
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));
554 const auto HandleHipStdPar = [=, &DriverArgs, &CC1Args]() {
559 if (!HIPStdParPathArg.empty() ||
560 !FS.exists(Inc +
"/thrust/system/hip/hipstdpar/hipstdpar_lib.hpp")) {
561 D.
Diag(diag::err_drv_no_hipstdpar_lib);
564 if (!HasRocThrustLibrary && !FS.exists(Inc +
"/thrust")) {
565 D.
Diag(diag::err_drv_no_hipstdpar_thrust_lib);
568 if (!HasRocPrimLibrary && !FS.exists(Inc +
"/rocprim")) {
569 D.
Diag(diag::err_drv_no_hipstdpar_prim_lib);
572 const char *ThrustPath;
573 if (HasRocThrustLibrary)
574 ThrustPath = DriverArgs.MakeArgString(HIPRocThrustPathArg);
576 ThrustPath = DriverArgs.MakeArgString(Inc +
"/thrust");
578 const char *HIPStdParPath;
580 HIPStdParPath = DriverArgs.MakeArgString(HIPStdParPathArg);
582 HIPStdParPath = DriverArgs.MakeArgString(StringRef(ThrustPath) +
583 "/system/hip/hipstdpar");
585 const char *PrimPath;
586 if (HasRocPrimLibrary)
587 PrimPath = DriverArgs.MakeArgString(HIPRocPrimPathArg);
589 PrimPath = DriverArgs.MakeArgString(
getIncludePath() +
"/rocprim");
591 CC1Args.append({
"-idirafter", ThrustPath,
"-idirafter", PrimPath,
592 "-idirafter", HIPStdParPath,
"-include",
593 "hipstdpar_lib.hpp"});
596 if (DriverArgs.hasArg(options::OPT_nogpuinc)) {
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 {
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");
629 Args.AddAllArgs(CmdArgs, options::OPT_L);
630 getToolChain().AddFilePathLibArgs(Args, CmdArgs);
632 if (
C.getDriver().isUsingLTO()) {
633 addLTOOptions(getToolChain(), Args, CmdArgs, Output, Inputs[0],
635 }
else if (Args.hasArg(options::OPT_mcpu_EQ)) {
636 CmdArgs.push_back(Args.MakeArgString(
637 "-plugin-opt=mcpu=" +
639 Args.getLastArgValue(options::OPT_mcpu_EQ))));
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");
870 CC1Args.push_back(
"-Werror=atomic-alignment");
876 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
881 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
882 if (TargetID.empty())
883 return {std::nullopt, std::nullopt, std::nullopt};
885 llvm::StringMap<bool> FeatureMap;
887 if (!OptionalGpuArch)
888 return {TargetID.str(), std::nullopt, std::nullopt};
890 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
894 const llvm::opt::ArgList &DriverArgs)
const {
896 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
898 << *PTID.OptionalTargetID;
906 if (Arg *A = Args.getLastArg(options::OPT_amdgpu_arch_tool_EQ))
907 Program = A->getValue();
913 return StdoutOrErr.takeError();
916 for (StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
918 GPUArchs.push_back(Arch.str());
920 if (GPUArchs.empty())
921 return llvm::createStringError(std::error_code(),
922 "No AMD GPU detected in the system");
924 return std::move(GPUArchs);
928 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
931 DeviceOffloadingKind);
936 DriverArgs.hasArg(options::OPT_nostdlib))
939 if (DriverArgs.hasArg(options::OPT_nogpulib))
943 const StringRef GpuArch =
getGPUArch(DriverArgs);
944 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
945 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
953 bool Wave64 =
isWave64(DriverArgs, Kind);
957 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
959 bool FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only);
962 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations);
963 bool FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math);
965 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt);
973 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
974 FastRelaxedMath, CorrectSqrt, ABIVer,
false));
977 CC1Args.push_back(
"-mlink-bitcode-file");
981 for (StringRef BCFile : BCLibs) {
982 CC1Args.push_back(
"-mlink-builtin-bitcode");
983 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
988 StringRef GPUArch, StringRef LibDeviceFile,
990 if (!hasDeviceLibrary()) {
991 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
994 if (LibDeviceFile.empty()) {
995 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
999 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString();
1007 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
bool Wave64,
1008 bool DAZ,
bool FiniteOnly,
bool UnsafeMathOpt,
bool FastRelaxedMath,
1012 auto AddBCLib = [&](StringRef BCFile) { BCLibs.push_back(BCFile.str()); };
1014 AddBCLib(getOCMLPath());
1016 AddBCLib(getOCKLPath());
1017 AddBCLib(getDenormalsAreZeroPath(DAZ));
1018 AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath));
1019 AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath));
1020 AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt));
1021 AddBCLib(getWavefrontSize64Path(Wave64));
1022 AddBCLib(LibDeviceFile);
1023 auto ABIVerPath = getABIVersionPath(ABIVer);
1024 if (!ABIVerPath.empty())
1025 AddBCLib(ABIVerPath);
1032 const std::string &GPUArch,
1033 bool isOpenMP)
const {
1034 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1035 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1047 bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
1048 options::OPT_fno_gpu_flush_denormals_to_zero,
1050 bool FiniteOnly = DriverArgs.hasFlag(
1051 options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only,
false);
1052 bool UnsafeMathOpt =
1053 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
1054 options::OPT_fno_unsafe_math_optimizations,
false);
1055 bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,
1056 options::OPT_fno_fast_math,
false);
1057 bool CorrectSqrt = DriverArgs.hasFlag(
1058 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
1059 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt,
true);
1060 bool Wave64 =
isWave64(DriverArgs, Kind);
1063 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
1064 FastRelaxedMath, CorrectSqrt, ABIVer, isOpenMP);
1068 const ToolChain &TC,
const llvm::opt::ArgList &DriverArgs,
1069 StringRef TargetID,
const llvm::opt::Arg *A)
const {
1071 if (TargetID.empty())
1073 Option O = A->getOption();
1074 if (!O.matches(options::OPT_fsanitize_EQ))
1077 if (!DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
1078 options::OPT_fno_gpu_sanitize,
true))
1085 if (K != SanitizerKind::Address)
1088 llvm::StringMap<bool> FeatureMap;
1091 assert(OptionalGpuArch &&
"Invalid Target ID");
1092 (void)OptionalGpuArch;
1093 auto Loc = FeatureMap.find(
"xnack");
1094 if (
Loc == FeatureMap.end() || !
Loc->second) {
1096 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1097 << 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...
std::string SysRoot
sysroot, if present
DiagnosticsEngine & getDiags() const
const char * getClangProgramPath() const
Get the path to the main clang executable.
DiagnosticBuilder Diag(unsigned DiagID) const
const llvm::opt::OptTable & getOpts() const
std::string ResourceDir
The path to the compiler resource directory.
llvm::vfs::FileSystem & getVFS() const
std::string Dir
The path the driver executable was in, as invoked from the command line.
StringRef getIncludePath() const
Get the detected path to Rocm's bin directory.
RocmInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args, bool DetectHIPRuntime=true, bool DetectDeviceLib=false)
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.
bool hasHIPRuntime() const
Check whether we detected a valid HIP runtime.
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.
void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const
void detectDeviceLibrary()
void print(raw_ostream &OS) const
Print information about the detected ROCm installation.
The JSON file list parser is used to communicate input to InstallAPI.
std::optional< llvm::StringRef > parseTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch, llvm::StringMap< bool > *FeatureMap)
Parse a target ID to get processor and feature map.
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.
SanitizerMask parseSanitizerValue(StringRef Value, bool AllowGroups)
Parse a single value from a -fsanitize= or -fno-sanitize= value list.
ABI version of device library.
static DeviceLibABIVersion fromCodeObjectVersion(unsigned CodeObjectVersion)
bool requiresLibrary()
Whether ABI version bc file is requested.
static constexpr ResponseFileSupport AtFileCurCP()