11#include "clang/Config/config.h"
17#include "llvm/ADT/StringExtras.h"
18#include "llvm/Option/ArgList.h"
19#include "llvm/Support/Error.h"
20#include "llvm/Support/LineIterator.h"
21#include "llvm/Support/Path.h"
22#include "llvm/Support/Process.h"
23#include "llvm/Support/VirtualFileSystem.h"
24#include "llvm/TargetParser/Host.h"
25#include "llvm/TargetParser/TargetParser.h"
27#include <system_error>
35RocmInstallationDetector::CommonBitcodeLibsPreferences::
36 CommonBitcodeLibsPreferences(
const Driver &D,
37 const llvm::opt::ArgList &DriverArgs,
40 const bool NeedsASanRT)
43 const auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
44 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
48 const bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
50 !HasWave32 || DriverArgs.hasFlag(options::OPT_mwavefrontsize64,
51 options::OPT_mno_wavefrontsize64,
false);
58 const bool DefaultDAZ =
59 (Kind == llvm::AMDGPU::GK_NONE)
61 : !((ArchAttr &
llvm::
AMDGPU::FEATURE_FAST_FMA_F32) &&
62 (ArchAttr &
llvm::
AMDGPU::FEATURE_FAST_DENORMAL_F32));
65 DAZ = IsKnownOffloading
66 ? DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
67 options::OPT_fno_gpu_flush_denormals_to_zero,
69 : DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) || DefaultDAZ;
71 FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only) ||
72 DriverArgs.hasFlag(options::OPT_ffinite_math_only,
73 options::OPT_fno_finite_math_only,
false);
76 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations) ||
77 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
78 options::OPT_fno_unsafe_math_optimizations,
false);
80 FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math) ||
81 DriverArgs.hasFlag(options::OPT_ffast_math,
82 options::OPT_fno_fast_math,
false);
84 const bool DefaultSqrt = IsKnownOffloading ?
true :
false;
86 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt) ||
88 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
89 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, DefaultSqrt);
92 GPUSan = (DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
93 options::OPT_fno_gpu_sanitize,
true) &&
97void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) {
98 assert(!Path.empty());
100 const StringRef Suffix(
".bc");
101 const StringRef Suffix2(
".amdgcn.bc");
104 for (llvm::vfs::directory_iterator LI = D.getVFS().dir_begin(Path, EC), LE;
105 !EC && LI != LE; LI = LI.increment(EC)) {
106 StringRef FilePath = LI->path();
107 StringRef
FileName = llvm::sys::path::filename(FilePath);
113 BaseName =
FileName.drop_back(Suffix2.size());
114 else if (
FileName.ends_with(Suffix))
115 BaseName =
FileName.drop_back(Suffix.size());
117 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
118 if (BaseName ==
"ocml") {
120 }
else if (BaseName ==
"ockl") {
122 }
else if (BaseName ==
"opencl") {
124 }
else if (BaseName ==
"asanrtl") {
126 }
else if (BaseName ==
"oclc_finite_only_off") {
127 FiniteOnly.Off = FilePath;
128 }
else if (BaseName ==
"oclc_finite_only_on") {
129 FiniteOnly.On = FilePath;
130 }
else if (BaseName ==
"oclc_daz_opt_on") {
131 DenormalsAreZero.On = FilePath;
132 }
else if (BaseName ==
"oclc_daz_opt_off") {
133 DenormalsAreZero.Off = FilePath;
134 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_on") {
135 CorrectlyRoundedSqrt.On = FilePath;
136 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_off") {
137 CorrectlyRoundedSqrt.Off = FilePath;
138 }
else if (BaseName ==
"oclc_unsafe_math_on") {
139 UnsafeMath.On = FilePath;
140 }
else if (BaseName ==
"oclc_unsafe_math_off") {
141 UnsafeMath.Off = FilePath;
142 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
143 WavefrontSize64.On = FilePath;
144 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
145 WavefrontSize64.Off = FilePath;
146 }
else if (BaseName.starts_with(ABIVersionPrefix)) {
147 unsigned ABIVersionNumber;
148 if (BaseName.drop_front(ABIVersionPrefix.size())
149 .getAsInteger(0, ABIVersionNumber))
151 ABIVersionMap[ABIVersionNumber] = FilePath.str();
155 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
156 if (!BaseName.starts_with(DeviceLibPrefix))
159 StringRef IsaVersionNumber =
160 BaseName.drop_front(DeviceLibPrefix.size());
162 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
165 std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
172bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
173 SmallVector<StringRef, 4> VersionParts;
174 V.split(VersionParts,
'\n');
175 unsigned Major = ~0U;
176 unsigned Minor = ~0U;
177 for (
auto Part : VersionParts) {
178 auto Splits = Part.rtrim().split(
'=');
179 if (Splits.first ==
"HIP_VERSION_MAJOR") {
180 if (Splits.second.getAsInteger(0, Major))
182 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
183 if (Splits.second.getAsInteger(0, Minor))
185 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
186 VersionPatch = Splits.second.str();
188 if (Major == ~0U || Minor == ~0U)
190 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
192 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
198const SmallVectorImpl<RocmInstallationDetector::Candidate> &
199RocmInstallationDetector::getInstallationPathCandidates() {
202 if (!ROCmSearchDirs.empty())
203 return ROCmSearchDirs;
205 auto DoPrintROCmSearchDirs = [&]() {
206 if (PrintROCmSearchDirs)
207 for (
auto Cand : ROCmSearchDirs) {
208 llvm::errs() <<
"ROCm installation search path: " << Cand.Path <<
'\n';
214 if (!RocmPathArg.empty()) {
215 ROCmSearchDirs.emplace_back(RocmPathArg.str());
216 DoPrintROCmSearchDirs();
217 return ROCmSearchDirs;
218 }
else if (std::optional<std::string> RocmPathEnv =
219 llvm::sys::Process::GetEnv(
"ROCM_PATH")) {
220 if (!RocmPathEnv->empty()) {
221 ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));
222 DoPrintROCmSearchDirs();
223 return ROCmSearchDirs;
228 StringRef InstallDir = D.Dir;
233 auto DeduceROCmPath = [](StringRef ClangPath) {
235 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
236 StringRef ParentName = llvm::sys::path::filename(ParentDir);
239 if (ParentName ==
"bin") {
240 ParentDir = llvm::sys::path::parent_path(ParentDir);
241 ParentName = llvm::sys::path::filename(ParentDir);
246 if (ParentName ==
"llvm" || ParentName.starts_with(
"aomp"))
247 ParentDir = llvm::sys::path::parent_path(ParentDir);
249 return Candidate(ParentDir.str(),
true);
254 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
258 llvm::SmallString<256> RealClangPath;
259 llvm::sys::fs::real_path(D.getClangProgramPath(), RealClangPath);
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);
270 ROCmSearchDirs.emplace_back(D.ResourceDir,
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 llvm::replace(VerStr,
'-',
'.');
290 for (llvm::vfs::directory_iterator
291 File = D.getVFS().dir_begin(D.SysRoot +
"/opt", EC),
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(options::OPT_rocm_path_EQ);
327 PrintROCmSearchDirs = Args.hasArg(options::OPT_print_rocm_search_dirs);
328 RocmDeviceLibPathArg =
329 Args.getAllArgValues(options::OPT_rocm_device_lib_path_EQ);
330 HIPPathArg = Args.getLastArgValue(options::OPT_hip_path_EQ);
331 HIPStdParPathArg = Args.getLastArgValue(options::OPT_hipstdpar_path_EQ);
332 HasHIPStdParLibrary =
333 !HIPStdParPathArg.empty() && D.getVFS().exists(HIPStdParPathArg +
334 "/hipstdpar_lib.hpp");
335 HIPRocThrustPathArg =
336 Args.getLastArgValue(options::OPT_hipstdpar_thrust_path_EQ);
337 HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&
338 D.getVFS().exists(HIPRocThrustPathArg +
"/thrust");
339 HIPRocPrimPathArg = Args.getLastArgValue(options::OPT_hipstdpar_prim_path_EQ);
340 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
341 D.getVFS().exists(HIPRocPrimPathArg +
"/rocprim");
343 if (
auto *A = Args.getLastArg(options::OPT_hip_version_EQ)) {
344 HIPVersionArg = A->getValue();
345 unsigned Major = ~0U;
346 unsigned Minor = ~0U;
347 SmallVector<StringRef, 3> Parts;
348 HIPVersionArg.split(Parts,
'.');
350 Parts[0].getAsInteger(0, Major);
351 if (Parts.size() > 1)
352 Parts[1].getAsInteger(0, Minor);
353 if (Parts.size() > 2)
354 VersionPatch = Parts[2].str();
355 if (VersionPatch.empty())
357 if (Major != ~0U && Minor == ~0U)
359 if (Major == ~0U || Minor == ~0U)
360 D.Diag(diag::err_drv_invalid_value)
361 << A->getAsString(Args) << HIPVersionArg;
363 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
365 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
367 VersionPatch = DefaultVersionPatch;
369 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
370 DetectedVersion = (Twine(DefaultVersionMajor) +
"." +
371 Twine(DefaultVersionMinor) +
"." + VersionPatch)
375 if (DetectHIPRuntime)
382 assert(LibDevicePath.empty());
384 if (!RocmDeviceLibPathArg.empty())
385 LibDevicePath = RocmDeviceLibPathArg[RocmDeviceLibPathArg.size() - 1];
386 else if (std::optional<std::string> LibPathEnv =
387 llvm::sys::Process::GetEnv(
"HIP_DEVICE_LIB_PATH"))
388 LibDevicePath = std::move(*LibPathEnv);
390 auto &FS = D.getVFS();
391 if (!LibDevicePath.empty()) {
395 if (!FS.exists(LibDevicePath))
398 scanLibDevicePath(LibDevicePath);
399 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
404 auto CheckDeviceLib = [&](StringRef Path,
bool StrictChecking) {
405 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
406 if (CheckLibDevice && !FS.exists(Path))
409 scanLibDevicePath(Path);
411 if (!NoBuiltinLibs) {
413 if (!allGenericLibsValid())
418 if (LibDeviceMap.empty())
425 LibDevicePath = D.ResourceDir;
426 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
427 "amdgcn",
"bitcode");
428 HasDeviceLibrary = CheckDeviceLib(LibDevicePath,
true);
429 if (HasDeviceLibrary)
434 auto &ROCmDirs = getInstallationPathCandidates();
435 for (
const auto &Candidate : ROCmDirs) {
436 LibDevicePath = Candidate.Path;
437 llvm::sys::path::append(LibDevicePath,
"amdgcn",
"bitcode");
438 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);
439 if (HasDeviceLibrary)
446 if (!HIPPathArg.empty())
447 HIPSearchDirs.emplace_back(HIPPathArg.str());
448 else if (std::optional<std::string> HIPPathEnv =
449 llvm::sys::Process::GetEnv(
"HIP_PATH")) {
450 if (!HIPPathEnv->empty())
451 HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));
453 if (HIPSearchDirs.empty())
454 HIPSearchDirs.append(getInstallationPathCandidates());
455 auto &FS = D.getVFS();
457 for (
const auto &Candidate : HIPSearchDirs) {
458 InstallPath = Candidate.Path;
459 if (InstallPath.empty() || !FS.exists(InstallPath))
462 BinPath = InstallPath;
463 llvm::sys::path::append(BinPath,
"bin");
464 IncludePath = InstallPath;
465 llvm::sys::path::append(IncludePath,
"include");
466 LibPath = InstallPath;
467 llvm::sys::path::append(LibPath,
"lib");
468 SharePath = InstallPath;
469 llvm::sys::path::append(SharePath,
"share");
472 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
473 llvm::sys::path::append(ParentSharePath,
"share");
476 const Twine &
c =
"",
const Twine &d =
"") {
478 llvm::sys::path::append(newpath, a,
b,
c, d);
482 std::vector<SmallString<0>> VersionFilePaths = {
483 Append(SharePath,
"hip",
"version"),
484 InstallPath != D.SysRoot +
"/usr/local"
485 ?
Append(ParentSharePath,
"hip",
"version")
487 Append(BinPath,
".hipVersion")};
489 for (
const auto &VersionFilePath : VersionFilePaths) {
490 if (VersionFilePath.empty())
492 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
493 FS.getBufferForFile(VersionFilePath);
496 if (HIPVersionArg.empty() && VersionFile)
497 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
500 HasHIPRuntime =
true;
505 if (!Candidate.StrictChecking) {
506 HasHIPRuntime =
true;
510 HasHIPRuntime =
false;
515 OS <<
"Found HIP installation: " << InstallPath <<
", version "
516 << DetectedVersion <<
'\n';
520 ArgStringList &CC1Args)
const {
521 bool UsesRuntimeWrapper = VersionMajorMinor > llvm::VersionTuple(3, 5) &&
522 !DriverArgs.hasArg(options::OPT_nohipwrapperinc);
523 bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
525 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
540 if (UsesRuntimeWrapper)
541 llvm::sys::path::append(P,
"include",
"cuda_wrappers");
542 CC1Args.push_back(
"-internal-isystem");
543 CC1Args.push_back(DriverArgs.MakeArgString(P));
546 const auto HandleHipStdPar = [=, &DriverArgs, &CC1Args]() {
548 auto &FS = D.getVFS();
551 if (!HIPStdParPathArg.empty() ||
552 !FS.exists(Inc +
"/thrust/system/hip/hipstdpar/hipstdpar_lib.hpp")) {
553 D.Diag(diag::err_drv_no_hipstdpar_lib);
556 if (!HasRocThrustLibrary && !FS.exists(Inc +
"/thrust")) {
557 D.Diag(diag::err_drv_no_hipstdpar_thrust_lib);
560 if (!HasRocPrimLibrary && !FS.exists(Inc +
"/rocprim")) {
561 D.Diag(diag::err_drv_no_hipstdpar_prim_lib);
564 const char *ThrustPath;
565 if (HasRocThrustLibrary)
566 ThrustPath = DriverArgs.MakeArgString(HIPRocThrustPathArg);
568 ThrustPath = DriverArgs.MakeArgString(Inc +
"/thrust");
570 const char *HIPStdParPath;
572 HIPStdParPath = DriverArgs.MakeArgString(HIPStdParPathArg);
574 HIPStdParPath = DriverArgs.MakeArgString(StringRef(ThrustPath) +
575 "/system/hip/hipstdpar");
577 const char *PrimPath;
578 if (HasRocPrimLibrary)
579 PrimPath = DriverArgs.MakeArgString(HIPRocPrimPathArg);
581 PrimPath = DriverArgs.MakeArgString(
getIncludePath() +
"/rocprim");
583 CC1Args.append({
"-idirafter", ThrustPath,
"-idirafter", PrimPath,
584 "-idirafter", HIPStdParPath,
"-include",
585 "hipstdpar_lib.hpp"});
588 if (!DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
597 D.Diag(diag::err_drv_no_hip_runtime);
601 CC1Args.push_back(
"-idirafter");
603 if (UsesRuntimeWrapper)
604 CC1Args.append({
"-include",
"__clang_hip_runtime_wrapper.h"});
613 const char *LinkingOutput)
const {
615 ArgStringList CmdArgs;
616 if (!Args.hasArg(options::OPT_r)) {
617 CmdArgs.push_back(
"--no-undefined");
618 CmdArgs.push_back(
"-shared");
621 if (
C.getDriver().isUsingLTO()) {
622 const bool ThinLTO = (
C.getDriver().getLTOMode() ==
LTOK_Thin);
624 }
else if (Args.hasArg(options::OPT_mcpu_EQ)) {
625 CmdArgs.push_back(Args.MakeArgString(
626 "-plugin-opt=mcpu=" +
628 Args.getLastArgValue(options::OPT_mcpu_EQ))));
632 Args.AddAllArgs(CmdArgs, options::OPT_L);
636 std::vector<StringRef> Features;
639 if (!Features.empty()) {
641 Args.MakeArgString(
"-plugin-opt=-mattr=" + llvm::join(Features,
",")));
644 if (Args.hasArg(options::OPT_stdlib))
645 CmdArgs.append({
"-lc",
"-lm"});
646 if (Args.hasArg(options::OPT_startfiles)) {
647 std::optional<std::string> IncludePath =
getToolChain().getStdlibPath();
649 IncludePath =
"/lib";
651 llvm::sys::path::append(P,
"crt1.o");
652 CmdArgs.push_back(Args.MakeArgString(P));
655 CmdArgs.push_back(
"-o");
657 C.addCommand(std::make_unique<Command>(
659 CmdArgs, Inputs, Output));
663 const llvm::Triple &Triple,
664 const llvm::opt::ArgList &Args,
665 std::vector<StringRef> &Features) {
669 if (Args.hasArg(options::OPT_mcpu_EQ))
670 TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
671 else if (Args.hasArg(options::OPT_march_EQ))
672 TargetID = Args.getLastArgValue(options::OPT_march_EQ);
673 if (!TargetID.empty()) {
674 llvm::StringMap<bool> FeatureMap;
675 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
676 if (OptionalGpuArch) {
677 StringRef GpuArch = *OptionalGpuArch;
683 auto Pos = FeatureMap.find(
Feature);
684 if (Pos == FeatureMap.end())
686 Features.push_back(Args.MakeArgStringRef(
687 (Twine(Pos->second ?
"+" :
"-") +
Feature).str()));
692 if (Args.hasFlag(options::OPT_mwavefrontsize64,
693 options::OPT_mno_wavefrontsize64,
false))
694 Features.push_back(
"+wavefrontsize64");
696 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
697 options::OPT_mno_amdgpu_precise_memory_op,
false))
698 Features.push_back(
"+precise-memory");
701 options::OPT_m_amdgpu_Features_Group);
709 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
725 DerivedArgList *DAL =
731 DAL =
new DerivedArgList(Args.getBaseArgs());
737 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
738 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
739 DAL->eraseArg(options::OPT_mcpu_EQ);
743 << llvm::Triple::getArchTypeName(
getArch())
744 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
746 auto &GPUs = *GPUsOrErr;
747 if (GPUs.size() > 1) {
749 << llvm::Triple::getArchTypeName(
getArch())
750 << llvm::join(GPUs,
", ") <<
"-mcpu";
752 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
753 Args.MakeArgString(GPUs.front()));
759 if (Args.getLastArgValue(options::OPT_x) !=
"cl")
763 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
764 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
766 : options::OPT_m32));
770 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
772 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
780 llvm::AMDGPU::GPUKind Kind) {
783 if (Kind == llvm::AMDGPU::GK_NONE)
786 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
790 const bool BothDenormAndFMAFast =
791 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
792 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
793 return !BothDenormAndFMAFast;
797 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
798 const llvm::fltSemantics *FPType)
const {
800 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
801 return llvm::DenormalMode::getIEEE();
806 auto Kind = llvm::AMDGPU::parseArchAMDGCN(
Arch);
807 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
808 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
809 options::OPT_fno_gpu_flush_denormals_to_zero,
811 return llvm::DenormalMode::getPreserveSign();
813 return llvm::DenormalMode::getIEEE();
816 const StringRef GpuArch =
getGPUArch(DriverArgs);
817 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
821 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
826 return DAZ ? llvm::DenormalMode::getPreserveSign() :
827 llvm::DenormalMode::getIEEE();
831 llvm::AMDGPU::GPUKind Kind) {
832 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
833 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
835 return !HasWave32 || DriverArgs.hasFlag(
836 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
848 const llvm::opt::ArgList &DriverArgs,
849 llvm::opt::ArgStringList &CC1Args,
853 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
854 options::OPT_fvisibility_ms_compat)) {
855 CC1Args.push_back(
"-fvisibility=hidden");
856 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
870 !DriverArgs.hasArg(options::OPT_disable_llvm_optzns))
871 CC1Args.push_back(
"-disable-llvm-optzns");
880 CC1Args.push_back(
"-Werror=atomic-alignment");
884 ArgStringList &CC1Args)
const {
885 if (DriverArgs.hasArg(options::OPT_nostdinc) ||
886 DriverArgs.hasArg(options::OPT_nostdlibinc))
896 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
901 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
902 if (TargetID.empty())
903 return {std::nullopt, std::nullopt, std::nullopt};
905 llvm::StringMap<bool> FeatureMap;
907 if (!OptionalGpuArch)
908 return {TargetID.str(), std::nullopt, std::nullopt};
910 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
914 const llvm::opt::ArgList &DriverArgs)
const {
916 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
918 << *PTID.OptionalTargetID;
926 if (Arg *A = Args.getLastArg(options::OPT_offload_arch_tool_EQ))
927 Program = A->getValue();
933 return StdoutOrErr.takeError();
936 for (StringRef
Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
938 GPUArchs.push_back(
Arch.str());
940 if (GPUArchs.empty())
941 return llvm::createStringError(std::error_code(),
942 "No AMD GPU detected in the system");
944 return std::move(GPUArchs);
948 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
951 DeviceOffloadingKind);
956 DriverArgs.hasArg(options::OPT_nostdlib))
959 if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
964 const StringRef GpuArch =
getGPUArch(DriverArgs);
965 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
966 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
980 DriverArgs, LibDeviceFile, GpuArch, DeviceOffloadingKind,
983 for (
auto [BCFile, Internalize] : BCLibs) {
985 CC1Args.push_back(
"-mlink-builtin-bitcode");
987 CC1Args.push_back(
"-mlink-bitcode-file");
988 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
993 StringRef GPUArch, StringRef LibDeviceFile,
996 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
999 if (LibDeviceFile.empty()) {
1000 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
1007 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString() << 0;
1009 D.Diag(diag::err_drv_no_rocm_device_lib)
1010 << 2 << ABIVer.
toString() << 1 <<
"6.3";
1018 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
1020 const bool NeedsASanRT)
const {
1023 CommonBitcodeLibsPreferences Pref{D, DriverArgs, GPUArch,
1024 DeviceOffloadingKind, NeedsASanRT};
1027 bool Internalize =
true) {
1029 BCLibs.emplace_back(BCLib);
1031 auto AddSanBCLibs = [&]() {
1040 else if (Pref.GPUSan && Pref.IsOpenMP)
1047 AddBCLib(LibDeviceFile);
1049 if (!ABIVerPath.empty())
1050 AddBCLib(ABIVerPath);
1057 const llvm::opt::ArgList &DriverArgs,
const std::string &GPUArch,
1059 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1060 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1070 DriverArgs, LibDeviceFile, GPUArch, DeviceOffloadingKind,
1075 const ToolChain &TC,
const llvm::opt::ArgList &DriverArgs,
1076 StringRef TargetID,
const llvm::opt::Arg *A)
const {
1078 if (TargetID.empty())
1080 Option O = A->getOption();
1082 if (!O.matches(options::OPT_fsanitize_EQ))
1085 if (!DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
1086 options::OPT_fno_gpu_sanitize,
true))
1093 if (K != SanitizerKind::Address)
1097 llvm::StringRef Processor =
1099 auto ProcKind = TC.
getTriple().isAMDGCN()
1100 ? llvm::AMDGPU::parseArchAMDGCN(Processor)
1101 : llvm::AMDGPU::parseArchR600(Processor);
1102 auto Features = TC.
getTriple().isAMDGCN()
1103 ? llvm::AMDGPU::getArchAttrAMDGCN(ProcKind)
1104 : llvm::AMDGPU::getArchAttrR600(ProcKind);
1105 if (Features & llvm::AMDGPU::FEATURE_XNACK_ALWAYS)
1109 llvm::StringMap<bool> FeatureMap;
1111 assert(OptionalGpuArch &&
"Invalid Target ID");
1112 (void)OptionalGpuArch;
1113 auto Loc = FeatureMap.find(
"xnack");
1114 if (Loc == FeatureMap.end() || !Loc->second) {
1116 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1117 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
static void Append(char *Start, char *End, char *&Buffer, unsigned &BufferSize, unsigned &BufferCapacity)
__device__ __2f16 float c
const char * getOffloadingArch() const
OffloadKind getOffloadingDeviceKind() const
Compilation - A set of tasks to perform for a single driver invocation.
Driver - Encapsulate logic for constructing compilation processes from a set of gcc-driver-like comma...
DiagnosticsEngine & getDiags() const
llvm::Expected< std::unique_ptr< llvm::MemoryBuffer > > executeProgram(llvm::ArrayRef< llvm::StringRef > Args) const
DiagnosticBuilder Diag(unsigned DiagID) const
const llvm::opt::OptTable & getOpts() const
StringRef getCorrectlyRoundedSqrtPath(bool Enabled) const
StringRef getIncludePath() const
Get the detected path to Rocm's bin directory.
StringRef getUnsafeMathPath(bool Enabled) const
StringRef getOCMLPath() const
StringRef getAsanRTLPath() const
Returns empty string of Asan runtime library is not available.
RocmInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args, bool DetectHIPRuntime=true, bool DetectDeviceLib=false)
StringRef getOCKLPath() const
StringRef getFiniteOnlyPath(bool Enabled) const
bool hasDeviceLibrary() const
Check whether we detected a valid ROCm device library.
StringRef getDenormalsAreZeroPath(bool Enabled) const
bool checkCommonBitcodeLibs(StringRef GPUArch, StringRef LibDeviceFile, DeviceLibABIVersion ABIVer) const
Check file paths of default bitcode libraries common to AMDGPU based toolchains.
bool hasHIPStdParLibrary() const
Check whether we detected a valid HIP STDPAR Acceleration library.
StringRef getABIVersionPath(DeviceLibABIVersion ABIVer) const
llvm::SmallVector< ToolChain::BitCodeLibraryInfo, 12 > getCommonBitcodeLibs(const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, StringRef GPUArch, const Action::OffloadKind DeviceOffloadingKind, const bool NeedsASanRT) const
Get file paths of default bitcode libraries common to AMDGPU based toolchains.
bool hasHIPRuntime() const
Check whether we detected a valid HIP runtime.
void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const
void detectDeviceLibrary()
StringRef getWavefrontSize64Path(bool Enabled) const
void print(raw_ostream &OS) const
Print information about the detected ROCm installation.
SmallVector< InputInfo, 4 > InputInfoList
The JSON file list parser is used to communicate input to InstallAPI.
if(T->getSizeExpr()) TRY_TO(TraverseStmt(const_cast< Expr * >(T -> getSizeExpr())))
std::optional< llvm::StringRef > parseTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch, llvm::StringMap< bool > *FeatureMap)
Parse a target ID to get processor and feature map.
llvm::StringRef getProcessorFromTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch)
Get processor name from target ID.
llvm::SmallVector< llvm::StringRef, 4 > getAllPossibleTargetIDFeatures(const llvm::Triple &T, llvm::StringRef Processor)
Get all feature strings that can be used in target ID for Processor.
SanitizerMask parseSanitizerValue(StringRef Value, bool AllowGroups)
Parse a single value from a -fsanitize= or -fno-sanitize= value list.
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()