12#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>
41RocmInstallationDetector::findSPACKPackage(
const Candidate &Cand,
42 StringRef PackageName) {
46 std::string Prefix = Twine(PackageName +
"-" + Cand.SPACKReleaseStr).str();
48 for (llvm::vfs::directory_iterator
File = D.
getVFS().dir_begin(Cand.Path, EC),
50 File != FileEnd && !EC;
File.increment(EC)) {
51 llvm::StringRef FileName = llvm::sys::path::filename(
File->path());
52 if (FileName.startswith(Prefix)) {
53 SubDirs.push_back(FileName);
54 if (SubDirs.size() > 1)
58 if (SubDirs.size() == 1) {
59 auto PackagePath = Cand.Path;
60 llvm::sys::path::append(PackagePath, SubDirs[0]);
63 if (SubDirs.size() == 0 && Verbose) {
64 llvm::errs() <<
"SPACK package " << Prefix <<
" not found at " << Cand.Path
69 if (SubDirs.size() > 1 && Verbose) {
70 llvm::errs() <<
"Cannot use SPACK package " << Prefix <<
" at " << Cand.Path
71 <<
" due to multiple installations for the same version\n";
76void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) {
77 assert(!Path.empty());
79 const StringRef Suffix(
".bc");
80 const StringRef Suffix2(
".amdgcn.bc");
83 for (llvm::vfs::directory_iterator LI = D.
getVFS().dir_begin(Path, EC), LE;
84 !EC && LI != LE; LI = LI.increment(EC)) {
85 StringRef FilePath = LI->path();
86 StringRef FileName = llvm::sys::path::filename(FilePath);
87 if (!FileName.endswith(Suffix))
91 if (FileName.endswith(Suffix2))
92 BaseName = FileName.drop_back(Suffix2.size());
93 else if (FileName.endswith(Suffix))
94 BaseName = FileName.drop_back(Suffix.size());
96 const StringRef ABIVersionPrefix =
"oclc_abi_version_";
97 if (BaseName ==
"ocml") {
99 }
else if (BaseName ==
"ockl") {
101 }
else if (BaseName ==
"opencl") {
103 }
else if (BaseName ==
"hip") {
105 }
else if (BaseName ==
"asanrtl") {
107 }
else if (BaseName ==
"oclc_finite_only_off") {
108 FiniteOnly.Off = FilePath;
109 }
else if (BaseName ==
"oclc_finite_only_on") {
110 FiniteOnly.On = FilePath;
111 }
else if (BaseName ==
"oclc_daz_opt_on") {
112 DenormalsAreZero.On = FilePath;
113 }
else if (BaseName ==
"oclc_daz_opt_off") {
114 DenormalsAreZero.Off = FilePath;
115 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_on") {
116 CorrectlyRoundedSqrt.On = FilePath;
117 }
else if (BaseName ==
"oclc_correctly_rounded_sqrt_off") {
118 CorrectlyRoundedSqrt.Off = FilePath;
119 }
else if (BaseName ==
"oclc_unsafe_math_on") {
120 UnsafeMath.On = FilePath;
121 }
else if (BaseName ==
"oclc_unsafe_math_off") {
122 UnsafeMath.Off = FilePath;
123 }
else if (BaseName ==
"oclc_wavefrontsize64_on") {
124 WavefrontSize64.On = FilePath;
125 }
else if (BaseName ==
"oclc_wavefrontsize64_off") {
126 WavefrontSize64.Off = FilePath;
127 }
else if (BaseName.startswith(ABIVersionPrefix)) {
128 unsigned ABIVersionNumber;
129 if (BaseName.drop_front(ABIVersionPrefix.size())
130 .getAsInteger(0, ABIVersionNumber))
132 ABIVersionMap[ABIVersionNumber] = FilePath.str();
136 const StringRef DeviceLibPrefix =
"oclc_isa_version_";
137 if (!BaseName.startswith(DeviceLibPrefix))
140 StringRef IsaVersionNumber =
141 BaseName.drop_front(DeviceLibPrefix.size());
143 llvm::Twine GfxName = Twine(
"gfx") + IsaVersionNumber;
146 std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
153bool RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef
V) {
155 V.split(VersionParts,
'\n');
156 unsigned Major = ~0
U;
157 unsigned Minor = ~0
U;
158 for (
auto Part : VersionParts) {
159 auto Splits = Part.rtrim().split(
'=');
160 if (Splits.first ==
"HIP_VERSION_MAJOR") {
161 if (Splits.second.getAsInteger(0, Major))
163 }
else if (Splits.first ==
"HIP_VERSION_MINOR") {
164 if (Splits.second.getAsInteger(0, Minor))
166 }
else if (Splits.first ==
"HIP_VERSION_PATCH")
167 VersionPatch = Splits.second.str();
169 if (Major == ~0
U || Minor == ~0
U)
171 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
173 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
180RocmInstallationDetector::getInstallationPathCandidates() {
183 if (!ROCmSearchDirs.empty())
184 return ROCmSearchDirs;
186 auto DoPrintROCmSearchDirs = [&]() {
187 if (PrintROCmSearchDirs)
188 for (
auto Cand : ROCmSearchDirs) {
189 llvm::errs() <<
"ROCm installation search path";
191 llvm::errs() <<
" (Spack " << Cand.SPACKReleaseStr <<
")";
192 llvm::errs() <<
": " << Cand.Path <<
'\n';
198 if (!RocmPathArg.empty()) {
199 ROCmSearchDirs.emplace_back(RocmPathArg.str());
200 DoPrintROCmSearchDirs();
201 return ROCmSearchDirs;
202 }
else if (std::optional<std::string> RocmPathEnv =
203 llvm::sys::Process::GetEnv(
"ROCM_PATH")) {
204 if (!RocmPathEnv->empty()) {
205 ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));
206 DoPrintROCmSearchDirs();
207 return ROCmSearchDirs;
217 auto DeduceROCmPath = [](StringRef ClangPath) {
219 StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);
220 StringRef ParentName = llvm::sys::path::filename(ParentDir);
223 if (ParentName ==
"bin") {
224 ParentDir = llvm::sys::path::parent_path(ParentDir);
225 ParentName = llvm::sys::path::filename(ParentDir);
233 if (ParentName.startswith(
"llvm-amdgpu-")) {
235 ParentName.drop_front(strlen(
"llvm-amdgpu-")).split(
'-');
236 auto SPACKReleaseStr = SPACKPostfix.first;
237 if (!SPACKReleaseStr.empty()) {
238 ParentDir = llvm::sys::path::parent_path(ParentDir);
239 return Candidate(ParentDir.str(),
true,
246 if (ParentName ==
"llvm" || ParentName.startswith(
"aomp"))
247 ParentDir = llvm::sys::path::parent_path(ParentDir);
249 return Candidate(ParentDir.str(),
true);
254 ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));
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);
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 std::replace(VerStr.begin(), VerStr.end(),
'-',
'.');
290 for (llvm::vfs::directory_iterator
293 File != FileEnd && !EC;
File.increment(EC)) {
294 llvm::StringRef FileName = llvm::sys::path::filename(
File->path());
295 if (!FileName.startswith(
"rocm-"))
297 if (LatestROCm.empty()) {
298 LatestROCm = FileName.str();
299 LatestVer = GetROCmVersion(LatestROCm);
302 auto Ver = GetROCmVersion(FileName);
303 if (LatestVer < Ver) {
304 LatestROCm = FileName.str();
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);
332 if (
auto *A = Args.getLastArg(clang::driver::options::OPT_hip_version_EQ)) {
333 HIPVersionArg = A->getValue();
334 unsigned Major = ~0
U;
335 unsigned Minor = ~0
U;
337 HIPVersionArg.split(Parts,
'.');
339 Parts[0].getAsInteger(0, Major);
340 if (Parts.size() > 1)
341 Parts[1].getAsInteger(0, Minor);
342 if (Parts.size() > 2)
343 VersionPatch = Parts[2].str();
344 if (VersionPatch.empty())
346 if (Major != ~0
U && Minor == ~0
U)
348 if (Major == ~0
U || Minor == ~0
U)
349 D.
Diag(diag::err_drv_invalid_value)
350 << A->getAsString(Args) << HIPVersionArg;
352 VersionMajorMinor = llvm::VersionTuple(Major, Minor);
354 (Twine(Major) +
"." + Twine(Minor) +
"." + VersionPatch).str();
356 VersionPatch = DefaultVersionPatch;
358 llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
359 DetectedVersion = (Twine(DefaultVersionMajor) +
"." +
360 Twine(DefaultVersionMinor) +
"." + VersionPatch)
364 if (DetectHIPRuntime)
371 assert(LibDevicePath.empty());
373 if (!RocmDeviceLibPathArg.empty())
374 LibDevicePath = RocmDeviceLibPathArg[RocmDeviceLibPathArg.size() - 1];
375 else if (std::optional<std::string> LibPathEnv =
376 llvm::sys::Process::GetEnv(
"HIP_DEVICE_LIB_PATH"))
377 LibDevicePath = std::move(*LibPathEnv);
380 if (!LibDevicePath.empty()) {
384 if (!FS.exists(LibDevicePath))
387 scanLibDevicePath(LibDevicePath);
388 HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
393 auto CheckDeviceLib = [&](StringRef Path,
bool StrictChecking) {
394 bool CheckLibDevice = (!NoBuiltinLibs || StrictChecking);
395 if (CheckLibDevice && !FS.exists(Path))
398 scanLibDevicePath(Path);
400 if (!NoBuiltinLibs) {
402 if (!allGenericLibsValid())
407 if (LibDeviceMap.empty())
415 llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,
416 "amdgcn",
"bitcode");
417 HasDeviceLibrary = CheckDeviceLib(LibDevicePath,
true);
418 if (HasDeviceLibrary)
423 auto &ROCmDirs = getInstallationPathCandidates();
424 for (
const auto &Candidate : ROCmDirs) {
425 LibDevicePath = Candidate.Path;
426 llvm::sys::path::append(LibDevicePath,
"amdgcn",
"bitcode");
427 HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);
428 if (HasDeviceLibrary)
435 if (!HIPPathArg.empty())
436 HIPSearchDirs.emplace_back(HIPPathArg.str());
437 else if (std::optional<std::string> HIPPathEnv =
438 llvm::sys::Process::GetEnv(
"HIP_PATH")) {
439 if (!HIPPathEnv->empty())
440 HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));
442 if (HIPSearchDirs.empty())
443 HIPSearchDirs.append(getInstallationPathCandidates());
446 for (
const auto &Candidate : HIPSearchDirs) {
447 InstallPath = Candidate.Path;
448 if (InstallPath.empty() || !FS.exists(InstallPath))
452 auto SPACKPath = findSPACKPackage(Candidate,
"hip");
453 InstallPath = SPACKPath.empty() ? InstallPath : SPACKPath;
455 BinPath = InstallPath;
456 llvm::sys::path::append(BinPath,
"bin");
457 IncludePath = InstallPath;
458 llvm::sys::path::append(IncludePath,
"include");
459 LibPath = InstallPath;
460 llvm::sys::path::append(LibPath,
"lib");
461 SharePath = InstallPath;
462 llvm::sys::path::append(SharePath,
"share");
465 SmallString<0> ParentSharePath = llvm::sys::path::parent_path(InstallPath);
466 llvm::sys::path::append(ParentSharePath,
"share");
469 const Twine &
c =
"",
const Twine &d =
"") {
471 llvm::sys::path::append(newpath, a,
b,
c, d);
475 for (
const auto &VersionFilePath :
476 {
Append(SharePath,
"hip",
"version"),
477 Append(ParentSharePath,
"hip",
"version"),
478 Append(BinPath,
".hipVersion")}) {
479 llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
480 FS.getBufferForFile(VersionFilePath);
483 if (HIPVersionArg.empty() && VersionFile)
484 if (parseHIPVersionFile((*VersionFile)->getBuffer()))
487 HasHIPRuntime =
true;
492 if (!Candidate.StrictChecking) {
493 HasHIPRuntime =
true;
497 HasHIPRuntime =
false;
502 OS <<
"Found HIP installation: " << InstallPath <<
", version "
503 << DetectedVersion <<
'\n';
507 ArgStringList &CC1Args)
const {
508 bool UsesRuntimeWrapper = VersionMajorMinor > llvm::VersionTuple(3, 5) &&
509 !DriverArgs.hasArg(options::OPT_nohipwrapperinc);
511 if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
526 if (UsesRuntimeWrapper)
527 llvm::sys::path::append(
P,
"include",
"cuda_wrappers");
528 CC1Args.push_back(
"-internal-isystem");
529 CC1Args.push_back(DriverArgs.MakeArgString(
P));
532 if (DriverArgs.hasArg(options::OPT_nogpuinc))
536 D.
Diag(diag::err_drv_no_hip_runtime);
540 CC1Args.push_back(
"-idirafter");
542 if (UsesRuntimeWrapper)
543 CC1Args.append({
"-include",
"__clang_hip_runtime_wrapper.h"});
550 const char *LinkingOutput)
const {
552 std::string
Linker = getToolChain().GetProgramPath(getShortName());
553 ArgStringList CmdArgs;
554 CmdArgs.push_back(
"--no-undefined");
555 CmdArgs.push_back(
"-shared");
558 Args.AddAllArgs(CmdArgs, options::OPT_L);
560 if (
C.getDriver().isUsingLTO())
561 addLTOOptions(getToolChain(), Args, CmdArgs, Output, Inputs[0],
563 else if (Args.hasArg(options::OPT_mcpu_EQ))
564 CmdArgs.push_back(Args.MakeArgString(
565 "-plugin-opt=mcpu=" + Args.getLastArgValue(options::OPT_mcpu_EQ)));
566 CmdArgs.push_back(
"-o");
568 C.addCommand(std::make_unique<Command>(
570 CmdArgs, Inputs, Output));
574 const llvm::Triple &Triple,
575 const llvm::opt::ArgList &Args,
576 std::vector<StringRef> &Features) {
579 StringRef TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);
580 if (!TargetID.empty()) {
581 llvm::StringMap<bool> FeatureMap;
582 auto OptionalGpuArch =
parseTargetID(Triple, TargetID, &FeatureMap);
583 if (OptionalGpuArch) {
584 StringRef GpuArch = *OptionalGpuArch;
590 auto Pos = FeatureMap.find(Feature);
591 if (Pos == FeatureMap.end())
593 Features.push_back(Args.MakeArgStringRef(
594 (Twine(Pos->second ?
"+" :
"-") + Feature).str()));
599 if (Args.hasFlag(options::OPT_mwavefrontsize64,
600 options::OPT_mno_wavefrontsize64,
false))
601 Features.push_back(
"+wavefrontsize64");
604 options::OPT_m_amdgpu_Features_Group);
612 {{options::OPT_O,
"3"}, {options::OPT_cl_std_EQ,
"CL1.2"}}) {
628 DerivedArgList *DAL =
634 DAL =
new DerivedArgList(Args.getBaseArgs());
640 Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);
641 if (LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==
"native") {
642 DAL->eraseArg(options::OPT_mcpu_EQ);
646 << llvm::Triple::getArchTypeName(
getArch())
647 << llvm::toString(GPUsOrErr.takeError()) <<
"-mcpu";
649 auto &GPUs = *GPUsOrErr;
650 if (GPUs.size() > 1) {
652 << llvm::Triple::getArchTypeName(
getArch())
653 << llvm::join(GPUs,
", ") <<
"-mcpu";
655 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_mcpu_EQ),
656 Args.MakeArgString(GPUs.front()));
662 if (!Args.getLastArgValue(options::OPT_x).equals(
"cl"))
666 if (Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {
667 DAL->AddFlagArg(
nullptr, Opts.getOption(
getTriple().isArch64Bit()
669 : options::OPT_m32));
673 if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
675 DAL->AddJoinedArg(
nullptr, Opts.getOption(options::OPT_O),
683 llvm::AMDGPU::GPUKind Kind) {
686 if (Kind == llvm::AMDGPU::GK_NONE)
689 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
693 const bool BothDenormAndFMAFast =
694 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
695 (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
696 return !BothDenormAndFMAFast;
700 const llvm::opt::ArgList &DriverArgs,
const JobAction &JA,
701 const llvm::fltSemantics *FPType)
const {
703 if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
704 return llvm::DenormalMode::getIEEE();
709 auto Kind = llvm::AMDGPU::parseArchAMDGCN(Arch);
710 if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
711 DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
712 options::OPT_fno_gpu_flush_denormals_to_zero,
714 return llvm::DenormalMode::getPreserveSign();
716 return llvm::DenormalMode::getIEEE();
719 const StringRef GpuArch =
getGPUArch(DriverArgs);
720 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
724 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
729 return DAZ ? llvm::DenormalMode::getPreserveSign() :
730 llvm::DenormalMode::getIEEE();
734 llvm::AMDGPU::GPUKind Kind) {
735 const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
736 bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);
738 return !HasWave32 || DriverArgs.hasFlag(
739 options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,
false);
751 const llvm::opt::ArgList &DriverArgs,
752 llvm::opt::ArgStringList &CC1Args,
756 if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
757 options::OPT_fvisibility_ms_compat)) {
758 CC1Args.push_back(
"-fvisibility=hidden");
759 CC1Args.push_back(
"-fapply-global-visibility-to-externs");
766 getTriple(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));
771 StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
772 if (TargetID.empty())
773 return {std::nullopt, std::nullopt, std::nullopt};
775 llvm::StringMap<bool> FeatureMap;
777 if (!OptionalGpuArch)
778 return {TargetID.str(), std::nullopt, std::nullopt};
780 return {TargetID.str(), OptionalGpuArch->str(), FeatureMap};
784 const llvm::opt::ArgList &DriverArgs)
const {
786 if (PTID.OptionalTargetID && !PTID.OptionalGPUArch) {
788 << *PTID.OptionalTargetID;
796 if (Arg *A = Args.getLastArg(options::OPT_amdgpu_arch_tool_EQ))
797 Program = A->getValue();
803 return StdoutOrErr.takeError();
806 for (StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(),
"\n"))
808 GPUArchs.push_back(Arch.str());
810 if (GPUArchs.empty())
811 return llvm::createStringError(std::error_code(),
812 "No AMD GPU detected in the system");
814 return std::move(GPUArchs);
818 const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
821 DeviceOffloadingKind);
826 DriverArgs.hasArg(options::OPT_nostdlib))
829 if (DriverArgs.hasArg(options::OPT_nogpulib))
833 const StringRef GpuArch =
getGPUArch(DriverArgs);
834 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
835 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
843 bool Wave64 =
isWave64(DriverArgs, Kind);
847 bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
849 bool FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only);
852 DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations);
853 bool FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math);
855 DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt);
863 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
864 FastRelaxedMath, CorrectSqrt, ABIVer,
false));
866 for (StringRef BCFile : BCLibs) {
867 CC1Args.push_back(
"-mlink-builtin-bitcode");
868 CC1Args.push_back(DriverArgs.MakeArgString(BCFile));
873 StringRef GPUArch, StringRef LibDeviceFile,
875 if (!hasDeviceLibrary()) {
876 D.Diag(diag::err_drv_no_rocm_device_lib) << 0;
879 if (LibDeviceFile.empty()) {
880 D.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;
884 D.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.
toString();
892 const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,
bool Wave64,
893 bool DAZ,
bool FiniteOnly,
bool UnsafeMathOpt,
bool FastRelaxedMath,
897 auto AddBCLib = [&](StringRef BCFile) { BCLibs.push_back(BCFile.str()); };
899 AddBCLib(getOCMLPath());
900 AddBCLib(getOCKLPath());
901 AddBCLib(getDenormalsAreZeroPath(DAZ));
902 AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath));
903 AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath));
904 AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt));
905 AddBCLib(getWavefrontSize64Path(Wave64));
906 AddBCLib(LibDeviceFile);
907 auto ABIVerPath = getABIVersionPath(ABIVer);
908 if (!ABIVerPath.empty())
909 AddBCLib(ABIVerPath);
916 const std::string &GPUArch,
917 bool isOpenMP)
const {
918 auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
919 const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
931 bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,
932 options::OPT_fno_gpu_flush_denormals_to_zero,
934 bool FiniteOnly = DriverArgs.hasFlag(
935 options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only,
false);
937 DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
938 options::OPT_fno_unsafe_math_optimizations,
false);
939 bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,
940 options::OPT_fno_fast_math,
false);
941 bool CorrectSqrt = DriverArgs.hasFlag(
942 options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
943 options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt,
true);
944 bool Wave64 =
isWave64(DriverArgs, Kind);
947 DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,
948 FastRelaxedMath, CorrectSqrt, ABIVer, isOpenMP);
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...
std::string SysRoot
sysroot, if present
const char * getClangProgramPath() const
Get the path to the main clang executable.
DiagnosticBuilder Diag(unsigned DiagID) const
const llvm::opt::OptTable & getOpts() const
std::string ResourceDir
The path to the compiler resource directory.
const char * getInstalledDir() const
Get the path to where the clang executable was installed.
llvm::vfs::FileSystem & getVFS() 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 hasHIPRuntime() const
Check whether we detected a valid HIP runtime.
llvm::SmallVector< std::string, 12 > getCommonBitcodeLibs(const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, bool Wave64, bool DAZ, bool FiniteOnly, bool UnsafeMathOpt, bool FastRelaxedMath, bool CorrectSqrt, DeviceLibABIVersion ABIVer, bool isOpenMP) const
Get file paths of default bitcode libraries common to AMDGPU based toolchains.
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.
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.
@ C
Languages that the frontend can parse and compile.
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.
ABI version of device library.
static DeviceLibABIVersion fromCodeObjectVersion(unsigned CodeObjectVersion)
bool requiresLibrary()
Whether ABI version bc file is requested.
static constexpr ResponseFileSupport AtFileCurCP()