11#include "clang/Config/config.h"
17#include "llvm/ADT/StringExtras.h"
18#include "llvm/Option/ArgList.h"
19#include "llvm/Support/Error.h"
20#include "llvm/Support/LineIterator.h"
21#include "llvm/Support/Path.h"
22#include "llvm/Support/Process.h"
23#include "llvm/Support/VirtualFileSystem.h"
24#include "llvm/TargetParser/Host.h"
25#include "llvm/TargetParser/TargetParser.h"
27#include <system_error>
35RocmInstallationDetector::CommonBitcodeLibsPreferences::
36 CommonBitcodeLibsPreferences(
const Driver &D,
37 const llvm::opt::ArgList &DriverArgs,
40 const bool NeedsASanRT)
43 const auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
44 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
48 const bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
50 !HasWave32 || DriverArgs.hasFlag(options::OPT_mwavefrontsize64,
51 options::OPT_mno_wavefrontsize64,
false);
58 const bool DefaultDAZ =
59 (Kind == llvm::AMDGPU::GK_NONE)
61 : !((ArchAttr &
llvm::
AMDGPU::FEATURE_FAST_FMA_F32) &&
62 (ArchAttr &
llvm::
AMDGPU::FEATURE_FAST_DENORMAL_F32));
65 DAZ = IsKnownOffloading
66 ? DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
67 options::OPT_fno_gpu_flush_denormals_to_zero,
69 : DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) || DefaultDAZ;
71 FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only) ||
72 DriverArgs.hasFlag(options::OPT_ffinite_math_only,
73 options::OPT_fno_finite_math_only,
false);
76 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations) ||
77 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
78 options::OPT_fno_unsafe_math_optimizations,
false);
80 FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math) ||
81 DriverArgs.hasFlag(options::OPT_ffast_math,
82 options::OPT_fno_fast_math,
false);
84 const bool DefaultSqrt = IsKnownOffloading ?
true :
false;
86 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt) ||
88 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
89 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, DefaultSqrt);
92 GPUSan = (DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
93 options::OPT_fno_gpu_sanitize,
true) &&
97void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) {
98 assert(!Path.empty());
100 const StringRef Suffix(
".bc");
101 const StringRef Suffix2(
".amdgcn.bc");
104 for (llvm::vfs::directory_iterator LI = D.getVFS().dir_begin(Path, EC), LE;
105 !EC && LI != LE; LI = LI.increment(EC)) {
106 StringRef FilePath = LI->path();
107 StringRef
FileName = llvm::sys::path::filename(FilePath);
113 BaseName =
FileName.drop_back(Suffix2.size());
114 else if (
FileName.ends_with(Suffix))
115 BaseName =
FileName.drop_back(Suffix.size());
117 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
118 if (BaseName ==
"ocml") {
120 }
else if (BaseName ==
"ockl") {
122 }
else if (BaseName ==
"opencl") {
124 }
else if (BaseName ==
"asanrtl") {
126 }
else if (BaseName ==
"oclc_finite_only_off") {
127 FiniteOnly.Off = FilePath;
128 }
else if (BaseName ==
"oclc_finite_only_on") {
129 FiniteOnly.On = FilePath;
130 }
else if (BaseName ==
"oclc_daz_opt_on") {
131 DenormalsAreZero.On = FilePath;
132 }
else if (BaseName ==
"oclc_daz_opt_off") {
133 DenormalsAreZero.Off = FilePath;
134 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_on") {
135 CorrectlyRoundedSqrt.On = FilePath;
136 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_off") {
137 CorrectlyRoundedSqrt.Off = FilePath;
138 }
else if (BaseName ==
"oclc_unsafe_math_on") {
139 UnsafeMath.On = FilePath;
140 }
else if (BaseName ==
"oclc_unsafe_math_off") {
141 UnsafeMath.Off = FilePath;
142 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
143 WavefrontSize64.On = FilePath;
144 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
145 WavefrontSize64.Off = FilePath;
146 }
else if (BaseName.starts_with(ABIVersionPrefix)) {
147 unsigned ABIVersionNumber;
148 if (BaseName.drop_front(ABIVersionPrefix.size())
149 .getAsInteger(0, ABIVersionNumber))
151 ABIVersionMap[ABIVersionNumber] = FilePath.str();
155 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
156 if (!BaseName.starts_with(DeviceLibPrefix))
159 StringRef IsaVersionNumber =
160 BaseName.drop_front(DeviceLibPrefix.size());
162 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
165 std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
172bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
173 SmallVector<StringRef, 4> VersionParts;
174 V.split(VersionParts,
'\n');
175 unsigned Major = ~0U;
176 unsigned Minor = ~0U;
177 for (
auto Part : VersionParts) {
178 auto Splits = Part.rtrim().split(
'=');
179 if (Splits.first ==
"HIP_VERSION_MAJOR") {
180 if (Splits.second.getAsInteger(0, Major))
182 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
183 if (Splits.second.getAsInteger(0, Minor))
185 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
186 VersionPatch = Splits.second.str();
188 if (Major == ~0U || Minor == ~0U)
190 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
192 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
198const SmallVectorImpl<RocmInstallationDetector::Candidate> &
199RocmInstallationDetector::getInstallationPathCandidates() {
202 if (!ROCmSearchDirs.empty())
203 return ROCmSearchDirs;
205 auto DoPrintROCmSearchDirs = [&]() {
206 if (PrintROCmSearchDirs)
207 for (
auto Cand : ROCmSearchDirs) {
208 llvm::errs() <<
"ROCm installation search path: " << Cand.Path <<
'\n';
214 if (!RocmPathArg.empty()) {
215 ROCmSearchDirs.emplace_back(RocmPathArg.str());
216 DoPrintROCmSearchDirs();
217 return ROCmSearchDirs;
218 }
else if (std::optional<std::string> RocmPathEnv =
219 llvm::sys::Process::GetEnv(
"ROCM_PATH")) {
220 if (!RocmPathEnv->empty()) {
221 ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));
222 DoPrintROCmSearchDirs();
223 return ROCmSearchDirs;
228 StringRef InstallDir = D.Dir;
233 auto DeduceROCmPath = [](StringRef ClangPath) {
235 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
236 StringRef ParentName = llvm::sys::path::filename(ParentDir);
239 if (ParentName ==
"bin") {
240 ParentDir = llvm::sys::path::parent_path(ParentDir);
241 ParentName = llvm::sys::path::filename(ParentDir);
246 if (ParentName ==
"llvm" || ParentName.starts_with(
"aomp"))
247 ParentDir = llvm::sys::path::parent_path(ParentDir);
249 return Candidate(ParentDir.str(),
true);
254 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
258 llvm::SmallString<256> RealClangPath;
259 llvm::sys::fs::real_path(D.getClangProgramPath(), RealClangPath);
260 auto ParentPath = llvm::sys::path::parent_path(RealClangPath);
261 if (ParentPath != InstallDir)
262 ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));
265 auto ClangRoot = llvm::sys::path::parent_path(InstallDir);
266 auto RealClangRoot = llvm::sys::path::parent_path(ParentPath);
267 ROCmSearchDirs.emplace_back(ClangRoot.str(),
true);
268 if (RealClangRoot != ClangRoot)
269 ROCmSearchDirs.emplace_back(RealClangRoot.str(),
true);
270 ROCmSearchDirs.emplace_back(D.ResourceDir,
273 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/rocm",
278 std::string LatestROCm;
279 llvm::VersionTuple LatestVer;
281 auto GetROCmVersion = [](StringRef DirName) {
282 llvm::VersionTuple
V;
283 std::string VerStr = DirName.drop_front(strlen(
"rocm-")).str();
286 llvm::replace(VerStr,
'-',
'.');
290 for (llvm::vfs::directory_iterator
291 File = D.getVFS().dir_begin(D.SysRoot +
"/opt", EC),
293 File != FileEnd && !EC;
File.increment(EC)) {
294 llvm::StringRef
FileName = llvm::sys::path::filename(
File->path());
297 if (LatestROCm.empty()) {
299 LatestVer = GetROCmVersion(LatestROCm);
302 auto Ver = GetROCmVersion(
FileName);
303 if (LatestVer < Ver) {
308 if (!LatestROCm.empty())
309 ROCmSearchDirs.emplace_back(D.SysRoot +
"/opt/" + LatestROCm,
312 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr/local",
314 ROCmSearchDirs.emplace_back(D.SysRoot +
"/usr",
317 DoPrintROCmSearchDirs();
318 return ROCmSearchDirs;
322 const Driver &D,
const llvm::Triple &HostTriple,
323 const llvm::opt::ArgList &Args,
bool DetectHIPRuntime,
bool DetectDeviceLib)
325 Verbose = Args.hasArg(options::OPT_v);
326 RocmPathArg = Args.getLastArgValue(clang::driver::options::OPT_rocm_path_EQ);
327 PrintROCmSearchDirs =
328 Args.hasArg(clang::driver::options::OPT_print_rocm_search_dirs);
329 RocmDeviceLibPathArg =
330 Args.getAllArgValues(clang::driver::options::OPT_rocm_device_lib_path_EQ);
331 HIPPathArg = Args.getLastArgValue(clang::driver::options::OPT_hip_path_EQ);
333 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_path_EQ);
334 HasHIPStdParLibrary =
335 !HIPStdParPathArg.empty() && D.getVFS().exists(HIPStdParPathArg +
336 "/hipstdpar_lib.hpp");
337 HIPRocThrustPathArg =
338 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_thrust_path_EQ);
339 HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&
340 D.getVFS().exists(HIPRocThrustPathArg +
"/thrust");
342 Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_prim_path_EQ);
343 HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&
344 D.getVFS().exists(HIPRocPrimPathArg +
"/rocprim");
346 if (
auto *A = Args.getLastArg(clang::driver::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,
967 const StringRef GpuArch =
getGPUArch(DriverArgs);
968 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
969 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
983 DriverArgs, LibDeviceFile, GpuArch, DeviceOffloadingKind,
986 for (
auto [BCFile, Internalize] : BCLibs) {
988 CC1Args.push_back(
"-mlink-builtin-bitcode");
990 CC1Args.push_back(
"-mlink-bitcode-file");
991 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
996 StringRef GPUArch, StringRef LibDeviceFile,
999 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
1002 if (LibDeviceFile.empty()) {
1003 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
1010 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString() << 0;
1012 D.Diag(diag::err_drv_no_rocm_device_lib)
1013 << 2 << ABIVer.
toString() << 1 <<
"6.3";
1021 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
1023 const bool NeedsASanRT)
const {
1026 CommonBitcodeLibsPreferences Pref{D, DriverArgs, GPUArch,
1027 DeviceOffloadingKind, NeedsASanRT};
1030 bool Internalize =
true) {
1032 BCLibs.emplace_back(BCLib);
1034 auto AddSanBCLibs = [&]() {
1043 else if (Pref.GPUSan && Pref.IsOpenMP)
1050 AddBCLib(LibDeviceFile);
1052 if (!ABIVerPath.empty())
1053 AddBCLib(ABIVerPath);
1060 const llvm::opt::ArgList &DriverArgs,
const std::string &GPUArch,
1062 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
1063 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
1073 DriverArgs, LibDeviceFile, GPUArch, DeviceOffloadingKind,
1078 const ToolChain &TC,
const llvm::opt::ArgList &DriverArgs,
1079 StringRef TargetID,
const llvm::opt::Arg *A)
const {
1081 if (TargetID.empty())
1083 Option O = A->getOption();
1085 if (!O.matches(options::OPT_fsanitize_EQ))
1088 if (!DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
1089 options::OPT_fno_gpu_sanitize,
true))
1096 if (K != SanitizerKind::Address)
1100 llvm::StringRef Processor =
1102 auto ProcKind = TC.
getTriple().isAMDGCN()
1103 ? llvm::AMDGPU::parseArchAMDGCN(Processor)
1104 : llvm::AMDGPU::parseArchR600(Processor);
1105 auto Features = TC.
getTriple().isAMDGCN()
1106 ? llvm::AMDGPU::getArchAttrAMDGCN(ProcKind)
1107 : llvm::AMDGPU::getArchAttrR600(ProcKind);
1108 if (Features & llvm::AMDGPU::FEATURE_XNACK_ALWAYS)
1112 llvm::StringMap<bool> FeatureMap;
1114 assert(OptionalGpuArch &&
"Invalid Target ID");
1115 (void)OptionalGpuArch;
1116 auto Loc = FeatureMap.find(
"xnack");
1117 if (Loc == FeatureMap.end() || !Loc->second) {
1119 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1120 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
static void Append(char *Start, char *End, char *&Buffer, unsigned &BufferSize, unsigned &BufferCapacity)
__device__ __2f16 float c
const char * getOffloadingArch() const
OffloadKind getOffloadingDeviceKind() const
Compilation - A set of tasks to perform for a single driver invocation.
Driver - Encapsulate logic for constructing compilation processes from a set of gcc-driver-like comma...
DiagnosticsEngine & getDiags() const
llvm::Expected< std::unique_ptr< llvm::MemoryBuffer > > executeProgram(llvm::ArrayRef< llvm::StringRef > Args) const
DiagnosticBuilder Diag(unsigned DiagID) const
const llvm::opt::OptTable & getOpts() const
StringRef getCorrectlyRoundedSqrtPath(bool Enabled) const
StringRef getIncludePath() const
Get the detected path to Rocm's bin directory.
StringRef getUnsafeMathPath(bool Enabled) const
StringRef getOCMLPath() const
StringRef getAsanRTLPath() const
Returns empty string of Asan runtime library is not available.
RocmInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args, bool DetectHIPRuntime=true, bool DetectDeviceLib=false)
StringRef getOCKLPath() const
StringRef getFiniteOnlyPath(bool Enabled) const
bool hasDeviceLibrary() const
Check whether we detected a valid ROCm device library.
StringRef getDenormalsAreZeroPath(bool Enabled) const
bool checkCommonBitcodeLibs(StringRef GPUArch, StringRef LibDeviceFile, DeviceLibABIVersion ABIVer) const
Check file paths of default bitcode libraries common to AMDGPU based toolchains.
bool hasHIPStdParLibrary() const
Check whether we detected a valid HIP STDPAR Acceleration library.
StringRef getABIVersionPath(DeviceLibABIVersion ABIVer) const
llvm::SmallVector< ToolChain::BitCodeLibraryInfo, 12 > getCommonBitcodeLibs(const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, StringRef GPUArch, const Action::OffloadKind DeviceOffloadingKind, const bool NeedsASanRT) const
Get file paths of default bitcode libraries common to AMDGPU based toolchains.
bool hasHIPRuntime() const
Check whether we detected a valid HIP runtime.
void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const
void detectDeviceLibrary()
StringRef getWavefrontSize64Path(bool Enabled) const
void print(raw_ostream &OS) const
Print information about the detected ROCm installation.
SmallVector< InputInfo, 4 > InputInfoList
The JSON file list parser is used to communicate input to InstallAPI.
if(T->getSizeExpr()) TRY_TO(TraverseStmt(const_cast< Expr * >(T -> getSizeExpr())))
std::optional< llvm::StringRef > parseTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch, llvm::StringMap< bool > *FeatureMap)
Parse a target ID to get processor and feature map.
llvm::StringRef getProcessorFromTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch)
Get processor name from target ID.
llvm::SmallVector< llvm::StringRef, 4 > getAllPossibleTargetIDFeatures(const llvm::Triple &T, llvm::StringRef Processor)
Get all feature strings that can be used in target ID for Processor.
SanitizerMask parseSanitizerValue(StringRef Value, bool AllowGroups)
Parse a single value from a -fsanitize= or -fno-sanitize= value list.
Diagnostic wrappers for TextAPI types for error reporting.
ABI version of device library.
unsigned getAsCodeObjectVersion() const
static DeviceLibABIVersion fromCodeObjectVersion(unsigned CodeObjectVersion)
bool requiresLibrary()
Whether ABI version bc file is requested.
static constexpr ResponseFileSupport AtFileCurCP()