Home - News ( United States | United Kingdom | Italy | Germany ) - Football scores

Showing content from https://clang.llvm.org/doxygen/Driver_2ToolChains_2AMDGPU_8cpp_source.html below:

clang: lib/Driver/ToolChains/AMDGPU.cpp Source File

;

RocmInstallationDetector::findSPACKPackage(

Candidate &Cand,

StringRef PackageName) {

std::string Prefix = Twine(PackageName +

+ Cand.SPACKReleaseStr).str();

(llvm::vfs::directory_iterator

= D.

().dir_begin(Cand.Path, EC),

!= FileEnd && !EC;

.increment(EC)) {

llvm::StringRef

= llvm::sys::path::filename(

->path());

(

.starts_with(Prefix)) {

(SubDirs.size() > 1)

(SubDirs.size() == 1) {

PackagePath = Cand.Path;

llvm::sys::path::append(PackagePath, SubDirs[0]);

(SubDirs.size() == 0 && Verbose) {

llvm::errs() <<

<< Prefix <<

<< Cand.Path

(SubDirs.size() > 1 && Verbose) {

llvm::errs() <<

<< Prefix <<

<< Cand.Path

<<

;

RocmInstallationDetector::scanLibDevicePath(llvm::StringRef

) {

assert(!

.empty());

StringRef Suffix(

);

StringRef Suffix2(

);

(llvm::vfs::directory_iterator LI = D.

().dir_begin(

, EC), LE;

!EC && LI != LE; LI = LI.increment(EC)) {

StringRef FilePath = LI->path();

StringRef

= llvm::sys::path::filename(FilePath);

BaseName =

.drop_back(Suffix2.size());

(

.ends_with(Suffix))

BaseName =

.drop_back(Suffix.size());

StringRef ABIVersionPrefix =

;

(BaseName ==

) {

}

(BaseName ==

) {

}

(BaseName ==

) {

}

(BaseName ==

) {

}

(BaseName ==

) {

}

(BaseName ==

) {

FiniteOnly.Off = FilePath;

}

(BaseName ==

) {

FiniteOnly.On = FilePath;

}

(BaseName ==

) {

DenormalsAreZero.On = FilePath;

}

(BaseName ==

) {

DenormalsAreZero.Off = FilePath;

}

(BaseName ==

) {

CorrectlyRoundedSqrt.On = FilePath;

}

(BaseName ==

) {

CorrectlyRoundedSqrt.Off = FilePath;

}

(BaseName ==

) {

UnsafeMath.On = FilePath;

}

(BaseName ==

) {

UnsafeMath.Off = FilePath;

}

(BaseName ==

) {

WavefrontSize64.On = FilePath;

}

(BaseName ==

) {

WavefrontSize64.Off = FilePath;

}

(BaseName.starts_with(ABIVersionPrefix)) {

ABIVersionNumber;

(BaseName.drop_front(ABIVersionPrefix.size())

.getAsInteger(

0, ABIVersionNumber))

ABIVersionMap[ABIVersionNumber] = FilePath.str();

StringRef DeviceLibPrefix =

;

(!BaseName.starts_with(DeviceLibPrefix))

StringRef IsaVersionNumber =

BaseName.drop_front(DeviceLibPrefix.size());

llvm::Twine GfxName = Twine(

) + IsaVersionNumber;

std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));

RocmInstallationDetector::parseHIPVersionFile(llvm::StringRef

) {

.split(VersionParts,

);

Major = ~0

;

Minor = ~0

;

(

Part : VersionParts) {

Splits = Part.rtrim().split(

);

(Splits.first ==

) {

(Splits.second.getAsInteger(0, Major))

}

(Splits.first ==

) {

(Splits.second.getAsInteger(0, Minor))

}

(Splits.first ==

)

VersionPatch = Splits.second.str();

(Major == ~0

|| Minor == ~0

)

VersionMajorMinor = llvm::VersionTuple(Major, Minor);

(Twine(Major) +

+ Twine(Minor) +

+ VersionPatch).str();

RocmInstallationDetector::getInstallationPathCandidates() {

(!ROCmSearchDirs.empty())

ROCmSearchDirs;

DoPrintROCmSearchDirs = [&]() {

(PrintROCmSearchDirs)

(

Cand : ROCmSearchDirs) {

llvm::errs() <<

;

llvm::errs() <<

<< Cand.SPACKReleaseStr <<

;

llvm::errs() <<

<< Cand.Path <<

;

(!RocmPathArg.empty()) {

ROCmSearchDirs.emplace_back(RocmPathArg.str());

DoPrintROCmSearchDirs();

ROCmSearchDirs;

}

(std::optional<std::string> RocmPathEnv =

llvm::sys::Process::GetEnv(

)) {

(!RocmPathEnv->empty()) {

ROCmSearchDirs.emplace_back(std::move(*RocmPathEnv));

DoPrintROCmSearchDirs();

ROCmSearchDirs;

StringRef InstallDir = D.

;

DeduceROCmPath = [](StringRef ClangPath) {

StringRef ParentDir = llvm::sys::path::parent_path(ClangPath);

StringRef ParentName = llvm::sys::path::filename(ParentDir);

(ParentName ==

) {

ParentDir = llvm::sys::path::parent_path(ParentDir);

ParentName = llvm::sys::path::filename(ParentDir);

(ParentName.starts_with(

)) {

ParentName.drop_front(strlen(

)).split(

);

SPACKReleaseStr = SPACKPostfix.first;

(!SPACKReleaseStr.empty()) {

ParentDir = llvm::sys::path::parent_path(ParentDir);

Candidate(ParentDir.str(),

,

(ParentName ==

|| ParentName.starts_with(

))

ParentDir = llvm::sys::path::parent_path(ParentDir);

Candidate(ParentDir.str(),

);

ROCmSearchDirs.emplace_back(DeduceROCmPath(InstallDir));

ParentPath = llvm::sys::path::parent_path(RealClangPath);

(ParentPath != InstallDir)

ROCmSearchDirs.emplace_back(DeduceROCmPath(ParentPath));

ClangRoot = llvm::sys::path::parent_path(InstallDir);

RealClangRoot = llvm::sys::path::parent_path(ParentPath);

ROCmSearchDirs.emplace_back(ClangRoot.str(),

);

(RealClangRoot != ClangRoot)

ROCmSearchDirs.emplace_back(RealClangRoot.str(),

);

ROCmSearchDirs.emplace_back(D.

+

,

std::string LatestROCm;

llvm::VersionTuple LatestVer;

GetROCmVersion = [](StringRef DirName) {

llvm::VersionTuple

;

std::string VerStr = DirName.drop_front(strlen(

)).str();

std::replace(VerStr.begin(), VerStr.end(),

,

);

(llvm::vfs::directory_iterator

!= FileEnd && !EC;

.increment(EC)) {

llvm::StringRef

= llvm::sys::path::filename(

->path());

(!

.starts_with(

))

(LatestROCm.empty()) {

LatestVer = GetROCmVersion(LatestROCm);

Ver = GetROCmVersion(

);

(LatestVer < Ver) {

(!LatestROCm.empty())

ROCmSearchDirs.emplace_back(D.

+

+ LatestROCm,

ROCmSearchDirs.emplace_back(D.

+

,

ROCmSearchDirs.emplace_back(D.

+

,

DoPrintROCmSearchDirs();

ROCmSearchDirs;

&

,

llvm::Triple &HostTriple,

llvm::opt::ArgList &Args,

DetectHIPRuntime,

DetectDeviceLib)

Verbose = Args.hasArg(options::OPT_v);

RocmPathArg = Args.getLastArgValue(clang::driver::options::OPT_rocm_path_EQ);

PrintROCmSearchDirs =

Args.hasArg(clang::driver::options::OPT_print_rocm_search_dirs);

RocmDeviceLibPathArg =

Args.getAllArgValues(clang::driver::options::OPT_rocm_device_lib_path_EQ);

HIPPathArg = Args.getLastArgValue(clang::driver::options::OPT_hip_path_EQ);

Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_path_EQ);

HasHIPStdParLibrary =

!HIPStdParPathArg.empty() &&

.getVFS().exists(HIPStdParPathArg +

);

HIPRocThrustPathArg =

Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_thrust_path_EQ);

HasRocThrustLibrary = !HIPRocThrustPathArg.empty() &&

.getVFS().exists(HIPRocThrustPathArg +

);

Args.getLastArgValue(clang::driver::options::OPT_hipstdpar_prim_path_EQ);

HasRocPrimLibrary = !HIPRocPrimPathArg.empty() &&

.getVFS().exists(HIPRocPrimPathArg +

);

(

*A = Args.getLastArg(clang::driver::options::OPT_hip_version_EQ)) {

HIPVersionArg = A->getValue();

Major = ~0

;

Minor = ~0

;

HIPVersionArg.split(Parts,

);

Parts[0].getAsInteger(0, Major);

(Parts.size() > 1)

Parts[1].getAsInteger(0, Minor);

(Parts.size() > 2)

VersionPatch = Parts[2].str();

(VersionPatch.empty())

(Major != ~0

&& Minor == ~0

)

(Major == ~0

|| Minor == ~0

)

.Diag(diag::err_drv_invalid_value)

<< A->getAsString(Args) << HIPVersionArg;

VersionMajorMinor = llvm::VersionTuple(Major, Minor);

(Twine(Major) +

+ Twine(Minor) +

+ VersionPatch).str();

VersionPatch = DefaultVersionPatch;

llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);

DetectedVersion = (Twine(DefaultVersionMajor) +

+

Twine(DefaultVersionMinor) +

+ VersionPatch)

(DetectHIPRuntime)

assert(LibDevicePath.empty());

(!RocmDeviceLibPathArg.empty())

LibDevicePath = RocmDeviceLibPathArg[RocmDeviceLibPathArg.size() - 1];

(std::optional<std::string> LibPathEnv =

llvm::sys::Process::GetEnv(

))

LibDevicePath = std::move(*LibPathEnv);

(!LibDevicePath.empty()) {

(!FS.exists(LibDevicePath))

scanLibDevicePath(LibDevicePath);

HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();

CheckDeviceLib = [&](StringRef

,

StrictChecking) {

CheckLibDevice = (!NoBuiltinLibs || StrictChecking);

(CheckLibDevice && !FS.exists(

))

scanLibDevicePath(

);

(!NoBuiltinLibs) {

(!allGenericLibsValid())

(LibDeviceMap.empty())

llvm::sys::path::append(LibDevicePath, CLANG_INSTALL_LIBDIR_BASENAME,

,

);

HasDeviceLibrary = CheckDeviceLib(LibDevicePath,

);

(HasDeviceLibrary)

&ROCmDirs = getInstallationPathCandidates();

(

&Candidate : ROCmDirs) {

LibDevicePath = Candidate.Path;

llvm::sys::path::append(LibDevicePath,

,

);

HasDeviceLibrary = CheckDeviceLib(LibDevicePath, Candidate.StrictChecking);

(HasDeviceLibrary)

(!HIPPathArg.empty())

HIPSearchDirs.emplace_back(HIPPathArg.str());

(std::optional<std::string> HIPPathEnv =

llvm::sys::Process::GetEnv(

)) {

(!HIPPathEnv->empty())

HIPSearchDirs.emplace_back(std::move(*HIPPathEnv));

(HIPSearchDirs.empty())

HIPSearchDirs.append(getInstallationPathCandidates());

(

&Candidate : HIPSearchDirs) {

InstallPath = Candidate.Path;

(InstallPath.empty() || !FS.exists(InstallPath))

SPACKPath = findSPACKPackage(Candidate,

);

InstallPath = SPACKPath.empty() ? InstallPath : SPACKPath;

BinPath = InstallPath;

llvm::sys::path::append(BinPath,

);

IncludePath = InstallPath;

llvm::sys::path::append(IncludePath,

);

LibPath = InstallPath;

llvm::sys::path::append(LibPath,

);

SharePath = InstallPath;

llvm::sys::path::append(SharePath,

);

ParentSharePath = llvm::sys::path::parent_path(InstallPath);

llvm::sys::path::append(ParentSharePath,

);

Twine &

=

,

Twine &d =

) {

llvm::sys::path::append(newpath, a,

,

, d);

std::vector<SmallString<0>> VersionFilePaths = {

(SharePath,

,

),

InstallPath != D.

+

?

(ParentSharePath,

,

)

(BinPath,

)};

(

&VersionFilePath : VersionFilePaths) {

(VersionFilePath.empty())

llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =

FS.getBufferForFile(VersionFilePath);

(HIPVersionArg.empty() && VersionFile)

(parseHIPVersionFile((*VersionFile)->getBuffer()))

HasHIPRuntime =

;

(!Candidate.StrictChecking) {

HasHIPRuntime =

;

HasHIPRuntime =

;

OS <<

<< InstallPath <<

<< DetectedVersion <<

;

ArgStringList &CC1Args)

{

UsesRuntimeWrapper = VersionMajorMinor > llvm::VersionTuple(3, 5) &&

!DriverArgs.hasArg(options::OPT_nohipwrapperinc);

HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);

(!DriverArgs.hasArg(options::OPT_nobuiltininc)) {

(UsesRuntimeWrapper)

llvm::sys::path::append(

,

,

);

CC1Args.push_back(

);

CC1Args.push_back(DriverArgs.MakeArgString(

));

HandleHipStdPar = [=, &DriverArgs, &CC1Args]() {

(!HIPStdParPathArg.empty() ||

!FS.exists(Inc +

)) {

D.

(diag::err_drv_no_hipstdpar_lib);

(!HasRocThrustLibrary && !FS.exists(Inc +

)) {

D.

(diag::err_drv_no_hipstdpar_thrust_lib);

(!HasRocPrimLibrary && !FS.exists(Inc +

)) {

D.

(diag::err_drv_no_hipstdpar_prim_lib);

*ThrustPath;

(HasRocThrustLibrary)

ThrustPath = DriverArgs.MakeArgString(HIPRocThrustPathArg);

ThrustPath = DriverArgs.MakeArgString(Inc +

);

*HIPStdParPath;

HIPStdParPath = DriverArgs.MakeArgString(HIPStdParPathArg);

HIPStdParPath = DriverArgs.MakeArgString(StringRef(ThrustPath) +

);

*PrimPath;

(HasRocPrimLibrary)

PrimPath = DriverArgs.MakeArgString(HIPRocPrimPathArg);

PrimPath = DriverArgs.MakeArgString(

() +

);

CC1Args.append({

, ThrustPath,

, PrimPath,

, HIPStdParPath,

,

});

(DriverArgs.hasArg(options::OPT_nogpuinc)) {

D.

(diag::err_drv_no_hip_runtime);

CC1Args.push_back(

);

(UsesRuntimeWrapper)

CC1Args.append({

,

});

*LinkingOutput)

{

std::string

= getToolChain().GetLinkerPath();

ArgStringList CmdArgs;

(!Args.hasArg(options::OPT_r)) {

CmdArgs.push_back(

);

CmdArgs.push_back(

);

Args.AddAllArgs(CmdArgs, options::OPT_L);

getToolChain().AddFilePathLibArgs(Args, CmdArgs);

(

.getDriver().isUsingLTO()) {

(getToolChain(), Args, CmdArgs, Output, Inputs[0],

}

(Args.hasArg(options::OPT_mcpu_EQ)) {

CmdArgs.push_back(Args.MakeArgString(

+

Args.getLastArgValue(options::OPT_mcpu_EQ))));

std::vector<StringRef> Features;

(!Features.empty()) {

Args.MakeArgString(

+ llvm::join(Features,

)));

(Args.hasArg(options::OPT_stdlib))

CmdArgs.append({

,

});

(Args.hasArg(options::OPT_startfiles)) {

std::optional<std::string> IncludePath = getToolChain().getStdlibPath();

IncludePath =

;

llvm::sys::path::append(

,

);

CmdArgs.push_back(Args.MakeArgString(

));

CmdArgs.push_back(

);

.addCommand(std::make_unique<Command>(

CmdArgs, Inputs, Output));

llvm::Triple &Triple,

llvm::opt::ArgList &Args,

std::vector<StringRef> &Features) {

(Args.hasArg(options::OPT_mcpu_EQ))

TargetID = Args.getLastArgValue(options::OPT_mcpu_EQ);

(Args.hasArg(options::OPT_march_EQ))

TargetID = Args.getLastArgValue(options::OPT_march_EQ);

(!TargetID.empty()) {

llvm::StringMap<bool> FeatureMap;

OptionalGpuArch =

(Triple, TargetID, &FeatureMap);

(OptionalGpuArch) {

StringRef GpuArch = *OptionalGpuArch;

Pos = FeatureMap.find(Feature);

(Pos == FeatureMap.end())

Features.push_back(Args.MakeArgStringRef(

(Twine(Pos->second ?

:

) + Feature).str()));

(Args.hasFlag(options::OPT_mwavefrontsize64,

options::OPT_mno_wavefrontsize64,

))

Features.push_back(

);

(Args.hasFlag(options::OPT_mamdgpu_precise_memory_op,

options::OPT_mno_amdgpu_precise_memory_op,

))

Features.push_back(

);

options::OPT_m_amdgpu_Features_Group);

{{options::OPT_O,

}, {options::OPT_cl_std_EQ,

}}) {

DerivedArgList *DAL =

DAL =

DerivedArgList(Args.getBaseArgs());

Arg *LastMCPUArg = DAL->getLastArg(options::OPT_mcpu_EQ);

(LastMCPUArg && StringRef(LastMCPUArg->getValue()) ==

) {

DAL->eraseArg(options::OPT_mcpu_EQ);

<< llvm::Triple::getArchTypeName(

())

<< llvm::toString(GPUsOrErr.takeError()) <<

;

&GPUs = *GPUsOrErr;

(GPUs.size() > 1) {

<< llvm::Triple::getArchTypeName(

())

<< llvm::join(GPUs,

) <<

;

DAL->AddJoinedArg(

, Opts.getOption(options::OPT_mcpu_EQ),

Args.MakeArgString(GPUs.front()));

(Args.getLastArgValue(options::OPT_x) !=

)

(Args.hasArg(options::OPT_c) && Args.hasArg(options::OPT_emit_llvm)) {

DAL->AddFlagArg(

, Opts.getOption(

().isArch64Bit()

: options::OPT_m32));

(!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,

DAL->AddJoinedArg(

, Opts.getOption(options::OPT_O),

llvm::AMDGPU::GPUKind Kind) {

(Kind == llvm::AMDGPU::GK_NONE)

ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);

BothDenormAndFMAFast =

(ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&

(ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);

!BothDenormAndFMAFast;

llvm::opt::ArgList &DriverArgs,

&JA,

llvm::fltSemantics *FPType)

{

(!FPType || FPType != &llvm::APFloat::IEEEsingle())

llvm::DenormalMode::getIEEE();

Kind = llvm::AMDGPU::parseArchAMDGCN(Arch);

(FPType && FPType == &llvm::APFloat::IEEEsingle() &&

DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,

options::OPT_fno_gpu_flush_denormals_to_zero,

llvm::DenormalMode::getPreserveSign();

llvm::DenormalMode::getIEEE();

StringRef GpuArch =

(DriverArgs);

Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);

DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||

DAZ ? llvm::DenormalMode::getPreserveSign() :

llvm::DenormalMode::getIEEE();

llvm::AMDGPU::GPUKind Kind) {

ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);

HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32);

!HasWave32 || DriverArgs.hasFlag(

options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64,

);

llvm::opt::ArgList &DriverArgs,

llvm::opt::ArgStringList &CC1Args,

(!DriverArgs.hasArg(options::OPT_fvisibility_EQ,

options::OPT_fvisibility_ms_compat)) {

CC1Args.push_back(

);

CC1Args.push_back(

);

CC1Args.push_back(

);

(), DriverArgs.getLastArgValue(options::OPT_mcpu_EQ));

StringRef TargetID = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);

(TargetID.empty())

{std::nullopt, std::nullopt, std::nullopt};

llvm::StringMap<bool> FeatureMap;

(!OptionalGpuArch)

{TargetID.str(), std::nullopt, std::nullopt};

{TargetID.str(), OptionalGpuArch->str(), FeatureMap};

llvm::opt::ArgList &DriverArgs)

{

(PTID.OptionalTargetID && !PTID.OptionalGPUArch) {

<< *PTID.OptionalTargetID;

(Arg *A = Args.getLastArg(options::OPT_amdgpu_arch_tool_EQ))

Program = A->getValue();

StdoutOrErr.takeError();

(StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(),

))

GPUArchs.push_back(Arch.str());

(GPUArchs.empty())

llvm::createStringError(std::error_code(),

);

std::move(GPUArchs);

llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,

DeviceOffloadingKind);

DriverArgs.hasArg(options::OPT_nostdlib))

(DriverArgs.hasArg(options::OPT_nogpulib))

StringRef GpuArch =

(DriverArgs);

Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);

StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);

std::tuple<bool, const SanitizerArgs> GPUSan(

DriverArgs.hasFlag(options::OPT_fgpu_sanitize,

options::OPT_fno_gpu_sanitize,

),

Wave64 =

(DriverArgs, Kind);

DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||

FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only);

DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations);

FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math);

DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt);

DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,

FastRelaxedMath, CorrectSqrt, ABIVer, GPUSan,

));

(

[BCFile, Internalize] : BCLibs) {

CC1Args.push_back(

);

CC1Args.push_back(

);

CC1Args.push_back(DriverArgs.MakeArgString(BCFile));

StringRef GPUArch, StringRef LibDeviceFile,

(!hasDeviceLibrary()) {

.Diag(diag::err_drv_no_rocm_device_lib) << 0;

(LibDeviceFile.empty()) {

.Diag(diag::err_drv_no_rocm_device_lib) << 1 << GPUArch;

(ABIVer.

() && getABIVersionPath(ABIVer).empty()) {

.Diag(diag::err_drv_no_rocm_device_lib) << 2 << ABIVer.

();

llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile,

Wave64,

DAZ,

FiniteOnly,

UnsafeMathOpt,

FastRelaxedMath,

std::tuple<bool, const SanitizerArgs> &GPUSan,

isOpenMP =

)

{

GPUSanEnabled = [GPUSan]() {

std::get<bool>(GPUSan); };

Internalize =

) {

BCLibs.emplace_back(BCLib);

AddSanBCLibs = [&]() {

(GPUSanEnabled()) {

SanArgs = std::get<const SanitizerArgs>(GPUSan);

(SanArgs.needsAsanRt())

AddBCLib(getAsanRTLPath(),

);

AddBCLib(getOCMLPath());

AddBCLib(getOCKLPath());

(GPUSanEnabled() && isOpenMP)

AddBCLib(getOCKLPath(),

);

AddBCLib(getDenormalsAreZeroPath(DAZ));

AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath));

AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath));

AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt));

AddBCLib(getWavefrontSize64Path(Wave64));

AddBCLib(LibDeviceFile);

ABIVerPath = getABIVersionPath(ABIVer);

(!ABIVerPath.empty())

AddBCLib(ABIVerPath);

std::string &GPUArch,

isOpenMP)

{

Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);

StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);

std::tuple<bool, const SanitizerArgs> GPUSan(

DriverArgs.hasFlag(options::OPT_fgpu_sanitize,

options::OPT_fno_gpu_sanitize,

),

DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,

options::OPT_fno_gpu_flush_denormals_to_zero,

FiniteOnly = DriverArgs.hasFlag(

options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only,

);

UnsafeMathOpt =

DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,

options::OPT_fno_unsafe_math_optimizations,

);

FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,

options::OPT_fno_fast_math,

);

CorrectSqrt = DriverArgs.hasFlag(

options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,

options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt,

);

Wave64 =

(DriverArgs, Kind);

DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt,

FastRelaxedMath, CorrectSqrt, ABIVer, GPUSan, isOpenMP);

&TC,

llvm::opt::ArgList &DriverArgs,

StringRef TargetID,

llvm::opt::Arg *A)

{

(TargetID.empty())

Option O = A->getOption();

(!O.matches(options::OPT_fsanitize_EQ))

(!DriverArgs.hasFlag(options::OPT_fgpu_sanitize,

options::OPT_fno_gpu_sanitize,

))

(K != SanitizerKind::Address)

llvm::StringMap<bool> FeatureMap;

assert(OptionalGpuArch &&

);

(void)OptionalGpuArch;

= FeatureMap.find(

);

(

== FeatureMap.end() || !

->second) {

clang::diag::warn_drv_unsupported_option_for_offload_arch_req_feature)

<< A->getAsString(DriverArgs) << TargetID <<

;

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

DiagnosticsEngine & getDiags() const

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.

llvm::vfs::FileSystem & getVFS() const

std::string Dir

The path the driver executable was in, as invoked from the command line.

InputInfo - Wrapper for information about an input source.

const char * getFilename() const

StringRef getIncludePath() const

Get the detected path to Rocm's bin directory.

llvm::SmallVector< ToolChain::BitCodeLibraryInfo, 12 > getCommonBitcodeLibs(const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, bool Wave64, bool DAZ, bool FiniteOnly, bool UnsafeMathOpt, bool FastRelaxedMath, bool CorrectSqrt, DeviceLibABIVersion ABIVer, const std::tuple< bool, const SanitizerArgs > &GPUSan, bool isOpenMP) const

Get file paths of default bitcode libraries common to AMDGPU based toolchains.

RocmInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args, bool DetectHIPRuntime=true, bool DetectDeviceLib=false)

bool checkCommonBitcodeLibs(StringRef GPUArch, StringRef LibDeviceFile, DeviceLibABIVersion ABIVer) const

Check file paths of default bitcode libraries common to AMDGPU based toolchains.

bool hasHIPStdParLibrary() const

Check whether we detected a valid HIP STDPAR Acceleration library.

bool hasHIPRuntime() const

Check whether we detected a valid HIP runtime.

void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const

void detectDeviceLibrary()

void print(raw_ostream &OS) const

Print information about the detected ROCm installation.

The JSON file list parser is used to communicate input to InstallAPI.

std::optional< llvm::StringRef > parseTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch, llvm::StringMap< bool > *FeatureMap)

Parse a target ID to get processor and feature map.

llvm::StringRef getProcessorFromTargetID(const llvm::Triple &T, llvm::StringRef OffloadArch)

Get processor name from target ID.

llvm::SmallVector< llvm::StringRef, 4 > getAllPossibleTargetIDFeatures(const llvm::Triple &T, llvm::StringRef Processor)

Get all feature strings that can be used in target ID for Processor.

SanitizerMask parseSanitizerValue(StringRef Value, bool AllowGroups)

Parse a single value from a -fsanitize= or -fno-sanitize= value list.

ABI version of device library.

static DeviceLibABIVersion fromCodeObjectVersion(unsigned CodeObjectVersion)

bool requiresLibrary()

Whether ABI version bc file is requested.

static constexpr ResponseFileSupport AtFileCurCP()

RetroSearch is an open source project built by @garambo | Open a GitHub Issue

Search and Browse the WWW like it's 1997 | Search results from DuckDuckGo

HTML: 3.2 | Encoding: UTF-8 | Version: 0.7.4