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"
26#include <system_error>
34RocmInstallationDetector::CommonBitcodeLibsPreferences::
35 CommonBitcodeLibsPreferences(
const Driver &
D,
36 const llvm::opt::ArgList &DriverArgs,
39 const bool NeedsASanRT)
42 const auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
43 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
47 const bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
49 !HasWave32 || DriverArgs.hasFlag(options::OPT_mwavefrontsize64,
50 options::OPT_mno_wavefrontsize64,
false);
57 const bool DefaultDAZ =
58 (
Kind == llvm::AMDGPU::GK_NONE)
60 : !((ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
61 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32));
64 DAZ = IsKnownOffloading
65 ? DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
66 options::OPT_fno_gpu_flush_denormals_to_zero,
68 : DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) || DefaultDAZ;
70 FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only) ||
71 DriverArgs.hasFlag(options::OPT_ffinite_math_only,
72 options::OPT_fno_finite_math_only,
false);
75 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations) ||
76 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
77 options::OPT_fno_unsafe_math_optimizations,
false);
79 FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math) ||
80 DriverArgs.hasFlag(options::OPT_ffast_math,
81 options::OPT_fno_fast_math,
false);
83 const bool DefaultSqrt = IsKnownOffloading ?
true :
false;
85 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt) ||
87 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
88 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, DefaultSqrt);
91 GPUSan = (DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
92 options::OPT_fno_gpu_sanitize,
true) &&
96void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef
Path) {
97 assert(!
Path.empty());
99 const StringRef Suffix(
".bc");
100 const StringRef Suffix2(
".amdgcn.bc");
103 for (llvm::vfs::directory_iterator LI =
D.getVFS().dir_begin(
Path, EC), LE;
104 !EC && LI != LE; LI = LI.increment(EC)) {
105 StringRef FilePath = LI->path();
106 StringRef
FileName = llvm::sys::path::filename(FilePath);
112 BaseName =
FileName.drop_back(Suffix2.size());
113 else if (
FileName.ends_with(Suffix))
114 BaseName =
FileName.drop_back(Suffix.size());
116 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
117 if (BaseName ==
"ocml") {
119 }
else if (BaseName ==
"ockl") {
121 }
else if (BaseName ==
"opencl") {
123 }
else if (BaseName ==
"asanrtl") {
125 }
else if (BaseName ==
"oclc_finite_only_off") {
126 FiniteOnly.Off = FilePath;
127 }
else if (BaseName ==
"oclc_finite_only_on") {
128 FiniteOnly.On = FilePath;
129 }
else if (BaseName ==
"oclc_daz_opt_on") {
130 DenormalsAreZero.On = FilePath;
131 }
else if (BaseName ==
"oclc_daz_opt_off") {
132 DenormalsAreZero.Off = FilePath;
133 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_on") {
134 CorrectlyRoundedSqrt.On = FilePath;
135 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_off") {
136 CorrectlyRoundedSqrt.Off = FilePath;
137 }
else if (BaseName ==
"oclc_unsafe_math_on") {
138 UnsafeMath.On = FilePath;
139 }
else if (BaseName ==
"oclc_unsafe_math_off") {
140 UnsafeMath.Off = FilePath;
141 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
142 WavefrontSize64.On = FilePath;
143 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
144 WavefrontSize64.Off = FilePath;
145 }
else if (BaseName.starts_with(ABIVersionPrefix)) {
146 unsigned ABIVersionNumber;
147 if (BaseName.drop_front(ABIVersionPrefix.size())
148 .getAsInteger(0, ABIVersionNumber))
150 ABIVersionMap[ABIVersionNumber] = FilePath.str();
154 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
155 if (!BaseName.starts_with(DeviceLibPrefix))
158 StringRef IsaVersionNumber =
159 BaseName.drop_front(DeviceLibPrefix.size());
161 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
164 std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
171bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
173 V.split(VersionParts,
'\n');
174 unsigned Major = ~0
U;
175 unsigned Minor = ~0
U;
176 for (
auto Part : VersionParts) {
177 auto Splits = Part.rtrim().split(
'=');
178 if (Splits.first ==
"HIP_VERSION_MAJOR") {
179 if (Splits.second.getAsInteger(0, Major))
181 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
182 if (Splits.second.getAsInteger(0, Minor))
184 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
185 VersionPatch = Splits.second.str();
187 if (Major == ~0
U || Minor == ~0
U)
189 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
191 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
198RocmInstallationDetector::getInstallationPathCandidates() {
201 if (!ROCmSearchDirs.empty())
202 return ROCmSearchDirs;
204 auto DoPrintROCmSearchDirs = [&]() {
205 if (PrintROCmSearchDirs)
206 for (
auto Cand : ROCmSearchDirs) {
207 llvm::errs() <<
"ROCm installation search path: " << Cand.Path <<
'\n';
213 if (!RocmPathArg.empty()) {
214 ROCmSearchDirs.emplace_back(RocmPathArg.str());
215 DoPrintROCmSearchDirs();
216 return ROCmSearchDirs;
217 }
else if (std::optional<std::string> RocmPathEnv =
218 llvm::sys::Process::GetEnv(
"ROCM_PATH")) {
219 if (!RocmPathEnv->empty()) {
220 ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));
221 DoPrintROCmSearchDirs();
222 return ROCmSearchDirs;
227 StringRef InstallDir =
D.Dir;
232 auto DeduceROCmPath = [](StringRef ClangPath) {
234 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
235 StringRef ParentName = llvm::sys::path::filename(ParentDir);
238 if (ParentName ==
"bin") {
239 ParentDir = llvm::sys::path::parent_path(ParentDir);
240 ParentName = llvm::sys::path::filename(ParentDir);
245 if (ParentName ==
"llvm" || ParentName.starts_with(
"aomp"))
246 ParentDir = llvm::sys::path::parent_path(ParentDir);
248 return Candidate(ParentDir.str(),
true);
253 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
258 llvm::sys::fs::real_path(
D.getClangProgramPath(), RealClangPath);
259 auto ParentPath = llvm::sys::path::parent_path(RealClangPath);
260 if (ParentPath != InstallDir)
261 ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));
264 auto ClangRoot = llvm::sys::path::parent_path(InstallDir);
265 auto RealClangRoot = llvm::sys::path::parent_path(ParentPath);
266 ROCmSearchDirs.emplace_back(ClangRoot.str(),
true);
267 if (RealClangRoot != ClangRoot)
268 ROCmSearchDirs.emplace_back(RealClangRoot.str(),
true);
269 ROCmSearchDirs.emplace_back(
D.ResourceDir,
272 ROCmSearchDirs.emplace_back(
D.SysRoot +
"/opt/rocm",
277 std::string LatestROCm;
278 llvm::VersionTuple LatestVer;
280 auto GetROCmVersion = [](StringRef DirName) {
281 llvm::VersionTuple
V;
282 std::string VerStr = DirName.drop_front(strlen(
"rocm-")).str();
285 llvm::replace(VerStr,
'-',
'.');
289 for (llvm::vfs::directory_iterator
290 File =
D.getVFS().dir_begin(
D.SysRoot +
"/opt", EC),
292 File != FileEnd && !EC;
File.increment(EC)) {
293 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
296 if (LatestROCm.empty()) {
298 LatestVer = GetROCmVersion(LatestROCm);
301 auto Ver = GetROCmVersion(
FileName);
302 if (LatestVer < Ver) {
307 if (!LatestROCm.empty())
308 ROCmSearchDirs.emplace_back(
D.SysRoot +
"/opt/" + LatestROCm,
311 ROCmSearchDirs.emplace_back(
D.SysRoot +
"/usr/local",
313 ROCmSearchDirs.emplace_back(
D.SysRoot +
"/usr",
316 DoPrintROCmSearchDirs();
317 return ROCmSearchDirs;
321 const Driver &
D,
const llvm::Triple &HostTriple,
322 const llvm::opt::ArgList &Args,
bool DetectHIPRuntime,
bool DetectDeviceLib)
324 Verbose = Args.hasArg(options::OPT_v);
325 RocmPathArg = Args.getLastArgValue(clang::driver::options::OPT_rocm_path_EQ);
326 PrintROCmSearchDirs =
327 Args.hasArg(clang::driver::options::OPT_print_rocm_search_dirs);
328 RocmDeviceLibPathArg =
329 Args.getAllArgValues(clang::driver::options::OPT_rocm_device_lib_path_EQ);
330 HIPPathArg = Args.getLastArgValue(clang::driver::options::OPT_hip_path_EQ);
332 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_path_EQ);
333 HasHIPStdParLibrary =
334 !HIPStdParPathArg.empty() &&
D.getVFS().exists(HIPStdParPathArg +
335 "/hipstdpar_lib.hpp");
336 HIPRocThrustPathArg =
337 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_thrust_path_EQ);
338 HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&
339 D.getVFS().exists(HIPRocThrustPathArg +
"/thrust");
341 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_prim_path_EQ);
342 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
343 D.getVFS().exists(HIPRocPrimPathArg +
"/rocprim");
345 if (
auto *A = Args.getLastArg(clang::driver::options::OPT_hip_version_EQ)) {
346 HIPVersionArg = A->getValue();
347 unsigned Major = ~0
U;
348 unsigned Minor = ~0
U;
350 HIPVersionArg.split(Parts,
'.');
352 Parts[0].getAsInteger(0, Major);
353 if (Parts.size() > 1)
354 Parts[1].getAsInteger(0, Minor);
355 if (Parts.size() > 2)
356 VersionPatch = Parts[2].str();
357 if (VersionPatch.empty())
359 if (Major != ~0
U && Minor == ~0
U)
361 if (Major == ~0
U || Minor == ~0
U)
362 D.Diag(diag::err_drv_invalid_value)
363 << A->getAsString(Args) << HIPVersionArg;
365 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
367 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
369 VersionPatch = DefaultVersionPatch;
371 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
372 DetectedVersion = (Twine(DefaultVersionMajor) +
"." +
373 Twine(DefaultVersionMinor) +
"." + VersionPatch)
377 if (DetectHIPRuntime)
384 assert(LibDevicePath.empty());
386 if (!RocmDeviceLibPathArg.empty())
387 LibDevicePath = RocmDeviceLibPathArg[RocmDeviceLibPathArg.size() - 1];
388 else if (std::optional<std::string> LibPathEnv =
389 llvm::sys::Process::GetEnv(
"HIP_DEVICE_LIB_PATH"))
390 LibDevicePath = std::move(*LibPathEnv);
392 auto &FS =
D.getVFS();
393 if (!LibDevicePath.empty()) {
397 if (!FS.exists(LibDevicePath))
400 scanLibDevicePath(LibDevicePath);
401 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
406 auto CheckDeviceLib = [&](StringRef
Path,
bool StrictChecking) {
407 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
408 if (CheckLibDevice && !FS.exists(
Path))
411 scanLibDevicePath(
Path);
413 if (!NoBuiltinLibs) {
415 if (!allGenericLibsValid())
420 if (LibDeviceMap.empty())
427 LibDevicePath =
D.ResourceDir;
428 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
429 "amdgcn",
"bitcode");
430 HasDeviceLibrary = CheckDeviceLib(LibDevicePath,
true);
431 if (HasDeviceLibrary)
436 auto &ROCmDirs = getInstallationPathCandidates();
437 for (
const auto &Candidate : ROCmDirs) {
438 LibDevicePath = Candidate.Path;
439 llvm::sys::path::append(LibDevicePath,
"amdgcn",
"bitcode");
440 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);
441 if (HasDeviceLibrary)
448 if (!HIPPathArg.empty())
449 HIPSearchDirs.emplace_back(HIPPathArg.str());
450 else if (std::optional<std::string> HIPPathEnv =
451 llvm::sys::Process::GetEnv(
"HIP_PATH")) {
452 if (!HIPPathEnv->empty())
453 HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));
455 if (HIPSearchDirs.empty())
456 HIPSearchDirs.append(getInstallationPathCandidates());
457 auto &FS =
D.getVFS();
459 for (
const auto &Candidate : HIPSearchDirs) {
460 InstallPath = Candidate.Path;
461 if (InstallPath.empty() || !FS.exists(InstallPath))
464 BinPath = InstallPath;
465 llvm::sys::path::append(BinPath,
"bin");
466 IncludePath = InstallPath;
467 llvm::sys::path::append(IncludePath,
"include");
468 LibPath = InstallPath;
469 llvm::sys::path::append(LibPath,
"lib");
470 SharePath = InstallPath;
471 llvm::sys::path::append(SharePath,
"share");
474 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
475 llvm::sys::path::append(ParentSharePath,
"share");
478 const Twine &
c =
"",
const Twine &d =
"") {
480 llvm::sys::path::append(newpath, a,
b,
c, d);
484 std::vector<SmallString<0>> VersionFilePaths = {
485 Append(SharePath,
"hip",
"version"),
486 InstallPath !=
D.SysRoot +
"/usr/local"
487 ?
Append(ParentSharePath,
"hip",
"version")
489 Append(BinPath,
".hipVersion")};
491 for (
const auto &VersionFilePath : VersionFilePaths) {
492 if (VersionFilePath.empty())
494 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
495 FS.getBufferForFile(VersionFilePath);
498 if (HIPVersionArg.empty() && VersionFile)
499 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
502 HasHIPRuntime =
true;
507 if (!Candidate.StrictChecking) {
508 HasHIPRuntime =
true;
512 HasHIPRuntime =
false;
517 OS <<
"Found HIP installation: " << InstallPath <<
", version "
518 << DetectedVersion <<
'\n';
522 ArgStringList &CC1Args)
const {
523 bool UsesRuntimeWrapper = VersionMajorMinor > llvm::VersionTuple(3, 5) &&
524 !DriverArgs.hasArg(options::OPT_nohipwrapperinc);
525 bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
527 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
542 if (UsesRuntimeWrapper)
543 llvm::sys::path::append(
P,
"include",
"cuda_wrappers");
544 CC1Args.push_back(
"-internal-isystem");
545 CC1Args.push_back(DriverArgs.MakeArgString(
P));
548 const auto HandleHipStdPar = [=, &DriverArgs, &CC1Args]() {
550 auto &FS =
D.getVFS();
553 if (!HIPStdParPathArg.empty() ||
554 !FS.exists(Inc +
"/thrust/system/hip/hipstdpar/hipstdpar_lib.hpp")) {
555 D.Diag(diag::err_drv_no_hipstdpar_lib);
558 if (!HasRocThrustLibrary && !FS.exists(Inc +
"/thrust")) {
559 D.Diag(diag::err_drv_no_hipstdpar_thrust_lib);
562 if (!HasRocPrimLibrary && !FS.exists(Inc +
"/rocprim")) {
563 D.Diag(diag::err_drv_no_hipstdpar_prim_lib);
566 const char *ThrustPath;
567 if (HasRocThrustLibrary)
568 ThrustPath = DriverArgs.MakeArgString(HIPRocThrustPathArg);
570 ThrustPath = DriverArgs.MakeArgString(Inc +
"/thrust");
572 const char *HIPStdParPath;
574 HIPStdParPath = DriverArgs.MakeArgString(HIPStdParPathArg);
576 HIPStdParPath = DriverArgs.MakeArgString(StringRef(ThrustPath) +
577 "/system/hip/hipstdpar");
579 const char *PrimPath;
580 if (HasRocPrimLibrary)
581 PrimPath = DriverArgs.MakeArgString(HIPRocPrimPathArg);
583 PrimPath = DriverArgs.MakeArgString(
getIncludePath() +
"/rocprim");
585 CC1Args.append({
"-idirafter", ThrustPath,
"-idirafter", PrimPath,
586 "-idirafter", HIPStdParPath,
"-include",
587 "hipstdpar_lib.hpp"});
590 if (!DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
599 D.Diag(diag::err_drv_no_hip_runtime);
603 CC1Args.push_back(
"-idirafter");
605 if (UsesRuntimeWrapper)
606 CC1Args.append({
"-include",
"__clang_hip_runtime_wrapper.h"});
615 const char *LinkingOutput)
const {
616 std::string
Linker = getToolChain().GetLinkerPath();
617 ArgStringList CmdArgs;
618 if (!Args.hasArg(options::OPT_r)) {
619 CmdArgs.push_back(
"--no-undefined");
620 CmdArgs.push_back(
"-shared");
623 if (
C.getDriver().isUsingLTO()) {
624 const bool ThinLTO = (
C.getDriver().getLTOMode() ==
LTOK_Thin);
625 addLTOOptions(getToolChain(), Args, CmdArgs, Output, Inputs, ThinLTO);
626 }
else if (Args.hasArg(options::OPT_mcpu_EQ)) {
627 CmdArgs.push_back(Args.MakeArgString(
628 "-plugin-opt=mcpu=" +
630 Args.getLastArgValue(options::OPT_mcpu_EQ))));
633 getToolChain().AddFilePathLibArgs(Args, CmdArgs);
634 Args.AddAllArgs(CmdArgs, options::OPT_L);
638 std::vector<StringRef> Features;
641 if (!Features.empty()) {
643 Args.MakeArgString(
"-plugin-opt=-mattr=" + llvm::join(Features,
",")));
646 if (Args.hasArg(options::OPT_stdlib))
647 CmdArgs.append({
"-lc",
"-lm"});
648 if (Args.hasArg(options::OPT_startfiles)) {
649 std::optional<std::string> IncludePath = getToolChain().getStdlibPath();
651 IncludePath =
"/lib";
653 llvm::sys::path::append(
P,
"crt1.o");
654 CmdArgs.push_back(Args.MakeArgString(
P));
657 CmdArgs.push_back(
"-o");
659 C.addCommand(std::make_unique<Command>(
661 CmdArgs, Inputs, Output));
665 const llvm::Triple &Triple,
666 const llvm::opt::ArgList &Args,
667 std::vector<StringRef> &Features) {
671 if (Args.hasArg(options::OPT_mcpu_EQ))
672 TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
673 else if (Args.hasArg(options::OPT_march_EQ))
674 TargetID = Args.getLastArgValue(options::OPT_march_EQ);
675 if (!TargetID.empty()) {
676 llvm::StringMap<bool> FeatureMap;
677 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
678 if (OptionalGpuArch) {
679 StringRef GpuArch = *OptionalGpuArch;
685 auto Pos = FeatureMap.find(
Feature);
686 if (Pos == FeatureMap.end())
688 Features.push_back(Args.MakeArgStringRef(
689 (Twine(Pos->second ?
"+" :
"-") +
Feature).str()));
694 if (Args.hasFlag(options::OPT_mwavefrontsize64,
695 options::OPT_mno_wavefrontsize64,
false))
696 Features.push_back(
"+wavefrontsize64");
698 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
699 options::OPT_mno_amdgpu_precise_memory_op,
false))
700 Features.push_back(
"+precise-memory");
703 options::OPT_m_amdgpu_Features_Group);
711 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
727 DerivedArgList *DAL =
733 DAL =
new DerivedArgList(Args.getBaseArgs());
739 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
740 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
741 DAL->eraseArg(options::OPT_mcpu_EQ);
745 << llvm::Triple::getArchTypeName(
getArch())
746 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
748 auto &GPUs = *GPUsOrErr;
749 if (GPUs.size() > 1) {
751 << llvm::Triple::getArchTypeName(
getArch())
752 << llvm::join(GPUs,
", ") <<
"-mcpu";
754 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
755 Args.MakeArgString(GPUs.front()));
761 if (Args.getLastArgValue(options::OPT_x) !=
"cl")
765 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
766 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
768 : options::OPT_m32));
772 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
774 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
782 llvm::AMDGPU::GPUKind Kind) {
785 if (Kind == llvm::AMDGPU::GK_NONE)
788 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
792 const bool BothDenormAndFMAFast =
793 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
794 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
795 return !BothDenormAndFMAFast;
799 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
800 const llvm::fltSemantics *FPType)
const {
802 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
803 return llvm::DenormalMode::getIEEE();
808 auto Kind = llvm::AMDGPU::parseArchAMDGCN(
Arch);
809 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
810 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
811 options::OPT_fno_gpu_flush_denormals_to_zero,
813 return llvm::DenormalMode::getPreserveSign();
815 return llvm::DenormalMode::getIEEE();
818 const StringRef GpuArch =
getGPUArch(DriverArgs);
819 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
823 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
828 return DAZ ? llvm::DenormalMode::getPreserveSign() :
829 llvm::DenormalMode::getIEEE();
833 llvm::AMDGPU::GPUKind Kind) {
834 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
835 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
837 return !HasWave32 || DriverArgs.hasFlag(
838 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
850 const llvm::opt::ArgList &DriverArgs,
851 llvm::opt::ArgStringList &CC1Args,
855 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
856 options::OPT_fvisibility_ms_compat)) {
857 CC1Args.push_back(
"-fvisibility=hidden");
858 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
868 CC1Args.push_back(
"-Werror=atomic-alignment");
874 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
879 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
880 if (TargetID.empty())
881 return {std::nullopt, std::nullopt, std::nullopt};
883 llvm::StringMap<bool> FeatureMap;
885 if (!OptionalGpuArch)
886 return {TargetID.str(), std::nullopt, std::nullopt};
888 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
892 const llvm::opt::ArgList &DriverArgs)
const {
894 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
896 << *PTID.OptionalTargetID;
904 if (Arg *A = Args.getLastArg(options::OPT_offload_arch_tool_EQ))
905 Program = A->getValue();
911 return StdoutOrErr.takeError();
914 for (StringRef
Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
916 GPUArchs.push_back(
Arch.str());
918 if (GPUArchs.empty())
919 return llvm::createStringError(std::error_code(),
920 "No AMD GPU detected in the system");
922 return std::move(GPUArchs);
926 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
929 DeviceOffloadingKind);
934 DriverArgs.hasArg(options::OPT_nostdlib))
937 if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
942 const StringRef GpuArch =
getGPUArch(DriverArgs);
943 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
944 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
958 DriverArgs, LibDeviceFile, GpuArch, DeviceOffloadingKind,
961 for (
auto [BCFile, Internalize] : BCLibs) {
963 CC1Args.push_back(
"-mlink-builtin-bitcode");
965 CC1Args.push_back(
"-mlink-bitcode-file");
966 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
971 StringRef GPUArch, StringRef LibDeviceFile,
973 if (!hasDeviceLibrary()) {
974 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
977 if (LibDeviceFile.empty()) {
978 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
985 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString() << 0;
987 D.Diag(diag::err_drv_no_rocm_device_lib)
988 << 2 << ABIVer.
toString() << 1 <<
"6.3";
996 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
998 const bool NeedsASanRT)
const {
1001 CommonBitcodeLibsPreferences Pref{
D, DriverArgs, GPUArch,
1002 DeviceOffloadingKind, NeedsASanRT};
1005 bool Internalize =
true) {
1007 BCLibs.emplace_back(BCLib);
1009 auto AddSanBCLibs = [&]() {
1011 AddBCLib(getAsanRTLPath(),
false);
1015 AddBCLib(getOCMLPath());
1017 AddBCLib(getOCKLPath());
1018 else if (Pref.GPUSan && Pref.IsOpenMP)
1019 AddBCLib(getOCKLPath(),
false);
1020 AddBCLib(getDenormalsAreZeroPath(Pref.DAZ));
1021 AddBCLib(getUnsafeMathPath(Pref.UnsafeMathOpt || Pref.FastRelaxedMath));
1022 AddBCLib(getFiniteOnlyPath(Pref.FiniteOnly || Pref.FastRelaxedMath));
1023 AddBCLib(getCorrectlyRoundedSqrtPath(Pref.CorrectSqrt));
1024 AddBCLib(getWavefrontSize64Path(Pref.Wave64));
1025 AddBCLib(LibDeviceFile);
1026 auto ABIVerPath = getABIVersionPath(Pref.ABIVer);
1027 if (!ABIVerPath.empty())
1028 AddBCLib(ABIVerPath);
1035 const llvm::opt::ArgList &DriverArgs,
const std::string &GPUArch,
1037 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1038 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1048 DriverArgs, LibDeviceFile, GPUArch, DeviceOffloadingKind,
1053 const ToolChain &TC,
const llvm::opt::ArgList &DriverArgs,
1054 StringRef TargetID,
const llvm::opt::Arg *A)
const {
1056 if (TargetID.empty())
1058 Option O = A->getOption();
1060 if (!O.matches(options::OPT_fsanitize_EQ))
1063 if (!DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
1064 options::OPT_fno_gpu_sanitize,
true))
1071 if (K != SanitizerKind::Address)
1074 llvm::StringMap<bool> FeatureMap;
1077 assert(OptionalGpuArch &&
"Invalid Target ID");
1078 (void)OptionalGpuArch;
1079 auto Loc = FeatureMap.find(
"xnack");
1080 if (
Loc == FeatureMap.end() || !
Loc->second) {
1082 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1083 << 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 getIncludePath() const
Get the detected path to Rocm's bin directory.
RocmInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args, bool DetectHIPRuntime=true, bool DetectDeviceLib=false)
bool checkCommonBitcodeLibs(StringRef GPUArch, StringRef LibDeviceFile, DeviceLibABIVersion ABIVer) const
Check file paths of default bitcode libraries common to AMDGPU based toolchains.
bool hasHIPStdParLibrary() const
Check whether we detected a valid HIP STDPAR Acceleration library.
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()
void print(raw_ostream &OS) const
Print information about the detected ROCm installation.
The JSON file list parser is used to communicate input to InstallAPI.
std::optional< llvm::StringRef > parseTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch, llvm::StringMap< bool > *FeatureMap)
Parse a target ID to get processor and feature map.
llvm::StringRef getProcessorFromTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch)
Get processor name from target ID.
llvm::SmallVector< llvm::StringRef, 4 > getAllPossibleTargetIDFeatures(const llvm::Triple &T, llvm::StringRef Processor)
Get all feature strings that can be used in target ID for Processor.
SanitizerMask parseSanitizerValue(StringRef Value, bool AllowGroups)
Parse a single value from a -fsanitize= or -fno-sanitize= value list.
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()