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) {
172 SmallVector<StringRef, 4> VersionParts;
173 V.split(VersionParts,
'\n');
174 unsigned Major = ~0U;
175 unsigned Minor = ~0U;
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 == ~0U || Minor == ~0U)
189 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
191 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
197const SmallVectorImpl<RocmInstallationDetector::Candidate> &
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));
257 llvm::SmallString<256> RealClangPath;
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 = ~0U;
348 unsigned Minor = ~0U;
349 SmallVector<StringRef, 3> Parts;
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 != ~0U && Minor == ~0U)
361 if (Major == ~0U || Minor == ~0U)
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 {
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);
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))));
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,
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 = [&]() {
1018 else if (Pref.GPUSan && Pref.IsOpenMP)
1025 AddBCLib(LibDeviceFile);
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 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()