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_correctly_rounded_sqrt_on") {
131 CorrectlyRoundedSqrt.On = FilePath;
132 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_off") {
133 CorrectlyRoundedSqrt.Off = FilePath;
134 }
else if (BaseName ==
"oclc_unsafe_math_on") {
135 UnsafeMath.On = FilePath;
136 }
else if (BaseName ==
"oclc_unsafe_math_off") {
137 UnsafeMath.Off = FilePath;
138 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
139 WavefrontSize64.On = FilePath;
140 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
141 WavefrontSize64.Off = FilePath;
142 }
else if (BaseName.starts_with(ABIVersionPrefix)) {
143 unsigned ABIVersionNumber;
144 if (BaseName.drop_front(ABIVersionPrefix.size())
145 .getAsInteger(0, ABIVersionNumber))
147 ABIVersionMap[ABIVersionNumber] = FilePath.str();
151 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
152 if (!BaseName.starts_with(DeviceLibPrefix))
155 StringRef IsaVersionNumber =
156 BaseName.drop_front(DeviceLibPrefix.size());
158 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
161 std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
168bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
169 SmallVector<StringRef, 4> VersionParts;
170 V.split(VersionParts,
'\n');
171 unsigned Major = ~0U;
172 unsigned Minor = ~0U;
173 for (
auto Part : VersionParts) {
174 auto Splits = Part.rtrim().split(
'=');
175 if (Splits.first ==
"HIP_VERSION_MAJOR") {
176 if (Splits.second.getAsInteger(0, Major))
178 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
179 if (Splits.second.getAsInteger(0, Minor))
181 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
182 VersionPatch = Splits.second.str();
184 if (Major == ~0U || Minor == ~0U)
186 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
188 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
194const SmallVectorImpl<RocmInstallationDetector::Candidate> &
195RocmInstallationDetector::getInstallationPathCandidates() {
198 if (!ROCmSearchDirs.empty())
199 return ROCmSearchDirs;
201 auto DoPrintROCmSearchDirs = [&]() {
202 if (PrintROCmSearchDirs)
203 for (
auto Cand : ROCmSearchDirs) {
204 llvm::errs() <<
"ROCm installation search path: " << Cand.Path <<
'\n';
210 if (!RocmPathArg.empty()) {
211 ROCmSearchDirs.emplace_back(RocmPathArg.str());
212 DoPrintROCmSearchDirs();
213 return ROCmSearchDirs;
214 }
else if (std::optional<std::string> RocmPathEnv =
215 llvm::sys::Process::GetEnv(
"ROCM_PATH")) {
216 if (!RocmPathEnv->empty()) {
217 ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));
218 DoPrintROCmSearchDirs();
219 return ROCmSearchDirs;
224 StringRef InstallDir = D.Dir;
229 auto DeduceROCmPath = [](StringRef ClangPath) {
231 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
232 StringRef ParentName = llvm::sys::path::filename(ParentDir);
235 if (ParentName ==
"bin") {
236 ParentDir = llvm::sys::path::parent_path(ParentDir);
237 ParentName = llvm::sys::path::filename(ParentDir);
242 if (ParentName ==
"llvm" || ParentName.starts_with(
"aomp")) {
243 ParentDir = llvm::sys::path::parent_path(ParentDir);
244 ParentName = llvm::sys::path::filename(ParentDir);
248 if (ParentName ==
"lib")
249 ParentDir = llvm::sys::path::parent_path(ParentDir);
252 return Candidate(ParentDir.str(),
true);
257 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
261 llvm::SmallString<256> RealClangPath;
262 llvm::sys::fs::real_path(D.getClangProgramPath(), RealClangPath);
263 auto ParentPath = llvm::sys::path::parent_path(RealClangPath);
264 if (ParentPath != InstallDir)
265 ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));
268 auto ClangRoot = llvm::sys::path::parent_path(InstallDir);
269 auto RealClangRoot = llvm::sys::path::parent_path(ParentPath);
270 ROCmSearchDirs.emplace_back(ClangRoot.str(),
true);
271 if (RealClangRoot != ClangRoot)
272 ROCmSearchDirs.emplace_back(RealClangRoot.str(),
true);
273 ROCmSearchDirs.emplace_back(D.ResourceDir,
276 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/rocm",
281 std::string LatestROCm;
282 llvm::VersionTuple LatestVer;
284 auto GetROCmVersion = [](StringRef DirName) {
285 llvm::VersionTuple
V;
286 std::string VerStr = DirName.drop_front(strlen(
"rocm-")).str();
289 llvm::replace(VerStr,
'-',
'.');
293 for (llvm::vfs::directory_iterator
294 File = D.getVFS().dir_begin(D.SysRoot +
"/opt", EC),
296 File != FileEnd && !EC;
File.increment(EC)) {
297 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
300 if (LatestROCm.empty()) {
302 LatestVer = GetROCmVersion(LatestROCm);
305 auto Ver = GetROCmVersion(
FileName);
306 if (LatestVer < Ver) {
311 if (!LatestROCm.empty())
312 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/" + LatestROCm,
315 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr/local",
317 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr",
320 DoPrintROCmSearchDirs();
321 return ROCmSearchDirs;
325 const Driver &D,
const llvm::Triple &HostTriple,
326 const llvm::opt::ArgList &Args,
bool DetectHIPRuntime,
bool DetectDeviceLib)
328 Verbose = Args.hasArg(options::OPT_v);
329 RocmPathArg = Args.getLastArgValue(options::OPT_rocm_path_EQ);
330 PrintROCmSearchDirs = Args.hasArg(options::OPT_print_rocm_search_dirs);
331 RocmDeviceLibPathArg =
332 Args.getAllArgValues(options::OPT_rocm_device_lib_path_EQ);
333 HIPPathArg = Args.getLastArgValue(options::OPT_hip_path_EQ);
334 HIPStdParPathArg = Args.getLastArgValue(options::OPT_hipstdpar_path_EQ);
335 HasHIPStdParLibrary =
336 !HIPStdParPathArg.empty() && D.getVFS().exists(HIPStdParPathArg +
337 "/hipstdpar_lib.hpp");
338 HIPRocThrustPathArg =
339 Args.getLastArgValue(options::OPT_hipstdpar_thrust_path_EQ);
340 HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&
341 D.getVFS().exists(HIPRocThrustPathArg +
"/thrust");
342 HIPRocPrimPathArg = Args.getLastArgValue(options::OPT_hipstdpar_prim_path_EQ);
343 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
344 D.getVFS().exists(HIPRocPrimPathArg +
"/rocprim");
346 if (
auto *A = Args.getLastArg(options::OPT_hip_version_EQ)) {
347 HIPVersionArg = A->getValue();
348 unsigned Major = ~0U;
349 unsigned Minor = ~0U;
350 SmallVector<StringRef, 3> Parts;
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 != ~0U && Minor == ~0U)
362 if (Major == ~0U || Minor == ~0U)
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);
393 auto &FS = D.getVFS();
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())
428 LibDevicePath = D.ResourceDir;
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());
458 auto &FS = D.getVFS();
460 for (
const auto &Candidate : HIPSearchDirs) {
461 InstallPath = Candidate.Path;
462 if (InstallPath.empty() || !FS.exists(InstallPath))
465 BinPath = InstallPath;
466 llvm::sys::path::append(BinPath,
"bin");
467 IncludePath = InstallPath;
468 llvm::sys::path::append(IncludePath,
"include");
469 LibPath = InstallPath;
470 llvm::sys::path::append(LibPath,
"lib");
471 SharePath = InstallPath;
472 llvm::sys::path::append(SharePath,
"share");
475 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
476 llvm::sys::path::append(ParentSharePath,
"share");
479 const Twine &
c =
"",
const Twine &d =
"") {
481 llvm::sys::path::append(newpath, a,
b,
c, d);
485 std::vector<SmallString<0>> VersionFilePaths = {
486 Append(SharePath,
"hip",
"version"),
487 InstallPath != D.SysRoot +
"/usr/local"
488 ?
Append(ParentSharePath,
"hip",
"version")
490 Append(BinPath,
".hipVersion")};
492 for (
const auto &VersionFilePath : VersionFilePaths) {
493 if (VersionFilePath.empty())
495 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
496 FS.getBufferForFile(VersionFilePath);
499 if (HIPVersionArg.empty() && VersionFile)
500 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
503 HasHIPRuntime =
true;
508 if (!Candidate.StrictChecking) {
509 HasHIPRuntime =
true;
513 HasHIPRuntime =
false;
518 OS <<
"Found HIP installation: " << InstallPath <<
", version "
519 << DetectedVersion <<
'\n';
523 ArgStringList &CC1Args)
const {
524 bool UsesRuntimeWrapper = VersionMajorMinor > llvm::VersionTuple(3, 5) &&
525 !DriverArgs.hasArg(options::OPT_nohipwrapperinc);
526 bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
528 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
543 if (UsesRuntimeWrapper)
544 llvm::sys::path::append(P,
"include",
"cuda_wrappers");
545 CC1Args.push_back(
"-internal-isystem");
546 CC1Args.push_back(DriverArgs.MakeArgString(P));
549 const auto HandleHipStdPar = [=, &DriverArgs, &CC1Args]() {
551 auto &FS = D.getVFS();
554 if (!HIPStdParPathArg.empty() ||
555 !FS.exists(Inc +
"/thrust/system/hip/hipstdpar/hipstdpar_lib.hpp")) {
556 D.Diag(diag::err_drv_no_hipstdpar_lib);
559 if (!HasRocThrustLibrary && !FS.exists(Inc +
"/thrust")) {
560 D.Diag(diag::err_drv_no_hipstdpar_thrust_lib);
563 if (!HasRocPrimLibrary && !FS.exists(Inc +
"/rocprim")) {
564 D.Diag(diag::err_drv_no_hipstdpar_prim_lib);
567 const char *ThrustPath;
568 if (HasRocThrustLibrary)
569 ThrustPath = DriverArgs.MakeArgString(HIPRocThrustPathArg);
571 ThrustPath = DriverArgs.MakeArgString(Inc +
"/thrust");
573 const char *HIPStdParPath;
575 HIPStdParPath = DriverArgs.MakeArgString(HIPStdParPathArg);
577 HIPStdParPath = DriverArgs.MakeArgString(StringRef(ThrustPath) +
578 "/system/hip/hipstdpar");
580 const char *PrimPath;
581 if (HasRocPrimLibrary)
582 PrimPath = DriverArgs.MakeArgString(HIPRocPrimPathArg);
584 PrimPath = DriverArgs.MakeArgString(
getIncludePath() +
"/rocprim");
586 CC1Args.append({
"-idirafter", ThrustPath,
"-idirafter", PrimPath,
587 "-idirafter", HIPStdParPath,
"-include",
588 "hipstdpar_lib.hpp"});
591 if (!DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
600 D.Diag(diag::err_drv_no_hip_runtime);
604 CC1Args.push_back(
"-idirafter");
606 if (UsesRuntimeWrapper)
607 CC1Args.append({
"-include",
"__clang_hip_runtime_wrapper.h"});
616 const char *LinkingOutput)
const {
618 ArgStringList CmdArgs;
619 if (!Args.hasArg(options::OPT_r)) {
620 CmdArgs.push_back(
"--no-undefined");
621 CmdArgs.push_back(
"-shared");
624 if (
C.getDriver().isUsingLTO()) {
625 const bool ThinLTO = (
C.getDriver().getLTOMode() ==
LTOK_Thin);
627 }
else if (Args.hasArg(options::OPT_mcpu_EQ)) {
628 CmdArgs.push_back(Args.MakeArgString(
629 "-plugin-opt=mcpu=" +
631 Args.getLastArgValue(options::OPT_mcpu_EQ))));
635 Args.AddAllArgs(CmdArgs, options::OPT_L);
639 std::vector<StringRef> Features;
642 if (!Features.empty()) {
644 Args.MakeArgString(
"-plugin-opt=-mattr=" + llvm::join(Features,
",")));
647 if (Args.hasArg(options::OPT_stdlib))
648 CmdArgs.append({
"-lc",
"-lm"});
649 if (Args.hasArg(options::OPT_startfiles)) {
650 std::optional<std::string> IncludePath =
getToolChain().getStdlibPath();
652 IncludePath =
"/lib";
654 llvm::sys::path::append(P,
"crt1.o");
655 CmdArgs.push_back(Args.MakeArgString(P));
658 CmdArgs.push_back(
"-o");
660 C.addCommand(std::make_unique<Command>(
662 CmdArgs, Inputs, Output));
666 const llvm::Triple &Triple,
667 const llvm::opt::ArgList &Args,
668 std::vector<StringRef> &Features) {
672 if (Args.hasArg(options::OPT_mcpu_EQ))
673 TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
674 else if (Args.hasArg(options::OPT_march_EQ))
675 TargetID = Args.getLastArgValue(options::OPT_march_EQ);
676 if (!TargetID.empty()) {
677 llvm::StringMap<bool> FeatureMap;
678 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
679 if (OptionalGpuArch) {
680 StringRef GpuArch = *OptionalGpuArch;
686 auto Pos = FeatureMap.find(
Feature);
687 if (Pos == FeatureMap.end())
689 Features.push_back(Args.MakeArgStringRef(
690 (Twine(Pos->second ?
"+" :
"-") +
Feature).str()));
695 if (Args.hasFlag(options::OPT_mwavefrontsize64,
696 options::OPT_mno_wavefrontsize64,
false))
697 Features.push_back(
"+wavefrontsize64");
699 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
700 options::OPT_mno_amdgpu_precise_memory_op,
false))
701 Features.push_back(
"+precise-memory");
704 options::OPT_m_amdgpu_Features_Group);
712 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
728 DerivedArgList *DAL =
734 DAL =
new DerivedArgList(Args.getBaseArgs());
740 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
741 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
742 DAL->eraseArg(options::OPT_mcpu_EQ);
746 << llvm::Triple::getArchTypeName(
getArch())
747 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
749 auto &GPUs = *GPUsOrErr;
750 if (GPUs.size() > 1) {
752 << llvm::Triple::getArchTypeName(
getArch())
753 << llvm::join(GPUs,
", ") <<
"-mcpu";
755 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
756 Args.MakeArgString(GPUs.front()));
762 if (Args.getLastArgValue(options::OPT_x) !=
"cl")
766 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
767 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
769 : options::OPT_m32));
773 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
775 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
783 llvm::AMDGPU::GPUKind Kind) {
786 if (Kind == llvm::AMDGPU::GK_NONE)
789 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
793 const bool BothDenormAndFMAFast =
794 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
795 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
796 return !BothDenormAndFMAFast;
800 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
801 const llvm::fltSemantics *FPType)
const {
803 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
804 return llvm::DenormalMode::getIEEE();
809 auto Kind = llvm::AMDGPU::parseArchAMDGCN(
Arch);
810 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
811 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
812 options::OPT_fno_gpu_flush_denormals_to_zero,
814 return llvm::DenormalMode::getPreserveSign();
816 return llvm::DenormalMode::getIEEE();
819 const StringRef GpuArch =
getGPUArch(DriverArgs);
820 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
824 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
829 return DAZ ? llvm::DenormalMode::getPreserveSign() :
830 llvm::DenormalMode::getIEEE();
834 llvm::AMDGPU::GPUKind Kind) {
835 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
836 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
838 return !HasWave32 || DriverArgs.hasFlag(
839 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
851 const llvm::opt::ArgList &DriverArgs,
852 llvm::opt::ArgStringList &CC1Args,
856 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
857 options::OPT_fvisibility_ms_compat)) {
858 CC1Args.push_back(
"-fvisibility=hidden");
859 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
873 !DriverArgs.hasArg(options::OPT_disable_llvm_optzns))
874 CC1Args.push_back(
"-disable-llvm-optzns");
883 CC1Args.push_back(
"-Werror=atomic-alignment");
887 ArgStringList &CC1Args)
const {
888 if (DriverArgs.hasArg(options::OPT_nostdinc) ||
889 DriverArgs.hasArg(options::OPT_nostdlibinc))
899 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
904 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
905 if (TargetID.empty())
906 return {std::nullopt, std::nullopt, std::nullopt};
908 llvm::StringMap<bool> FeatureMap;
910 if (!OptionalGpuArch)
911 return {TargetID.str(), std::nullopt, std::nullopt};
913 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
917 const llvm::opt::ArgList &DriverArgs)
const {
919 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
921 << *PTID.OptionalTargetID;
929 if (Arg *A = Args.getLastArg(options::OPT_offload_arch_tool_EQ))
930 Program = A->getValue();
936 return StdoutOrErr.takeError();
939 for (StringRef
Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
941 GPUArchs.push_back(
Arch.str());
943 if (GPUArchs.empty())
944 return llvm::createStringError(std::error_code(),
945 "No AMD GPU detected in the system");
947 return std::move(GPUArchs);
951 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
954 DeviceOffloadingKind);
959 DriverArgs.hasArg(options::OPT_nostdlib))
962 if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
971 const StringRef GpuArch =
getGPUArch(DriverArgs);
972 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
973 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
987 DriverArgs, LibDeviceFile, GpuArch, DeviceOffloadingKind,
990 for (
auto [BCFile, Internalize] : BCLibs) {
992 CC1Args.push_back(
"-mlink-builtin-bitcode");
994 CC1Args.push_back(
"-mlink-bitcode-file");
995 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
1000 StringRef GPUArch, StringRef LibDeviceFile,
1003 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
1006 if (LibDeviceFile.empty()) {
1007 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
1014 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString() << 0;
1016 D.Diag(diag::err_drv_no_rocm_device_lib)
1017 << 2 << ABIVer.
toString() << 1 <<
"6.3";
1025 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
1027 const bool NeedsASanRT)
const {
1030 CommonBitcodeLibsPreferences Pref{D, DriverArgs, GPUArch,
1031 DeviceOffloadingKind, NeedsASanRT};
1034 bool Internalize =
true) {
1036 BCLibs.emplace_back(BCLib);
1038 auto AddSanBCLibs = [&]() {
1047 else if (Pref.GPUSan && Pref.IsOpenMP)
1053 AddBCLib(LibDeviceFile);
1055 if (!ABIVerPath.empty())
1056 AddBCLib(ABIVerPath);
1063 const llvm::opt::ArgList &DriverArgs,
const std::string &GPUArch,
1065 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1066 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1076 DriverArgs, LibDeviceFile, GPUArch, DeviceOffloadingKind,
1081 const ToolChain &TC,
const llvm::opt::ArgList &DriverArgs,
1082 StringRef TargetID,
const llvm::opt::Arg *A)
const {
1084 bool IsExplicitDevice =
1085 A->getBaseArg().getOption().matches(options::OPT_Xarch_device);
1088 llvm::StringRef Processor =
1090 auto ProcKind = TC.
getTriple().isAMDGCN()
1091 ? llvm::AMDGPU::parseArchAMDGCN(Processor)
1092 : llvm::AMDGPU::parseArchR600(Processor);
1093 auto Features = TC.
getTriple().isAMDGCN()
1094 ? llvm::AMDGPU::getArchAttrAMDGCN(ProcKind)
1095 : llvm::AMDGPU::getArchAttrR600(ProcKind);
1096 if (Features & llvm::AMDGPU::FEATURE_XNACK_ALWAYS)
1100 llvm::StringMap<bool> FeatureMap;
1102 assert(OptionalGpuArch &&
"Invalid Target ID");
1103 (void)OptionalGpuArch;
1104 auto Loc = FeatureMap.find(
"xnack");
1105 if (Loc == FeatureMap.end() || !Loc->second) {
1106 if (IsExplicitDevice) {
1108 clang::diag::err_drv_unsupported_option_for_offload_arch_req_feature)
1109 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
1112 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1113 << 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.
bool checkCommonBitcodeLibs(StringRef GPUArch, StringRef LibDeviceFile, DeviceLibABIVersion ABIVer) const
Check file paths of default bitcode libraries common to AMDGPU based toolchains.
bool hasHIPStdParLibrary() const
Check whether we detected a valid HIP STDPAR Acceleration library.
StringRef getABIVersionPath(DeviceLibABIVersion ABIVer) const
llvm::SmallVector< ToolChain::BitCodeLibraryInfo, 12 > getCommonBitcodeLibs(const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, StringRef GPUArch, const Action::OffloadKind DeviceOffloadingKind, const bool NeedsASanRT) const
Get file paths of default bitcode libraries common to AMDGPU based toolchains.
bool hasHIPRuntime() const
Check whether we detected a valid HIP runtime.
void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const
void detectDeviceLibrary()
StringRef getWavefrontSize64Path(bool Enabled) const
void print(raw_ostream &OS) const
Print information about the detected ROCm installation.
SmallVector< InputInfo, 4 > InputInfoList
The JSON file list parser is used to communicate input to InstallAPI.
if(T->getSizeExpr()) TRY_TO(TraverseStmt(const_cast< Expr * >(T -> getSizeExpr())))
std::optional< llvm::StringRef > parseTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch, llvm::StringMap< bool > *FeatureMap)
Parse a target ID to get processor and feature map.
llvm::StringRef getProcessorFromTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch)
Get processor name from target ID.
llvm::SmallVector< llvm::StringRef, 4 > getAllPossibleTargetIDFeatures(const llvm::Triple &T, llvm::StringRef Processor)
Get all feature strings that can be used in target ID for Processor.
Diagnostic wrappers for TextAPI types for error reporting.
ABI version of device library.
unsigned getAsCodeObjectVersion() const
static DeviceLibABIVersion fromCodeObjectVersion(unsigned CodeObjectVersion)
bool requiresLibrary()
Whether ABI version bc file is requested.
static constexpr ResponseFileSupport AtFileCurCP()