12#include "clang/Config/config.h"
17#include "llvm/ADT/StringExtras.h"
18#include "llvm/Option/ArgList.h"
19#include "llvm/Support/Error.h"
20#include "llvm/Support/LineIterator.h"
21#include "llvm/Support/Path.h"
22#include "llvm/Support/Process.h"
23#include "llvm/Support/VirtualFileSystem.h"
24#include "llvm/TargetParser/Host.h"
26#include <system_error>
41RocmInstallationDetector::findSPACKPackage(
const Candidate &Cand,
42 StringRef PackageName) {
46 std::string Prefix = Twine(PackageName +
"-" + Cand.SPACKReleaseStr).str();
48 for (llvm::vfs::directory_iterator
File = D.
getVFS().dir_begin(Cand.Path, EC),
50 File != FileEnd && !EC;
File.increment(EC)) {
51 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
54 if (SubDirs.size() > 1)
58 if (SubDirs.size() == 1) {
59 auto PackagePath = Cand.Path;
60 llvm::sys::path::append(PackagePath, SubDirs[0]);
63 if (SubDirs.size() == 0 && Verbose) {
64 llvm::errs() <<
"SPACK package " << Prefix <<
" not found at " << Cand.Path
69 if (SubDirs.size() > 1 && Verbose) {
70 llvm::errs() <<
"Cannot use SPACK package " << Prefix <<
" at " << Cand.Path
71 <<
" due to multiple installations for the same version\n";
76void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) {
77 assert(!Path.empty());
79 const StringRef Suffix(
".bc");
80 const StringRef Suffix2(
".amdgcn.bc");
83 for (llvm::vfs::directory_iterator LI = D.
getVFS().dir_begin(Path, EC), LE;
84 !EC && LI != LE; LI = LI.increment(EC)) {
85 StringRef FilePath = LI->path();
86 StringRef
FileName = llvm::sys::path::filename(FilePath);
92 BaseName =
FileName.drop_back(Suffix2.size());
94 BaseName =
FileName.drop_back(Suffix.size());
96 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
97 if (BaseName ==
"ocml") {
99 }
else if (BaseName ==
"ockl") {
101 }
else if (BaseName ==
"opencl") {
103 }
else if (BaseName ==
"hip") {
105 }
else if (BaseName ==
"asanrtl") {
107 }
else if (BaseName ==
"oclc_finite_only_off") {
108 FiniteOnly.Off = FilePath;
109 }
else if (BaseName ==
"oclc_finite_only_on") {
110 FiniteOnly.On = FilePath;
111 }
else if (BaseName ==
"oclc_daz_opt_on") {
112 DenormalsAreZero.On = FilePath;
113 }
else if (BaseName ==
"oclc_daz_opt_off") {
114 DenormalsAreZero.Off = FilePath;
115 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_on") {
116 CorrectlyRoundedSqrt.On = FilePath;
117 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_off") {
118 CorrectlyRoundedSqrt.Off = FilePath;
119 }
else if (BaseName ==
"oclc_unsafe_math_on") {
120 UnsafeMath.On = FilePath;
121 }
else if (BaseName ==
"oclc_unsafe_math_off") {
122 UnsafeMath.Off = FilePath;
123 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
124 WavefrontSize64.On = FilePath;
125 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
126 WavefrontSize64.Off = FilePath;
127 }
else if (BaseName.starts_with(ABIVersionPrefix)) {
128 unsigned ABIVersionNumber;
129 if (BaseName.drop_front(ABIVersionPrefix.size())
130 .getAsInteger(0, ABIVersionNumber))
132 ABIVersionMap[ABIVersionNumber] = FilePath.str();
136 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
137 if (!BaseName.starts_with(DeviceLibPrefix))
140 StringRef IsaVersionNumber =
141 BaseName.drop_front(DeviceLibPrefix.size());
143 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
146 std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
153bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
155 V.split(VersionParts,
'\n');
156 unsigned Major = ~0
U;
157 unsigned Minor = ~0
U;
158 for (
auto Part : VersionParts) {
159 auto Splits = Part.rtrim().split(
'=');
160 if (Splits.first ==
"HIP_VERSION_MAJOR") {
161 if (Splits.second.getAsInteger(0, Major))
163 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
164 if (Splits.second.getAsInteger(0, Minor))
166 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
167 VersionPatch = Splits.second.str();
169 if (Major == ~0
U || Minor == ~0
U)
171 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
173 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
180RocmInstallationDetector::getInstallationPathCandidates() {
183 if (!ROCmSearchDirs.empty())
184 return ROCmSearchDirs;
186 auto DoPrintROCmSearchDirs = [&]() {
187 if (PrintROCmSearchDirs)
188 for (
auto Cand : ROCmSearchDirs) {
189 llvm::errs() <<
"ROCm installation search path";
191 llvm::errs() <<
" (Spack " << Cand.SPACKReleaseStr <<
")";
192 llvm::errs() <<
": " << Cand.Path <<
'\n';
198 if (!RocmPathArg.empty()) {
199 ROCmSearchDirs.emplace_back(RocmPathArg.str());
200 DoPrintROCmSearchDirs();
201 return ROCmSearchDirs;
202 }
else if (std::optional<std::string> RocmPathEnv =
203 llvm::sys::Process::GetEnv(
"ROCM_PATH")) {
204 if (!RocmPathEnv->empty()) {
205 ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));
206 DoPrintROCmSearchDirs();
207 return ROCmSearchDirs;
212 StringRef InstallDir = D.
Dir;
217 auto DeduceROCmPath = [](StringRef ClangPath) {
219 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
220 StringRef ParentName = llvm::sys::path::filename(ParentDir);
223 if (ParentName ==
"bin") {
224 ParentDir = llvm::sys::path::parent_path(ParentDir);
225 ParentName = llvm::sys::path::filename(ParentDir);
233 if (ParentName.starts_with(
"llvm-amdgpu-")) {
235 ParentName.drop_front(strlen(
"llvm-amdgpu-")).split(
'-');
236 auto SPACKReleaseStr = SPACKPostfix.first;
237 if (!SPACKReleaseStr.empty()) {
238 ParentDir = llvm::sys::path::parent_path(ParentDir);
239 return Candidate(ParentDir.str(),
true,
246 if (ParentName ==
"llvm" || ParentName.starts_with(
"aomp"))
247 ParentDir = llvm::sys::path::parent_path(ParentDir);
249 return Candidate(ParentDir.str(),
true);
254 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
260 auto ParentPath = llvm::sys::path::parent_path(RealClangPath);
261 if (ParentPath != InstallDir)
262 ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));
265 auto ClangRoot = llvm::sys::path::parent_path(InstallDir);
266 auto RealClangRoot = llvm::sys::path::parent_path(ParentPath);
267 ROCmSearchDirs.emplace_back(ClangRoot.str(),
true);
268 if (RealClangRoot != ClangRoot)
269 ROCmSearchDirs.emplace_back(RealClangRoot.str(),
true);
273 ROCmSearchDirs.emplace_back(D.
SysRoot +
"/opt/rocm",
278 std::string LatestROCm;
279 llvm::VersionTuple LatestVer;
281 auto GetROCmVersion = [](StringRef DirName) {
282 llvm::VersionTuple
V;
283 std::string VerStr = DirName.drop_front(strlen(
"rocm-")).str();
286 std::replace(VerStr.begin(), VerStr.end(),
'-',
'.');
290 for (llvm::vfs::directory_iterator
293 File != FileEnd && !EC;
File.increment(EC)) {
294 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
297 if (LatestROCm.empty()) {
299 LatestVer = GetROCmVersion(LatestROCm);
302 auto Ver = GetROCmVersion(
FileName);
303 if (LatestVer < Ver) {
308 if (!LatestROCm.empty())
309 ROCmSearchDirs.emplace_back(D.
SysRoot +
"/opt/" + LatestROCm,
312 ROCmSearchDirs.emplace_back(D.
SysRoot +
"/usr/local",
314 ROCmSearchDirs.emplace_back(D.
SysRoot +
"/usr",
317 DoPrintROCmSearchDirs();
318 return ROCmSearchDirs;
322 const Driver &D,
const llvm::Triple &HostTriple,
323 const llvm::opt::ArgList &Args,
bool DetectHIPRuntime,
bool DetectDeviceLib)
325 Verbose = Args.hasArg(options::OPT_v);
326 RocmPathArg = Args.getLastArgValue(clang::driver::options::OPT_rocm_path_EQ);
327 PrintROCmSearchDirs =
328 Args.hasArg(clang::driver::options::OPT_print_rocm_search_dirs);
329 RocmDeviceLibPathArg =
330 Args.getAllArgValues(clang::driver::options::OPT_rocm_device_lib_path_EQ);
331 HIPPathArg = Args.getLastArgValue(clang::driver::options::OPT_hip_path_EQ);
333 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_path_EQ);
334 HasHIPStdParLibrary =
335 !HIPStdParPathArg.empty() && D.
getVFS().exists(HIPStdParPathArg +
336 "/hipstdpar_lib.hpp");
337 HIPRocThrustPathArg =
338 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_thrust_path_EQ);
339 HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&
340 D.
getVFS().exists(HIPRocThrustPathArg +
"/thrust");
342 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_prim_path_EQ);
343 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
344 D.
getVFS().exists(HIPRocPrimPathArg +
"/rocprim");
346 if (
auto *A = Args.getLastArg(clang::driver::options::OPT_hip_version_EQ)) {
347 HIPVersionArg = A->getValue();
348 unsigned Major = ~0
U;
349 unsigned Minor = ~0
U;
351 HIPVersionArg.split(Parts,
'.');
353 Parts[0].getAsInteger(0, Major);
354 if (Parts.size() > 1)
355 Parts[1].getAsInteger(0, Minor);
356 if (Parts.size() > 2)
357 VersionPatch = Parts[2].str();
358 if (VersionPatch.empty())
360 if (Major != ~0
U && Minor == ~0
U)
362 if (Major == ~0
U || Minor == ~0
U)
363 D.
Diag(diag::err_drv_invalid_value)
364 << A->getAsString(Args) << HIPVersionArg;
366 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
368 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
370 VersionPatch = DefaultVersionPatch;
372 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
373 DetectedVersion = (Twine(DefaultVersionMajor) +
"." +
374 Twine(DefaultVersionMinor) +
"." + VersionPatch)
378 if (DetectHIPRuntime)
385 assert(LibDevicePath.empty());
387 if (!RocmDeviceLibPathArg.empty())
388 LibDevicePath = RocmDeviceLibPathArg[RocmDeviceLibPathArg.size() - 1];
389 else if (std::optional<std::string> LibPathEnv =
390 llvm::sys::Process::GetEnv(
"HIP_DEVICE_LIB_PATH"))
391 LibDevicePath = std::move(*LibPathEnv);
394 if (!LibDevicePath.empty()) {
398 if (!FS.exists(LibDevicePath))
401 scanLibDevicePath(LibDevicePath);
402 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
407 auto CheckDeviceLib = [&](StringRef Path,
bool StrictChecking) {
408 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
409 if (CheckLibDevice && !FS.exists(Path))
412 scanLibDevicePath(Path);
414 if (!NoBuiltinLibs) {
416 if (!allGenericLibsValid())
421 if (LibDeviceMap.empty())
429 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
430 "amdgcn",
"bitcode");
431 HasDeviceLibrary = CheckDeviceLib(LibDevicePath,
true);
432 if (HasDeviceLibrary)
437 auto &ROCmDirs = getInstallationPathCandidates();
438 for (
const auto &Candidate : ROCmDirs) {
439 LibDevicePath = Candidate.Path;
440 llvm::sys::path::append(LibDevicePath,
"amdgcn",
"bitcode");
441 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);
442 if (HasDeviceLibrary)
449 if (!HIPPathArg.empty())
450 HIPSearchDirs.emplace_back(HIPPathArg.str());
451 else if (std::optional<std::string> HIPPathEnv =
452 llvm::sys::Process::GetEnv(
"HIP_PATH")) {
453 if (!HIPPathEnv->empty())
454 HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));
456 if (HIPSearchDirs.empty())
457 HIPSearchDirs.append(getInstallationPathCandidates());
460 for (
const auto &Candidate : HIPSearchDirs) {
461 InstallPath = Candidate.Path;
462 if (InstallPath.empty() || !FS.exists(InstallPath))
466 auto SPACKPath = findSPACKPackage(Candidate,
"hip");
467 InstallPath = SPACKPath.empty() ? InstallPath : SPACKPath;
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]() {
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.hasArg(options::OPT_nogpuinc)) {
603 D.
Diag(diag::err_drv_no_hip_runtime);
607 CC1Args.push_back(
"-idirafter");
609 if (UsesRuntimeWrapper)
610 CC1Args.append({
"-include",
"__clang_hip_runtime_wrapper.h"});
619 const char *LinkingOutput)
const {
621 std::string
Linker = getToolChain().GetProgramPath(getShortName());
622 ArgStringList CmdArgs;
623 CmdArgs.push_back(
"--no-undefined");
624 CmdArgs.push_back(
"-shared");
627 Args.AddAllArgs(CmdArgs, options::OPT_L);
628 getToolChain().AddFilePathLibArgs(Args, CmdArgs);
630 if (
C.getDriver().isUsingLTO())
631 addLTOOptions(getToolChain(), Args, CmdArgs, Output, Inputs[0],
633 else if (Args.hasArg(options::OPT_mcpu_EQ))
634 CmdArgs.push_back(Args.MakeArgString(
635 "-plugin-opt=mcpu=" + Args.getLastArgValue(options::OPT_mcpu_EQ)));
636 CmdArgs.push_back(
"-o");
638 C.addCommand(std::make_unique<Command>(
640 CmdArgs, Inputs, Output));
644 const llvm::Triple &Triple,
645 const llvm::opt::ArgList &Args,
646 std::vector<StringRef> &Features) {
649 StringRef TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
650 if (!TargetID.empty()) {
651 llvm::StringMap<bool> FeatureMap;
652 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
653 if (OptionalGpuArch) {
654 StringRef GpuArch = *OptionalGpuArch;
660 auto Pos = FeatureMap.find(Feature);
661 if (Pos == FeatureMap.end())
663 Features.push_back(Args.MakeArgStringRef(
664 (Twine(Pos->second ?
"+" :
"-") + Feature).str()));
669 if (Args.hasFlag(options::OPT_mwavefrontsize64,
670 options::OPT_mno_wavefrontsize64,
false))
671 Features.push_back(
"+wavefrontsize64");
673 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
674 options::OPT_mno_amdgpu_precise_memory_op,
false))
675 Features.push_back(
"+precise-memory");
678 options::OPT_m_amdgpu_Features_Group);
686 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
702 DerivedArgList *DAL =
708 DAL =
new DerivedArgList(Args.getBaseArgs());
714 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
715 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
716 DAL->eraseArg(options::OPT_mcpu_EQ);
720 << llvm::Triple::getArchTypeName(
getArch())
721 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
723 auto &GPUs = *GPUsOrErr;
724 if (GPUs.size() > 1) {
726 << llvm::Triple::getArchTypeName(
getArch())
727 << llvm::join(GPUs,
", ") <<
"-mcpu";
729 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
730 Args.MakeArgString(GPUs.front()));
736 if (!Args.getLastArgValue(options::OPT_x).equals(
"cl"))
740 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
741 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
743 : options::OPT_m32));
747 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
749 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
757 llvm::AMDGPU::GPUKind Kind) {
760 if (Kind == llvm::AMDGPU::GK_NONE)
763 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
767 const bool BothDenormAndFMAFast =
768 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
769 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
770 return !BothDenormAndFMAFast;
774 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
775 const llvm::fltSemantics *FPType)
const {
777 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
778 return llvm::DenormalMode::getIEEE();
783 auto Kind = llvm::AMDGPU::parseArchAMDGCN(Arch);
784 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
785 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
786 options::OPT_fno_gpu_flush_denormals_to_zero,
788 return llvm::DenormalMode::getPreserveSign();
790 return llvm::DenormalMode::getIEEE();
793 const StringRef GpuArch =
getGPUArch(DriverArgs);
794 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
798 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
803 return DAZ ? llvm::DenormalMode::getPreserveSign() :
804 llvm::DenormalMode::getIEEE();
808 llvm::AMDGPU::GPUKind Kind) {
809 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
810 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
812 return !HasWave32 || DriverArgs.hasFlag(
813 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
825 const llvm::opt::ArgList &DriverArgs,
826 llvm::opt::ArgStringList &CC1Args,
830 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
831 options::OPT_fvisibility_ms_compat)) {
832 CC1Args.push_back(
"-fvisibility=hidden");
833 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
840 CC1Args.push_back(
"-Werror=atomic-alignment");
846 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
851 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
852 if (TargetID.empty())
853 return {std::nullopt, std::nullopt, std::nullopt};
855 llvm::StringMap<bool> FeatureMap;
857 if (!OptionalGpuArch)
858 return {TargetID.str(), std::nullopt, std::nullopt};
860 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
864 const llvm::opt::ArgList &DriverArgs)
const {
866 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
868 << *PTID.OptionalTargetID;
876 if (Arg *A = Args.getLastArg(options::OPT_amdgpu_arch_tool_EQ))
877 Program = A->getValue();
883 return StdoutOrErr.takeError();
886 for (StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
888 GPUArchs.push_back(Arch.str());
890 if (GPUArchs.empty())
891 return llvm::createStringError(std::error_code(),
892 "No AMD GPU detected in the system");
894 return std::move(GPUArchs);
898 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
901 DeviceOffloadingKind);
906 DriverArgs.hasArg(options::OPT_nostdlib))
909 if (DriverArgs.hasArg(options::OPT_nogpulib))
913 const StringRef GpuArch =
getGPUArch(DriverArgs);
914 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
915 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
923 bool Wave64 =
isWave64(DriverArgs, Kind);
927 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
929 bool FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only);
932 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations);
933 bool FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math);
935 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt);
943 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
944 FastRelaxedMath, CorrectSqrt, ABIVer,
false));
946 for (StringRef BCFile : BCLibs) {
947 CC1Args.push_back(
"-mlink-builtin-bitcode");
948 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
953 StringRef GPUArch, StringRef LibDeviceFile,
955 if (!hasDeviceLibrary()) {
956 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
959 if (LibDeviceFile.empty()) {
960 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
964 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString();
972 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
bool Wave64,
973 bool DAZ,
bool FiniteOnly,
bool UnsafeMathOpt,
bool FastRelaxedMath,
977 auto AddBCLib = [&](StringRef BCFile) { BCLibs.push_back(BCFile.str()); };
979 AddBCLib(getOCMLPath());
981 AddBCLib(getOCKLPath());
982 AddBCLib(getDenormalsAreZeroPath(DAZ));
983 AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath));
984 AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath));
985 AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt));
986 AddBCLib(getWavefrontSize64Path(Wave64));
987 AddBCLib(LibDeviceFile);
988 auto ABIVerPath = getABIVersionPath(ABIVer);
989 if (!ABIVerPath.empty())
990 AddBCLib(ABIVerPath);
997 const std::string &GPUArch,
998 bool isOpenMP)
const {
999 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1000 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1012 bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
1013 options::OPT_fno_gpu_flush_denormals_to_zero,
1015 bool FiniteOnly = DriverArgs.hasFlag(
1016 options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only,
false);
1017 bool UnsafeMathOpt =
1018 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
1019 options::OPT_fno_unsafe_math_optimizations,
false);
1020 bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,
1021 options::OPT_fno_fast_math,
false);
1022 bool CorrectSqrt = DriverArgs.hasFlag(
1023 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
1024 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt,
true);
1025 bool Wave64 =
isWave64(DriverArgs, Kind);
1028 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
1029 FastRelaxedMath, CorrectSqrt, ABIVer, isOpenMP);
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
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.
ABI version of device library.
static DeviceLibABIVersion fromCodeObjectVersion(unsigned CodeObjectVersion)
bool requiresLibrary()
Whether ABI version bc file is requested.
static constexpr ResponseFileSupport AtFileCurCP()