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,
",")));
637 if (Args.hasArg(options::OPT_stdlib))
638 CmdArgs.append({
"-lc",
"-lm"});
639 if (Args.hasArg(options::OPT_startfiles)) {
640 std::optional<std::string> IncludePath =
getToolChain().getStdlibPath();
642 IncludePath =
"/lib";
644 llvm::sys::path::append(P,
"crt1.o");
645 CmdArgs.push_back(Args.MakeArgString(P));
648 CmdArgs.push_back(
"-o");
650 C.addCommand(std::make_unique<Command>(
652 CmdArgs, Inputs, Output));
656 const llvm::Triple &Triple,
657 const llvm::opt::ArgList &Args,
658 std::vector<StringRef> &Features) {
662 if (Args.hasArg(options::OPT_mcpu_EQ))
663 TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
664 else if (Args.hasArg(options::OPT_march_EQ))
665 TargetID = Args.getLastArgValue(options::OPT_march_EQ);
666 if (!TargetID.empty()) {
667 llvm::StringMap<bool> FeatureMap;
668 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
669 if (OptionalGpuArch) {
670 StringRef GpuArch = *OptionalGpuArch;
676 auto Pos = FeatureMap.find(
Feature);
677 if (Pos == FeatureMap.end())
679 Features.push_back(Args.MakeArgStringRef(
680 (Twine(Pos->second ?
"+" :
"-") +
Feature).str()));
685 if (Args.hasFlag(options::OPT_mwavefrontsize64,
686 options::OPT_mno_wavefrontsize64,
false))
687 Features.push_back(
"+wavefrontsize64");
689 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
690 options::OPT_mno_amdgpu_precise_memory_op,
false))
691 Features.push_back(
"+precise-memory");
694 options::OPT_m_amdgpu_Features_Group);
702 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
718 DerivedArgList *DAL =
724 DAL =
new DerivedArgList(Args.getBaseArgs());
730 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
731 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
732 DAL->eraseArg(options::OPT_mcpu_EQ);
736 << llvm::Triple::getArchTypeName(
getArch())
737 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
739 auto &GPUs = *GPUsOrErr;
740 if (llvm::SmallSet<std::string, 1>(GPUs.begin(), GPUs.end()).size() > 1)
742 << llvm::Triple::getArchTypeName(
getArch())
743 << llvm::join(GPUs,
", ") <<
"-mcpu";
744 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
745 Args.MakeArgString(GPUs.front()));
751 if (Args.getLastArgValue(options::OPT_x) !=
"cl")
755 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
756 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
758 : options::OPT_m32));
762 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
764 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
772 llvm::AMDGPU::GPUKind Kind) {
775 if (Kind == llvm::AMDGPU::GK_NONE)
778 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
782 const bool BothDenormAndFMAFast =
783 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
784 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
785 return !BothDenormAndFMAFast;
789 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
790 const llvm::fltSemantics *FPType)
const {
792 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
793 return llvm::DenormalMode::getIEEE();
798 auto Kind = llvm::AMDGPU::parseArchAMDGCN(
Arch);
799 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
800 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
801 options::OPT_fno_gpu_flush_denormals_to_zero,
803 return llvm::DenormalMode::getPreserveSign();
805 return llvm::DenormalMode::getIEEE();
808 const StringRef GpuArch =
getGPUArch(DriverArgs);
809 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
813 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
818 return DAZ ? llvm::DenormalMode::getPreserveSign() :
819 llvm::DenormalMode::getIEEE();
823 llvm::AMDGPU::GPUKind Kind) {
824 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
825 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
827 return !HasWave32 || DriverArgs.hasFlag(
828 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
836 if (Triple.getEnvironment() != llvm::Triple::LLVM)
841 const llvm::opt::ArgList &DriverArgs,
842 llvm::opt::ArgStringList &CC1Args,
846 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
847 options::OPT_fvisibility_ms_compat)) {
848 CC1Args.push_back(
"-fvisibility=hidden");
849 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
863 !DriverArgs.hasArg(options::OPT_disable_llvm_optzns))
864 CC1Args.push_back(
"-disable-llvm-optzns");
873 CC1Args.push_back(
"-Werror=atomic-alignment");
877 ArgStringList &CC1Args)
const {
878 if (DriverArgs.hasArg(options::OPT_nostdinc) ||
879 DriverArgs.hasArg(options::OPT_nostdlibinc))
889 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
894 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
895 if (TargetID.empty())
896 return {std::nullopt, std::nullopt, std::nullopt};
898 llvm::StringMap<bool> FeatureMap;
900 if (!OptionalGpuArch)
901 return {TargetID.str(), std::nullopt, std::nullopt};
903 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
907 const llvm::opt::ArgList &DriverArgs)
const {
909 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
911 << *PTID.OptionalTargetID;
919 if (Arg *A = Args.getLastArg(options::OPT_offload_arch_tool_EQ))
920 Program = A->getValue();
926 return StdoutOrErr.takeError();
929 for (StringRef
Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
931 GPUArchs.push_back(
Arch.str());
933 if (GPUArchs.empty())
934 return llvm::createStringError(std::error_code(),
935 "No AMD GPU detected in the system");
937 return std::move(GPUArchs);
941 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
944 DeviceOffloadingKind);
949 DriverArgs.hasArg(options::OPT_nostdlib))
952 if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
964 if (TT.getEnvironment() == llvm::Triple::LLVM)
968 const StringRef GpuArch =
getGPUArch(DriverArgs);
969 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
970 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
984 DriverArgs, LibDeviceFile, GpuArch, DeviceOffloadingKind,
987 for (
auto [BCFile, Internalize] : BCLibs) {
989 CC1Args.push_back(
"-mlink-builtin-bitcode");
991 CC1Args.push_back(
"-mlink-bitcode-file");
992 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
997 StringRef GPUArch, StringRef LibDeviceFile,
1000 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
1003 if (LibDeviceFile.empty()) {
1004 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
1011 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString() << 0;
1013 D.Diag(diag::err_drv_no_rocm_device_lib)
1014 << 2 << ABIVer.
toString() << 1 <<
"6.3";
1022 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
1024 const bool NeedsASanRT)
const {
1027 CommonBitcodeLibsPreferences Pref{D, DriverArgs, GPUArch,
1028 DeviceOffloadingKind, NeedsASanRT};
1031 bool Internalize =
true) {
1032 if (!BCLib.
Path.empty()) {
1034 BCLibs.emplace_back(BCLib);
1037 auto AddSanBCLibs = [&]() {
1046 else if (Pref.GPUSan && Pref.IsOpenMP)
1051 AddBCLib(LibDeviceFile);
1053 if (!ABIVerPath.empty())
1054 AddBCLib(ABIVerPath);
1061 const llvm::opt::ArgList &DriverArgs,
const std::string &GPUArch,
1063 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1064 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1074 DriverArgs, LibDeviceFile, GPUArch, DeviceOffloadingKind,
1079 const ToolChain &TC,
const llvm::opt::ArgList &DriverArgs,
1080 StringRef TargetID,
const llvm::opt::Arg *A)
const {
1082 bool IsExplicitDevice =
1083 A->getBaseArg().getOption().matches(options::OPT_Xarch_device);
1086 llvm::StringRef Processor =
1088 auto ProcKind = TC.
getTriple().isAMDGCN()
1089 ? llvm::AMDGPU::parseArchAMDGCN(Processor)
1090 : llvm::AMDGPU::parseArchR600(Processor);
1091 auto Features = TC.
getTriple().isAMDGCN()
1092 ? llvm::AMDGPU::getArchAttrAMDGCN(ProcKind)
1093 : llvm::AMDGPU::getArchAttrR600(ProcKind);
1094 if (Features & llvm::AMDGPU::FEATURE_XNACK_ALWAYS)
1098 llvm::StringMap<bool> FeatureMap;
1100 assert(OptionalGpuArch &&
"Invalid Target ID");
1101 (void)OptionalGpuArch;
1102 auto Loc = FeatureMap.find(
"xnack");
1103 if (Loc == FeatureMap.end() || !Loc->second) {
1104 if (IsExplicitDevice) {
1106 clang::diag::err_drv_unsupported_option_for_offload_arch_req_feature)
1107 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
1110 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1111 << 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
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()