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);
86 GPUSan = (DriverArgs.hasFlag(options::OPT_fgpu_sanitize,
87 options::OPT_fno_gpu_sanitize,
true) &&
91void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) {
92 assert(!Path.empty());
94 const StringRef Suffix(
".bc");
95 const StringRef Suffix2(
".amdgcn.bc");
98 for (llvm::vfs::directory_iterator LI = D.getVFS().dir_begin(Path, EC), LE;
99 !EC && LI != LE; LI = LI.increment(EC)) {
100 StringRef FilePath = LI->path();
101 StringRef
FileName = llvm::sys::path::filename(FilePath);
107 BaseName =
FileName.drop_back(Suffix2.size());
108 else if (
FileName.ends_with(Suffix))
109 BaseName =
FileName.drop_back(Suffix.size());
111 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
112 if (BaseName ==
"ocml") {
114 }
else if (BaseName ==
"ockl") {
116 }
else if (BaseName ==
"opencl") {
118 }
else if (BaseName ==
"asanrtl") {
120 }
else if (BaseName ==
"oclc_finite_only_off") {
121 FiniteOnly.Off = FilePath;
122 }
else if (BaseName ==
"oclc_finite_only_on") {
123 FiniteOnly.On = FilePath;
124 }
else if (BaseName ==
"oclc_unsafe_math_on") {
125 UnsafeMath.On = FilePath;
126 }
else if (BaseName ==
"oclc_unsafe_math_off") {
127 UnsafeMath.Off = FilePath;
128 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
129 WavefrontSize64.On = FilePath;
130 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
131 WavefrontSize64.Off = FilePath;
132 }
else if (BaseName.starts_with(ABIVersionPrefix)) {
133 unsigned ABIVersionNumber;
134 if (BaseName.drop_front(ABIVersionPrefix.size())
135 .getAsInteger(0, ABIVersionNumber))
137 ABIVersionMap[ABIVersionNumber] = FilePath.str();
141 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
142 if (!BaseName.starts_with(DeviceLibPrefix))
145 StringRef IsaVersionNumber =
146 BaseName.drop_front(DeviceLibPrefix.size());
148 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
151 std::make_pair(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[RocmDeviceLibPathArg.size() - 1];
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,
",")));
635 if (Args.hasArg(options::OPT_stdlib))
636 CmdArgs.append({
"-lc",
"-lm"});
637 if (Args.hasArg(options::OPT_startfiles)) {
638 std::optional<std::string> IncludePath =
getToolChain().getStdlibPath();
640 IncludePath =
"/lib";
642 llvm::sys::path::append(P,
"crt1.o");
643 CmdArgs.push_back(Args.MakeArgString(P));
646 CmdArgs.push_back(
"-o");
648 C.addCommand(std::make_unique<Command>(
650 CmdArgs, Inputs, Output));
654 const llvm::Triple &Triple,
655 const llvm::opt::ArgList &Args,
656 std::vector<StringRef> &Features) {
660 if (Args.hasArg(options::OPT_mcpu_EQ))
661 TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
662 else if (Args.hasArg(options::OPT_march_EQ))
663 TargetID = Args.getLastArgValue(options::OPT_march_EQ);
664 if (!TargetID.empty()) {
665 llvm::StringMap<bool> FeatureMap;
666 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
667 if (OptionalGpuArch) {
668 StringRef GpuArch = *OptionalGpuArch;
674 auto Pos = FeatureMap.find(
Feature);
675 if (Pos == FeatureMap.end())
677 Features.push_back(Args.MakeArgStringRef(
678 (Twine(Pos->second ?
"+" :
"-") +
Feature).str()));
683 if (Args.hasFlag(options::OPT_mwavefrontsize64,
684 options::OPT_mno_wavefrontsize64,
false))
685 Features.push_back(
"+wavefrontsize64");
687 if (Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,
688 options::OPT_mno_amdgpu_precise_memory_op,
false))
689 Features.push_back(
"+precise-memory");
692 options::OPT_m_amdgpu_Features_Group);
700 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
716 DerivedArgList *DAL =
722 DAL =
new DerivedArgList(Args.getBaseArgs());
728 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
729 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
730 DAL->eraseArg(options::OPT_mcpu_EQ);
734 << llvm::Triple::getArchTypeName(
getArch())
735 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
737 auto &GPUs = *GPUsOrErr;
738 if (GPUs.size() > 1) {
740 << llvm::Triple::getArchTypeName(
getArch())
741 << llvm::join(GPUs,
", ") <<
"-mcpu";
743 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
744 Args.MakeArgString(GPUs.front()));
750 if (Args.getLastArgValue(options::OPT_x) !=
"cl")
754 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
755 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
757 : options::OPT_m32));
761 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
763 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
771 llvm::AMDGPU::GPUKind Kind) {
774 if (Kind == llvm::AMDGPU::GK_NONE)
777 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
781 const bool BothDenormAndFMAFast =
782 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
783 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
784 return !BothDenormAndFMAFast;
788 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
789 const llvm::fltSemantics *FPType)
const {
791 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
792 return llvm::DenormalMode::getIEEE();
797 auto Kind = llvm::AMDGPU::parseArchAMDGCN(
Arch);
798 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
799 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
800 options::OPT_fno_gpu_flush_denormals_to_zero,
802 return llvm::DenormalMode::getPreserveSign();
804 return llvm::DenormalMode::getIEEE();
807 const StringRef GpuArch =
getGPUArch(DriverArgs);
808 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
812 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
817 return DAZ ? llvm::DenormalMode::getPreserveSign() :
818 llvm::DenormalMode::getIEEE();
822 llvm::AMDGPU::GPUKind Kind) {
823 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
824 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
826 return !HasWave32 || DriverArgs.hasFlag(
827 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
835 if (Triple.getEnvironment() != llvm::Triple::LLVM)
840 const llvm::opt::ArgList &DriverArgs,
841 llvm::opt::ArgStringList &CC1Args,
845 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
846 options::OPT_fvisibility_ms_compat)) {
847 CC1Args.push_back(
"-fvisibility=hidden");
848 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
862 !DriverArgs.hasArg(options::OPT_disable_llvm_optzns))
863 CC1Args.push_back(
"-disable-llvm-optzns");
872 CC1Args.push_back(
"-Werror=atomic-alignment");
876 ArgStringList &CC1Args)
const {
877 if (DriverArgs.hasArg(options::OPT_nostdinc) ||
878 DriverArgs.hasArg(options::OPT_nostdlibinc))
888 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
893 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
894 if (TargetID.empty())
895 return {std::nullopt, std::nullopt, std::nullopt};
897 llvm::StringMap<bool> FeatureMap;
899 if (!OptionalGpuArch)
900 return {TargetID.str(), std::nullopt, std::nullopt};
902 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
906 const llvm::opt::ArgList &DriverArgs)
const {
908 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
910 << *PTID.OptionalTargetID;
918 if (Arg *A = Args.getLastArg(options::OPT_offload_arch_tool_EQ))
919 Program = A->getValue();
925 return StdoutOrErr.takeError();
928 for (StringRef
Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
930 GPUArchs.push_back(
Arch.str());
932 if (GPUArchs.empty())
933 return llvm::createStringError(std::error_code(),
934 "No AMD GPU detected in the system");
936 return std::move(GPUArchs);
940 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
943 DeviceOffloadingKind);
948 DriverArgs.hasArg(options::OPT_nostdlib))
951 if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
963 if (TT.getEnvironment() == llvm::Triple::LLVM)
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) {
1031 if (!BCLib.
Path.empty()) {
1033 BCLibs.emplace_back(BCLib);
1036 auto AddSanBCLibs = [&]() {
1045 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 bool IsExplicitDevice =
1082 A->getBaseArg().getOption().matches(options::OPT_Xarch_device);
1085 llvm::StringRef Processor =
1087 auto ProcKind = TC.
getTriple().isAMDGCN()
1088 ? llvm::AMDGPU::parseArchAMDGCN(Processor)
1089 : llvm::AMDGPU::parseArchR600(Processor);
1090 auto Features = TC.
getTriple().isAMDGCN()
1091 ? llvm::AMDGPU::getArchAttrAMDGCN(ProcKind)
1092 : llvm::AMDGPU::getArchAttrR600(ProcKind);
1093 if (Features & llvm::AMDGPU::FEATURE_XNACK_ALWAYS)
1097 llvm::StringMap<bool> FeatureMap;
1099 assert(OptionalGpuArch &&
"Invalid Target ID");
1100 (void)OptionalGpuArch;
1101 auto Loc = FeatureMap.find(
"xnack");
1102 if (Loc == FeatureMap.end() || !Loc->second) {
1103 if (IsExplicitDevice) {
1105 clang::diag::err_drv_unsupported_option_for_offload_arch_req_feature)
1106 << A->getAsString(DriverArgs) << TargetID <<
"xnack+";
1109 clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)
1110 << 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.
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()