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 {
620 std::string
Linker = getToolChain().GetLinkerPath();
621 ArgStringList CmdArgs;
622 CmdArgs.push_back(
"--no-undefined");
623 CmdArgs.push_back(
"-shared");
626 Args.AddAllArgs(CmdArgs, options::OPT_L);
627 getToolChain().AddFilePathLibArgs(Args, CmdArgs);
629 if (
C.getDriver().isUsingLTO())
630 addLTOOptions(getToolChain(), Args, CmdArgs, Output, Inputs[0],
632 else if (Args.hasArg(options::OPT_mcpu_EQ))
633 CmdArgs.push_back(Args.MakeArgString(
634 "-plugin-opt=mcpu=" + Args.getLastArgValue(options::OPT_mcpu_EQ)));
635 CmdArgs.push_back(
"-o");
637 C.addCommand(std::make_unique<Command>(
639 CmdArgs, Inputs, Output));
643 const llvm::Triple &Triple,
644 const llvm::opt::ArgList &Args,
645 std::vector<StringRef> &Features) {
648 StringRef TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
649 if (!TargetID.empty()) {
650 llvm::StringMap<bool> FeatureMap;
651 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
652 if (OptionalGpuArch) {
653 StringRef GpuArch = *OptionalGpuArch;
659 auto Pos = FeatureMap.find(Feature);
660 if (Pos == FeatureMap.end())
662 Features.push_back(Args.MakeArgStringRef(
663 (Twine(Pos->second ?
"+" :
"-") + Feature).str()));
668 if (Args.hasFlag(options::OPT_mwavefrontsize64,
669 options::OPT_mno_wavefrontsize64,
false))
670 Features.push_back(
"+wavefrontsize64");
672 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
673 options::OPT_mno_amdgpu_precise_memory_op,
false))
674 Features.push_back(
"+precise-memory");
677 options::OPT_m_amdgpu_Features_Group);
685 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
701 DerivedArgList *DAL =
707 DAL =
new DerivedArgList(Args.getBaseArgs());
713 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
714 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
715 DAL->eraseArg(options::OPT_mcpu_EQ);
719 << llvm::Triple::getArchTypeName(
getArch())
720 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
722 auto &GPUs = *GPUsOrErr;
723 if (GPUs.size() > 1) {
725 << llvm::Triple::getArchTypeName(
getArch())
726 << llvm::join(GPUs,
", ") <<
"-mcpu";
728 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
729 Args.MakeArgString(GPUs.front()));
735 if (!Args.getLastArgValue(options::OPT_x).equals(
"cl"))
739 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
740 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
742 : options::OPT_m32));
746 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
748 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
756 llvm::AMDGPU::GPUKind Kind) {
759 if (Kind == llvm::AMDGPU::GK_NONE)
762 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
766 const bool BothDenormAndFMAFast =
767 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
768 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
769 return !BothDenormAndFMAFast;
773 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
774 const llvm::fltSemantics *FPType)
const {
776 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
777 return llvm::DenormalMode::getIEEE();
782 auto Kind = llvm::AMDGPU::parseArchAMDGCN(Arch);
783 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
784 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
785 options::OPT_fno_gpu_flush_denormals_to_zero,
787 return llvm::DenormalMode::getPreserveSign();
789 return llvm::DenormalMode::getIEEE();
792 const StringRef GpuArch =
getGPUArch(DriverArgs);
793 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
797 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
802 return DAZ ? llvm::DenormalMode::getPreserveSign() :
803 llvm::DenormalMode::getIEEE();
807 llvm::AMDGPU::GPUKind Kind) {
808 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
809 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
811 return !HasWave32 || DriverArgs.hasFlag(
812 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
824 const llvm::opt::ArgList &DriverArgs,
825 llvm::opt::ArgStringList &CC1Args,
829 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
830 options::OPT_fvisibility_ms_compat)) {
831 CC1Args.push_back(
"-fvisibility=hidden");
832 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
839 CC1Args.push_back(
"-Werror=atomic-alignment");
845 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
850 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
851 if (TargetID.empty())
852 return {std::nullopt, std::nullopt, std::nullopt};
854 llvm::StringMap<bool> FeatureMap;
856 if (!OptionalGpuArch)
857 return {TargetID.str(), std::nullopt, std::nullopt};
859 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
863 const llvm::opt::ArgList &DriverArgs)
const {
865 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
867 << *PTID.OptionalTargetID;
875 if (Arg *A = Args.getLastArg(options::OPT_amdgpu_arch_tool_EQ))
876 Program = A->getValue();
882 return StdoutOrErr.takeError();
885 for (StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
887 GPUArchs.push_back(Arch.str());
889 if (GPUArchs.empty())
890 return llvm::createStringError(std::error_code(),
891 "No AMD GPU detected in the system");
893 return std::move(GPUArchs);
897 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
900 DeviceOffloadingKind);
905 DriverArgs.hasArg(options::OPT_nostdlib))
908 if (DriverArgs.hasArg(options::OPT_nogpulib))
912 const StringRef GpuArch =
getGPUArch(DriverArgs);
913 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
914 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
922 bool Wave64 =
isWave64(DriverArgs, Kind);
926 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
928 bool FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only);
931 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations);
932 bool FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math);
934 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt);
942 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
943 FastRelaxedMath, CorrectSqrt, ABIVer,
false));
945 for (StringRef BCFile : BCLibs) {
946 CC1Args.push_back(
"-mlink-builtin-bitcode");
947 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
952 StringRef GPUArch, StringRef LibDeviceFile,
954 if (!hasDeviceLibrary()) {
955 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
958 if (LibDeviceFile.empty()) {
959 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
963 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString();
971 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
bool Wave64,
972 bool DAZ,
bool FiniteOnly,
bool UnsafeMathOpt,
bool FastRelaxedMath,
976 auto AddBCLib = [&](StringRef BCFile) { BCLibs.push_back(BCFile.str()); };
978 AddBCLib(getOCMLPath());
980 AddBCLib(getOCKLPath());
981 AddBCLib(getDenormalsAreZeroPath(DAZ));
982 AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath));
983 AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath));
984 AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt));
985 AddBCLib(getWavefrontSize64Path(Wave64));
986 AddBCLib(LibDeviceFile);
987 auto ABIVerPath = getABIVersionPath(ABIVer);
988 if (!ABIVerPath.empty())
989 AddBCLib(ABIVerPath);
996 const std::string &GPUArch,
997 bool isOpenMP)
const {
998 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
999 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1011 bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
1012 options::OPT_fno_gpu_flush_denormals_to_zero,
1014 bool FiniteOnly = DriverArgs.hasFlag(
1015 options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only,
false);
1016 bool UnsafeMathOpt =
1017 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
1018 options::OPT_fno_unsafe_math_optimizations,
false);
1019 bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,
1020 options::OPT_fno_fast_math,
false);
1021 bool CorrectSqrt = DriverArgs.hasFlag(
1022 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
1023 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt,
true);
1024 bool Wave64 =
isWave64(DriverArgs, Kind);
1027 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
1028 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()