11#include "clang/Config/config.h"
17#include "llvm/ADT/SmallSet.h"
18#include "llvm/ADT/StringExtras.h"
19#include "llvm/Option/ArgList.h"
20#include "llvm/Support/Error.h"
21#include "llvm/Support/LineIterator.h"
22#include "llvm/Support/Path.h"
23#include "llvm/Support/Process.h"
24#include "llvm/Support/VirtualFileSystem.h"
25#include "llvm/TargetParser/Host.h"
26#include "llvm/TargetParser/TargetParser.h"
28#include <system_error>
36RocmInstallationDetector::CommonBitcodeLibsPreferences::
37 CommonBitcodeLibsPreferences(
const Driver &D,
38 const llvm::opt::ArgList &DriverArgs,
41 const bool NeedsASanRT)
44 const auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
45 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
49 const bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
51 !HasWave32 || DriverArgs.hasFlag(options::OPT_mwavefrontsize64,
52 options::OPT_mno_wavefrontsize64,
false);
59 const bool DefaultDAZ =
60 (Kind == llvm::AMDGPU::GK_NONE)
62 : !((ArchAttr &
llvm::
AMDGPU::FEATURE_FAST_FMA_F32) &&
63 (ArchAttr &
llvm::
AMDGPU::FEATURE_FAST_DENORMAL_F32));
66 DAZ = IsKnownOffloading
67 ? DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
68 options::OPT_fno_gpu_flush_denormals_to_zero,
70 : DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) || DefaultDAZ;
72 FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only) ||
73 DriverArgs.hasFlag(options::OPT_ffinite_math_only,
74 options::OPT_fno_finite_math_only,
false);
77 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations) ||
78 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
79 options::OPT_fno_unsafe_math_optimizations,
false);
81 FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math) ||
82 DriverArgs.hasFlag(options::OPT_ffast_math,
83 options::OPT_fno_fast_math,
false);
87 GPUSan = (DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
88 options::OPT_fno_gpu_sanitize,
true) &&
92void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) {
93 assert(!Path.empty());
95 const StringRef Suffix(
".bc");
96 const StringRef Suffix2(
".amdgcn.bc");
99 for (llvm::vfs::directory_iterator LI = D.getVFS().dir_begin(Path, EC), LE;
100 !EC && LI != LE; LI = LI.increment(EC)) {
101 StringRef FilePath = LI->path();
102 StringRef
FileName = llvm::sys::path::filename(FilePath);
108 BaseName =
FileName.drop_back(Suffix2.size());
109 else if (
FileName.ends_with(Suffix))
110 BaseName =
FileName.drop_back(Suffix.size());
112 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
113 if (BaseName ==
"ocml") {
115 }
else if (BaseName ==
"ockl") {
117 }
else if (BaseName ==
"opencl") {
119 }
else if (BaseName ==
"asanrtl") {
121 }
else if (BaseName ==
"oclc_finite_only_off") {
122 FiniteOnly.Off = FilePath;
123 }
else if (BaseName ==
"oclc_finite_only_on") {
124 FiniteOnly.On = FilePath;
125 }
else if (BaseName ==
"oclc_unsafe_math_on") {
126 UnsafeMath.On = FilePath;
127 }
else if (BaseName ==
"oclc_unsafe_math_off") {
128 UnsafeMath.Off = FilePath;
129 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
130 WavefrontSize64.On = FilePath;
131 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
132 WavefrontSize64.Off = FilePath;
133 }
else if (BaseName.starts_with(ABIVersionPrefix)) {
134 unsigned ABIVersionNumber;
135 if (BaseName.drop_front(ABIVersionPrefix.size())
136 .getAsInteger(0, ABIVersionNumber))
138 ABIVersionMap[ABIVersionNumber] = FilePath.str();
142 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
143 if (!BaseName.starts_with(DeviceLibPrefix))
146 StringRef IsaVersionNumber =
147 BaseName.drop_front(DeviceLibPrefix.size());
149 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
151 LibDeviceMap.insert({GfxName.toStringRef(Tmp), FilePath.str()});
158bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
159 SmallVector<StringRef, 4> VersionParts;
160 V.split(VersionParts,
'\n');
161 unsigned Major = ~0U;
162 unsigned Minor = ~0U;
163 for (
auto Part : VersionParts) {
164 auto Splits = Part.rtrim().split(
'=');
165 if (Splits.first ==
"HIP_VERSION_MAJOR") {
166 if (Splits.second.getAsInteger(0, Major))
168 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
169 if (Splits.second.getAsInteger(0, Minor))
171 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
172 VersionPatch = Splits.second.str();
174 if (Major == ~0U || Minor == ~0U)
176 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
178 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
184const SmallVectorImpl<RocmInstallationDetector::Candidate> &
185RocmInstallationDetector::getInstallationPathCandidates() {
188 if (!ROCmSearchDirs.empty())
189 return ROCmSearchDirs;
191 auto DoPrintROCmSearchDirs = [&]() {
192 if (PrintROCmSearchDirs)
193 for (
auto Cand : ROCmSearchDirs) {
194 llvm::errs() <<
"ROCm installation search path: " << Cand.Path <<
'\n';
200 if (!RocmPathArg.empty()) {
201 ROCmSearchDirs.emplace_back(RocmPathArg.str());
202 DoPrintROCmSearchDirs();
203 return ROCmSearchDirs;
204 }
else if (std::optional<std::string> RocmPathEnv =
205 llvm::sys::Process::GetEnv(
"ROCM_PATH")) {
206 if (!RocmPathEnv->empty()) {
207 ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));
208 DoPrintROCmSearchDirs();
209 return ROCmSearchDirs;
214 StringRef InstallDir = D.Dir;
219 auto DeduceROCmPath = [](StringRef ClangPath) {
221 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
222 StringRef ParentName = llvm::sys::path::filename(ParentDir);
225 if (ParentName ==
"bin") {
226 ParentDir = llvm::sys::path::parent_path(ParentDir);
227 ParentName = llvm::sys::path::filename(ParentDir);
232 if (ParentName ==
"llvm" || ParentName.starts_with(
"aomp")) {
233 ParentDir = llvm::sys::path::parent_path(ParentDir);
234 ParentName = llvm::sys::path::filename(ParentDir);
238 if (ParentName ==
"lib")
239 ParentDir = llvm::sys::path::parent_path(ParentDir);
242 return Candidate(ParentDir.str(),
true);
247 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
251 llvm::SmallString<256> RealClangPath;
252 llvm::sys::fs::real_path(D.getClangProgramPath(), RealClangPath);
253 auto ParentPath = llvm::sys::path::parent_path(RealClangPath);
254 if (ParentPath != InstallDir)
255 ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));
258 auto ClangRoot = llvm::sys::path::parent_path(InstallDir);
259 auto RealClangRoot = llvm::sys::path::parent_path(ParentPath);
260 ROCmSearchDirs.emplace_back(ClangRoot.str(),
true);
261 if (RealClangRoot != ClangRoot)
262 ROCmSearchDirs.emplace_back(RealClangRoot.str(),
true);
263 ROCmSearchDirs.emplace_back(D.ResourceDir,
266 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/rocm",
271 std::string LatestROCm;
272 llvm::VersionTuple LatestVer;
274 auto GetROCmVersion = [](StringRef DirName) {
275 llvm::VersionTuple
V;
276 std::string VerStr = DirName.drop_front(strlen(
"rocm-")).str();
279 llvm::replace(VerStr,
'-',
'.');
283 for (llvm::vfs::directory_iterator
284 File = D.getVFS().dir_begin(D.SysRoot +
"/opt", EC),
286 File != FileEnd && !EC;
File.increment(EC)) {
287 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
290 if (LatestROCm.empty()) {
292 LatestVer = GetROCmVersion(LatestROCm);
295 auto Ver = GetROCmVersion(
FileName);
296 if (LatestVer < Ver) {
301 if (!LatestROCm.empty())
302 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/" + LatestROCm,
305 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr/local",
307 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr",
310 DoPrintROCmSearchDirs();
311 return ROCmSearchDirs;
315 const Driver &D,
const llvm::Triple &HostTriple,
316 const llvm::opt::ArgList &Args,
bool DetectHIPRuntime)
318 Verbose = Args.hasArg(options::OPT_v);
319 RocmPathArg = Args.getLastArgValue(options::OPT_rocm_path_EQ);
320 PrintROCmSearchDirs = Args.hasArg(options::OPT_print_rocm_search_dirs);
321 RocmDeviceLibPathArg =
322 Args.getAllArgValues(options::OPT_rocm_device_lib_path_EQ);
323 HIPPathArg = Args.getLastArgValue(options::OPT_hip_path_EQ);
324 HIPStdParPathArg = Args.getLastArgValue(options::OPT_hipstdpar_path_EQ);
325 HasHIPStdParLibrary =
326 !HIPStdParPathArg.empty() && D.getVFS().exists(HIPStdParPathArg +
327 "/hipstdpar_lib.hpp");
328 HIPRocThrustPathArg =
329 Args.getLastArgValue(options::OPT_hipstdpar_thrust_path_EQ);
330 HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&
331 D.getVFS().exists(HIPRocThrustPathArg +
"/thrust");
332 HIPRocPrimPathArg = Args.getLastArgValue(options::OPT_hipstdpar_prim_path_EQ);
333 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
334 D.getVFS().exists(HIPRocPrimPathArg +
"/rocprim");
336 if (
auto *A = Args.getLastArg(options::OPT_hip_version_EQ)) {
337 HIPVersionArg = A->getValue();
338 unsigned Major = ~0U;
339 unsigned Minor = ~0U;
340 SmallVector<StringRef, 3> Parts;
341 HIPVersionArg.split(Parts,
'.');
343 Parts[0].getAsInteger(0, Major);
344 if (Parts.size() > 1)
345 Parts[1].getAsInteger(0, Minor);
346 if (Parts.size() > 2)
347 VersionPatch = Parts[2].str();
348 if (VersionPatch.empty())
350 if (Major != ~0U && Minor == ~0U)
352 if (Major == ~0U || Minor == ~0U)
353 D.Diag(diag::err_drv_invalid_value)
354 << A->getAsString(Args) << HIPVersionArg;
356 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
358 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
360 VersionPatch = DefaultVersionPatch;
362 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
363 DetectedVersion = (Twine(DefaultVersionMajor) +
"." +
364 Twine(DefaultVersionMinor) +
"." + VersionPatch)
368 if (DetectHIPRuntime)
373 assert(LibDevicePath.empty());
375 if (!RocmDeviceLibPathArg.empty())
376 LibDevicePath = RocmDeviceLibPathArg.back();
377 else if (std::optional<std::string> LibPathEnv =
378 llvm::sys::Process::GetEnv(
"HIP_DEVICE_LIB_PATH"))
379 LibDevicePath = std::move(*LibPathEnv);
381 auto &FS = D.getVFS();
382 if (!LibDevicePath.empty()) {
386 if (!FS.exists(LibDevicePath))
389 scanLibDevicePath(LibDevicePath);
390 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
395 auto CheckDeviceLib = [&](StringRef Path,
bool StrictChecking) {
396 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
397 if (CheckLibDevice && !FS.exists(Path))
400 scanLibDevicePath(Path);
402 if (!NoBuiltinLibs) {
404 if (!allGenericLibsValid())
409 if (LibDeviceMap.empty())
416 LibDevicePath = D.ResourceDir;
417 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
418 "amdgcn",
"bitcode");
419 HasDeviceLibrary = CheckDeviceLib(LibDevicePath,
true);
420 if (HasDeviceLibrary)
425 auto &ROCmDirs = getInstallationPathCandidates();
426 for (
const auto &Candidate : ROCmDirs) {
427 LibDevicePath = Candidate.Path;
428 llvm::sys::path::append(LibDevicePath,
"amdgcn",
"bitcode");
429 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);
430 if (HasDeviceLibrary)
437 if (!HIPPathArg.empty())
438 HIPSearchDirs.emplace_back(HIPPathArg.str());
439 else if (std::optional<std::string> HIPPathEnv =
440 llvm::sys::Process::GetEnv(
"HIP_PATH")) {
441 if (!HIPPathEnv->empty())
442 HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));
444 if (HIPSearchDirs.empty())
445 HIPSearchDirs.append(getInstallationPathCandidates());
446 auto &FS = D.getVFS();
448 for (
const auto &Candidate : HIPSearchDirs) {
449 InstallPath = Candidate.Path;
450 if (InstallPath.empty() || !FS.exists(InstallPath))
453 BinPath = InstallPath;
454 llvm::sys::path::append(BinPath,
"bin");
455 IncludePath = InstallPath;
456 llvm::sys::path::append(IncludePath,
"include");
457 LibPath = InstallPath;
458 llvm::sys::path::append(LibPath,
"lib");
459 SharePath = InstallPath;
460 llvm::sys::path::append(SharePath,
"share");
463 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
464 llvm::sys::path::append(ParentSharePath,
"share");
467 const Twine &
c =
"",
const Twine &d =
"") {
469 llvm::sys::path::append(newpath, a,
b,
c, d);
473 std::vector<SmallString<0>> VersionFilePaths = {
474 Append(SharePath,
"hip",
"version"),
475 InstallPath != D.SysRoot +
"/usr/local"
476 ?
Append(ParentSharePath,
"hip",
"version")
478 Append(BinPath,
".hipVersion")};
480 for (
const auto &VersionFilePath : VersionFilePaths) {
481 if (VersionFilePath.empty())
483 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
484 FS.getBufferForFile(VersionFilePath);
487 if (HIPVersionArg.empty() && VersionFile)
488 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
491 HasHIPRuntime =
true;
496 if (!Candidate.StrictChecking) {
497 HasHIPRuntime =
true;
501 HasHIPRuntime =
false;
506 OS <<
"Found HIP installation: " << InstallPath <<
", version "
507 << DetectedVersion <<
'\n';
511 ArgStringList &CC1Args)
const {
512 bool UsesRuntimeWrapper = VersionMajorMinor > llvm::VersionTuple(3, 5) &&
513 !DriverArgs.hasArg(options::OPT_nohipwrapperinc);
514 bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
516 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
531 if (UsesRuntimeWrapper)
532 llvm::sys::path::append(P,
"include",
"cuda_wrappers");
533 CC1Args.push_back(
"-internal-isystem");
534 CC1Args.push_back(DriverArgs.MakeArgString(P));
537 const auto HandleHipStdPar = [=, &DriverArgs, &CC1Args]() {
539 auto &FS = D.getVFS();
542 if (!HIPStdParPathArg.empty() ||
543 !FS.exists(Inc +
"/thrust/system/hip/hipstdpar/hipstdpar_lib.hpp")) {
544 D.Diag(diag::err_drv_no_hipstdpar_lib);
547 if (!HasRocThrustLibrary && !FS.exists(Inc +
"/thrust")) {
548 D.Diag(diag::err_drv_no_hipstdpar_thrust_lib);
551 if (!HasRocPrimLibrary && !FS.exists(Inc +
"/rocprim")) {
552 D.Diag(diag::err_drv_no_hipstdpar_prim_lib);
555 const char *ThrustPath;
556 if (HasRocThrustLibrary)
557 ThrustPath = DriverArgs.MakeArgString(HIPRocThrustPathArg);
559 ThrustPath = DriverArgs.MakeArgString(Inc +
"/thrust");
561 const char *HIPStdParPath;
563 HIPStdParPath = DriverArgs.MakeArgString(HIPStdParPathArg);
565 HIPStdParPath = DriverArgs.MakeArgString(StringRef(ThrustPath) +
566 "/system/hip/hipstdpar");
568 const char *PrimPath;
569 if (HasRocPrimLibrary)
570 PrimPath = DriverArgs.MakeArgString(HIPRocPrimPathArg);
572 PrimPath = DriverArgs.MakeArgString(
getIncludePath() +
"/rocprim");
574 CC1Args.append({
"-idirafter", ThrustPath,
"-idirafter", PrimPath,
575 "-idirafter", HIPStdParPath,
"-include",
576 "hipstdpar_lib.hpp"});
579 if (!DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
588 D.Diag(diag::err_drv_no_hip_runtime);
592 CC1Args.push_back(
"-idirafter");
594 if (UsesRuntimeWrapper)
595 CC1Args.append({
"-include",
"__clang_hip_runtime_wrapper.h"});
604 const char *LinkingOutput)
const {
606 ArgStringList CmdArgs;
607 if (!Args.hasArg(options::OPT_r)) {
608 CmdArgs.push_back(
"--no-undefined");
609 CmdArgs.push_back(
"-shared");
612 if (
C.getDriver().isUsingLTO()) {
613 const bool ThinLTO = (
C.getDriver().getLTOMode() ==
LTOK_Thin);
615 }
else if (Args.hasArg(options::OPT_mcpu_EQ)) {
616 CmdArgs.push_back(Args.MakeArgString(
617 "-plugin-opt=mcpu=" +
619 Args.getLastArgValue(options::OPT_mcpu_EQ))));
623 Args.AddAllArgs(CmdArgs, options::OPT_L);
627 std::vector<StringRef> Features;
630 if (!Features.empty()) {
632 Args.MakeArgString(
"-plugin-opt=-mattr=" + llvm::join(Features,
",")));
638 if (Args.hasArg(options::OPT_stdlib))
639 CmdArgs.append({
"-lc",
"-lm"});
640 if (Args.hasArg(options::OPT_startfiles)) {
641 std::optional<std::string> IncludePath =
getToolChain().getStdlibPath();
643 IncludePath =
"/lib";
645 llvm::sys::path::append(P,
"crt1.o");
646 CmdArgs.push_back(Args.MakeArgString(P));
649 CmdArgs.push_back(
"-o");
651 C.addCommand(std::make_unique<Command>(
653 CmdArgs, Inputs, Output));
657 const llvm::Triple &Triple,
658 const llvm::opt::ArgList &Args,
659 std::vector<StringRef> &Features) {
663 if (Args.hasArg(options::OPT_mcpu_EQ))
664 TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
665 else if (Args.hasArg(options::OPT_march_EQ))
666 TargetID = Args.getLastArgValue(options::OPT_march_EQ);
667 if (!TargetID.empty()) {
668 llvm::StringMap<bool> FeatureMap;
669 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
670 if (OptionalGpuArch) {
671 StringRef GpuArch = *OptionalGpuArch;
677 auto Pos = FeatureMap.find(
Feature);
678 if (Pos == FeatureMap.end())
680 Features.push_back(Args.MakeArgStringRef(
681 (Twine(Pos->second ?
"+" :
"-") +
Feature).str()));
686 if (Args.hasFlag(options::OPT_mwavefrontsize64,
687 options::OPT_mno_wavefrontsize64,
false))
688 Features.push_back(
"+wavefrontsize64");
690 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
691 options::OPT_mno_amdgpu_precise_memory_op,
false))
692 Features.push_back(
"+precise-memory");
695 options::OPT_m_amdgpu_Features_Group);
703 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
704 loadMultilibsFromYAML(Args, D);
721 DerivedArgList *DAL =
727 DAL =
new DerivedArgList(Args.getBaseArgs());
733 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
734 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
735 DAL->eraseArg(options::OPT_mcpu_EQ);
739 << llvm::Triple::getArchTypeName(
getArch())
740 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
742 auto &GPUs = *GPUsOrErr;
743 if (llvm::SmallSet<std::string, 1>(GPUs.begin(), GPUs.end()).size() > 1)
745 << llvm::Triple::getArchTypeName(
getArch())
746 << llvm::join(GPUs,
", ") <<
"-mcpu";
747 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
748 Args.MakeArgString(GPUs.front()));
754 if (Args.getLastArgValue(options::OPT_x) !=
"cl")
758 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
759 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
761 : options::OPT_m32));
765 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
767 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
775 llvm::AMDGPU::GPUKind Kind) {
778 if (Kind == llvm::AMDGPU::GK_NONE)
781 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
785 const bool BothDenormAndFMAFast =
786 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
787 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
788 return !BothDenormAndFMAFast;
792 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
793 const llvm::fltSemantics *FPType)
const {
795 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
796 return llvm::DenormalMode::getIEEE();
801 auto Kind = llvm::AMDGPU::parseArchAMDGCN(
Arch);
802 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
803 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
804 options::OPT_fno_gpu_flush_denormals_to_zero,
806 return llvm::DenormalMode::getPreserveSign();
808 return llvm::DenormalMode::getIEEE();
811 const StringRef GpuArch =
getGPUArch(DriverArgs);
812 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
816 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
821 return DAZ ? llvm::DenormalMode::getPreserveSign() :
822 llvm::DenormalMode::getIEEE();
826 llvm::AMDGPU::GPUKind Kind) {
827 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
828 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
830 return !HasWave32 || DriverArgs.hasFlag(
831 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
839 if (Triple.getEnvironment() != llvm::Triple::LLVM)
844 const llvm::opt::ArgList &DriverArgs,
845 llvm::opt::ArgStringList &CC1Args,
850 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
851 options::OPT_fvisibility_ms_compat) &&
853 CC1Args.push_back(
"-fvisibility=hidden");
854 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
868 !DriverArgs.hasArg(options::OPT_disable_llvm_optzns))
869 CC1Args.push_back(
"-disable-llvm-optzns");
878 CC1Args.push_back(
"-Werror=atomic-alignment");
882 ArgStringList &CC1Args)
const {
883 if (DriverArgs.hasArg(options::OPT_nostdinc) ||
884 DriverArgs.hasArg(options::OPT_nostdlibinc))
893 llvm::sys::path::append(Dir, M.includeSuffix());
906 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
911 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
912 if (TargetID.empty())
913 return {std::nullopt, std::nullopt, std::nullopt};
915 llvm::StringMap<bool> FeatureMap;
917 if (!OptionalGpuArch)
918 return {TargetID.str(), std::nullopt, std::nullopt};
920 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
924 const llvm::opt::ArgList &DriverArgs)
const {
926 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
928 << *PTID.OptionalTargetID;
936 if (Arg *A = Args.getLastArg(options::OPT_offload_arch_tool_EQ))
937 Program = A->getValue();
943 return StdoutOrErr.takeError();
946 for (StringRef
Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
948 GPUArchs.push_back(
Arch.str());
950 if (GPUArchs.empty())
951 return llvm::createStringError(std::error_code(),
952 "No AMD GPU detected in the system");
954 return std::move(GPUArchs);
958 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
961 DeviceOffloadingKind);
966 DriverArgs.hasArg(options::OPT_nostdlib))
969 if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
981 if (TT.getEnvironment() == llvm::Triple::LLVM)
985 const StringRef GpuArch =
getGPUArch(DriverArgs);
986 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
987 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1001 DriverArgs, LibDeviceFile, GpuArch, DeviceOffloadingKind,
1004 for (
auto [BCFile, Internalize] : BCLibs) {
1006 CC1Args.push_back(
"-mlink-builtin-bitcode");
1008 CC1Args.push_back(
"-mlink-bitcode-file");
1009 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
1014 StringRef GPUArch, StringRef LibDeviceFile,
1017 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
1020 if (LibDeviceFile.empty()) {
1021 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
1028 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString() << 0;
1030 D.Diag(diag::err_drv_no_rocm_device_lib)
1031 << 2 << ABIVer.
toString() << 1 <<
"6.3";
1039 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
1041 const bool NeedsASanRT)
const {
1044 CommonBitcodeLibsPreferences Pref{D, DriverArgs, GPUArch,
1045 DeviceOffloadingKind, NeedsASanRT};
1048 bool Internalize =
true) {
1049 if (!BCLib.
Path.empty()) {
1051 BCLibs.emplace_back(BCLib);
1054 auto AddSanBCLibs = [&]() {
1063 else if (Pref.GPUSan && Pref.IsOpenMP)
1068 AddBCLib(LibDeviceFile);
1070 if (!ABIVerPath.empty())
1071 AddBCLib(ABIVerPath);
1078 const llvm::opt::ArgList &DriverArgs,
const std::string &GPUArch,
1080 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1081 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1091 DriverArgs, LibDeviceFile, GPUArch, DeviceOffloadingKind,
1096 const ToolChain &TC,
const llvm::opt::ArgList &DriverArgs,
1097 StringRef TargetID,
const llvm::opt::Arg *A)
const {
1099 bool IsExplicitDevice =
1100 A->getBaseArg().getOption().matches(options::OPT_Xarch_device);
1103 llvm::StringRef Processor =
1105 auto ProcKind = TC.
getTriple().isAMDGCN()
1106 ? llvm::AMDGPU::parseArchAMDGCN(Processor)
1107 : llvm::AMDGPU::parseArchR600(Processor);
1108 auto Features = TC.
getTriple().isAMDGCN()
1109 ? llvm::AMDGPU::getArchAttrAMDGCN(ProcKind)
1110 : llvm::AMDGPU::getArchAttrR600(ProcKind);
1111 if (Features & llvm::AMDGPU::FEATURE_XNACK_ALWAYS)
1115 llvm::StringMap<bool> FeatureMap;
1117 assert(OptionalGpuArch &&
"Invalid Target ID");
1118 (void)OptionalGpuArch;
1119 auto Loc = FeatureMap.find(
"xnack");
1120 if (Loc == FeatureMap.end() || !Loc->second) {
1121 if (IsExplicitDevice) {
1123 clang::diag::err_drv_unsupported_option_for_offload_arch_req_feature)
1124 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
1127 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1128 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
static StringRef getTriple(const Command &Job)
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
This corresponds to a single GCC Multilib, or a segment of one controlled by a command line flag.
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)
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()